When NVIDIA unveiled its next-generation Blackwell GPU architecture, the headline numbers promised historic leaps in LLM training and inference performance. At the center of the pitch was NVFP4 — a hyper-dense 4-bit floating-point format meant to shrink LLM VRAM footprints by roughly 4× without catastrophic accuracy loss. What went largely unpublicized: not all Blackwell chips are created equal, and the gap between them reshapes how NVFP4 actually has to be programmed.
Through systematic, programmatic brute-forcing of NVIDIA's compiler tool (ptxas) across hundreds of iterations, hardware researcher Jetha Chan exposed a deep hardware divide between data center Blackwell silicon and consumer/edge Blackwell silicon. This piece walks through that architectural split, the software crisis it caused for consumer AI deployment, and the engineering workarounds — including a custom native A4Q attention kernel — built to bridge the gap.
The Architectural Divide SM100 vs. SM12x
Historically, NVIDIA kept its consumer and enterprise execution architectures aligned at the instruction-set level. With Blackwell, that continuity breaks. The lineage splits into two distinct programming models:
| Feature | Datacenter · SM100 | Consumer / Edge · SM12x |
|---|---|---|
| Target processors | B100, B200, GB200 Superchip | RTX 5090 (SM120), RTX PRO (SM121) |
| Primary tensor ISA | tcgen05.mma (5th-gen) |
Extended mma.sync |
| Execution model | Single-thread async dispatch | Warp-synchronous, 32 threads lock-step |
| Accumulator storage | TMEM — dedicated 256 KB | Registers — shared with CUDA vector units |
| Warpgroup MMA | Native hardware support | Software emulation only |
In chips like the B200, Tensor Cores stop relying on the standard register file entirely. TMEM, a completely separate physical memory layer, pairs with tcgen05.mma so a single thread can asynchronously command the Tensor Cores to fetch, compute, and write back — unburdening the register file and isolating execution lines.
To cut cost and maximize die space, NVIDIA stripped TMEM entirely out of the RTX 5090 sm120 and workstation sm121. Without it, these cards can't parse tcgen05 at all — they fall back to mma.sync, which forces all 32 threads in a warp to synchronize and serialize data through local registers.
Consumer Blackwell GPUs possess the raw hardware structures to understand NVFP4 and FP6 — but they lack the dedicated memory highway to route that data autonomously. They must compute modern low-bit math using legacy scheduling tracks.
The Problem Shared Memory Brick Wall
This split caused immediate chaos for open-source AI deployment stacks trying to run compressed models natively on consumer GPUs. Enterprise software paths — NVIDIA TensorRT-LLM, FlashAttention-4 — were compiled on the assumption that the underlying hardware had TMEM and tcgen05 support.
On an RTX 5090, the software instead fell back to slow software-emulation: ingest NVFP4 from global memory, unpack it via bit-manipulation up to FP8, then hand it to the Tensor Cores. That created two compounding bottlenecks.
Modern architectures like Gemma 4 use exceptionally wide attention layers — 512-wide heads. Staging an entire 512-wide matrix block into Shared Memory after unpacking to FP8 blew straight through the 99 KB SM120 allowance. Context length was hard-capped; long prompts triggered unrecoverable out-of-memory faults.
The Breakthrough The A4Q Kernel
To route around the enterprise code barrier, Jetha Chan engineered a custom native attention kernel — A4Q (Attention with 4-bit Q) — designed to bypass shared-memory scaling entirely by managing data at the register layer instead.
Rather than decompressing sequences into Shared Memory before computation, A4Q streams packed NVFP4 data directly into registers and performs in-loop dequantization — extracting and scaling exactly 16 low-bit values at a time, right as the Tensor Core requests them. That keeps the active footprint well under the 99 KB ceiling and enables 128k+ context windows on a standard desktop card.
Preserving precision — dual-tier scale handling
An NVFP4 value is extremely narrow — 1 sign bit, 2 exponent bits, 1 mantissa bit — so its native dynamic range is tight. To stop outputs from decaying into gibberish over long sequences, NVFP4 leans on dual-tier scaling: a global tensor scale plus a localized block-scale for every 16 numbers. A4Q integrates that math natively into the register file, managing distinct asymmetric scaling blocks for Keys (K) and Values (V) on the fly.
// Conceptual loop inside A4Q register-level dequantization #pragma unroll for (int i = 0; i < 16; ++i) { uint32_t packed_val = fetch_nvfp4_from_reg(reg_array, i); float scale_k = compute_asymmetric_scale_k(block_scale_array, i); float dequantized_k = dequantize_fp4_to_float(packed_val) * scale_k; // Immediate injection into the mma.sync accumulator loop asm volatile("mma.sync.aligned.m16n8k32..." : ...); }
By contributing these architecture-specific patches directly to open-source engines — vLLM, FlashInfer, SGLang — this register-level approach let an RTX 5090 host heavily compressed models like Qwen 2.5 35B MoE or Gemma 4 locally at speeds pushing past 160–175+ tokens per second.
Alternative Pathways what else the ecosystem tried
Register-level assembly gives maximum throughput, but it isn't the only route around the consumer Blackwell constraints.
Hybrid split precision
Base model weights stay heavily compressed in NVFP4 (via Marlin or CUTLASS) to save VRAM, while the volatile attention layers and KV cache stay in FP8 or BF16 — sidestepping the 99 KB wall entirely, since the attention loop never has to unpack multi-layered 4-bit tensors.
SageAttention 3 & Quantization-Aware Distillation
Common in image and video diffusion pipelines: models are trained to structurally expect the precision limits of consumer Blackwell cores, shifting the burden from runtime assembly to training time. No low-level register manipulation required.
Triton fallback compilers
Developers write high-level Python describing NVFP4 tensor manipulations; Triton handles block tiling, memory scheduling, and register mapping for sm120 automatically. Flexible, but typically leaves some hardware efficiency on the table versus a hand-tuned kernel.
Dual-chunk FlashAttention execution
Segments the prompt into uneven, asymmetric chunks of Key/Value pairs, iteratively computing and recombining partial attention scores — capping shared-memory overhead per block to a safe, static size regardless of sequence length.
Conclusion
The Blackwell split is a reminder that software optimization can't treat a GPU generation as a monolith. Removing TMEM from consumer silicon created a hardware profile that demands a distinct programming paradigm. Enterprise environments get autonomous hardware loops via tcgen05; the open-source community has shown — through developments like A4Q — that meticulous, register-level engineering can still get consumer hardware to stellar local inference speeds.