Qwen3.8-Flash-Next on one 24 GB GPU

Qwen3.8-Flash-Next is a 176B-parameter mixture-of-experts model (Qwen4-Exp architecture) with 6B active parameters per token. This repository is the code side of serving it on a single RTX PRO 4000 Blackwell (24 GB, sm_120) with 32 GB of host RAM: quantized with AutoRound to 2.572 bits per weight (2-bit experts with group size 128, INT8 dense projections, 16-bit routers and norms) and served through a patched SGLang. It contains the serving patch, the layered patch scripts it was built from, the kernels, the two serving assets, the measurement tools, the design documents with their outcomes and the dated engineering log.

Weights: Hugging Face Hub · Patch: sglang/qwen4exp-serving-73a255206f.patch · License: Apache-2.0 (code, LICENSE) / Qwen Community License 1.0 (weights, on the Hub)

Headline numbers

Sources, row by row: docs/logs/tiers-validate.log "streaming bench" and assets/phase1_state.json accepted.decode_all (CAMPAIGN.md:413); the same log; docs/logs/night.log "258k spike check" (CAMPAIGN.md:448); patches/kv_tiers.py docstring (CAMPAIGN.md:413); docs/logs/elastic.ctl.status (CAMPAIGN.md:320); docs/logs/tiers-validate.log (CAMPAIGN.md:396, :413).

CAMPAIGN.md:N cites line N of docs/CAMPAIGN.md, the append-only engineering log; line numbers stay valid as the log grows. The accepted state is the flag set, environment, patch list and checkpoint recorded in assets/phase1_state.json, on which the published numbers were measured. S is the number of experts per layer resident in VRAM. All measurements were taken on one machine (one RTX PRO 4000 Blackwell, 32 GB host RAM); see Reproducing the measurements.

Weights. The checkpoint (text-only, 222,856 tensors, 38,755,352,600 bytes) and the PLE n-gram table (ple/ple.f8_e4m3.bin, 51,200,245,760 bytes, plus ple/ple.json) are published on the Hugging Face Hub under the Qwen Community License 1.0: https://huggingface.co/HaberstrohSystems/Qwen3.8-Flash-Next-int2-mixed-AutoRound-24GB-SGLang. The model card there describes the quantization recipe and the serving configuration.

Quick start

Requirements: one 24 GB GPU of compute capability 12.0 (the tuned Triton configs are for the RTX PRO 4000 Blackwell), 32 GB host RAM, an NVMe disk for the PLE table, Python 3.12 (the virtualenv's lib/python3.12/site-packages path is hard-wired in scripts/serve.sh), the CUDA 13 toolkit inside the virtualenv. The serving stack that was measured: torch 2.13.0+cu130, triton 3.7.1, flash-attn 2.8.3.post1 built for sm_120, transformers 5.12.1 (pinned by SGLang at 73a255206f), SGLang at commit 73a255206f with the serving patch.

# 1. this repository and SGLang at the base commit, with the serving patch applied
mkdir -p ~/quant && cd ~/quant
git clone https://github.com/HaberstrohSystems/qwen3.8-flash-next-24gb-sglang.git
git clone https://github.com/sgl-project/sglang.git && cd sglang
git checkout 73a255206f916366c8d26d4022f82ddfb0ab558d
git apply --check ../qwen3.8-flash-next-24gb-sglang/sglang/qwen4exp-serving-73a255206f.patch
git apply         ../qwen3.8-flash-next-24gb-sglang/sglang/qwen4exp-serving-73a255206f.patch

# 2. a virtualenv with torch for CUDA 13, SGLang, the CUDA 13 compiler and FlashAttention 2. nvcc and the
#    runtime headers are installed under site-packages/nvidia/cu13/, which scripts/serve.sh puts on PATH and
#    CUDA_HOME (SGLang JIT-compiles kernels for compute_120a with that nvcc).
python3.12 -m venv ~/quant/venv-sglang && . ~/quant/venv-sglang/bin/activate
pip install torch==2.13.0 --index-url https://download.pytorch.org/whl/cu130
pip install -e python              # SGLang (the patched checkout is the current directory)
pip install nvidia-cuda-nvcc==13.3.73 nvidia-cuda-runtime==13.3.29   # compiler and headers, same minor version
$VIRTUAL_ENV/lib/python3.12/site-packages/nvidia/cu13/bin/nvcc --version   # release 13.3, V13.3.73
export CUDA_HOME=$VIRTUAL_ENV/lib/python3.12/site-packages/nvidia/cu13
export PATH=$CUDA_HOME/bin:$PATH
FLASH_ATTN_CUDA_ARCHS=120 MAX_JOBS=4 pip install --no-build-isolation --no-deps flash-attn==2.8.3.post1

# 3. the weights (about 90 GB: checkpoint shards plus ple/)
hf download HaberstrohSystems/Qwen3.8-Flash-Next-int2-mixed-AutoRound-24GB-SGLang --local-dir ~/quant/model

# 4. launch (the published configuration: full 262,144-token context, tiered KV cache, one request at a time)
cd ~/quant/qwen3.8-flash-next-24gb-sglang
SGLANG=~/quant/sglang VENV=~/quant/venv-sglang MODEL=~/quant/model PLE=~/quant/model/ple scripts/serve.sh

# 5. verify
python3 tools/bench_speed.py

The nvcc and the CUDA runtime headers inside the venv must be the same minor version (tilelang checks; docs/WRITEUP.md section 7). nvidia-cuda-nvcc depends on nvidia-cuda-runtime without a version pin, and neither SGLang nor its tilelang dependency installs the compiler, so step 2 pins both to what the measured venv carries: nvidia-cuda-nvcc 13.3.73 and nvidia-cuda-runtime 13.3.29. A system CUDA 12 cannot target compute_120a.

scripts/serve.sh waits for /health, sends one warm-up request and POST /freeze_gc; the comment block at its end explains every flag and environment variable. The server speaks the OpenAI chat API on 127.0.0.1:30000 with thinking in reasoning_content and parsed tool calls. tools/bench_speed.py prints one line per context (tokens, prefill seconds, prefill tok/s, decode tok/s, decoded tokens); expected on the measured machine: decode 54-57 tok/s at all five contexts and about 2,270 tok/s prefill at 10,001 tokens (docs/logs/tiers-validate.log lines 13-18).

Two flags to know: --cpu-offload-gb 19 is a budget for expert weights only once the patch is applied. --max-running-requests 1 --max-mamba-cache-size 1 serve one request at a time; on a 32 GB host that is the configuration that keeps memory pressure low (CAMPAIGN.md:454); higher concurrency needs more host RAM.

Pointers: the patch is in sglang/, reproduction of every headline number is described in Reproducing the measurements, the license is LICENSE.

Repository layout

What the patch does

sglang/qwen4exp-serving-73a255206f.patch, the serving patch, is the full difference between SGLang commit 73a255206f and the tree that served every number above (34 files, +4,155 / -89 lines, SHA-256 92f669b2…744bb5), produced from the served tree and verified in a clean worktree: base commit + patch reproduces all 34 files byte for byte (sglang/PATCH_NOTES.md section 2). It is the flattened result of the base patch plus every layer in patches/; the layered scripts remain useful for peeling features off one at a time. Terms used below: Gated DeltaNet (GDN) and Qwen sparse attention (QSA) are the two layer types of the model (36 + 12 of 48 layers); PLE is SGLang's name for the model's 320M-row n-gram embedding table (ple.json, --no-ple-offload-embedding); VMM is CUDA virtual memory management. Mechanisms, in the order they were built:

• 2-bit MoE path (base patch, docs/WRITEUP.md). Stock SGLang and vLLM cap moe_wna16 at 4 and 8 bits. The patch adds a 2-bit unpack to the fused MoE Triton kernel and a 2-bit loader and addresses nine findings that only surface with CPU offload on sm_120: eight correctness fixes (the worst one silent: a conv_weights view that went stale when the offloader replaced param.data) and, for the ninth, whole-layer offloading replaced by an ExpertStreamer that moves only the ten selected experts per layer per token (26 GB -> 0.31 GB of PCIe traffic per token). The 51.2 GB fp8 PLE table is memory-mapped from NVMe instead of pinned in host RAM.
• Read experts where they are (patches/ncontig_gemv.py, placement.py). Experts are re-laid after loading to an N-contiguous [E, K/16, N] int32-word layout; the batch-1 decode GEMV reads each expert in place through an int64 address table, from HBM at 320 GB/s or from pinned host memory at 51 GB/s (PCIe line rate, measured inside the kernel; CAMPAIGN.md:196-204). Residency is an address, so placement is a table write: the hottest S = 184 experts of every layer (routing-mass histogram assets/expert_freq.pt; the top 32 / 64 / 128 / 171 / 256 of 512 experts carry 37 / 53 / 73 / 82 / 93 % of the mass, CAMPAIGN.md:206-216) live on the GPU, the rest in pinned host slots.
• Breakable CUDA graphs for decode (host_fixes.py items bcg, bcg2): the mmap PLE lookup is the one eager break; capture 3.97 s / 0.12 GB (CAMPAIGN.md:102, :291).
• Elastic expert cache (gemv/row_arena.py, gemv/expert_elastic.py, patches/elastic.py). Every (layer, tensor kind) keeps its expert rows in a CUDA VMM arena in routing-mass rank order; growing gathers rows from host slots, shrinking copies the tail back and unmaps 4 MiB chunks (aligned to the 2 MiB device granularity), so VRAM really returns to the driver while every address stays fixed and captured graphs stay valid. S is a live dial through the SGLANG_MOE_ELASTIC_CTL file (tools/elastic_sweep.py).
• Lazy VMM KV cache (patches/kv_lazy.py). The KV cache is reserved as address space for 262,144 tokens and backed in 2,048-token steps as pages are handed out; at idle the backing beyond a 4,096-token floor is unmapped. A watermark keeps 1,536 MB driver-free after every commit by shrinking the elastic expert cache first; prompts whose KV cannot be backed are refused at admission (capacity = min(requested, 0.77 x profiled): 366,976 profiled -> 262,144 admitted, docs/logs/tiers-validate.log lines 3-4).
• Quantized KV cache formats for the QSA layers (patches/kv_fp8.py < kv_int8.py < kv_int4.py < kv_tiers.py). The sparse-attention kernels read bf16 only, so quantized storage gets a dequantizing read path: a fused quantize+scatter Triton write and dequantization inside the two existing gather sites (decode compaction into the FlashAttention scratch, prefix-chunk row gather). The attention kernels are untouched; scales live in the same lazy VMM arena as the payload. The default, the tiered KV cache (int8ring_int4), dual-writes every token INT8 into a ring of the last 8,192 slots and INT4 into the full-context pool; readers test ownership on the device.

The full mechanism descriptions are in docs/ELASTIC_MEMORY.md and the per-layer table in patches/README.md.

KV cache modes

Reproducing the measurements

Served proof of the published files (2026-09-03, CAMPAIGN.md entry 13:36; docs/logs/hfcheck4.log): the SHA-256 of every weight file on the Hub matches the local release copy, and that copy, launched with scripts/serve.sh (one request, int8ring_int4), answered /health after 164 s and measured decode 57.4 / 57.6 / 58.2 / 57.6 / 57.1 tok/s at 101 / 421 / 1,701 / 6,821 / 10,001 tokens of context and 2,338 tok/s prefill at 10,001 tokens (4.28 s).

All numbers were measured on one RTX PRO 4000 Blackwell (24 GB) with 32 GB host RAM against a server launched with the flag set and environment of scripts/serve.sh (the campaign harness scripts/sweep.sh passed the same set; the concurrency flags were 1 / 1 for every headline log except night4.log and elastic.ctl.status, which were recorded at restart #23 with --max-running-requests 4 --max-mamba-cache-size 8). The tools talk to 127.0.0.1:30000 and never start a server. Expected values are those of the accepted state; a different GPU, disk or host will differ.

The validation sequence a KV mode has to pass, and the thresholds, are in CONTRIBUTING.md.

Upstream status

The patch is prepared for SGLang as five reviewable commits (correctness fixes and breakable graphs under CPU offload; 2-bit moe_wna16 with expert streaming and the in-place GEMV; VMM-backed elasticity; quantized KV pools with dequant-on-gather; NGRAM speculation on PLE models). The plan, what is general and what is Qwen4-Exp-specific, and the hunks that are measurement aids and do not go upstream are in sglang/UPSTREAM.md and sglang/PATCH_NOTES.md section 5. The series itself (git am patches against the head of qwen4-main-squashed and against the served base), the RFC issue text and the five PR descriptions are in sglang/upstream/.

Hardware and software requirements

The checkpoint is auto_round:auto_gptq, bits 2, group_size 128, sym, with 2,342 extra_config entries (its config.json); 85 dense projection tensors are re-packed to INT8 g128 (scripts/requant_int8.py, CAMPAIGN.md:298-300); the PLE table is 320,001,536 rows x 160 fp8 (ple.json). The AutoRound output measures 2.572 bits per weight (docs/WRITEUP.md section 2).

The path from 2.24 to 56 tok/s

Every exact step was gated by a teacher-forced logprob oracle (450 fixed continuation tokens; same-config noise MAX 0.09 / MEAN 0.002 nats; accept MEAN <= 0.01, MAX <= 0.5 for host-side changes, MEAN <= 0.05 for kernel-class changes after a per-layer identical-input A/B, CAMPAIGN.md:92-96, :247). Approximate steps were gated by held-out NLL. Steps that did not improve the number were reverted; they are listed with their reasons in docs/HISTORY.md, and the dated sequence is in docs/TIMELINE.md.

The KV precision ladder that led to the default (all-position test, bf16 pool as reference, CAMPAIGN.md:406): noise +0.0002 / 0.059; fake INT4-G32 +0.0078 / 0.139 (rerun +0.0059 / 0.137, CAMPAIGN.md:416); real INT4-G32 +0.0088 / 0.138; fake INT2-G16 +0.3004 / 0.619 (the cliff); K/V := 0 control +3.699 / 3.923 (the metric is sharp). No valid all-position INT8-vs-bf16 figure exists; the INT8 evidence is the 512-window comparison in the KV table (see the ERRATUM note in docs/HISTORY.md).

Speculative decoding was evaluated and left opt-in: patches/ngram_ple.py makes NGRAM speculation work with the PLE embedding (linear draft chain, PLE n-gram history committed after verify) and is lossless within the decode-vs-prefill floor, but own-history n-gram drafts accept only 1.11-1.27 tokens per step, so decode lands at 22-25 tok/s with eager verify vs 56 without (CAMPAIGN.md:387). The MTP head is bf16 (4.85 GiB, docs/WRITEUP.md section 3; its experts alone 4.7 GiB, CAMPAIGN.md:379) and does not fit next to the expert cache.

Limits

• Host RAM is the binding limit. 31 offloaded layers pin ~24 GB of 32 GB; a runtime cudaHostAlloc of 4 MB took 22 s and triggered systemd-oomd; a 315 MB pinned reserve at startup killed the server at load (CAMPAIGN.md:346-347). The elastic expert cache therefore never pins after load, its floor is S = 184, and the published configuration serves one request at a time (CAMPAIGN.md:454). The server runs in a systemd-run scope with MemoryMax=30G (27G caused reclaim storms at ~55k-token prefills, CAMPAIGN.md:345). Do not exempt the scope from systemd-oomd; see the host-RAM section of docs/ELASTIC_MEMORY.md.
• fp8_e4m3 mode crashes on a 1-token prompt (CAMPAIGN.md:355); the default mode does not.
• Random text churns the PLE page cache (fresh 4 KB page per n-gram lookup); ple_random.py drops pages after bulk gathers, and prefill on random text runs at ~920 tok/s vs ~1,500 on prose at 248k (CAMPAIGN.md:408-409).
• Long prefills are O(prefix) in the model: the paged prefix-chunk kernel was built, verified and rejected on timing (6.1 ms vs 1.4 ms per head-layer at prefix 60k, CAMPAIGN.md:439; docs/KV_PAGED_PREFIX_PLAN.md "Outcome"); the slope is the QSA indexer.
• The tuned Triton configs and the placement histogram are specific to the RTX PRO 4000 Blackwell (Triton 3.7.1) and to the routing probe's workload mix; another GPU needs a tuning run, another workload may want its own histogram (assets/README.md).
• No standard-benchmark score is published for the 2-bit model; the internal quality protocol (NLL ladder, logprob oracle, needle retrieval) is what the numbers above rest on (docs/HISTORY.md).

Packaging notes on the patch (which environment-gated measurement hooks are in the serving patch, what has and has not been verified about the layered scripts) are in patches/README.md.

License

• The code in this repository (sglang/, patches/, gemv/, tools/, scripts/, docs/) is licensed under the Apache License 2.0, see LICENSE.
• The serving patch and the patch scripts contain SGLang diff context and replacement anchors; SGLang is Apache-2.0, (c) SGLang contributors.
• The quantization pipeline drives AutoRound (Apache-2.0, (c) Intel Corporation).
• The weights are not in this repository. They are published under the Qwen Community License 1.0 at https://huggingface.co/HaberstrohSystems/Qwen3.8-Flash-Next-int2-mixed-AutoRound-24GB-SGLang; base model Qwen3.8-Flash-Next (c) Qwen Team, Alibaba Group.

Citation and contact

Maximilian Roland Haberstroh, Haberstroh Systems, 2026 — info@haberstroh-systems.de. Citation metadata is in CITATION.cff (GitHub renders it as "Cite this repository"); the Hub model card carries the matching BibTeX entry.