The Blackwell Frontier: Deploying DeepSeek-V4-Flash on RTX PRO 6000 GPUs

Introduction

The deployment of cutting-edge large language models on novel GPU architectures is one of the most technically demanding challenges in modern AI infrastructure. This article synthesizes a multi-phase engineering effort to bring up DeepSeek-V4-Flash—a state-of-the-art Mixture-of-Experts (MoE) model with FP4 quantization and sparse attention—on NVIDIA's RTX PRO 6000 Blackwell GPUs (compute capability sm_120). The work spans environment setup, codebase archaeology, deployment orchestration, and systematic performance optimization, ultimately revealing a hard ceiling imposed by architecture-specific kernel availability.

The journey documented in this session is a case study in the realities of hardware bring-up: the gap between what a GPU architecture can do and what the software ecosystem supports is often the decisive factor in deployment success. The RTX PRO 6000 Blackwell, with its sm_120 compute capability, represents a distinct branch of NVIDIA's Blackwell family—separate from the sm_100 datacenter GPUs (B200) for which most kernel optimizations were written. This architectural divergence creates a landscape where the fast fused kernel stack (DeepGEMM, trtllm-gen, FP4 indexer) is arch-gated to sm_100, forcing sm_120 to rely on Triton and PyTorch fallbacks that run on CUDA cores rather than tensor cores, achieving a fraction of the potential throughput.

Part I: Establishing the Foundation

Environment Setup and Driver Installation

The session begins with a blank slate: a remote machine running Ubuntu 24.04 with two NVIDIA RTX PRO 6000 Blackwell GPUs. The first task is to establish a complete ML development environment—a process that, despite being routine, reveals the complexities of working with hardware at the frontier of the software ecosystem.

The assistant installs NVIDIA drivers version 590.48.01 and CUDA Toolkit 13.1, configuring environment paths and verifying that both GPUs are operational. A Python virtual environment is created using uv, and core packages including PyTorch are installed. This establishes a solid foundation for the development stack [1].

The Flash-Attn Bottleneck

The installation of flash-attn becomes a major bottleneck, highlighting the challenges of building CUDA extensions for new architectures. The system's CUDA 13.1 conflicts with PyTorch's CUDA 12.8 base, requiring the installation of a secondary CUDA 12.8 toolkit. The build process repeatedly exhausts system memory, forcing a collaborative trial-and-error reduction of parallel compilation jobs (MAX_JOBS) from 128 down to 20. This is only resolved after the machine is rebooted with an expanded 432 GB of RAM [1].

Further dependency conflicts arise when vLLM downgrades PyTorch, breaking the compiled flash-attn binary. This necessitates a targeted rebuild against the correct PyTorch version (2.9.1). Ultimately, the environment is stabilized with a fully compatible stack including PyTorch 2.9.1, flash-attn 2.8.3, and vLLM 0.15.1 [1].

Hardware Upgrade and Model Deployment

The machine is then upgraded to 8 GPUs, and the user tasks the assistant with deploying the GLM-5-NVFP4 model using a nightly build of SGLang. This pivot from infrastructure setup to application deployment underscores the iterative nature of the work, moving from resolving deep system-level build issues to preparing for high-performance model serving [1].

Part II: Codebase Archaeology

Mapping the Blackwell Frontier

With the environment established, the assistant embarks on a systematic exploration of the SGLang codebase to understand how DeepSeek-V4-Flash interacts with sm_120 hardware. The user's request is extraordinarily detailed: a six-point questionnaire demanding file paths, line numbers, symbol names, and version strings for model requirements, server arguments, sm_120 handling, MoE quantization paths, dependency versions, and NIXL disaggregation setup [2].

The assistant begins with reconnaissance: checking the git commit (735a256 on main), listing top-level directories, and identifying the model files. The DeepSeek V4 model file (deepseek_v4.py) is 2,399 lines long, while the NEXTN/MTP module (deepseek_v4_nextn.py) is a more focused 283 lines [7]. This sizing-first approach is a hallmark of experienced investigators—measuring cognitive load before committing to deep reading.

The Architecture of DeepSeek-V4

Reading the model files reveals that DeepSeek-V4 uses a sophisticated sparse attention mechanism called DSA (DeepSeek Sparse Attention) with a Compressor and C4Indexer, not the more commonly discussed NSA (Native Sparse Attention) pattern [38]. The model employs a hybrid quantization scheme: FP8 for attention layers and FP4-packed (mxfp4) format for the MoE expert weights. The attention mechanism is Multi-head Latent Attention (MLA) with a split of nope_dim=448 and rope_dim=64 dimensions, using a paged KV cache [4].

The codebase includes specialized attention backends under dsv4/ and dsa/ directories, with separate CUDA and HIP (AMD) implementations. JIT-compiled kernels live under jit_kernel/dsv4/ and include compressors, indexers, dequantizers, and GEMM operations [26].

The Blackwell Architecture Divide

A critical discovery comes when the assistant traces the hardware detection utilities in common.py. The codebase distinguishes between three Blackwell-related GPU classes:

The FP4 Quantization Path

Tracing the quantization configuration reveals a non-obvious architectural choice: a stock FP4 checkpoint for DeepSeek V4 does not use a dedicated Mxfp4Config class. Instead, it loads as an Fp8Config with an is_fp4_experts boolean flag set to True [17]. The MoE layer's quantization method selection then branches on this flag combined with the MoE runner backend type.

On SM120 hardware, the standard Marlin CUDA kernel produces NaN values, forcing a fallback to a Triton-based backend called mxfp4_moe_sm120_triton [15]. This is a hardware-specific workaround baked into the quantization layer—the codebase explicitly accounts for the architectural discontinuity between sm_100 and sm_120.

The Attention Backend Divide

The attention mechanism presents a similar story. On SM120, the flash_mla CUDA kernel is unavailable, replaced by a custom Triton kernel and a pure-PyTorch fallback implemented in flash_mla_sm120.py [20]. The _create_flashmla_metadata function returns None on SM120 because the CUDA kernel for Flash MLA metadata generation simply isn't available on that architecture [8].

The prefill path introduces a further complication. A threshold constant _LARGE_INDEXER_QUERY_THRESHOLD = 11673 determines when prefill operations switch to a sparse kernel (flash_mla_sparse_fwd from sgl_kernel). This kernel lacks an SM120 implementation, meaning prefill chunks exceeding ~11,673 tokens could fall through to an unsupported path on Blackwell hardware [18].

The Indexer Fallback Chain

The indexer—responsible for computing attention logits against the KV cache—reveals a multi-layer fallback chain. The highest-performance path uses DeepGEMM for FP4 operations, but DeepGEMM is only available on sm_100. The fallback chain then checks for TileLang, then AITER (for AMD HIP), and ultimately routes to a pure PyTorch implementation activated by the SGLANG_FP8_PAGED_MQA_LOGITS_TORCH flag, which the DeepSeek V4 hook sets to True on SM120 [4].

NIXL Disaggregation Infrastructure

For prefill-decode (PD) disaggregation, the assistant traces the NIXL transfer backend. NIXL is not listed in pyproject.toml dependencies, meaning it must be installed separately [25]. The backend selection defaults to UCX (via SGLANG_DISAGGREGATION_NIXL_BACKEND = EnvStr("UCX")), and the server argument disaggregation_transfer_backend defaults to "mooncake"—not NIXL [21]. This means that deploying with NIXL requires explicit configuration changes [23].

The assistant discovers that NIXL is the appropriate backend for non-datacenter hardware because mooncake validates InfiniBand devices during startup, which may not be present on workstation-class RTX PRO 6000 systems [23]. NIXL delegates device selection to UCX, which supports multiple transports (TCP, shared memory, RDMA) and can gracefully degrade.

The Cookbook Discovery

After extensive code archaeology, the assistant discovers a dedicated cookbook file containing verified launch commands, RTX 6000 configuration notes, and deployment instructions for DeepSeek-V4 [1]. This marks a pivot from bottom-up code exploration to top-down documentation consultation. The cookbook reveals the authoritative deployment configuration: single-node setup, DeepSeek-V4-Flash with FP4 quantization, low-latency strategy, tensor parallelism across 4 GPUs, and the Marlin MoE runner backend [33].

The cookbook also documents critical constraints: the RTX PRO 6000 with 96 GB can only run Flash inference (the Pro variant doesn't fit across 8×96 GB), HiCache and MegaMoE aren't supported on this hardware, and a nightly development image is required since SM120 support hasn't made it into the stable release [33].

Part III: Deployment and Optimization

Achieving Prefill-Decode Disaggregation

The assistant successfully deploys DeepSeek-V4-Flash on SGLang main across 8× RTX PRO 6000 GPUs. After stopping the previous service, freeing disk space, and downloading the 146 GB checkpoint, a fresh SGLang environment is built with all required dependencies (flashinfer 0.6.12, sglang-kernel 0.4.3, tilelang, NIXL, etc.). The model loads and generates correctly on a single-node TP4 configuration.

The headline deliverable—prefill-decode disaggregation—is 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 [chunk summary].

The Performance Wall

However, performance falls far short of expectations: approximately 10 tok/s at batch size 1 and ~25 tok/s at concurrency 16, versus the user's target of ~1000 tok/s. The assistant systematically exhausts every configuration lever:

Profiling the Bottleneck

A definitive GPU profile traces 63% of decode time to a single kernel: _tiled_sparse_decode_kernel, the sm_120 Triton fallback for sparse MLA attention. This kernel 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 earlier work [chunk summary].

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 sm_100, and no amount of configuration tuning can close the ~40× gap to 1000 tok/s.

MTP Speculative Decoding and NVFP4 Quantization

The assistant pulls the latest upstream SGLang, generates SM120 FP8 GEMM autotune configs, and enables MTP/EAGLE speculative decoding. The FP8 configs give only ~6% improvement (since FP8 GEMM is only 6% of decode time). MTP boosts single-request throughput by 47% but provides zero gain at concurrency due to verifier saturation [chunk summary].

A C=16 profile reveals the decisive bottleneck: the MoE slot-GEMV and sparse-decode attention kernels consume 39% and 38% of GPU time respectively, both running on CUDA cores (SIMT) rather than the sm_120a FP4/FP8 tensor cores. This explains the user's observation of <1% tensor-pipe utilization.

The assistant then pivots to the highest-leverage fix: switching from the stock MXFP4 checkpoint to the official NVIDIA NVFP4 quantization, which routes MoE execution through tensor-core paths (Marlin W4A16 or native cutlass FP4 grouped GEMM). PR #25820 is fetched and applied for correct NVFP4 auto-detection, and the 149 GB checkpoint is downloaded. Both NVFP4 backends deliver identical throughput (~28 tok/s at C=16, a ~24% improvement over baseline), confirming that MoE is now on tensor cores but attention remains the dominant bottleneck [chunk summary].

The MTP Memory Overhead

Crucially, the user notes that the hardware should be capable of 300–600 tok/s at C=16 based on roofline analysis (1.9 TB/s VRAM bandwidth, 13B active parameters). The assistant discovers that the MTP verifier consumes enough GPU memory to halve the effective batch size from 16 to 8, explaining a large portion of the throughput gap.

Part IV: The Hard Ceiling

The Structural Bottleneck

The investigation culminates in a sobering conclusion: configuration tuning and even checkpoint format changes can only deliver incremental gains against a structural bottleneck. The sm_120 fallback kernels for sparse-MLA decode and MoE slot-GEMV are fundamentally inefficient, running on CUDA cores instead of tensor cores. The NVFP4 checkpoint fixes the MoE path, but attention remains the primary limiter.

The path forward requires either:

  1. Disabling MTP to restore batch capacity and re-measuring
  2. Building a custom split-K tensor-core sparse-attention kernel to attack the remaining 38% bottleneck—the same playbook that delivered 3–6× gains in earlier K2.6 work [chunk summary]

The Architecture Gate

The fundamental issue is that NVIDIA's Blackwell architecture has two distinct compute capabilities: sm_100 (datacenter B200) and sm_120 (workstation RTX PRO 6000). The software ecosystem—DeepGEMM, trtllm-gen, flash_mla CUDA kernels, the FP4 indexer—was optimized for sm_100. On sm_120, these kernels either don't exist or produce incorrect results (NaN from Marlin), forcing fallbacks to Triton and PyTorch implementations that run on general-purpose CUDA cores rather than specialized tensor cores.

This is not a bug that can be fixed with configuration changes. It is a fundamental limitation of the current software stack's support for the sm_120 architecture. 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.

Conclusion

The deployment of DeepSeek-V4-Flash on NVIDIA RTX PRO 6000 Blackwell GPUs is a story of systematic engineering triumph and architectural limitation. The assistant successfully established a complex ML environment, navigated a sprawling codebase to map every sm_120-specific code path, deployed the model with prefill-decode disaggregation, and executed a methodical optimization campaign spanning MTP speculative decoding, NVFP4 quantization, and FP8 autotuning.

Yet the performance ceiling—approximately 28 tok/s at C=16 against a target of 1000 tok/s—was ultimately determined not by configuration choices but by the availability of optimized kernels for the sm_120 architecture. The fast fused kernel stack that powers DeepSeek-V4 on datacenter Blackwell GPUs is simply not available on the workstation variant, and no amount of tuning can compensate for the absence of tensor-core-accelerated sparse attention kernels.

This session serves as a powerful case study in the realities of hardware bring-up: the gap between what a GPU architecture can theoretically deliver and what the software ecosystem actually supports is often the decisive factor. For engineers deploying cutting-edge models on new hardware, the lesson is clear: before investing in optimization, first map the kernel availability landscape. The ceiling may be lower than expected, and only targeted kernel development—not configuration tuning—can raise it.## References

[1] The Cookbook Discovery: A Pivot from Code Archaeology to Documentation ([chunk 67.0]) [2] Mapping the Blackwell Frontier: A Deep Dive into the DeepSeek-V4-Flash SGLang Bring-Up Request ([chunk 67.0]) [3] Tracing the FP4 Expert Detection: A Deep Dive into DeepSeek V4 Quantization Configuration on Blackwell GPUs ([chunk 67.0]) [4] Tracing the SM120 Fallback Chain: How SGLang Routes Around Missing Hardware Support for DeepSeek V4 on Blackwell Desktop GPUs ([chunk 67.0]) [5] Tracing the SM120 Sparse Prefill Boundary in DeepSeek-V4's Attention Backend ([chunk 67.0]) [6] Verifying the NIXL Backend: A Methodical Verification in the DeepSeek V4 SM120 Codebase ([chunk 67.0]) [7] The First Cut: How a Single wc -l Shaped the DeepSeek-V4 Exploration ([chunk 67.0]) [8] Tracing the SM120 Decode Path: A Deep Dive into GPU Architecture Adaptation ([chunk 67.0]) [9] The Verification of a Backend: Tracing NIXL's Place in SGLang's Disaggregation Architecture ([chunk 67.0]) [10] Navigating the SM120 Frontier: A Deep Dive into SGLang's DeepSeek V4 Server Argument Exploration ([chunk 67.0]) [11] Reading the Configuration Hooks: A Pivotal Data-Gathering Step in Mapping SGLang for DeepSeek-V4 on Blackwell ([chunk 67.0]) [12] The Blackwell Architecture Divide: Tracing sm_120 vs sm_100 GPU Support in SGLang's DeepSeek-V4 Serving Path ([chunk 67.0]) [13] The Fork That Had SM120 All Along: A Pivotal Realization in the DeepSeek-V4 Blackwell Bring-Up ([chunk 67.0]) [14] Tracing the Configuration Labyrinth: How SGLang Routes DeepSeek V4 Through SM120-Specific MoE Paths ([chunk 67.0]) [15] Tracing the NaN Fault Line: How SM120 Blackwell Redirects DeepSeek V4's MoE from Marlin to Triton ([chunk 67.0]) [16] Tracing the FP4 Quantization Path: How SGLang Routes DeepSeek V4's MoE Layers on Blackwell GPUs ([chunk 67.0]) [17] Tracing the FP4 Quantization Path: A Deep Dive into SM120 Kernel Selection for DeepSeek V4 ([chunk 67.0]) [18] The 11,673-Token Threshold: Uncovering a Coverage Gap in Blackwell's DeepSeek V4 Attention Backend ([chunk 67.0]) [19] Mapping the NIXL Disaggregation Backend: A Targeted Code Exploration for DeepSeek V4 on SM120 ([chunk 67.0]) [20] Tracing the SM120 Sparse Decode Path: A Pivot Point in DeepSeek-V4 Bring-Up on Blackwell ([chunk 67.0]) [21] Mapping the NIXL Disaggregation Backend: A Methodical Search Through SGLang's Configuration ([chunk 67.0]) [22] The Configuration Hook: Tracing DeepSeek V4's Blackwell Configuration Logic ([chunk 67.0]) [23] The NIXL Dispatch: Uncovering the Right Disaggregation Backend for Blackwell Desktop ([chunk 67.0]) [24] The Architecture Gate: Tracing SM Detection in SGLang's DeepSeek-V4 Deployment ([chunk 67.0]) [25] Tracing the Dependency Trail: How an AI Agent Discovered NIXL's Missing Link in the SGLang Codebase ([chunk 67.0]) [26] Mapping the Blackwell Frontier: How an AI Agent Dissected SGLang's DeepSeek-V4 Support for sm_120 GPUs ([chunk 67.0]) [27] Tracing the SM120 Code Path: A Deep Dive into SGLang's DeepSeek V4 Indexer Exploration ([chunk 67.0]) [28] Tracing the FP4 Quantization Path: How SGLang Routes DeepSeek V4's MoE Layers on Blackwell GPUs ([chunk 67.0]) [29] The Pivot Point: Tracing SM120 Configuration Logic in SGLang's DeepSeek V4 Server Arguments ([chunk 67.0]) [30] A Read That Missed Its Mark: Tracing the NIXL Validation in SGLang's Disaggregation Configuration ([chunk 67.0]) [31] Tracing the NaN Fault Line: How SM120 Blackwell Redirects DeepSeek V4's MoE from Marlin to Triton ([chunk 67.0]) [32] The Authoritative Configuration: A Pivotal Discovery in the DeepSeek-V4 SM120 Bring-Up ([chunk 67.0]) [33] The Authoritative Configuration: Decoding SM120 Deployment Constraints for DeepSeek-V4 on Blackwell Desktop GPUs ([chunk 67.0]) [34] Tracing the FP4 Quantization Path: A Deep Dive into SM120 Kernel Selection for DeepSeek V4 ([chunk 67.0]) [35] The Opening Move: Reconnaissance in a Deep-Sea Codebase ([chunk 67.0]) [36] The Final Verification: Confirming SM120 Kernel Support in SGLang's DeepSeek-V4 Pipeline ([chunk 67.0]) [37] The Architecture Gate: Confirming SM120's Hard Ceiling in SGLang's DeepSeek-V4 Stack ([chunk 67.0]) [38] Navigating the SGLang Codebase: A Methodical Deep Dive into DeepSeek-V4 on Blackwell ([chunk 67.0]) [39] The Final Verification: Tracing Attention Backend Aliases and Distributed Init in SGLang's DeepSeek V4 Deployment ([chunk 67.0]) [40] The Final Verification Pass: Confirming DeepSeek-V4 Attention Backend Registration on Blackwell SM120 ([chunk 67.0]) [41] The Pivot Point: From Kernel Verification to Disaggregation Architecture in SM120 Bring-Up ([chunk 67.0]) [42] The Pivot Point: Tracing SM120 Hardware Gating in SGLang's DeepSeek-V4 Backend ([chunk 67.0]) [43] Tracing FP4 Experts Through the SGLang Stack: A Subagent's Final Verification ([chunk 67.0]) [44] When Tools Deceive: Tracing FP4 Configuration Through Redacted Outputs in an SGLang Debugging Session ([chunk 67.0]) [45] The Architecture of Inference: Mapping DeepSeek-V4-Flash onto Blackwell's sm_120 ([chunk 67.0])