Skip to content

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

Open

Description

@Tonoken3

Summary

In the MoE path of ggml_cuda_mul_mat_q (ggml/src/ggml-cuda/mmq.cu), the tail padding of the src1_q8_1 buffer is computed from ne11, which is always 1 for MUL_MAT_ID (src1 is [ne10, 1, n_tokens]). ggml_cuda_mmq_get_J_max(type, fallback, cc, 1) returns 0, so the buffer gets no tail padding, while mul_mat_q unconditionally loads a full J-column tile of y for the last tile of the last used expert. The kernel therefore reads up to (J-1) * sizeof(block_q8_1_mmq) bytes past the end of the allocation (≈18 KB for J=128).

Whether this is observable depends on how much slack the VMM pool happens to have after that allocation: usually the read lands inside mapped pool memory and nothing happens; for some ubatch sizes the slack is a few KB and the kernel dies with CUDA error: an illegal memory access was encountered.

This is still present on master (fc35562ba4, "cuda: unblock mmq for MoE on sm_60"), line 206:

    const size_t nbytes_src1_q8_1 = ne12*n_expert_used*ne10_padded * y_block_size/y_values_per_block +
        ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, ne11) * sizeof(block_q8_1_mmq);   // ne11 == 1 for mul_mat_id

The dense path (line 137) is fine because there ne11 is the real column count.

Environment

  • RTX PRO 2000 Blackwell (sm_120), CUDA 13, Linux. Build: GGML_CUDA=ON, CMAKE_CUDA_ARCHITECTURES=120.
  • MoE model with experts kept on the host (--cpu-moe, i.e. -ot exps=CPU), so for ubatches ≥ 32 tokens the scheduler copies the used experts to the GPU and runs MUL_MAT_ID with MMQ (Q4_K experts, 512 experts, top-10, ne00=2560, ne01=640).
  • Observed on a Qwen4-exp checkpoint (PR model: add Qwen3.8-Flash-Next (qwen4exp) #27742 tree), but the buffer math is model-independent; any MoE model that takes the MMQ mul_mat_id path with host-resident experts should be affected. Which ubatch sizes crash depends on the pool layout for that model.

Reproduction (what we measured)

  1. Start llama-server with the MoE model and --cpu-moe (fresh process).
  2. Send a prompt whose first ubatch is 508 tokens (our fork splits a 512-token prompt into [508, 4]; on master use -ub 508 with a longer prompt, or otherwise arrange a 508-token ubatch). Result: ggml-cuda.cu:107: CUDA error: an illegal memory access was encountered in MUL_MAT_ID of the first MoE layer. Same with 1016 tokens. 507, 509, 1021, 2044 etc. are fine.
  3. The crash is independent of the prompt content (any 512-token slice crashes) and of flash attention.
  4. Pool test: in a fresh process, first process a 1021-token ubatch, then the 512-token prompt → no crash. The extra earlier allocation grows the VMM pool so the out-of-bounds read stays inside mapped memory. This is what makes the bug look size-specific.
  5. Keeping the experts on the CPU for prefill (GGML_OP_OFFLOAD_MIN_BATCH=999999) avoids the path entirely (and, on this box, costs no prefill speed).

Arithmetic for the 508 case (Q4_K, n_expert_used=10): ne_get_rows = 5080, src1_q8_1 = 5080·2560·144/128 = 14,630,400 B; pool allocations before it (ids_src1, ids_dst, expert_bounds) put the buffer end 6,784 B below the 2 MiB pool granularity boundary, while the last tile of the highest used expert reads up to 127·144 = 18,288 B past the end. For 507 the slack is 35,584 B and nothing is observed. compute-sanitizer --tool memcheck should report the invalid global read for both sizes.

Fix (one line)

Size the padding from the actual number of y columns the kernel can touch (ne_get_rows), never from ne11:

--- a/ggml/src/ggml-cuda/mmq.cu
+++ b/ggml/src/ggml-cuda/mmq.cu
@@ -203,7 +203,7 @@ void ggml_cuda_mul_mat_q(
     const size_t nbytes_src1_q8_1 = ne12*n_expert_used*ne10_padded * y_block_size/y_values_per_block +
-        ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, ne11) * sizeof(block_q8_1_mmq);
+        ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, std::max<int64_t>(ne_get_rows, 128)) * sizeof(block_q8_1_mmq);

With this patch the 508/1016-token cases pass and the outputs of the previously working sizes are unchanged (the padding is never read into results). Happy to open a PR if you prefer.

Activity

  1. vkunde commented on Aug 29, 2026

    @vkunde

    Retracted — this was not the same bug. Sorry for the noise.

    I originally reported an illegal memory access with qwen4exp (Qwen3.8-Flash-Next) on Pascal (GTX 1060, sm_61) and attributed it to this issue. That attribution was wrong.

    I have since built llama.cpp from the exact same commit the failing binary was built from (18443257), targeting sm_61 with CUDA 12.9, and it runs fine — both --cpu-moe and -ngl 0, no crash:

    build --cpu-moe -ngl 0
    third-party redistribution (CUDA 11, patched fork) crash crash
    same, with GGML_CUDA_NO_PINNED=1 crash crash
    upstream 18443257, CUDA 12.9, sm_61 ok (57.6 pp / 9.4 tg) ok (146.8 pp / 5.6 tg)

    So upstream at that commit is not affected on Pascal, and pinned host memory is not the trigger either. The problem is in the third-party build I was using, not in mmq.cu. I am taking it to that vendor instead.

    The one thing here that may still be worth keeping: upstream qwen4exp + --cpu-moe works on sm_61, in case a Pascal data point is useful to anyone.

  2. manunicholasjacob commented on Sep 2, 2026

    @manunicholasjacob

    Ran the ids path under compute-sanitizer with an exact-size pool on a P100, T4, A100 and L4; results, and two gaps the one-line fix leaves (rows below 8, and the ids_dst tile load), are in #27044 (comment).

  3. pradeep1883 commented on Sep 24, 2026

    @pradeep1883

    Confirming this on different hardware and weight type, still present on current master.

    Environment

    • 2× RTX A4000 16 GB (Ampere, sm_86), driver 580.173.02, CUDA 12.8 toolkit, Linux (Pop!_OS)
    • master 8212c78 (2026-09-24) and release b11054 (7d4b92b), built with GGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=86
    • Model: Qwen3.6-35B-A3B, Q8_0 (unsloth GGUF), llama-server --fit on -c 65536 (fit keeps part of the experts on the host)

    Symptom
    CUDA error: an illegal memory access was encountered during prefill of one specific 777-token prompt, every time, on a fresh server. Other prompts (e.g. 1666 tokens) are fine, matching the pool-slack explanation above. Reproduces with one GPU (-sm none, --n-cpu-moe 40) and with two; with -ctk/-ctv q8_0 and with f16 KV; with -ub 512.

    compute-sanitizer (memcheck), same prompt:

    Invalid __global__ read of size 4 bytes
        at void mul_mat_q<(ggml_type)8, (int)128, (bool)0>(...)+0x2220
        by thread (0,3,0) in block (1016,0,0)
        Address 0x139e800150 is out of bounds
        and is 337 bytes after the nearest allocation at 0x139e600000 of size 2,097,152 bytes
    Host Frame: launch_mul_mat_q<(ggml_type)8, 128, false>
    Host Frame: ggml_cuda_mul_mat_q
    Host Frame: ggml_backend_cuda_graph_compute
    

    So it is not specific to Q4_K / Blackwell: Q8_0 experts on sm_86 hit the same unpadded src1_q8_1 tail read.

    Workaround that works here: building with -DGGML_CUDA_FORCE_CUBLAS=ON. With it the same prompt and all our benchmarks run cleanly (flash attention on, q8_0 KV), ~300 tok/s prefill and ~40 tok/s generation for this model. Happy to test #27044 on this hardware if useful.

  4. pwasiewi commented on Oct 2, 2026

    @pwasiewi

    Same bug on Blackwell; our numbers match the arithmetic in the issue.

    RTX 5070 Ti 16 GB (sm_120), driver 615.71.09 open, CUDA 13, Ryzen 9 5950X, 64 GB. Build: master 25747b0 (b11284) plus PR #27861 v3 (expert cache; no ggml-cuda changes, crash reproduces with it off). Models: Qwen3.5/3.6/3.8 35B-A3B (n_embd 2048, 8 of 256 experts), Q4_K_M and IQ4_XS, --n-cpu-moe 20..32 -b 2048 -ub 2048 --load-mode none.

    Symptom: about one cold start in eight dies on the first ~1K-token prompt, 30-40 ms after the slot launches, with CUDA error: an illegal memory access was encountered in ggml_backend_cuda_synchronize. dmesg has 20 identical Xid 31 ... FAULT_PDE ACCESS_TYPE_VIRT_READ entries at VA 0x3_03200000 = VMM pool base (0x302000000 in /proc/pid/maps) + 18 MiB, the first unmapped 2 MiB page past the pool. A server that survives one 2048-token ubatch (pool at 38 MiB) never faults again.

    Arithmetic for this family: a gate/up ubatch takes 18,432 B of q8_1 per token (8 x 2048 x 144/128) plus 64 B of ids, so alloc_end(n) = 18,496 n + 1,152; the pool maps in 2 MiB steps. n = 1020 ends 7,296 B below the 18 MiB boundary while the unpadded J=128 y-tile reads up to 18,288 B past the end; 1019 leaves 25,728 B, 1021 grows the pool to 20 MiB. Our prompt lands on 1019..1028 tokens (a uuid line varies), hence the ~1/8.

    Deterministic version (907 tokens leave 128 B of slack):

    llama-server -m Ornith-1.5-35B-Q4_K_M.gguf -ngl 99 --n-cpu-moe 28 -c 131072 -np 1 \
      -fa on -ctk q8_0 -ctv q8_0 --load-mode none -b 2048 -ub 907
    # then one prompt >= 907 tokens: 3/3 cold starts died
    

    With the mmq.cu hunk from #27044 applied: 0/3 at -ub 907, then 0/8 cold starts at -ub 2048. 49K-token pp 2240 -> 2292 t/s, tg 71.3 -> 70.4, no measurable change.

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

Metadata

Metadata

Assignees

No one assigned

    Labels

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions