ThunderKittens Ported to NVIDIA Vera Rubin NVL72: NVFP4 GEMM Breaks 22 PFLOPS

The deep learning kernel library ThunderKittens has received a critical upgrade — the team has ported it to NVIDIA's latest Vera Rubin NVL72 platform and rebuilt its NVFP4 GEMM kernel around the new hardware. The results are impressive: performance climbed from just 42% of roofline (theoretical peak) all the way past 22 PFLOPS, putting it in direct competition with highly optimized official solutions like cuBLAS and CuTe DSL.

This article breaks down the technical changes behind this port, along with the opportunities and challenges introduced by the new hardware ISA (Instruction Set Architecture).

What Is ThunderKittens and Why Does This Port Matter

ThunderKittens is a lightweight GPU kernel authoring framework for deep learning, with a core goal of enabling researchers and engineers to write high-performance kernels close to hardware limits using clean, minimal abstractions. Rather than writing raw CUDA by hand or relying on sprawling template libraries, it emphasizes a "small but elegant" programming model organized around the core abstraction of tiles — structured blocks of computation and data movement.

Porting such a framework to a brand-new architecture is far more than just switching a compile target. Each generation of NVIDIA hardware brings significant evolution in tensor cores, memory hierarchies, and asynchronous data movement instructions. A kernel library that wants to fully exploit new hardware often requires targeted rewrites of critical paths. The leap from 42% roofline to 22 PFLOPS is a clear demonstration of the enormous gap between "recompiling for a new architecture" and "actually rebuilding for it."

Why NVFP4 GEMM Is the Core Performance Battleground

This upgrade centers on NVFP4 GEMM. GEMM (General Matrix Multiplication) is the performance cornerstone of virtually every deep learning workload — from attention mechanisms to feed-forward networks, the core computation reduces to matrix multiplication. Whoever can push GEMM close to theoretical peak effectively controls the efficiency ceiling of the entire inference and training stack.

NVFP4 represents the cutting edge of low-precision computing. As model sizes balloon, the 4-bit floating point (FP4) format dramatically reduces memory bandwidth and compute requirements while maintaining acceptable accuracy. NVIDIA's latest hardware provides dedicated tensor core support for these low-precision formats, but actually harnessing those units requires the software to precisely match the hardware's expectations around data layout, scale factors, and accumulation precision.

The team notes that rebuilding NVFP4 GEMM specifically around the new hardware features of the Vera Rubin NVL72 is what pushed performance from 42% of roofline to a competitive level. This implies the old kernel had a significant "impedance mismatch" on the new architecture — and the rewrite closed that gap.

FP4 (4-bit floating point) sits at the frontier of low-precision computing. Compared to the widely used FP16/BF16 (16-bit) formats, FP4 compresses each value to 4 bits, theoretically reducing memory bandwidth requirements to one quarter while allowing tensor cores to pack more parallel compute units into the same silicon area. However, the extremely low bit-width also means an extremely narrow dynamic range — standard FP4 can only represent a handful of distinct values, so a "scale factor" mechanism is required to cover the distribution of real model weights. NVIDIA's NVFP4, introduced in Blackwell and extended in Vera Rubin, is an FP4 variant with block-wise scaling, with native hardware support for reading and fusing scale factors into computation. This makes NVFP4 GEMM software far more complex than ordinary matrix multiplication: the kernel must not only manage tiled loading of the main matrices, but also synchronously handle the scale factor for each data block and ensure both are precisely aligned in the tensor core pipeline — otherwise both accuracy and performance degrade severely.

Opportunities from ISA Changes

The article highlights that the key to this performance leap was exploiting ISA-level changes in the new architecture. While not all details are spelled out, the general patterns of GPU architecture evolution suggest these improvements typically come from several directions:

Stronger Tensor Cores and Low-Precision Pathways

Next-generation tensor cores typically offer higher throughput and more flexible scaling mechanisms for formats like FP4/FP8. By correctly mapping to these instructions, software can gain multiplicative compute throughput with no additional precision loss.

Asynchronous Memory and Data Movement Optimizations

The performance bottleneck in GEMM often lies not in computation itself, but in whether data can reach compute units in time. New ISA support for async copies, shared memory scheduling, and cluster-level cooperation directly determines whether memory latency can be hidden and roofline can be approached.

NVIDIA introduced asynchronous data movement instructions (such as cp.async) starting with the Ampere architecture. Hopper went further with the TMA (Tensor Memory Accelerator) hardware unit, allowing SMs to fully offload data movement from global memory or the NVLink domain to shared memory onto a dedicated engine — the main compute pipeline no longer needs to wait. Vera Rubin continues and extends this approach. For GEMM, this means constructing "double buffering" or even deeper pipelines: while tensor cores consume data tile N, TMA is already prefetching tile N+1 in the background, ideally hiding memory latency entirely. ThunderKittens' tile abstraction is a natural fit for this pipeline pattern — the framework can encapsulate async prefetching into scheduling primitives, letting developers avoid manual synchronization barrier management while still achieving memory efficiency close to hand-written assembly.

NVL72's Large-Scale Interconnect

Vera Rubin NVL72 is itself a high-density interconnected platform — the number 72 hints at large-scale GPU domain collaboration. At this scale, kernels must be efficient not just on a single card, but must maintain scalability across multi-GPU cooperation. The 22 PFLOPS figure is built on top of platform-level compute aggregation.

What It Means to Compete with cuBLAS and CuTe DSL

cuBLAS is NVIDIA's official highly optimized linear algebra library; CuTe DSL is the domain-specific language in the CUTLASS ecosystem for describing tensor layouts and computation — both represent the top tier of kernel optimization in the industry. A relatively lightweight open-source framework achieving competitive performance against them on new hardware is a powerful validation of its programming model and abstraction design.

For developers, this sends a positive signal: you don't necessarily need to dive into official closed-source libraries or complex template labyrinths to write near-peak-performance kernels on the latest hardware. Frameworks like ThunderKittens lower the barrier to extracting performance from cutting-edge hardware — which matters especially for the research community.

CuTe DSL is the core subsystem within NVIDIA's CUTLASS library, using C++ template metaprogramming to provide a unified description of tensors' logical layout and physical stride, and generating highly specialized kernel code at compile time. Its design philosophy is "zero-cost abstraction": developers describe computations in a high-level layout language, and the compiler automatically derives optimal data movement and instruction selection strategies. This mechanism is extremely powerful but has a steep learning curve and complex debugging. ThunderKittens differs by prioritizing "readability and research iteration speed," trading some optimality in extreme edge cases for the ability to have a working custom kernel running within hours. The fact that it now approaches CuTe DSL performance on NVL72 shows that this tradeoff does not carry an unacceptable cost on the latest hardware.

Conclusion: The Endless Race Between Hardware Iteration and Software Adaptation

This port case once again confirms a fundamental truth: theoretical hardware peak performance never automatically translates into usable performance — there is a mountain of targeted software engineering in between. The gap from 42% to 22 PFLOPS lies not in the chip, but in whether the software has truly internalized the new ISA.

As low-precision formats and large-scale interconnected platforms become mainstream directions in AI infrastructure, open-source projects like ThunderKittens — willing to closely track the latest architectures and rapidly rebuild core kernels — will play an increasingly important role in the performance ecosystem. For engineers paying attention to the AI compute stack, this kind of work is well worth following.