The Final Lever: Expert Parallelism and the Verification of a Working Hypothesis

Introduction

In the methodical, often grueling work of deploying large language models on non-reference hardware, progress rarely comes in leaps. It comes in the accumulation of small, measured steps—each one ruling out a dead end, confirming a constraint, or narrowing the search space toward a viable configuration. Message 12436 in this opencode session captures one such step: the moment when an assistant, having systematically exhausted nearly every configurable lever in the SGLang inference stack, launches an Expert Parallelism (EP4) deployment of DeepSeek-V4-Flash and waits to see if it will even start.

The message is deceptively simple—a bash polling loop that checks a log file every 25 seconds until the server announces readiness. But beneath this mundane operation lies the culmination of an extended optimization campaign, a carefully reasoned decision tree, and the last remaining hope for a meaningful throughput improvement on a hardware platform that has proven stubbornly resistant to conventional tuning.

The Message

The assistant issues a single bash command:

for i in $(seq 1 12); do
  sleep 25
  st=$(timeout 12 ssh -o StrictHostKeyChecking=no root@10.1.230.171 'tail -1 /root/dsv4_ep4.log' 2>/dev/null)
  echo "[$((i*25))s] ${st: -130}"
  echo "$st" | grep -qiE "fired up|ready to roll" && { echo ">>> READY"; break; }
  echo "$st" | grep -qiE "out of memory|Traceback|CUDA error|Killed|kill_process_tree" && { echo ">>> DIED"; timeout 12 ssh root@10.1.230.171 'grep -iE "error|exception|out of memory|not support|assert|a2a|expert" /root/dsv4_ep4.log | grep -viE "INFO|disabled because"|tail -5'; break; }
done

The output confirms success:

[25s] e cache initialized: source=default impl=SWARadixCache hybrid_swa=True hybrid_ssm=False hierarchical=False streaming_wrapped=False
[50s] [2026-06-17 16:58:34] The server is fired up and ready to roll!
>>> READY

The EP4 server for DeepSeek-V4-Flash has started successfully on the 8× RTX PRO 6000 Blackwell (sm_120) machine. The model loaded, the KV cache initialized with SWA (Sliding Window Attention) radix caching, and the server entered its ready state after approximately 50 seconds. This is a working deployment—but whether it will deliver the throughput gains the assistant is searching for remains an open question.

The Road to This Point: Why This Message Was Written

To understand why this particular message matters, one must trace the optimization journey that preceded it. The assistant had been tasked with deploying DeepSeek-V4-Flash—a massive Mixture-of-Experts (MoE) model with 256 experts, using the NVFP4 (4-bit floating point) quantization format—on NVIDIA RTX PRO 6000 Blackwell GPUs (compute capability sm_120). The target was ambitious: approximately 1000 tokens per second at a concurrency of 16 requests.

The initial deployment achieved roughly 10 tok/s at batch size 1 and 25 tok/s at C=16—a staggering 40× short of the target. The assistant then embarked on a systematic optimization campaign, testing every available configuration lever:

  1. NCCL tuning (LL protocol, Ring algorithm, channel counts): negligible effect, confirming that inter-GPU communication was not the bottleneck (only ~2% of decode time).
  2. CUDA graphs: already enabled by default, providing no additional headroom.
  3. Tilelang indexer fusion: failed at JIT compilation time on sm_120 with CUDA 13, producing an unhandled exception that crashed the server.
  4. Alternative MoE backends: the assistant confirmed through code analysis that for FP4 experts on sm_120, the marlin backend (dispatching to sm_120 Triton kernels) is the only valid path. Any other backend would misinterpret the FP4-packed weight format, producing garbage output or crashes.
  5. Non-marlin MoE runners: ruled out by the FP4 weight format constraint. Each lever either failed outright, produced no measurable improvement, or was structurally incompatible with the hardware or model format. The assistant's reasoning in the preceding messages reveals a growing awareness of the fundamental problem: DeepSeek-V4-Flash's DSA (Dynamic Speculative Architecture) relies on fused kernels (lightning indexer, mHC operations) that are arch-gated to NVIDIA's SM100 compute capability. On sm_120 (Blackwell), SGLang falls back to latency-bound torch and Triton implementations that cannot approach the throughput of native fused kernels. With the obvious levers exhausted, the assistant turned to the one remaining high-value option identified in prior research: Expert Parallelism (EP). During earlier work with the Kimi K2.6 model on the same hardware, EP had demonstrated a clear throughput advantage over Tensor Parallelism (TP) for autoregressive MoE inference on PCIe-connected GPUs—approximately 1500 tok/s for EP versus 1291 tok/s for TP. The reasoning was sound: EP distributes the 256 experts across GPUs, reducing the per-GPU MoE compute load, at the cost of all-to-all communication. On PCIe, where the all-to-all overhead is manageable at moderate batch sizes, this tradeoff can be net positive. Message 12436 is the verification step for this hypothesis. Before the assistant can benchmark EP4 throughput, the server must first start successfully—and this message confirms that it does.

Decisions Made in This Message

While the message itself is a straightforward polling operation, several implicit decisions shaped its design:

The polling strategy: The assistant chose a 25-second polling interval with a maximum of 12 iterations (5 minutes total). This reflects an expectation that the server initialization—model loading, weight quantization, KV cache allocation, and CUDA graph warmup—would complete within a few minutes. The 25-second granularity avoids excessive SSH overhead while providing reasonably prompt detection of both success and failure.

The readiness detection pattern: The assistant checks for the phrases "fired up" or "ready to roll" in the log output. These are SGLang's standard server-ready messages. The choice of two patterns provides robustness against minor logging variations across SGLang versions.

The failure detection pattern: The assistant monitors for "out of memory", "Traceback", "CUDA error", "Killed", and "kill_process_tree". This covers the most likely failure modes: OOM (given the 146 GB model on 8 GPUs with 0.7 memory fraction), Python exceptions, CUDA runtime errors, OS-level process termination, and SGLang's own process tree cleanup (which would indicate a crash during initialization).

The error detail extraction: On failure, the assistant fetches additional log context filtered for error-related patterns, excluding INFO-level messages and known harmless warnings. This demonstrates a learned awareness of SGLang's logging verbosity—the logs contain many informational messages that could obscure real errors.

The decision to proceed with EP4 despite uncertainty: The assistant had not yet confirmed that EP4 would improve throughput. The earlier K2.6 results suggested EP could be beneficial, but DeepSeek-V4-Flash has a different architecture (DSA, different expert count, FP4 quantization). The assistant chose to invest the iteration cost of testing EP4 rather than reporting the pessimistic conclusion that the throughput target was unreachable. This reflects a methodological commitment to empirical verification over theoretical reasoning.

Assumptions Embedded in This Message

Several assumptions underpin the assistant's approach:

EP4 is compatible with FP4 experts on sm_120: The assistant verified through code analysis that marlin is the only valid MoE backend for FP4 experts. EP4 with marlin should work because marlin handles the per-GPU expert computation, while the all-to-all communication is handled by SGLang's EP infrastructure. This assumption proved correct—the server started successfully.

The server will not OOM during initialization: The model is 146 GB in FP4 format. With 8 GPUs at 0.7 memory fraction (~33 GB usable per GPU on a 48 GB card), TP4 would require ~36.5 GB per GPU for the full model. EP4 distributes only the experts across GPUs, potentially reducing per-GPU memory. The assistant assumed this would fit, and it did.

CUDA graphs remain compatible with EP: The EP4 launch script does not explicitly disable CUDA graphs (--cuda-graph-max-bs 32 is present, which enables them). The assistant assumed that CUDA graph capture would succeed with EP, which is not guaranteed—EP introduces dynamic all-to-all communication that can complicate graph capture.

The polling approach is sufficient: The assistant checks only the last line of the log file. This assumes that the server-ready message appears as the final log entry, which is true for clean startups. However, if the server crashed after printing "fired up" (e.g., during the first inference request), this would not be detected by the polling loop. The assistant accepts this limitation, treating server startup as the primary validation gate.

Mistakes and Incorrect Assumptions

The broader context reveals several incorrect assumptions that led to this point:

The tilelang indexer would work on sm_120: The assistant attempted to enable the tilelang-based fused indexer, which promised significant speedups for the DSA attention path. It failed at JIT compilation time on sm_120 with CUDA 13, producing an unhandled exception. This was a reasonable assumption—tilelang is designed to be architecture-agnostic—but the reality is that its kernel generation pipeline has not been validated on Blackwell hardware.

NCCL tuning would improve throughput: The assistant invested significant effort in NCCL configuration (LL protocol, Ring algorithm, channel counts, buffer sizes). The result was negligible improvement because inter-GPU communication accounts for only ~2% of decode time. The bottleneck is compute-bound (fallback kernels on sm_120), not communication-bound.

The throughput target was achievable through configuration tuning: The assistant's systematic exploration of every config lever reflects an implicit assumption that the 1000 tok/s target was reachable through software configuration alone. The evidence increasingly suggests that reaching this target requires custom sm_120 kernels—a multi-week engineering effort analogous to the earlier K2.6 custom kernel work.

The pkill command would not self-match: In the previous message, the assistant's heredoc containing the port number caused the pkill pattern to match its own shell process. This is a subtle but instructive mistake—the assistant learned to separate script writing from process management in subsequent commands.

Input Knowledge Required

To fully understand this message, one needs:

Knowledge of SGLang's architecture: The server launch command uses flags like --tp 4 --ep-size 4 --moe-runner-backend marlin --cuda-graph-max-bs 32 --num-continuous-decode-steps 4. Understanding these requires familiarity with SGLang's parallelism strategies (tensor parallelism, expert parallelism), MoE backend dispatch, and CUDA graph optimization.

Knowledge of DeepSeek-V4-Flash: The model uses DSA with 256 experts, FP4 quantization, and SWA (Sliding Window Attention). The "SWARadixCache" initialization message confirms the SWA radix caching is active. The model's architecture determines which optimization levers are available and which are structurally constrained.

Knowledge of Blackwell (sm_120) hardware: The RTX PRO 6000 Blackwell GPUs have compute capability sm_120, which differs from the Hopper (sm_90) and future Blackwell Ultra (sm_100) architectures. Many fused kernels (DeepGEMM, trtllm-gen, FP4 indexer) are arch-gated to sm_100, leaving sm_120 with fallback implementations.

Knowledge of PCIe vs NVLink topology: The machine has 8 GPUs connected via PCIe, not NVLink. This makes all-to-all communication for EP more expensive than on NVLink-connected systems, but the assistant's prior research suggested EP still wins for MoE-heavy workloads on PCIe.

Knowledge of the earlier K2.6 work: The assistant references prior results showing EP beating TP for K2.6 on the same hardware (~1500 vs 1291 tok/s). This precedent motivates the EP4 experiment.

Output Knowledge Created

This message produces several concrete outputs:

EP4 server startup confirmed: The DeepSeek-V4-Flash model loads successfully with EP4 configuration on 8× RTX PRO 6000 GPUs. This validates the compatibility of EP with FP4 experts and marlin backend on sm_120.

Startup time measured: The server reaches ready state in approximately 50 seconds. This is consistent with model loading and initialization for a 146 GB checkpoint across 8 GPUs.

SWA radix cache confirmed: The log shows "SWARadixCache" initialization, confirming that the sliding window attention radix caching is active. This is relevant for throughput because SWA reduces the effective KV cache size for long contexts.

A working baseline established: The EP4 server is now available for benchmarking. The assistant can proceed to measure throughput at various batch sizes and concurrency levels, comparing against the TP4 baseline of ~25 tok/s at C=16.

A methodological precedent: The message demonstrates a pattern for server deployment verification that can be reused: poll the log with timeout, check for readiness signals, detect failures with specific error patterns, and extract relevant diagnostic information on failure. This pattern is applicable across SGLang deployments.

The Thinking Process Visible in Reasoning

The assistant's reasoning, visible in the preceding messages, reveals a methodical and self-aware optimization process:

Hypothesis-driven exploration: Each config lever is tested with a clear hypothesis about what it should improve. NCCL tuning targets communication overhead; tilelang targets indexer latency; EP targets MoE compute distribution. When a lever fails, the assistant understands why it failed and incorporates that knowledge into the next decision.

Cost-aware iteration: The assistant is acutely aware of the iteration cost. Each server restart takes 50+ seconds, and each benchmark run adds more time. The decision to try EP4 reflects a cost-benefit calculation: EP4 is the last high-value lever that can be tested without custom kernel development, and the iteration cost is justified by the potential upside.

Architectural reasoning: The assistant reasons from first principles about the model architecture and hardware constraints. The FP4 weight format locks the MoE backend to marlin. The sm_120 compute capability gates the fused kernels. The PCIe topology makes all-to-all communication expensive but potentially worthwhile. This architectural understanding prevents wasted effort on impossible configurations.

Learning from mistakes: The pkill self-matching bug in the previous message is handled cleanly in this one. The assistant separates script writing from process management, uses standalone kill commands, and verifies GPU memory is free before launching. This demonstrates adaptive learning within the session.

Honesty about limitations: While the assistant does not explicitly state it in this message, the preceding reasoning shows a growing recognition that the 1000 tok/s target may be unreachable through configuration tuning alone. The EP4 experiment is the last test before that conclusion becomes unavoidable. This intellectual honesty—willingness to confront negative results—is a hallmark of rigorous engineering.

Broader Significance

Message 12436 captures a universal pattern in systems engineering: the moment when a practitioner, having exhausted the obvious solutions, tests the last remaining hypothesis before confronting a fundamental limitation. The EP4 server starts successfully, but whether it delivers the hoped-for throughput gains is a question that will be answered in subsequent messages.

The deeper lesson is about the nature of performance optimization in the AI inference stack. The bottlenecks are increasingly architectural and hardware-specific, not configurable. The sm_120 fallback kernel problem—where optimized fused kernels are gated to newer hardware generations—is a structural constraint that no amount of configuration tuning can overcome. The path to higher throughput requires either hardware that supports the fused kernels (sm_100), or custom kernel development targeting the specific hardware (as was done for K2.6).

This message, then, is a pivot point. The EP4 server is running. The assistant will benchmark it, and the results will determine whether the optimization campaign continues or whether the conclusion—that the throughput target requires custom kernel development—becomes unavoidable. The polling loop ticks, the server reports readiness, and the next chapter begins.