From 235 Watts to a Ranked Roadmap: The GLM-5-NVFP4 Optimization Campaign on Blackwell GPUs

Introduction

In the high-stakes world of large language model inference, the difference between a system that crawls and one that screams often comes down to a single question: where is the bottleneck? In this segment of an extended opencode session, an AI assistant and its user grapple with exactly that question while deploying the GLM-5-NVFP4 model—a 744-billion-parameter Mixture-of-Experts (MoE) transformer with 256 experts—across eight NVIDIA RTX PRO 6000 Blackwell GPUs. What unfolds is a masterclass in systematic performance engineering: benchmarking to confirm the bottleneck, deep kernel analysis to understand its root cause, server parameter tuning to extract a 28% throughput improvement, parallel research campaigns to explore radical alternatives, and disciplined documentation to capture everything for posterity.

This article synthesizes the key messages and discoveries from this segment, tracing the arc from hardware characterization through actionable optimization roadmap to the beginning of documentation as structured improvement files.

Part 1: The TP4+PP2 Benchmark — Confirming the Compute-Bound Hypothesis

The segment opens with a critical diagnostic pivot. In the previous segment, the assistant had been operating under the assumption that allreduce communication was the primary bottleneck. The 8 GPUs are split across two NUMA sockets with no NVLink, communicating solely over PCIe Gen5 at ~40-54 GB/s. Each forward pass requires 156 allreduce operations (two per layer across 78 layers). The hypothesis was that reducing the allreduce group size from 8 to 4 would cut communication overhead in half and improve throughput.

The user crystallizes this hypothesis in message 835 with a topology-aware proposal: "Can we slice the model up such that TP is 4, on each socket, and we infer half of the layers on one socket, half of the layers on the second socket?" [msg 835]. The assistant confirms this is Pipeline Parallelism (PP) and elaborates the specific mapping: PP stage 0 runs layers 0-38 on GPUs 0-3 (NUMA node 0) using TP4, and PP stage 1 runs layers 39-77 on GPUs 4-7 (NUMA node 1) using TP4 [msg 836].

The launch of TP4+PP2 is itself a saga of debugging. The assistant first encounters an ambiguous --pp flag that Python's argparse rejects because it could match --pp-size, --pp-max-micro-batch-size, or --pp-async-batch-depth [msg 842]. After correcting to --pp-size, the same error persists. The assistant traces the problem to a shell quoting nightmare: the launch command is a deeply nested SSH invocation with bash -c, nohup, line continuations, and escaped variable references, any of which could mangle the flag <msg id=845-846>. The solution is a methodological breakthrough: instead of continuing to fight with inline quoting, the assistant writes a standalone shell script to the remote machine using a heredoc, then executes it with a clean command [msg 847]. This script pivot—eliminating layers of indirection to create a reproducible artifact—is a textbook example of recognizing when the approach itself is flawed and pivoting to a fundamentally better strategy.

The TP4+PP2 server launches successfully, confirmed by PP0 TP0 and PP1 TP0 prefixes in the log <msg id=849-850>. But the benchmarks tell a stark story. At 256 concurrency, TP4+PP2 achieves 727 total tok/s versus TP8's 1,867 tok/s. At 512 concurrency, it achieves 1,318 tok/s versus 2,800 tok/s <msg id=855-857>. The user immediately recognizes the implication: "Doesn't the halving of perf sort of imply we're compute bound -> kernels not tuned to the GPU?" [msg 858].

The math is decisive. With TP8, each GPU computes 1/8th of each layer. With TP4+PP2, each GPU computes 1/4th of each layer (but only half the layers). Per-token compute per GPU is the same total, but the per-layer GEMM sizes double. If the model were communication-bound, halving the allreduce participants from 8 to 4 and keeping them NUMA-local should have improved throughput. Instead, throughput roughly halved. The ~2x slowdown directly correlates with the 2x larger GEMMs per GPU, confirming that the bottleneck is in the GPU's arithmetic units, not in the interconnect [msg 859].

Part 2: Deep Dive into FP4 GEMM Kernel Efficiency on SM120

With the compute-bound diagnosis confirmed, the assistant pivots to a deep investigation of the FP4 GEMM kernels on Blackwell's SM120 architecture. The investigation spans multiple dimensions: micro-benchmarking the raw kernel performance, analyzing the CUTLASS tile configurations, understanding the shared memory constraints, and measuring actual GPU power draw during inference.

The Micro-Benchmark Results

The assistant writes a series of Python scripts to benchmark flashinfer.mm_fp4 and sgl_kernel.cutlass_scaled_fp4_mm across a range of matrix sizes representative of the GLM-5 model's decode workload <msg id=867-875>. The results are sobering:

| M (batch) | N (output) | K (input) | TFLOPS | % of ~1,850 Peak | |---|---|---|---|---| | 1 | 768 | 6144 | 0.30 | 0.02% | | 64 | 768 | 6144 | 15.5 | 0.8% | | 256 | 768 | 6144 | 81.3 | 4.4% | | 512 | 768 | 6144 | 160.9 | 8.7% | | 1024 | 2048 | 6144 | 849 | 46% | | 4096 | 2048 | 6144 | 1,189 | 64% | | 8192 | 8192 | 8192 | 1,313 | 71% |

The kernels plateau around 1,200-1,300 TFLOPS for large matrices—about 65-70% of the estimated 1,850 TFLOPS dense peak. But during actual decode, each MoE expert processes at most batch_size × num_experts_per_token / total_experts tokens. With 256 experts, 8 activated per token, and 512 concurrent requests, each expert gets roughly 16 tokens on average. That puts the per-expert GEMM in the 0.8-55 TFLOPS range—only 0.02-3% of peak [msg 868].

The 99 Kilobyte Ceiling

The investigation into why larger CUTLASS tile configurations fail reveals a fundamental hardware constraint. The assistant queries CUDA device attributes and finds that SM120 provides 100 KB of shared memory per SM, but a CUDA kernel can only opt-in to a maximum of 99 KB (101,376 bytes) via cudaFuncSetAttribute [msg 912]. The assistant builds a Python estimation script that calculates shared memory requirements for different tile configurations, modeling FP4 packed storage (0.5 bytes per element), scale factors (0.0625 bytes per element), mainloop pipeline staging buffers, and epilogue storage [msg 913].

The results are revealing. The failing tiles—M128×N256×K128 and M256×N128×K128—require approximately 142 KB at 2 pipeline stages with finalize epilogue, far exceeding the 99 KB ceiling. The working tiles (128×128×128, 128×128×256) fit within the limit, but barely—the 128×128×128 tile with finalize epilogue and 2 stages uses exactly 100 KB, right at the boundary. The assistant confirms by examining the generated CUTLASS kernel files, finding that only 128×128×128 and 128×128×256 tile configurations survived the autotune; the 128×256×128 and 256×128×128 tiles failed with "Error Internal" during initializeMoeGroupedGemm <msg id=883-886>.

This analysis transforms an opaque error into a transparent, predictable hardware constraint. It also closes one optimization path: simply "fixing" the tiles is not feasible without a CUTLASS redesign for SM120, which would require either reducing pipeline stages (hurting latency hiding) or finding a way to use the full 100 KB per SM (which CUDA's API doesn't allow).

The Power Draw Revelation

While investigating kernel efficiency, the assistant makes a startling discovery. During a live benchmark at 1024 concurrency, the assistant queries nvidia-smi and finds that each of the eight RTX PRO 6000 GPUs is drawing only ~235W—roughly 39% of their 600W thermal design power—despite reporting 100% utilization <msg id=903-904>.

This is a classic symptom of low arithmetic intensity: the SMs are never idle (there is always work scheduled), but the work they are doing is not keeping the FP4 tensor core pipelines full. The kernels are spending most of their time moving data—loading weights, fetching activations, writing results—rather than performing actual floating-point operations. The 100% utilization metric from nvidia-smi measures whether the SMs have work scheduled, not whether the arithmetic pipelines are fully occupied.

The power draw revelation provides a concrete target for optimization: there is 2.5× more power headroom available if kernel efficiency can be improved. The GPUs have enormous compute capacity that the current software stack cannot effectively utilize.

Part 3: The 28% Breakthrough — Server Parameter Tuning

With the kernel-level optimization paths blocked by hardware constraints, the assistant pivots to a different lever: increasing the effective per-expert batch size through server parameter tuning. The core insight is elegant and data-driven. With 256 experts and 8 active per token, at 1024 concurrent tokens each expert processes only ~32 tokens. At 2048 concurrent tokens, that doubles to ~64 tokens per expert. Larger batches mean the GPU's tensor cores can amortize their overhead more effectively.

The assistant tunes two SGLang server parameters: --max-running-requests (raised from 1024 to 2048) and --num-continuous-decode-steps (raised from 4 to 8). The results are striking:

| Concurrency | Old (max_rr=1024, cds=4) | New (max_rr=2048, cds=8) | Change | |---|---|---|---| | 256 | 1,383 tot / 705 out | 1,384 tot / 705 out | Same | | 512 | 2,205 tot / 1,106 out | 2,216 tot / 1,111 out | Same | | 1024 | 3,190 tot / 1,601 out | 3,250 tot / 1,631 out | +2% | | 2048 | 3,249 tot / 1,640 out | 4,151 tot / 2,095 out | +28% |

The 28% improvement at 2048 concurrency validates the compute-bound bottleneck analysis and demonstrates that for MoE models on Blackwell GPUs, the most impactful optimization lever may not be in the kernel code at all, but in the scheduling parameters that control how many tokens the GPU processes at once [msg 950].

The assistant then tests num-continuous-decode-steps=16 and finds it produces essentially identical results to 8. This negative result is valuable: it establishes that scheduler overhead is negligible at this scale, and that the KV cache capacity (~495K tokens) is the throughput-limiting resource for concurrency scaling.

Part 4: The Pivot to "Wild and Ambitious Ideas"

Despite the 28% improvement, the user is not satisfied. In a pivotal message, the user writes: "We don't want to game the numbers, we want legit maximum throughput, in and out, at 1/2/10/1000 parallel; Strat research agents to explore wild and ambitious ideas to improve compute ops and to use all available PFLops."

This directive fundamentally shifts the trajectory of the work. The user rejects incremental parameter tuning and demands architecture-level innovation. The assistant's response is to launch six parallel research agents, each tasked with exploring a different dimension of the optimization space:

  1. Expert Parallelism (EP) — Could distributing experts across GPUs differently improve GEMM utilization? Under EP, each GPU owns a subset of experts and processes the full hidden dimension, yielding much better pipeline filling for FP4 tensor cores. Expected impact: 1.3-1.8× on MoE layers.
  2. Piecewise CUDA Graphs & SM120 Options — Could piecewise CUDA graph capture reduce kernel launch overhead for attention layers while leaving MoE layers dynamic? Expected impact: 10-20%.
  3. SM120 FP4 Kernel Optimization — What are the most cutting-edge approaches to maximize FP4 tensor core utilization on Blackwell? This included investigating persistent grouped GEMM kernels that fuse the gate-up projection, SiLU activation, and down-projection into a single persistent kernel. Expected impact: 2-5× speedup for small-batch MoE.
  4. Compute/Communication Overlap and Fusion — Could allreduce be overlapped with computation to hide latency on PCIe-only GPUs? This included MSCCLPP allreduce (Microsoft's low-level communication library offering IPC-based one-shot allreduce that is 2-4× faster than NCCL's ring pattern for small messages) and Single Batch Overlap (SBO) which overlaps the MoE combine-scatter operation with the down-projection GEMM by partitioning SMs.
  5. Alternative Model Formats — Would FP8 or other quantizations perform better on SM120? The investigation found that FP8 would require 2× memory bandwidth, making it worse for this memory-bandwidth-bound scenario.
  6. Memory Bandwidth Optimization — How could the 1.6 TB/s GDDR7 bandwidth be better utilized? This included L2 Cache Pinning for hot experts, where the 128 MB L2 cache could hold frequently-accessed expert weights, potentially achieving 50-80% hit rates and 1.5-3× effective bandwidth for MoE layers. This parallel research campaign represents a sophisticated use of the tool-calling infrastructure. Each research agent operates independently, and the parent session blocks until all six complete. The wall-clock time is roughly the duration of the longest-running agent, compressing what could be hours of sequential investigation into a single round.

Part 5: The Synthesis — A Ranked Optimization Plan

The six research agents return comprehensive findings, and the assistant synthesizes them into a ranked optimization plan. The plan is organized into three tiers based on expected impact and engineering effort:

Tier 1: Flag Flips (minutes each)

Part 6: The Documentation Pivot — glb5improvement-xx.md Files

The user's response to this comprehensive plan is not "execute it" but "document it." The instruction "Write down glb5improvement-xx.md for each" shifts the assistant's priority from execution to documentation. The assistant begins creating a series of files, starting with glb5improvement-01-piecewise-cuda-graphs.md.

This documentation pivot is a deliberate engineering discipline. Each file captures the motivation, technical details, implementation steps, risk assessment, and expected impact of a specific optimization approach. The numbering scheme creates a canonical ordering that maps directly to the execution priority. These files transform ephemeral conversation content into persistent, structured knowledge that can be consulted during execution, shared with collaborators, or revisited weeks later.

The assistant also attempts to test the piecewise CUDA graphs approach immediately, launching a server with --enable-piecewise-cuda-graph and --disable-cuda-graph [msg 970]. The server crashes during graph capture because torch dynamo traces into flashinfer/fp4_quantization.py, which calls subprocess to check the CUDA version—a function that dynamo cannot trace [msg 980]. The assistant identifies the root cause and explores fixes, including patching is_cuda_version_at_least to return True directly since CUDA 12.8 is confirmed, or wrapping fp4_quantize with torch.compiler.disable() to make it opaque to dynamo <msg id=990-1010>. However, the piecewise graph compiler uses fullgraph=True which doesn't allow graph breaks, creating a fundamental incompatibility.

Conclusion: The Arc of Optimization

This segment of the opencode session tells a compelling story about the nature of ML inference optimization. It begins with a hypothesis test (TP4+PP2 benchmarking) that decisively confirms the model is compute-bound. It moves through deep kernel analysis that reveals the 99KB SMEM ceiling, the 235W power draw, and the catastrophic utilization collapse at small batch sizes. It achieves a concrete victory through server parameter tuning—the 28% breakthrough at 2048 concurrency. It then pivots to radical exploration through six parallel research agents, synthesizes their findings into a ranked optimization plan, and begins documenting each approach as structured improvement files.

Several themes emerge. First, the importance of understanding the bottleneck at multiple levels: the model is compute-bound overall, but the individual GEMM kernels are memory-bandwidth-bound at decode batch sizes. Second, the value of negative results: knowing that larger CUTLASS tiles don't fit, that cuBLASLt isn't faster, that CDS=16 doesn't help, and that TP4+PP2 is 2× slower, prevents wasted effort and sharpens the focus on what might work. Third, the power of parallel exploration: launching six research agents simultaneously compressed weeks of sequential investigation into a single round.

The 28% throughput improvement is a real achievement, but the deeper value of this session lies in the knowledge it generates. The ranked optimization plan, the negative-space map, and the glb5improvement-xx.md documentation series together form a comprehensive roadmap for extracting maximum performance from Blackwell GPUs for MoE inference. Whether the "wild and ambitious ideas" succeed or fail, the systematic methodology—diagnose, research, rank, document, execute—provides a template that can be applied to any complex optimization problem.

The arc from the TP4+PP2 benchmark to the glb5improvement files is a complete engineering narrative: hypothesis formation, experimental validation, deep analysis, parameter tuning, radical exploration, synthesis, and documentation. It is a case study in how systematic engineering discipline enables progress in the most challenging of technical domains.## References

[1] "From Knowledge Artifact to Pipeline Parallelism: A Strategic Pivot in Blackwell Inference Optimization" — Chunk article covering the TP4+PP2 launch saga, from knowledge artifact through shell quoting nightmare to successful server launch.

[2] "From 235 Watts to 28%: The GLM-5-NVFP4 Optimization Journey on Blackwell GPUs" — Chunk article covering the compute-bound confirmation, FP4 GEMM kernel analysis, 28% throughput breakthrough, and parallel research campaign.