Skip to content

cuda/spec: qwen4exp decode improvements (Q8_0 MMVQ, MTP drafts, small-batch verify, HIP graph keys) - #107

Closed
gaetan-puleo wants to merge 26 commits into
halo-box:masterfrom
gaetan-puleo:feat/qwen4exp-decode-opt
Closed

gaetan-puleo wants to merge 26 commits into
halo-box:masterfrom
gaetan-puleo:feat/qwen4exp-decode-opt

Conversation

@gaetan-puleo

@gaetan-puleo gaetan-puleo commented Sep 29, 2026 •

Copy link
Copy Markdown
Collaborator

Summary

Decode-throughput improvements for qwen4exp (Qwen3.8-Flash-Next UD-IQ4_XS) on gfx1151, in two groups: byte-exact Q8_0 MMVQ fixes that raise plain single-token decode, and speculative-decoding (MTP) improvements that raise both single-stream and multi-sequence decode. Nothing here changes model math for the default path except where noted.

All measurements on the same machine, llama-bench/server interleaved A/B, medians.

Results

area before after delta
plain decode tg128 (3 Q8_0 MMVQ kernels) 25.71 t/s 26.39 t/s +2.65 %
MTP single-stream 36.31 t/s 41.58 t/s +14.5 %
MTP multi-sequence (-np 3) 37.05 t/s 44.3 t/s +19.6 % (now level with plain batched decode 45.2)

Changes

Plain decode (byte-exact Q8_0 MMVQ, logits SHA unchanged c283a8f4…)

  1. Rows-per-wave for short-K Q8_0 (hc_{attn,ffn}_up, ncols<=512).
  2. Deep prefetch for long-K Q8_0 without fusion (ssm_out/attn_output, attn_q).
  3. Deep prefetch for the grouped long-K Q8_0 path (attn_qkv+attn_gate).

Speculative decoding
4. Even-length MTP drafts across sequences (STRIX_SPEC_EVEN_DRAFTS, default on). The target's hybrid memory splits a batch into ubatches that need the same token count per sequence, so drafts of different lengths turned one verify into extra whole-model passes (each a fresh graph). All drafting sequences now pad to the longest confident run. Multi-sequence aggregate: 31.4 -> 36.3 t/s (adaptive), fixed-length np=3 unchanged (already even). Single-sequence byte-identical.
5. Vector-kernel chunks for small-batch MoE router and experts (STRIX_F32_VEC_CHUNK_MAX=32, STRIX_MOE_VEC_CHUNK=4 / _MAX_T=16). A 5-16 token verify step's F32 router [2560x512] and its routed experts (past the vector limits: IQ3_S 4, IQ4_NL 6) fell to tiny GEMM/MMQ tiles; they now run the vector kernel over column chunks. Multi-sequence -np 3: 37.0 -> 42.7 t/s.
6. Small products of a verify step off one-workgroup kernels (STRIX_Q_VEC_CHUNK_MAX/ROWS, STRIX_CONCAT_T_MIN). Quant weights of <=1024 output rows (hc_down [10240->320], attn k/v, shared-expert gate/up) chunked through the vector kernel, and the GDN conv concat uses the tiled transpose from 2 tokens. Multi-sequence np=3: +2.6 %.
7. HIP graphs keyed on the graph shape as well as its first node (STRIX_GRAPH_KEY_SHAPE). Every draft graph replays instead of resetting its neighbour (warmup_reset 416 -> 5). Single-stream MTP: 36.45 -> 39.86 t/s (+9.4 %), byte-exact.
8. IQ3_S vec-dot for the HIP MMVQ path. The sign mask is one v_cmp instead of a per-half unpack; byte-exact. Single-stream MTP: 39.86 -> 42.02 t/s (+5.4 %).

Correctness

  • test-backend-ops test -b ROCm0 -o MUL_MAT,MUL_MAT_ID: 14577/14577 on the final build.
  • The Q8_0 MMVQ changes (1-3), the HIP graph-key change (7) and the IQ3_S vec-dot (8) are byte-exact (identical final logits SHA / identical MTP content sha1).
  • The speculative multi-sequence changes (4-6) change which kernel small batches take, not the arithmetic; they are byte-exact on single-sequence and on the measured multi-sequence outputs (identical content sha1).

Not included / could not port

From the same optimization set, these did not apply to this fork (diverged plumbing/backends) and are left out: GDN deferred rollback (lazy recurrent-state records), QSA decode gather + small-batch mask (gated on a direct_indices path we do not expose), and the MMB Q8_0 hyper-connection / MoE GLU3 fusions (our mmb.cu is a different backend). Measured but reverted: Q6_K MMQ at all widths (noise), BF16 MMB from 9 columns (regressed prefill/tg).

Notes

  • These kernels are shared with other architectures; measurements here are on qwen4exp.
  • STRIX_* env knobs disable each change (default is the new behaviour) for A/B.

gaetan-puleo and others added 22 commits September 27, 2026 23:12
Measured on qwen3.8-flash-next / gfx1151 with -p 2048 -d 0,12000,32000,64000:
+3.6% d0, +4.7% d12k, +4.0% d32k, +1.9% d64k vs THRESH=128. Byte-identical
greedy output, test-backend-ops 29643/29643. MMB_THRESH env overrides.

Assisted-by: opencode (qwen3.8-flash-next)
…oE fusions

Prefill (pp4096, -b/-ub 4096) of Qwen3.8-Flash-Next on Strix Halo (gfx1151, ROCm 10):
1229 -> ~1640 t/s (lazy-mode off or on-direct + LLAMA_PLE_PRELOAD=1).

- gated_delta_net: chunked prefill path (chunk prep + sequential scan on WMMA, f16 q/k/v)
  with the causal conv + SiLU, the q/k L2 norm and the gated output RMS norm fused in.
  fp16 WMMA operands: not bit-identical to the sequential kernel (NMSE ~2e-7).
  GGML_GDN_CHUNK=0 disables it, GGML_GDN_CHUNK_MIN sets the token threshold.
- PLE: the host-side row gather runs on up to 32 threads (LLAMA_PLE_HOST_THREADS);
  LLAMA_PLE_PRELOAD=1 keeps the on-direct table in RAM; LLAMA_PLE_STATS prints timings.
- mmb: HC down GEMM reads Q8_0 directly (no BF16 shadow, MMB_HCD_SHADOW=1 restores it),
  f32split prefetch + narrow tile, ssm_alpha/ssm_beta as one dual F32 GEMM; fix the
  MMB_BN=256 routed GLU instance, whose waves covered only 128 of the 256 columns.
- qsa3_rows: bitonic sort instead of the rank sort (GGML_QSA_RANKSORT=1 restores it).
- shared expert: sigmoid(gate) * out + moe_out fused in prefill (GGML_CUDA_DISABLE_SGMA=1).
- tests: MOE_PREFILL (512 experts, top-10, skewed routing measured on a 4096-token
  ubatch), eval at n=512; GGML_PERF_MOE=1 runs it at n=4096 in perf mode.
…efill)

Branchless IQ3_S grid/sign decode in mmb_tile_gemm_glu (XOR sign bits, v_perm
packing) and mmb_pack2 via __builtin_amdgcn_perm. Same output (md5 bafa569b,
MOE_PREFILL 2/2).
…t add

moe_weighted_reduction_sgma_v4 computes sum_k(e_k*scale_k)*w_k + shexp*sigmoid(gate)
in one pass, bit-identical to the two kernels. graph_optimize moves the reduction
next to SIGMOID->MUL->ADD so the fusion matches. GGML_CUDA_DISABLE_MOE_SGMA=1
disables it. pp4096 1639 -> 1669 t/s on gfx1151, PPL 1.3851 unchanged.
Router f32split kernel 87 -> 77 ms per pp4096 on gfx1151 (about -0.4% GPU time),
same output (PPL 1.3851 per chunk, md5 bafa569b, MUL_MAT f32 217/217).
n_embd/4 = 640 for n_embd 2560: one block per token and no idle lanes (256 left
128 of 768 idle). 66 -> 64 ms per pp4096 on gfx1151, same output.
…usion

Review fixes for the fused reduction + shared-expert merge:
- the graph_optimize reorder also rejects in-between nodes that write (in place /
  through a view) memory the reduction reads, not only nodes that read it
- reorder and fusion use one predicate (ggml_cuda_moe_sgma_tail, including the
  kernel type/shape checks) and honour GGML_CUDA_DISABLE_FUSION, so nodes are only
  moved when the fusion can fire
- drop the unreachable BF16-output instantiation: a BF16-only sum stays unfused

Same GPU time and output (PPL 1.3851 per chunk, md5 bafa569b, MOE_PREFILL 2/2).
…E_SGMA_THREADS

640 threads only fit n_embd 2560. Use n_embd/4 (one block per token, no idle
lanes) when it is a whole number of warps <= 1024, else 256. The env override is
rounded to whole warps within [32, 1024]; 0 or garbage no longer divides by zero.
Same 640 threads, GPU time and output on qwen4exp (md5 bafa569b).
The MMB_F32_TILE=5 default applied to every F32-weight GEMM with M > 64. Above
M 1024 the narrow tile reloads the activations ~3x more per output, so those keep
master's 128x128 tile. The M=512 router keeps 32x64 (77 ms per pp4096, about
-0.4% GPU time, not -0.5% as the comment said). MMB_F32_TILE still overrides.
Same output (MUL_MAT f32 217/217, PPL 1.3851 per chunk, md5 bafa569b).
The other __builtin_amdgcn_perm uses in mmb.cu are unguarded (device code in
this file is AMD-only), so the #if __HIP_DEVICE_COMPILE__ split with a scalar
fallback was dead and inconsistent. Same code on device (MMB_QUANT 140/140, md5 bafa569b).
MOE_SHEXP_MERGE builds reduction + shexp * sigmoid(gate) with the merge adjacent
(fusion only) and with the shared-expert GEMMs in between, allocated through the
scheduler so graph_optimize reorders and buffer reuse follows the new order. F32
and IQ4_NL / Q8_0 (MUL_MAT_ID) experts, n_embd 2560 / 2048 / 4608 (640, 512 and
256 threads). Two hazard cases put an in-place writer of the experts / router
weights in between and check no reader of that memory ends up after it; both fail
if the reorder's hazard check is removed. 7 of 8 cases run the fused kernel.
The chunked GDN computes q/k/v convolution at the GDN node after eliding the earlier conv/SiLU. Keep its activation, history and taps allocated through that later read; otherwise buffer reuse corrupts Qwen3.6 prefill. Extend dependencies only when the fused prefill pattern matches. Add focused backend coverage for the scheduler-allocated fused path.
Master now gates MMB per backend context. Pass the current CUDA context at the three Qwen4Exp fusion checks introduced by this branch, so the optimized path follows that gate.

Assisted-by: OpenCode (gpt-6-sol)
On gfx1151, the Q4_K routed gate/up kernel takes 476 ms instead of 510 ms per pp4096 trace. Retain the one-lane kernel behind MMB_Q4K_HALF=0 and cover Q4_K routing boundaries plus Q5_1 down weights.

Assisted-by: OpenCode (gpt-6-sol)
…ss per row

At depth the QSA block selection dominated prefill growth. The F32 score GEMM now applies relu, the 4-head sum and the compact visibility in its epilogue, and the head-sum kernel absorbs the visibility chain when the GEMM is not fused; both reproduce the unfused arithmetic bit for bit. Top-k uses one workgroup per row with a 12-bit LDS histogram and an LDS candidate bucket instead of four histogram passes, and selects the same set. pp4096 at depth 128k on gfx1151 goes from 1074 to 1304 t/s (UD-Q4_K_XL, -b/-ub 4096). GGML_CUDA_DISABLE_IDX_GEMM, GGML_CUDA_DISABLE_IDX_VIS and GGML_CUDA_TOPK_WG=0 restore the previous kernels.

Assisted-by: OpenCode (gpt-6-sol)
The fused scorer re-read its F32 key tile for each of the 16 128-column tiles of the (query, head) columns. At 250k depth the ~32 MB of pooled keys per layer no longer fit the Infinity Cache, so the scorer took about 329 ms per 4096-token ubatch on gfx1151. Each workgroup now rounds 64 key rows to F16 in LDS once and walks every column, staging 128-column chunks of the queries (rounded once per call); same WMMAs, fragments and K order, so the result is bit-identical. 176 ms at 250k; unchanged at 128k. MMB_IDX_SCORE_V2=0 selects the previous kernel.

Assisted-by: OpenCode (gpt-6-sol)
…flight

For 512 rows of 62500 scores (QSA block selection at 250k depth) the per-row kernel takes 748 us instead of 1107 us on gfx1151 (radix path: 4596 us). The selected set is unchanged. Adds perf cases for the 128k and 250k shapes.

Assisted-by: OpenCode (gpt-6-sol)
log(step(d)) is exactly 0 or -inf and the head sums are >= +0, so a select gives the same bits as the logf. With 64-column query steps the scorer takes about 1450 us instead of 1710 us for 62500 blocks x 512 queries on gfx1151. Adds INDEXER_SCORE shapes for 128k and 250k depth.

Assisted-by: OpenCode (gpt-6-sol)
…rger cost

The pack copies the whole cache per ubatch whatever n_q is, so short prompts at
depth pay O(n_kv) for nothing. qsa3_attn_kernel gains a DIRECT variant that reads
the same fragments from the cache rows (K: 32 B at key*nb1 + (2w+t)*32; V: four
scalar 2 B loads packed into the same uint2), bit-identical to the packed path.
The selector is size-based (n_q*32 < nk) with QSA_DIRECT to force either side.
Measured interleaved (r=1, on-direct): pp1024 d240k +22%, pp2048 d240k +9.6%,
pp1024 d128k +4.1%, pp4096 d240k +5.0%; pp4096 d124k keeps the pack (-1.8% if
forced direct). QSA backend 32/32 and numeric-corpus logits byte-identical in
both modes.
Short-K Q8_0 single-row matvecs (ncols<=512, e.g. the [320 -> 10240] hc_attn/hc_ffn up projections)
expose the full DRAM latency of one ~340 B row per wave. Put 2 output rows per wave for this shape so
a wave keeps independent weight streams in flight. Byte-exact: one wave still reduces each full row
with the unchanged lane->K mapping and per-lane accumulation order; the fused scale->silu/sigmoid
activation prologue is replicated with the same expressions. Env knobs: MMVQ_FQ_RPW (1 disables,
2 default), MMVQ_FQ_NW / MMVQ_FQ_PFLONG for geometry sweeps.

Verified: saved-logits SHA b363cf26...752bcd identical (default and disabled), test-backend-ops
MUL_MAT/MUL_MAT_ID 14577/14577. Perf tg128 d0: 24.28 -> 24.50 t/s (+0.9%).
The hyper-connection stream mix (sigmoid -> mul -> 3x add -> scale, ~6 tiny kernels per mix per layer)
was fused only for prefill; decode forced the unfused path. hc_mix_reduce_f32 replays the unfused
mul/add roundings exactly (explicitly rounded v_mul/v_add in stream order 0..hc-1) and the hc==4
n_tokens<=32 variant uses the same in-order sum, so the fused reduction is bit-identical at decode.
Enabling it saves ~6 launches x 2 mixes x 47 layers per token.

Verified: saved-logits identical on numeric/code/structured corpora (b363cf26/e4d17c1c/eb837cb3),
test-backend-ops MUL_MAT,MUL_MAT_ID,MOE_PREFILL 14583/14583. Perf tg128 d0: 24.49 -> 25.33 t/s
(+3.5%, 3 interleaved pairs, sigma <= 0.05). HC_NO_MIX_FUSE=1 restores the old decode path.
@gaetan-puleo gaetan-puleo changed the title cuda: Q8_0 MMVQ decode improvements for qwen4exp (rows-per-wave + deep prefetch) cuda/spec: qwen4exp decode improvements (Q8_0 MMVQ, MTP drafts, small-batch verify, HIP graph keys) Sep 29, 2026
Extend long_rows to long-K Q8_0 shapes (ncols>=640) with no fusion and
a grid larger than 1024 rows, so the case 16/8 dispatch uses the
PF_LONG (pf=8) prefetch. Only the load distance changes; the per-lane
K order and accumulation are unchanged, so logits stay byte-exact
(SHA c283a8f4...). ncols<640 and fused shapes are untouched (hc_up
keeps its rows-per-wave). test-backend-ops MUL_MAT,MUL_MAT_ID
14577/14577. tg128 +0.83% (26.05 vs 25.835).
In ggml_cuda_mmv_group, dispatch the case-16 grouped matvec to
mul_mat_vec_fq_group<16, MMVQ_FQ_PF_LONG> when the shared K is long
((ncols/QK8_0) >= 20), targeting attn_qkv+attn_gate [2560 x 10240+6144]
Q8_0. Only the load distance changes, so the per-lane accumulation
order and results are unchanged (logits SHA c283a8f4...). pf=8 beats
pf=4 (26.415 vs 26.18) and the pre-change binary (26.075). Env knob
MMVQ_FQ_GROUPPF=4 selects pf=4. test-backend-ops MUL_MAT,MUL_MAT_ID
14577/14577. tg128 +1.30%.
One context alternates between batch shapes on the same first node (the MTP
draft context runs its catch-up at 1+n_draft tokens then its draft steps at 1,
the target runs 1+n_draft). Keying on the first node alone made every switch
look like a property change: warmup reset, the next graph run direct, the one
after captured again. Fold the shape (n_nodes and every node's op and ne) into
the key so each shape keeps its own captured instance. STRIX_GRAPH_KEY_SHAPE=0
restores the first-node key.

Measured, single sequence MTP fixed drafting, 192 tokens, same build A/B:
  off: 36.45 tok/s, warmup_complete 416, warmup_reset 416
  on : 39.86 tok/s, warmup_complete  12, warmup_reset   5
Byte-identical output (same content sha1s, draft counts and acceptance).
Port of strixllama apply_iq3s_vecdot. Sign handling uses a 16-entry byte-mask
table and a single v_cmp on HIP instead of the per-half unpack; the per-element
products and their order are unchanged, so output is byte-identical.

Measured, single sequence MTP fixed drafting, 192 tokens:
  before: 39.86 tok/s
  after : 42.05 / 42.00 tok/s
Same content sha1s, draft counts and acceptance. test-backend-ops
MUL_MAT,MUL_MAT_ID OK on ROCm0.
@gaetan-puleo
gaetan-puleo force-pushed the feat/qwen4exp-decode-opt branch from 7054c4b to d3f4d59 Compare September 29, 2026 10:52
@gaetan-puleo

Copy link
Copy Markdown
Collaborator Author

Superseded: this duplicated the decode branch of #106 and carried MTP commits. Reorganized — general decode (Q8_0 deep-prefetch, HIP graph keys, IQ3_S vec-dot) now stacks on #106; the speculative-decoding changes (even-length MTP drafts, small-batch verify) now live on #105.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant