The Benchmark That Confirmed the Ceiling: NVFP4 Cutlass on sm_120
Introduction
In the high-stakes world of large language model inference optimization, progress is often measured in single-digit percentage gains. Every kernel rewrite, every quantization format change, every parallelism strategy shuffle yields a few more tokens per second—until you hit a wall. Message <msg id=12512> captures the moment when the assistant, after an exhaustive optimization campaign spanning multiple days, launches what should have been the decisive benchmark: a native FP4×FP4 tensor-core MoE backend for DeepSeek-V4-Flash on NVIDIA RTX PRO 6000 Blackwell GPUs. The result would confirm that even the most theoretically optimal quantization path cannot escape the deeper structural bottleneck of sm_120 fallback kernels.
This article examines that single message in detail: the reasoning that motivated it, the decisions embedded within it, the assumptions it carried, and the knowledge it ultimately produced. It is a story about the difference between theoretical hardware capability and practical kernel availability—and about how methodical measurement, not speculation, reveals the truth.
The Message in Full
The message is a bash command executed over SSH on a remote machine (CT200, root@10.1.230.171). It launches a benchmark script in the background, then enters a polling loop that checks every 30 seconds for completion, printing progress and finally dumping the full results:
timeout 30 ssh -o StrictHostKeyChecking=no root@10.1.230.171 '
cd /root && nohup bash /root/run_measure.sh NVFP4_cutlass > /root/dsv4_bench/run_nvfp4c.log 2>&1 & echo "PID $!"'
for i in $(seq 1 22); do
sleep 30
d=$(timeout 10 ssh -o StrictHostKeyChecking=no root@10.1.230.171 'grep -c ALL_DONE_NVFP4_cutlass /root/dsv4_bench/measure_NVFP4_cutlass.log 2>/dev/null')
c=$(timeout 10 ssh -o StrictHostKeyChecking=no root@10.1.230.171 'tail -1 /root/dsv4_bench/measure_NVFP4_cutlass.log' 2>/dev/null)
echo "[$((i*30))s] $c"
[ "$d" = "1" ] && break
done
echo "===== NVFP4+cutlass(native FP4) RESULTS ====="; timeout 12 ssh -o StrictHostKeyChecking=no root@10.1.230.171 'cat /root/dsv4_bench/measure_NVFP4_cutlass.log'
The output shows the benchmark progressing through concurrency levels C=1, C=8, and C=16, with the final results truncated at the end of the captured output.
Why This Message Was Written: The Reasoning and Motivation
To understand why this particular benchmark was launched, we must trace the optimization journey that preceded it. The assistant had been systematically working through a hierarchy of performance levers for DeepSeek-V4-Flash on sm_120 hardware (8× RTX PRO 6000 Blackwell GPUs, CUDA capability 12.0).
The Optimization Hierarchy
The campaign began with the stock MXFP4 checkpoint, which delivered a disappointing ~23 tok/s at C=16. A GPU profile revealed the root cause: two kernels consumed nearly 80% of decode time—the _mxfp4_slot_gemv_kernel (MoE, 39%) and the _tiled_sparse_decode_kernel (sparse attention, 38%)—both running on CUDA cores (SIMT) rather than tensor cores. The user's observation of less than 1% tensor-pipe utilization confirmed the diagnosis.
The assistant then worked through every available configuration lever:
- FP8 autotune configs: Generated five sm_120-specific block-GEMM tuning configurations, yielding only ~6% improvement at batch size 1 (since FP8 GEMM accounts for only 6% of decode time).
- NCCL LL+Ring tuning: Negligible effect (communication is only 2% of the profile).
- MTP/EAGLE speculative decoding: Boosted single-request throughput by 47%, but provided zero gain at concurrency due to verifier saturation.
- Expert parallelism (EP4): Made performance worse due to PCIe all-to-all overhead.
- Tilelang indexer fusion: Failed at JIT-compile time on sm_120/CUDA 13. None of these moved the needle meaningfully at C=16. The ceiling was structural: the fast fused DSA/MoE stack (DeepGEMM, trtllm-gen, FP4 indexer) is architecture-gated to SM100, and sm_120 has only Triton fallbacks.
The NVFP4 Pivot
The highest-leverage intervention identified was switching from the stock MXFP4 checkpoint to the official NVIDIA NVFP4 quantization (nvidia/DeepSeek-V4-Flash-NVFP4). The MXFP4 format (group-32, UE8M0) forces the slow CUDA-core _mxfp4_slot_gemv_kernel because Marlin produces NaNs on MXFP4 weights. NVFP4 (group-16, E4M3) unlocks tensor-core paths: either Marlin W4A16 (dequantize FP4→BF16, run on BF16 tensor cores) or the native cutlass FP4 grouped GEMM (true FP4×FP4 on sm_120a tensor cores).
The assistant applied PR #25820 as a patch to enable correct NVFP4 auto-detection from the checkpoint's hf_quant_config.json, downloaded the 149 GB checkpoint, and launched the first NVFP4 benchmark with the Marlin backend. The result: C=16 throughput improved from 23.2 to 28.7 tok/s (+24%), and time-to-first-token roughly halved. The MoE was now on tensor cores—but attention remained the dominant bottleneck at 38% of decode time.
The Cutlass Hypothesis
The reasoning behind message <msg id=12512> is captured in the preceding message <msg id=12509>. The assistant reasoned that Marlin's W4A16 approach still incurs dequantization overhead—it streams FP4 weights from memory and dequantizes them to BF16 before the tensor-core GEMM. The native FP4 path via the Triton backend (--moe-runner-backend triton) routes to cutlass_fp4_group_mm, a true FP4×FP4 grouped GEMM that operates entirely in FP4 on the sm_120a tensor cores, avoiding the dequantization step entirely.
The assistant's reasoning chain was:
"The triton backend with cutlass native FP4 should work on sm120 — the MoeRunner gets constructed fine and routes to cutlass_fp4_group_mm regardless. Let me test that path since it could outperform marlin."
This hypothesis was grounded in a sound understanding of the hardware: sm_120a tensor cores support native FP4 matrix multiplication, and the cutlass kernel (nvfp4_blockwise_moe.cuh) is already in the SGLang source tree. If the kernel executed efficiently, it could potentially outperform the Marlin W4A16 path by avoiding the bandwidth overhead of streaming and dequantizing weights.
How Decisions Were Made
The decision to benchmark NVFP4+cutlass was the product of a systematic, measurement-driven methodology. The assistant did not guess or speculate—it followed a clear pattern:
- Profile to identify the bottleneck: The torch profiler showed MoE at 39% and attention at 38% of decode time, both on CUDA cores.
- Identify the highest-leverage fix: Switching the MoE from CUDA-core slot-GEMV to tensor-core execution.
- Select the best available implementation: NVFP4 checkpoint + explicit MoE backend choice.
- Test the most promising option first: Marlin (W4A16) was the safer, known-working path.
- Then test the theoretically optimal option: Cutlass native FP4, which could be faster but was less proven on sm_120.
- Measure, don't assume: Launch a rigorous benchmark across multiple concurrency levels (C=1, C=8, C=16) with multiple samples each (n=8, n=32, n=48 respectively). The message itself encodes several design decisions about the measurement framework: - The benchmark script
run_measure.shis parameterized by a label (NVFP4_cutlass) for organized result tracking. - Results go to a structured log file (measure_NVFP4_cutlass.log) with a completion marker (ALL_DONE_NVFP4_cutlass). - The polling loop uses a 30-second interval, balancing responsiveness against overhead. - The timeout of 22 iterations (~11 minutes) provides an upper bound for the benchmark run.
Assumptions Embedded in the Message
Every benchmark carries assumptions, and this one is no exception:
1. The cutlass backend is functionally correct on sm_120
The assistant had verified correctness with a single test query (17×23 = 391) in <msg id=12511>, but the benchmark assumes this correctness generalizes across all decoding scenarios. This is a reasonable assumption—if the kernel produced correct output for one query, it likely works for many—but it's not a formal verification.
2. The benchmark results will be comparable to the Marlin results
The assistant assumes that the measurement methodology (same run_measure.sh script, same concurrency levels, same prompt lengths) produces comparable numbers. This is a sound engineering practice, but it assumes no confounding variables like GPU thermal state, memory fragmentation, or scheduler behavior differences between runs.
3. Native FP4 should outperform W4A16 Marlin
This is the core hypothesis being tested. The assumption is that avoiding the dequantization overhead will yield higher throughput. As we will see, this assumption turned out to be incorrect—both backends delivered essentially identical throughput.
4. The benchmark will complete within 11 minutes
The 22-iteration polling loop assumes the benchmark finishes in a predictable timeframe. This is based on prior runs with the Marlin backend, which completed within similar bounds.
5. The server is stable and won't crash during the benchmark
The assistant had just verified the server was "fired up and ready to roll" and produced correct output. The benchmark assumes continued stability under sustained load.
Mistakes and Incorrect Assumptions
The most significant incorrect assumption was that native FP4×FP4 tensor-core execution would outperform Marlin's W4A16 dequantization path. The results, which became visible in the following messages, showed both backends delivering nearly identical throughput (~28.5–28.7 tok/s at C=16). This is a fascinating result that reveals something fundamental about the hardware:
On sm_120, the tensor-core FP4 path and the BF16 path have similar effective throughput for this workload. The cutlass FP4 grouped GEMM, while theoretically more elegant, does not provide a meaningful advantage over the Marlin dequantization approach. This could be because:
- The FP4 tensor-core path on sm_120 may have lower effective utilization than the BF16 path.
- The dequantization overhead in Marlin is small relative to the overall kernel time.
- Both paths are limited by the same underlying memory bandwidth or instruction throughput constraints. This is not a mistake in the engineering sense—testing the hypothesis was the correct thing to do. But it reveals that the assistant's mental model of the performance hierarchy was incomplete. The real bottleneck was not the MoE backend at all, but the attention kernel, which remained at 38% of decode time regardless of which MoE backend was chosen. A more subtle issue is the assumption that benchmarking at C=16 captures the full performance picture. The user later pointed out that the hardware should be capable of 300–600 tok/s at C=16 based on roofline analysis (1.9 TB/s VRAM bandwidth, 13B active parameters), and the assistant discovered that the MTP verifier was consuming enough GPU memory to halve the effective batch size from 16 to 8. The benchmark was measuring throughput under a hidden capacity constraint, not the raw kernel performance.
Input Knowledge Required
To fully understand this message, one needs:
Hardware Knowledge
- NVIDIA RTX PRO 6000 Blackwell (sm_120): 188 SMs, 99 KB shared memory per block, CUDA capability 12.0, ~102 GB GDDR7 at ~1.5 TB/s, no NVLink, PCIe Gen5.
- sm_120 vs sm_100: The Blackwell architecture has two tiers. SM100 (B200/GB300) has the full fused kernel stack (DeepGEMM, trtllm-gen, FP4 indexer). SM120 (RTX PRO 6000) has tensor cores but relies on Triton/CUTLASS fallbacks for many operations.
- Tensor core types: sm_120a supports FP4, FP8, BF16, and INT8 tensor operations, but the highest-performance fused kernels are architecture-gated to SM100.
Software Knowledge
- SGLang: The serving framework, with its MoE runner abstraction supporting multiple backends (marlin, triton, flashinfer_trtllm_routed).
- NVFP4 quantization: NVIDIA's ModelOpt FP4 format (group-16, E4M3), distinct from the stock MXFP4 format (group-32, UE8M0).
- PR #25820: The SGLang pull request that adds NVFP4 auto-detection from
hf_quant_config.json. - Cutlass FP4 grouped GEMM: The
nvfp4_blockwise_moe.cuhkernel in the SGLang source tree for native FP4×FP4 tensor-core MoE computation.
Model Knowledge
- DeepSeek-V4-Flash: 284B total parameters, ~13B active (MoE with 256 experts, top-6 routing), 43 layers, hidden size 4096, MLA attention with 64 query heads and 512 top-k tokens.
- The NVFP4 checkpoint: 149 GB, stored at
/root/models/DeepSeek-V4-Flash-NVFP4.
Prior Optimization Context
- The assistant had already tested and measured: MXFP4 baseline, FP8 autotune, NCCL tuning, MTP, EP4, tilelang indexer, and NVFP4+Marlin.
- Each lever was benchmarked at C=1, C=8, and C=16 with results logged in
/root/dsv4_bench/.
Output Knowledge Created
This message, combined with the results it produced, created several pieces of actionable knowledge:
1. Native FP4 cutlass ≈ Marlin W4A16 on sm_120
The most important finding: both tensor-core MoE backends deliver essentially identical throughput (~28.5–28.7 tok/s at C=16). The theoretically superior native FP4 path does not provide a practical advantage. This is a significant data point for anyone deploying DeepSeek-V4 on sm_120 hardware.
2. The MoE bottleneck is resolved; attention remains
With either NVFP4 backend, the MoE computation is now on tensor cores and no longer the primary bottleneck. The attention kernel (_tiled_sparse_decode_kernel) at 38% of decode time is now the dominant constraint. Any further gains must come from optimizing the sparse MLA attention path.
3. The ceiling is structural, not configurable
The benchmark confirms that no amount of MoE backend tuning can break through the ~28 tok/s ceiling at C=16. The remaining bottleneck is the sm_120 fallback attention kernel, which runs on CUDA cores with low occupancy (64 blocks on ~170 SMs). This is the same pathology that plagued the earlier K2.6 verify kernel.
4. The measurement framework is validated
The run_measure.sh script, the polling loop, and the structured log format all proved reliable across multiple benchmark runs. The assistant could trust the numbers and focus on the next bottleneck.
5. The path forward is clear
With the MoE bottleneck addressed, the next step is to build a split-K tensor-core sparse-attention kernel—the same playbook that delivered 3–6× gains in the earlier K2.6 work. The assistant documented this roadmap in the following messages.
The Thinking Process Visible in Reasoning
The assistant's reasoning, visible in the agent reasoning blocks of surrounding messages, reveals a methodical and self-critical thought process:
Systematic Elimination
The assistant does not jump to conclusions. Each lever is tested independently, measured, and either adopted or discarded based on data. The progression from MXFP4 → FP8 configs → NCCL → MTP → EP4 → NVFP4+Marlin → NVFP4+Cutlass is a textbook example of the scientific method applied to systems optimization.
Self-Correction
When the Marlin results came in at only +24%, the assistant did not declare victory. Instead, it immediately analyzed why the gain was limited: "The gains are meaningful but not the dramatic speedup I was hoping for... Marlin uses FP4 weights dequantized to BF16 for tensor-core operations, which means it still has to stream and dequantize weights from memory—it's not a native FP4×FP4 computation."
This led to the hypothesis that native FP4 cutlass could be faster—a hypothesis that was then tested in this message.
Roofline Awareness
The assistant consistently reasons about the hardware's theoretical capability. In <msg id=12515>, after the cutlass results came in, it performs a detailed roofline analysis: "With 13B active parameters in FP4 (0.5 bytes per param), a decode step should read roughly 6.5 GB of weights... at 1.9 TB/s, that's only ~0.86 ms per GPU for weight loading." This awareness of the gap between theoretical and actual performance drives the investigation deeper.
Honesty About Limitations
The assistant acknowledges when it hits a wall. In the "Blocked" section of the goal summary (<msg id=12517>), it states: "~1000 t/s @C=16 is unreachable on sm_120 for DSv4... Realistic ceiling ~120-300 t/s agg with full kernel work; 1000 needs B200/SM100." This is not defeatism—it's a realistic assessment that guides the next steps.
Conclusion
Message <msg id=12512> appears, on its surface, to be a routine benchmark launch—a bash command, a polling loop, a results dump. But in the context of the broader optimization campaign, it represents a pivotal moment: the test of the theoretically optimal MoE backend, and the discovery that even the best available implementation cannot escape the structural bottleneck of sm_120 fallback kernels.
The message teaches us several lessons about systems optimization:
- Measure everything: The assistant's systematic benchmarking across multiple concurrency levels, with multiple samples each, produced trustworthy data that guided decision-making.
- Test your hypotheses: The assumption that native FP4 would outperform W4A16 Marlin was reasonable but wrong. Only measurement could reveal the truth.
- Know your bottleneck hierarchy: Fixing the MoE bottleneck merely exposed the attention bottleneck as the next constraint. Optimization is a process of peeling layers.
- Understand your hardware's capabilities and limitations: The sm_120 vs sm_100 distinction is not just a marketing difference—it determines which kernels are available and how fast they run. The benchmark confirmed that the ceiling was real, but it also clarified the path forward. The attention kernel, not the MoE backend, was now the target. And the assistant had already proven, in the earlier K2.6 work, that it could build custom kernels that deliver 3–6× speedups on sm_120. The stage was set for the next round of the optimization campaign.