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.

Pennyroyal v2.4.0 — maintenance fixes and faster Flash-Next prefill.

Reduced Flash-Next prefill overhead, with approximately 4–10% higher throughput observed in cold-prefill tests at 64K and 490K. Cached-prefix extension and NIXL restart restoration remain verified. Both profiles retain their full context, token pools, speculative decoding, and HiCache/NIXL configurations.

The update also corrects disconnected-request cleanup, explicit Chat reasoning effort overrides, and bare/undeclared tool-markup handling. This is an optional update, not an urgent upgrade or a new decode-speedup claim. Changes and upstream credits · Build/update instructions.

Verified models

These model configurations have been verified with HiCache and NIXL on a single RTX PRO 6000. The links point directly to the model downloads:

Flash-Next performance at a glance

Measured September 9, 2026, with v2.4.0 and FR-Spec on one RTX PRO 6000, TP1, keeping 524,288-token context, 824,384 KV tokens, and HiCache/NIXL enabled. These are dated measurements, not guaranteed speeds.

Prefill rates divide input tokens by the server's initial-prefill time; client time to first token also includes other request processing. The 4–10% gain compares that same server timing window before and after this update, not the TTFT-based rates in older tables. These are single cold-prefill observations; warm-prefill speedup was not separately measured. Decode varies with workload and acceptance, and no decode improvement is claimed. The four-request figure is combined throughput, not each response's speed. See RESULTS.md for the comparison; earlier v2.3 measurements remain there as dated history.

In real agentic use: a session on September 6 sustained 155 tok/s across nine responses of at least 1,024 output tokens, with roughly 203K–280K input context. Those responses generated 22,535 tokens; the rate excludes prefill and time between requests. Short replies are excluded from this sustained-generation figure. This is an observed session, not a controlled benchmark. Session details.

Expected compatibility with HiCache/NIXL

The standard Qwen/Qwen3.8-27B-FP8 target is expected to work with the existing 27B DFlash2 + HiCache/NIXL recipe, but has not been actively verified on the v2.3 release line. Its architecture, quantization configuration, tokenizer, and chat template match the verified 27B FP8 configuration.

Broader model support without HiCache/NIXL

Without HiCache and NIXL, many more SGLang-supported models and speculative configurations can run and may benefit from applicable performance optimizations. Potential speed improvements come from those optimizations, not automatically from disabling HiCache/NIXL.

HiCache/NIXL requires model- and speculative-method-specific cache layouts and complete state-restoration support. Different MTP or draft designs can require additional integration to save and restore target KV, draft KV, and recurrent or other model-specific state together.

The verified configurations above are therefore not an exhaustive model support list. Other models still need compatible target/draft configurations, supported kernels, sufficient memory, and appropriate launch settings.

See RUN.md for the launchers and the configuration matrix for dtypes, context, cache capacity, and backends.

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.

v2.3 — Faster Flash-Next decoding with FR-Spec

Pennyroyal v2.3 adds FR-Spec to Flash-Next, improving decode throughput while keeping 524,288-token context, 824,384 KV tokens, CUDA graphs, and HiCache/NIXL persistence. Qwen3.8-27B with DFlash2 remains supported and unchanged.

Thanks to Gabriel's gabrielolympie/sglang-flashnext-sm120 for the reduced draft-vocabulary optimization used here. Pennyroyal uses SGLang's FR-Spec support with a 65,536-token draft map generated for this release. The draft model scores fewer tokens; the target model keeps its full vocabulary, verification, and acceptance policy.

Against two v2.1.2 baseline runs on the same Flash-Next checkpoint, v2.3 measured 9.7–16.8% faster single-request decode and 4.5–7.0% higher four-request aggregate throughput. Each request generated 1,024 tokens; the comparison uses medians from six v2.3 samples and three samples per baseline run. The baselines were on separate boots, so the range reflects observed variation, not a guaranteed speedup. Cold 64K/490K prefill was essentially unchanged by FR-Spec.

The performance table above gives the v2.3 speeds directly. RESULTS.md preserves the complete baseline comparison and measurement definitions.

Reasoning, tools, vision, agent workflows, long-context continuation, and NIXL restart restoration were tested. Detailed scores and evaluator limitations are in the validation report.

v2.3 uses the same executable and dependencies as v2.1.2. Start with the FR-Spec launch guide; the non-FR Flash-Next launcher remains available.

Previous v2.1.2 — maintenance release

v2.1.2 added two correctness fixes for both supported profiles—27B FP8/DFlash2 and Flash-Next NVFP4/native MTP:

• #37962: skip token-ID synchronization when the selected group has only one GPU. This avoids an unnecessary NCCL operation and its possible late GPU-memory allocation, including during grammar-constrained JSON output.
• #37408: keep separator whitespace between streamed Qwen3 Coder tool calls from overtaking pending JSON arguments. Genuine prose and its word spaces are preserved.

That update kept launch settings and dependencies unchanged. Test coverage is in CHANGES.md.

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

Install v2.4.0 from the updated source using BUILD.md; retaining an older wheel alone does not apply these changes. Dependency versions and launch settings are unchanged. Source lineage and upstream relationships 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. The 27B model is dense; Flash-Next's MoE routing corrections do not apply to it. No separate local DeepGEMM SM120 patch is claimed.

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

The Flash-Next table above shows the v2.4.0 observations. The earlier Flash-Next campaign, including its source, clock settings, and quality results, remains in RESULTS.md as historical evidence. The separate 27B and third-party results below are not a matched A/B comparison with Flash-Next.

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 FP8/DFlash2 — August 24, 2026 campaign

Measured for the qwen38-dflash2-pro6000-20260824 release on one RTX PRO 6000 at TP1. These are retained 27B measurements, not a new v2.4.0 benchmark.

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 August 26 confirmation of the 24-slot/five-state configuration 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.

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 two supported model profiles. Flash-Next provides FR-Spec and non-FR launchers:

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-frspec.sh
# Flash-Next without FR-Spec:
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 35 runtime commits above its integration base, excluding documentation and recipe-only commits. 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.
• v2.1.2 maintenance: skip one-rank sampler synchronization (#37962) and preserve streamed tool-call framing around separator whitespace (#37408). The graph-lifetime trial is reverted and is not part of this release's runtime behavior.
• v2.3: FR-Spec reduces draft-vocabulary work for Flash-Next. The token map and launcher use the same runtime and kernels as v2.1.2.
• v2.3.1 maintenance: align paired Triton GDN verification and accepted-state recovery with ordinary decode's gate rounding, adapted from #36014.
• v2.4.0: reduce Flash-Next prefill preparation, bound short QSA extensions, correct routing dependency order, and improve shared request cancellation, explicit reasoning-effort handling, and tool-markup parsing.
• Project upstream submissions include #36520, #36524, and #35584; dated status records are in CHANGES.md.
• 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 earlier Radix 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.
• v2.1.2 received focused sampler/parser tests and both-profile JSON, incremental parallel-tool, prefill, C1/C4 decode, and NIXL restart-restore regressions. Full reasoning, tools, and vision were not repeated for that maintenance update; its historical performance tables were unchanged.
• v2.3 testing covered the full Flash-Next suite and both profiles' startup, prefill, single/concurrent decode, and NIXL restoration. Reasoning scored 98.52/100; tools completed 29 clean workflows plus one redundant read-only call. Detailed results distinguish tool execution from response-quality and evaluator limitations. No 27B speed improvement is claimed.
• 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.