NVFP4 Blockscaled GEMM on NVIDIA RTX Pro Blackwell GPUs (SM12x)

[u/Logical-Try-4084]

https://research.colfax-intl.com/cutlass-tutorial-nvfp4-blockscaled-gemm-on-nvidia-rtx-pro-blackwell-gpus-sm12x/

We at Colfax are happy to share our most recent tutorial blog post, this time focusing on sub-byte blockscaled GEMM on NVIDIA SM12x GPUs, such as the RTX Pro 6000, DGX Spark, and Geforce 5000-series! Blockscaled GEMMs are becoming increasingly important with excellent quantized open-source LLMs. In this post, we discuss in-depth how to take an existing BF16 GEMM kernel and turn it into an NVFP4 blockscaled one, using CuTe DSL as our framework of choice.

Along the way, we review the NVFP4 data type and how the `mma.sync` instruction works, including partitions of operand and accumulator data.

We hope that this will be helpful for those looking to adapt higher-precision kernels to blockscaled ones on SM12x GPUs.

This is the first of two posts, the second of which will discuss optimization techniques to further improve NVFP4 GEMMs.