CUDA MMQ: mul_mat_id reads past src1_q8_1 when the column count is not a multiple of J #28151

$\textcolor{red}{\textbf{(USER WAS BANNED FOR THIS POST)}}$

Summary

In the CUDA MMQ path for GGML_OP_MUL_MAT_ID, the y tile is loaded over a full J-column
tile regardless of how many columns actually remain. tile_y_max_j is computed but only
reaches write_back, never the load. When the column count is not a multiple of the
selected J, the final tile reads past the end of src1_q8_1.

The tail padding of that buffer does not cover the overhang: it is sized from ne11,
which in the mul_mat_id path is n_expert_used, while mul_mat_q_switch_J picks J
from ncols_max = ne12.

This is a read past the buffer. Results are unaffected — write_back is bounded by
tile_y_max_j, so the extra columns are computed and discarded. The only symptom is a
fault when the read crosses an unmapped page.

Reproducible with test-backend-ops alone — no model, no GGUF.

Where

ggml/src/ggml-cuda/mmq.cuh, mul_mat_q_process_tile:

load_tiles(x, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x);   // x: bounded by rows
for (int l0 = 0; l0 < J * MMQ_TILE_Y_K; l0 += nwarps*warp_size) {
    tile_y[l] = by0[l];                                              // y: full J, unbounded
}
...
write_back(sum, ids_dst, dst, y_scale, stride_col_dst, tile_x_max_i, tile_y_max_j);
//                                                                  ^^^^^^^^^^^^ only here

The x side receives its bound at load time (tile_x_max_i); the y side only at write-back.

ggml/src/ggml-cuda/mmq.cu (mul_mat_id path):

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/src/ggml-cuda/mmq.cuh:

static __host__ int ggml_cuda_mmq_get_J_max(..., const int64_t ne11) {
    int ret = std::min(ne11, int64_t(512));
    ret -= ret % 8;
    ...

In this path ne11 == n_expert_used (the function asserts ne1 == n_expert_used a few
lines above). With n_expert_used = 10 the padding is J_max = 8 → 8 * 144 = 1152
bytes, while J = 128 is selected from ncols_max = ne12 = 257.

Measured against three allocations, with a log of the pool allocations:

The difference is exactly J_max * sizeof(block_q8_1_mmq) = 8 * 144.

In the non-mul_mat_id path ne11 is the column count and the two agree; in the
mul_mat_id path the axes swap roles and the padding formula was not adjusted.

Reproduction

Test cases only, no library changes. Patch below.

./build/bin/test-backend-ops perf -o MUL_MAT_ID -b CUDA0

MUL_MAT_ID(type_a=iq4_nl, type_b=f32, n_mats=512, n_used=10, b=0, m=2560, n=257, k=640)
  -> CUDA error: an illegal memory access was encountered

Reproducible 3/3 runs, aborting at the same point each time (35 cases pass, 36 shapes
reached). The full case list is part of the reproduction — see the note below.

The preceding cases are part of the reproduction. Measured with -p filtering:

-p 'n_mats=512'      36 cases   exit 134   fault
-p 'm=640,n=257'      1 case    exit 0     passes
-p 'm=2560,n=257'     1 case    exit 0     passes

Either shape run alone passes. The fault requires the preceding cases in the same
process — the out-of-bounds read otherwise lands in mapped memory and is discarded
silently. A shortened case list gives a clean run that proves nothing.

compute-sanitizer --tool memcheck:

Invalid __global__ read of size 4 bytes
  at void mul_mat_q<(ggml_type)12, (int)128, (bool)0>
  Host Frame: launch_mul_mat_q<(ggml_type)12, 128, false>
  Address 0x... is out of bounds

(ggml_type)12 is GGML_TYPE_Q4_K. Note that the case name printed in the run log is
the iq4_nl line above: CUDA reports asynchronously, so the log line is the next
synchronisation point, not necessarily the originating call. We have not narrowed which
case issues the faulting launch; the kernel identity above is what compute-sanitizer
reports.

CUDA_LAUNCH_BLOCKING=1 does not move the report to a different case: the log line stays
on the iq4_nl case while compute-sanitizer names a Q4_K kernel. The blocking flag
makes the launch synchronous, not the point at which cudaGetLastError is called, so a
sticky error can still surface on the next test. We have not resolved which case issues
the faulting launch.

The overrun size lies between the 1152-byte padding and one full tile
(128 * 144 = 18432 bytes).

Scope

Not specific to a quantization type or to k:

• three different k (640 / 768 / 1024) fail at the same case
• n = 255 and n = 256 pass (two full tiles); n = 257 fails (third tile, one real column)
• n = 129 has the same overhang (129 % 128 == 1) and does not fault — the read
  still goes past the buffer, it just lands in mapped memory

The overrun happens whenever the column count is not a multiple of the selected J; it
only faults when it crosses an unmapped page.

Suggested fix

Two levels:

1. One line — size the padding from the quantity that drives the J selection:

-        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, ne12) * sizeof(block_q8_1_mmq);

Padding goes from 1152 to 18432 bytes — 0.23 % of a 7.4 MB buffer. This makes the
overrun harmless; it does not remove it.

2. The cause — bound the y tile load by tile_y_max_j, the way the x load is bounded
   by tile_x_max_i.

Environment

GPU     NVIDIA RTX PRO 4500 Blackwell, compute capability 12.0, 82 SMs
Driver  610.43.02
CUDA    nvcc 12.8.93, CMAKE_CUDA_ARCHITECTURES=120
Commit  0eadefebd3f8f92a86d634a0e5b8fffc9dc792c0

diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 996b88db2..d4de4b899 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -9229,6 +9229,38 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_Q8_0, GGML_TYPE_F32, 128, 128, false, 8192, 1, 5120)); // Llama-4-Maverick-17B-128E-PAB-Q8_0
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 8192, 1, 5120, {128, 1}, {1, 1}));
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 8192, 512, 5120, {128, 1}, {1, 1}));
+
+    // MUL_MAT_ID shapes with 512 experts, 10 used.
+    //
+    // NOTE ON THE n SERIES: the failure needs the *preceding* cases.  Running only the
+    // failing shape in isolation does not crash -- the out-of-bounds read then lands in
+    // memory that is still mapped and silently returns garbage instead of faulting.
+    // The full series below is therefore part of the reproduction, not decoration.
+    // Shortening it produces a clean run that proves nothing.
+    //
+    // Observed on RTX PRO 4500 Blackwell (sm_120, 82 SMs), CUDA 12.8:
+    //   IQ4_NL m=2560 n=257 k=640/768/1024 -> illegal memory access, same case each time
+    //   IQ4_NL m=2560 n=255 and n=256      -> pass
+    //   Q4_K   m= 640 n=257 k=2560         -> pass
+    //   compute-sanitizer: "Invalid __global__ read of size 4 bytes" in
+    //                      mul_mat_q<(ggml_type)12, (int)128, (bool)0>
+    //
+    // n = 257 is the first column count above 256, where ntx goes from 2 to 3.  But ntx
+    // makes the same jump in the passing Q4_K case, so the column count alone is not it.
+    // k has been excluded by measurement (three different k, same failing case).
+    //
+    // The last shape dimension still entangled with the type is m: the failing cases all
+    // have m = 2560 (nty = 20), the passing ones m = 640 (nty = 5).  The Q4_K line with
+    // m = 2560 below crosses that -- it is valid because Q4_K needs k % 256 == 0, which
+    // 2560 satisfies.
+    for (int64_t n : {1, 8, 9, 16, 30, 31, 32, 33, 48, 63, 64, 65, 127, 128, 129, 255, 256, 257, 258, 320, 508, 511, 512, 513}) {
+        test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_IQ4_NL, GGML_TYPE_F32, 512, 10, false, 2560, n,  640)); // reported failure
+        test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_Q4_K,   GGML_TYPE_F32, 512, 10, false,  640, n, 2560)); // gate/up shape, passes
+    }
+    for (int64_t n : {255, 256, 257, 258, 320, 512}) {
+        test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_IQ4_NL, GGML_TYPE_F32, 512, 10, false, 2560, n, 1024)); // k varied, still fails
+        test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_Q4_K,   GGML_TYPE_F32, 512, 10, false, 2560, n, 2560)); // m crossed against the type
+    }
 #endif
 
     for (ggml_type type_a : all_types) {

Possibly related, different symptom: #19672 (assert in mm_ids_helper, same expert
configuration). Reverting #27978 does not change the behaviour reported here.