CUDA: size MMQ ids-path tail padding from the flattened row count, not ne11

[glennneuber]

Overview

The ids branch of ggml_cuda_mul_mat_q() sizes the src1_q8_1 allocation from two different row counts. The data term uses ne12*n_expert_used. The tail-padding term uses ne11. The correct value is already computed a few lines below as ne11_flat.

For MoE gate/up projections the activations are broadcast across experts, so ne11 == 1. ggml_cuda_mmq_get_J_max() with ne11 == 1 computes min(1,512) = 1, then 1 - 1%8 = 0, skips its loop and returns 0. The buffer gets no tail padding at all, while MMQ reads in tiles of up to J_max = 512 rows and runs past the end.

ffn_down is affected too but less: it passes ne11 = 8, so padding is sized for 8 rows instead of ne12*n_expert_used.

The !ids branch above is correct, because there ne11 really is the buffer's row count.

This patch uses ne12*n_expert_used for the padding term.

Additional information

Hit this on a MoE vision model (256 experts, 8 used, q4_K gate/up) on sm_120 / CUDA 13.0 when a whole image was submitted as one 2040-token ubatch:

decoding image batch 1/1, n_tokens_batch = 2040
find_slot: non-consecutive token position 4 after 3 for sequence 0 with 2040 new tokens
ggml-cuda.cu:106: CUDA error
ggml_cuda_compute_forward: MUL_MAT_ID failed
CUDA error: an illegal memory access was encountered
  current device: 0, in function ggml_cuda_compute_forward at ggml-cuda.cu:2374

The failing node is ffn_moe_gate, with ne2=2040 ne02=256 ne11=1 ne12=2040, type q4_K, taking the mmq branch. The find_slot line also shows up on runs that do not crash, so it marks the path rather than the fault.

Notes for testing:

• test-backend-ops does not catch it. The over-read lands in padding rows the kernel discards, so output is unchanged and NMSE is unaffected. compute-sanitizer --tool memcheck does flag it.
• A passing run does not prove much. The overrun is a fixed size past the end, so it only faults when it crosses an unmapped page. The same request crashes on a fresh pool and passes after the pool has served a larger allocation.
• Extra memory is at most 512 * sizeof(block_q8_1_mmq) = 72 KB per call, and does not grow with ne12.

Tested with 4 cold runs on the patched build (no fault) against 2 unpatched controls from the same tree (both fault).

May be related to #24399, #19705 and #18331. I have not reproduced those configurations, so this is a guess, but the GGML_CUDA_FORCE_CUBLAS=ON workaround in #24399 skips MMQ entirely and requantising changes the allocation size, which would both hide this.

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: YES - I used AI to investigate the crash, locate the cause and produce the patch. I have reviewed the change and can explain it.

[ggml-gh-bot]

Hi @glennneuber, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

• PR Template not respected: Please respect the template when creating a new pull request. Make sure to fill out all required sections.

• AI-generated content: While code is allowed to be generated by AI, please write the PR description and commit messages on your own without the help of AI.

---

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

[glennneuber]

This is a regression imho. I was able to pin it to a commit.

It was introduced by 6eddde0 ("CUDA: refactor MMQ kernel configuration", #24127), build b9992. That commit changed the padding term in both allocation branches:

-            get_mmq_x_max_host(cc)*sizeof(block_q8_1_mmq);
+            ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, ne11) * sizeof(block_q8_1_mmq);
-        get_mmq_x_max_host(cc)*sizeof(block_q8_1_mmq);
+        ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, ne11) * sizeof(block_q8_1_mmq);

That is correct for the !ids branch, where ne11 is the row count. In the ids branch the row count is ne12*n_expert_used, so passing ne11 is wrong there.

The difference matters because the two functions have different failure modes. get_mmq_x_max_host(cc) only looks at the architecture and returns 128 or 64, so the old code always allocated some padding no matter what the shape was. ggml_cuda_mmq_get_J_max() looks at ne11 as well, and for ne11 == 1 it computes min(1, 512) = 1, then 1 - 1 % 8 = 0, skips its loop and returns 0. So the broadcast case gets no padding at all rather than a bit too little.

Last clean build is b9990 (259ae1d), first affected is b9992. b9991 is not tagged and the only other commit in that range is Vulkan-only.

Bisected by source, then checked at runtime on both sides of the boundary. Same GPU, same request, cold server, first request each time, n_ubatch = 2048 and n_tokens_batch = 2040 in every log:

build	result
b9888 (ollama 0.32.1)	no fault, 3/3 cold runs
b10069 (ollama 0.32.2)	illegal memory access
b10353	illegal memory access

For anyone hitting this through ollama: v0.32.1 is the last release on a clean llama.cpp, v0.32.2 is the first affected.

[manunicholasjacob]

I ran the ids path under compute-sanitizer on four architectures to see what this patch covers, since whether the over-read shows up depends on the allocator. Build at b96806d with a diagnostic pool (every ggml_cuda_pool_alloc is its own exact-size cudaMalloc, GGML_CUDA_NO_VMM=ON), then test-backend-ops test -b CUDA0 -o MUL_MAT_ID -p type_a=<q> under compute-sanitizer --tool memcheck --padding 65536 --destroy-on-device-error kernel, master versus this PR's one line. Buffers are identified by allocation size: ne_get_rows*4 is ids_dst, ne_get_rows*ne10_padded*144/128 is src1_q8_1. The stock legacy pool with no padding reports 0 errors and 74/74 on every card, which is why nobody sees this.

P100 (sm_60) and T4 (sm_75) on Kaggle, A100 (sm_80) and L4 (sm_89) on Colab, CUDA 12.8:

arch	filter	binary	memcheck errors	reads past src1_q8_1	reads past ids_dst
sm_60	q4_0	master	9,668	64-row buffer, up to 221 B past, mul_mat_q<Q4_0, J=32>	yes
sm_60	q4_0	this PR	20,815	none	yes, up to 61 B
sm_60	q8_0	master	18,020	5-row buffer (2,880 B), up to 461 B past, J=8	yes
sm_60	q8_0	this PR	17,375	5-row buffer, up to 317 B past, still there	yes
sm_75	q4_0	master	95,063	64-row buffer, up to 989 B past	yes
sm_75	q4_0	this PR	11,937	none	yes, up to 117 B
sm_80	q4_0	master	78,497	64-row buffer, up to 1,373 B past	yes
sm_80	q4_0	this PR	12,469	none	yes, up to 109 B
sm_89	q4_0	master	86,154	64-row buffer, up to 957 B past	yes
sm_89	q4_0	this PR	13,298	none	yes, up to 121 B
all	q8_0, q4_K	both	6,500 to 18,000	none except the sm_60 5-row case above	yes, J = 16 to 80

Three things fall out.

1. The patch does what it says: the 64-row src1 over-read is gone on all four cards.
2. ne_get_rows < 8 stays uncovered wherever that many rows reach MMQ (sm_60 here; the other three send 5 rows to MMVQ). ggml_cuda_mmq_get_J_max() does ret -= ret % 8 before its loop, so 5 rows gives 0 padding with or without the patch, while the launched kernel is the J=8 instance and loads 8 rows. The dense path has the same arithmetic for any ne11 below 8 that the per-arch tables send to MMQ (Ada: Q2_K at 5..8; Blackwell: K-quants at 6..8), which I have not measured. Padding by the J actually launched, or by a flat 512 * sizeof(block_q8_1_mmq) as the description already contemplates, closes both.
3. ids_dst is read past its end by every MMQ instance on every card and both binaries, by up to (J-1)*4 bytes: mmq.cuh loads ids_dst[col_low + jt*J + j] for j < J with no bound on col_diff, so the last tile of the last expert runs off the array. The values are discarded by write_back, same class as the src1 tail. Same allocation site, so it probably belongs in this PR: pad ids_dst by the same J, or clamp the load.

The FAIL counts under --destroy-on-device-error kernel are the sanitizer terminating the offending kernel, not wrong numerics; with the stock pool every case passes on both binaries. Two q4_K launches on the T4 report reads gigabytes from any allocation; they appear only after an earlier kernel in the same case was terminated and are excluded. Raw logs, the pool patch and the parser are available if useful.

[1jeffchristensen]

Confirming this fix resolves a reproducible illegal memory access on sm_120.

Setup: Windows, llama.cpp master 465e49b9c plus the qwen4exp PR stack, CUDA 13.3, 2x RTX 5060 Ti (sm_120), -sm layer -ts 3,1. Model sh0wie/Qwen3.8-Flash-Next-REAP-288-GGUF Q4_K_M (qwen4exp, 288 experts, top-10, Q4_K), served with -ncmoe 27 so 27 layers of experts are host resident and take the MMQ ids path.

Before this patch, -ub 1024 -b 4096 runs correctly for a while and then dies mid-generation:

CUDA error: an illegal memory access was encountered
  current device: 0, in function ggml_backend_cuda_buffer_set_tensor at ggml/src/ggml-cuda/ggml-cuda.cu:792
  cudaStreamSynchronize(((cudaStream_t)0x2))

It killed the server about 1,100 generated tokens into the second task of a benchmark run, having completed the first task and a 1,380-token generation before that, which matches the "depends on how much slack the pool happens to have" description in #27792.

With this one-line change and no other difference (same build tree, same flags, same model), the identical benchmark runs to completion: ten tasks, 9.34/10, including the task that previously crashed the process. The default -ub 512 never crashed for us either way, so the exposure here really is ubatch-size dependent.

Worth noting for anyone hitting this: the same configuration is also a meaningful speedup, so the workaround of staying at -ub 512 has a real cost. A 96k-token prefill measured 140.4 tok/s at -ub 512 against 183.8 tok/s at -ub 1024, same placement.

[1jeffchristensen]

Follow-up with a correction, and a better argument for this patch than the one I gave.

Above I cited 140.4 tok/s at -ub 512 against 183.8 at -ub 1024 as the cost of the "just stay at 512" workaround. Both of those numbers were measured on the pre-patch binary. I had rebuilt in place at the same path, so my own run records could not tell the two builds apart, and I only caught it going back through the logs.

Separated by build, same model, same placement (-sm layer -ts 3,1, -ncmoe 26 at 512 and 27 at 1024, ctx 163840, q8_0 KV), same 95,904 token prompt, 2x RTX 5060 Ti:

binary	-ub 512	-ub 1024
master + qwen4exp stack, without this patch	140.4	183.8
same tree, with this patch	142.3	269.8

So this is not only a crash fix. It is worth about +47% prefill at -ub 1024 (183.8 to 269.8 tok/s), while being neutral at -ub 512 (+1.4%, inside run to run noise).

That shape matches the mechanism: the mis-sized src1_q8_1 tail padding is on the ids path that every host-resident expert takes, and the damage scales with the ubatch, which is also why the illegal access only ever appeared at 1024 and never at 512. Whatever the kernel is doing with the undersized buffer, it is not only unsafe, it is slow.

One caveat so this is not overclaimed. Page-cache warmth is not perfectly controlled between those two runs, because the expert weights are host resident and mmap'd. Each run was preceded by a multi-minute run over the same weights, and the -ub 512 row moved only 1.4% across the same rebuild, which bounds how much warmth can be worth on this setup. I would treat the 47% as approximate but well outside noise.

[1jeffchristensen]

Retracting the performance claim in my previous comment. This patch is performance-neutral on my setup; the +47% I reported was my own benchmarking error. The correctness finding stands and is unaffected.

What went wrong: the run I quoted at 269.8 tok/s was the only run in my whole series that executed a multi-minute generation workload against the same server before the throughput probe. With ~34 GB of expert weights host-resident on a 63 GB box, that pre-run is what makes the difference, not the patch. I attributed it to the rebuild because the rebuild happened between the two runs.

I have since measured the same configuration (-ub 1024 -b 4096 -ncmoe 27 -sm layer -ts 3,1, ctx 163840, q8_0 KV, same model, same 95,904-token prompt) five more times on the patched binary, varying only unrelated knobs:

patched, cold:  180.7, 189.5, 191.4, 189.4, 191.1 tok/s
patched, with a generation run first:            269.8 tok/s   <- the number I quoted
unpatched, cold:                                 183.8 tok/s

Unpatched cold 183.8 against patched cold ~187 is within run-to-run spread. The honest reading is no measurable throughput effect either way. Every arm above 195 tok/s in my data has the warm-up confound and no arm without it exceeded 194.

The related claim in my first comment, that staying at -ub 512 has a large cost, also needs correcting downward: measured cold on the patched binary it is 142.3 vs ~187 tok/s, so about +31%, not the ~+90% implied by comparing against the warmed run.

Sorry for the noise. The reason I am still glad this landed: without it, -ub 1024 dies with the illegal memory access described above, so the ubatch increase is not usable at all on this model. That is the argument for the patch, and it does not depend on any throughput number.

[homeofe]

Tested on Turing: v0.5.0 + this patch fixes a deterministic illegal memory access with MiMo-V2.6-Flash (host-resident experts, -ub 2048) on an RTX 2080 Ti, sm_75, CUDA 12.0. Stock crashed 3 of 3 times; patched ran 4 of 4 prefills (3.5K and 28K tokens) at 192-208 tok/s, vs 68 tok/s with the -ub 512 workaround. Details: #27792

[NullExpert]

Independent confirmation from another Windows / Blackwell setup.
Environment

• Windows 11 Pro for Workstations
• single RTX 5060 Ti 16 GB (sm_120)
• Qwen3.8-Flash-Next / qwen4exp GGUF, 512 experts / 10 active
• --cpu-moe --op-offload
• -b 4800 -ub 4800
• flash attention enabled
• llama.cpp tested across b11342 (f1cee99) and b11344 (ec7630a)
  On the unpatched b11342 build I had a deterministic first-request failure from a fresh server / empty prompt cache. The request contained 2219 tokens in total; its first processed prefill batch contained 2178 tokens. The process died during that initial prefill with a CUDA illegal-memory-access error. Other, substantially larger prompts could complete normally, which initially made this look unrelated to prompt size and is consistent with the pool-slack / ubatch-shape behavior described here and in CUDA MMQ mul_mat_id: src1_q8_1 tail padding computed from ne11 (== 1 for MoE) -> out-of-bounds read, illegal memory access for some ubatch sizes #27792.
  I first tested the padding fix proposed in CUDA MMQ mul_mat_id: src1_q8_1 tail padding computed from ne11 (== 1 for MoE) -> out-of-bounds read, illegal memory access for some ubatch sizes #27792 rather than applying this PR verbatim:
  ggml_cuda_mmq_get_J_max(..., std::max<int64_t>(ne_get_rows, 128))
  I want to be explicit about that distinction. For my actual failing case, the first processed batch was 2178 tokens and the model uses 10 experts, so the flattened ids-path row count is 2178 × 10 = 21780. The max(..., 128) therefore has no effect for this reproducer: the padding calculation is based on the same effective flattened row count (ne12 * n_expert_used) used by this PR.
  With that change applied, the previously failing request completes normally. After updating to b11344 (ec7630a), the vulnerable ne11 padding pattern was still present, so I rebuilt the CUDA backend with the same change and repeated the test.
  On b11344 the formerly failing cold request completed with:
• 2219 total prompt tokens, empty cache
• first processed 2178-token prefill batch: 395.81 tok/s
• complete 2219-token prompt evaluation: 338.35 tok/s
• generation: 27.87 tok/s
• normal slot release, no CUDA error
  I then ran multiple subsequent inference cycles, including contexts growing beyond 10k tokens, without reproducing the illegal memory access. Generation remained around 27–28 tok/s.
  As a basic regression check, I also smoke-tested the same patched CUDA backend with four other local Qwen-family GGUF models. All loaded, completed prompt evaluation and generation normally, and showed no CUDA errors. This is only a smoke test, not exhaustive validation.
  So while I have not tested this PR's exact one-line diff in isolation, I can independently confirm the underlying ids-path padding bug on a single RTX 5060 Ti / Windows configuration, and for the failing shape I tested, the CUDA MMQ mul_mat_id: src1_q8_1 tail padding computed from ne11 (== 1 for MoE) -> out-of-bounds read, illegal memory access for some ubatch sizes #27792 variant and this PR use the same effective flattened row count.
  Happy to provide the relevant llama-server logs if they would be useful for review.

[hclsys]

Same bug as #29847, which comes with a test-backend-ops reproducer: test_mul_mat_id(GGML_TYPE_Q4_0, GGML_TYPE_F32, 512, 10, b, 640, 508, 2560) for b = false and true.

On a GB10 (sm_121a, CUDA 13.0, master 11fe021) that case aborts with an illegal memory access without this change, and memcheck reports 485 errors. With the same change (ne12*n_expert_used for the padding) both cases pass and memcheck reports 0 errors. The full MUL_MAT_ID (933 cases, these two included) and MUL_MAT (1304) suites pass.

Maybe add those two cases here. They fault without the sanitizer on this box and on the reporter's, so test-backend-ops would catch it on those setups.

[JohannesGaessler]

Sorry for the long radio silence, I took a break from llama.cpp and didn't see this in my backlog. The fix in this PR looks 90% correct to me, can you check whether #29941 also works?

[hclsys]

Yes, #29941 works here. On a GB10 (sm_121a, CUDA 13.0) with ne12 in the J_max call, both reproducer cases pass and memcheck reports 0 errors (without the change: abort, 485 errors). MUL_MAT_ID 933/933 and MUL_MAT 1304/1304 pass. Same result as the ne12*n_expert_used version.

[JohannesGaessler]

I'm closing this PR in favor of the other one, assuming the issue is now fixed.

[glennneuber]

@JohannesGaessler, thanks for looking into this. Here are our test results:

Table for #27044 (sm_120, RTX PRO 6000 Blackwell, CUDA 13.0, llama.cpp master 0504396; compute-sanitizer --tool memcheck, src1 buffer in its own exact-size allocation; memcheck errors, ✓ test passed, ✗ aborted):

case	master (ne11)	#29941 (ne12)	#27044 (ne12*n_expert_used)
65 tokens, 576 rows (fallback), q4_K, 256 experts, 8 used	3,041 ✗	3,229 ✗	0 ✓
100 tokens, 512 rows, q4_K, 256 experts, 8 used	6,693 ✗	557 ✗	0 ✓
100 tokens, #29847's shape (q4_0, 512 experts, 10 used)	7,341 ✗	53 ✗	0 ✓
2040 tokens, the original fault's shape	101 ✗	0 ✓	0 ✓
508 tokens, #29847, b=0	913 ✗	0 ✓	0 ✓
508 tokens, #29847, b=1	2,397 ✗	0 ✓	0 ✓

Host-only check, 2,439,360 shapes on sm_75/86/89/120: #29941 short in 268,440 (all below 128 tokens), #27044 in none.
Reproduction: MaxusAI/ollama#445
GPU-free test: https://github.com/MaxusAI/ollama/blob/main/docs/maxusai/tasks/mmq-ids-padding-test.cu

[JohannesGaessler]

I'm not seeing any compute-sanitizer errors with

    for (bool b : {false, true}) {
        test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_Q4_0, GGML_TYPE_F32, 512, 10, b, 640, 100, 2560));
    }

[hclsys]

Reproduced on a GB10 (sm_121a) with src1_q8_1 in its own exact-size cudaMalloc, memcheck, master dd26678 (includes #29941). q4_0, 512 experts, 10 used, m=640, k=2560:

tokens	master (ne12)	ne12*n_expert_used
100, b=0	1194 errors	0
100, b=1	274 errors	0
508, b=0 and b=1	0	0

q4_K, 256 experts, 8 used, n=65 and n=100: 0 with both here, I did not hit the numbers in the table above for those shapes.

Cause for n=100: mul_mat_q_switch_J picks J=112 (104 is skipped), while ggml_cuda_mmq_get_J_max(ne12=100) returns 96, so the padding is 16 columns short. J_max rounds down to a valid J, the picker rounds up. With ne_get_rows it is 128.

Also 0 errors on all 8 cases with J_max(..., 128), a flat max tile width that does not depend on ne12. GGML_PAD(ne12, 8) is not enough, J_max stays at 96.

[JohannesGaessler]

Yes, looking at the code in ggml_cuda_mmq_get_J_max I've found that I implemented in a way inconsistent with the use for kernel selection.

[JohannesGaessler]

Since I am unable to reproduce the issue locally, can either one of you please check #29953 ?