The Last Lever: Benchmarking Expert Parallelism on DeepSeek-V4-Flash

In the relentless pursuit of inference throughput, there comes a moment when every obvious knob has been turned, every configuration flag flipped, and the only remaining question is whether the final untested lever can deliver the breakthrough. Message [msg 12437] captures precisely such a moment in an extended optimization campaign for DeepSeek-V4-Flash running on NVIDIA RTX PRO 6000 Blackwell GPUs (sm_120 architecture). After exhausting NCCL tuning, CUDA graphs, tilelang kernel fusion, MoE backend selection, and even checkpoint format changes (NVFP4 quantization), the assistant turns to Expert Parallelism (EP) as the last major configuration lever before confronting the hard architectural ceiling beneath it all.

The Message

The message is deceptively simple—a single bash command dispatched to a remote server running SGLang:

`` [assistant] ## Agent Reasoning The EP4 server is running, so I'll do a quick decode rate check at C=1 first, then run the main benchmark at C=16 to see if EP improves throughput—I'll execute both through the bench script with appropriate sample sizes and let them run in the background. EP4 is up. Benchmark it (C=1 and C=16) vs the TP4 baseline. [bash] timeout 25 ssh -o StrictHostKeyChecking=no root@10.1.230.171 ' cat > /root/run_ep_bench.sh <<"XEOF" #!/bin/bash M=/root/models/DeepSeek-V4-Flash for spec in "1:8" "16:48"; do C=${spec%%:*}; N=${spec##*:} echo "===== EP4 concurrency=$C (n=$N) =====" /root/venv_sglang211/bin/python -m sglang.bench_serving \ --backend sglang --host 127.0.0.1 --port 30000 --model $M \ --dataset-name random --random-input-len 256 --random-output-len 256 \ --num-prompts $N --max-concurrency $C 2>&1 | \ grep -iE "Successful requests|Output token throughput|Total token throughput|Mean TPOT|Median TPOT|Median TTFT" done echo ALL_DONE XEOF cd /root && nohup bash /root/run_ep_bench.sh > /root/dsv4_bench/ep4.log 2>&1 & echo "bench PID $!"' bench PID 81421 ``

Beneath this straightforward benchmarking invocation lies a rich tapestry of prior investigation, architectural constraint, and hard-won systems knowledge. To understand why this particular command matters, one must trace the optimization path that led to it.

The Optimization Trajectory

The assistant had spent the preceding rounds methodically working through every available performance lever. The DeepSeek-V4-Flash model, with its Mixture-of-Experts architecture featuring 256 experts and FP4-quantized weights, presents unique challenges on Blackwell GPUs. The model's DSA (Dynamic Sparse Attention) design relies on fused kernels—lightning indexer, mHC (multi-head caching), and specialized MoE dispatch—that are optimized for NVIDIA's SM100 architecture (the Blackwell compute capability targeted by NVIDIA's own kernels). On sm_120 (the RTX PRO 6000 Blackwell variant), these fused kernels are unavailable, forcing SGLang to fall back to generic PyTorch and Triton implementations.

The assistant's earlier attempts told a story of incremental gains against a structural ceiling. NCCL tuning (LL protocol, Ring algorithm, increased channel counts) produced negligible improvement because communication accounts for only ~2% of decode time on this configuration. CUDA graphs were already enabled and provided no additional headroom. The tilelang indexer, which promised fused kernel execution for the sparse attention indexer, failed to JIT-compile on sm_120 with CUDA 13, crashing during initialization. The MoE backend was locked to Marlin with sm_120 Triton fallback—the only path that correctly handles FP4-packed weights; any alternative would misinterpret the quantization format and produce garbage outputs.

A definitive GPU profile had revealed the bottleneck's anatomy: 63% of decode time was consumed by a single kernel, _tiled_sparse_decode_kernel, the sm_120 Triton fallback for sparse MLA attention. This kernel launches only 64 blocks on ~170 SMs, serially iterating over all 512 top-k tokens—a classic low-occupancy pathology identical to what had plagued an earlier K2.6 deployment. The NVFP4 checkpoint switch had successfully moved MoE execution onto tensor cores (Marlin W4A16 path), yielding a ~24% throughput improvement to ~28 tok/s at concurrency 16, but attention remained the dominant bottleneck.

Why Expert Parallelism?

Expert Parallelism distributes the 256 experts across multiple GPUs, with each GPU responsible for computing a subset of expert feed-forward networks. For MoE models on PCIe-interconnected GPUs, EP can outperform Tensor Parallelism (TP) because it reduces the per-GPU expert computation while the all-to-all communication overhead is amortized across the batch. The assistant's research had shown this clearly: on the earlier K2.6 deployment, EP4 achieved ~1500 tok/s versus TP4's ~1291 tok/s—a meaningful advantage.

The decision to benchmark at both concurrency 1 (C=1) and concurrency 16 (C=16) reflects a nuanced understanding of where EP's benefits manifest. At C=1, single-request decode latency is dominated by the per-step overhead of the attention kernel and MoE dispatch—EP may help or hurt depending on whether the all-to-all communication latency exceeds the per-GPU compute savings. At C=16, the batch size is large enough that the MoE computation dominates, and EP's distribution of expert computation across GPUs should provide its greatest benefit. The assistant's choice of 8 prompts for C=1 (enough for statistical significance) and 48 prompts for C=16 (three full batches at the configured max-concurrency) follows standard benchmarking methodology.

Assumptions and Implicit Knowledge

The message rests on several assumptions worth examining. First, that the EP4 server is genuinely serving requests correctly—the assistant verified it started ("The server is fired up and ready to roll!") but did not run a quick sanity generation before launching the full benchmark. Given the history of crashes (the tilelang indexer failure, earlier OOM issues), this is a non-trivial assumption. Second, that the benchmark parameters (256-token input and output lengths, random dataset) are representative of the target workload. The assistant is benchmarking against a "random" dataset rather than a realistic prompt distribution, which may not reflect production serving patterns. Third, that the background execution will complete without issues—the nohup approach means the assistant cannot observe intermediate failures or progress.

The input knowledge required to understand this message is substantial. One must grasp the distinction between Tensor Parallelism (splitting individual operations across GPUs) and Expert Parallelism (assigning different experts to different GPUs), the role of all-to-all communication in EP, the peculiarities of FP4 quantization and its constraints on MoE backend selection, and the sm_120 vs SM100 architectural divergence that forces fallback kernel paths. The SGLang benchmarking tool's parameters (--num-prompts, --max-concurrency, --random-input-len, --random-output-len) and their impact on measurement validity are also essential context.

The Knowledge Being Produced

This message creates critical output knowledge: the EP4 throughput figures that will determine whether Expert Parallelism is a viable optimization for this hardware-software stack. If EP4 delivers a significant improvement over the TP4 baseline (~28 tok/s at C=16), it would validate the research hypothesis and open a new optimization path. If it matches or underperforms TP4, it would confirm that the sm_120 fallback kernel bottleneck is the dominant constraint, regardless of parallelism strategy.

The benchmark also produces secondary knowledge: the breakdown of throughput, TPOT (time per output token), and TTFT (time to first token) metrics that reveal where time is spent. The Median TTFT at C=16 would indicate whether the prefill stage scales acceptably under EP. The Mean and Median TPOT would show whether EP introduces latency variability due to all-to-all communication jitter.

The Broader Significance

This message represents the inflection point in a larger narrative. The assistant has systematically tested every configuration-level optimization available within the SGLang framework. Expert Parallelism is the last such lever before the fundamental constraint—the sm_120 fallback kernel bottleneck—must be confronted directly. The benchmark results will determine whether the assistant continues down the configuration-tuning path or pivots to the much heavier investment of writing custom CUDA kernels, as was done successfully for the earlier K2.6 deployment where custom sm_120 verify attention kernels delivered 3–6× speedups.

The message's structure also reveals the assistant's operational style: methodical, measurement-driven, and conservative about iteration cost. Rather than speculating about EP's impact, the assistant deploys it, benchmarks it, and waits for data. The background execution with nohup is a practical concession to the long-running nature of LLM serving benchmarks—a single C=16 run with 48 prompts at 256 output tokens each can take several minutes.

In the end, this message is about the moment before knowledge is gained. The assistant has set the experiment in motion, and the results—whether triumph or disappointment—will determine the next phase of the optimization campaign. It is a testament to the disciplined, empirical approach that characterizes effective systems engineering: formulate a hypothesis, design a test, execute cleanly, and let the data speak.