The Last Config Lever: Expert Parallelism and the Architecture Ceiling on Blackwell

Introduction

In the high-stakes world of large language model inference optimization, there comes a moment when every configuration knob has been turned, every environment variable set, every backend swapped—and the performance needle still refuses to budge. Message 12433 in this opencode session captures exactly that moment. The assistant, having spent dozens of rounds deploying and tuning DeepSeek-V4-Flash on eight NVIDIA RTX PRO 6000 Blackwell GPUs (sm_120 architecture), arrives at a sobering diagnosis: the model's DSA (Dynamic Sparse Attention) architecture relies on fused kernels that only exist for SM100 hardware, and on sm_120, every critical component falls back to slow, latency-bound torch and Triton operations. The target of ~1000 tokens per second at concurrency 16 remains out of reach by a factor of forty.

But rather than simply report this conclusion, the assistant does something more valuable: it pauses, re-examines its assumptions, identifies one remaining untested lever—Expert Parallelism (EP)—and pivots decisively to try it. This message is a masterclass in structured reasoning under uncertainty, the discipline of exhausting all options before conceding a bottleneck, and the art of knowing when to keep turning knobs versus when to accept that the hardware architecture itself is the constraint.

The Message in Context

To understand message 12433, one must appreciate the optimization campaign that precedes it. The assistant has been working with DeepSeek-V4-Flash, a 256-expert Mixture-of-Experts model with a sophisticated DSA attention mechanism, deployed across eight RTX PRO 6000 GPUs connected via PCIe. The hardware is formidable—1.9 TB/s of aggregate VRAM bandwidth, 170+ streaming multiprocessors—but the software stack is a patchwork of fallbacks.

The critical architectural fact is that DeepSeek-V4-Flash's DSA pipeline was designed for NVIDIA's SM100 architecture (the Blackwell B200/B300 datacenter GPUs). On SM100, the "lightning indexer" (which selects which tokens participate in attention), the manifold-constrained hyper-connections (mHC) with Sinkhorn normalization, and the Mixture-of-Experts dispatch all run through fused, hand-tuned kernels that achieve high occupancy and tensor-core utilization. On sm_120 (the RTX PRO 6000 Blackwell workstation GPU), these fused kernels are absent. SGLang's hook code explicitly disables them, forcing fallback to torch operations and Triton kernels that are fundamentally latency-bound.

Previous messages in this segment document the systematic exploration: NCCL protocol tuning (LL, Ring) had negligible effect because communication accounts for only ~2% of decode time. CUDA graphs were already enabled. The tilelang-based fused indexer kernel failed to JIT-compile on sm_120 with CUDA 13, confirming why the hook had disabled it in the first place. The MoE backend was locked to marlin (the only valid path for FP4-packed experts on sm_120). Each lever produced at best a few percent improvement, and the measured throughput hovered around 23–28 tok/s at concurrency 16—a far cry from the 1000 tok/s target.

The Reasoning Process: A Deep Dive

Message 12433 opens with the assistant examining the MoE method routing code in SGLang's fp8.py:

Looking at the MoE method routing for FP4 experts, I see that marlin uses the SM120 triton implementation, flashinfer_mxfp4 routes to either SM90 cutlass or SM100 trtllm, and any other backend falls back to the FP8 method while keeping the experts in FP4-packed format.

>

The core issue is that FP8MoEMethod would misinterpret the FP4-packed weights, causing either garbage output or crashes. This means for FP4 experts on SM120, marlin with SM120 triton is the only viable path—anything else breaks.

This is a critical piece of systems knowledge. The assistant is reading the actual source code to confirm a constraint, rather than relying on assumptions. The MoE dispatch logic in fp8.py (lines 283–305) routes FP4 experts through a specific path: if is_fp4_experts and the backend is marlin, it uses Mxfp4MarlinMoE; otherwise, it falls through to Fp8MoEMethod, which treats the weights as FP8 and produces garbage. The assistant correctly identifies that "anything else breaks" and that the user's suggestion to try "anything other than marlin" is not viable for this checkpoint format.

This moment of code-reading is emblematic of the assistant's methodology throughout the session. Rather than speculating about what might work, it traces the actual execution paths to establish hard constraints. The reasoning is grounded in the source code, not in guesswork.

The Architecture-Hardware Mismatch Diagnosis

The assistant then articulates the core diagnosis with remarkable clarity:

Given the constraints, I'm recognizing that the real bottleneck isn't the MoE backend but the underlying architecture: DeepSeek-V4-Flash relies on DSA with lightning indexer and mHC, which have fused kernels only for SM100, while SM120 runs these components in torch (latency-bound) alongside the MoE in triton. With decode at bs=1 taking 94ms and bs=16 steps around 700ms, and neither NCCL nor cuda-graph providing meaningful gains, hitting 1000 tps at C=16 would require fundamentally faster fused SM120 kernels across the entire DSA-indexer-mHC-MoE pipeline—essentially a multi-week custom kernel effort.

This is the moment of reckoning. The assistant quantifies the gap precisely: 94ms per decode step at batch size 1, ~700ms at batch size 16. To achieve 1000 tok/s at concurrency 16, each decode step would need to complete in ~16ms—a 43× speedup. The assistant correctly identifies that no configuration-level tuning can close this gap; it would require writing custom fused kernels for sm_120, analogous to the multi-week effort that delivered 3–6× speedups for the earlier Kimi K2.6 model.

But notice the intellectual honesty: the assistant does not simply declare defeat. It immediately questions its own diagnosis:

But I should double-check whether I'm missing something. The user's confidence in "much faster than 25 T/s" might stem from comparing against the FP8 checkpoint path, which avoids the FP4 marlin constraint, or from a configuration where DSA fused kernels actually work.

This self-doubt is a hallmark of rigorous engineering reasoning. The assistant considers alternative explanations: perhaps the user's expectation is based on a different checkpoint format (FP8 vs FP4), or perhaps there is a configuration where the DSA fused kernels do work on sm_120. It even questions whether the tilelang compile failure might be fixable with a different CUDA version or compiler flags.

The Pivot to Expert Parallelism

The most interesting part of the message is the internal debate about whether to try Expert Parallelism:

Before I settle on that conclusion, I should profile the actual bottleneck... Both the indexer/mHC (fixed overhead, dominates at bs=1) and the MoE (scales with batch, dominates at bs=16) are sm120 fallbacks, so I'm confident in the diagnosis now. Before giving the honest assessment, I should revert to the working baseline and confirm the ~23 tps number with NCCL. One more lever worth exploring is expert parallelism—at C=16, distributing the 256 experts across GPUs with all-to-all communication could speed up the MoE significantly at higher batch sizes...

The assistant is torn. On one hand, the diagnosis seems clear: the bottleneck is architectural, not configurable. On the other hand, Expert Parallelism (EP) is a genuine throughput lever that hasn't been tested yet. The K2.6 benchmarks showed EP outperforming TP on PCIe-connected GPUs (1500 tok/s vs 1291 tok/s). The assistant recognizes that giving up without trying EP would be premature:

I'm also second-guessing whether the bottleneck is truly unfixable or if EP could unlock a meaningful jump—let me verify that before deciding whether to keep iterating or report the honest assessment that reaching 1000 tps requires fused kernels, not just config-level tuning. I'm torn between trying EP4 as the last real config lever versus just reporting back honestly. The research clearly shows EP is the throughput multiplier for PCIe MoE, and I haven't actually tested it yet—that would be premature to give up without that attempt.

This internal conflict is the heart of the message. The assistant has enough evidence to conclude that the target is unreachable, but it refuses to stop iterating until every plausible lever has been tried. It chooses action over reporting, deciding to "revert the broken tilelang flag, try EP4 with whatever a2a backend is available on sm120, and see if it moves the needle on throughput."

The Execution: Reverting and Reconfiguring

The message concludes with concrete action: reverting the tilelang environment variable that caused the server crash, and creating a new launch script for EP4:

cat > /root/dsv4_nccl_env.sh <<"EOF"
export CUDA_HOME=/usr/local/cuda-13.0
export PATH=/usr/local/cuda-13.0/bin:$PATH
export SGLANG_DSV4_FP4_EXPERTS=1
export PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True
export NCCL_IB_DISABLE=1
export NCCL_P2P_LEVEL=5
export NCCL_MIN_NCHANNELS=8
export NCCL_PROTO=LL
export NCCL_ALGO=Ring
export NCCL_MAX_NCHANNELS=16
export NCCL_BUFFSIZE=16777216
export NCCL_NTHREADS=512
export OMP_NUM_THREADS=8
EOF

The EP4 script uses --tp 4 --ep-size 4, which configures 4-way tensor parallelism and 4-way expert parallelism across the 8 GPUs. This is a sensible configuration: TP4 handles the non-MoE layers (attention, embeddings) across 4 GPUs, while EP4 distributes the 256 experts across 4 groups of 2 GPUs each. The all-to-all communication between TP and EP groups happens via NCCL.

The assistant also kills any existing server processes and verifies GPU memory is freed before relaunching—a careful operational discipline that prevents subtle bugs from stale processes.

Assumptions and Potential Blind Spots

The assistant makes several assumptions in this message that deserve scrutiny:

Assumption 1: The MoE backend is locked to marlin. This is confirmed by reading the source code, but it assumes that the FP4-packed weight format is immutable. If the weights could be dequantized to FP8 or converted to a different format, other MoE backends would become available. However, dequantizing 256 experts from FP4 to FP8 would double memory usage and might not fit in VRAM, so this assumption is reasonable.

Assumption 2: Expert Parallelism will work on sm_120 with the flashinfer all-to-all backend. The assistant acknowledges this risk: "though there were issues with flashinfer graph capture and DeepEP on sm120 before." EP requires efficient all-to-all communication, and on sm_120 with PCIe, the communication backend may not support CUDA graph capture, which could negate any compute gains.

Assumption 3: The tilelang indexer failure is fundamental, not fixable. The assistant accepts that tilelang cannot JIT-compile on sm_120 with CUDA 13, but does not investigate whether a different CUDA version or compiler flags could resolve this. Given the time constraints and the multi-week effort required for custom kernels, this is a pragmatic assumption, but it leaves a potential avenue unexplored.

Assumption 4: The batch scaling behavior is understood. The assistant notes that bs=1 takes 94ms and bs=16 takes ~700ms, and attributes the difference to MoE compute scaling with batch size. However, this is an inference from aggregate measurements, not a detailed profile. Without a kernel-level breakdown (which would require disabling CUDA graphs and running a profiler), the assistant is reasoning from summary statistics, which could mask unexpected bottlenecks.

Input Knowledge Required

To fully understand this message, the reader needs knowledge of:

  1. DeepSeek-V4-Flash architecture: The model uses 256 experts with FP4 quantization, DSA attention with a lightning indexer that selects top-512 tokens, and manifold-constrained hyper-connections with 20 Sinkhorn iterations per layer.
  2. SGLang's MoE dispatch system: The fp8.py file contains routing logic that selects between marlin, flashinfer, and fallback MoE backends based on expert quantization format and hardware architecture.
  3. Blackwell GPU variants: SM100 (B200/B300 datacenter GPUs) has full fused kernel support for DSA operations. SM120 (RTX PRO 6000 workstation GPUs) lacks these fused kernels and falls back to torch/Triton implementations.
  4. Expert Parallelism vs Tensor Parallelism: TP splits each operation across GPUs (requiring communication every layer), while EP distributes experts across GPUs and uses all-to-all communication only at MoE boundaries. EP is generally more efficient for autoregressive decoding on PCIe-connected GPUs.
  5. NCCL tuning parameters: The environment variables (NCCL_PROTO=LL, NCCL_ALGO=Ring, NCCL_MIN_NCHANNELS=8, etc.) control how NVIDIA's collective communication library handles GPU-to-GPU transfers.

Output Knowledge Created

This message produces several valuable pieces of knowledge:

  1. Confirmed constraint: FP4 experts on sm_120 can only use the marlin MoE backend. Any other backend would misinterpret the FP4-packed weights and produce incorrect results.
  2. Diagnosis of the throughput ceiling: The 23–28 tok/s ceiling is caused by sm_120 fallback kernels for the DSA indexer, mHC, and MoE operations, which run on CUDA cores (SIMT) rather than tensor cores. This is an architectural constraint, not a configuration problem.
  3. Quantified gap: To reach 1000 tok/s at C=16, each decode step would need to complete in ~16ms, requiring a ~43× speedup over the current ~700ms step time. This cannot be achieved through configuration tuning alone.
  4. EP4 configuration: A concrete launch script for Expert Parallelism with TP4+EP4 on 8 GPUs, including NCCL tuning parameters optimized for PCIe connectivity.
  5. Methodological precedent: The assistant demonstrates a rigorous approach to optimization—exhaust all configurable levers before concluding that a bottleneck is architectural, but also know when to stop tuning and accept the hardware constraint.

The Thinking Process: A Window into Engineering Judgment

The most valuable aspect of this message is the visible thinking process. The assistant does not present a polished conclusion; instead, it shows the messy, iterative process of engineering judgment:

  1. Code reading: The assistant reads fp8.py to confirm the MoE backend constraint, rather than relying on memory or assumptions.
  2. Hypothesis testing: It considers multiple explanations for the throughput gap—MoE backend, indexer latency, mHC overhead, batch scaling—and weighs their plausibility.
  3. Self-correction: The assistant explicitly questions its own diagnosis ("But I should double-check whether I'm missing something") and considers alternative explanations.
  4. Cost-benefit analysis: It weighs the cost of trying EP4 (another server restart, potential failure) against the benefit (possibly unlocking a meaningful throughput gain).
  5. Decision under uncertainty: Despite being "confident in the diagnosis," the assistant chooses to try EP4 because "that would be premature to give up without that attempt." This is the essence of good engineering: exhaust all options before conceding.

Conclusion

Message 12433 is a turning point in the optimization campaign. The assistant has systematically eliminated every configurable lever—NCCL tuning, CUDA graphs, tilelang fusion, MoE backends—and arrived at a clear diagnosis: the throughput ceiling is architectural, not configurable. But rather than report this conclusion immediately, it pauses, re-examines its reasoning, identifies one remaining untested lever (Expert Parallelism), and pivots to try it.

This message exemplifies the discipline of rigorous optimization: ground every conclusion in source code, quantify every gap, question every assumption, and exhaust every plausible lever before declaring a bottleneck fundamental. Whether EP4 succeeds or fails, the assistant will have earned the right to deliver its honest assessment, backed by evidence from every angle of attack.

The article that follows this message—whether it reports a breakthrough from EP or confirms the architectural ceiling—will be built on the foundation laid in this moment of careful reasoning and decisive action.