Repository navigation
Port Qwen3.8-27B (EXL3 3.00bpw + DFlash2) to the native CUDA engine on Ampere sm_86 (RTX 3090) — proposal and gate #603
Description
Proposal
Port the dense Qwen3.8-27B family (model_type qwen3_5) onto the native Zig CUDA engine, serving
turboderp/Qwen3.8-27B-exl3@3.00bpw with the z-lab/Qwen3.8-27B-DFlash2 drafter, on Ampere sm_86
(RTX 3090 24 GiB). Target command:
tensorfold-native serve <qwen27b-exl3> --drafter <z-lab DFlash2>
following CONTRIBUTING: this issue comes before the arithmetic lands; each PR is one change.
Where it sits
- The native CUDA registry (
zig/src/native/cuda.zig) serves Nemotron only. The native Qwen3.8-27B
engine (1.0.3, core: backend-neutral EXL3 trellis reference decoder with an oracle fixture #496-era) is the Metal familyzig/src/families/qwen27/with its own DFlash2
runtime (dflash/);qwen35is Metal-only. - Shared device kernels already live in
zig/kernels/cuda/as device-only extractions of the Python
0.6 engine's extensions (gdn, prefill_attention, sample, qmm_group/prefill, experts*, lane_gemv),
checked byzig/tests/cuda/copies.py. The Pythonfamilies/qwen3_5/cuda/engine (branch
python-0.6) is the working reference: same model, same kernels, served on GB10. zig/src/core/exl3_format.zig(core: backend-neutral EXL3 trellis reference decoder with an oracle fixture #496) is the backend-neutral EXL3 trellis reference decoder, and
cuda: one quantized-weight interface (qlinear.zig); Nemotron's 4-bit projections route through it #587'szig/src/cuda/qlinear.zigis the one quantized-projection interface (affine4, fp8g today).
EXL3 joins as a third variant there, on the core: backend-neutral EXL3 trellis reference decoder with an oracle fixture #496 decoder's semantics.
Verified ground (a working prototype, patched Python 0.6.6 on this RTX 3090)
- The FP8 (8.9) capability floor is not needed for EXL3 serving. Stock 0.6.6 refuses sm_86
globally ("CUDA kernels need compute capability 8.9 or newer (FP8 MMA)"), but the EXL3 path's
kernels (cuda/exl3/linear.cu,experts_cb0-2.cu,cuda/kernels/gdn.cu,gdn_prefill.cu,
prefill_attention.cu,families/qwen3_5/cuda/b16.cu) contain zero FP8/cluster/TMA/wgmma
intrinsics; they compile for sm_86 and run correctly. FP8 lives only in the MLX-4bit lane matmul
(qmm*.cu,qmm_prefill8.cu) andcuda/nvfp4/*— formats this port does not load. - DFlash2 on sm_86: the official drafter's fused 4-bit path builds
tensorfold_qmm_v5(FP8) —
unusable on sm_86. The bf16 path (DFlash2fast=False,self.q4left empty so_linuses
F.linear) works and stays exact: 60/60 concurrent==alone, 10/10 drafted==serial on the patched
engine. The Mia-AiLab DFlash2-EXL3-5.0bpw drafter cannot substitute (its projections are EXL3
trellis parts the 0.6.6 loader does not read). - Memory on 24 GiB (startup estimate charges host staging against the device budget): weights
resident 12.17 GiB + 7.10 GiB staging vs ~19.2 GiB budget at the 4 GiB default reserve;
TENSORFOLD_MEMORY_RESERVE_GIB=3fixes it. Fitted windows at--context 0: 80,971 with the
drafter, 98,308 without (reserve 3). Marginal cache cost ~69.7 KB/token; the drafter costs
~17.3k tokens of window. - Python-engine perf on the same 3090 (the bar for acceptance (c)): decode 136.1/70.4/113.8/69.3
tok/s (code-sampled/chat-sampled/code-greedy/chat-greedy), cold prefill ~1,000-1,070 tok/s at
2k-8k — vs the GB10 receipt's 83.4/44.2/64.5/39.9 and 880-970.
Plan (one PR each)
- Kernels:
exl3linear (row-invariant 1-128 rows, any codebook/width, mma.m16n8k16 fragments
— plainmma.syncbf16, no FP8) +b16plain projection, extracted intozig/kernels/cuda/the
waycopies.pydocuments, with nvcc flags from the Python build. No FP8/cluster intrinsics
anywhere in the set this port loads. - Runtime:
zig/src/cuda/exl3.zig— trellis/scale upload, group scan from safetensors headers
(theexl3_format.zigsemantics), module/symbol binding, decode/prompt dispatch behind
qlinear.zig'sWeightunion. Host tests on theexl3_oracle.binfixture. - Family:
zig/src/families/qwen27/cuda/— decode rounds, GDN tree/replay (gdn.cu symbols),
attention, prompt path on the gdn_prefill/prefill_attention kernels, on the lane core's backend
interface (core/lanes/backend.zig), with DFlash2 as the lane drafter (bf16 projections through
cuBLASLt bf16 GEMM — the provenfast=Falsepath, no q4 packing). Registry entry in
zig/src/native/cuda.zigwithformats = exl3(the MLX 4-bit lane stays off for sm_86 until its
FP8 floor question is settled separately). - Gate: the three exactness contracts (drafted ==
\"draft\": falsetoken-for-token; resumed ==
fresh; concurrent == solo viatools/bench_concurrent.py --alone --serial), token_sha at multiple
context lengths; row-exact oracle before speed;lean_check.pyzero problems; then receipts
(bench_openai/bench_concurrent/prefill_cold) with context ceilings and memory admission documented.
The sm_86 floor decision to record
Format-conditional, never lowered for FP8 formats. The engine refuses a checkpoint whose formats
need arithmetic the device cannot run (FP8 lane/nvfp4 on <8.9/12.x, as the __CUDA_ARCH__ guards
already do inside the kernels), and admits EXL3 + bf16 on 8.6. The refusal message names the missing
capability and the format that needs it. This mirrors what the Python 0.6.6 prototype proved: the
global 8.9 gate is a build-level convenience there, not a property of the EXL3 arithmetic.
Hardware and toolchain
RTX 3090 24 GiB (sm_86), Threadripper 3970X, 128 GB RAM, CachyOS, driver 595.71.05, power limit
settable 225/350 W (I will report which was used — see #592 for the 225 vs 350 W comparison on this
same card). Builds in containers (nvidia/cuda:13.0-devel-ubuntu24.04, nvcc target sm_86 via
-Dsm=86, the Zig pinned at 0.17.0 by .zig-version). The Python 0.6.6 engine on the same card is
the token oracle during bring-up.
Acceptance
a) tensorfold-native serve loads EXL3 27B + DFlash2 on sm_86, greedy and sampled, drafting on.
b) The three exactness gates pass with token_sha evidence at multiple context lengths.
c) Decode at or above the Python-engine numbers above on the same card, or a written bandwidth-vs-
compute analysis of why not.
d) A receipt: context ceilings, memory admission, and the sm_86 floor decision.
Activity
ashhart commentedon Oct 10, 2026
Thank you, @briepala, for asking before the arithmetic lands, and for the 3090 prototype. Timing is good: #608 arrived today with the qwen27 family's engine for NVIDIA GPUs, every stream drafted, so that's the base for the 27B on CUDA rather than a second port. Once it's in, your pieces fit as follow-ups: EXL3 as a third format behind #587's
qlinearinterface on the #496 decoder's semantics, then sm_86 kernels for the paths EXL3 needs. Each needs an FP64 reference per projection class and the usual contracts (drafted equal to"draft": false, resumed equal to fresh, concurrent equal to solo), measured against llama.cpp on the same card. Let's keep this open as the tracker for the 3090 and EXL3 part.alexderz commentedon Oct 14, 2026
A second RTX 3090 receipt for this tracker, in case it helps the sm_86 work. It is from TensorFold 0.6.6 (the Python line) with a
patch, and I know that line is frozen, so I'm offering numbers rather than a pull request.
Setup. RTX 3090 24 GB, default power limit, Linux, TensorFold v0.6.6 (
cb2ebf05), torch 2.14.1, triton 3.8.0,CUDA toolkit 13.0.3.
TensorFold/Qwen3.8-27B-MLX-4bit@22d8d538withz-lab/Qwen3.8-27B-DFlash2@50307d4c,16k context, thinking off, temperature 0, 400 tokens, 5 prompts per task, 3 rounds with a fresh server each round.
Patch. The same split as #311: floor at 8.6, FP8 named as an 8.9 level, the one FP8-only 4-bit prefill file
replaced below 8.9 by a stub under its own extension name,
--prefill-fp8refused at startup on 8.6, NVFP4 and FP8checkpoints still 8.9.
Decode, tok/s after the first token (mean, range over rounds):
reserve and 73,308 with
TENSORFOLD_MEMORY_RESERVE_GIB=2. At 73k, long prompts logged CUDA OOM warnings but everyrequest still answered correctly.
test_qwen27_prefill15 passed / 12 skipped (FP8 prompts).Five host-test failures are unrelated to the patch: three expect the older torch build-lock behaviour, two call
Flash Next, which needs SM 12.0.
For the native port, the window looks like the constraint on a 24 GB card more than the speed; #799's opt-in 4-bit
attention cache would matter here. I'm happy to run the native build on this 3090 against the gate once there is a
branch to try.
ashhart commentedon Oct 16, 2026
Thank you, @briepala, for the RTX 3090 proposal, and @alexderz for the second receipt, including the 17,744 MiB footprint and the long-window limits.
The selected native Qwen CUDA base is now #836. #849 and #859 add dense mixed widths, while EXL3 is coordinated in #832. The image integration in #885 also belongs with that base. Your patched Python run is useful comparison evidence; the native 3090 gate still needs to run on the combined implementation.
Please share the 15 prompt fixtures from the receipt so the native comparison can use the same inputs, checkpoint and DFlash2 revision.
Posted by Ash's AI Assistant Synthena
alexderz commentedon Oct 16, 2026
Hi Synthena — the 15 prompt fixtures from the second 3090 receipt are in the public harness repo. No replacement harness.
Existing runner and all 15 original prompts: https://github.com/alexderz/agentic-model-testing/blob/0b9187ec944aca3df8a7703be469acff0477e374/harness/tools/bench_tasks.py
Receipt: https://github.com/alexderz/agentic-model-testing/blob/0b9187ec944aca3df8a7703be469acff0477e374/harness/engines/tensorfold/README.md
Start here: https://github.com/alexderz/agentic-model-testing/blob/d8f631e71887a9820a4229c7f390a99a6a8d88bc/docs/repro/tensorfold-603/README-agent-start-here.md
Fixtures: https://github.com/alexderz/agentic-model-testing/blob/d8f631e71887a9820a4229c7f390a99a6a8d88bc/docs/repro/tensorfold-603/fixtures.json
Offline checker: https://github.com/alexderz/agentic-model-testing/blob/d8f631e71887a9820a4229c7f390a99a6a8d88bc/docs/repro/tensorfold-603/verify_fixtures.py
SM86 patch the existing build already references: https://github.com/alexderz/agentic-model-testing/blob/d8f631e71887a9820a4229c7f390a99a6a8d88bc/harness/engines/tensorfold/patches/0001-sm86-floor-and-fp8-stub.patch
Pins:
cb2ebf0540f42604e2759b2ddef497861e928248, SM86-patched Python engine, torch 2.14.1, triton 3.8.0, CUDA toolkit 13.0.3With that server already up:
The original three rounds interleaved engines and started a fresh server for each arm. Synthetic prompts and configuration only. The native 3090 comparison still has to be run on the combined implementation; these inputs are for that comparison, not evidence it passed.
CanadaApollo6 commentedon Oct 16, 2026
Thank you, @alexderz, yes please. #836 is the Qwen CUDA base now (#856 adds the MoE on it), and it hasn't run on Ampere yet. I'm bringing it up on my RTX 3090 first. When there's a branch that builds and serves on sm_86, I'll post it here, and a second 3090 run on your 15 fixtures would make a good cross-check. The 4-bit attention cache from #799 was accepted as a value of
--kv-cache, off by default, and will follow on that base.alexderz commentedon Oct 16, 2026
Second RTX 3090 cross-check (24 GB, sm_86, Linux). No source change on either pull request.
#836
26e9b686, Zig 0.17.0, CUDA 13.0 nvcc,zig build native install -Dsm=86. The build succeeded and the fatbins are sm_86.tf-cuda-testinfo and smoke passed.tensorfold-native1.0.12 then servedTensorFold/Qwen3.8-27B-MLX-4bit@22d8d538154e0e5b6f8dcb1fcc74a60a0798104cwithz-lab/Qwen3.8-27B-DFlash2@50307d4c4cde6860d4eee73e2547cd786fe8e8a4, context 16384, one stream, greedy, thinking off, drafts on. Load was 25 s, weights 16.01 GiB.The 15 public fixtures from
alexderz/agentic-model-testingharness/tools/bench_tasks.pyat0b9187ec(code, prose, structured; temperature 0; 400 tokens) all returned text. Decode after the first token, mean of five: code 150.1 tok/s, prose 76.1, structured 163.0. Fourteen runs hit 400 tokens; one structured run stopped at 343. Time to first token was 0.20–0.24 s. This is not the old patched-Python receipt.#856
3d6f96ee(still the pre-1.0.12 tip), same toolchain and-Dsm=86. The build succeeded.tf-cuda-testsmoke and moe passed; the SwiGLU pairs were within 9.8e-7 of the fp64 reference.tensorfold-native1.0.11 servedTensorFold/Qwen3.6-35B-A3B-MLX-4bit-MTP@f84b054c5677b6e59bdb969ae61915b0983d68bbwith its MTP drafter, context 4096, greedy. Load was 19 s, weights 19.09 GiB, and one short completion came back. I did not run the 15 fixtures on this MoE model.No Ampere-only patch was required to build or serve either head.
alexderz commentedon Oct 16, 2026
Follow-up on the MoE head, same card and the same 15 public fixtures.
#856 at 3d6f96e, TensorFold/Qwen3.6-35B-A3B-MLX-4bit-MTP@f84b054c5677b6e59bdb969ae61915b0983d68bb, its own MTP drafter, context 8192, one stream, greedy, thinking off. Load 19 s, weights 19.10 GiB. bench_tasks.py at 0b9187ec, temperature 0, 400 tokens: all 15 returned text. Decode after the first token, mean of five: code 216.3 tok/s, prose 152.4, structured 243.4. Thirteen runs hit 400 tokens; one code run stopped at 362 and one structured run at 322. Time to first token was 0.05-0.07 s.
Still no source change.
ashhart commentedon Oct 17, 2026
Thank you, @briepala, and @alexderz for the native RTX3090 cross-check and the 15 original public fixtures with exact checkpoint pins.
We checked the supplied follow-up against the chosen #836 base. The new 26e9b68 run is native, unlike the earlier patched-Python arm; it returned all 15 requests, with one natural stop at 343 tokens. That is useful new Ampere evidence.
We will retain those fixtures for the final Qwen CUDA landing, including full-token draft/cache/concurrency and owned-image checks. The requested dense EXL3 path remains with #832 and is separate from these MLX-4-bit results or #972's MoE width. Thanks as well for the #856 MoE run at 3d6f96e (code 216.3, prose 152.4, structured 243.4 tok/s, all 15 fixtures answered); we will keep that receipt for #856 on the same base. Please keep the native logs and original receipt pinned for that integrated comparison.
Posted by Ash's AI Assistant Synthena