Skip to content

Qwen3.8-27B on CUDA: the qwen27 family's engine for NVIDIA GPUs, every stream drafted - #608

Closed
philip-pentatonic wants to merge 3 commits into
ashhart:mainfrom
machinegenieorg:see-4347-qwen27-cuda
Closed

philip-pentatonic wants to merge 3 commits into
ashhart:mainfrom
machinegenieorg:see-4347-qwen27-cuda

Conversation

@philip-pentatonic

@philip-pentatonic philip-pentatonic commented Oct 10, 2026 •

Copy link
Copy Markdown
Contributor

What this changes

Adds a CUDA engine for Qwen3.8-27B to the qwen27 family (zig/src/families/qwen27/cuda), beside its Metal engine: a qwen3_5 checkpoint with the 27B's shapes opens it in tensorfold-native. It runs the Python 0.6.6 engine's forward with up to 16 concurrent streams, each drafted by DFlash2, every token through the lane rounds.

Three commits:

  • Lane core: requests that arrive together share their prompt passes (optional prefill_many on the backend, addStreams in the core; other backends unchanged).
  • Nemotron: the checkpoint Source copies to a Target, so the 27B's loader shares it.
  • The 27B's CUDA engine, its CLI check (tf-qwen27-cuda) and its capture scripts (zig/tests/cuda/qwen27).

Receipt

  • Environment: this branch on main at 9041e20 (TensorFold 1.0.5 and later); the speed and memory receipts below were measured on its earlier base, e7b4a58. Rebasing changed two things in the 27B's code: its weight loader takes the affine-4 repack kernels from affine4_pack (cuda: the affine-4 repack beside its format (qlinear.zig, affine4_pack.cu); Nemotron's loader packs through it #602), and the server passes 1.0.5's --slide fields only to families that have them. After the rebase the checks were re-run on the GB10: 4 of 4 serial replies, 54 of 54 windows, every drafted shared-round reply equal, and bench_concurrent.py --alone --serial at 1 and 4 streams, greedy and sampled, every reply equal. Measured on one DGX Spark (GB10, driver 580.173.02, CUDA 13.3.73); Python 0.6.6 for comparison (tensorfold:v0.6.6: PyTorch 2.13.0a0+9186a08b2c.nv26.07, Triton 3.7.1); Vontra/Qwen3.8-27B-MLX-4bit (70ae7fac), drafter z-lab/Qwen3.8-27B-DFlash2 (50307d4c); both servers --parallel 16 --context 32768 for decode, --parallel 1 --context 70000 for prefill. The Triton glue set was captured on the same GB10 from 0.6.6 (tools/zig/triton_aot_manifest.py, aot_pack.py). The engine was also checked on an RTX PRO 6000 (sm_120).

  • Exactness, tools/bench_concurrent.py --alone --serial at 1, 4 and 16 streams, code and chat, temperature 0 and 1.0: 12 of 12 cells, every reply equal to its solo run and to its "draft": false decode, none failed. Against captures of 0.6.6 (tf-qwen27-cuda): 4 of 4 serial replies equal token for token (a 4722-token prompt among them); 54 of 54 verify windows (chains of 1 to 16 rows, trees, shared rounds of 2 to 16 streams and up to 256 rows, every alignment of the packed inputs) equal in every row's logit bytes; 4 of 4 drafted replies and 6 of 6 drafted shared-round replies equal 0.6.6's, with each stream's round count equal.

  • Decode, tools/bench_openai.py (256 tokens, 3 reps, median tok/s), 0.6.6 then this branch:

    prompt temperature 0.6.6 this branch
    fibonacci-raw 0 69.2 69.5
    gpu-chat-no-think 0 51.4 49.9
    fibonacci-raw 0.8 83.1 80.2
    gpu-chat-no-think 0.8 46.6 46.8

    tools/bench_concurrent.py, aggregate tok/s (steady in brackets), 0.6.6 then this branch:

    prompt, temperature 4 streams 16 streams
    code, 0 202.1 (195.8) / 199.3 (201.0) 386.2 (386.5) / 373.6 (373.8)
    chat, 0 148.3 (147.2) / 141.0 (145.9) 303.6 (298.6) / 278.1 (287.4)
    code, 1.0 205.8 (244.7) / 209.9 (231.7) 398.1 (407.2) / 380.3 (438.3)
    chat, 1.0 128.1 (142.6) / 129.5 (146.7) 279.9 (290.6) / 265.5 (288.2)

    The slowest first token at 16 streams is 0.66 to 0.86 s here against 0.21 to 0.26 s on 0.6.6.

  • Prompt speed, tools/prefill_cold.py (3 reps, median time to first token), 0.6.6 then this branch: 2k 1.86 / 1.50 s, 8k 6.51 / 5.85 s, 16k 12.71 / 12.21 s, 32k 26.86 / 26.11 s, 64k 57.91 / 57.68 s.

  • Memory, the host's used memory at peak (GB10, unified; about 30 GiB of it held by other processes before each server started): decode run 56 GiB on 0.6.6, 54 GiB here; prefill run to 64k 76 GiB on 0.6.6, 56 GiB here.

  • Tests: zig build test on main 9041e20 plus this branch, 39 of 39 steps, every test passed, none failed or skipped (Linux host). tools/zig/lean_check.py reports the same 542 problems as main, none in changed lines. zig fmt --check build.zig zig passes.

  • Not run, and why: Metal (no change there, and no Mac in this test); prefill at 128k and the native 262k window (the receipts ran at a 70000-token window on a shared Spark); vLLM on the same Spark (not installed there); two ranks (one GPU only). The first-token gap at 16 streams and a drafter rounding difference are follow-ups: four of six streams' DFlash2 rows equal 0.6.6's through every layer; in the other two a few context values differ, which changes the order of low tree nodes but not any reply or round count.

Checklist

  • Every token still goes through the lane rounds. No serial path, and nothing that needs drafts off.
  • Drafted output equals "draft": false, a resumed prompt equals a fresh one, and concurrent equals solo. (No prompt reuse in this engine yet, so no resumed path.)
  • No precision traded for speed. The 0.6.6 kernels and their sums, byte for byte.
  • Prompt processing is no slower than the last release.
  • The description says which platforms I ran, Metal M1 to M5 and CUDA, and which I could not.
  • New tests fail before the change, pass after it, and skip cleanly without their dependency.
  • Comments and docstrings are one line. No measurements or history in the source.
  • No personal data, machine names, internal hosts or local paths. No AI attribution lines.

The lane host admits every queued request at once: a lone one still goes to
the backend's driver, the rest to the core's addStreams (lanes/admit.zig).
It runs the backend's optional prefill_many when there is more than one,
draws each stream's first token, then asks for every stream's first drafts
in one request, as multi.py's _fill_batch admits a batch. addStream is its
one-stream case, with the same calls in the same order; a backend without
prefill_many prefills each stream as before.
…ilies share it

Source needed Nemotron's kernel handles only for the driver, the stream and
whether the card is discrete; it now takes those as a Target (Target.of(ops)
for Nemotron's own loader). cuda.zig exports the source module and the torch
ops for the 27B's CUDA engine.
…y stream drafted

A qwen3_5 checkpoint with the 27B's shapes opens zig/src/families/qwen27/cuda
in tensorfold-native, beside the family's Metal engine. It reads the family's
config.zig and refuses other shapes (cuda/shape.zig).

The forward is the Python 0.6.6 engine's: its CUDA kernels (qmm_group,
qmm_prefill at prompt tile 9, the GDN tree, replay and prompt chain,
prefill attention) built from zig/kernels/cuda, and its Triton glue from a
set captured on the chip (TENSORFOLD_CUDA_KERNELS). Verify rounds take every
stream's window together, up to 16 streams of 16 rows, as multi_tree_forward
lays them out, and commit only each stream's accepted path.

DFlash2 (--drafter, else the pulled z-lab/Qwen3.8-27B-DFlash2) drafts every
stream in one batch: each stream's kept taps absorbed together, every block
in one forward, a best-first tree per stream, or a copied continuation where
the context repeats. The lane core verifies the trees and trims them by each
node's chance against a cost curve timed at load. Prompts that arrive
together share one prefill pass (prefill_rows).

Checked against captures of the Python engine (zig/tests/cuda/qwen27), on an
RTX PRO 6000 and on one DGX Spark (GB10): serial replies equal token for
token (a 4722-token prompt among them); 54 verify windows (chains, trees,
shared rounds of 2 to 16 streams and up to 256 rows, every alignment of the
packed inputs) equal row for row in their logits' bytes; drafted replies and
each stream's rounds equal Python's; served replies of 1, 4 and 16 concurrent
requests, greedy and sampled, equal the same request alone and its
"draft": false decode.
@jregan-beasley-smc

Copy link
Copy Markdown

What this changes

I investigated the first-token follow-up on 99bcc5b29805e29a4baf2cd4beade0f7fe1a5fe8. A tested partial fix is available as a cherry-pickable commit, with validation material on the same branch.

Qwen CUDA currently prepares the initial speculative draft before handing the already-selected first token to the HTTP stream. The patch commits and delivers that token first, then drafts for unfinished streams. The capability defaults off for other backends. It preserves the initial draft context, cancellation, terminal-token handling and job ownership. No kernels, weight precision or sampling arithmetic change.

This builds on Philip's Qwen CUDA engine in this PR and the existing lane core. The engine implementation and its Python parity work remain his contribution.

Receipt

  • Single DGX Spark / GB10, native Zig 0.17 and CUDA, the same pinned Qwen3.8-27B MLX4 checkpoint (70ae7fac) and DFlash2 (50307d4c) used here. Both variants use identical CUDA/Triton kernel assets. --parallel 16 --context 32768, prompt cache zero, thinking off, 128 output tokens.
  • Two fresh base→candidate pairs. Each cell has a first-use burst and three warm bursts, with drafted/plain solo references. Normal builds cover coding at 1/4/16 streams and sampled chat at 4 streams; coding at 1/16 is repeated in the second pair. All 372 requests passed completion, cold-cache, length and full-output/token-fingerprint checks.
  • With normal calibration, sixteen-stream coding's warm median TTFT was 635→598 ms and 630→580 ms. Median full-batch times were 5.691→5.730 s and 5.704→5.724 s. Four-stream coding was essentially flat; sampled chat and first-use timings were variable.
  • One normal single-stream pair took about 7.2% longer to finish and used 26 verification rounds rather than 24; the repeat used 24 on both sides and was flat. I retained that result and ran a separate diagnostic control with identical initial planning-cost tables across all four processes. Normal GPU warmup and online planner adaptation remained enabled. The diagnostic override is not in the fix.

The calibration control produced these warm medians:

Streams / pair Median TTFT, base → fix Full response/batch time, base → fix
1 / A 126.9 → 120.8 ms 2.072 → 2.081 s
1 / B 126.8 → 120.2 ms 2.082 → 2.083 s
16 / A 637.2 → 598.4 ms 5.741 → 5.729 s
16 / B 644.9 → 592.2 ms 5.731 → 5.757 s

All 288 control requests passed the same output checks. Both single-stream arms used 24 rounds; some sixteen-stream speculative work still varied. This supports a warm first-token responsiveness improvement, not an overall throughput gain. These median figures are not the slowest-first-token metric in the PR description, and they do not resolve the full Python-versus-Zig gap. The evidence is limited to short synthetic prompts and two same-order A/B pairs, not long-context or general quality validation.

Linux CUDA-host tests: baseline 284/284, candidate 287/287. Independent host checks pass eight cases on the fix; the original passes the two controls and fails the six desired early-delivery/edge-case checks. Source review, formatting and output-parity checks pass. The original service was restored after testing.

Authored by Jonathan Regan-Beasley in a private capacity. This contribution does not represent the views or positions of any organisation.

@ashhart

ashhart commented Oct 11, 2026

Copy link
Copy Markdown
Owner

Thank you, @philip-pentatonic. This is a serious piece of work: exact against 0.6.6 in every row we looked at (54/54 windows, 12/12 concurrency cells) and level with 0.6.6 on decode. Two things keep it out of 1.0.6. Both are rules of ours, not problems with your engine:

  • Its glue kernels come from a Triton set captured from the Python engine (TENSORFOLD_CUDA_KERNELS). The native release has to run a model with nothing recorded by Python, so the archive can't serve the 27B on its own. Nemotron's CUDA path shows the shape we need: its own glue kernels in zig/kernels/cuda/nemotron_*.cu, each checked against a host reference (glue_ref.zig). The 27B needs the same for each kernel it takes from the Triton set now.
  • zig/tests/cuda/qwen27/capture*.py import the Python engine. Since 10 Oct no tooling that imports the engine lands; the captures can stay as fixture files, recorded once outside the tree.

If you'd like to take the glue kernels on, we'll gate each one against your captures on a GB10 and land the engine with them. If you'd rather we write them, we'll do it on your branch with your commits kept under your name.

@jregan-beasley-smc

jregan-beasley-smc commented Oct 11, 2026 •

Copy link
Copy Markdown

What this changes

I've opened a focused follow-up against this branch: perf(qwen27 cuda): reduce TTFT with prompt caching and state reuse. It combines the earlier first-token handoff fix with opt-in exact-prompt caching, target/drafter allocation reuse, raw-logit reuse and guarded omission of redundant initialization clears.

This builds on Philip's engine and parity work here. The PR leads with the repeated short-prompt TTFT gap and includes a fresh matched stock PR608 / Python / candidate comparison on one DGX Spark.

Receipt

Using this PR's maximum-TTFT statistic over three primed bursts, sixteen-stream chat improves from 850.7 ms stock / 247.4 ms Python to 204.8 ms; coding improves from 689.7 / 299.0 ms to 255.3 ms. Warm single-stream TTFT is 3.4–3.8 ms, versus 128.3–135.5 ms stock / 111.3–114.2 ms Python. Warm sixteen-stream medians are 126.2 ms coding / 126.8 ms chat, versus 219.5 / 233.0 ms Python.

Two fresh process orders cover 624 requests, all with exact paired input/output and token-fingerprint agreement. The prompts/settings and TTFT boundary match the original comparison; the uncached first-use request is extended from 16 to 256 output tokens so cold TTFT, decode and full-response time are also retained. The portable receipt and offline checker preserve the measurements and complete source/model/build pins.

The trade-offs are explicit: uncached chat TTFT is 162.0 ms versus 131.4 stock / 123.5 Python, with full response 0.7% / 2.5% longer. Warm sixteen-stream chat finishes 1.35% later than Python but 4.16% earlier than stock; per-stream client decode is 1.3% below Python. This is a large first-token responsiveness improvement for identical cached short prompts, with roughly similar decode performance. Native restores N prompt tokens versus Python's N−1, and speculative work differs; general-prefix reuse and general quality are not claimed.

All switches remain default-off for review. Ash's native-glue and capture-tooling release requirements remain separate. The original production service was restored after testing.

Authored by Jonathan Regan-Beasley in a private capacity. This contribution does not represent the views or positions of any organisation.

@jschmied

Copy link
Copy Markdown
Contributor

On the first point, the glue kernels: I have generic versions of the 27B's set on a branch, written as shared zig/kernels/cuda modules rather than qwen27 files, so Nemotron, Kolibri and the 27B can take the same ones. Each has a float64 host reference in core/, a tf-cuda-test command, bits checked equal at 1, 3, 17 and 64 rows, and a deliberate mutation it catches. All were tested at the 27B's shapes.

Triton kernel replacement
_add_rmsnorm norms.addRmsH (bf16 residual, bf16 weights, 64-value group sums)
_gated_norm norms.gatedRms
_gate_mul, _swiglu, _group_sums act_rows
_embed embed_rows (MLX affine-4, bf16)
_attn_prep qk_rows.rowsFrom (strided [q | gate] rows, rotary 64 of 256)
_gdn_pre gdn_prep (each row's conv window, so trees work)
_paths, _tail, _merge tree_attention at head dim 256

Branch: https://github.com/jschmied/TensorFold/tree/cuda-qwen-glue (stacked on the shared-layer pieces claimed in #548).

Two gaps. I have not run them inside this engine or against your captures, because I don't have the 27B on my GB10. And _shared (a stream's committed keys read once for all of its draft rows) has no counterpart yet; the tree kernel reads them per row, which is exact but slower at long context with wide trees.

@philip-pentatonic, if they're useful, take them into your branch, or tell me and I'll wire them into qwen27/cuda there. If you or @ashhart would rather write your own, say so and I'll leave it.

BobClawblaw added a commit to BobClawblaw/TensorFold that referenced this pull request Oct 12, 2026
Conflicts resolved toward main's admission (open.zig, chunked prompt
passes); the PR's shared prefill batching (admit.zig, addStreams) is
dropped, prefill_many stays as an unused optional vtable entry. Living
Weights setup in native/cuda.zig is guarded by @hasDecl(F, "slide").
@BobClawblaw

Copy link
Copy Markdown
Contributor

@jschmied I ran your cuda-qwen-glue branch (at 0c2d20f3) inside this engine against the Python captures, on an RTX 5090 (sm_120). It's merged onto #726, which is this PR plus image input and an fp8 KV cache. The harness is on https://github.com/BobClawblaw/TensorFold/tree/qwen27-glue-gate:

  • TF_QWEN27_GLUE sends the Triton launches it names to your modules.
  • tf-qwen27-cuda glue-diff runs a native kernel and the Triton kernel it replaces on the same random rows.

All your synthetic tests pass on sm_120: norms, qk-rows, tree, block, gdn-prep, act-rows, window, dense-rows and moe-glue.

Your kernels as they are, one at a time. The baseline is all Triton: serial 4/4, windows 54/54 cases with 0 rows differing, drafted 4/4, and multi with 0 replies differing. The baseline's multi run already has 112 drafter windows that differ from Python's while every reply is equal, so the drafter kernels' counts are shown against 112.

Triton kernel → native serial windows: rows that differ drafted multi: replies that differ
_embed → embed_rows.affine4 4/4 0 4/4 0
_group_sums → act_rows.groupSums 4/4 0 4/4 0
_paths → tree_attention.paths 4/4 0 4/4 0
_dconv_kernel → act_rows.dconv 4/4 0 4/4 0
_block_attention → block_attention 4/4 0 4/4 0 (118 drafter windows vs 112)
_prep_kernel → qk_rows.rows 4/4 0 4/4 0 (127 vs 112)
_add_rmsnorm → norms.addRmsH 3/4 2,049 3/4 2
_gdn_pre → gdn_prep 3/4 2,049 3/4 3
_gated_norm → norms.gatedRms 3/4 2,049 3/4 3
_swiglu → act_rows.swiglu 3/4 2,049 3/4 3
_attn_prep → qk_rows.rowsFrom 3/4 2,049 3/4 3
_gate_mul → act_rows.outputGate 3/4 1,409 3/4 2
_shared/_tail/_merge → tree_attention 3/4 2,049 3/4 3

The failing replies change at a near-tie, for example token 38 of one serial prompt. Every native output is within about 1 ulp of Triton's. Here is where the differences come from, measured with glue-diff (64 rows at the 27B's shapes):

  1. RoPE input rounding. _attn_prep and DFlash2's _prep_kernel round the normed q and k to bf16 before the rotation ((x * inv * w).to(bf16).to(f32)); qk_rows keeps them in fp32. About 8% of _attn_prep's elements and 30% of _prep_kernel's differ.

    • Rounding y and yp in both tf_qk_rows and tf_qk_rows_w leaves 4 or 5 elements in 400k.
    • Your angles from inv already match our rotary tables exactly.
    • The qk-rows float64 reference would need the same rounding; with the change it fails.
  2. exp, log and sigmoid. On this Triton build:

    • tl.exp compiles to ex2.approx, the same as __expf;
    • tl.log is the precise logf;
    • tl.sigmoid divides with div.full, the same as __fdividef.

    With those in gdn_prep, g goes from 68% of values differing to bit-identical, and beta from 31%. With the same sigmoid, swiglu's output becomes identical.

  3. The fused 64-value sums follow each Triton kernel's layout.

    • _gate_mul (2 warps over 256 values) and _gated_norm (1 warp over 128) sum 4 values a thread, then 16 lanes from the widest offset down.
    • group_sum8 (8 values a thread, 8 lanes) is the order of _group_sums. It also matches _swiglu in all but 7 of 17,408 groups; I haven't pinned down those 7.
    • A 4x2 version of group_sum8 cuts the differing sums from 44 of 6,144 to 1 for _gate_mul, and from 37 to 1 for _gated_norm.
  4. What's left. Single elements 1 ulp off, from the norms' sum of squares and rsqrtf vs Triton's 1.0 / tl.sqrt(...) (add_rmsnorm, gated_norm, gdn_prep's q and k), and from the attention kernels' accumulation. About 14% of block attention's outputs are 1 ulp off.

The same gate with changes 1–3 applied (experimental commit 74ba9d73 on that branch):

  • _gate_mul now passes everything: 0 window rows differ.
  • _prep_kernel stays exact on serial and drafted, with 124 drafter windows vs 112.
  • _swiglu drops from 2,049 differing rows to 257, but still changes replies.
  • _attn_prep, _gdn_pre and _gated_norm still differ: over 64 layers, one element 1 ulp off is enough to change the logits.

So today, bit-exact against the captures are _embed, _group_sums, _paths and _dconv_kernel, plus _gate_mul with the sum-order change. Not covered here: _shared8/_tail8 (#726's fp8 cache) and _append, which is captured but never launched in this engine.

@jschmied

Copy link
Copy Markdown
Contributor

@BobClawblaw, thank you, that harness is exactly what was missing. I reproduced it on a GB10 (sm_121) against our own
captured set: the same counts, 5 of 26 outputs bit-identical.

Reading the captured kernels' PTX and TTGIR, the differences that remain come in two kinds:

  1. Arithmetic, the same in every kernel:

    • exp is ex2.approx of x * log2e;
    • sigmoid and / are div.full;
    • 1 / sqrt is sqrt.approx.ftz then div.full, and DFlash2's rsqrt is rsqrt.approx.ftz;
    • normed q and k are rounded to bf16 before RoPE;
    • in a*b ± c*d, LLVM fuses the first product and rounds the second.
  2. Reduction trees, one per kernel. Each sum follows that kernel's blocked layout:

    • _add_rmsnorm: 8 warps × 8 values a thread per pass, then 3 shuffle steps across the warps;
    • _attn_prep: 2 warps × 4;
    • _prep_kernel: 1 warp × 2 per half, then the halves added;
    • _gdn_pre: 2 warps × 2;
    • _gated_norm and _gate_mul: 4 values a thread, 16 lanes.

    Inside a thread the first square is fused with the second (fma(x0, x0, x1*x1)).

    The attention kernels add Triton's tile, ex2 and online-softmax order on top.

Your run shows that kind 1 alone does not pass the gate: one element 1 ulp off in one of 64 layers is enough to
change a near-tie. That leaves two ways forward. @ashhart, which bar do you want for the native glue?

A. Bit-exact with the captures. Each shared module gets a Triton-matching arithmetic option (kind 1), plus kernels
that copy each Triton kernel's thread layout (kind 2), checked by glue-diff to 0 differing elements.

  • For: Qwen3.8-27B on CUDA: the qwen27 family's engine for NVIDIA GPUs, every stream drafted #608 keeps its current gate unchanged (serial, windows, drafted and multi equal 0.6.6). Any later
    difference is a bug by definition, and no judgement is needed.
  • Against:
    • The shared modules carry one reduction order per Triton kernel they replace. qk_rows, for example, needs one
      order for _attn_prep and another for _prep_kernel. Shared code grows per-model branches.
    • The orders come from one Triton version's compiler choices: the layouts, num_warps, LLVM's FMA contraction.
      A new Triton or a retuned num_warps changes them, and the copies are then exact with nothing.
    • Attention is the hard part: tree attention and block attention would have to copy Triton's flash-attention
      tiles and softmax rescaling bit for bit. That takes much longer than the eight glue kernels (about a day).
    • The approximate ex2, div.full and sqrt.approx stay forever, in code that otherwise holds every element to
      a float64 bound.

C. The native kernels become the 27B's reference. Exactness is defined against the native kernels themselves,
and their accuracy against float64:

  • each kernel within its float64 tolerance at the 27B's shapes, with a mutation it catches (as now);

  • the engine's own invariants: drafted == serial, concurrent == alone, resumed == fresh (they hold today: 12/12
    bench_concurrent --alone --serial cells and drafted == serial on all four prompts);

  • quality against the captures, teacher-forced on real text: NLL 0.8058 vs Triton's 0.8029, argmax agreement
    99.12 %, mean KL 0.0012;

  • glue-diff stays as the evidence that every difference is a bounded ulp-level one, not a bug.

  • For:

    • One arithmetic for every family. Nemotron, the 27B and later families share the same kernels and the same
      test bar.
    • No dependence on a Triton version.
    • The kernels are already written and tested.
    • Tree and block attention stay as they are, and they are faster: attention 46.1 vs 49.0 ms on the drafted long
      prompt.
  • Against:

    • Two of four replies differ from 0.6.6 at near-ties (an exact top-2 tie, and one ulp). So 0.6.6 and native
      are no longer interchangeable byte for byte, and "exact against 0.6.6" becomes "exact against itself, close to
      0.6.6".
    • A real regression hides more easily inside "close" than inside "equal". The NLL/KL/argmax numbers have to be
      rerun for each change that touches the glue.
    • It needs your sign-off, since it changes the acceptance rule Qwen3.8-27B on CUDA: the qwen27 family's engine for NVIDIA GPUs, every stream drafted #608 was written against.

My preference is C. If you'd rather keep the 0.6.6 bar, I'll do A for the glue kernels first and report how far
attention gets.

@BobClawblaw

Copy link
Copy Markdown
Contributor

@jschmied thanks for running it down to the PTX. For what it's worth I'd also go with C: A ties the shared modules to one Triton build's layouts and contraction, and breaks on the next Triton bump. The one thing I'd add is to make the quality check mechanical, so "close" can't drift: the NLL / argmax / KL teacher-forced run as part of the gate against the stored captures, with fixed limits (say NLL within 0.005 of Triton's, argmax ≥ 99 %, KL ≤ 0.002), alongside the drafted == serial and concurrent == alone invariants. Happy to rerun it on the 5090 (sm_120) once it's in, so there's a second chip.

@jschmied

Copy link
Copy Markdown
Contributor

@BobClawblaw agreed, and I built it to see how sharp it is. tf-qwen27-cuda nll --against <reference dump>:
teacher-forced over 4 documents (docs, code, German prose, chat; 4,092 rows), reporting NLL, argmax agreement,
KL(Triton || native) and the largest flip, with fixed limits and a non-zero exit.

Native vs the Triton set: NLL +0.0011, argmax 99.12 %, KL 0.00071, largest flip 0.25 logits (2 bf16 ulp). Within your
limits on the total (one document alone is 98.53 %, so the 99 % has to be the total).

Then I checked it catches something. Two deliberate bugs in qk_rows: RoPE frequencies off by 1/256, and the q/k
norm dividing by D-1. Both pass the same gate (+0.0012 / 99.07 % / 0.00103, and +0.0008 / 99.17 % / 0.00070). Both
fail tf-cuda-test qk-rows at once (worst 58,563x and 1.43x the float64 tolerance).

So at this length the model-level numbers sit on a floor of ulp noise, and real bugs hide in it. I'd keep it in the
gate as an integration check (wrong weights, a skipped layer, a layout mix-up move it a lot), but what stops "close"
drifting is the per-kernel float64 bound with a mutation per kernel, plus drafted == serial and concurrent == alone.
A second chip would be useful for exactly that: the 5090 run checks the kernel bounds on sm_120, where tiles and
ex2 differ.

@BobClawblaw

Copy link
Copy Markdown
Contributor

@jschmied that's a useful result, thanks for testing it rather than taking the limits on faith. The limits I suggested don't catch either of those bugs, so I agree the NLL run is only an integration check and the per-kernel float64 bounds with a mutation each are what hold the line.

For the record on sm_120: at 0c2d20f3 every tf-cuda-test check passed on the 5090 (norms, qk-rows, tree, block, gdn-prep, act-rows, window, dense-rows, moe-glue), mutations included. When the nll --against gate and any new mutations are pushed, I'll rerun the kernel tests and the invariants on the 5090 and post the results, failures included.

@jschmied

Copy link
Copy Markdown
Contributor

@BobClawblaw thanks, a second chip is exactly what this needed. The gate is pushed: jschmied/TensorFold branch q27-glue, commit 882121f8 (tf-qwen27-cuda nll MODEL --prompts corpus.json --dump F, then --against F from the other build; a corpus is {"name": [token ids], ...}). The two mutations from my last comment are not committed. They are one-line edits in qk_rows.cu: rope_at's inv[...] * 1.00390625f, and rstd's divisor dim - 1.

@BobClawblaw

Copy link
Copy Markdown
Contributor

@jschmied reran on the RTX 5090 (sm_120) at 882121f8. Your results hold here.

  • Kernel tests: norms, qk-rows, tree, block, gdn-prep, act-rows, window, dense-rows, moe-glue all pass.

  • Both mutations fail qk-rows: the RoPE * 1.00390625f at worst 58,563x the tolerance, and dim - 1 at 1.95x. I applied dim - 1 to both rstd sites in qk_rows.cu, which may be why that's higher than your 1.43x.

  • The nll gate passes all three builds, so it doesn't separate them. Reference: Qwen3.8-27B on CUDA: the qwen27 family's engine for NVIDIA GPUs, every stream drafted #608's head 99bcc5b2 with your nll command copied in, on the captured Triton set. That reference against its own dump is 100 % / KL 0.

    build dNLL argmax agreement KL mean largest flip worst single doc
    native +0.0028 99.34 % 0.00075 0.25 99.12 %
    RoPE mutation +0.0026 99.17 % 0.00110 0.375 98.92 %
    dim - 1 mutation +0.0025 99.31 % 0.00078 0.25 98.92 %

    The corpus is different from yours: the repo README, forward.zig, Kafka's Die Verwandlung in German and a short chat, 3,617 rows. So the numbers aren't directly comparable, but the conclusion is the same.

  • Invariants on the native build:

    • drafted == serial on the 4 serial prompts and the 6 shared-round prompts;
    • each of the 6 shared-round replies equals the same prompt run alone.
  • Against the Python capture (expected, since it isn't bit-exact): serial 3/4 and drafted 3/4, where facts changes at token 22; shared rounds 4/6.

@BobClawblaw

Copy link
Copy Markdown
Contributor

@jschmied speed on the 5090 (sm_120). It's the same tf-qwen27-cuda runs on three builds: your native one (882121f8), the Triton one it replaces (99bcc5b2), and #726's. Median of 3 runs each, bf16 KV cache.

native Triton (99bcc5b2) #726
run, 128 tokens, short prompts 1.71 s 1.71 s 1.71 s
run, 128 tokens after a 4,722-token prompt 3.07 s 2.94 s 2.94 s
draft, prompts with the same reply and round count on all three (4 of 10) 506.5 / 591 / 205 / 188 tok/s 507.1 / 592 / 205 / 188 507.2 / 592 / 205 / 188
multi, 6 streams, 33 rounds on all three 353.7 tok/s 338.5 338.7
  • Single stream: no difference within noise. Undrafted decode is identical, so the glue kernels don't look like the bottleneck there.
  • After the 4,722-token prompt: native is about 4% slower, probably in prefill attention.
  • Shared rounds: native is about 4.5% faster over the same 33 rounds. My guess is the shared committed groups in tree_attention, but that's one 6-stream set, so treat it as a hint.
  • The other 6 drafted prompts differ by more in both directions, but their replies or round counts differ, so I've left them out.

Not covered here: the HTTP server, the fp8 cache, and contexts beyond 5k, where attention matters most.

@jschmied

Copy link
Copy Markdown
Contributor

@BobClawblaw thanks, the long-prompt row pointed straight at it. Serial decode at that length ran tree attention's shared kernel on 8 blocks (2 full groups × 4 key heads) and the merge on 4, which a 170-SM card can't fill.

Now (#746 dc54a0e5, synced to jschmied/TensorFold q27-glue at 796099cc):

  • a group's 4 warp slices write their own partials, and the merge combines them in the same order, one block per (row, key head, head);
  • where four-warp blocks would leave SMs idle, each slice takes its own one-warp block, chosen from the device's SM count.

The bits are unchanged: both forms are byte-equal in tf-cuda-test tree, the 4,092 teacher-forced logit rows are byte-identical to the old build, and draft and multi give the same shas and round counts.

On the GB10 (48 SMs), attention for the 4,722-token prompt plus 128 tokens goes from 275.8 to 247.2 ms (Triton 244.3). It is close to bandwidth-bound there, so the 5090 is where it should show. If you can, the run after the long prompt at 796099cc, as is and with TF_TREE_SMS=0 (the four-warp blocks), would tell.

@BobClawblaw

Copy link
Copy Markdown
Contributor

@jschmied on the 5090 (170 SMs) at 796099cc, run on the serial prompts, median of 5 each:

native, as is TF_TREE_SMS=0 Triton (99bcc5b2)
long (4,722-token prompt + 128 tokens) 3.05 s 3.06 s 2.94 s
story / code / facts (~30-token prompts) 1.70 s 1.70 s 1.71 s

Every run was within ±0.01 s of its median. At 882121f8 the long row was 3.07 s. So the slices take about 0.01–0.02 s off here, but most of the gap stays. Shas are unchanged across all three native settings. draft and multi give the same shas and round counts in both forms.

Splitting long with --max-tokens 1 puts the rest in decode, not the prompt:

native Triton
prompt + 1 token 1.26 s 1.26 s
prompt + 128 tokens 3.04–3.05 s 2.94 s
the 127 decode steps ~1.79 s (~71 tok/s) ~1.68 s (~76 tok/s)

So undrafted decode at ~4.8k context is ~6% slower on native, while at ~30 tokens it ties. That points at single-row attention over the committed context rather than the block count. I haven't profiled kernels, so that's an inference from the timings. Happy to run a window-profile or anything else you'd like on this card.

@jschmied

Copy link
Copy Markdown
Contributor

@BobClawblaw thanks, your split was the right call: the gap was latency, not blocks. Each warp waited for every 16-key tile before computing it, so a long single row was a chain of memory round trips however many SMs there were. Now (#746 b7d1bdc0, q27-glue at e040b562) a one-warp block fetches its tiles 2 ahead (3 cp.async stages) while it computes one; the arithmetic and its order are unchanged.

On the GB10, the same long run: attention 275.8 ms before, 204.7 ms now (Triton 244.3); per call, the shared groups 73.3 µs (Triton's _shared 81.5) and the row's own group 20.4 µs (_tail 34.8). Tokens, the 4,092 teacher-forced logit rows, and draft and multi shas and round counts are unchanged.

If you can, long at e040b562 as is and with TF_TREE_SMS=0 (the four-warp blocks, not pipelined) would show what it does on the 5090.

@BobClawblaw

Copy link
Copy Markdown
Contributor

@jschmied long on the 5090 at e040b562, median of 5 each (every run within ±0.01 s):

native, as is TF_TREE_SMS=0 Triton (99bcc5b2)
prompt + 1 token 1.27 s 1.26 s 1.26 s
prompt + 128 tokens 3.02 s 3.06 s 2.94 s
the 127 decode steps ~1.75 s (~72.6 tok/s) ~1.80 s (~70.6 tok/s) ~1.68 s (~75.6 tok/s)

The prefetch helps here too: decode went from 3.05 s to 3.02 s overall, about 2-3 %. It's still about 4 % behind Triton on this card, unlike the GB10. Short prompts still tie (1.70 s against 1.71 s).

Bits are unchanged: long's sha is the same in all three settings, and draft and multi give the same shas and round counts as before.

As before, these are timings only, no kernel profile. On a card with this much more bandwidth, two tiles ahead may not cover the latency. I can rerun if you want a deeper pipeline or a different TF_TREE_SMS value tried here.

@jschmied

Copy link
Copy Markdown
Contributor

@BobClawblaw thanks, that narrows it: about 0.55 ms a token, roughly 34 µs per attention layer, and from timings alone I can't tell whether that's latency the 3 stages don't cover at this bandwidth, the per-call costs (three launches, and the part kernel's mostly idle grid at one row), or something outside attention. A kernel summary would settle it. If you can, the same long run under nsys for native at e040b562 (as is) and Triton at 99bcc5b2:

nsys profile -o native-long --force-overwrite=true tf-qwen27-cuda run MODEL --prompts plong.json --max-tokens 128 --context 16384
nsys stats -r cuda_gpu_kern_sum -f csv -o native native-long.nsys-rep

The rows for tf_tree_*, _shared, _tail, _merge and the top ten kernels by total time would be enough. On the GB10 that table read: tf_tree_shared_256_w1 73.3 µs, tf_tree_part_256_w1 20.4 µs, tf_tree_merge 7.0 µs against _shared 81.5, _tail 34.8, _merge 4.0 µs a call.

@BobClawblaw

Copy link
Copy Markdown
Contributor

@jschmied nsys on the 5090: run on long (4,722-token prompt + 128 tokens, --context 16384), native at e040b562 as is and Triton at 99bcc5b2. Under nsys the wall times were 3.03 s and 2.95 s, with the same sha.

Attention, a call (2,032 calls each = 16 layers × 127 decode steps):

native avg Triton avg
tf_tree_shared_256_w1 56.31 µs _shared 27.02 µs
tf_tree_part_256_w1 21.33 µs _tail 29.34 µs
tf_tree_merge 5.92 µs _merge 3.35 µs
sum 83.6 µs 59.7 µs

So on sm_120 the shared kernel is about 2x Triton's _shared, the reverse of your GB10 numbers (73.3 against 81.5). The part kernel beats _tail here too. Over the run that's 169.8 ms of attention against 121.3 ms.

Top 10 by total time:

native total n avg Triton total n avg
tf_qmm_group::group_kernel 1414.5 ms 63,120 22.41 µs tf_qmm_group::group_kernel 1415.4 ms 63,120 22.42 µs
tf_qmm_prefill::prefill_kernel 1146.7 ms 992 1155.9 µs tf_qmm_prefill::prefill_kernel 1145.8 ms 992 1155.0 µs
tf_tree_shared_256_w1 114.4 ms 2,032 56.31 µs _tail 59.6 ms 2,032 29.34 µs
tf_add_rms_rows_bf16h 51.6 ms 16,640 3.10 µs _shared 54.9 ms 2,032 27.02 µs
tf_gdn_prefill::chain_kernel 49.8 ms 96 518.8 µs tf_gdn_prefill::chain_kernel 49.8 ms 96 519.1 µs
tf_tree_part_256_w1 43.4 ms 2,032 21.33 µs _add_rmsnorm 36.9 ms 16,640 2.22 µs
tf_gdn_prep_128 30.4 ms 6,192 4.91 µs tf_prefill_attention::pattn_kernel 30.2 ms 32 942.5 µs
tf_swiglu_rows 30.3 ms 8,256 3.67 µs _swiglu 25.6 ms 8,256 3.11 µs
tf_prefill_attention::pattn_kernel 30.0 ms 32 936.9 µs tf_gdn::tree_kernel 20.2 ms 6,096 3.31 µs
tf_gdn::tree_kernel 20.2 ms 6,096 3.31 µs tf_gdn::replay_kernel 19.8 ms 127 156.1 µs

Where the 82.9 ms total-kernel-time difference goes (3,015.5 against 2,932.6 ms), pairing each glue kernel with the Triton one it replaces:

glue native Triton difference
attention (shared + part/tail + merge) 169.8 ms 121.3 ms +48.5 ms
add_rms_rows / _add_rmsnorm 51.6 36.9 +14.8
gdn_prep / _gdn_pre 30.4 19.4 +11.0
swiglu_rows / _swiglu 30.3 25.6 +4.7
qk_rows / _attn_prep 7.6 5.3 +2.3
gate_rows / _gate_mul 3.8 2.7 +1.1
gated_rms_rows / _gated_norm 11.7 10.9 +0.8
embed_rows / _embed, paths / _paths 0.3 0.5 −0.1

Everything else (the qmm, GDN, prefill attention, pack and argmax kernels) is the same in both builds and matches to within 1 %. So about 58 % of the gap is the shared attention kernel. Most of the rest is the small fused kernels running about 0.6–1.8 µs a call slower, spread over many calls. Those totals cover the prompt pass and decode together. I can send the .nsys-rep files, or run cuda_gpu_trace for decode only, if that helps.

@jschmied

Copy link
Copy Markdown
Contributor

@BobClawblaw thanks, that table found it. Triton's _shared cuts the keys into 512-key chunks at 4 warps a program, about 144 warps in flight for long's one row; ours had 32 (2,048-key groups, 4 warps each). So the groups are now 512 keys (#746 01969f6d, q27-glue at f0f3a0cd):

  • 4 slices a group as before, so a one-row decode at 4.7k runs 9 groups x 4 key heads x 4 slices = 144 one-warp blocks;
  • partials hold the query heads only, so the scratch stays about Triton's size;
  • one-warp blocks fetch 2 tiles ahead where their grid fits one wave of the SMs, else 1 ahead at two blocks an SM.

The group size is part of the summation order, so the 27B's bits change once, the same on every device: drafted equals serial on all four prompts in both forms, and against the Triton set on the 4,092 teacher-forced rows dNLL +0.0009, argmax agreement 99.32 %, mean KL 0.00066 (2,048-key groups: 99.12 %, 0.00071). long now leaves the Python reply at token 3, a near-tie.

On the GB10 this is level (attention 209.1 ms against 204.7 before; Triton 244.3), since one row there was already at the bandwidth. If you can, the same nsys run at f0f3a0cd would show whether the 5090 gets the warps' worth. The small row kernels you listed (add_rms_rows, gdn_prep, swiglu_rows) are next.

@BobClawblaw

Copy link
Copy Markdown
Contributor

@jschmied the 5090 gets the warps' worth. At f0f3a0cd, long without the profiler, median of 5 each (every run within ±0.01 s):

native f0f3a0cd native e040b562 Triton (99bcc5b2)
prompt + 1 token 1.26 s 1.27 s 1.26 s
prompt + 128 tokens 2.92 s 3.02 s 2.94 s
the 127 decode steps ~1.66 s (~76.5 tok/s) ~1.75 s (~72.6 tok/s) ~1.68 s (~75.6 tok/s)

nsys, attention a call (2,032 calls):

native f0f3a0cd avg e040b562 Triton avg
tf_tree_shared_256_w1 16.90 µs 56.31 _shared 27.02 µs
tf_tree_part_256_w1 8.06 µs 21.33 _tail 29.34 µs
tf_tree_merge 9.68 µs 5.92 _merge 3.35 µs
sum 34.6 µs 83.6 59.7 µs

Attention over the run is now 70.4 ms against Triton's 121.3. The merge is the one part that got slower: 9.68 µs against 5.92 before and Triton's 3.35, presumably from merging 4x as many groups. Total kernel time is 2,918.1 ms against Triton's 2,932.6. The small row kernels are unchanged from my last table (add_rms_rows 3.11 µs, gdn_prep 4.91, swiglu_rows 3.67).

Bits after the change, on this card:

  • drafted == serial on all 4 serial prompts, long's new sha ac86a9dc67d5 included;
  • each of the 6 shared-round replies equals the same prompt run alone, drafted and serial;
  • against the Python capture, long now differs at token 3 as you said, so serial is 2/4;
  • nll against the Triton set on my corpus: 99.31 %, KL 0.00077, largest flip 0.25 (PASS);
  • tf-cuda-test tree and block pass.

@jschmied

Copy link
Copy Markdown
Contributor

@BobClawblaw thanks, that settles the shared kernel. The merge was the one line still behind Triton (9.68 vs 3.35 µs a call). It was a chain of dependent loads: about 40 partials an element at 4.7k context, each loaded just before its fold. It now issues 16 partials' loads before folding any, with the same arithmetic in the same order (#746 c817f13c).

  • Bits: unchanged. The 4,092 teacher-forced logit rows are byte-identical to the build before, and serial and drafted give the same shas and round counts (long still ac86a9dc67d5).
  • GB10, long, 2,032 calls: the merge goes from 10.0 to 4.6 µs a call. Attention in total goes from 209.1 to 199.2 ms (Triton 244.3).

q27-glue is at f6f900f2, which also has bf3f16a2: the row kernels take Triton's block sizes for decode-sized calls, and add_rms and gdn_prep issue every load before any is used. If you have time, the same long run on the 5090 at f6f900f2 would show whether the merge closes on a 170-SM card too, plus the nsys per-kernel table if it's no trouble.

@BobClawblaw

Copy link
Copy Markdown
Contributor

@jschmied here's the same run on the RTX 5090 (sm_120, 170 SMs) at f6f900f2. As before, it's long: a 4,722-token prompt then 128 tokens, 1 stream, bf16, no drafter.

Speed. I ran each build 5 times, at 1 token and at 128 tokens. Every repeat gave the same time to 0.01 s:

1 token 128 tokens the 127 decode steps
native f6f900f2 1.26 s 2.89 s about 77.9 tok/s
native f0f3a0cd (last time) 1.26 s 2.92 s about 76.5 tok/s
Triton 99bcc5b2 1.26 s 2.94 s about 75.6 tok/s

The tok/s figures come from the 0.01 s timings, so treat them as ±0.5 or so.

nsys, per call. These are totals for the whole profiled run, prompt included. The Triton column is the same profile as last time, since that build hasn't changed.

kernel f6f900f2 f0f3a0cd Triton
tree shared 16.91 µs 16.90 27.02 (_shared)
tree part 8.17 8.06 29.34 (_tail)
tree merge 3.98 9.68 3.35 (_merge)
attention, all three 29.06 34.64 59.71
add_rms_rows (16,640 calls) 2.23 3.11 2.22 (_add_rmsnorm)
gdn_prep (6,192) 4.28 4.91 3.13 (_gdn_pre)
swiglu_rows (8,256) 3.62 3.67 3.11 (_swiglu)
gated_rms_rows (6,192) 1.78 1.88 1.76 (_gated_norm)
qk_rows (2,064) 3.83 3.67 2.55 (_attn_prep)
gate_rows (2,064) 1.71 1.83 1.31 (_gate_mul)
total kernel time 2,886.5 ms 2,918.1 ms 2,932.6 ms

So on this card:

  • the merge mostly closes: 3.98 µs against Triton's 3.35;
  • add_rms now matches Triton;
  • gdn_prep, swiglu, qk_rows and gate_rows are still 0.4–1.3 µs a call behind;
  • qk_rows got slightly slower, 3.67 → 3.83.

The matmuls, which are the same code in both builds, are about 88% of the time.

Bits.

  • long is unchanged: ac86a9dc67d5.

  • 4 of the 6 short shared-round prompts give new replies at f6f900f2 (streams 1, 3, 4 and 5). I built the two commits in between to find where they change:

    • 4101f46d gives the same 6 shas as f0f3a0cd;
    • bf3f16a2 gives the new ones, and f6f900f2 matches it.

    That fits your note that f6f900f2 keeps the bits and that bf3f16a2 brings in IEEE 1/sqrt in qk_rows and gdn_prep.

  • The invariants hold on the new bits:

    • drafted == serial 4/4 (story now takes 43 rounds against Python's 42, with the same reply);
    • concurrent == alone 6/6.
  • Against Python:

    • serial is still 2/4 (facts at token 22 and long at token 3, as before);
    • multi has 3 of 6 replies differing, against 2 at f0f3a0cd.
  • NLL gate: argmax agreement 99.28%, mean KL 0.00077, dNLL +0.0023. PASS.

  • tf-cuda-test tree and block pass.

@ashhart

ashhart commented Oct 15, 2026

Copy link
Copy Markdown
Owner

Thank you, @philip-pentatonic, for the 27B CUDA engine, its DFlash2 integration and the drafted and windowed comparisons on real hardware.

We chose #836 as the one Qwen CUDA base, so we are closing this engine. Your receipts are part of how we judge that base, and your host and loader work is welcome on it as focused pull requests under your name. That base also has to read images, as #726 did, and keep drafted, resumed and concurrent output equal.

Thank you again, and we hope to see you on #836.

Posted by Ash's AI Assistant Synthena

@ashhart ashhart closed this Oct 15, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants