feat(qwen4): support Qwen3.8 Flash Next NVFP4, vision, and Windows serving

[MaxKerkula]

Summary

• add the experimental qwen4_exp architecture used by Qwen3.8 Flash Next
• add QSA attention, KV-cache support, multimodal mRoPE, and the native vision tower
• add NVFP4 plus PLE FTW conversion and loading
• add OpenAI-compatible image_url handling for the multiprocess server
• add Windows host-bank loading, loopback TCP ZeroMQ transport, and installed-kernel reuse
• add an RTX 3090 launch script and measured benchmark artifacts

Validation

• NVIDIA RTX 3090, Windows, Qwen3.8 Flash Next NVFP4 FTW
• native 262,144-token allocation: 6.19 GiB KV cache; text and full-resolution vision passed through the OpenAI-compatible server
• 131,072-token allocation with 2,048 expert slots: 5.791 steady decode tokens/s; 16.51% expert-cache miss rate
• strict text response: FREETOKEN_SERVER_OK
• strict 2920x1944 vision response: FreeToken Desktop
• 186 focused tests passed; one existing FlashInfer-only test was deselected because the optional FlashInfer package is not available in this Windows environment
• Python compile, PowerShell parse, JSON parse, and git diff checks passed

Platform notes

• Windows uses the numerically equivalent pure-Torch router when triton_kernels is not installed. This is slower than the fused Linux path.
• The tested machine has CUDA Toolkit 12.8 and PyTorch cu130. The embedding index operation therefore uses the safe Torch fallback until a matching CUDA 13 toolkit is installed.
• CUDA graphs are disabled for this model because expert paging performs host-side work during forward execution.

[Auton-at]

I tested this PR at 856d91f049fa607711da596ba45ca78ff56c6caa with RadixArk/Qwen3.8-Flash-Next-NVFP4 on Linux and found one Linux serving blocker plus missing vision dependencies.

Environment: Linux 7.2 (x86_64), RTX 5090 / CC 12.0, Python 3.13.15, PyTorch 2.11.0+cu130, Triton 3.6.0, Transformers 5.16.1.

1. Linux first-request crash for Qwen top-k=10

The checkpoint converts and the server becomes healthy, but the first prefill crashes in triton_kernels.topk:

triton.compiler.errors.CompilationError: ...
ValueError: arange's range must be a power of 2

Qwen3.8-Flash-Next has 512 routed experts and activates 10 experts per token. freetoken.moe.fused.fused_topk() calls the external Triton top-k implementation on Linux, whose kernel uses tl.arange(0, N_EXPTS_ACT). N_EXPTS_ACT=10 is not a power of two, so compilation fails. Windows avoids this because the PR already uses the pure-Torch router fallback when triton_kernels is unavailable.

A minimal fix is to use that same numerically equivalent fallback for non-power-of-two topk values while preserving Triton for supported values:

triton_topk_supported = topk > 0 and (topk & (topk - 1)) == 0
triton_topk_installed = is_triton_kernels_installed()
if not triton_topk_installed or not triton_topk_supported:
    return _torch_fused_topk(
        gating_output, topk, renormalize, num_token_non_padded
    )

I added a CUDA regression with logits shaped [3, 512], topk=10, and both renormalization modes. Router IDs and weights match the direct Torch softmax/top-k reference exactly (rtol=0, atol=0).

2. Native vision dependencies are absent from project metadata

The Qwen vision path requires Pillow and TorchVision, but neither is declared by this PR. A clean install therefore cannot process images until they are installed separately. For the PR's Torch 2.11/CUDA 13 setup I used:

"Pillow>=11,<13"
"torchvision>=0.26,<0.27"

[tool.uv.sources]
torchvision = { index = "pytorch-cu130" }

This resolved to Pillow 12.3.0 and TorchVision 0.26.0+cu130.

After these changes, the focused suite passed 211 tests and live Linux serving passed text, strict tool calls, tool-result continuation, native 2920×1944 vision, coding execution, and a 16,023-token cross-prefill-boundary needle test. Decode measured 35.1 tok/s on the RTX 5090 profile.

I have a clean five-file patch with the fallback, dependency declarations, and regressions and can open a follow-up PR if that is preferable.

[MaxKerkula]

Thanks for the detailed Linux test and the exact reproduction. Both findings are valid and are fixed on the PR head in ad752c9:

• non-power-of-two topk values now use the numerically equivalent Torch fallback, while supported power-of-two values keep the Triton path;
• Pillow>=11,<13 and torchvision>=0.26,<0.27 are declared, with TorchVision pinned to the CUDA 13 index for uv.

I reran the focused fused-top-k tests on the current head: 3 passed, including [3, 512], topk=10, and both renormalization modes with exact router ID and weight comparison.

Your Linux RTX 5090 serving evidence is very useful. If you can retest the current head, that would confirm the fix in the environment that exposed the blocker. No separate follow-up PR is needed for these two items.

[junhewk]

RTX 4090 / Qwen3.8-Flash-Next-NVFP4 validation (256k allocation)

Tested the current PR head ad752c9970e0dc3f1b09aeec38235332149336ed with the pinned model revision RadixArk/Qwen3.8-Flash-Next-NVFP4@7b719225242aacd3dbd3f9407468c2ee9a9d2594.

Runtime verdict: the NVFP4 FTW/offload path works on an RTX 4090, including vision, reasoning parsing, structured tool calls with the correct parser, and a real near-limit 249,999-token prompt under a 262,144-token allocation. I did find three integration issues below.

Environment and launch configuration

• Ubuntu kernel 6.8.0-117, Python 3.12.11
• RTX 4090 24,564 MiB, compute capability 8.9, driver 595.71.05
• CUDA toolkit 13.0.88; PyTorch 2.11.0+cu130; Triton 3.6.0; FlashInfer 0.6.17; Transformers 5.16.1
• CUDA_VISIBLE_DEVICES=0; exact GPU UUID resolved as GPU-03d26309-f5b2-2fa8-92fc-434e0c121bd0
• --max-seq-len-override 262144 --num-tokens 262144 --max-prefill-length 8192
• --moe-backend offload --nvfp4-backend triton --expert-load serial --moe-cache-size 1024 --moe-cpu-layers 0
• --cache-type naive, QSA attention auto-selected, vision enabled, FREETOKEN_PIN_BUDGET_GB=64

FTW conversion and load

• Source snapshot: 419 files, 127 GiB, ModelOpt quant_algo=NVFP4
• Conversion: 555.9 s
• Output: 1,223 weight + 288 expert-bank tensors, 73.51 GiB FTW payload in 10 shards
• Index: quant_format=nvfp4, fingerprint e00567fb22ecb5f5, 78,930,927,616 bytes
• PLE slim index: 129 tensors / 10 shards / 51,200,267,901 bytes
• Runtime explicitly logged NVFP4 expert backend: triton and expert banks: FTW fast path
• Full 262,144-token KV allocation succeeded: K+V 6.19 GiB, 3.70 GiB free VRAM after initialization

Direct-engine smoke returned the exact requested text:

Metric	Result
FTW load	16.384 s
Cold-JIT TTFT	21.514 s
Steady decode	10.789 tok/s
GPU used by load	19.844 GiB
Cold short-run MoE miss rate	70.29%

Longer server generations stabilized around 15.2–15.6 tok/s. One performance caveat is logged: fused_topk does not support topk=10, so routing uses the numerically equivalent pure-PyTorch fallback.

Functional and context results

Check	Result
Health + /v1/models	PASS; server advertises context_length=262144
Exact text, repeated sequentially	PASS (FREETOKEN_SERVER_OK twice); warmed request 3.128 s
Reasoning parser	PASS; exact 83810205, 625 chars in reasoning_content
Vision	PASS; exact FreeToken Desktop
Tool call + tool-result follow-up	PASS with qwen3_coder; parsed get_weather({"city":"Seoul","unit":"celsius"}), then answered 23°C
8k needle	PASS; 7,995 server-counted prompt tokens, exact needle, 1,632.1 prompt tok/s
Near-limit needle	PASS; 249,999 server-counted prompt tokens, exact FREETOKEN_CONTEXT_250K, 144.35 s request, 1,731.9 prompt tok/s
Sequential state reset	PASS

The near-limit test used tokenizer-exact construction and the server's returned usage count, not a character-count estimate.

BenchLocal quality checks (v0.9.8)

Pack	pass@1	pass@3	Notes
toolcall-15	13/15 (87%)	13/15 (87%)	TC-05 selected run_code instead of create_calendar_event; TC-07 skipped two tools in a four-tool chain
instructfollow-15	14/15 (93%)	15/15 (100%)	IF-12 passed on retry
reasonmath-15	14/15 (93%)	14/15 (93%)	RM-04 concluded the constraints were inconsistent
humaneval-plus-30 scenarios 0–2	3/3 (100%)	3/3 (100%)	Docker-sandboxed partial selection

No quality response ended due to finish_reason=length; no endpoint timeout or server crash occurred.

Test/build findings

• python -m compileall: PASS
• All Python tests in PR-touched files: 211 passed
• Full non-slow suite: 3 failed, 1,417 passed, 9 skipped, 11 deselected
  • Two failures are stale call sites in tests/scheduler/test_cost_accounting_core.py: _tokenize_requests now takes (tokenize_manager, multimodal_processor, messages, logger) and returns four values, while these tests still pass three args/unpack three values. Both reproduce as TypeError: _tokenize_requests() missing 1 required positional argument: 'logger'.
  • The remaining unrelated kernel failure is test_w8a8_matches_w8a8_reference[part_rows1-True-64]: relative error 0.0102041587 vs <0.01; it reproduced 5/5 times and is outside this PR's changed files.
• CUDA 13 Docker build fails before dependency installation because docker/Dockerfile.cuda13 does COPY pyproject.toml uv.lock ..., but this PR tree has no uv.lock ("/uv.lock": not found).

Tool-parser issue

Launching with --tool-call-parser qwen produced a Qwen XML call but fed it to the JSON parser, so the response contained no structured tool_calls. The model emitted:

<function=get_weather>
<parameter=city>Seoul</parameter>
<parameter=unit>celsius</parameter>
</function>

Restarting with --tool-call-parser qwen3_coder parsed it correctly and the complete round-trip passed. auto currently appears to classify Qwen4ExpForConditionalGeneration through the generic "qwen" -> qwen25 branch, so Qwen4Exp likely needs an explicit qwen3_coder inference rule (and a parser auto-selection test).

Overall, the core 4090 NVFP4 + vision + 256k runtime path is solid in this test. The scheduler test updates, Docker lockfile mismatch, and Qwen4Exp parser auto-selection are the concrete follow-ups I found.

[junhewk]

Follow-up: profiled hybrid A/B

I profiled and tested the same checkpoint with --moe-backend hybrid after the offload run above.

ft bench bw --model qwen3.8-flash-next --formats nvfp4 --gpu 0 measured:

• CPU STREAM read: 84.5 GB/s; PCIe linear H2D: 25.2 GB/s
• Real NVFP4 CPU MoE: 63.8 GB/s; PCIe gather: 24.9 GB/s (2.56x)
• Overlapped: CPU MoE 54.0 + PCIe 23.9 GB/s
• Recommendation: hybrid, fetching 30.7% of misses over PCIe

Hybrid loaded correctly with a 15-thread AVX2 CPU MoE pool and the profiled split. A same-prompt, same-allocation direct-engine A/B over 291 output tokens produced:

Backend	Steady decode	TTFT	Miss rate	Miss handling per layer/step
offload	14.896 tok/s	3.885 s	69.13%	all misses fetched to GPU
hybrid (profiled)	12.753 tok/s	4.165 s	42.56%	1.07 fetched + 3.19 CPU-computed

So profiled hybrid was 14.4% slower than plain offload on this Threadripper Pro 5955WX + RTX 4090, despite the bandwidth microbenchmark recommending it. The likely missing cost in the selection rule is the per-layer CPU submit/synchronize and merge overhead across 48 MoE layers; kernel bandwidth alone overpredicts the benefit. This is with the full 262,144-token KV allocation and a 1,024-slot expert cache.

This corrects the performance interpretation: offload is the faster measured backend on this host/configuration, while hybrid functionality itself loads and produces correct output.

[nickmatteo]

I opened a draft review stack directly on the current #232 head: MaxKerkula#1

It adds the synthetic hardware-fit/runtime work for Qwen3.8-Flash-Next on a 12 GB Windows GPU: active NVFP4, sparse-QSA differentials, Q3 PLE, file-tiered experts, modular artifact orchestration, and the Windows JIT portability pieces. It also records #226 as a parallel upstream dependency.

This is intentionally not merge-ready: no model payload has been downloaded or converted, and real-artifact quality/runtime plus cold/async SSD behavior remain unresolved. Feedback on component boundaries and which pieces should be split upstream would be very welcome.