The previous commit excluded IQ1_S/IQ1_M outright, based on regressions at
n=4/n=8. That was the wrong reading: token generation runs at ncols_dst==1,
where IQ1_S is 5.2% FASTER, so excluding it would have removed a win from the
path that actually matters and made the model slower.
Re-measured with base/exp alternating per type and a discarded warm-up
invocation (the earlier sweep ran all-base-then-all-exp, so the first base
invocation of each pass caught the 645 MHz idle clock -- a bias favouring the
experiment, which is also what produced a bogus -78% for iq2_xxs).
iq1_s n=1 -5.2% n=2 -11.1% n=4 +7.7% n=8 +2.2% spreads <= 0.2%
iq2_s n=1 -9.2% n=2 -16.2% n=4 -13.3% n=8 -11.6% (8 KiB control, wins)
So restrict IQ1 to ncols_dst <= 2 and keep every other grid type at all n.
Both share the 8 KiB iq1s_grid_gpu. At 128 threads/block that drops blocks per
SM from 8 to 7, and measured they get slower: iq1_s +8.0% at n=4, iq1_m +13.5%
at n=8. IQ2_S is also 8 KiB and still wins, so size alone is not the cause --
the IQ1 kernels issue fewer codebook loads per call, so there is less LG-pipe
pressure to relieve and the lost occupancy dominates.
Kept: IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ3_S (1-8 KiB), all measured -8% to
-18% at n=1..8.
Extends the IQ2_XXS result (-13.6% at n=1) to every codebook type. Nsight shows
these kernels are lg_throttle-bound (19.81/issue vs mio_throttle 0.06) with L1
already at 98%, so the win is issue slots, not cache: an LDG costs an LG slot
even on a hit, and LDS uses the idle MIO pipe.
iq3xxs 1 KiB, iq2xxs 2 KiB, iq3s 2 KiB, iq2xs 4 KiB, iq2s 8 KiB,
iq1s_grid_gpu 8 KiB (IQ1_S and IQ1_M)
At 128 threads/block the 8 KiB grids cost one block of occupancy per SM
(8 -> 7), so IQ2_S / IQ1_S / IQ1_M are the ones to watch when measuring.
Nsight (via the CMP unlock) on mul_mat_vec_q<IQ2_XXS,ncols=1>, SM75:
lg_throttle 19.81/issue <- dominant, 43x the next stall
wait 1.43
long_scoreboard 0.46
mio_throttle 0.06 <- shared/MIO pipe idle
L1 hit 98.04%, DRAM 22.99%, SM 28.54%, occupancy 92.82%
Not bandwidth, not cache, not occupancy: the global load/store issue queue is
full. A vec_dot issues ~16 loads (4x16-bit from get_int_b2, 4x64-bit grid, 8x
32-bit activations); a load costs an LG issue slot even when it hits L1.
Move the four codebook lookups to LDS, which uses the MIO pipe that is sitting
idle. 2 KiB of shared per block, loaded once before the k-loop.
Measured +0.27% at n=2..4 against a -0.07% noise floor, i.e. no gain, so the
change was reverted. Records why the IQ4_XS technique does not transfer:
iq2xxs_grid is 2 KB and L1-resident so the eliminated lookups were already
cheap, while the decoded state is ~2x larger and costs occupancy. Also records
that load widening is impossible for IQ2_XXS (66-byte block, qs at offset 2).
Records every measurement with the exact launch command, plus the hypotheses
that were tested and rejected so they are not retried. Key entries:
- FORCE_CUBLAS is a stock upstream option worth +24.9% dense prefill; it is not
fork code and upstream gets it too.
- Fork AVX1 Q4_K CPU path is worth +7.1% MoE tg (4 alternating loads per arm).
- OMP_PROC_BIND=close + --threads-batch 40 is a free +4.6% MoE prefill.
- OMP_WAIT_POLICY=passive costs -47% tg: the libgomp spinning is load-bearing.
- mmap costs 5-10% vs --load-mode none, because numactl --interleave only
stripes anonymous memory, not page-cache-backed mappings.
- Method: alternate arms across model LOADS; reps within one load are correlated
and produced two contradictory conclusions that had to be withdrawn.
CUDA builds with CCCL < 3.2 have no cub::DeviceTopK, so ggml_top_k falls back to a full argsort + copy (~12 bytes of temp per cell). For sparse-attention top-k over long contexts this OOMs small cards (200K cols x 128 rows ~= 190MB+ of temp per layer). Reuse the exact radix selection from #27466 (merged for HIP) on CUDA when CUB_TOP_K_AVAILABLE is off: O(nrows) temp, same selected set. Also gate next_power_of_2, unused on that path now.
The thresholds were measured with GGML_CUDA_FORCE_CUBLAS=OFF, where dropping
out of mmvq lands in MMQ. With the option ON, ggml_cuda_should_use_mmq()
returns false unconditionally, so the same traffic would land in cuBLAS
instead, which dequantises the whole weight to F16 to multiply a handful of
columns. The measured 1.7x-2.6x MMQ advantage says nothing about that path,
so the thresholds are restricted to builds where MMQ actually exists.
Turing keeps the previous ne11 <= 8 behaviour when FORCE_CUBLAS is set.
Turing had no entry in ggml_cuda_should_use_mmvq and fell through to the
generic ne11 <= 8, while Ada, Blackwell, DGX Spark and Orin all have measured
switch points. MMQ overtakes MMVQ far earlier than n=8 on SM75.
Measured on a CMP 50HX (cc 7.5, CUDA 13.3) with test-backend-ops perf, medians
of 4 reps after a discarded warmup, forcing each path via an instrumentation
build. MMQ speedup over MMVQ at n=8:
q2_K 2.57x q3_K 2.13x q4_K 1.96x iq4_xs 1.93x q5_K 1.79x q6_K 1.74x
End to end on Qwen3.5-9B-UD-Q4_K_XL, llama-bench, medians of 3 interleaved
rounds at -r 20:
pp1 -0.2% pp2 0.0% pp3 0.0% pp4 +3.2% pp5 +12.1% pp8 +47.6%
q6_K is deliberately 3 rather than the 2 the synthetic sweep suggested. That
sweep had a single shape for q6_K and MMQ won it by only 1.02x, the weakest
margin measured. With the threshold at 2 the model regressed 2% at pp3,
reproducibly across three rounds: its 24 q6_K tensors were the only ones
switching path at that batch size. The synthetic margin did not survive on the
model's real q6_K shapes, so mmvq keeps n=3 for that type.
q4_K was measured across five m/k shapes; at n=3 mmvq wins 1.19x/1.09x/1.08x on
three and loses 1.01x on two, geometric mean 1.065 in mmvq's favour. At n=4 mmq
wins on all five.
The n=512 cases act as a control: above MMVQ_MAX_BATCH_SIZE neither mode can
select mmvq, and both measure within 1.00x-1.01x of each other.
Only the measured types are listed; everything else keeps the default. q2_K and
q3_K rest on a single shape each and are not present in the model used for the
end-to-end check, so they are less well evidenced than the rest.
Same story as the packed scale loads. The Q4_K reuse path decoded a block once
and applied it across ncols_dst columns, which beat the old branchy generic
path because that path re-ran the scale unpack per column. Upstream's
branchless rewrite hoists the unpack out of the loop, and the generic path is
now the faster of the two.
Measured on a CMP 50HX (cc 7.5), test-backend-ops perf, medians of 5 reps after
a discarded warmup, two independent runs:
reuse disabled vs enabled, q4_K n=2..4: -0.52% and -0.54% (faster without)
n=5..8, identical code in both builds: -0.02% (noise check)
widening the gate to n<=8 instead: +0.80% (worse)
14 of 15 cases at n=2..4 favour removal, so the effect is consistent rather
than drift.
The IQ4_XS reuse path is deliberately kept: upstream's change only touched
Q4_K/Q5_K scale unpacking, IQ4_XS decodes via kvalues_iq4nl table lookups and
its generic path did not improve. Measured with GGML_SM75_IQ4_REUSE=0, removing
it costs +1.14% at n=2..4, so it stays.
Default build: this narrow automatic dispatch gives about 9-10% prefill
gain in the tested 9B models. No extra build or launch flag is required.
For the larger measured 24-25% long-prompt gain, configure with
-DGGML_CUDA_FORCE_CUBLAS=ON and rebuild. This is an existing CMake option,
not a runtime environment variable or a new feature in this commit.
With it enabled, Ornith-1.0-9B UD-Q4_K_XL gained 25.24% prefill and
Qwen3.5-9B UD-Q4_K_XL gained 24.26% on 16384-token server requests,
with TG near 66 t/s. These results are specific to one CMP50HX 10 GiB,
the tested models and batch settings; they are not a universal 25% gain
or a full-model 27B measurement.
Global cuBLAS is not enabled by default: the tested N256 matrix operations
lost 12.83-28.20%. That is an operation-level result, not a measured
short-request server regression. Other shapes, workloads and peak memory
still need validation before broadening the default dispatch.
Select the existing FP16 cuBLAS path for contiguous, device-backed Q4_K
and IQ4_XS FFN weights at batches of 1024 through 2048 tokens on SM75.
Limit the rule to 4096x12288 and 5120x17408 matrices in either orientation
to bound the FP16 weight workspace and avoid unmeasured shapes.
Keep MMVQ decode, smaller batches, batched matrices, host-backed weights,
MoE dispatch, other quant types and other architectures unchanged.
Honor GGML_CUDA_FORCE_MMQ. No new kernel or quantization format is added.
Validation on one CMP50HX 10 GiB:
- 52 existing and 24 added numerical matrix cases passed.
- Three paired timing runs: selected shapes gained 12.4-43.9%; small-batch
controls ranged from -0.14% to +1.32%.
- Real 16384-token requests, three measured runs after warmup, no cached
prompt tokens: Ornith 9B UD-Q4_K_XL gained 9.77% prefill and Qwen3.5 9B
UD-Q4_K_XL gained 9.14% against the final baseline repeat. TG stayed
within 0.2% of that baseline.
Compute Sanitizer remains unavailable on this device; numerical
checks are not a substitute for sanitizer coverage.
Place IQ4_XS in the existing two-warp Turing dispatch for one output token column. Keep the arithmetic, tensor layout and multi-token launch choices unchanged. Exclude test-backend-ops tracing and deterministic-test modifications.
The full K=5120 IQ4 matrix gained 4.24% in the isolated comparison; another shape's smaller change had overlapping ranges. Correctness checks passed. The MTP model trace did not use N=1, and the no-MTP trace did; neither model comparison showed a clear gain. No-MTP text differed between builds, so do not claim bitwise model output equivalence. Retain as a separate experiment, not a production speed claim.
Separate Q4_K fragment decoding from the Q8_1 dot product and reuse decoded nibbles, scales and minima across N=2..4. Preserve DP4A input-sum correction and floating-point operation order. Limit the new path to Turing and keep standard tensor layouts unchanged.
Baseline and candidate passed 72 matrix cases each. Warm six-run comparisons gained roughly 4-7% on larger matrix shapes, peaking at 7.41%. A separate warm prefill recheck did not reproduce the first series' 1.01% outlier. Model PP/TG ranges overlap (+0.051% PP, -0.162% TG medians). No demonstrated end-to-end gain.
Decode an IQ4_XS fragment once and reuse it across two to four token columns. Use aligned int2 loads and guard the layout assumptions with static assertions. Share the fusion eligibility predicate between graph selection, launch dispatch and entry checks. Limit multi-column fusion to IQ4_XS on Turing; leave GGUF and other architecture paths unchanged.
Passed 52 matrix and 36 fused correctness cases. Nsight confirmed the fused path. Warm six-run graph comparisons gained 1.32%, 12.29% and 9.85% at N=2,3,4 for full K=5120 gate/up. The six-run CPU-heavy model comparison had overlapping PP/TG ranges, with +0.080% PP and -0.133% TG medians. Kernel gain only; no demonstrated model speedup.
Use 128-bit sign application and four-row transpose helpers for AVX1 panel construction. Enable IQ3_XXS panels only with GGML_AVX1_IQ3_PANEL=1; keep default format selection and standard GGUF layouts unchanged. Do not include the audit-only GGML_IQP_VERIFY additions.
The combined AVX1 verification build passed 173 numerical cases, with 33 unsupported cases separate. Synthetic IQ3_XXS matrix throughput improved 124-230%. The installed Flash file has different tensor types, so this is not a measured Flash speedup. No model gain or production promotion is claimed.
Use 128-bit unpack operations to transpose four quantized rows into the existing Q8_Kx4 layout. Build signed block sums with packed integer operations. Keep input quantization and the AVX2 path unchanged; share the panel helpers with the separate IQ3 experiment.
Measured CPU Q4_K matrix throughput gains: 35.17% at N=8, 12.10% at N=32 and 6.01% at N=256. Helper checks and the combined AVX1 numerical verification passed. Dense-model PP/TG ranges overlap: no demonstrated model speedup. Keep this as an isolated experiment on base 609be054c.
Add AVX1 CPU kernels for the Q4_K and IQ4_XS quantized matrix paths.
The Q4_K changes add panel repacking and prefill support. The IQ4_XS changes decode eight-row panels with AVX1 intrinsics and select the path only when AVX is available. Existing fallback paths remain in place for CPUs without AVX and for other quant types.
Measured on Qwen3.8-Flash-Next UD-Q4_K_XL (512 experts, 10 routed, 28 expert
layers pinned to host RAM by -ot): expert routing has strong temporal locality
(LRU-64 ~67% hit rate over a mixed workload) even though the long-run expert
distribution is near-uniform, so a per-layer LRU cache of expert slices in VRAM
removes most of the per-token host-RAM streaming that bounds decode.
Mechanism (no custom CUDA kernels):
- companion tensors [ne0, ne1, K+1] per cached layer in the device buffer of
that layer's router; slot K stays all-zero (dummy)
- I32 id->slot tables, one device copy (get_rows remaps ids for a second
mul_mat_id chain over the cache) and one host copy (src[3] of the CPU
mul_mat_id, which skips cached ids and zeroes their dst rows)
- the two down outputs are summed; the split is exact by construction
- decode-only (n_tokens == 1); batch/prefill builds the stock graph
- throttled async uploads: evictions are published at a decode-boundary sync
point, slices are copied by a worker thread, and the new mapping is only
published after the upload completed, so a running graph never reads a torn
slot
Enable with LLAMA_MOE_CACHE_SLOTS=<K> (+ LLAMA_MOE_CACHE_INSERTS, _DEBUG).
Inert without the env var.
Measured decode, this box (2x3090, 1 RAM channel/socket, with numactl
--interleave=all): 15.5 -> 19-20 tok/s warm at K=48-64.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01A7KRyLzuBiGfVn3bdczjka
Each cold PLE row demand-faults a 4 KiB page for ~90 bytes of data,
capping cold diverse-text prefill at 218-360 tok/s vs ~785 warm on
GB10. All n-gram row indices of a ubatch are known host-side before
the graph runs, so under the new LLAMA_LAZY_MODE_DIRECT
(--lazy-mode on-direct) they are staged into an input tensor with
sorted, deduplicated, parallel pread()s and dequantized exactly like
ggml_get_rows; downstream kernels unchanged, table stays on disk.
Cold diverse prefill on qwen3.8-flash-next: 542-741 tok/s (2.0-3.1x,
within ~6% of warm); warm, decode and greedy outputs bit-identical.
* llama-context : report graph inputs and input tensors during sched reserve
- fix the tg (token generation) graph bs label to use n_seqs instead of a hardcoded 1
- report the number of graph inputs from llm_graph_result::inputs for both the pp and tg graphs
- report the number of input tensors (nodes and their src tensors flagged with GGML_TENSOR_FLAG_INPUT)
- log a warning when an input tensor has an op other than GGML_OP_NONE
- log a trace line for each input tensor and the nodes (name and op) that use it
Assisted-by: llama.cpp:DeepSeek-v4-Flash-0731
* cont : count input tensors before reserving the sched
* wip
* llama-graph : name the unnamed graph input tensors
- name the kv-cache idxs input tensors (attn_inp_k_idxs, attn_inp_v_idxs)
- name the recurrent state copy idxs input tensor (rs_s_copy)
- report the input tensor shape in the sched_reserve trace
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* llama-context : rename "graph inputs" to "graph input objects"
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* llama-context : report the sched reserve graph stats on a single line
- print nodes, splits, input objects and input tensors in one line
- when the pp and tg graphs differ, print each value as 'pp / tg'
and annotate the line with the batch sizes used for each graph
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* cont : pad logs
The 5-argument load_ldmatrix added in 1884824fd only defines tile<16,8>, so the Volta tile<8,4> does not match. See https://github.com/ggml-org/llama.cpp/issues/29222 for details. Building on 1884824fd, generalize the tile shape of the 5-argument load_ldmatrix from <16,8> to <I,J>, so the non-swizzle branch forwards to the 3-argument loader for any shape. Local compilation and testing passed.
Assisted-by: DeepSeek V4.1 Flash (OpenCode)
convert_unary handles the contiguous case through the general strided kernel,
one element per thread: each lane reads 4 bytes and writes 2. Converting the
activations for a bf16 matrix multiplication that way moves 126 MB in 1021 us
on gfx1151, about 65% of what the memory system can do.
Give the contiguous path its own kernel that takes four elements per thread
through a vector type, so a warp loads 512 bytes at a time instead of 128. It
is used only when the element count is a multiple of four and both pointers
carry the alignment the vector type needs, and falls back to the strided
kernel otherwise.
Model level, Qwen3.8-Next-Flash IQ3_XXS on gfx1151, llama-bench -ub 2048 -r 6,
mean of the last 3 reps, ABBA counterbalanced:
pp2048 688.0 680.0 -> 694.3 691.1 +1.26%
tg128 24.8 24.8 -> 24.8 24.8 +0.14%
Every conversion in a prefill takes the new kernel (kernel trace: 1146
convert_unary_cont_vec4, no convert_unary). Output is bit identical; MUL_MAT,
MUL_MAT_ID, CPY, CONT, GET_ROWS and SET_ROWS pass.
Assisted-by: Claude Opus 5
Ilya Maslennikov
Georgi Gerganov
xrip
rhonstin
Rhonstin
Denis Artyomov
csantiago78andClaude Fable 5
Josh Leverette
Foad Abo Dahood
lingyezhixing
Piotr Wilkin (ilintar)