One SGLang runtime, two Qwen3.8 configurations, one RTX PRO 6000

One optimized SGLang runtime for a single RTX PRO 6000, with qualified launch configurations for both Qwen3.8-27B/DFlash2 and Qwen3.8 Flash-Next.

This repository contains the complete SGLang-derived source used on one NVIDIA RTX PRO 6000 Blackwell Workstation Edition (96 GB, SM120, TP=1). It is not two builds: both model configurations run from the same patched source.

• Qwen3.8-27B pairs an FP8 target with incoai/Qwen3.8-27B-DFlash2.
• Qwen3.8 Flash-Next pairs an NVFP4 target with its native NEXTN MTP layer.

Flash-Next is day-one engineering. It has extensive qualification on this exact machine—including 524K context, multimodal input, reasoning, tools, agentic workloads, CUDA-graph recovery, and persistent prefix restoration—but it may still contain rough edges or hardware/model-specific assumptions.

Validation suite

The canonical test suite used to qualify these runtimes—including reasoning, tool calling, long-context needles, vision, and concurrent decode—is maintained in jpezzulli/pennyroyal-validation. This repository contains the runtime source, configuration, and measured results; pennyroyal-validation contains the reusable tests and result catalog.

Why this runtime exists

The important work is architectural, not merely a collection of launch flags:

• DFlash2 integration for Qwen3.8-27B, with the selector inside the draft CUDA graph, fused draft KV materialization, and fixed-width TRTLLM-MHA/XQA target verification on SM120.
• Native NEXTN MTP for Flash-Next with QSA index sharing, recovery graphs, and gdn_mtp_cache_mode=none.
• FlashInfer GDN decode and prefill for Flash-Next, plus a deliberately narrow SM120 FlashInfer WY output-only RecoverSSM path.
• QSA sparse attention: Triton sparse-GQA prefill, the FlashInfer QSA wrapper resolving to XQA on SM120, and sgl-kernel top-k.
• FlashInfer CUTLASS MoE for the Flash-Next target and native-MTP layer.
• Correct three-axis multimodal mRoPE in the fused QK RMSNorm+RoPE kernel.
• Target, draft, verify, prefill, and accepted-state recovery CUDA graphs.
• FP8 KV and 524,288-token factor-2 YaRN for both launch configurations.
• HiCache/NIXL persistence for complete hybrid state—not KV alone—including DFlash2 side state or Flash-Next GDN, Qwen4 PLE, and compressed QSA state.
• Restart restoration, representation-specific FILE namespaces, and radix- prefix reuse followed by long continuation.
• Measured reasoning, tool-calling, vision, long-context needle, controlled decode, and real agentic behavior.

Published identity

The current documentation commit sits above the executable source. It does not change python/, rust/, or the kernel trees. Exact source lineage and current upstream status are in CHANGES.md and PROVENANCE.md.

Qualified configuration matrix

Rows intentionally separate representation, datatype, backend, and capacity. The BF16 runtime dtype is not presented as a claim that every quantized kernel executes or accumulates entirely in BF16.

The current 27B launcher uses the 24-slot/five-state setting qualified on 2026-08-26. The earlier 2026-08-24 performance release used 16 slots, a three-state path cap, and a 1,194,496-token KV pool. Both campaigns are retained in RESULTS.md rather than silently merging their allocations.

Resolved backends: Qwen3.8-27B/DFlash2

The selected 27B target keeps lm_head in BF16. The quantized-head selector from upstream PR #35496 is present in the base but is not the source of this checkpoint's measured speed. No separate local DeepGEMM SM120 patch is claimed; this FP8 configuration resolves MoE to Triton.

Resolved backends: Qwen3.8 Flash-Next

SM120 enablement was scoped, not indiscriminate

Several upstream checks treated SM100 as the only eligible architecture even when an underlying kernel already compiled for SM120. This runtime distinguishes three cases:

1. Explicit phase override. Automatic GDN selection used an SM100 helper and therefore did not select FlashInfer on SM120. FlashInfer 0.6.17 already supported the BF16 decode and prefill operations used here. The launcher selects only those phases with --linear-attn-decode-backend flashinfer and --linear-attn-prefill-backend flashinfer. Long prefill, exact needles, decode, and vision/reasoning qualification cover the resulting runtime.

2. Narrow source integration. RecoverSSM needed FlashInfer WY output-only state output and accepted-state recovery. FlashInfer already contained an sm_120a implementation, but SGLang did not route SM120 into it. Commit 280825c3e2 enables only the output/recovery route used by gdn_mtp_cache_mode=none and preserves Qwen4 PLE accepted-state commit ordering. Direct tests cover all accepted lengths for four draft tokens, batch sizes 1-4, mixed acceptance, positions 63/64/65, extra_buffer, prefix restoration, retraction, long continuation, and captured-graph replay.

3. QSA wrapper enabled, backend distinguished. Commit c1da0eef56 lets SM120 enter FlashInfer's page-aligned QSA decode wrapper, while 8de07a058f adopts upstream #36806's exact (12, 0) gate so SM121/GB10 is not admitted. On SM120 the wrapper resolves to XQA, an existing supported implementation; it does not select TRTLLM-Gen. No upstream SM100 throughput claim is attributed to Penny's SM120 results.

4. Incompatible gates deliberately retained. Forced TRTLLM-Gen fails with Unsupported architecture. Bypassing the guard and loading the exact BF16-Q/FP8-KV/H256/P64 cubin directly returns CUDA_ERROR_NO_BINARY_FOR_GPU (209): the binary targets SM100 and contains the sm10x_tcgen05 instruction family unavailable on SM120. Supporting it would require kernel/compiler adaptation upstream of SGLang; see TensorRT-LLM #11799 and FlashInfer #3628. Ordinary FlashInfer state-writing GDN verification also remains gated and uses Triton, while FlashInfer HyperConnection Mix remains SM100-only and falls back to the qualified persistent Triton Mix.

Other bounded corrections follow the same rule. 0f159cd545 supplies XQA's packed mask for fixed-width DFlash2 verification rather than enabling a new global attention backend. 64ecd64924 extends the fused rotary kernel to real three-axis mRoPE rather than dropping the fused path or ignoring H/W positions.

Memory recovery and automatic KV sizing

The sequence matters:

1. BF16 reduced recurrent-state bytes per slot.
2. Automatic Mamba sizing then expanded from 21 to 49 slots and consumed much of that saving.
3. The workload needed far fewer slots; 24 remained above observed demand and returned the rest of the budget to KV.
4. RecoverSSM removed the separate 1.05 GiB intermediate speculative SSM pool.
5. With no forced KV-token count, the measured KV pool increased by 78,784 tokens—from 745,600 to 824,384, or 10.6%.

The physical pool is removed in 280825c3e2, but the current kv_cache_configurator.py::_handle_max_mamba_cache estimator only omits that reserve for ReplaySSM, not RecoverSSM none. There is no separate allocator correction commit in this runtime. The qualified launcher increases --mem-fraction-static from .97 to .981, approximately the recovered GPU fraction, so automatic KV sizing can consume the physically freed memory. The 824,384-token result was measured, not hard-coded, but the estimator debt is real and documented.

Performance and qualification

These rows are observations from different model configurations and are not a single A/B benchmark.

Flash-Next final campaign — source 64ecd64924

The published campaign below used lactd with the workstation card's fan curve active and this relevant power/clock profile:

power_cap: 450.0
min_core_clock: 210
max_core_clock: 2750
gpu_clock_offsets:
  0: 1000
mem_clock_offsets:
  0: 2000

Per-stream post-first-token rates do not sum to the synchronized aggregate because each stream uses its own interval. The aggregate includes TTFT and the batch tail. These results were measured with QSA sparse decode resolving to XQA. No matched SM120 end-to-end A/B supports a percentage claim against another QSA backend.

Independent TP=2 FP8 validation

H3PO independently validated this runtime on a dual-GPU TP=2/EP=2 FP8 deployment with native MTP after removing the optional overlap-plan-stream setting. The raw third-party table, corpus description, capacity result, and scope limits are preserved in RESULTS.md.

Qwen3.8-27B/DFlash2 dated performance campaign

The cited public TP1 official-FP8/MTP3 community capture is directional rather than a strict A/B because checkpoint, runtime, speculation, power, harness, output duration, and cache configuration differ.

The current 24-slot/five-state confirmation retained exact long-context needles, reached a 96.92 reasoning score across 50,986 completion tokens, and used at most 10 of 24 Mamba entries.

27B real agentic context behavior

The dated 124-request sample covered 85,156 output tokens and inputs from 183 to 350,195 tokens. RESULTS.md preserves the context-band tables, overlap treatment, and distinction between completed-request and instantaneous telemetry measurements.

Flash-Next real agentic context behavior

A separate decontaminated 96-request window covered 100,666 output tokens. RESULTS.md preserves its context-band and aggregate tables. The first sample after a large prefill was excluded because its telemetry interval mixed prefill or idle time with decode.

Detailed definitions, complete 27B context bands, both 27B campaigns, and persistence evidence are in RESULTS.md.

HiCache/NIXL persistence

HiCache/NIXL does not cause GPU-resident decode speed. It provides disposable, reusable prefix state across GPU/host eviction and service restart:

GPU radix state
  -> page-first HiCache host RAM, kernel I/O, write-through
  -> NIXL POSIX FILE storage, io_uring + O_DIRECT

Flash-Next restart evidence:

• 63,808 of 63,864 input tokens restored; 56 recomputed.
• 489,856 of 489,879 restored; 23 recomputed.
• all three needles exact after restart.

The earlier 27B work also demonstrated 518,528 restored tokens with a six-token tail in 14.64 seconds, 60,032-token namespace reuse after restart, A→B→A page- size namespace rollback, three concurrent ~60K restores, and a 775,168-token three-request concurrent restore with only 320 tail/page-rounding tokens computed.

Two portable NIXL fixes from this project remain open upstream: PR #36520 isolates overlapping path registrations, and PR #36524 bounds bounce-backed hybrid transfers. Flash-Next adds complete PLE/GDN/QSA sibling-state persistence and load/COW ordering. MEMORY-AND-PERSISTENCE.md documents the representation namespace and actual failure boundaries.

Build and launch

There is one native build procedure and exactly two canonical launch recipes:

uv python install 3.12.13
uv venv --python 3.12.13 .venv
source .venv/bin/activate

export CUDA_HOME=/usr/local/cuda
export CC=/usr/bin/gcc-15 CXX=/usr/bin/g++-15 CUDAHOSTCXX=/usr/bin/g++-15
export MAX_JOBS=24 CMAKE_BUILD_PARALLEL_LEVEL=24
export FLASHINFER_NINJA_JOBS=24 FLASHINFER_NVCC_THREADS=4
export TORCHINDUCTOR_COMPILE_THREADS=24

uv pip install --prerelease=allow --index-strategy unsafe-best-match \
  --extra-index-url https://docs.sglang.ai/whl/cu130/ \
  --no-build-isolation -e python

Set REPO_ROOT, CACHE_BASE, NIXL_STORAGE_BASE, and the checkpoint paths, then choose one recipe:

configs/pennyroyal/serve-qwen38-27b-dflash2.sh
configs/pennyroyal/serve-flash-next.sh

Do not run them simultaneously on one GPU. BUILD.md records exact dependencies and RUN.md provides startup assertions, ordinary OpenAI-compatible smoke requests, and cold/radix/NIXL cache distinctions.

Source changes and upstream work

The cumulative history starts with the 2026-08-24 27B release, then layers the unified runtime and Flash-Next work on the same source line. The current active stack contains 24 commits above its integration base. Major groups are:

• Qwen3.8-27B/DFlash2: independent target/draft overrides, fixed-width XQA mask, NVCC host-compiler identity, request-span observability, and NIXL correctness.
• Flash-Next: Qwen4 model support, QSA/HC/PLE/native-MTP integration, SM120 QSA decode, FP8 QSA prefill, RecoverSSM, complete hybrid persistence, and mRoPE.
• v2.1 upstream sync: exact-SM120 QSA routing from #36806 and the Mamba radix ghost-node/speculative tracking correction from #35821, adapted to Penny's fused CUDA and KDA accepted-state paths.
• v2.1.1 maintenance: accumulated additive sampling penalties are applied correctly during 27B DFlash2 verification; shared HiCache JIT transfers are bound to their torch copy streams and load-back waits for in-flight forward writes for both supported models.
• Open project PRs: #36520, #36524, and #35584.
• Closed project submissions retained in runtime history: #35583; transient ragged/DSpARK PR #35586 is documented but not in the active source.
• Related upstream work: #30967, #35371, #35496, #35744, #35821, #36497, #36644, and #36806.

CHANGES.md lists every material current commit, the earlier dated release hashes, exact PR links/status/heads, affected execution paths, and test coverage.

Scope and limitations

• One RTX PRO 6000, TP=1, exact checkpoint families, and the recorded CUDA, PyTorch, FlashInfer, compiler, NIXL, and SGLang source.
• SM120 paths were qualified narrowly; unsupported-by-default is not treated as proof of incompatibility, but no architecture gate was globally deleted.
• Benchmark observations are not guarantees for another system.
• Flash-Next and 27B results remain separate; neither model's features or numbers are silently attributed to the other.
• The full published Flash-Next reasoning/tool/vision campaign ran at 64ecd64924. The v2.1.1 source at fb1216c6c4 received focused CPU/GPU DFlash, HiCache, hybrid/NIXL, and unified-load-back tests plus ordinary 64K/490K/1K/C4 and post-restart NIXL restoration on both supported models. The full reasoning and tool suites were not rerun for this narrow update.
• NIXL cleaner thresholds are whole-filesystem occupancy percentages, not an absolute directory byte quota.

See LIMITATIONS.md for measurement and reproducibility detail.

Documentation map

• BUILD.md — one native build and dependency identity.
• RUN.md — the two canonical launch configurations.
• BACKENDS.md — source-backed resolved backend and SM120 tables.
• RESULTS.md — controlled, reasoning, tools, context, agentic, and persistence measurements for both models.
• MEMORY-AND-PERSISTENCE.md — GPU allocation and hybrid state representation.
• CHANGES.md — cumulative release and upstream history.
• PROVENANCE.md — source, package, checkpoint, and tag identity.
• LIMITATIONS.md — what the evidence does and does not prove.
• jpezzulli/pennyroyal-validation — maintained public validation harness and result catalog.

Engineering history

The experimental path—including DFlash2 work and the HiCache/Mooncake/NIXL detour—is described at msoexpert.com. Mooncake is not a fallback in this runtime; it was rejected and removed.

License

The SGLang-derived source remains under Apache-2.0. Model checkpoints and third-party runtimes retain their own licenses.