To Infinity and Beyond: ThunderKittens Now on NVIDIA Vera Rubin NVL72!

To Infinity and Beyond: ThunderKittens Now on NVIDIA Vera Rubin NVL72! Webflow Analyze/Optimize tracking bridge -->

🚀 DeepSeek V4 Pro 0813 vs. GPT-5.6 Sol on DeepSWE →

📈 GLM-5.3 vs. GLM 5.3 Flash on DeepSWE →

⚡ On-demand B200s now available on Together GPU Clusters →

🚀 Now serving MiniMax-M3 for efficient inference →

All blog posts

Research

Published 9/10/2026

To Infinity and Beyond: ThunderKittens Now on NVIDIA Vera Rubin NVL72!

Authors

Dylan Lim, Xinyi Li, Dan Fu, Simran Arora

Table of contents

40+ Models Chosen for Production...40+ Models Chosen for Production...40+ Models Chosen for Production...

The kernels team at Together recently received access to the NVIDIA Vera Rubin NVL72 platform. We spent the past few days digging through the new ISA and poking the chip with micros. There are many fun new features! We finished adding some functionality into ThunderKittens to write NVFP4 and FP8 GEMMs on Vera Rubin, as well as help fellow kittens explore the stars.

Astronomer Kittens. So cute! Before diving into what Vera Rubin delivers, we start our journey with a quick refresher of a Blackwell GPU GEMM. Starting Point: An NVIDIA HGX B200 GEMM NVIDIA Blackwell architecture’s fifth-generation tensor cores fundamentally changed the GEMM programming model. While NVIDIA Hopper architecture’s wgmma instruction was issued collectively by a warpgroup, Blackwell’s tcgen05 instruction is issued by a single thread, enabling one small producer warp to drive the tensor cores. The accumulator also moved out of registers into Tensor Memory and operands are read directly from shared memory, enabling a single MMA to span two CTAs across two SMs. To reach competitive performance on Blackwell, our GEMM: Launches threadblock clusters so each CTA pair can share operands by TMA multicast, cutting memory traffic from HBM by half. Specializes warps within clusters: loaders bring A and B into shared memory over TMA, a single MMA warp drives the tensor cores, and a consumer warpgroup carries finished accumulators from tensor memory to HBM. Runs persistently, with one tile's inputs streaming in while the previous tile's outputs are still draining.

Through these efforts, we realized the following results.

Read more about these kernels and their optimizations in our earlier Together blog post or the ThunderKittens 2.0 release! As Rubin preserves the Blackwell programming model, our old GEMMs still work. However, naively running them on Vera Rubin, we observe that our NVFP4 and FP8 kernels only achieve around 42.1% and 44.4% of the roofline — plenty of room for optimization!

The rest of this post is in two parts. First, we cover the new Rubin features that matter for a GEMM and how to use them in ThunderKittens. Then, we slowly integrate these features into our existing Blackwell NVFP4 kernel, taking it to over 22 PFLOPS and competitive with cuBLAS and CuTE DSL. The core issue we find is that while Rubin enables tensor cores to consume operands twice as fast, our old Blackwell kernel does not feed them fast enough. To reach the compute ceiling, we need tiles to squeeze more reuse out of the data they already have on-chip. What’s new with the NVIDIA Vera Rubin Platform? Comparing vendor specifications, we see the following improvements from Blackwell to Vera Rubin. NVIDIA HGX B200 NVIDIA Vera Rubin NVL72 NVFP4 Tensor Core 9 PFLOPS / GPU 35 PFLOPS / GPU FP8 Tensor Core 4.5 PFLOPS / GPU 17.5 PFLOPS / GPU FP16 / BF16 Tensor Core 2.25 PFLOPS / GPU 4 PFLOPS / GPU Memory Bandwidth 8 TB/s / GPU 22 TB/s / GPU SMs 148 / GPU 224 / GPU Peak Power 1000W / GPU 2300W / GPU

As it relates to writing performant GEMMs, we specifically take note of the subsequent features. 1. Tensor cores take twice the K For review, a tcgen05.mma computes C = A@B + C over a MxNxK tile, consuming a fixed number of bytes along K per step. On Blackwell, this step is 32 bytes, but on Vera Rubin it can be increased to 64 bytes. The MMA itself still takes the same number of cycles, so doubled-K allows us to pack twice as much work into the same instruction window.

In ThunderKittens, we express this through a new template parameter to our existing mma operation. mma_ABt (...); // Blackwell Default: 32-byte K step mma_ABt (...); // Vera Rubin: 64-byte K step 2. Tensor memory grows to 576 Columns Blackwell introduced the concept of tensor memory, a 128 lane x 512 column x 32-bit space which tensor cores could directly read and write to. On Vera Rubin, this space increases to 576 columns, providing an extra 32 KiB of tensor memory to play with. Note that these extra columns are reachable only through the .exclusive qualifier, a PTX 9.4 addition that ensures there’s only one live tensor memory allocation on an SM. Non-exclusive allocations are still capped at 512 and must be a power of two. In ThunderKittens users may request this with a template parameter to our tensor memory allocator, specifying that the allocation is exclusive. template<int _nblocks_per_sm, int _ncta, bool _managed = true, bool _exclusive = false&gt struct tensor_allocator { .... }

tensor_allocator tm; // Blackwell default: 512 columns tensor_allocator tm; // Rubin: Up to 576 columns 3. Shared memory increases to 328 KiB While Hopper and Blackwell provided 228 KiB of shared memory, Vera Rubin introduces an oversized shared memory mode that can be dynamically increased to 328 KiB. This is a host-side specification that can be called as below. CUfunction function = nullptr; cudaGetFuncBySymbol(&function,reinterpret_cast (kernel)); cuFuncSetAttribute(function, CU_FUNC_ATTRIBUTE_SHARED_MEMORY_MODE, CU_SHARED_MEMORY_MODE_ALLOW_OVERSIZED_SHARED_MEMORY); 4. B Side Collector Blackwell introduced the concept of a collector buffer, a small MMA staging buffer that could latch onto an A tile so the next instruction takes it from there instead of from shared memory. Vera Rubin extends this functionality to the B tile as well with .collector::b::* .

To utilize this, we annotate each MMA’s operands with one of four labels that describe its action on the collector buffer. “FILL” reads the operand from shared memory and latches it “USE” reads from the buffer “LASTUSE” reads from the buffer and releases it “DISCARD” is the default and skips the latch

These labels are permission qualifiers for reuse, not guarantees. This means that the Tensor Core could still reload a matrix even if it has permission to reuse. Now that either operand can be resident in the collector...