Chunk 67.0

In this chunk, the assistant successfully deployed **DeepSeek-V4-Flash (FP4/nvfp4)** on SGLang main across 8× RTX PRO 6000 (sm_120). After stopping the previous K2.6 service, freeing disk, and downloading the 146 GB checkpoint, the assistant built a fresh SGLang environment with all required dependencies (flashinfer 0.6.12, sglang-kernel 0.4.3, tilelang, NIXL, etc.). The model loaded and generated correctly on a single-node TP4 configuration, and the headline deliverable—**prefill-decode disaggregation**—was achieved: prefill TP4 on GPU0-3/NUMA0, decode TP4 on GPU4-7/NUMA1, with KV transfer via NIXL/UCX and a router on port 8000, all verified end-to-end. However, performance fell far short of expectations: only ~10 tok/s at bs=1 and ~25 tok/s at C=16, versus the user's target of ~1000 tok/s. The assistant systematically exhausted every config lever—NCCL LL+Ring tuning, CUDA graphs (already on), tilelang indexer fusion (JIT-compile failure on sm_120), non-marlin MoE backends (invalid for FP4 experts), and expert parallelism (worse due to PCIe all-to-all overhead)—none moved the needle. A definitive GPU profile traced 63% of decode time to a single kernel: `_tiled_sparse_decode_kernel`, the sm_120 Triton fallback for sparse MLA attention, which launches only 64 blocks (1 batch × 64 heads) on ~170 SMs, serially iterating all 512 top-k tokens. This is the same low-occupancy pathology that plagued the earlier K2.6 verify kernel. The core theme is the **hard ceiling of sm_120 fallback kernels**: the fast fused DSA/MoE stack (DeepGEMM, trtllm-gen, FP4 indexer) is arch-gated to SM100, and no amount of config tuning or even a single-kernel rewrite (which could yield ~2-3× overall) can close the ~40× gap to 1000 tok/s. The assistant documented the precise bottleneck, quantified the ceiling, and concluded that reaching the user's throughput target on this hardware would require a multi-week custom kernel effort akin to the earlier K2.6 work. The PD disaggregation orchestration itself works correctly, but decode remains the bottleneck regardless of resource isolation.

The Blackwell Ceiling: Deploying DeepSeek-V4-Flash with Prefill-Decode Disaggregation and Confronting the sm_120 Fallback Kernel Bottleneck 2900 words

Message Articles

Subagent Sessions