HIP : RDNA3.5 MMVQ/MMQ kernels and per-type dispatch thresholds for quantised matmul - #56
Open
SimonTeixidor wants to merge 15 commits into
Open
SimonTeixidor wants to merge 15 commits into
SimonTeixidor wants to merge 15 commits into
Conversation
The RDNA3.5 short-circuit forces IQ2_XS and IQ2_S onto MMQ at every batch size. On gfx1151 hipBLAS is faster than MMQ for both types well below the ubatch sizes used for prefill. At ne11 <= 128 MMQ still wins, so that threshold is kept. The MMQ/hipBLAS time ratio is monotone in ne11 for every quantised type, so the "ne11 <= T" form is right; only T was wrong for these two. Mechanism: IQ2_XS/IQ2_S carry one scale per 16 weights (IQ4_XS: per 32), forcing a doubled WMMA epilogue that no config change can remove, while their dequant kernels are the cheapest of any quantised type. Worst MMQ plus cheapest dequant gives the earliest crossover of any type. This reverts one clause of PR ggml-org#18666 to what PR ggml-org#18537 shipped four days earlier. The IQ2_XS/IQ2_S config rows in this tree are byte-identical to upstream's; this tree additionally carries an IQ2_S prefetch that makes its MMQ faster than upstream's. Assisted-by: Claude Opus 5 Claude-Session: https://claude.ai/code/session_01VdFZgAf86WFrMQjttNagfW
__vsub4 was forwarded to __vsubss4, which saturates; CUDA's __vsub4 wraps. The two disagree on every input where a byte subtraction overflows. Assisted-by: Claude Opus 5 Claude-Session: https://claude.ai/code/session_01VdFZgAf86WFrMQjttNagfW (cherry picked from commit ae8d32eebe49107d7b0d6f6b53013ddbf706e0be)
RDNA has no packed-byte compare or subtract, so hip.h emulates both per byte. The IQ2_XXS/IQ2_XS/IQ2_S/IQ3_XXS/IQ3_S unpacks call that pair once per 4 weights. Replacing it removes the emulation from every call site in the MMVQ and MMQ kernels. The replacement is bit-identical to CUDA's __vcmpne4/__vsub4 for every input, checked exhaustively over grid words and sign nibbles on the host and on gfx1151. Adding the two's complement +1 only into the low 7 bits of each byte, then folding bit 7 back with an XOR, keeps the carry inside its byte. Note hip.h routes __vsub4 to the saturating __vsubss4, so master's HIP path already differs from CUDA whenever a grid byte is 0x80. This follows CUDA's semantics on both vendors and removes that divergence; it is unreachable for the five grids in any case, whose bytes are 0x01..0x3e. unpack_ksigns() now returns the 8 sign bits as a byte rather than broadcast, and is renamed unpack_ksigns8 so stale callers cannot compile. (cherry picked from commit 81dd22156a1d2ae2aebb74de9e6619d6f482bb6c) Assisted-by: Claude Opus 5
The upstream F1 change only covers the call sites that exist upstream. This fork carries three more copies of the same __vcmpne4/__vsub4 sign idiom, in code that does not exist upstream: mmvq.cu mul_mat_vec_q4_columns_rdna3_5 (IQ3_S, 4-column RDNA3.5 kernel) mmvq.cu vec_dot_iq3_s_q8_1_grid vecdotq.cuh vec_dot_iq3_s_q8_1_pair Left alone they would keep the emulated pair in the hottest IQ3_S decode kernels, i.e. lose most of the win on this fork. The nibble mapping is the same one apply_signs4 takes: ((s & 0x03) << 7) | ((s & 0x0C) << 21) makes byte i of the selector word nonzero exactly for bit i of s & 0x0f, and ((s & 0x30) << 3) | ((s & 0xC0) << 17) does the same for s >> 4. apply_signs4 is bit-exact with that idiom under CUDA's __vsub4, so vec_dot_iq3_s_q8_1_pair stays bit-identical to two scalar vec_dot_iq3_s_q8_1 calls, which is the invariant its comment states. Assisted-by: Claude Opus 5 Claude-Session: https://claude.ai/code/session_01VdFZgAf86WFrMQjttNagfW
ggml_cuda_should_use_mmvq falls through to ne11 <= MMVQ_MAX_BATCH_SIZE (8) for every quantized type on AMD; the per-type tables are only filled in for Ada, Blackwell, DGX Spark and CDNA. On gfx1151 8 is too high for most types. MMVQ re-reads the src1 column and redoes the dot per output column while the weight unpack is hoisted out of the ncols_dst loop, so its cost is A + ncols_dst*D with D dominated by the type's vec_dot. MMQ's smallest instantiated tile on RDNA3.5 is J = 16, so one column tile covers the whole ne11 = 1..16 range and MMQ's cost is flat across it. The crossover is therefore a per-type constant. Both paths were timed at the same ne11 on real model shapes with a temporary dispatch override, so MMQ's cost below ne11 = 9 is measured rather than extrapolated. The crossover does not move with m or k once MMQ fills a wave. Q1_0, Q2_0, IQ2_XXS, IQ2_XS, IQ2_S and NVFP4 keep the existing bound of 8. This replaces an earlier table that covered only four k-quants and extrapolated the other types from an assumption that turned out to be false. This fork has GGML_HINT_EXACT_BATCH. The hint promises that a batched mul_mat equals n single-column mul_mats exactly, which any ne11-dependent kernel switch inside [1, MMVQ_MAX_BATCH_SIZE] breaks, so the exact_batch bypass is threaded through mmvq.cuh / ggml-cuda.cu. Assisted-by: Claude Opus 5 Claude-Session: https://claude.ai/code/session_01VdFZgAf86WFrMQjttNagfW
calc_rows_per_block returns 1 unconditionally for MMVQ_PARAMETERS_RDNA3_5, while GENERIC, GCN, TURING and GB10 all return 2 for ncols_dst 2..8. The RDNA3.5 entry was added alongside the nwarps = 1 entry and inherited its value without a separate measurement. On RDNA3.5 nwarps is 1, so a thread block is a single wave and every lane re-reads the whole q8_1 activation block for the one row the block owns. At k = 5120 and ncols_dst = 6 that is 6 * 5120 * 1.125 B = 34.5 KB of activation traffic per output row, through the vector memory unit, as 4-byte loads (block_q8_1 is 36 B so qs is only 4-byte aligned). With two rows per block the same loaded activation words feed two rows, halving that traffic per row. Bit-exact: rows_per_cuda_block only changes which rows a block owns and how many accumulators it carries. The per-lane K decomposition (kbx stride blocks_per_iter, kqs) and therefore the dot summation order per row is unchanged, and the cross-lane reduction is the same warp_reduce_sum. The rest of the kernel is already written for rows_per_cuda_block > 1 and is exercised there by the four other tables: calc_launch_params divides nrows_x by rpb, tmp/tmp_gate are [ncols_dst][rpb], the fusion bias prefetch and the write-back both guard with threadIdx.x < rpb and row0 + i < stride_col_dst, and the tmp_shared reduction is indexed by rpb. The RDNA3.5-specific specialisations (the q4_columns kernels at ncols_dst = 4, the IQ3_S ncols_dst = 1 kernels, the MUL_MAT_ID multi-token kernel) compute their own launch geometry and do not go through calc_rows_per_block. Register pressure: at ncols_dst = 6 the accumulator array grows from 6 to 12 floats, and to 24 when a fused gate is present. Assisted-by: Claude Opus 5
The MMVQ/MMQ and MMVF/MMF thresholds are per-type ne11 crossovers measured at m >= 5120, where MMQ fills the device. They say nothing about small m. The smallest MMQ tile instantiated on RDNA3.5 is I = 64 rows, so at ne01 = 48 -- the GDN ssm_alpha and ssm_beta projections, 96 dispatches per decode step on the 27B models -- MMQ launches ceil(48/64) = 1 block and 39 of the 40 CUs sit idle for the whole call, while MMVQ launches a block per row (or row pair) and uses the device. The same holds for MMF against MMVF. So: on RDNA3.5, when ne01 is below one tile, take the vector kernel for any ne11 the vector kernel supports. ggml_cuda_should_use_mmvq did not receive ne01, so it is threaded in from ggml_cuda_mul_mat; ggml_cuda_should_use_mmvf already had src0_ne. Guarded to RDNA3.5 because that is where these tile sizes apply. This changes which kernel runs, so the result can differ in the last ulp the same way the existing per-type thresholds already make it depend on ne11. Under GGML_HINT_EXACT_BATCH the RDNA3.5 MMVQ path already returns ne11 <= MMVQ_MAX_BATCH_SIZE, so the new rule agrees with it. Assisted-by: Claude Opus 5
These three vec_dots finish by multiplying an accumulated dp4a sum by its
integer block scale: sumi0*ls0 + sumi1*ls1 for iq2_xs and iq2_s, sumi *= 1+2*sc
for iq3_s. On RDNA3 a 32-bit integer multiply is quarter rate; the 24-bit form
v_mul_i32_i24 is full rate.
The compiler does not find this on its own: in a built gfx1151 library the
MMVQ kernels for these types multiply each block scale with v_mul_lo_u32. It
does use the unsigned 24-bit form elsewhere in the same file when it can prove
the operands are small, but it cannot bound a signed dp4a accumulator.
The multiply must be signed. The codebook bytes reach the dp4a with their signs
already applied, so sumi is genuinely negative about half the time and
__umul24 would be wrong.
Range proof, so the 24-bit form is bit-exact. __mul24 sign-extends the low 24
bits of each operand, so it agrees with a full multiply when both operands lie
in [-2^23, 2^23) and the product fits in int.
iq2_xs, iq2_s: sumi0 and sumi1 each accumulate 4 dp4a, i.e. 16 int8 products.
The iq2xs and iq2s codebook byte alphabet is {8, 25, 43}, so a signed grid
byte is at most 43 in absolute value; a q8_1 byte is at most 128. So
|sumi| <= 16*43*128 = 88064. The scales are nibbles, 0..15. Product at most
1320960.
iq3_s: sumi accumulates 8 dp4a, i.e. 32 int8 products, over the iq3s codebook
alphabet {1, 3, ..., 15}, so |sumi| <= 32*15*128 = 61440. The scale is
1 + 2*nibble, at most 31. Product at most 1904640.
Both are within 2^23 = 8388608 by more than a factor of 90, and both products
are far inside int. The remaining epilogue arithmetic ((sumi0 + sumi1)/2, the
/4) is adds and shifts and is untouched.
Applied at every site these three types reach: vec_dot_iq2_xs_q8_1,
vec_dot_iq2_s_q8_1, vec_dot_iq3_s_q8_1 and vec_dot_iq3_s_q8_1_pair in
vecdotq.cuh, and vec_dot_iq3_s_q8_1_grid and the IQ3_S branch of
mul_mat_vec_q4_columns_rdna3_5 in mmvq.cu.
Guarded to HIP: __mul24 is full rate on RDNA3 but emulated, and slower than a
plain multiply, on every NVIDIA architecture since Fermi.
Not applied to MMQ: load_tiles for all three types folds the integer scale into
the float block scale (x_df) instead of multiplying the accumulated sum, so
there is no integer multiply on that path to convert.
Assisted-by: Claude Opus 5
mul_mat_vec_q4_columns_rdna3_5 and mul_mat_vec_q4_columns compute their own launch geometry (row = blockIdx.x) and never consult calc_rows_per_block, so the two-rows-per-block change misses them and ncols_dst == 4 keeps re-reading the four q8_1 activation columns once per output row. Give both kernels the same treatment directly: a block owns rows row0 and row0 + 1, each row's weights are decoded into its own registers, and the four activation columns are loaded once and dotted against both. Launch geometry becomes ceil(nrows_x / rows_per_block) blocks, the accumulator and the cross-warp staging array gain a row dimension, and the write-back moves from lane 0 to lane r under the same row0 + r < stride_col_dst guard the generic kernel uses. The per-lane K decomposition, the dp4a order within a column and the cross-lane reduction are unchanged for every (row, column) pair, so the results are bit-exact. The ids and fused-gate paths never reach these kernels: the dispatch site takes them only for ids == nullptr with no fusion arguments. Assisted-by: Claude Opus 5
calc_rows_per_block returned 2 for every type at ncols_dst 2..8 on the RDNA3.5 parameter table. Four rows per block lowers the per-column cost for the codebook types, does nothing for q1_0/q2_0/q8_0, and regresses q5_1. A blanket 4 also drives mul_mat_vec_q<q2_K,8>, <NVFP4,7> and <NVFP4,8> into register spills. Take the subset that measured well and leave the rest alone: iq1_s iq1_m iq2_xxs iq2_xs iq2_s iq3_s iq4_xs 4 rows, ncols_dst <= 6 q2_K 4 rows, ncols_dst <= 4 everything else 2 rows, as before q2_K stops at 4 because that is also its MMVQ threshold on this arch, so 5..8 is reachable only through the ne01 < 64 path, and it is the type with the worst register growth. No type is promoted at ncols_dst 7 or 8, which is where the spills are. calc_rows_per_block gains a ggml_type parameter; both call sites already have type as a template parameter inherited from the same mul_mat_vec_q_switch_ncols_dst<type> instantiation, so the host grid calculation and the device kernel cannot see different types. The type argument is read only inside the RDNA3.5 arm; every other parameter table returns exactly what it returned before, pinned by static_asserts. Also assert rows_per_cuda_block <= warp_size in mul_mat_vec_q: the epilogue writes row row0+i from lane i, an invariant that was never close at two rows. Assisted-by: Claude Opus 5
… quants
Two changes to the MMVQ vector-dot path on the IQ codebook types, both aimed at
mul_mat_vec_q<IQ2_S, 6> and <IQ3_S, 6>, the largest consumers of GPU time in a
six-token decode step on models built from those types.
L-EPI: IQ2_XXS / IQ2_XS / IQ2_S / IQ3_XXS / IQ3_S scaled the integer dot product with a
truncating signed integer divide inside the ncols_dst loop. Folding the block scale and the
divisor into the float scale moves all of it to per-weight-block work that hoists out of that
loop. It is also exact where the integer form truncated.
The IQ3_S four-column kernel and the LDS-codebook variant carry the same transform so the
n-column and single-column paths still agree (test_mul_mat_exact_batch checks that to 1e-7).
L-SGN: apply_signs4 spends three instructions containing a carry that cannot occur for any
grid apply_signs4 is used with -- iq2*_grid bytes are {8, 25, 43}, iq3xxs_grid {4..62},
iq3s_grid {1..15}, none zero, so 0xff ^ g <= 0xfe and the +1 never leaves its byte. MMVQ
gets apply_signs4_nz; the MMQ callers keep the general form.
IQ4_XS and IQ1_S kernels are unchanged, and no kernel gains a register spill.
Assisted-by: Claude Opus 5
Claude-Session: https://claude.ai/code/session_016inZWVhdrrrkrjh2djvH7N
GGML_MMVQ_THR="<ggml type id>:<max ne11 kept on MMVQ>[,...]" overrides ggml_cuda_should_use_mmvq's per-type RDNA3.5 thresholds at runtime, so one binary can serve both arms of a crossover comparison. Unset (the default) changes nothing, and the override is placed after the GGML_HINT_EXACT_BATCH early return so exact-batch semantics are untouched. The existing thresholds predate the rows-per-block and epilogue changes, which lowered the MMVQ per-column cost, so they need re-measuring. Assisted-by: Claude Opus 5 Claude-Session: https://claude.ai/code/session_016inZWVhdrrrkrjh2djvH7N
…/Q6_K J = 16 is the only MMQ tile instantiated for ne11 <= 16 on RDNA3.5, so it is the tile every decode, small-batch decode and speculative-verification step runs on, and it was not in the sweep the original whitelist was tuned on. The motivation for prefetching does not depend on J: Q4_K at J = 16 needs 22080 B of LDS against 26816 B at the already-whitelisted J = 48, so both sit at exactly 2 blocks resident per 64 KiB CU, while staging the activation tile costs fewer registers at J = 16 than at the larger tiles. No kernel gains a spill. IQ4_XS and IQ3_S at J = 16 are left out: they showed no per-call effect and would bring their kernels close to the register limit. Assisted-by: Claude Opus 5 Claude-Session: https://claude.ai/code/session_016inZWVhdrrrkrjh2djvH7N
Every bound in this table except IQ3_XXS was set before the rows_per_block and
epilogue work lowered MMVQ's per-column cost. This re-measures all of it, using
GGML_MMVQ_THR so one binary serves both arms, at the (ne00 x ne01) pairs each
type takes in Qwen3.8-27B-UD-IQ4_XS.
IQ3_S 6 -> 8 Q6_K 3 -> 4 Q3_K 3 (re-measured, unchanged)
IQ4_XS 6 -> 8 Q8_0 4 -> 5 Q4_K 2 (re-measured, unchanged)
IQ3_XXS 6 -> 8 Q5_K 2 (re-measured, unchanged)
IQ2_S 8 (re-measured, unchanged)
The k-quant bounds stay low because their per-column decode cost is several
times that of the IQ types at the same shape, and MMVQ redoes it per column.
For IQ3_S, IQ4_XS and IQ3_XXS MMVQ stays faster than MMQ through ne11 = 8, so
8 here is MMVQ_MAX_BATCH_SIZE rather than a measured meeting point. These three
types are most of a UD-IQ4_XS file's matmul weights, and all three left the
vector path at the same width, which moved nearly the whole file onto MMQ at a
speculative verification width of 7.
Not measured, left alone, and said so in the comment: Q2_K, Q4_0, Q4_1, Q5_0,
Q5_1, MXFP4, IQ4_NL, IQ1_S, IQ2_XS, IQ2_XXS, NVFP4, Q1_0, Q2_0.
Known incomplete, and recorded in the comment: the crossover is a per-(type,
ne01) constant, not a per-type one. At ne01 = 1024 (attn_k, attn_v) MMQ does
not fill the device and the crossovers differ. The bounds here are the
ne01 >= 5120 values, which is the conservative choice; expressing the other
regime needs a code change.
The default MUL_MAT test grid cannot reach this code: its cases are m = 16 and
should_use_mmvq returns early when ne01 < 64.
Assisted-by: Claude Opus 5
Claude-Session: https://claude.ai/code/session_016inZWVhdrrrkrjh2djvH7N
The previous commit re-measured the nine types a local GGUF contains and left
the other eleven alone. This measures those eleven the same way, at generic 27B
shapes (5120x17408, 17408x5120, 5120x6144).
Q5_1 5 -> 2 Q2_K 4 -> 6 IQ2_XS 8 (re-measured, unchanged)
NVFP4 8 -> 4 Q4_1 5 -> 6 IQ2_XXS 8 (re-measured, unchanged)
MXFP4 5 -> 6
IQ4_NL 5 -> 6
Q4_0 6 -> 8
Q5_0 6 -> 8
IQ1_S 6 -> 8
These bounds come from representative shapes, not from a real file, and have no
end-to-end check, because no model on the tuning box contains these types. The
comment says so as well.
The two lowerings matter most, because the old bounds sent these types to MMVQ
where it is slower:
Q5_1 has one of the largest per-column costs in this function, above Q4_K and
Q5_K, and crosses between ne11 = 2 and 3. Its bound of 5 kept it on MMVQ at
ne11 = 3..5.
NVFP4 crosses between ne11 = 4 and 5. Its bound of 8 kept it on MMVQ at
ne11 = 5..8. The previous comment's "NVFP4 does not cross below 9" does not
reproduce. This fits rdna3_5_rows4_max_ncols_dst already excluding NVFP4
because its wide kernels spill.
Q1_0 and Q2_0 are now the only two types in this function that have never been
measured on gfx1151.
Assisted-by: Claude Opus 5
Claude-Session: https://claude.ai/code/session_016inZWVhdrrrkrjh2djvH7N
dzannotti
requested changes
Sep 15, 2026
Collaborator
There was a problem hiding this comment.
Validation: a Release HIP gfx1151 build completed. test-backend-ops MUL_MAT for q4_0 passed 1273/1273 ROCm cases.
Benchmark attempt: against current master on the locally available Qwen3.6-35B-A3B Q4_K model, pp128 measured 177.69 versus 177.84 t/s at ubatch 6 (+0.08%) and 64.89 versus 65.17 t/s at ubatch 1 (+0.44%), six samples per arm with alternating builds. This does not reproduce the reported large gain, but the PR used a different Qwen3.8 IQ4_XS model and quant mix, so it is not a direct refutation.
Please split by coherent kernel/type family, add boundary tests around every changed crossover, run all changed quant types, and provide a reproducible gfx1151 benchmark using the published fixture or an accessible equivalent.
Scope/reviewability: 15 commits and broad dispatch changes are too much for one acceptance decision. Much of the edited CUDA/HIP shared backend is generic; keep any Strix-only policy narrowly gated or stage it upstream.
comment generated by my clanker Codex
Author
|
Hi! To validate the speedup, please try an IQ-type quant. This is the specific model I tested with: https://huggingface.co/unsloth/Qwen3.8-27B-GGUF/blob/main/Qwen3.8-27B-UD-IQ4_XS.gguf. I am trying to upstream these things but I'm getting little response so far: |
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
2 participants
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.
Branch
pr/rdna35-quant-matmul(79194d304), 15 commits onhalo/master69946438a.Overview
Quantised matmul at small batch widths (decode, speculative verify, 2-16 columns).
2ae2ee95f,66efda4f6,e1430b151: branch-free SWAR__vsub4/__vcmpne4/__vcmpeq4and IQ2/IQ3 sign path (CUDA semantics;__vsub4no longer saturates).856d62f47,9a08368a9,dada19e0c: two output rows per thread block (four for codebook types) in MMVQ, bit-exact.c6e682ce0: 24-bit signed multiply for the IQ2_XS/IQ2_S/IQ3_S scale epilogue, bit-exact.7ed89d9e5: fold block scale and divisor into the float scale for IQ2_XXS/IQ2_XS/IQ2_S/IQ3_XXS/IQ3_S, removing a truncating divide (changes results); carry-freeapply_signs4_nz.151cdf1f7: MMQ prefetch for the J = 16 tile on Q4_K/Q5_K/Q6_K.df3e4d79b,16730c2c8,79194d304: re-measured MMVQ/MMQ crossover thresholds, all types but Q1_0/Q2_0 (e.g. IQ3_S/IQ4_XS/IQ3_XXS 6 -> 8, Q6_K 3 -> 4, Q5_1 5 -> 2, NVFP4 8 -> 4).b9bbaa7ab: vector kernel whenne01is below one MMQ tile (e.g. 48-row GDN alpha/beta projections).b35847506: IQ2_XS/IQ2_S use hipBLAS abovene11 = 128.c4d485be3:GGML_MMVQ_THRenv override of the threshold table (unset = no change).Measurements
Stock:
654803517, same flags, image and session. Mirrored blocks, n = 6 (width sweep) or 4 per build.All width-sweep (ub 1-512), TG and batched-decode cells resolve; pp2048 at ub 1024 is unresolved at every depth.
Correctness:
test-backend-ops -b ROCm0: MUL_MAT 7438/7438, MUL_MAT_ID 7097/7097, GET_ROWS 215/215. Full suite: 22599 OK, 0 FAIL, then a SOLVE_TRI abort stock shares; remaining ops 352/352.7ed89d9e5by design) by no more KLD than stock's own ub 8 vs ub 512, with PPL deltas inside the error bars (prose +/- 0.355):-c 512 -b/-ub 512, 4 chunks)Raw results
Batch width:
llama-bench -p 128 -b 128 -ub 1,2,3,4,6,8,12,16 -n 0 -r 3and-p 512 -b 512 -ub 64,512 -r 3, both with-ngl 99 -fa on -ctk f16 -ctv f16 --load-mode none; blocks stock, PR, PR, stock.Token generation:
llama-bench -p 0 -n 128 -d 0,32000 -b 1024 -ub 1024 -r 2.Batched decode:
llama-batched-bench -c 8192 -b 4096 -ub 512 -npp 512 -ntg 128 -npl 1,2,4,6,8, decode fromspeed_tg, one sample per invocation (A/C: stock, gdn PR, this PR; B/D reversed).Prompt processing:
llama-bench -p 2048 -d 0,12000,32000,64000 -b 1024 -ub 1024 -n 0 -r 2.Correctness:
llama-perplexity -c 512 -b/-ub 512, 4 chunks,--kl-divergence-base/--kl-divergence(base file quantises log-probs to 16 bits);llama-debug --save-logits -b 16384 -ub 512;llama-servergreedy/completion, 256 tokens (128 on the long prompt),n_probs 10. Same-b/-ub, context, KV types and FA on both builds.At
-ub 8(LLAMA_GRAPH_REUSE_DISABLE=1, 2 chunks), PR vs stock: prose KLD 0.000937, top-1 98.24 %, PPL 5.5569 -> 5.5752; code KLD 0.000610, top-1 99.61 %, PPL 2.2002 -> 2.2055.Greedy decode: top-10 log-probs differ from step 0 on all prompts; max |d logprob| of the chosen token before divergence 0.028 .. 0.53.
Threshold case files (explicit crossing cases) passed 6960/6960 and 1160/1160 on earlier builds with a fixed
test-backend-ops(results/2026-09-10-crossover-retune/); the default grid'sMUL_MATcases usem = 16, below thene01 < 64early return.Full suite log: SOLVE_TRI abort at
SOLVE_TRI(type=f32,ne_lhs=[64,64,2,2],...); stock aborts the same way on-o SOLVE_TRI. Ops run by name:-o ADD1,ADD_REL_POS,CROSS_ENTROPY_LOSS,...,GATED_DELTA_NET,...,WIN_UNPART.Additional information
Requirements
654803517. The branch was then rebased onto69946438a, whose 6 new commits touch onlyggml/rocmfpx, one Vulkan shader andtests/CMakeLists.txt; not re-measured.ne111-8,min {66, 67, 130, 1026}): not re-run,--test-filefalse FAIL on this tree.pr/hip-rebase-fixes) and SOLVE_TRI (aborts on stock too).-r 2/-r 3over mirrored blocks, not one-r 4.vecdotq.cuh/mmvq.cu: not compiled.https://claude.ai/code/session_01MkuKwgPNgpAg7EbxKR2iWC