The Nightly Pivot: Deploying Qwen3.5-397B-A17B-NVFP4 on Blackwell GPUs Through Strategic Stack Upheaval
Introduction
In the high-stakes world of deploying large language models on cutting-edge hardware, the difference between a working inference server and a NaN-spewing disaster often comes down to a single flag, a kernel version mismatch, or a missing patch. This article synthesizes a pivotal chunk of a sprawling opencode coding session in which an AI assistant and its user collaborate to deploy the Qwen3.5-397B-A17B-NVFP4 model — a 397-billion-parameter Mixture-of-Experts model quantized to NVIDIA's NVFP4 format — on a machine equipped with 8× NVIDIA RTX PRO 6000 Blackwell Server Edition GPUs (compute capability SM120). What unfolds is a masterclass in strategic engineering: a deliberate, calculated pivot from a stable, known-working software stack to bleeding-edge nightly builds, driven by the hard reality that only the latest code contains the Blackwell-specific FP4 kernels needed for correct inference.
The chunk spans approximately 140 messages (indices 5856–5993) and covers the full arc from diagnosis to production deployment. It begins with the model producing garbage output (repeated exclamation marks and NaN values), proceeds through a user-ordered "update all to nightly" directive, executes a multi-component stack upgrade (PyTorch → FlashInfer → sgl-kernel → SGLang), conducts exhaustive backend testing to find a working FP4 GEMM configuration, fixes a critical FP8 KV cache accuracy issue, and culminates in a production systemd service achieving ~172 tok/s single-request and over 2100 tok/s at high concurrency. This article examines the key decisions, reasoning patterns, assumptions, and technical insights that define this journey.
The Crisis: Garbage Output on Blackwell GPUs
The chunk opens at a moment of frustration. The assistant had successfully deployed Kimi-K2.5 INT4 with EAGLE-3 speculative decoding, achieving impressive throughput numbers ([msg 5856]). But the pivot to Qwen3.5-397B-A17B-NVFP4 hit a wall. When the assistant started the server with --attention-backend triton --quantization modelopt_fp4 and sent a test prompt, the model responded with repeated exclamation marks — pure garbage. A second test returned null content with a finish_reason of "NaN happened" ([msg 5843], [msg 5844]). The FP4 kernels were producing numerical garbage on the SM120 (Blackwell) hardware.
The root cause was identified through a critical external reference: a GitHub Gist by developer "catid" at https://gist.github.com/catid/87cca824963f17fe7479a0ed26221397 ([msg 5853]). The gist documented that the default auto MoE runner backend selects a kernel that produces NaN on SM120, and that the fix requires explicitly pinning --moe-runner-backend flashinfer_cutlass and --fp4-gemm-runner-backend flashinfer_cudnn. But there was a deeper concern: the pre-built sgl-kernel 0.3.21+cu130 wheel might not contain SM120-specific FP4 kernels at all, because it was compiled against PyTorch 2.9.1 without Blackwell architecture support.
The State Dump: Consolidating Knowledge Before Action
Before diving into the fix, the assistant produced one of the most remarkable messages in the entire session: a comprehensive state document at [msg 5856] that consolidated everything it knew about the deployment environment. This "state dump" spanned hardware configuration (8× RTX PRO 6000 Blackwell GPUs, PCIe Gen5, no NVLink), software stack versions (PyTorch 2.9.1+cu130, sgl-kernel 0.3.21+cu130, flashinfer 0.6.4, CUDA 13.0.1), ABI compatibility notes, NCCL tuning parameters, GPU switching workflows between LXC containers and VMs, benchmark results from the completed Kimi-K2.5 work, model characteristics of Qwen3.5, and a detailed action plan with four prioritized steps.
This message is remarkable because it represents the assistant's metacognitive awareness of its own knowledge boundaries. The assistant recognized that it was about to undertake a complex, multi-step process and needed a consolidated reference to avoid losing track of version numbers, file paths, and configuration parameters across dozens of subsequent messages. The document served as an external memory system — a cognitive artifact that the assistant could refer back to throughout the upgrade process [1].
The state dump also revealed the assistant's assumptions and risk awareness. It explicitly noted that "torch 2.10.0+cu130 and nightly 2.12.0+cu130 have ABI break with sgl-kernel 0.3.21" and cautioned that building sgl-kernel from source "will likely conflict with torch 2.9.1 since catid uses nightly 2.12.0." The assistant's conservative plan was to try the simpler backend flag fix first, and only escalate to a source build if needed.
The User's Directive: "Update All to Nightly"
The user's response at [msg 5857] cut through this cautious, multi-step plan with a single, unambiguous command: "Update all to nightly, the patches are more recent than our build." This was not a suggestion — it was a directive that fundamentally altered the trajectory of the session.
The reasoning embedded in this seven-word message is worth unpacking. The user had absorbed the same information the assistant had: catid's gist was the authoritative reference for making NVFP4 inference work on Blackwell, and that gist explicitly built sgl-kernel from source against PyTorch nightly. The user recognized a logical inconsistency in the assistant's plan: if the patches (the CMakeLists.txt modifications, the CUDA 13 cccl include path fixes, the FA3 fallback import) were designed for the nightly build, then trying to apply them to the stable stack would be fighting against the grain. The ABI incompatibility the assistant worried about was not a bug to be worked around — it was a signal that the entire stack needed to move forward together [2].
The user followed up at [msg 5858] with a second message that defined the operational boundaries: "Continue if you have next steps, or stop and ask for clarification if you are unsure how to proceed." This "permission slip" simultaneously granted autonomy, established a fallback protocol, and validated the assistant's expertise. In 16 words, it accomplished what many project managers struggle to achieve in paragraphs [3].
The Pivot: From Conservative to Aggressive
The assistant's response at [msg 5859] was the acceptance and execution plan for this directive. The message announced: "I'll update the entire stack to PyTorch nightly and rebuild sgl-kernel from source with SM120 support. The rationale is that the pre-built sgl-kernel 0.3.21+cu130 was built against older torch, and catid's gist (which works) uses nightly."
This was followed by a structured todowrite with four high-priority tasks:
- Upgrade PyTorch to nightly (2.12.0+cu130 or latest)
- Upgrade flashinfer to latest nightly/compatible version
- Build sgl-kernel from source with SM120 FP4 support (TORCH_CUDA_ARCH_LIST=12.0a)
- Apply catid's CMakeLists.txt patches to sgl-kernel The assistant's reasoning revealed a sophisticated understanding of the dependency chain. The key insight was that the ABI incompatibility between nightly PyTorch and sgl-kernel 0.3.21 existed only because the pre-built wheel was compiled against an older torch version. If you build sgl-kernel from source against the same nightly PyTorch, the ABI mismatch disappears. The constraint was not the torch version itself, but the pre-built binary [4]. This is a classic systems-engineering tradeoff: stability versus capability. The stable stack (torch 2.9.1) was safe but couldn't run the model correctly. The nightly stack (torch 2.12.0) was risky but was the only path demonstrated to work on SM120. The assistant correctly identified that the risk could be mitigated by rebuilding sgl-kernel from source.
The Reconnaissance: Measuring Before Changing
Before executing the upgrade, the assistant paused to gather intelligence. At [msg 5860], it checked the current state of the container — confirming that all 8 GPUs showed 0 MiB memory usage and no SGLang or Python inference processes were running. This was standard operational procedure: never make changes to a system you don't understand [5].
At [msg 5861], the assistant issued two parallel tool calls: a webfetch to retrieve catid's gist for the exact build steps, and a bash command to query the current software versions. The results confirmed torch==2.9.1+cu130, sgl_kernel==0.3.21, and flashinfer==0.6.4 — establishing the baseline for the upgrade. The webfetch returned garbled HTML rather than raw markdown, a limitation that would force the assistant to adapt its approach in subsequent messages [6].
At [msg 5862], the assistant acknowledged having the full gist and announced the next steps: checking the latest nightly torch cu130 available and examining the sgl-kernel source. The todo list was updated to mark the PyTorch upgrade as "in_progress," signaling the transition from planning to execution [7].
Step 1: Upgrading PyTorch to Nightly
The first concrete action came at [msg 5863], where the assistant executed a dry-run upgrade using uv pip install --dry-run against the PyTorch nightly index for CUDA 13.0. The --dry-run flag was a deliberate choice — it allowed the assistant to preview what would change without actually modifying the environment. The output revealed 33 packages would change, including CUDA bindings jumping from 12.9.4 to 13.0.3 [8].
At [msg 5864], the assistant filtered the dry-run output to the critical line: torch==2.12.0.dev20260307+cu130 would replace torch==2.9.1+cu130. This single data point — confirming that the nightly version existed and was accessible — was the green light needed to proceed [9].
The actual installation at [msg 5865] used uv pip install --upgrade --force-reinstall against the nightly index. The output showed a dramatic version jump: torch from 2.9.1 to 2.12.0.dev, torchaudio from 2.9.1 to 2.11.0.dev, torchvision from 0.24.1 to 0.26.0.dev, and triton from 3.5.1 to 3.6.0+git. The NVIDIA CUDA libraries were also upgraded — nvidia-nvjitlink from 13.0.39 to 13.0.88, nvidia-nvshmem-cu13 from 3.3.24 to 3.4.5. The assistant labeled this "Step 1" and proceeded methodically [10].
The verification at [msg 5866] was minimal but decisive: a three-line Python script confirming torch.__version__ as 2.12.0.dev20260307+cu130, torch.cuda.is_available() as True, and the GPU as NVIDIA RTX PRO 6000 Blackwell Server Edition. The nightly PyTorch was operational on the target hardware [11].
The todo list was updated at [msg 5867]: PyTorch upgrade marked "completed," flashinfer upgrade moved to "in_progress." The structured checklist provided continuity across the multi-step process [12].
Step 2: The FlashInfer Gauntlet
The flashinfer upgrade proved far more challenging than PyTorch. At [msg 5868], the assistant tried uv pip install --dry-run flashinfer and got a hard failure: "No solution found when resolving dependencies — Because flashinfer was not found in the package registry." The package simply didn't exist on PyPI under that name [13].
At [msg 5869], the assistant tried flashinfer-python (the actual PyPI name). The dry-run revealed a critical conflict: installing flashinfer-python from PyPI would downgrade torch from 2.12.0.dev20260307+cu130 to 2.10.0. This was a catastrophic regression — the entire point of the nightly upgrade was to get Blackwell support, and torch 2.10.0 doesn't even target CUDA 13. The assistant's cautious dry-run approach saved it from silently undoing the entire stack upgrade [14].
At [msg 5870], the assistant tried the FlashInfer nightly wheel index at https://flashinfer.ai/whl/cu130/torch2.12/nightly/. The resolver returned "No solution found" — the package wasn't in that registry either. The assistant noted: "As expected — flashinfer-python from PyPI would downgrade torch to 2.10.0. We need a nightly flashinfer build" [15].
At [msg 5871], the assistant probed the URL directly with curl and discovered a GitHub Pages 404 page. The nightly index for torch 2.12 simply didn't exist. This was a definitive answer: the FlashInfer project had not published pre-built wheels for this torch version [16].
At [msg 5872], the assistant broadened the search. It checked the root CUDA 130 index at https://flashinfer.ai/whl/cu130/ and found a single subdirectory: flashinfer-jit-cache/. This was the breakthrough. The JIT-cache approach is torch-version-independent — kernels are compiled at runtime against whatever torch is installed. The assistant also checked for a torch2.11 index (also a 404), confirming that no pre-built wheels existed for any torch version under CUDA 130 [17].
At [msg 5873], the assistant confirmed that the flashinfer-jit-cache index was live and that the current FlashInfer installation was still importable after the PyTorch upgrade. The JIT-cache path was viable [18].
The resolution came at [msg 5874] through a flash of insight. The assistant realized that flashinfer-jit-cache wheels are torch-version-independent — they contain pre-compiled CUDA kernels that don't link against specific torch binaries. By upgrading just the JIT-cache and cubin components (which are decoupled from the torch version), the assistant could get FlashInfer 0.6.5 working with the nightly PyTorch without any dependency conflict. The subsequent messages confirmed this approach: flashinfer-jit-cache 0.6.5+cu130 was installed, and the package remained functional [19][20][21][22][23][24][25].
The Status Checkpoint: Consolidating Progress
At [msg 5881], the assistant produced another comprehensive status update, documenting the completed PyTorch and FlashInfer upgrades and outlining the remaining steps: updating SGLang to the latest main branch, building sgl-kernel from source with SM120 support, and deploying the model. The todo list served as a decision record, making the assistant's progress transparent to the user [26].
Step 3: Building sgl-kernel from Source with SM120 Support
The sgl-kernel build was the most technically complex step in the entire upgrade. The assistant had to navigate a multi-layered build system, apply patches from catid's gist, and compile CUDA kernels specifically for the Blackwell (SM120) architecture.
At [msg 5882], the assistant checked the current state of the SGLang repository, noting that it was on the main branch with custom SM120 patches applied to all_reduce_utils.py and torch_symm_mem.py. The assistant recognized that pulling the latest upstream changes could overwrite these patches, so it planned to stash them first [27][28].
At [msg 5884], the assistant executed git stash to save the custom patches, then pulled the latest upstream changes from SGLang's main branch. The output showed a fast-forward merge with changes to multiple files. The assistant then popped the stash to reapply the SM120 patches on top of the updated code [29][30].
At [msg 5886], the assistant checked whether upstream SGLang had added native SM120 support. A grep for 12_0 in the sgl-kernel source revealed that upstream had not yet added Blackwell support — the patches were still necessary [31][32].
The build process began at [msg 5897] with the assistant examining the sgl-kernel build system. The pyproject.toml revealed that the build used scikit-build-core with CMake, and the CMakeLists.txt contained the architecture targeting logic. The assistant identified the need to set TORCH_CUDA_ARCH_LIST=12.0a to enable Blackwell-specific code paths [33][34][35][36][37][38][39][40].
The first build attempt at [msg 5902] failed with a Python interpreter mismatch — the build system was using the system Python 3.12 instead of the virtual environment's Python. The assistant diagnosed this and set the correct PYTHON path [47][48][49][50].
The second attempt at [msg 5906] progressed further but failed with a CMake policy error related to CMP0169 and CMP0177. This was exactly the issue catid's patches addressed. The assistant applied the patches: adding CMake policy guards, fixing CUDA 13 cccl include paths, and adding FA3 import fallback [51][52][53].
The successful build at [msg 5908] used:
export TORCH_CUDA_ARCH_LIST=12.0a
export CMAKE_ARGS='-DSGL_KERNEL_COMPILE_THREADS=16 -DENABLE_BELOW_SM90=OFF -DSGL_KERNEL_ENABLE_FP4=ON -DSGL_KERNEL_ENABLE_FA3=OFF'
The compilation completed, producing a custom sgl-kernel with Blackwell FP4 support [53].
The verification at [msg 5909] confirmed that the built wheel contained SM120 kernels. The assistant checked the compiled .so files for Blackwell architecture markers and ran a quick import test to ensure the module loaded correctly [54][55][56].
The assistant then rebuilt SGLang itself against the nightly PyTorch stack at [msg 5912], ensuring that the entire inference engine was ABI-compatible with the new PyTorch and sgl-kernel versions [57].
The Verification: Everything Loads Correctly
At [msg 5913], the assistant performed the definitive verification: importing all critical packages and confirming they loaded without errors. The output showed torch==2.12.0.dev20260307+cu130, sgl_kernel==0.3.21 (custom build), flashinfer==0.6.5, and SGLang importing cleanly. This was the moment when the nightly stack was confirmed operational [58][59][60].
Step 4: Backend Testing — Finding the Working Configuration
With the stack upgraded, the assistant turned to the critical question: which FP4 GEMM and MoE backends actually produce correct output on SM120? The assistant conducted extensive testing, launching the server with different backend combinations and checking for NaN/garbage output.
At [msg 5929], the assistant investigated the model's architecture and discovered that Qwen3.5-397B-A17B-NVFP4 has built-in Multi-Token Prediction (MTP) heads — a form of speculative decoding that SGLang supports via the --speculative-algorithm NEXTN flag. This was a significant discovery: the model already contains the drafter heads, so no external EAGLE-3 drafter is needed [74][75][76][77][78][79][80][81][82][83][84][85][86].
The backend testing at [msg 5942] through [msg 5972] was exhaustive. The assistant systematically tested each combination:
flashinfer_trtllm— crashed or produced NaN/garbage on SM120flashinfer_cutedsl— produced garbage output (silent numerical corruption)flashinfer_cutlassfor MoE — worked correctlyflashinfer_cudnnfor FP4 GEMM — worked correctly The working configuration was--moe-runner-backend flashinfer_cutlass --fp4-gemm-runner-backend flashinfer_cudnn. The assistant documented this finding and codified it into the production service file [87][88][89][90][91][92][93][94][95][96][97][98][99][100][101][102][103][104][105][106][107][108][109][110][111][112][113][114][115][116][117][118][119][120].
The FP8 KV Cache Accuracy Fix
During backend testing, the assistant identified a critical accuracy issue that could have silently degraded model quality. The Qwen3.5 checkpoint's default KV cache dtype was FP8, but the checkpoint lacked proper FP8 scaling factors. Without these scaling factors, the FP8 KV cache would introduce numerical errors that could compound over long sequences — particularly problematic for the agentic coding tasks the model was intended for.
The assistant proactively forced the KV cache to BF16 using --kv-cache-dtype bf16. This provided approximately 1.57 million tokens of high-precision cache — more than sufficient for the intended workload. The fix was documented in the production configuration, and the assistant noted that this decision prioritized correctness over the marginal memory savings of FP8 [61][62][63][64][65][66][67][68][69][70][71][72][73].
Production Deployment
The final production configuration was codified into a systemd service file at /etc/systemd/system/sglang-qwen.service. The service included all the critical flags discovered through the testing process:
--attention-backend triton(required for hybrid GDN models)--moe-runner-backend flashinfer_cutlass(working MoE backend on SM120)--fp4-gemm-runner-backend flashinfer_cudnn(working FP4 backend on SM120)--kv-cache-dtype bf16(accuracy fix for missing FP8 scaling factors)--quantization modelopt_fp4(NVFP4 quantization format)--reasoning-parser qwen3 --tool-call-parser qwen3_coder(model-specific parsers)--disable-custom-all-reduce(required for PCIe-connected GPUs without NVLink) The service was deployed at [msg 5918] viascpto the remote container, and the server was launched at [msg 5920]. The smoke test at [msg 5922] confirmed correct output: the model produced coherent, meaningful text rather than the garbage output that plagued the initial attempts [62][63][64][65][66][67][68][69][70][71][72][73].
Performance Results
The final production deployment achieved impressive throughput:
- Single-request concurrency (C=1): ~172 tok/s
- High concurrency (C=32): over 2100 tok/s aggregate The assistant also tested built-in MTP speculative decoding (
--speculative-algorithm NEXTN), which successfully loaded the model's native MTP heads. However, on synthetic benchmark workloads, NEXTN showed no throughput gain — the baseline performance was already optimal for the benchmark's short-context, high-throughput pattern. The assistant noted that NEXTN might still provide benefits on real-world workloads with longer contexts and more complex generation patterns [73][74][75][76][77][78][79][80][81][82][83][84][85][86][120][121][122][123][124][125][126][127][128][129][130][131][132][133][134][135][136][137][138].
Themes and Insights
Several overarching themes emerge from this chunk:
Strategic Pivoting: The user's directive to "update all to nightly" was a decisive strategic intervention that overrode the assistant's conservative plan. The user recognized that the patches were designed for the nightly ecosystem and that trying to apply them to a stable stack would create more problems than it solved. This illustrates the value of human judgment in AI-assisted engineering — knowing when to abandon caution for coherence.
Methodical Execution: Despite the aggressive strategic direction, the assistant's execution was remarkably methodical. Every upgrade was preceded by a dry-run, every installation was verified, and every configuration change was tested. The use of --dry-run flags, structured todo lists, and incremental verification steps prevented catastrophic failures at multiple points — most notably when the FlashInfer PyPI package would have silently downgraded PyTorch.
Dependency Graph Navigation: The upgrade path revealed the complex dependency graph of modern ML infrastructure. PyTorch → FlashInfer → sgl-kernel → SGLang forms a chain where each component must be ABI-compatible with its neighbors. The assistant's decision to build sgl-kernel from source against the nightly PyTorch was the key insight that made the entire upgrade viable — it resolved the ABI incompatibility that had blocked earlier attempts.
Empirical Debugging: When assumptions about package availability failed (the FlashInfer nightly index didn't exist), the assistant pivoted to empirical probing — using curl to directly inspect the wheel server, discovering the JIT-cache approach, and adapting the strategy accordingly. This willingness to probe the environment directly rather than relying on documentation or assumptions was critical to the success.
Accuracy Over Performance: The FP8 KV cache fix exemplifies the assistant's commitment to correctness over marginal performance gains. The checkpoint's default FP8 KV cache lacked proper scaling factors, which would degrade long-context inference quality. The assistant proactively forced BF16, accepting a slight memory increase in exchange for guaranteed numerical accuracy.
Conclusion
This chunk of the coding session represents a complete arc from failure to production deployment. The Qwen3.5-397B-A17B-NVFP4 model, initially producing garbage output on Blackwell GPUs, was transformed into a production-grade inference service through a strategic pivot to nightly builds, methodical stack upgrades, exhaustive backend testing, and careful accuracy tuning. The final result — ~172 tok/s single-request and over 2100 tok/s at high concurrency on 8× RTX PRO 6000 Blackwell GPUs — represents a significant engineering achievement.
The session demonstrates that deploying cutting-edge ML models on novel hardware requires not just technical skill but strategic judgment: knowing when to abandon stability for capability, when to probe rather than assume, and when to prioritize correctness over performance. The collaboration between human and AI — the user providing strategic direction and the assistant executing methodically — proved to be a powerful combination for navigating the complex landscape of nightly builds, custom kernel compilation, and architecture-specific optimization.