The Last Lever: Expert Parallelism and the Hard Ceiling of sm_120 Fallback Kernels

A Single Bash Command That Encapsulates an Entire Optimization Campaign

In the sprawling, multi-day effort to deploy DeepSeek-V4-Flash on 8× NVIDIA RTX PRO 6000 Blackwell GPUs (sm_120), there is a moment captured in message [msg 12435] that reads, on its surface, like a mundane operational step: the assistant launches a server process with an Expert Parallelism (EP4) configuration. The command is a single nohup bash invocation over SSH, and the output is a process ID and a log line about setting a model configuration parameter. Yet this message is the culmination of a long, methodical, and increasingly frustrated optimization campaign—a campaign that has systematically exhausted every configuration lever available and is now down to its last roll of the dice. Understanding why this particular command was issued, and what it represents, requires tracing the reasoning, assumptions, and technical constraints that led to this point.

The Optimization Campaign: A Trail of Exhausted Levers

The assistant had been tasked with deploying DeepSeek-V4-Flash—a Mixture-of-Experts (MoE) model with 256 experts, using FP4 (4-bit floating-point) quantization—on Blackwell sm_120 hardware. The target was ambitious: approximately 1,000 tokens per second at a concurrency of 16 (C=16). What the assistant discovered, however, was a fundamental architectural mismatch.

The DeepSeek-V4-Flash architecture relies on a DSA (Dynamic Sparse Attention) pipeline with a "lightning indexer" and "mHC" (multi-head coalesced) operations. These components have fused, highly optimized CUDA kernels—but only for NVIDIA's SM100 architecture (the previous generation). On sm_120 (Blackwell), SGLang falls back to torch and Triton implementations that are inherently latency-bound. The assistant's profiling revealed that at batch size 1, each decode step took ~94 milliseconds (roughly 10 tok/s), and at C=16, the aggregate throughput was only ~23–25 tok/s—a step time of ~700 milliseconds. To reach 1,000 tok/s, that step time would need to drop to 16 milliseconds—a 43× speedup.

The assistant tried everything. NCCL LL+Ring tuning had negligible effect because communication was only ~2% of decode time. CUDA graphs were already enabled and provided no further headroom. The tilelang indexer fusion, which promised to accelerate the attention indexer, failed to JIT-compile on sm_120 with CUDA 13. The user had suggested trying MoE backends "other than marlin," but the assistant read the source code in fp8.py ([msg 12433]) and confirmed that for FP4 experts, the marlin backend with its sm_120 Triton implementation was the only valid path—any other backend would fall through to Fp8MoEMethod, which would misinterpret the FP4-packed weights and produce garbage or crash.

The Reasoning Behind the EP4 Decision

By the time the assistant reached message [msg 12433], the internal reasoning shows a clear tension. The assistant had already concluded, based on profiling data and architectural analysis, that the bottleneck was structural: the sm_120 fallback kernels for sparse-MLA decode and MXFP4 MoE were fundamentally inefficient, running on CUDA cores (SIMT) rather than the sm_120a FP4/FP8 tensor cores. The NVFP4 checkpoint had fixed the MoE path (routing it through tensor cores), but attention remained the primary limiter. The assistant documented the precise bottleneck, quantified the ceiling, and concluded that reaching the throughput target on this hardware would require a multi-week custom kernel effort—the same playbook that had delivered 3–6× gains in an earlier K2.6 project.

Yet the assistant hesitated to deliver that conclusion. The reasoning block in [msg 12433] reveals a back-and-forth internal debate:

"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 is the voice of a rigorous engineer who wants to be certain before declaring defeat. Expert Parallelism (EP) distributes the 256 MoE experts across multiple GPUs, with all-to-all communication routing tokens to the GPU that holds the relevant experts. On PCIe-connected GPUs, EP had previously outperformed Tensor Parallelism (TP) for autoregressive decoding—the K2.6 deployment had hit ~1,500 tok/s with EP versus 1,291 with TP. The assistant reasoned that at C=16, distributing the expert computation could significantly speed up the MoE-dominated decode path.

But there were risks. The flashinfer all-to-all backend had shown issues with CUDA graph capture on sm_120. DeepEP (another communication backend) had also failed previously. And EP4 meant dedicating 4 GPUs to expert computation, which required careful memory management. The assistant was essentially taking one last shot at a config-level solution before conceding that only custom kernels could close the gap.

The Execution: Launching EP4

Message [msg 12434] shows the immediate prelude: the assistant had attempted to kill the old server and write the EP4 script in a single SSH command, but the heredoc containing the port number 30000 caused the pkill regex pattern to self-match—a classic shell scripting gotcha where the pattern launch_server.*3000[0] matched the command line itself. The assistant recognized the issue, separated the script-writing from the kill command, and verified that all 8 GPUs showed only 4 MiB of memory used—clean and ready.

Then, in [msg 12435], the assistant launches:

timeout 25 ssh -o StrictHostKeyChecking=no root@10.1.230.171 'cd /root && nohup bash /root/serve_dsv4_ep4.sh > /root/dsv4_ep4.log 2>&1 & echo "PID $!"; sleep 6; grep -iE "ep.size|expert.parall|error" /root/dsv4_ep4.log | tail -3; tail -1 /root/dsv4_ep4.log'

The command is carefully constructed. It uses nohup to detach the server from the SSH session, redirects both stdout and stderr to a log file, captures the process ID, sleeps 6 seconds to let initialization begin, then greps for signs of EP configuration or errors. The output is promising: PID 80503 and a log line about setting swa_full_tokens_ratio to 0.1 for DeepseekV4ForCausalLM. The server is alive and starting to load the model.

What This Message Reveals About the Assistant's Thinking

The subject message is deceptively simple, but it encodes several layers of decision-making:

Assumptions. The assistant assumes that EP4 is a viable configuration on sm_120—that the all-to-all communication backend will work, that CUDA graphs can be captured with the EP topology, and that the memory fraction of 0.70 is sufficient for the model with EP's additional communication buffers. These assumptions are informed by the K2.6 precedent, but they are not yet validated for DeepSeek-V4-Flash on this hardware.

Knowledge required. To understand this message, one needs to know: what Expert Parallelism is and how it differs from Tensor Parallelism; that DeepSeek-V4-Flash has 256 MoE experts; that PCIe bandwidth is a bottleneck for TP but EP can be more communication-efficient; that the sm_120 architecture lacks fused DSA kernels; and that the assistant has already exhausted NCCL tuning, CUDA graphs, tilelang fusion, and alternative MoE backends.

Knowledge created. The message itself creates minimal new knowledge—it's an execution step. But the fact that this step was taken creates knowledge about the optimization process: it tells us that EP was the last untested config lever, and that the assistant was methodically ruling out possibilities before concluding that custom kernels were necessary.

Mistakes and incorrect assumptions. The assistant's earlier attempt to combine script-writing with pkill in a single SSH command ([msg 12433]) failed due to self-matching—a mistake that was recognized and corrected in [msg 12434]. More broadly, the assumption that any config-level change could bridge the ~40× throughput gap was probably incorrect from the start, but the assistant needed to verify this empirically rather than assert it.

The Broader Significance

This message sits at a critical inflection point in the optimization campaign. It represents the boundary between configuration tuning and fundamental kernel engineering. The assistant has tried every knob available in the SGLang framework—NCCL protocols, CUDA graphs, indexer fusion, MoE backends, and now parallelism strategy. If EP4 fails to move the needle (and the assistant's own analysis suggests it won't close a 40× gap), the only remaining path is to write custom sm_120 fused kernels: a multi-week effort that the assistant has already scoped based on the earlier K2.6 work.

The message is also a testament to the value of systematic, measurement-driven optimization. Rather than guessing at the bottleneck, the assistant profiled the GPU, identified _tiled_sparse_decode_kernel as consuming 63% of decode time, traced the MoE slot-GEMV and sparse-decode attention kernels as running on CUDA cores rather than tensor cores, and documented the precise architectural ceiling. Every config change was benchmarked individually. The EP4 launch is the final entry in a ledger of experiments, each one ruled out with evidence.

In the end, [msg 12435] is a message about intellectual honesty and engineering rigor. The assistant could have declared defeat earlier based on the profiling data alone. Instead, it chose to exhaust every option—including the risky, unproven EP4 path—so that when it finally reported the bottleneck to the user, it could do so with complete confidence that no config-level solution existed. The server started, the model began loading, and the assistant was one step closer to the truth, whatever it might be.