Streaming Multiprocessor
The replicated compute unit containing CUDA cores, Tensor Cores, schedulers, registers and on-chip memories.
One-page technical reference
Blackwell GPU & FP4 Cheat Sheet
How threads become Tensor Core work, how tiles move through the GPU, and how native NVFP4/MXFP4 GEMMs apply block scales without materializing dequantized matrices.
Execution organization
Data movement
Tensor computation
On/off-chip memory
1. Execution hierarchy
GPUmany SMs
SMschedules CTAs and warps
CTA / blockcooperating warps
Warp32 threads
2. Essential jargon
Cooperative Thread Array
PTX name for a CUDA thread block. Its warps cooperate through shared memory and barriers, usually on one SM.
32 execution lanes
A group of 32 threads issued together. Warp specialization assigns different warps to load, compute and store jobs.
Matrix multiply–accumulate
The Tensor Core operation D = A×B + C, performed on small matrix tiles.
5th-gen Tensor Core ISA
Blackwell data-center PTX instruction family for asynchronous FP16/BF16/FP8/FP6/FP4 MMA operations.
Tensor Memory Accelerator
Asynchronously moves multidimensional tiles between HBM/global memory and shared memory. It moves; it does not multiply.
Tensor Memory
Blackwell on-chip storage specialized for Tensor Core state, especially MMA accumulators. It is not TMA or HBM.
Shared memory
A general on-chip CTA scratchpad, commonly holding A/B input tiles before Tensor Core computation.
3. Native Blackwell GEMM flow
Packed matrices in HBMFP4 payloads + scale tensors
→
TMAloads tiles asynchronously
→
Shared memorystages A, B and scales
→
tcgen05.mmanative block-scaled Tensor Core dot products
→
TMEMFP32 accumulation
→
Epiloguebias / activation / BF16 or FP16 output
TMA feeds data; tcgen05.mma computes; TMEM retains the accumulated answer.
4. Three different meanings of “tile”
| Level | Typical meaning | Purpose |
|---|---|---|
| CTA tile | A large output region such as 128×128 | Work assigned to one CTA or CTA cluster. |
| MMA tile | A smaller M×N×K Tensor Core instruction shape | One native matrix operation; many compose a CTA tile. |
| Quantization block | 16 values for NVFP4, 32 for MXFP4 | Values sharing one local scale. This is not itself an output tile. |
5. FP4 block scaling
Original value: A[m,k]
Stored approximation: A[m,k] ≈ global_scale × local_scale[m,block(k)] × FP4_payload[m,k]
D[m,n] ≈ C[m,n] + Σblocks b scaleA[m,b] × scaleB[b,n] × (Σk∈b qA[m,k] × qB[k,n])
FP4 payload dot product
inside each scale block
inside each scale block
→
Multiply by A/B block scales
→
Accumulate in FP32 TMEM
This is mathematically equivalent to dequantizing each element first, but the hardware integrates scaling into the MMA. It does not construct complete BF16 copies of A and B.
6. NVFP4 versus MXFP4
| Property | NVFP4 | MXFP4 |
|---|---|---|
| Four-bit payload | E2M1 | E2M1 |
| Local scale group | 16 values | 32 values |
| Local scale format | FP8 E4M3 | Unsigned E8M0 |
| Scale flexibility | Fractional floating-point | Power of two only |
| Second-level scale | Usually one FP32 tensor/global scale | Not intrinsic to standard MXFP4 |
| Approx. payload + scale metadata | ~4.5 bits/value | ~4.25 bits/value |
| Typical tradeoff | Better accuracy, more metadata | Less metadata, coarser quantization |
7. What the scales actually do
NVFP4: two-level adaptation
Global FP32 scale maps the tensor’s overall range. A local E4M3 scale per 16 values then follows local magnitude changes. The local scale is not limited to powers of two.
MXFP4: compact power-of-two scaling
Each group of 32 values shares an E8M0 scale, mathematically 2^e. Scaling is simple and compact, but its steps are coarser.
8. Common misconceptions
| “TMEM and TMA are related names, so they are the same.” | No. TMEM stores Tensor Core state; TMA is an asynchronous transfer engine. |
| “FP4 GEMM dequantizes the full weights to BF16.” | No. Native block-scaled MMA consumes packed FP4 and scales directly. |
| “A quantization block is the Tensor Core tile.” | No. An MMA tile contains many quantization blocks. |
| “The variable called scale always means multiplication by that number.” | No. APIs may store a scale, inverse scale or amax-derived factor. Verify the convention. |
Debugging rule: treat the FP4 payload layout, scale values, scale-factor layout and global-scale convention as four separate things. A kernel can have correct numbers but associate them with the wrong blocks.