You’ll often see “peak TOPS” advertised for accelerators like GPUs - but where does that number come from, and can we achieve it? Let’s break it down using the DGX Spark’s 1000 TOPS as an example.

Peak TOPS refers to trillions of floating point operations per second - but in the case of accelerators like NVIDIA GPUs, that number is almost always in reference to an MMA operation - performing matrix multiplication and accumulating the results.

It’s also worth noting that TOPS is always in the context of a specific precision - the data type that you’re operating on. Many vendors now report peak TOPS in a smaller data type used in quantized data formats such as FP8, INT8, or as low as FP4 (or per-block precision like NVFP4).

In addition, oftentimes TOPS are now reported using “sparse” TOPS - which assumes that the underlying matrix being multiplied has a pattern of zero or near-zero values that would effectively skip the multiplication entirely. It’s often ratios of 2:1 sparsity, so when we see reports of “X TOPS”, it’s X/2. So in the case of DGX, we should target roughly 500 TOPS.

Can we derive this number? For NVidia blackwell GPUs, the TOPS match this equation:

SM_COUNT * TENSOR_CORES_PER_SM * TENSOR_CORE_FLOPS_PER_CYCLE * CLOCK_SPEED

Substituting the Blackwell architecture parameters for the GB10 GPU:

  • SMs = 48
  • Tensor Cores per SM = 4
  • Tensor Core FP4 FLOPS per cycle = 2048
  • Clock Speed = 2.418 GHz

Substituting these values into the equation above:

Peak FP4 = 48 SMs * 4 Tensor Cores/SM * 2048 FLOPs/Core/Clock * 2.418 * 10^9 Hz Peak FP4 = 950,273,280,000 operations per second ~= 950.27 TFLOPS

But of course those are reported as sparse TOPS, so we halve the value:

Sparse FP4 = Peak FP4 / 2 Sparse FP4 = 475,136,640,000 operations per second ~= 475.14 TFLOPS

To achieve this, however, you need to make sure you’re maximizing the number of values you can fit in the registers of the tensor core as well. Specifically, you have to use the mma.sync.aligned.kind::mxf4nvf4.block_scale.scale_vec matrix multiply instruction, ensuring that you don’t pad additional bits in the register, using only the required 4 bits per element. Otherwise, for example with 8 bit padding, you’re effectively halving the utilization of the tensor core.

However, if you get all of those right, you can finally achieve max TOPS!

Thanks to Alfonso De Gregorio for the nvfp4bench tool (https://github.com/secYOUre/nvfp4bench), and for sharing the final insight of the register padding slowing things down in my own benchmarks: https://github.com/toumorokoshi/yft-ml-sandbox/blob/main/cuda/gemm/experiment.md

Additional Notes #

How to Query Hardware Parameters #

Method 1: Programmatically Querying SM Count and Clock Speed via CUDA API #

You can write a simple CUDA program querying the cudaGetDeviceProperties API to retrieve these hardware specifications directly:

#include <stdio.h>
#include <cuda_runtime.h>

int main() {
    cudaDeviceProp prop;
    int deviceId = 0; // Target GPU device ID

    if (cudaGetDeviceProperties(&prop, deviceId) == cudaSuccess) {
        printf("GPU Device: %s\n", prop.name);
        printf("SM Count (multiProcessorCount): %d\n", prop.multiProcessorCount);

        int clockRateKHz = 0;
        // Query clock rate attribute (cudaDevAttrClockRate)
        if (cudaDeviceGetAttribute(&clockRateKHz, cudaDevAttrClockRate, deviceId) == cudaSuccess) {
            printf("Clock Rate (cudaDevAttrClockRate): %d kHz (%.3f GHz)\n", clockRateKHz, clockRateKHz / 1e6);
        }
    }
    return 0;
}

Inspecting Compute Capability via nvidia-smi #

You can query the GPU model name and its CUDA Compute Capability from the command line:

nvidia-smi --query-gpu=name,compute_cap --format=csv

For example, this returns NVIDIA GB10, 12.1. Once you have the Compute Capability, you can reference the NVIDIA Developer CUDA GPUs Page and the NVIDIA CUDA C++ Programming Guide Architecture Specs to look up the exact SM counts and Tensor Core properties.

References #

  • Information about SM / Tensor Core counts: https://chipsandcheese.com/p/analyzing-nvidia-gb10s-gpu
  • NVIDIA Blackwell Architecture Whitepaper: NVIDIA Blackwell GPU Architecture Whitepaper - Core specifications for SM counts, fifth-generation Tensor Cores, and FP4 throughput metrics.