Skip to content
Work Writing About Contact

Work / qgemm-mx

active 2026

qgemm-mx

Recovering block-scaled FP4 throughput on GPUs that have no native FP4 multiply. The memory saving already ships; the speed is being left on the floor, and this measures exactly where it goes.

  • CUDA
  • Quantization
  • Hopper
  • Kernels
Code

The gap

Block-scaled 4-bit weight formats — MXFP4, NVFP4 — are shipping in production checkpoints today. Native tensor-core support for them arrived with Blackwell, which is one hardware generation ahead of most of the capacity actually deployed. On Hopper there is no native FP4 multiply. So weights get stored at four bits and dequantized inside the kernel, and the current production path therefore keeps the memory saving and gives the speed back.

In the bandwidth-bound decode regime, MXFP4 moves roughly 1.88× fewer bytes per weight than FP8. The headroom is known in advance from the byte count alone. Close to none of it is being realised. This project is about measuring where it goes and how much can be recovered with a wgmma/TMA kernel that never materialises the dequantized weights.

Where it actually stands

Early, and staged deliberately so that each rung has to be measured before the next one starts.

  • R0 — measurement harness and correctness floor. Done locally: benchmark harness, host-side packing and probes, an occupancy table, the full frozen shape set on sm_86 against both a stream baseline and cuBLAS, a naive-dequant performance floor, and a reference kernel that matches the host packer bit-for-bit on device. Build verification on sm_90 is in.
  • R1 — decompose the Hopper FP4 gap. Not started.
  • R2 — the wgmma/TMA kernel that closes it. The A-fragment layout table and host-side permute are done. The kernel itself is not written.
  • R3 — MoE / grouped variant and a DSL comparison. Not started.

The reference kernel that exists today is correctness only — one thread per output, sequential nibble order, no tensor cores. It proves the encodings are right on device. It is not a performance result and the repository forbids citing it as one.

There is a preregistration document holding the predicted numbers, written before the kernel exists, so the eventual result can be checked against what was expected rather than against what would have been convenient.

Stack

CUDA C++ · Hopper wgmma / TMA · MXFP4 / NVFP4 · cuBLAS (baseline) · Python harness