Repository navigation
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
Closed
cuda/spec: qwen4exp decode improvements (Q8_0 MMVQ, MTP drafts, small-batch verify, HIP graph keys)#107gaetan-puleo wants to merge 26 commits into
gaetan-puleo wants to merge 26 commits into
Conversation
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.
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
force-pushed
the
feat/qwen4exp-decode-opt
branch
from
September 29, 2026 10:52
7054c4b to
d3f4d59
Compare
Collaborator
Author
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Decode-throughput improvements for
qwen4exp(Qwen3.8-Flash-Next UD-IQ4_XS) on gfx1151, in two groups: byte-exactQ8_0MMVQ 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
tg128(3 Q8_0 MMVQ kernels)-np 3)Changes
Plain decode (byte-exact
Q8_0MMVQ, logits SHA unchangedc283a8f4…)Q8_0(hc_{attn,ffn}_up,ncols<=512).Q8_0without fusion (ssm_out/attn_output,attn_q).Q8_0path (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_reset416 -> 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_cmpinstead 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.Q8_0MMVQ 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).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_indicespath we do not expose), and the MMBQ8_0hyper-connection / MoE GLU3 fusions (ourmmb.cuis a different backend). Measured but reverted: Q6_K MMQ at all widths (noise), BF16 MMB from 9 columns (regressed prefill/tg).Notes
qwen4exp.STRIX_*env knobs disable each change (default is the new behaviour) for A/B.