The Synthesis: How Six Research Agents Forged a Roadmap to Blackwell MoE Performance

Introduction

In the high-stakes world of large-scale ML inference, the difference between a system that crawls and one that screams often comes down to a single moment of clarity—a pivot from blind experimentation to strategic, research-driven optimization. Message 954 in this opencode session represents exactly that pivot. After weeks of battling NVIDIA driver incompatibilities, CUDA toolkit mismatches, PCIe topology constraints, and elusive NaN crashes during decode, the assistant had finally achieved a working deployment of the GLM-5-NVFP4 model—a 744-billion-parameter Mixture-of-Experts (MoE) transformer—on eight RTX PRO 6000 Blackwell GPUs. The system was running, but it was running at only 235 watts out of a 600-watt TDP, achieving a mere 3% of peak FP4 tensor core throughput on the critical MoE GEMM operations.

This message is the turning point. It is the moment the assistant stopped "gaming the numbers"—as the user had sharply demanded in [msg 951]—and instead launched a coordinated research campaign, synthesized the results into a coherent model of the bottleneck, and produced a ranked, actionable optimization plan. It is a masterclass in systematic performance engineering: diagnose, research, rank, execute. Let us examine this message in depth.

The Context: Why This Message Was Written

To understand message 954, we must understand what came immediately before it. The conversation had reached an inflection point. The assistant had spent several rounds benchmarking the server with different concurrency levels and num-continuous-decode-steps values, achieving a respectable 28% improvement at 2048 concurrent requests (from 1,640 to 2,095 output tok/s). But the user's response in [msg 951] was a rebuke: "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 changed everything. The assistant had been optimizing within a narrow band—tweaking server parameters, adjusting concurrency, trying different decode step counts. The user was demanding a fundamentally different approach: stop optimizing the benchmark and start optimizing the physics. The assistant's response in [msg 952] was 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?
  2. CUDA Graphs & SM120 Options — Could CUDA graph capture reduce kernel launch overhead despite MoE's dynamic nature?
  3. SM120 FP4 Kernel Optimization — What are the most cutting-edge approaches to maximize FP4 tensor core utilization on Blackwell?
  4. Overlap and Fusion Strategies — Could compute/communication overlap help on PCIe-only GPUs?
  5. Alternative Model Formats — Are there other quantization schemes or model variants that would perform better?
  6. Memory Bandwidth Optimization — How to maximize GDDR7 bandwidth utilization for memory-bound operations? These agents ran in parallel (the task tool spawns subagents that execute independently), each producing a detailed research document. Message 954 is the synthesis of all six reports into a coherent, prioritized plan.

The Message Structure: A Blueprint for Action

The message is organized with remarkable clarity. It opens with a bold header—"Research Results: Ranked Optimization Plan"—and immediately establishes a two-dimensional prioritization framework: expected impact and effort. This is classic engineering decision-making: you want the highest-impact changes that cost the least to implement, and you sequence them to build momentum.

The plan is divided into three tiers, each representing a different level of engineering investment:

Tier 1: Flag Flips — Try Immediately (minutes each)

This tier contains three optimizations that require no code changes, only command-line flag modifications:

1a. Enable Piecewise CUDA Graphs. The server was running with --disable-cuda-graph because standard CUDA graphs require fixed tensor shapes, which MoE layers violate (different experts are activated each iteration). However, the research agent discovered that FlashInfer has a "piecewise" CUDA graph mode that captures the attention layers—which do have fixed shapes—as CUDA graph segments, while leaving the MoE layers dynamic. This eliminates kernel launch overhead for roughly half of each transformer layer. The assistant estimates 10-20% throughput improvement across all concurrencies, with low risk since the code path already exists in FlashInfer (layer.py:1272).

1b. Enable MSCCLPP. Microsoft's low-level communication library offers IPC-based one-shot allreduce that is 2-4x faster than NCCL's ring pattern for small messages (typical of decode). The assistant notes that MSCCLPP falls back to NCCL for unsupported cases, making this a safe flag change with potentially 10-30% improvement for low-concurrency scenarios.

1c. Enable Single Batch Overlap (SBO). This technique overlaps the MoE combine-scatter operation with the down-projection GEMM by partitioning SMs: 32 SMs handle communication while the rest compute. The risk is medium because it requires the flashinfer_cutedsl backend path, but the expected impact is 10-15% MoE layer speedup.

Tier 2: Small Code Changes (hours)

2a. Re-enable FlashInfer Allreduce Fusion for SM120. This is the most interesting entry in Tier 2 because it comes with a cautionary tale. The fused allreduce+residual+RMSNorm kernel is gated on (_is_sm90_supported or _is_sm100_supported), which explicitly excludes SM120 (Blackwell workstation). The assistant had attempted to patch this earlier and it "destroyed performance"—the cudaGridDependencySynchronize calls in trtllm_allreduce.cuh don't work on SM120, causing an 8x slowdown. However, the assistant now notes that MSCCLPP uses a completely different code path and might work where FlashInfer fusion failed. This is a nuanced insight: just because one approach failed doesn't mean all approaches will fail. The risk is high, but if it works, the payoff is 20-30% improvement.

2b. Expert Parallelism (EP8) instead of TP8. This is arguably the most architecturally significant proposal in the entire plan. Under Tensor Parallelism (TP8), each expert GEMM operates on a tiny slice of the hidden dimension (N=256, K=6144), resulting in only 2 K-tiles—pathologically small for FP4 tensor cores, which need many tiles to pipeline effectively. Under Expert Parallelism (EP8), each GPU owns a subset of experts and processes the full hidden dimension (N=2048, K=6144), yielding 16 K-tiles and much better pipeline filling. Moreover, EP actually moves less data than TP's allreduce. The expected impact is 1.3-1.8x on MoE layers, which dominate the compute path.

Tier 3: Significant Engineering (days)

3a. L2 Cache Pinning for Hot Experts. The RTX PRO 6000 has 128 MB of L2 cache, which can hold 13-85 experts (depending on quantization). With Zipf-distributed routing (common in MoE models), popular experts stay hot, potentially achieving 50-80% L2 hit rates. The cudaAccessPolicyWindow API could pin frequently-accessed expert weights in L2, and combined with column-major GEMM scheduling, could achieve 4.4x better L2 utilization. Expected impact: 1.5-3x effective bandwidth for MoE layers.

3b. Persistent Grouped GEMM Kernel. This would fuse the gate-up projection, SiLU activation, and down-projection into a single persistent kernel, using Grid Dependency Control (GDC) to chain layers without launch gaps. Research shows 2-5x speedup for small-batch MoE on H100. The assistant correctly notes this is the largest engineering effort with the largest potential for single-stream performance.

3c. FP4 2:4 Structured Sparsity. SM120 natively supports sparse FP4 via SM120_SPARSE_16x8x128_TN_VS, achieving 1,872 TFLOPS on RTX 5080 (2x dense). This would require pruning GLM-5 weights to a 2:4 pattern and fine-tuning. The risk is very high—accuracy loss from pruning on top of FP4 quantization could be catastrophic—but the potential compute uplift is 2x.

What Was Ruled Out: The Negative Space

A crucial feature of this message is the "What WON'T Help" table. This demonstrates intellectual rigor: the assistant doesn't just propose optimizations; it also explicitly closes off dead ends that were investigated. This prevents wasted effort and builds confidence in the plan.

The ruled-out approaches include:

The Thinking Process: How Decisions Were Made

The reasoning visible in this message reveals a sophisticated decision-making framework. The assistant is not merely listing options; it is evaluating them along multiple axes simultaneously:

Impact vs. Effort Tradeoff. Every proposal is rated on both dimensions. Tier 1 items are "minutes each" with 10-30% potential impact. Tier 2 items are "hours" with 1.3-1.8x potential. Tier 3 items are "days" with 1.5-5x potential. This allows the user to make informed choices about where to invest time.

Risk Assessment. Each proposal includes a risk rating (low, medium, high) and, crucially, why the risk exists. For example, the allreduce fusion re-attempt is rated high-risk not because the assistant is uncertain, but because a previous attempt caused an 8x slowdown and the root cause is not fully understood. This is honest engineering.

Dependency Awareness. The assistant understands that some optimizations are independent (can be tested in parallel) while others are sequential. The proposed execution order reflects this: start with the three independent flag flips (all can be tested within the hour), then move to EP8 (the biggest architectural change), then L2 pinning, then re-investigate allreduce fusion, and finally the persistent kernel.

Stacking vs. Diminishing Returns. The assistant notes that "any combination of them could stack," acknowledging that performance optimizations often have complex interactions. The expected impacts are rough estimates, and the only way to know the true combined effect is to measure.

Assumptions and Potential Mistakes

The message makes several assumptions that are worth examining critically:

Assumption: Piecewise CUDA graphs are safe. The assistant assumes that because the code path exists in FlashInfer (layer.py:1272), it will work correctly on SM120. However, the piecewise graph code path may have been developed and tested only on H100 (SM90) or B200 (SM100). SM120 (Blackwell workstation/consumer) has different hardware scheduling characteristics. The risk is rated low, but there is always a possibility of subtle correctness issues or performance cliffs.

Assumption: MSCCLPP falls back gracefully. The assistant states that MSCCLPP "falls back to NCCL for unsupported cases," which is true in theory, but the fallback mechanism may itself introduce overhead or bugs. The assistant's risk assessment of "low" may be optimistic.

Assumption: EP8 will work on PCIe without NVLink. The assistant notes that FlashInfer's A2A (all-to-all) backend uses NCCL and works on PCIe, but the NUMA topology (4 GPUs per socket, cross-socket SYS links) could introduce contention that negates the benefits. The risk is rated "medium," which seems appropriate, but the actual performance could vary wildly depending on the all-to-all implementation details.

Potential Mistake: Overestimating L2 cache pinning impact. The 128 MB L2 cache is shared across all 188 SMs. Pinning expert weights in L2 could starve other operations (attention, allreduce) of cache capacity. The 1.5-3x bandwidth improvement assumes that the bottleneck is purely memory bandwidth for MoE GEMMs, but if attention or communication becomes the bottleneck after pinning, the overall gain could be much smaller.

Potential Mistake: Underestimating the persistent kernel effort. The assistant notes that research shows 2-5x speedup for small-batch MoE on H100, but this is a research result, not a production implementation. Building a persistent grouped GEMM kernel that works correctly across all layers, all expert routing patterns, and all batch sizes is a significant software engineering project. The "days" estimate may be optimistic.

Input Knowledge Required

To fully understand this message, the reader needs knowledge in several domains:

MoE Architecture. The message assumes familiarity with Mixture-of-Experts models: how routers select experts, how tokens are dispatched and combined, and why expert parallelism differs from tensor parallelism. The reader must understand that in TP, each GPU holds a slice of every expert's weights; in EP, each GPU holds all weights for a subset of experts.

GPU Architecture (SM120). The message references SM120-specific constraints: the 99 KB shared memory limit, the cudaGridDependencySynchronize API, the cudaAccessPolicyWindow API, and the sparse FP4 tensor core support. The reader needs to know that SM120 is the Blackwell workstation/consumer architecture (RTX 5000 series, RTX PRO 6000), which uses mma.sync instructions (Ampere-era programming model) rather than the sm100 warp-group specialization.

CUDA Graph Semantics. The distinction between standard CUDA graphs (fixed tensor shapes) and piecewise CUDA graphs (dynamic segments for MoE) is critical. The reader must understand why MoE layers are incompatible with standard graph capture.

Communication Primitives. The message discusses NCCL ring allreduce, MSCCLPP IPC allreduce, and all-to-all (A2A) communication patterns. The reader needs to understand how these differ in latency, bandwidth utilization, and hardware requirements.

Quantization Formats. The message references NVFP4 (NVIDIA's 4-bit floating point format), FP8, and structured sparsity patterns (2:4). The reader needs to understand that FP4 halves the memory bandwidth requirement compared to FP8, which is why the FP8 model variant was ruled out.

Output Knowledge Created

This message creates several forms of valuable knowledge:

A Prioritized Optimization Roadmap. The primary output is a concrete, actionable plan with ranked items, expected impacts, risk assessments, and execution order. This is immediately useful for the engineering team.

A Negative-Space Map. The "What WON'T Help" table encodes hard-won negative results. This prevents future engineers from wasting time on dead ends and serves as a reference for similar deployments.

A Bottleneck Model. The message implicitly contains a model of the bottleneck: the workload is memory-bandwidth-bound at decode batch sizes, with per-expert GEMMs achieving only 3% of peak FP4 throughput due to tiny tile sizes. This model justifies the proposed optimizations and provides a framework for evaluating their success.

A Research Methodology. The way the assistant decomposed the problem into six parallel research dimensions—expert parallelism, CUDA graphs, kernel optimization, overlap/fusion, alternative formats, memory bandwidth—is itself a valuable output. Future optimization efforts can follow the same pattern.

Conclusion

Message 954 is a pivotal document in this coding session. It represents the transition from empirical tuning to systematic, research-driven optimization. The assistant took the user's directive—"explore wild and ambitious ideas"—and executed a coordinated research campaign across six dimensions, then synthesized the results into a coherent plan with ranked priorities, risk assessments, and a clear execution order.

What makes this message exceptional is not just the quality of the analysis, but the intellectual honesty. The assistant explicitly rules out approaches that won't work, acknowledges risks and uncertainties, and provides the reasoning behind each recommendation. The "What WON'T Help" table is particularly valuable because it encodes hard-won negative results that would otherwise be lost.

The proposed execution order—flag flips first, then EP8, then L2 pinning, then re-investigating allreduce fusion, then the persistent kernel—reflects a deep understanding of engineering economics: get the quick wins first to build momentum, then tackle the architecturally significant changes, and finally invest in the high-risk, high-reward research projects.

For anyone working on large-scale MoE inference, particularly on Blackwell GPUs, this message is a treasure trove of practical knowledge. It demonstrates how to systematically decompose a performance problem, research solutions across multiple dimensions, and synthesize the results into an actionable plan. It is a model of engineering decision-making under uncertainty.