u/Logical-Try-4084

Our second blog post on writing NVFP4 blockscaled GEMM kernels for the NVIDIA RTX PRO 6000 Blackwell GPU is out! The blog iteratively optimizes a basic working NVFP4 GEMM kernel written in CuTe DSL to take it to speed-of-light, reaching over 80% TFLOP/s utilization for 16k square matrix shape. We give a detailed treatment of important optimization techniques such as threadblock swizzling, async and warp-specialized epilogue, and retiling for favorable wave quantization. Specific to blockscaled GEMM with scales consumed from registers, we also explain how to solve for bank conflicts that arise from the default choices of interleaved scale factor layouts.

We include complete code in the form of CuTe DSL kernels for all the optimizations discussed in the blog.

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.

Persistent tile scheduling is crucial for hiding epilogue latency in various kernels. Simply, a threadblock (CTA) is assigned available worktiles continuously until the kernel has completed. In its most naive form, a persistent tile scheduler will assign to CTA X successive tiles of index n*X and "delinearize" this list to compute the proper logical grid coordinates. However, when one has imbalanced workloads, CTAs may be assigned inequal work, a problem that compounds to a phenomenon of "CTA drift"; in this case, the persistent scheduler may perform significantly worse thanan ordinary single tile scheduler.

The solution to this problem is dynamic persistent schedulers: instead of assigning a pre-determined list of worktiles to a CTA, have a CTA claim the next available tile and start computing on that immediately. The common approach to dynamic persistence is to keep a global memory semaphore; once a CTA has finished its computation, it reads the semaphore and then atomically adds to increment. This adds enormous engineering complexity and has a lot of room for error.

Cluster Launch Control (CLC), introduced on Blackwell, approaches dynamic persistent scheduling with hardware-level instructions, alleviating much of the engineering challenge associated with the semaphore-based approach.

Please enjoy our recent blog on CLC, with examples in CuTeDSL! While the post focuses on datacenter Blackwell (Sm100), CLC is available on consumer cards as well (Sm120).