The Methodical Campaign: How Evidence-Based Optimization Drove a 2x Throughput Breakthrough on Blackwell GPUs

Introduction

In the high-stakes world of large language model inference, optimization is rarely a single breakthrough moment. More often, it is a grinding, methodical process of forming hypotheses, implementing changes, running benchmarks, and confronting the data — whether it confirms your hopes or dashes them. This article examines a remarkable stretch of such work: a concentrated optimization campaign for the GLM-5-NVFP4 model (a 744-billion-parameter Mixture-of-Experts language model) running on eight NVIDIA RTX PRO 6000 Blackwell GPUs, served through the SGLang inference framework.

The chunk of conversation analyzed here — spanning messages from index 1063 to 1164 — represents a microcosm of the entire campaign. In this span, the assistant executed a series of optimization experiments that ranged from the mundane (a git pull that unlocked 2x throughput) to the ambitious (implementing a novel decode-time routing algorithm from a research paper). The results were mixed: one approach delivered a stunning breakthrough, another produced a clean null result, and a third crashed with a hardware limitation. But across all of them, the assistant maintained a consistent methodology: implement cleanly, benchmark rigorously against a proper baseline, document faithfully, and let the data guide the next move.

This article synthesizes the key themes, decisions, and discoveries of this chunk, drawing on the detailed message-level analyses that have been written for each individual step.

The Landscape Before the Push

To understand the significance of this chunk, one must appreciate the state of the campaign at its outset. The assistant had already spent many hours deploying the GLM-5-NVFP4 model, resolving environment issues (CUDA initialization errors, dependency conflicts, build failures), and testing a series of standard optimization approaches. The results had been sobering. Piecewise CUDA Graphs were blocked by torch.compile incompatibility with FlashInfer's FP4 JIT compilation. MSCCLPP and Single Batch Overlap showed negligible improvements (0-2%). TP4+PP2 was actually 2x slower than TP8, proving the bottleneck was compute-bound rather than communication-bound. Expert Parallelism (EP8) had crashed with an out-of-memory error at 256+ concurrency.

The assistant's message at index 1063 — a sprawling 3,000+ word status report — crystallized the situation. The core problem was now clear: the per-expert GEMM operations in GLM-5's MoE layers were tiny, achieving only 3-12% of peak FP4 throughput on the SM120 GPUs. With 256 experts and top-8 routing, each expert processed only batch_size × 8 / 256 tokens, resulting in memory-bandwidth-bound matrix multiplications that could not utilize the GPU's immense compute capability. The SM120 architecture compounded the problem: only 100KB of shared memory per SM (versus 228KB on B200), no Tensor Memory (TMEM), and warp-level mma.sync instructions rather than the autonomous tensor units of SM100.

This diagnosis was the foundation for everything that followed. The assistant knew it needed to either increase per-expert batch sizes (via Expert Parallelism or routing optimizations) or improve the efficiency of small GEMMs (via kernel engineering). The chunk that follows traces the pursuit of both paths.

The SGLang Update: A 2x Breakthrough from a Single Command

The first major action in this chunk was deceptively simple. After dispatching five parallel research agents to investigate the latest SGLang commits, SM120 kernel optimizations, GLM-5 deployment strategies, and Opportunistic Expert Activation, the assistant faced a strategic decision at message 1069: should it dive directly into implementing custom optimizations, or should it first check whether the upstream SGLang project had already solved some of its problems?

The decision to check first was a meta-optimization — optimizing the optimization process itself. The research agents had revealed that SGLang's main branch had accumulated numerous SM120-related fixes since the installed version (commit bba2fc4, only 9 commits behind origin/main). The assistant evaluated the gap, found it low-risk (a fast-forward merge with no conflicting changes), and executed the update at message 1076 with a simple git pull.

The result was extraordinary: the update alone yielded a 2x throughput improvement at 256 concurrency compared to earlier baselines. Where the assistant had been achieving approximately 353 tok/s before, the updated SGLang delivered ~718 tok/s — without any custom patches, without any algorithmic innovations, without any kernel engineering. The upstream SGLang team had fixed problems the assistant hadn't even identified.

This outcome validated a crucial engineering principle: before building custom solutions, ensure you are standing on the shoulders of the open-source community. The 2x gain from a 30-second git pull dwarfed every other optimization attempted in the entire campaign. It was a humbling reminder that in rapidly evolving open-source projects, the most impactful optimization is often "update to the latest version."

Opportunistic Expert Activation: A Clever Idea Confronts Reality

Emboldened by the SGLang update success, the assistant turned to a more ambitious optimization: Opportunistic Expert Activation (OEA), a decode-time routing technique described in arXiv paper 2511.02237. The idea was elegant. During autoregressive decoding, each token independently selects its top-k experts via the router. In a standard implementation, the inference engine must load up to batch_size × top_k unique expert weight matrices for each decoding step. OEA exploits the fact that many experts are shared across tokens: it identifies the most "popular" experts in the batch (the baseline set), then reroutes some tokens' remaining expert slots to already-loaded experts rather than introducing new ones. This reduces the total number of unique experts loaded per step, increasing per-expert batch size and improving GEMM efficiency.

The implementation, executed at message 1106, was careful and well-structured. The assistant added a new function opportunistic_expert_activation() to SGLang's topk.py, gated by the environment variable SGLANG_OEA_K0 (default 0, meaning disabled). The function was integrated as a post-processing step in select_experts(), applied only during decode mode (not prefill), with a try/except guard to gracefully fall back if the forward context was unavailable.

But the implementation required two critical bug fixes before it could be properly benchmarked. The first, discovered at message 1123, was that the initial version used raw router logits for weight gathering rather than sigmoid-transformed scores. Since sigmoid is monotonic, the ranking was preserved, but the renormalization step required actual probabilities. The fix was a single line: scores = router_logits.float().sigmoid().

The second bug, uncovered at message 1137-1140, was more subtle. The OEA function assumed that topk_ids[:, :k0] contained the k0 highest-scored experts for each token. But the fused top-k kernel in SGLang uses torch.topk with sorted=False for most routing paths, meaning the returned indices are in arbitrary order — not sorted by score. The assistant confirmed this by grepping the source code and finding sorted=False in three separate routing paths. The fix involved explicitly sorting the top-k indices by their sigmoid scores before selecting the baseline set.

With both fixes applied, the assistant ran a clean A/B benchmark at message 1157, alternating between OEA-enabled and baseline servers across four concurrency levels (10, 64, 256, 1024). The results were unambiguous:

| Concurrency | Baseline (tok/s) | OEA (tok/s) | Delta | |---|---|---|---| | 10 | 38.46 | 38.40 | -0.2% | | 64 | 218.74 | 218.36 | -0.2% | | 256 | 718.83 | 716.66 | -0.3% | | 1024 | 1,603.04 | 1,599.13 | -0.2% |

Average throughput was essentially identical — within measurement noise. Peak throughput showed a modest improvement at high concurrency (+3.6% at 256, +5.9% at 1024), but this did not translate to sustained gains.

The assistant's analysis was precise and honest: "This makes sense: with random input data, expert routing is approximately uniform. OEA's benefit depends on expert routing patterns having natural clustering (some experts being 'hotter' than others). With perfectly uniform random routing, all experts are equally loaded and there's minimal piggybacking opportunity."

This is a textbook example of how to handle a negative experimental result. The assistant did not search for post-hoc justifications, move the goalposts, or downplay the data. It accepted the result, explained the mechanism behind it, and documented it faithfully in the findings document. The OEA implementation was correct — the optimization simply could not help on a workload where the precondition (non-uniform expert routing) was absent.

Expert Parallelism Retry: A Hardware Constraint Bites

While the OEA benchmarks were running, the assistant also revisited Expert Parallelism (EP8) — an approach that had crashed with an OOM error earlier in the campaign. The theory behind EP8 was compelling: by sharding experts across GPUs (rather than sharding layers or hidden dimensions), each GPU would process 8x larger per-expert batches, potentially transforming memory-bandwidth-bound GEMMs into compute-bound ones.

The assistant retried EP8 at message 1115 with a memory-safe configuration: --mem-fraction-static 0.75 --max-running-requests 512. This reduced the memory reserved for KV cache and running state, leaving more headroom for the all-to-all communication buffers that had caused the earlier OOM crash.

The results were a mixed bag. The server loaded successfully, using 59.5GB of the available ~96GB per GPU — a comfortable margin with ~34GB free. But during warmup, the server crashed with a CUTLASS tile failure: the 128×256×128 tile configuration required more than SM120's 100KB shared memory limit (91KB for the mainloop plus 65KB for the epilogue). The server ultimately recovered and ran, but the crash during warmup signaled that EP8 was pushing against fundamental hardware constraints.

This episode illustrates a recurring theme in the campaign: the SM120 GPU architecture, while powerful, imposes hard limits that no amount of software optimization can fully circumvent. The 100KB shared memory per SM, the absence of TMEM, and the warp-level programming model are not bugs to be fixed — they are design constraints that must be worked around or accepted.

Single and Dual-Stream Benchmarking

Amidst the larger optimization experiments, the assistant also conducted focused benchmarks of single-stream and dual-stream throughput at messages 1154-1156. The results were excellent: 10.36 tok/s for a single concurrent request, and 19.29 tok/s for two concurrent requests — demonstrating near-perfect linear scaling. This was an important sanity check. It confirmed that the server was not bottlenecked on CPU-side scheduling or Python-level overhead at low concurrency, and that the GPU could handle multiple streams without interference.

These numbers also served as the foundation for the assistant's theoretical maximum performance analysis, which was in progress at the end of the chunk. By establishing the single-stream ceiling and measuring how throughput scales with concurrency, the assistant could estimate the HBM bandwidth-limited and PCIe allreduce-limited throughput ceilings for this exact model/hardware combination.

The Documentation Discipline

Throughout this chunk, the assistant maintained a meticulous documentation practice. At message 1135, it began writing the comprehensive glm5findings.md document — a 500+ line artifact covering environment setup, every optimization attempted (with benchmark data), configuration details, and lessons learned. The document was updated iteratively as new results came in: the OEA comparison table was added at message 1159, the baseline improvement from the SGLang update was recorded at message 1160, and the outstanding work items and lessons learned were updated at messages 1162-1163.

This documentation discipline serves multiple purposes. It creates a permanent record that prevents future wasted effort (no one will re-test OEA on random data without checking the findings document first). It captures negative results as faithfully as positive ones, ensuring that the knowledge generated by failed experiments is not lost. It provides a structured reference for the assistant's own future reasoning, reducing the cognitive load of maintaining a coherent mental model across dozens of messages. And it communicates progress transparently to the human collaborator, maintaining trust and enabling informed decision-making about next steps.

The todo list updates at message 1164 — marking research tasks as completed and closing the loop on the work cycle — are a final expression of this discipline. The assistant is not just doing work; it is managing work, maintaining awareness of what has been done, what remains, and how each piece fits into the larger puzzle.

The Theoretical Maximum Analysis

As the chunk drew to a close, the assistant began computing the theoretical maximum single-stream performance for this exact model/hardware combination. This analysis, initiated by gathering model config parameters from the GLM-5-NVFP4 configuration file at messages 1091-1094, aimed to establish the HBM bandwidth-limited and PCIe allreduce-limited throughput ceilings. By understanding the absolute upper bounds, the assistant could calibrate expectations: if the current throughput was already close to the theoretical maximum, further optimization would be futile. If there was a large gap, the gap itself would guide the search for the next bottleneck.

This kind of analysis — grounding empirical measurements in theoretical first principles — is characteristic of mature engineering practice. It prevents the common trap of optimizing in circles, improving one metric at the expense of another without ever approaching the true limit.

Themes and Lessons

Several overarching themes emerge from this chunk of work.

The primacy of methodology. The assistant's approach is consistently methodical: form a hypothesis, implement the change, benchmark against a clean baseline, document the results, and let the data guide the next step. This discipline is what enables the campaign to make progress even when individual experiments fail. The OEA null result is not a setback — it is knowledge gained.

The value of negative results. The OEA experiment produced a clean null result, and the assistant treated it as valuable information. The documentation of this result in glm5findings.md ensures that future researchers will know that OEA was tried and found ineffective on random data, and why (uniform expert routing leaves no clustering to exploit). In optimization work, negative results are as informative as positive ones — they rule out hypotheses and narrow the search space.

The humility of platform updates. The single biggest gain in this chunk — 2x throughput — came not from a custom optimization but from updating to the latest SGLang commit. This is a powerful reminder that in rapidly evolving open-source ecosystems, the most impactful action is often to ensure you are running the latest version. The assistant's decision to check the update feasibility before diving into custom implementations was strategically brilliant.

Hardware constraints are not negotiable. The EP8 retry and the CUTLASS tile failures repeatedly demonstrated that SM120's 100KB shared memory limit is a hard constraint. No amount of software ingenuity can make a 128×256×128 tile fit in 99KB of shared memory. The assistant's willingness to accept these constraints and work around them — rather than fighting them — is a sign of engineering maturity.

Documentation is not optional. The glm5findings.md document, the todo list updates, the careful A/B benchmark tables — these are not overhead. They are integral to the optimization process, enabling the assistant to maintain coherence across a complex, multi-hour campaign and to communicate effectively with the human collaborator.

Conclusion

This chunk of the GLM-5 optimization campaign is a microcosm of the entire endeavor: a methodical, evidence-based exploration of the optimization landscape for a massive MoE model on novel hardware. The assistant updated SGLang and got 2x throughput for free. It implemented a clever algorithmic optimization (OEA) and confronted a null result with intellectual honesty. It retried Expert Parallelism and hit a hardware wall. It benchmarked single-stream throughput and began computing theoretical maximums.

Through all of this, the assistant maintained a consistent approach: implement cleanly, benchmark rigorously, document faithfully, and let the data speak. The result is not just a set of optimized throughput numbers — it is a body of knowledge about what works, what doesn't, and why, for this specific model on this specific hardware. That knowledge, captured in the findings document and in the conversation itself, is perhaps the most valuable output of the entire campaign.