diff --git a/docs/maxusai/tasks/mmq-29953-guard-timing.sh b/docs/maxusai/tasks/mmq-29953-guard-timing.sh
new file mode 100755
index 00000000000..fb5ded4097f
--- /dev/null
+++ b/docs/maxusai/tasks/mmq-29953-guard-timing.sh
@@ -0,0 +1,34 @@
+#!/usr/bin/env bash
+# What the y-tile guard costs at compile time: mmq.cu on #29953's head (3070d927f) with the guard patch, nvcc 12.8,
+# sm_120, built REPS times with the guard on and REPS times with it compiled out, interleaved so drifting host load
+# hits both alike. Prints each time and the medians. One build is too noisy to quote (first run 9.8 vs 8.9 s, second
+# 11.2 vs 7.9 s).
+set -u
+cd "$(dirname "$0")" || exit 1
+HERE=$PWD
+REPS=${REPS:-5}
+WT=$HERE/wt-29953-guard-timing
+OUT=$HERE/guard-timing-results
+CUDA_HOME=${CUDA_HOME:-/usr/local/cuda-12.8}
+unset CPATH C_INCLUDE_PATH CPLUS_INCLUDE_PATH
+rm -rf "$WT" "$OUT"; mkdir -p "$OUT"; git -C llama.cpp worktree prune
+git -C llama.cpp worktree add --quiet --detach "$WT" 3070d927f6c172242d56a2c1fb897e838e7b2632 || exit 1
+trap 'cd "$HERE"; git -C llama.cpp worktree remove --force "$WT" 2>/dev/null' EXIT
+cd "$WT" && git apply "$HERE/mmq-29953-y-tile-guard.patch" || exit 1
+CU=ggml/src/ggml-cuda/mmq.cu
+cp "$CU" "$OUT/on.cu"
+sed 's/^#    define GGML_CUDA_MMQ_Y_TILE_GUARD_ON$/\/\/ off/; s/^#        define GGML_CUDA_MMQ_Y_TILE_GUARD_ON$/\/\/ off/' "$OUT/on.cu" > "$OUT/off.cu"
+cmp -s "$OUT/on.cu" "$OUT/off.cu" && { echo "the off variant is identical to on: sed matched nothing"; exit 1; }
+echo "host load at start (1/5/15 min): $(cut -d' ' -f1-3 /proc/loadavg), $(nproc) cores"
+build() { cp "$OUT/$1.cu" "$CU"; local t0 t1; t0=$(date +%s.%N)
+    "$CUDA_HOME/bin/nvcc" -Wno-deprecated-gpu-targets -std=c++17 -gencode arch=compute_120,code=sm_120 \
+        -I. -Iggml/include -Iggml/src -Iggml/src/ggml-cuda -c "$CU" -o "$OUT/mmq.o" > "$OUT/$1.log" 2>&1 || { echo "build $1 failed"; exit 1; }
+    t1=$(date +%s.%N); awk -v a="$t0" -v b="$t1" 'BEGIN{printf "%.2f", b-a}'; }
+on=(); off=()
+for r in $(seq 1 "$REPS"); do
+    a=$(build on); b=$(build off); on+=("$a"); off+=("$b")
+    printf 'rep %d: guard on %ss, off %ss\n' "$r" "$a" "$b"
+done
+med() { printf '%s\n' "$@" | sort -n | awk '{v[NR]=$1} END {print (NR%2 ? v[(NR+1)/2] : (v[NR/2]+v[NR/2+1])/2)}'; }
+echo "median: guard on $(med "${on[@]}")s, off $(med "${off[@]}")s, over $REPS interleaved reps"
+echo "host load at end (1/5/15 min): $(cut -d' ' -f1-3 /proc/loadavg)"
diff --git a/docs/maxusai/tasks/mmq-29953-guard-verify.sh b/docs/maxusai/tasks/mmq-29953-guard-verify.sh
new file mode 100755
index 00000000000..ca1c10f4bea
--- /dev/null
+++ b/docs/maxusai/tasks/mmq-29953-guard-verify.sh
@@ -0,0 +1,102 @@
+#!/usr/bin/env bash
+# A compile-time guard for MMQ's y tile on top of llama.cpp#29953's head (3070d927f): material for upstream, GPU-free.
+#
+# mmq-29953-y-tile-guard.patch routes the src1 padding and the shared-memory y tile through one helper,
+# ggml_cuda_mmq_get_nbytes_y_tile(), and asserts at compile time -- for every config table and type -- that the
+# helper covers the y load loop's extent, written from the loop with exact ceiling division. Five nvcc builds of
+# mmq.cu for sm_120, each scored by WHY it failed, not just whether:
+#   1. the head as is, unpatched          -> must compile (baseline time)
+#   2. the head with the patch            -> must compile
+#   3. helper shrunk to J blocks          -> must FAIL by the guard, naming every table (the published #29953's rule)
+#   4. helper rounded to one warp's ints  -> must FAIL by the guard, naming every table
+#   5. one blackwell config at 192 threads-> must FAIL by the guard naming ONLY blackwell: GGML_PAD is a bit mask
+#                                            that assumes a power-of-two granularity, the tables only assert
+#                                            nthreads % 32 == 0, and the guard's exact arithmetic sees the gap
+# nvcc's front end stops at 100 errors by default, which cut mutation 3's table list to the first five tables;
+# --error_limit lifts that so every table that fails is named.
+# Usage: mmq-29953-guard-verify.sh   (from its own directory, with a llama.cpp clone at ./llama.cpp that has
+#        3070d927f fetched; results go to ./verify-29953-guard-results)
+set -u
+cd "$(dirname "$0")" || exit 1
+HERE=$PWD
+PATCH=$HERE/mmq-29953-y-tile-guard.patch
+WT=$HERE/wt-29953-guard-verify
+OUT=$HERE/verify-29953-guard-results
+CUDA_HOME=${CUDA_HOME:-/usr/local/cuda-12.8}
+HEAD29953=3070d927f6c172242d56a2c1fb897e838e7b2632
+unset CPATH C_INCLUDE_PATH CPLUS_INCLUDE_PATH
+rm -rf "$OUT"; mkdir -p "$OUT"
+echo "host load at start (1/5/15 min): $(cut -d' ' -f1-3 /proc/loadavg), $(nproc) cores"
+
+rm -rf "$WT"; git -C llama.cpp worktree prune
+git -C llama.cpp worktree add --quiet --detach "$WT" "$HEAD29953" || exit 1
+trap 'cd "$HERE"; git -C llama.cpp worktree remove --force "$WT" 2>/dev/null' EXIT
+cd "$WT" || exit 1
+CU=ggml/src/ggml-cuda/mmq.cu; CUH=ggml/src/ggml-cuda/mmq.cuh; BW=ggml/src/ggml-cuda/mmq-config-blackwell.cuh
+
+build() { # <label> <expect ok|fail> [expected table list, space-separated, sorted]
+    local label=$1 expect=$2 want_tables=${3:-}
+    local log="$OUT/$label.log" t0 t1 rc
+    t0=$(date +%s.%N)
+    "$CUDA_HOME/bin/nvcc" -Wno-deprecated-gpu-targets -std=c++17 -gencode arch=compute_120,code=sm_120 \
+        -Xcudafe --error_limit=100000 \
+        -I. -Iggml/include -Iggml/src -Iggml/src/ggml-cuda -c "$CU" -o "$OUT/mmq.o" > "$log" 2>&1
+    rc=$?
+    t1=$(date +%s.%N)
+    local got=ok; [ $rc -ne 0 ] && got=fail
+    local tables; tables=$(grep -oE 'in the [a-z0-9_]+ table' "$log" | sed 's/in the //;s/ table//' | sort -u | tr '\n' ' ' | sed 's/ $//')
+    # count errors, not lines: nvcc prints each failed static_assert on two
+    local nerr; nerr=$(grep -cE ': error' "$log")
+    local nguard; nguard=$(grep -c ': error: static assertion failed with "MMQ: the y tile helper' "$log")
+    local verdict=UNEXPECTED
+    if [ "$got" = "$expect" ]; then
+        if [ "$expect" = ok ]; then verdict=as-expected
+        elif [ "$nguard" -gt 0 ] && [ "$nerr" -eq "$nguard" ] && [ "$tables" = "$want_tables" ]; then verdict=as-expected
+        else verdict=WRONG-REASON
+        fi
+    fi
+    printf '%-34s rc=%-3s %-4s want=%-4s %-13s %5.1fs\n' "$label" "$rc" "$got" "$expect" "$verdict" "$(awk -v a="$t0" -v b="$t1" 'BEGIN{print b-a}')"
+    [ "$got" = fail ] && printf '    errors=%s, from the guard=%s, tables named: %s\n' "$nerr" "$nguard" "${tables:-none}"
+    [ "$verdict" = as-expected ]
+}
+
+ALL="ampere blackwell cdna gcn pascal_dp4a pascal_older rdna2 rdna3 rdna3_5 rdna4"
+ok=1
+build "1-head-unpatched" ok || ok=0
+
+git apply "$PATCH" || { echo "FAIL: the patch does not apply to $HEAD29953"; exit 1; }
+echo "patch applies to a pristine $HEAD29953"
+cp "$CU" "$OUT/mmq.cu.patched"; cp "$CUH" "$OUT/mmq.cuh.patched"; cp "$BW" "$OUT/blackwell.orig"
+build "2-head-with-guard" ok || ok=0
+
+swap() { # <file> <old> <new>: exact one-occurrence string swap, so a mutation cannot fail for a typo of mine
+    python3 - "$1" "$2" "$3" <<'PY' || exit 1
+import sys
+p, old, new = sys.argv[1:4]
+s = open(p).read()
+if s.count(old) != 1:
+    sys.exit("mutation anchor found %d times in %s, expected once: %r" % (s.count(old), p, old))
+open(p, "w").write(s.replace(old, new))
+PY
+}
+HELPER='    return GGML_PAD(config.J*sizeof(block_q8_1_mmq), config.nthreads*sizeof(int));'
+
+swap "$CUH" "$HELPER" '    return config.J*sizeof(block_q8_1_mmq);'
+build "3-helper-shrunk-to-J-blocks" fail "$ALL" || ok=0
+cp "$OUT/mmq.cuh.patched" "$CUH"
+
+swap "$CUH" "$HELPER" '    return GGML_PAD(config.J*sizeof(block_q8_1_mmq), 32*sizeof(int));'
+build "4-helper-rounded-to-one-warp" fail "$ALL" || ok=0
+cp "$OUT/mmq.cuh.patched" "$CUH"
+
+# the whole first CASE line with 256 threads: the type alone repeats once per J, so it is not a unique anchor
+first=$(grep -m1 -E '^ *CASE\(GGML_TYPE_[A-Z0-9_]+, *256,' "$BW")
+swap "$BW" "$first" "${first/256,/192,}"
+echo "mutation 5 rewrote: $first -> ${first/256,/192,}"
+build "5-blackwell-config-at-192-threads" fail "blackwell" || ok=0
+cp "$OUT/blackwell.orig" "$BW"
+
+sed -i 's/^#    define GGML_CUDA_MMQ_Y_TILE_GUARD_ON$/\/\/ guard disabled for timing/; s/^#        define GGML_CUDA_MMQ_Y_TILE_GUARD_ON$/\/\/ guard disabled for timing/' "$CU"
+build "6-guard-compiled-out" ok || ok=0
+
+echo "VERIFY29953GUARD COMPLETE $(date -u +%FT%TZ) all_as_expected=$ok"
diff --git a/docs/maxusai/tasks/mmq-29953-y-tile-guard.patch b/docs/maxusai/tasks/mmq-29953-y-tile-guard.patch
new file mode 100644
index 00000000000..729f22c2727
--- /dev/null
+++ b/docs/maxusai/tasks/mmq-29953-y-tile-guard.patch
@@ -0,0 +1,129 @@
+diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu
+index 027a7e60..6923d84a 100644
+--- a/ggml/src/ggml-cuda/mmq.cu
++++ b/ggml/src/ggml-cuda/mmq.cu
+@@ -4,6 +4,7 @@
+ #include "mmid.cuh"
+ 
+ #include <cstdint>
++#include <utility>
+ 
+ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream, const ggml_prec prec_src1) {
+     switch (args.type_x) {
+@@ -187,8 +188,8 @@ void ggml_cuda_mul_mat_q(
+     const size_t y_block_size       = use_native_fp4 ? sizeof(block_fp4_mmq) : sizeof(block_q8_1_mmq);
+     const size_t y_values_per_block = use_native_fp4 ? QK_FP4_MMQ            : QK8_1_MMQ;
+ 
+-    int J_best        = 0;
+-    int nthreads_best = 0;
++    int    J_best             = 0;
++    size_t nbytes_y_tile_best = 0;
+     {
+         int64_t ncols_opt = ne11;
+         if (ids) {
+@@ -218,7 +219,7 @@ void ggml_cuda_mul_mat_q(
+ 
+             if (ntiles_x < ntiles_J_best) {
+                 J_best = J;
+-                nthreads_best = config.nthreads;
++                nbytes_y_tile_best = ggml_cuda_mmq_get_nbytes_y_tile(config);
+                 ntiles_J_best = ntiles_x;
+             }
+         }
+@@ -226,10 +227,9 @@ void ggml_cuda_mul_mat_q(
+     GGML_ASSERT(J_best > 0);
+ 
+     // A tile of size J can read in at most J - 1 extra columns.
+-    // For simplicity, round up the padding of a full tile to a multiple of the number of bytes that nthreads can load in parallel.
+-    const size_t src1_load_chunk_size = nthreads_best * sizeof(int);
+-    const size_t src1_q8_1_padding = ((J_best * sizeof(block_q8_1_mmq) + src1_load_chunk_size - 1) / src1_load_chunk_size)
+-        * src1_load_chunk_size;
++    // Its y load copies the whole padded shared-memory y tile, so pad by that tile: the same helper sizes it in
++    // mmq_get_nbytes_shared(), so the padding cannot drift from what the kernel reads.
++    const size_t src1_q8_1_padding = nbytes_y_tile_best;
+ 
+     if (!ids) {
+         const size_t nbytes_src1_q8_1 = ne13*ne12 * ne11*ne10_padded * y_block_size/y_values_per_block + src1_q8_1_padding;
+@@ -486,3 +486,58 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
+ 
+     return (!GGML_CUDA_CC_IS_CDNA(cc)) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
+ }
++
++// Compile-time guard on the y tile: for every config table and type, ggml_cuda_mmq_get_nbytes_y_tile() -- which
++// sizes both the shared-memory y tile and the global src1 padding -- must cover the y load in
++// mul_mat_q_process_tile, which runs l0 over [0, J*MMQ_TILE_Y_K) in steps of nthreads ints with no bound on l.
++// The requirement is written from that loop with exact ceiling division, not from the helper or GGML_PAD, so a
++// change to either side that makes them disagree fails the build for every architecture's table at once, with no
++// GPU. One static_assert per (table, type) keeps each constant evaluation small. HIP keys off
++// __HIP_DEVICE_COMPILE__ because vendors/hip.h defines __CUDA_ARCH__ in every HIP pass, host included.
++#if defined(GGML_USE_HIP)
++#    if !defined(__HIP_DEVICE_COMPILE__)
++#        define GGML_CUDA_MMQ_Y_TILE_GUARD_ON
++#    endif
++#elif !defined(GGML_USE_MUSA) && !defined(__CUDA_ARCH__)
++#    define GGML_CUDA_MMQ_Y_TILE_GUARD_ON
++#endif
++#ifdef GGML_CUDA_MMQ_Y_TILE_GUARD_ON
++#define GGML_CUDA_MMQ_Y_TILE_GUARD(TABLE)                                                                               \
++    static constexpr bool ggml_cuda_mmq_y_tile_ok_##TABLE(const int t) {                                                \
++        for (int fb = 0; fb < 2; ++fb) {                                                                                \
++            for (int J = 8; J <= 128; J += 8) {                                                                         \
++                const ggml_cuda_mmq_config c = ggml_cuda_mmq_get_config_##TABLE((ggml_type) t, J, fb);                  \
++                if (c.type == GGML_TYPE_COUNT) {                                                                        \
++                    continue;                                                                                           \
++                }                                                                                                       \
++                const size_t y_load = size_t((c.J*MMQ_TILE_Y_K + c.nthreads - 1) / c.nthreads) * c.nthreads * sizeof(int); \
++                if (ggml_cuda_mmq_get_nbytes_y_tile(c) < y_load) {                                                      \
++                    return false;                                                                                       \
++                }                                                                                                       \
++            }                                                                                                           \
++        }                                                                                                               \
++        return true;                                                                                                    \
++    }                                                                                                                   \
++    template <int t> struct ggml_cuda_mmq_y_tile_guard_##TABLE {                                                        \
++        static_assert(ggml_cuda_mmq_y_tile_ok_##TABLE(t),                                                               \
++            "MMQ: the y tile helper is smaller than the y load for a config in the " #TABLE " table");                 \
++    };                                                                                                                  \
++    template <int... ts> static constexpr size_t ggml_cuda_mmq_y_tile_all_##TABLE(std::integer_sequence<int, ts...>) {  \
++        return (sizeof(ggml_cuda_mmq_y_tile_guard_##TABLE<ts>) + ... + 0);                                              \
++    }                                                                                                                   \
++    static_assert(ggml_cuda_mmq_y_tile_all_##TABLE(std::make_integer_sequence<int, GGML_TYPE_COUNT>()) > 0, "");
++
++GGML_CUDA_MMQ_Y_TILE_GUARD(pascal_older)
++GGML_CUDA_MMQ_Y_TILE_GUARD(pascal_dp4a)
++GGML_CUDA_MMQ_Y_TILE_GUARD(ampere)
++GGML_CUDA_MMQ_Y_TILE_GUARD(blackwell)
++GGML_CUDA_MMQ_Y_TILE_GUARD(gcn)
++GGML_CUDA_MMQ_Y_TILE_GUARD(cdna)
++GGML_CUDA_MMQ_Y_TILE_GUARD(rdna2)
++GGML_CUDA_MMQ_Y_TILE_GUARD(rdna3)
++GGML_CUDA_MMQ_Y_TILE_GUARD(rdna3_5)
++GGML_CUDA_MMQ_Y_TILE_GUARD(rdna4)
++
++#undef GGML_CUDA_MMQ_Y_TILE_GUARD
++#undef GGML_CUDA_MMQ_Y_TILE_GUARD_ON
++#endif // GGML_CUDA_MMQ_Y_TILE_GUARD_ON
+diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh
+index 198aab4b..9f18c068 100644
+--- a/ggml/src/ggml-cuda/mmq.cuh
++++ b/ggml/src/ggml-cuda/mmq.cuh
+@@ -1388,11 +1388,17 @@ struct mmq_args {
+     int J_best; // Tile width in ne11(dense)/ne12(MoE) direction to use for optimal performance.
+ };
+ 
++// The bytes of src1 that one tile's y load copies, starting at the tile's first column. The load in
++// mul_mat_q_process_tile runs to the next multiple of nthreads ints with no bound on the column count, so it fills
++// the whole shared-memory y tile. One expression sizes that tile here and the global src1 padding in mmq.cu.
++static constexpr __host__ __device__ size_t ggml_cuda_mmq_get_nbytes_y_tile(const ggml_cuda_mmq_config & config) {
++    return GGML_PAD(config.J*sizeof(block_q8_1_mmq), config.nthreads*sizeof(int));
++}
++
+ static size_t mmq_get_nbytes_shared(const ggml_cuda_mmq_config & config, const int cc) {
+     const size_t nbs_ids = config.J*sizeof(int);
+     const size_t nbs_x = ggml_cuda_mmq_get_nbytes_shared_x(config, cc);
+-    const size_t nbs_y = config.J * (sizeof(block_q8_1_mmq));
+-    return nbs_ids + nbs_x + GGML_PAD(nbs_y, config.nthreads*sizeof(int));
++    return nbs_ids + nbs_x + ggml_cuda_mmq_get_nbytes_y_tile(config);
+ }
+ 
+ template <ggml_type type, int J, bool fallback, ggml_prec prec_src1 = GGML_PREC_Q8>
diff --git a/docs/maxusai/tasks/mmq-rules-check.cu b/docs/maxusai/tasks/mmq-rules-check.cu
index 7630b7cb44b..41ee0f7f061 100644
--- a/docs/maxusai/tasks/mmq-rules-check.cu
+++ b/docs/maxusai/tasks/mmq-rules-check.cu
@@ -20,22 +20,27 @@
 // memory, the first with the fewest tiles for ncols_opt. That loop is replicated below because it sits in a
 // template that queries the device; everything else is llama.cpp's own code.
 //
-// The rules, as the argument of ggml_cuda_mmq_get_J_max() unless stated:
+// The rules, as the argument of get_J_max() unless stated:
 //   #24127    : ne11                   (1 for broadcast gate/up, n_expert_used otherwise), b9992 .. dd266785c^
 //   #29941    : ne12                   master, merged as dd266785c
-//   #29953    : the launched J itself  (padding moved next to the tile-size choice, so the two cannot disagree)
-//   #27044    : ne12*n_expert_used     our compat 903 until now
+//   #29953    : the launched J itself  (padding moved next to the tile-size choice, so the two cannot disagree);
+//               this is #29953 AS PUBLISHED (5bd8b0013), not its head
+//   #27044    : ne12*n_expert_used     the fork's compat 903 before #448
 //   #448      : the widest tile that has a config, get_J_max(type, fallback, cc, 512)
 //   #448+pad  : the widest padded tile, max over J of GGML_PAD(J*block, nthreads_J*4)/block
+//   #29953 head (3070d927f): src1 by the launched config's padded tile, GGML_PAD(J*block, nthreads*4) BYTES, and
+//               ids_dst by J - 1 entries. Sufficient by construction -- it is the load loop's own extent for the
+//               config that is launched -- so the sweep checks that this model agrees, not that the head is right.
 //
 // Build from the root of a llama.cpp checkout, with the CUDA toolkit's nvcc. It never touches a GPU:
 //   nvcc -Wno-deprecated-gpu-targets -std=c++17 -I. -Iggml/include -Iggml/src -Iggml/src/ggml-cuda \
 //        mmq-rules-check.cu -o mmq-rules-check -lcublas -lcuda
 // It calls ggml_cuda_mmq_get_config() in its four-argument form, which exists at the fork's pin b11081 and is
-// master's signature with prec_src1 defaulted, so the same source builds against either.
+// master's signature with prec_src1 defaulted, and it carries its own copy of ggml_cuda_mmq_get_J_max(), which
+// llama.cpp#29953 deletes. So the same source builds against b11081, dd266785c and #29953's head.
 //   ./mmq-rules-check            the whole sweep
 //   ./mmq-rules-check --design   shapes for test_mul_mat_id where only ids_dst is short
-// The exit code is non-zero if the last rule leaves a shape short.
+// The exit code is non-zero if #448 amended (the fork's 903) or #29953's head leaves a shape short.
 #include "ggml/src/ggml-cuda/mmq.cuh"
 #include "ggml/src/ggml-cuda/mmvq.cuh"   // MMVQ_MAX_BATCH_SIZE
 #include <climits>
@@ -49,6 +54,19 @@ extern "C" void ggml_abort(const char *, int, const char *, ...) { abort(); }
 
 static const size_t BLOCK = sizeof(block_q8_1_mmq);
 
+// ggml_cuda_mmq_get_J_max() as it is at dd266785c (mmq.cuh), verbatim but for the name. The older rules are
+// defined by it, and #29953 deletes it, so the check keeps a copy rather than depend on the checkout having one.
+static int get_J_max(const ggml_type type, const bool fallback, const int cc, const int64_t ne11) {
+    int ret = std::min(ne11, int64_t(512));
+    ret -= ret % 8;
+    for (;ret > 0; ret -= 8) {
+        if (ggml_cuda_mmq_get_config(type, ret, fallback, cc).type != GGML_TYPE_COUNT) {
+            return ret;
+        }
+    }
+    return ret;
+}
+
 // The largest MUL_MAT_ID batch MMVQ takes (mmvq.cu, get_mmvq_mmid_max_batch*; defined outside its header, so
 // replicated for the ten types checked). A larger batch goes to MMQ.
 //
@@ -93,7 +111,7 @@ static void launched(ggml_type t, bool fb, int cc, size_t smpbo, int64_t ncols_o
 
 // #448 as written: the widest tile that has a config, in blocks. And amended: the widest padded tile.
 static void widest(ggml_type t, bool fb, int cc, int * J_out, int * pad_out) {
-    *J_out = ggml_cuda_mmq_get_J_max(t, fb, cc, 512);
+    *J_out = get_J_max(t, fb, cc, 512);
     size_t widest_bytes = 0;
     for (int J = 8; J <= 128; J += 8) {
         const ggml_cuda_mmq_config c = ggml_cuda_mmq_get_config(t, J, fb, cc);
@@ -135,10 +153,17 @@ struct tally {
     }
 };
 
-enum { R24127, R29941, R29953, R27044, R448, R448P, NR };
+enum { R24127, R29941, R29953, R27044, R448, R448P, R29953H, NR };
 static const char * const RN[NR] = {
-    "#24127 (ne11)", "#29941 (ne12), master", "#29953 (the launched tile)",
-    "#27044 (ne12*n_expert_used)", "#448 (widest tile)", "#448 amended (widest padded tile)"};
+    "#24127 (ne11)", "#29941 (ne12), master", "#29953 as published (launched tile)",
+    "#27044 (ne12*n_expert_used)", "#448 (widest tile)", "#448 amended (widest padded tile)",
+    "#29953 head (launched padded tile)"};
+
+// The head pads src1 by GGML_PAD(J*block, nthreads*4) BYTES. Scored as floor(bytes/block) blocks, which
+// under-counts it by less than a block and is still >= the requirement (floor(T/B) >= ceil(T/B) - 1).
+static int head_pad(int J, int nthreads) {
+    return J == 0 ? 0 : (int) (GGML_PAD(J*BLOCK, nthreads*sizeof(int)) / BLOCK);
+}
 
 struct arch { const char * name; int cc; size_t smpbo; };
 // smpbo is the opt-in shared memory per block; on HIP ggml uses sharedMemPerBlock, 64 KiB on these parts.
@@ -183,10 +208,10 @@ static void design(int cc, size_t smpbo, const char * arch_name) {
             if (J == 0) continue;
             const int need = src1_need(J, nth);
             const int p[5] = {
-                ggml_cuda_mmq_get_J_max(t, fb, cc, n_used),      // b = false, so ne11 = n_used
-                ggml_cuda_mmq_get_J_max(t, fb, cc, ne12),
+                get_J_max(t, fb, cc, n_used),      // b = false, so ne11 = n_used
+                get_J_max(t, fb, cc, ne12),
                 J,
-                ggml_cuda_mmq_get_J_max(t, fb, cc, ne12*n_used),
+                get_J_max(t, fb, cc, ne12*n_used),
                 J_wide,
             };
             bool all_ok = true;
@@ -236,18 +261,19 @@ int main(int argc, char ** argv) {
                 if (bcast && n_used == 1) continue;   // one expert per token has nothing to broadcast
                 ++ids_shapes;
                 const int pad[NR] = {
-                    ggml_cuda_mmq_get_J_max(t, fb, a.cc, bcast ? 1 : n_used),
-                    ggml_cuda_mmq_get_J_max(t, fb, a.cc, ne12),
+                    get_J_max(t, fb, a.cc, bcast ? 1 : n_used),
+                    get_J_max(t, fb, a.cc, ne12),
                     J,
-                    ggml_cuda_mmq_get_J_max(t, fb, a.cc, ne12*n_used),
+                    get_J_max(t, fb, a.cc, ne12*n_used),
                     J_wide,
                     pad_wide_amended,
+                    head_pad(J, nth),
                 };
                 const int u = n_used == 1 ? 0 : 1;
                 for (int r = 0; r < NR; ++r) {
                     ids_src1[r][u].add(need_src1 - pad[r], a.name, TN[ti], fb, ne12, n_used, ne02, J, nth, pad[r]);
-                    // only #448 pads ids_dst, by the same count as src1
-                    const int dst_pad = (r == R448 || r == R448P) ? pad[r] : 0;
+                    // #448 pads ids_dst by the same count as src1, #29953's head by J - 1; nothing else pads it
+                    const int dst_pad = (r == R448 || r == R448P) ? pad[r] : r == R29953H ? J - 1 : 0;
                     ids_dst[r][u].add(need_ids - dst_pad, a.name, TN[ti], fb, ne12, n_used, ne02, J, nth, dst_pad);
                 }
             }
@@ -263,8 +289,8 @@ int main(int argc, char ** argv) {
             const int64_t last = ne11 - ((ne11 + J - 1) / J - 1) * J;
             const int need = src1_need(J, nth) - (int) (last - 1);
             const int pad[NR] = {
-                ggml_cuda_mmq_get_J_max(t, fb, a.cc, ne11), ggml_cuda_mmq_get_J_max(t, fb, a.cc, ne11), J,
-                ggml_cuda_mmq_get_J_max(t, fb, a.cc, ne11), J_wide, pad_wide_amended,
+                get_J_max(t, fb, a.cc, ne11), get_J_max(t, fb, a.cc, ne11), J,
+                get_J_max(t, fb, a.cc, ne11), J_wide, pad_wide_amended, head_pad(J, nth),
             };
             for (int r = 0; r < NR; ++r) dense[r].add(need - pad[r], a.name, TN[ti], fb, ne11, 0, 0, J, nth, pad[r]);
         }
@@ -289,10 +315,14 @@ int main(int argc, char ** argv) {
     printf("dense branch: %ld shapes (same archs, types and fallbacks x every MMQ batch up to 1024 columns)\n",
            dense_shapes);
     for (int r = 0; r < NR; ++r) {
-        if (r == R24127 || r == R29953 || r == R448 || r == R448P) dense[r].print(RN[r], "blocks");
+        if (r == R24127 || r == R29953 || r == R448 || r == R448P || r == R29953H) dense[r].print(RN[r], "blocks");
     }
 
-    const bool ok = ids_src1[R448P][0].n == 0 && ids_src1[R448P][1].n == 0
-                 && ids_dst[R448P][0].n == 0 && ids_dst[R448P][1].n == 0 && dense[R448P].n == 0;
+    // non-zero if either fix this work stands behind -- the fork's 903 or upstream's head -- leaves a shape short
+    bool ok = true;
+    for (int r : {R448P, R29953H}) {
+        ok = ok && ids_src1[r][0].n == 0 && ids_src1[r][1].n == 0 && ids_dst[r][0].n == 0 && ids_dst[r][1].n == 0
+                && dense[r].n == 0;
+    }
     return ok ? 0 : 1;
 }
diff --git a/docs/maxusai/tasks/mmq-rules-gpu.sh b/docs/maxusai/tasks/mmq-rules-gpu.sh
index e497db3d9d8..a63bfc8aba7 100755
--- a/docs/maxusai/tasks/mmq-rules-gpu.sh
+++ b/docs/maxusai/tasks/mmq-rules-gpu.sh
@@ -34,6 +34,7 @@ MODES=${MODES:-stock guard guard_ids_dst guard_src1 memcheck exact}
 # One build directory per CUDA_ARCH, or the two architectures clobber each other's objects.
 BUILD_DIR=${BUILD_DIR:-build-mmq-successor}
 VARIANTS=${VARIANTS:-ne11 ne12 p27044 successor p29953}
+BUILD_KINDS=${BUILD_KINDS:-stock debug}   # a pass that only runs guarded modes needs only the debug build
 # CMake hands nvcc the toolkit's include directory, but a search path from a shell profile (CPATH and friends) is
 # searched first, and one naming another toolkit mixes headers: CUDA 12.1's crt/host_config.h refuses gcc 13.
 unset CPATH C_INCLUDE_PATH CPLUS_INCLUDE_PATH
@@ -59,7 +60,7 @@ cmake -S . -B "$BUILD_DIR" -G Ninja -DCMAKE_BUILD_TYPE=Release -DGGML_CUDA=ON -D
 
 # p29953 is the only variant that edits mmq.cuh, so every MMQ instance rebuilds for it. Build it last and the
 # other eight builds only recompile mmq.cu.
-for v in $VARIANTS; do for e in stock debug; do
+for v in $VARIANTS; do for e in $BUILD_KINDS; do
     [ -x "$OUT/bin/test-backend-ops-$v-$e" ] && continue
     git checkout --quiet -- "$F" "$H"
     python3 "$HERE/mmq-variant.py" . "$v" $([ $e = debug ] && echo --debug) > /dev/null || exit 1
diff --git a/docs/maxusai/tasks/mmq-rules-head29953-controls.sh b/docs/maxusai/tasks/mmq-rules-head29953-controls.sh
new file mode 100755
index 00000000000..7a635d92fd5
--- /dev/null
+++ b/docs/maxusai/tasks/mmq-rules-head29953-controls.sh
@@ -0,0 +1,19 @@
+#!/usr/bin/env bash
+# Positive controls for mmq-rules-head29953.sh's sm_120 pass B, in the same harness and output directory.
+# Pass A controlled guard:src1 only. Pass B also ran the head under guard (both buffers), guard:ids_dst and exact
+# (memcheck with exact-size allocations); a clean head there means something only if the same harness shows the
+# published #29953 failing in those modes. ids16 is deterministic for both buffers; one113 is the case where the
+# published rule gave the most memcheck errors in matrix.md (11,939).
+#
+# Runs from its own directory, after mmq-rules-head29953.sh (it reuses that pass's p29953-debug binary).
+set -u
+cd "$(dirname "$0")" || exit 1
+HERE=$PWD
+SRC=$HERE/llama.cpp-e2e
+trap 'pkill -P $$ 2>/dev/null; git -C "$SRC" checkout --quiet -- ggml/src/ggml-cuda/mmq.cu ggml/src/ggml-cuda/mmq.cuh 2>/dev/null' EXIT
+echo "=== $(date -u +%FT%TZ) E: sm_120 controls, published #29953"
+BUILD_KINDS=debug VARIANTS=p29953 CASES="ids16 one113" MODES="guard guard_ids_dst exact" REPS=3 \
+    CUDA_ARCH=120 GPU=0 BUILD_DIR=build-mmq-successor ./mmq-rules-gpu.sh "$SRC" "$HERE/mmq-head-results-sm120" \
+    || { echo "E FAILED"; exit 1; }
+nvidia-smi --query-compute-apps=pid,process_name --format=csv,noheader | grep -E 'test-backend-ops|compute-sanitizer' || echo "no leftover GPU processes"
+echo "=== $(date -u +%FT%TZ) CONTROLS COMPLETE"
diff --git a/docs/maxusai/tasks/mmq-rules-head29953.sh b/docs/maxusai/tasks/mmq-rules-head29953.sh
new file mode 100755
index 00000000000..c726e41c90b
--- /dev/null
+++ b/docs/maxusai/tasks/mmq-rules-head29953.sh
@@ -0,0 +1,42 @@
+#!/usr/bin/env bash
+# llama.cpp#29953 at its head (3070d927f), the exact tree, through the same harness as matrix.md and matrix-sm75.md.
+# The published #29953 (p29953.patch, 5bd8b0013) runs alongside as the positive control: a pass in which the
+# control does not fault on ids16 proves nothing about the head.
+#
+#   sm_120, GPU0  A: ids16, j100_b1, j100_b0 under guard:src1, 10 reps each, control and head
+#                 B: every case, head only: stock, guard (both buffers), guard:src1, guard:ids_dst x3; exact x1
+#   sm_75,  GPU1  C: ids16 and dense321, guard:src1 and guard:ids_dst x3, control
+#                 D: ids16 and dense321, stock, guard:src1 and guard:ids_dst x3, head
+#
+# Runs from its own directory, as the other mmq-rules-* scripts here: llama.cpp-e2e next to it is a dd266785c
+# checkout with mmq-tests.patch applied, and mmq-variant.py's head29953 swaps in the head's mmq.cu/mmq.cuh.
+# Sequential on purpose: every build rewrites mmq.cu/mmq.cuh in the one checkout, so two passes may not overlap.
+set -u
+cd "$(dirname "$0")" || exit 1
+HERE=$PWD
+SRC=$HERE/llama.cpp-e2e
+trap 'pkill -P $$ 2>/dev/null; git -C "$SRC" checkout --quiet -- ggml/src/ggml-cuda/mmq.cu ggml/src/ggml-cuda/mmq.cuh 2>/dev/null' EXIT
+
+stamp() { echo "=== $(date -u +%FT%TZ) $*"; }
+
+O120=$HERE/mmq-head-results-sm120
+stamp "A: sm_120 control + head, guard:src1 x10"
+BUILD_KINDS=debug VARIANTS="p29953 head29953" CASES="ids16 j100_b1 j100_b0" MODES=guard_src1 REPS=10 \
+    CUDA_ARCH=120 GPU=0 BUILD_DIR=build-mmq-successor ./mmq-rules-gpu.sh "$SRC" "$O120" || { stamp "A FAILED"; exit 1; }
+
+stamp "B: sm_120 head, every case"
+BUILD_KINDS="stock debug" VARIANTS=head29953 MODES="stock guard guard_src1 guard_ids_dst exact" REPS=3 \
+    CUDA_ARCH=120 GPU=0 BUILD_DIR=build-mmq-successor ./mmq-rules-gpu.sh "$SRC" "$O120" || { stamp "B FAILED"; exit 1; }
+
+O75=$HERE/mmq-head-results-sm75
+stamp "C: sm_75 control"
+BUILD_KINDS=debug VARIANTS=p29953 CASES="ids16 dense321" MODES="guard_src1 guard_ids_dst" REPS=3 \
+    CUDA_ARCH=75 GPU=1 MIN_FREE_MIB=4096 BUILD_DIR=build-mmq-sm75 ./mmq-rules-gpu.sh "$SRC" "$O75" || { stamp "C FAILED"; exit 1; }
+
+stamp "D: sm_75 head"
+BUILD_KINDS="stock debug" VARIANTS=head29953 CASES="ids16 dense321" MODES="stock guard_src1 guard_ids_dst" REPS=3 \
+    CUDA_ARCH=75 GPU=1 MIN_FREE_MIB=4096 BUILD_DIR=build-mmq-sm75 ./mmq-rules-gpu.sh "$SRC" "$O75" || { stamp "D FAILED"; exit 1; }
+
+stamp "leftover GPU processes from this run (expect none):"
+nvidia-smi --query-compute-apps=pid,process_name --format=csv,noheader | grep -E 'test-backend-ops|compute-sanitizer' || echo "none"
+stamp "HEAD29953 COMPLETE"
diff --git a/docs/maxusai/tasks/mmq-rules-table.py b/docs/maxusai/tasks/mmq-rules-table.py
index ff20f0300c1..50aca06f2e9 100755
--- a/docs/maxusai/tasks/mmq-rules-table.py
+++ b/docs/maxusai/tasks/mmq-rules-table.py
@@ -1,10 +1,13 @@
 #!/usr/bin/env python3
 """Render mmq-rules-gpu.sh's results.tsv as one markdown table per mode (ADR 0012: tables come from the generator).
 
-    mmq-rules-table.py results.tsv [results.tsv ...]
+    mmq-rules-table.py [--summary] results.tsv [results.tsv ...]
 
 A cell reads "p/n" for the runs that passed of the runs made. A mode run under compute-sanitizer also carries its
 error count, "E=<n>"; "abort" means the process died instead of finishing the case.
+
+--summary prints one row per (rule, mode) instead: cases, runs, runs passed, runs that aborted, and the memcheck
+errors summed over the runs made under the sanitizer.
 """
 import collections
 import csv
@@ -21,20 +24,24 @@
 VARIANT_LABEL = {
     "ne11":      "ne11 (#24127)",
     "ne12":      "ne12 (#29941, master)",
-    "p29953":    "#29953",
+    "p29953":    "#29953 as published",
     "p27044":    "ne12*n_expert_used (#27044)",
     "successor": "widest tile (#448)",
     "s448p":     "widest padded tile (#448 amended)",
     "p29953fix": "#29953 + amendment",
+    "head29953": "#29953 head (`3070d927f`)",
 }
-VARIANT_ORDER = ["ne11", "ne12", "p29953", "p27044", "successor", "s448p", "p29953fix"]
+VARIANT_ORDER = ["ne11", "ne12", "p29953", "p27044", "successor", "s448p", "p29953fix", "head29953"]
 
 
 def main():
-    if len(sys.argv) < 2:
+    args = sys.argv[1:]
+    summary = "--summary" in args
+    paths = [a for a in args if a != "--summary"]
+    if not paths:
         sys.exit(__doc__)
     rows = []
-    for path in sys.argv[1:]:
+    for path in paths:
         with open(path) as f:
             rows += [r for r in csv.DictReader(f, delimiter="\t") if r.get("case")]
 
@@ -62,6 +69,21 @@ def fmt(rs):
             out += " abort"
         return out
 
+    if summary:
+        print("| rule | mode | cases | runs | passed | aborted | memcheck errors |")
+        print("|---|---|---|---|---|---|---|")
+        for v in variants:
+            for m in modes:
+                rs = [r for r in rows if r["variant"] == v and r["mode"] == m]
+                if not rs:
+                    continue
+                aborted = sum(r["illegal_access"] not in ("", "0") for r in rs)
+                errs = [int(r["memcheck_errors"]) for r in rs if r["memcheck_errors"]]
+                print("| %s | %s | %d | %d | %d | %d | %s |" % (
+                    VARIANT_LABEL[v], m, len({r["case"] for r in rs}), len(rs),
+                    sum(r["passed"] == "1" for r in rs), aborted, sum(errs) if errs else "-"))
+        return
+
     for m in modes:
         print("### %s\n" % MODE_LABEL[m])
         print("| case | " + " | ".join(VARIANT_LABEL[v] for v in variants) + " |")
diff --git a/docs/maxusai/tasks/mmq-successor-results/check-rules-wide.txt b/docs/maxusai/tasks/mmq-successor-results/check-rules-wide.txt
index 881dc1fca27..37d24215865 100644
--- a/docs/maxusai/tasks/mmq-successor-results/check-rules-wide.txt
+++ b/docs/maxusai/tasks/mmq-successor-results/check-rules-wide.txt
@@ -8,7 +8,7 @@ ids branch: 9431816 shapes (sm_70/75/80/86/89/90/120, gfx1151, CDNA3 and RDNA4 x
   #29941 (ne12), master              SHORT in 40306 shapes, worst 63 blocks past the allocation
                                       worst case: sm_75, q4_0, fallback=1, 65/1/256: J=128, nthreads=256, padding=64
                                       by arch: sm_75 2258 sm_86 2368 sm_89 2360 sm_120 2360 gfx1151 2064 sm_70 2360 sm_80 2368 sm_90 2360 CDNA3 19840 RDNA4 1968
-  #29953 (the launched tile)         SHORT in 168520 shapes, worst 12 blocks past the allocation
+  #29953 as published (launched tile) SHORT in 168520 shapes, worst 12 blocks past the allocation
                                       worst case: CDNA3, q4_0, fallback=0, 9/1/256: J=16, nthreads=512, padding=16
                                       by arch: sm_75 4984 sm_86 4488 sm_89 4480 sm_120 4480 gfx1151 58208 sm_70 4480 sm_80 4488 sm_90 4480 CDNA3 20320 RDNA4 58112
   #27044 (ne12*n_expert_used)        SHORT in 40306 shapes, worst 63 blocks past the allocation
@@ -18,6 +18,7 @@ ids branch: 9431816 shapes (sm_70/75/80/86/89/90/120, gfx1151, CDNA3 and RDNA4 x
                                       worst case: CDNA3, q4_0, fallback=0, 49/1/256: J=64, nthreads=512, padding=64
                                       by arch: gfx1151 128 CDNA3 19200 RDNA4 128
   #448 amended (widest padded tile)  covered
+  #29953 head (launched padded tile) covered
  src1_q8_1, two or more experts per token:
   #24127 (ne11)                      SHORT in 9106552 shapes, worst 127 blocks past the allocation
                                       worst case: sm_75, q4_0, fallback=0, 113/2/256: J=128, nthreads=256, padding=0
@@ -25,7 +26,7 @@ ids branch: 9431816 shapes (sm_70/75/80/86/89/90/120, gfx1151, CDNA3 and RDNA4 x
   #29941 (ne12), master              SHORT in 1255564 shapes, worst 63 blocks past the allocation
                                       worst case: sm_75, q4_0, fallback=1, 65/2/256: J=128, nthreads=256, padding=64
                                       by arch: sm_75 67740 sm_86 71040 sm_89 70800 sm_120 70800 gfx1151 84920 sm_70 70800 sm_80 71040 sm_90 70800 CDNA3 595200 RDNA4 82424
-  #29953 (the launched tile)         SHORT in 3520756 shapes, worst 12 blocks past the allocation
+  #29953 as published (launched tile) SHORT in 3520756 shapes, worst 12 blocks past the allocation
                                       worst case: CDNA3, q4_0, fallback=0, 9/2/256: J=16, nthreads=512, padding=16
                                       by arch: sm_75 149520 sm_86 134640 sm_89 134400 sm_120 134400 gfx1151 978626 sm_70 134400 sm_80 134640 sm_90 134400 CDNA3 609600 RDNA4 976130
   #27044 (ne12*n_expert_used)        SHORT in 635468 shapes, worst 17 blocks past the allocation
@@ -35,6 +36,7 @@ ids branch: 9431816 shapes (sm_70/75/80/86/89/90/120, gfx1151, CDNA3 and RDNA4 x
                                       worst case: CDNA3, q4_0, fallback=0, 49/2/256: J=64, nthreads=512, padding=64
                                       by arch: gfx1151 25256 CDNA3 576000 RDNA4 25256
   #448 amended (widest padded tile)  covered
+  #29953 head (launched padded tile) covered
  ids_dst, any n_expert_used:
   #24127 (ne11)                      SHORT in 9431816 shapes, worst 127 int32 entries past the allocation
                                       worst case: sm_75, q4_0, fallback=0, 113/1/256: J=128, nthreads=256, padding=0
@@ -42,7 +44,7 @@ ids branch: 9431816 shapes (sm_70/75/80/86/89/90/120, gfx1151, CDNA3 and RDNA4 x
   #29941 (ne12), master              SHORT in 9431816 shapes, worst 127 int32 entries past the allocation
                                       worst case: sm_75, q4_0, fallback=0, 113/1/256: J=128, nthreads=256, padding=0
                                       by arch: sm_75 630168 sm_86 630168 sm_89 629920 sm_120 629920 gfx1151 2197152 sm_70 629920 sm_80 630168 sm_90 629920 CDNA3 629920 RDNA4 2194560
-  #29953 (the launched tile)         SHORT in 9431816 shapes, worst 127 int32 entries past the allocation
+  #29953 as published (launched tile) SHORT in 9431816 shapes, worst 127 int32 entries past the allocation
                                       worst case: sm_75, q4_0, fallback=0, 113/1/256: J=128, nthreads=256, padding=0
                                       by arch: sm_75 630168 sm_86 630168 sm_89 629920 sm_120 629920 gfx1151 2197152 sm_70 629920 sm_80 630168 sm_90 629920 CDNA3 629920 RDNA4 2194560
   #27044 (ne12*n_expert_used)        SHORT in 9431816 shapes, worst 127 int32 entries past the allocation
@@ -50,14 +52,16 @@ ids branch: 9431816 shapes (sm_70/75/80/86/89/90/120, gfx1151, CDNA3 and RDNA4 x
                                       by arch: sm_75 630168 sm_86 630168 sm_89 629920 sm_120 629920 gfx1151 2197152 sm_70 629920 sm_80 630168 sm_90 629920 CDNA3 629920 RDNA4 2194560
   #448 (widest tile)                 covered
   #448 amended (widest padded tile)  covered
+  #29953 head (launched padded tile) covered
 dense branch: 203200 shapes (same archs, types and fallbacks x every MMQ batch up to 1024 columns)
   #24127 (ne11)                      SHORT in 3650 shapes, worst 20 blocks past the allocation
                                       worst case: CDNA3, q4_0, fallback=0, 9/0/0: J=16, nthreads=512, padding=0
                                       by arch: sm_75 130 sm_86 130 sm_89 130 sm_120 130 gfx1151 185 sm_70 130 sm_80 130 sm_90 130 CDNA3 2370 RDNA4 185
-  #29953 (the launched tile)         SHORT in 2175 shapes, worst 7 blocks past the allocation
+  #29953 as published (launched tile) SHORT in 2175 shapes, worst 7 blocks past the allocation
                                       worst case: CDNA3, q4_0, fallback=0, 193/0/0: J=64, nthreads=512, padding=64
                                       by arch: sm_75 45 gfx1151 45 CDNA3 2040 RDNA4 45
   #448 (widest tile)                 SHORT in 2050 shapes, worst 7 blocks past the allocation
                                       worst case: CDNA3, q4_0, fallback=0, 193/0/0: J=64, nthreads=512, padding=64
                                       by arch: gfx1151 45 CDNA3 1960 RDNA4 45
   #448 amended (widest padded tile)  covered
+  #29953 head (launched padded tile) covered
diff --git a/docs/maxusai/tasks/mmq-successor-results/guard-timing.txt b/docs/maxusai/tasks/mmq-successor-results/guard-timing.txt
new file mode 100644
index 00000000000..b821199acd9
--- /dev/null
+++ b/docs/maxusai/tasks/mmq-successor-results/guard-timing.txt
@@ -0,0 +1,8 @@
+host load at start (1/5/15 min): 6.96 11.67 17.05, 32 cores
+rep 1: guard on 11.09s, off 8.15s
+rep 2: guard on 10.91s, off 8.27s
+rep 3: guard on 10.90s, off 7.86s
+rep 4: guard on 11.26s, off 8.59s
+rep 5: guard on 11.90s, off 7.95s
+median: guard on 11.09s, off 8.15s, over 5 interleaved reps
+host load at end (1/5/15 min): 8.10 10.59 16.12
diff --git a/docs/maxusai/tasks/mmq-successor-results/matrix-head29953-sm120.md b/docs/maxusai/tasks/mmq-successor-results/matrix-head29953-sm120.md
new file mode 100644
index 00000000000..2e0973fb4f8
--- /dev/null
+++ b/docs/maxusai/tasks/mmq-successor-results/matrix-head29953-sm120.md
@@ -0,0 +1,90 @@
+### stock, memory pool, no sanitizer
+
+| case | #29953 as published | #29953 head (`3070d927f`) |
+|---|---|---|
+| `ids16` | - | 3/3 |
+| `j100_b1` | - | 3/3 |
+| `j100_b0` | - | 3/3 |
+| `p29847_b0` | - | 3/3 |
+| `p29847_b1` | - | 3/3 |
+| `e120_b0` | - | 3/3 |
+| `e120_b1` | - | 3/3 |
+| `one113` | - | 3/3 |
+| `orig2040` | - | 3/3 |
+| `fb65` | - | 3/3 |
+| `t100` | - | 3/3 |
+| `ids64` | - | 3/3 |
+| `dense321` | - | 3/3 |
+
+### exact-size cudaMalloc under memcheck
+
+| case | #29953 as published | #29953 head (`3070d927f`) |
+|---|---|---|
+| `ids16` | 0/1 E=298 | 1/1 E=0 |
+| `j100_b1` | - | 1/1 E=0 |
+| `j100_b0` | - | 1/1 E=0 |
+| `p29847_b0` | - | 1/1 E=0 |
+| `p29847_b1` | - | 1/1 E=0 |
+| `e120_b0` | - | 1/1 E=0 |
+| `e120_b1` | - | 1/1 E=0 |
+| `one113` | 0/1 E=12885 | 1/1 E=0 |
+| `orig2040` | - | 1/1 E=0 |
+| `fb65` | - | 1/1 E=0 |
+| `t100` | - | 1/1 E=0 |
+| `ids64` | - | 1/1 E=0 |
+| `dense321` | - | 1/1 E=0 |
+
+### guard page on src1 and ids_dst, no sanitizer
+
+| case | #29953 as published | #29953 head (`3070d927f`) |
+|---|---|---|
+| `ids16` | 0/3 abort | 3/3 |
+| `j100_b1` | - | 3/3 |
+| `j100_b0` | - | 3/3 |
+| `p29847_b0` | - | 3/3 |
+| `p29847_b1` | - | 3/3 |
+| `e120_b0` | - | 3/3 |
+| `e120_b1` | - | 3/3 |
+| `one113` | 0/3 abort | 3/3 |
+| `orig2040` | - | 3/3 |
+| `fb65` | - | 3/3 |
+| `t100` | - | 3/3 |
+| `ids64` | - | 3/3 |
+| `dense321` | - | 3/3 |
+
+### guard page on ids_dst only, no sanitizer
+
+| case | #29953 as published | #29953 head (`3070d927f`) |
+|---|---|---|
+| `ids16` | 0/3 abort | 3/3 |
+| `j100_b1` | - | 3/3 |
+| `j100_b0` | - | 3/3 |
+| `p29847_b0` | - | 3/3 |
+| `p29847_b1` | - | 3/3 |
+| `e120_b0` | - | 3/3 |
+| `e120_b1` | - | 3/3 |
+| `one113` | 0/3 abort | 3/3 |
+| `orig2040` | - | 3/3 |
+| `fb65` | - | 3/3 |
+| `t100` | - | 3/3 |
+| `ids64` | - | 3/3 |
+| `dense321` | - | 3/3 |
+
+### guard page on src1 only, no sanitizer
+
+| case | #29953 as published | #29953 head (`3070d927f`) |
+|---|---|---|
+| `ids16` | 0/10 abort | 10/10 |
+| `j100_b1` | 7/10 abort | 10/10 |
+| `j100_b0` | 5/10 abort | 10/10 |
+| `p29847_b0` | - | 3/3 |
+| `p29847_b1` | - | 3/3 |
+| `e120_b0` | - | 3/3 |
+| `e120_b1` | - | 3/3 |
+| `one113` | - | 3/3 |
+| `orig2040` | - | 3/3 |
+| `fb65` | - | 3/3 |
+| `t100` | - | 3/3 |
+| `ids64` | - | 3/3 |
+| `dense321` | - | 3/3 |
+
diff --git a/docs/maxusai/tasks/mmq-successor-results/matrix-head29953-sm75.md b/docs/maxusai/tasks/mmq-successor-results/matrix-head29953-sm75.md
new file mode 100644
index 00000000000..bbcfa9fdd2f
--- /dev/null
+++ b/docs/maxusai/tasks/mmq-successor-results/matrix-head29953-sm75.md
@@ -0,0 +1,21 @@
+### stock, memory pool, no sanitizer
+
+| case | #29953 as published | #29953 head (`3070d927f`) |
+|---|---|---|
+| `ids16` | - | 3/3 |
+| `dense321` | - | 3/3 |
+
+### guard page on ids_dst only, no sanitizer
+
+| case | #29953 as published | #29953 head (`3070d927f`) |
+|---|---|---|
+| `ids16` | 0/3 abort | 3/3 |
+| `dense321` | 3/3 | 3/3 |
+
+### guard page on src1 only, no sanitizer
+
+| case | #29953 as published | #29953 head (`3070d927f`) |
+|---|---|---|
+| `ids16` | 0/3 abort | 3/3 |
+| `dense321` | 0/3 abort | 3/3 |
+
diff --git a/docs/maxusai/tasks/mmq-successor-results/results-head29953-sm120.tsv b/docs/maxusai/tasks/mmq-successor-results/results-head29953-sm120.tsv
new file mode 100644
index 00000000000..ad2ae777277
--- /dev/null
+++ b/docs/maxusai/tasks/mmq-successor-results/results-head29953-sm120.tsv
@@ -0,0 +1,235 @@
+case	variant	mode	rep	passed	illegal_access	memcheck_errors
+ids16	p29953	guard_src1	1	0	1	
+ids16	p29953	guard_src1	2	0	1	
+ids16	p29953	guard_src1	3	0	1	
+ids16	p29953	guard_src1	4	0	1	
+ids16	p29953	guard_src1	5	0	1	
+ids16	p29953	guard_src1	6	0	1	
+ids16	p29953	guard_src1	7	0	1	
+ids16	p29953	guard_src1	8	0	1	
+ids16	p29953	guard_src1	9	0	1	
+ids16	p29953	guard_src1	10	0	1	
+ids16	head29953	guard_src1	1	1	0	
+ids16	head29953	guard_src1	2	1	0	
+ids16	head29953	guard_src1	3	1	0	
+ids16	head29953	guard_src1	4	1	0	
+ids16	head29953	guard_src1	5	1	0	
+ids16	head29953	guard_src1	6	1	0	
+ids16	head29953	guard_src1	7	1	0	
+ids16	head29953	guard_src1	8	1	0	
+ids16	head29953	guard_src1	9	1	0	
+ids16	head29953	guard_src1	10	1	0	
+j100_b1	p29953	guard_src1	1	1	0	
+j100_b1	p29953	guard_src1	2	1	0	
+j100_b1	p29953	guard_src1	3	0	1	
+j100_b1	p29953	guard_src1	4	1	0	
+j100_b1	p29953	guard_src1	5	0	1	
+j100_b1	p29953	guard_src1	6	1	0	
+j100_b1	p29953	guard_src1	7	0	1	
+j100_b1	p29953	guard_src1	8	1	0	
+j100_b1	p29953	guard_src1	9	1	0	
+j100_b1	p29953	guard_src1	10	1	0	
+j100_b1	head29953	guard_src1	1	1	0	
+j100_b1	head29953	guard_src1	2	1	0	
+j100_b1	head29953	guard_src1	3	1	0	
+j100_b1	head29953	guard_src1	4	1	0	
+j100_b1	head29953	guard_src1	5	1	0	
+j100_b1	head29953	guard_src1	6	1	0	
+j100_b1	head29953	guard_src1	7	1	0	
+j100_b1	head29953	guard_src1	8	1	0	
+j100_b1	head29953	guard_src1	9	1	0	
+j100_b1	head29953	guard_src1	10	1	0	
+j100_b0	p29953	guard_src1	1	1	0	
+j100_b0	p29953	guard_src1	2	1	0	
+j100_b0	p29953	guard_src1	3	0	1	
+j100_b0	p29953	guard_src1	4	1	0	
+j100_b0	p29953	guard_src1	5	0	1	
+j100_b0	p29953	guard_src1	6	1	0	
+j100_b0	p29953	guard_src1	7	1	0	
+j100_b0	p29953	guard_src1	8	0	1	
+j100_b0	p29953	guard_src1	9	0	1	
+j100_b0	p29953	guard_src1	10	0	1	
+j100_b0	head29953	guard_src1	1	1	0	
+j100_b0	head29953	guard_src1	2	1	0	
+j100_b0	head29953	guard_src1	3	1	0	
+j100_b0	head29953	guard_src1	4	1	0	
+j100_b0	head29953	guard_src1	5	1	0	
+j100_b0	head29953	guard_src1	6	1	0	
+j100_b0	head29953	guard_src1	7	1	0	
+j100_b0	head29953	guard_src1	8	1	0	
+j100_b0	head29953	guard_src1	9	1	0	
+j100_b0	head29953	guard_src1	10	1	0	
+p29847_b0	head29953	stock	1	1	0	
+p29847_b0	head29953	stock	2	1	0	
+p29847_b0	head29953	stock	3	1	0	
+p29847_b0	head29953	guard	1	1	0	
+p29847_b0	head29953	guard	2	1	0	
+p29847_b0	head29953	guard	3	1	0	
+p29847_b0	head29953	guard_src1	1	1	0	
+p29847_b0	head29953	guard_src1	2	1	0	
+p29847_b0	head29953	guard_src1	3	1	0	
+p29847_b0	head29953	guard_ids_dst	1	1	0	
+p29847_b0	head29953	guard_ids_dst	2	1	0	
+p29847_b0	head29953	guard_ids_dst	3	1	0	
+p29847_b0	head29953	exact	1	1	0	0
+p29847_b1	head29953	stock	1	1	0	
+p29847_b1	head29953	stock	2	1	0	
+p29847_b1	head29953	stock	3	1	0	
+p29847_b1	head29953	guard	1	1	0	
+p29847_b1	head29953	guard	2	1	0	
+p29847_b1	head29953	guard	3	1	0	
+p29847_b1	head29953	guard_src1	1	1	0	
+p29847_b1	head29953	guard_src1	2	1	0	
+p29847_b1	head29953	guard_src1	3	1	0	
+p29847_b1	head29953	guard_ids_dst	1	1	0	
+p29847_b1	head29953	guard_ids_dst	2	1	0	
+p29847_b1	head29953	guard_ids_dst	3	1	0	
+p29847_b1	head29953	exact	1	1	0	0
+j100_b0	head29953	stock	1	1	0	
+j100_b0	head29953	stock	2	1	0	
+j100_b0	head29953	stock	3	1	0	
+j100_b0	head29953	guard	1	1	0	
+j100_b0	head29953	guard	2	1	0	
+j100_b0	head29953	guard	3	1	0	
+j100_b0	head29953	guard_ids_dst	1	1	0	
+j100_b0	head29953	guard_ids_dst	2	1	0	
+j100_b0	head29953	guard_ids_dst	3	1	0	
+j100_b0	head29953	exact	1	1	0	0
+j100_b1	head29953	stock	1	1	0	
+j100_b1	head29953	stock	2	1	0	
+j100_b1	head29953	stock	3	1	0	
+j100_b1	head29953	guard	1	1	0	
+j100_b1	head29953	guard	2	1	0	
+j100_b1	head29953	guard	3	1	0	
+j100_b1	head29953	guard_ids_dst	1	1	0	
+j100_b1	head29953	guard_ids_dst	2	1	0	
+j100_b1	head29953	guard_ids_dst	3	1	0	
+j100_b1	head29953	exact	1	1	0	0
+e120_b0	head29953	stock	1	1	0	
+e120_b0	head29953	stock	2	1	0	
+e120_b0	head29953	stock	3	1	0	
+e120_b0	head29953	guard	1	1	0	
+e120_b0	head29953	guard	2	1	0	
+e120_b0	head29953	guard	3	1	0	
+e120_b0	head29953	guard_src1	1	1	0	
+e120_b0	head29953	guard_src1	2	1	0	
+e120_b0	head29953	guard_src1	3	1	0	
+e120_b0	head29953	guard_ids_dst	1	1	0	
+e120_b0	head29953	guard_ids_dst	2	1	0	
+e120_b0	head29953	guard_ids_dst	3	1	0	
+e120_b0	head29953	exact	1	1	0	0
+e120_b1	head29953	stock	1	1	0	
+e120_b1	head29953	stock	2	1	0	
+e120_b1	head29953	stock	3	1	0	
+e120_b1	head29953	guard	1	1	0	
+e120_b1	head29953	guard	2	1	0	
+e120_b1	head29953	guard	3	1	0	
+e120_b1	head29953	guard_src1	1	1	0	
+e120_b1	head29953	guard_src1	2	1	0	
+e120_b1	head29953	guard_src1	3	1	0	
+e120_b1	head29953	guard_ids_dst	1	1	0	
+e120_b1	head29953	guard_ids_dst	2	1	0	
+e120_b1	head29953	guard_ids_dst	3	1	0	
+e120_b1	head29953	exact	1	1	0	0
+one113	head29953	stock	1	1	0	
+one113	head29953	stock	2	1	0	
+one113	head29953	stock	3	1	0	
+one113	head29953	guard	1	1	0	
+one113	head29953	guard	2	1	0	
+one113	head29953	guard	3	1	0	
+one113	head29953	guard_src1	1	1	0	
+one113	head29953	guard_src1	2	1	0	
+one113	head29953	guard_src1	3	1	0	
+one113	head29953	guard_ids_dst	1	1	0	
+one113	head29953	guard_ids_dst	2	1	0	
+one113	head29953	guard_ids_dst	3	1	0	
+one113	head29953	exact	1	1	0	0
+orig2040	head29953	stock	1	1	0	
+orig2040	head29953	stock	2	1	0	
+orig2040	head29953	stock	3	1	0	
+orig2040	head29953	guard	1	1	0	
+orig2040	head29953	guard	2	1	0	
+orig2040	head29953	guard	3	1	0	
+orig2040	head29953	guard_src1	1	1	0	
+orig2040	head29953	guard_src1	2	1	0	
+orig2040	head29953	guard_src1	3	1	0	
+orig2040	head29953	guard_ids_dst	1	1	0	
+orig2040	head29953	guard_ids_dst	2	1	0	
+orig2040	head29953	guard_ids_dst	3	1	0	
+orig2040	head29953	exact	1	1	0	0
+fb65	head29953	stock	1	1	0	
+fb65	head29953	stock	2	1	0	
+fb65	head29953	stock	3	1	0	
+fb65	head29953	guard	1	1	0	
+fb65	head29953	guard	2	1	0	
+fb65	head29953	guard	3	1	0	
+fb65	head29953	guard_src1	1	1	0	
+fb65	head29953	guard_src1	2	1	0	
+fb65	head29953	guard_src1	3	1	0	
+fb65	head29953	guard_ids_dst	1	1	0	
+fb65	head29953	guard_ids_dst	2	1	0	
+fb65	head29953	guard_ids_dst	3	1	0	
+fb65	head29953	exact	1	1	0	0
+t100	head29953	stock	1	1	0	
+t100	head29953	stock	2	1	0	
+t100	head29953	stock	3	1	0	
+t100	head29953	guard	1	1	0	
+t100	head29953	guard	2	1	0	
+t100	head29953	guard	3	1	0	
+t100	head29953	guard_src1	1	1	0	
+t100	head29953	guard_src1	2	1	0	
+t100	head29953	guard_src1	3	1	0	
+t100	head29953	guard_ids_dst	1	1	0	
+t100	head29953	guard_ids_dst	2	1	0	
+t100	head29953	guard_ids_dst	3	1	0	
+t100	head29953	exact	1	1	0	0
+ids16	head29953	stock	1	1	0	
+ids16	head29953	stock	2	1	0	
+ids16	head29953	stock	3	1	0	
+ids16	head29953	guard	1	1	0	
+ids16	head29953	guard	2	1	0	
+ids16	head29953	guard	3	1	0	
+ids16	head29953	guard_ids_dst	1	1	0	
+ids16	head29953	guard_ids_dst	2	1	0	
+ids16	head29953	guard_ids_dst	3	1	0	
+ids16	head29953	exact	1	1	0	0
+ids64	head29953	stock	1	1	0	
+ids64	head29953	stock	2	1	0	
+ids64	head29953	stock	3	1	0	
+ids64	head29953	guard	1	1	0	
+ids64	head29953	guard	2	1	0	
+ids64	head29953	guard	3	1	0	
+ids64	head29953	guard_src1	1	1	0	
+ids64	head29953	guard_src1	2	1	0	
+ids64	head29953	guard_src1	3	1	0	
+ids64	head29953	guard_ids_dst	1	1	0	
+ids64	head29953	guard_ids_dst	2	1	0	
+ids64	head29953	guard_ids_dst	3	1	0	
+ids64	head29953	exact	1	1	0	0
+dense321	head29953	stock	1	1	0	
+dense321	head29953	stock	2	1	0	
+dense321	head29953	stock	3	1	0	
+dense321	head29953	guard	1	1	0	
+dense321	head29953	guard	2	1	0	
+dense321	head29953	guard	3	1	0	
+dense321	head29953	guard_src1	1	1	0	
+dense321	head29953	guard_src1	2	1	0	
+dense321	head29953	guard_src1	3	1	0	
+dense321	head29953	guard_ids_dst	1	1	0	
+dense321	head29953	guard_ids_dst	2	1	0	
+dense321	head29953	guard_ids_dst	3	1	0	
+dense321	head29953	exact	1	1	0	0
+ids16	p29953	guard	1	0	1	
+ids16	p29953	guard	2	0	1	
+ids16	p29953	guard	3	0	1	
+ids16	p29953	guard_ids_dst	1	0	1	
+ids16	p29953	guard_ids_dst	2	0	1	
+ids16	p29953	guard_ids_dst	3	0	1	
+ids16	p29953	exact	1	0	0	298
+one113	p29953	guard	1	0	1	
+one113	p29953	guard	2	0	1	
+one113	p29953	guard	3	0	1	
+one113	p29953	guard_ids_dst	1	0	1	
+one113	p29953	guard_ids_dst	2	0	1	
+one113	p29953	guard_ids_dst	3	0	1	
+one113	p29953	exact	1	0	0	12885
diff --git a/docs/maxusai/tasks/mmq-successor-results/results-head29953-sm75.tsv b/docs/maxusai/tasks/mmq-successor-results/results-head29953-sm75.tsv
new file mode 100644
index 00000000000..da324208b8b
--- /dev/null
+++ b/docs/maxusai/tasks/mmq-successor-results/results-head29953-sm75.tsv
@@ -0,0 +1,31 @@
+case	variant	mode	rep	passed	illegal_access	memcheck_errors
+ids16	p29953	guard_src1	1	0	1	
+ids16	p29953	guard_src1	2	0	1	
+ids16	p29953	guard_src1	3	0	1	
+ids16	p29953	guard_ids_dst	1	0	1	
+ids16	p29953	guard_ids_dst	2	0	1	
+ids16	p29953	guard_ids_dst	3	0	1	
+dense321	p29953	guard_src1	1	0	1	
+dense321	p29953	guard_src1	2	0	1	
+dense321	p29953	guard_src1	3	0	1	
+dense321	p29953	guard_ids_dst	1	1	0	
+dense321	p29953	guard_ids_dst	2	1	0	
+dense321	p29953	guard_ids_dst	3	1	0	
+ids16	head29953	stock	1	1	0	
+ids16	head29953	stock	2	1	0	
+ids16	head29953	stock	3	1	0	
+ids16	head29953	guard_src1	1	1	0	
+ids16	head29953	guard_src1	2	1	0	
+ids16	head29953	guard_src1	3	1	0	
+ids16	head29953	guard_ids_dst	1	1	0	
+ids16	head29953	guard_ids_dst	2	1	0	
+ids16	head29953	guard_ids_dst	3	1	0	
+dense321	head29953	stock	1	1	0	
+dense321	head29953	stock	2	1	0	
+dense321	head29953	stock	3	1	0	
+dense321	head29953	guard_src1	1	1	0	
+dense321	head29953	guard_src1	2	1	0	
+dense321	head29953	guard_src1	3	1	0	
+dense321	head29953	guard_ids_dst	1	1	0	
+dense321	head29953	guard_ids_dst	2	1	0	
+dense321	head29953	guard_ids_dst	3	1	0	
diff --git a/docs/maxusai/tasks/mmq-successor-results/verify-29953-guard.txt b/docs/maxusai/tasks/mmq-successor-results/verify-29953-guard.txt
new file mode 100644
index 00000000000..29b2632f130
--- /dev/null
+++ b/docs/maxusai/tasks/mmq-successor-results/verify-29953-guard.txt
@@ -0,0 +1,13 @@
+host load at start (1/5/15 min): 7.73 12.80 17.77, 32 cores
+1-head-unpatched                   rc=0   ok   want=ok   as-expected     8.1s
+patch applies to a pristine 3070d927f6c172242d56a2c1fb897e838e7b2632
+2-head-with-guard                  rc=0   ok   want=ok   as-expected    11.2s
+3-helper-shrunk-to-J-blocks        rc=2   fail want=fail as-expected     3.0s
+    errors=220, from the guard=220, tables named: ampere blackwell cdna gcn pascal_dp4a pascal_older rdna2 rdna3 rdna3_5 rdna4
+4-helper-rounded-to-one-warp       rc=2   fail want=fail as-expected     3.0s
+    errors=220, from the guard=220, tables named: ampere blackwell cdna gcn pascal_dp4a pascal_older rdna2 rdna3 rdna3_5 rdna4
+mutation 5 rewrote:     CASE(GGML_TYPE_MXFP4, 256, 1, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); ->     CASE(GGML_TYPE_MXFP4, 192, 1, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true);
+5-blackwell-config-at-192-threads  rc=2   fail want=fail as-expected     4.0s
+    errors=1, from the guard=1, tables named: blackwell
+6-guard-compiled-out               rc=0   ok   want=ok   as-expected     7.9s
+VERIFY29953GUARD COMPLETE 2026-10-05T02:33:16Z all_as_expected=1
diff --git a/docs/maxusai/tasks/mmq-variant.py b/docs/maxusai/tasks/mmq-variant.py
index ee82a2afb5c..58ce35c273e 100755
--- a/docs/maxusai/tasks/mmq-variant.py
+++ b/docs/maxusai/tasks/mmq-variant.py
@@ -15,6 +15,9 @@
   p27044      get_J_max(ne12*n_expert_used): the line llama.cpp#27044 proposed, our compat 903 until now
   successor   #448 as published: both branches pad for the widest tile that has a config,
               get_J_max(type, fallback, cc, 512), and ids_dst and the NVFP4 src1_scale by the same count
+  head29953   llama.cpp#29953 at its head, 3070d927f, verbatim: mmq.cu and mmq.cuh replaced by the head's blobs.
+              The head commit touches only those two files and its parent is dd266785c, so on a dd266785c
+              checkout this is a tree byte-identical to 3070d927f (checked before writing, not assumed).
   s448p       #448 amended: pad for the widest *padded* tile instead. A tile's y load copies the whole padded
               shared-memory y tile from global memory, GGML_PAD(J*sizeof(block_q8_1_mmq), nthreads*sizeof(int))
               bytes, which is more than J blocks when nthreads*4 does not divide J*144
@@ -34,6 +37,8 @@
 import sys
 
 HERE = os.path.dirname(os.path.abspath(__file__))
+HEAD29953 = "3070d927f6c172242d56a2c1fb897e838e7b2632"
+MMQ_FILES = ["ggml/src/ggml-cuda/mmq.cu", "ggml/src/ggml-cuda/mmq.cuh"]
 
 IDS_LINE   = "        ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, ne12) * sizeof(block_q8_1_mmq);"
 DENSE_LINE = "            ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, ne11) * sizeof(block_q8_1_mmq);"
@@ -96,6 +101,17 @@ def variant(src, path, v):
     s = open(path).read()
     if v == "ne12":
         return s
+    if v == "head29953":
+        git = lambda *a: subprocess.run(["git", *a], cwd=src, check=True, capture_output=True).stdout
+        parent = git("rev-parse", HEAD29953 + "^").decode().strip()
+        if git("rev-parse", "HEAD").decode().strip() != parent:
+            sys.exit("head29953 needs a checkout at the head's parent %s" % parent[:9])
+        touched = sorted(git("diff", "--name-only", parent, HEAD29953).decode().split())
+        if touched != sorted(MMQ_FILES):
+            sys.exit("the head commit touches %s, not just mmq.cu/mmq.cuh" % touched)
+        for rel in MMQ_FILES:
+            open(os.path.join(src, rel), "wb").write(git("show", HEAD29953 + ":" + rel))
+        return open(path).read()
     if v == "head":
         return s      # the PR's own head as fetched: already pads the tile and ids_dst, but not the y scales
     if v == "headyscale":
diff --git a/docs/maxusai/upstream-mmq-29953-material.md b/docs/maxusai/upstream-mmq-29953-material.md
index e66422eb5e3..1a53e90516c 100644
--- a/docs/maxusai/upstream-mmq-29953-material.md
+++ b/docs/maxusai/upstream-mmq-29953-material.md
@@ -3,7 +3,9 @@
 Companion to [upstream-mmq-successor-material.md](upstream-mmq-successor-material.md). That file was written
 against master `dd266785c` (#29941 merged) and argued for padding by the widest tile. This file supersedes its
 arithmetic: **the required src1 padding is larger than every rule proposed so far assumed, including ours**, and
-#29953 is the right shape of fix but still short in two places.
+#29953 as first published was the right shape of fix but short in two places. **Its head, `3070d927f`, fixes
+both** -- see the Conclusion below. Elsewhere in this file, "#29953" without "head" means the published version,
+`5bd8b0013`.
 
 > [!IMPORTANT]
 > **This is raw material, not text to post.** ggml-org/llama.cpp prohibits AI-written posts (bug reports, pull
@@ -12,7 +14,37 @@ arithmetic: **the required src1 padding is larger than every rule proposed so fa
 > the checks, the harness) is AI-generated, which the project allows **with disclosure**: fill in the template's
 > `AI usage disclosure:` line.
 
-## Upstream got there first: #29953 now carries both amendments (2026-10-05, 18:55Z)
+## Conclusion
+
+**#29953's head, `3070d927f`, is correct, and nothing about its correctness is owed upstream.** It chooses
+`J_best` once, on the host, before allocating; launches exactly that tile (`mul_mat_q_switch_J` now switches on
+`args.J_best`); pads src1 by that config's padded y tile, `GGML_PAD(J_best*sizeof(block_q8_1_mmq),
+nthreads*sizeof(int))` bytes, which is the load loop's own extent; and pads `ids_dst` by `J_best - 1`. The padding
+is computed from the config that is launched, so no shape can outgrow it.
+
+Measured on the exact head, on four architectures:
+
+| where | by | how | result |
+|---|---|---|---|
+| sm_120, RTX PRO 6000 Blackwell | this host | 13 cases x stock, three guard-page modes and exact-size memcheck, the published #29953 alongside as the positive control | **190/190 runs pass**, no abort, 0 memcheck errors; the control fails in every mode it ran in |
+| sm_75, RTX 2080 Ti | this host | `ids16` and `dense321` x stock and two guard-page modes, same control | **18/18**; the control aborts every run that can fault |
+| GB10, sm_121a | hclsys, on #29953 | exact-size `cudaMalloc` per buffer under memcheck `--padding 65536`: 16 routing seeds at `n = 100`, a sweep of 48 shapes at 2 seeds, and `test-backend-ops` | 0 errors: 16/16 seeds, where the published version failed 10/16; 96/96 runs; 2291/2291 |
+| gfx1151, Radeon 8060S | the ROCm host, #449 | the twelve cases and the hand-routed shapes, the device reporting the head's padding | 72/72 and 36/36 runs |
+
+hclsys tested `20c31408`, and all of `ggml/src/ggml-cuda` is byte-identical between that and the head, so it is the
+head's result. The CPU sweep with the head's rule added is covered on all ten architectures, by construction (see
+the note under "Ten architectures"). Details: "The exact head, measured".
+
+**What upstream still lacks is a guard.** The head computes the y tile's size twice -- in
+`mmq_get_nbytes_shared()` and in its new padding line, spelled differently -- and those two drifting apart is what
+this bug was. On the stock pool nothing would notice it happen again. "A guard upstream could take" has a patch on
+the head (+70/-9 lines): one helper for both, and a compile-time check over all ten config tables, verified with
+nvcc. Whether to offer it is the maintainer's call.
+
+**For the fork: keep compat 903 until the pin includes #29953**, then retire its padding by the gate under "For
+the fork" and keep the guard.
+
+## Upstream got there first: #29953's head carries both amendments (pushed 2026-10-04 18:54Z)
 
 **Everything below about src1 and `ids_dst` is now fixed upstream, found independently.** On #29953, `hclsys`
 reported from a GB10 (sm_121a, CUDA 13.0) exactly the two gaps this file was written about, with the same
@@ -31,6 +63,10 @@ The author pushed both the same day. At head `3070d927f` ("CUDA: fix MMQ out-of-
 and `mmq-amend-29953.patch` is of historical interest only.** Our independent arrival at the same formula, from a
 different direction (the CPU sweep rather than a seed sweep), is a cross-check on both.
 
+hclsys's re-test was of `20c31408`, not of the head. The head is a squash of `5bd8b001` and `20c31408` onto
+`dd266785`, and all of `ggml/src/ggml-cuda` is byte-identical between `20c31408` and `3070d927f` (checked
+2026-10-05), so their GB10 result is the head's result.
+
 ### RETRACTED: the NVFP4 y scales are not read out of bounds
 
 This file previously claimed that `src1_scale` is read up to `J_best - 1` floats past its end, by the same
@@ -72,7 +108,9 @@ nothing goes past the buffer. Rather than guess at the reconciliation, the claim
 anywhere. There is therefore **nothing outstanding for upstream from this work**: #29953's head fixes src1 and
 `ids_dst`, and the third buffer turned out not to need it.
 
-## Where it stands (2026-10-05)
+## Where it stood: #29953 as published (`5bd8b0013`)
+
+Kept as the record of what the published version did. Its head fixes both gaps below.
 
 - **#29953** ("CUDA: fix MMQ ncols rounding direction for alloc", head `5bd8b0013`, base `46847e615`) moves the
   tile-size choice ahead of the allocation and pads by exactly the `J_best` it is about to launch, so the two can
@@ -134,10 +172,11 @@ The padding each rule gives src1, as the argument of `ggml_cuda_mmq_get_J_max()`
 |---|---|---|---|
 | #24127 (b9992 .. `dd266785c^`) | `get_J_max(ne11)` -- 0 for broadcast gate/up | no | no |
 | #29941 (master, `dd266785c`) | `get_J_max(ne12)` | no | no |
-| #29953 | the launched `J_best`, in blocks | no | no |
-| #27044 (our compat 903 until now) | `get_J_max(ne12*n_expert_used)` | no | no |
+| #29953 as published (`5bd8b0013`) | the launched `J_best`, in blocks | no | no |
+| #27044 (the fork's 903 before #448) | `get_J_max(ne12*n_expert_used)` | no | no |
 | #448 as published | `get_J_max(type, fallback, cc, 512)` -- the widest tile | yes | yes |
-| **#448 amended** (recommended) | the widest **padded** tile, `max_J ceil(T/B)` | yes | yes |
+| **#448 amended** (the fork's 903 now) | the widest **padded** tile, `max_J ceil(T/B)` | yes | yes |
+| **#29953 head** (`3070d927f`) | the launched config's padded tile, `T` bytes | yes, `J_best - 1` | no, and none needed (retracted above) |
 
 ## The CPU check: every rule, 4.7 million shapes
 
@@ -161,7 +200,7 @@ Shapes left short, from [tasks/mmq-successor-results/check-rules.txt](tasks/mmq-
 |---|---|---|---|---|
 | #24127 (`ne11`) | 162,672 (worst 127) | 4,554,640 (worst 127) | all 4,717,328 (worst 127) | 705 (worst 9) |
 | #29941 (`ne12`) | 11,410 (worst 63) | 365,300 (worst 63) | all 4,717,328 | - |
-| **#29953** | **76,640 (worst 6)** | **1,531,586 (worst 6)** | **all 4,717,328** | **90 (worst 5)** |
+| **#29953 as published** | **76,640 (worst 6)** | **1,531,586 (worst 6)** | **all 4,717,328** | **90 (worst 5)** |
 | #27044 | 11,410 (worst 63) | 29,944 (worst 17) | all 4,717,328 | - |
 | #448 as published | 128 (worst 5), gfx1151 q2_K only | 25,256 (worst 5), gfx1151 q2_K only | covered | 45 (worst 5), gfx1151 only |
 | **#448 amended** | covered | covered | covered | covered |
@@ -178,8 +217,9 @@ Shapes left short, from [tasks/mmq-successor-results/check-rules.txt](tasks/mmq-
 >    row all along and the earlier text said two types anyway.
 
 "worst" is blocks (or `int32` entries for `ids_dst`) past the allocation. Per-architecture counts are in the
-program's output; on sm_120 #29953 is short in 4,488 one-expert and 134,640 multi-expert shapes. gfx1151's column
-comes from the host-side config tables as nvcc compiles them, not from a measurement on AMD hardware.
+program's output; on sm_120 #29953 as published is short in 4,488 one-expert and 134,640 multi-expert shapes.
+gfx1151's column comes from the host-side config tables as nvcc compiles them, not from a measurement on AMD
+hardware.
 
 A count here is **shapes where the rule can be short**, not shapes that fault. The check asks the question a
 padding rule has to answer -- is the allocation large enough for the worst routing this shape admits, one row in
@@ -187,8 +227,10 @@ the last non-empty expert's last tile -- and `ne12*n_expert_used` rows over `ne0
 Whether a particular run reads past depends on its routing, on where the pool put the buffer, and on what the
 pool mapped after it.
 
-Note the direction of the trade: #29953 is short in **more** shapes than #29941 but by **less** -- at most 6
-blocks instead of 63 -- because it fixes the rounding and leaves only the padded-tile term.
+The head's rule is not in this five-architecture run; it is in the ten-architecture run below, covered.
+
+Note the direction of the trade: #29953 as published is short in **more** shapes than #29941 but by **less** --
+at most 6 blocks instead of 63 -- because it fixes the rounding and leaves only the padded-tile term.
 
 ## Ten architectures: the amended rule is sufficient by construction
 
@@ -204,10 +246,22 @@ ids-branch shapes ([tasks/mmq-successor-results/check-rules-wide.txt](tasks/mmq-
 | rule | src1, >=2 experts/token | worst | short on which architectures |
 |---|---|---|---|
 | #29941 (`ne12`), master | 1,255,564 | 63 blocks | all ten |
-| #29953 at `3070d927f` | 3,520,756 | **12 blocks** | all ten |
-| #27044, main's 903 | 635,468 | 17 blocks | all ten |
+| #29953 as published (`5bd8b0013`) | 3,520,756 | **12 blocks** | all ten |
+| #27044, the fork's 903 before #448 | 635,468 | 17 blocks | all ten |
 | #448 as published | 626,512 | **7 blocks** | **gfx1151 25,256, CDNA3 576,000, RDNA4 25,256; zero on all seven NVIDIA** |
 | **#448 amended** | **covered** | - | none |
+| **#29953 head (`3070d927f`)** | **covered** | - | none |
+
+> [!NOTE]
+> **Corrected 2026-10-05.** This table used to label the published row "#29953 at `3070d927f`". The checker's
+> #29953 rule is the published one -- pad by the launched `J_best` in blocks -- and the head's rule was not in it
+> at all. Read literally, the old label said upstream's current fix is short in 3.5 million shapes, which is
+> false. The head's rule (the launched config's padded tile, and `J_best - 1` for `ids_dst`) is now in
+> [tasks/mmq-rules-check.cu](tasks/mmq-rules-check.cu) and is covered in all four tallies on all ten
+> architectures: 9,431,816 ids-branch shapes and 203,200 dense ones. Built against `dd266785c` and against
+> `3070d927f` itself, the check prints byte-identical output; the config tables are the same files in both. That is by construction -- its padding is
+> the load loop's own extent for the launched config -- so the sweep shows the model agrees, not that the head
+> is right; the measurements in the Conclusion are the evidence for that.
 
 **This widens the case for the amendment well beyond gfx1151.** `nthreads` is 512 on CDNA, not 256, so the
 padded tile is larger: the worst case is CDNA3, q4_0, non-fallback, `J = 64` with `nthreads = 512`, where
@@ -264,14 +318,15 @@ the last non-empty expert holding well under one tile, so its last tile reads pa
 
 Run it three ways. The results below are from this host; the full matrix is under "Results".
 
-| `MMQ_DEBUG_ALLOC` | what it exposes | `ne11` | `ne12` | #29953 | #27044 | #448 |
+| `MMQ_DEBUG_ALLOC` | what it exposes | `ne11` | `ne12` | #29953 as published | #27044 | #448 |
 |---|---|---|---|---|---|---|
 | unset (stock pool) | nothing -- the pool hides both | pass | pass | pass | pass | pass |
 | `guard:src1` | the padded-tile term | abort | abort | **abort** | pass | pass |
 | `guard:ids_dst` | the unpadded `ids_dst` | abort | abort | **abort** | abort | pass |
 
 `#27044` passes `guard:src1` here only because `get_J_max(ne12*n_expert_used) = get_J_max(256) = 128`, far more
-than the 21 blocks this shape needs; it is short elsewhere (see the CPU check).
+than the 21 blocks this shape needs; it is short elsewhere (see the CPU check). At #29953's head both guarded
+rows pass: see "The exact head, measured".
 
 **Why this shape.** `need` is the worst case over all routings. What a given case actually reads past the data
 is `ceil(T/B) - r` blocks, where `r` is the last non-empty expert's row count, so the worst case needs `r = 1`.
@@ -284,10 +339,10 @@ Whether a case can realise it is set by the mean rows per expert, `ne12*n_expert
 | `t100` | 800 | 256 | 3.1 | 112 | sometimes |
 | `ids64` | 6,400 | 256 | 25.0 | 112 | no: reads about 89 blocks past, which `ne12`'s 96 covers |
 
-So `ids16` is the case to hand over: one row per expert, every run, and a 5-block shortfall against #29953 rather
-than one. `ids64` is the control that shows the routing matters -- it passes `guard:src1` under every rule but
-`ne11`. An over-read is a function of the routing as well as the shape, which is the third reason this bug is
-hard to hit on purpose, after the pool and the fixup buffer.
+So `ids16` is the case to hand over: one row per expert, every run, and a 5-block shortfall against the published
+#29953 rather than one. `ids64` is the control that shows the routing matters -- it passes `guard:src1` under
+every rule but `ne11`. An over-read is a function of the routing as well as the shape, which is the third reason
+this bug is hard to hit on purpose, after the pool and the fixup buffer.
 
 ## The amendment
 
@@ -299,7 +354,8 @@ hard to hit on purpose, after the pool and the fixup buffer.
   instead of `J_best * sizeof(block_q8_1_mmq)`, in both branches. It is the expression
   `mmq_get_nbytes_shared()` already uses for the tile the load copies into.
 - `ids_dst` gets `J_best` more entries.
-- For NVFP4 with native FP4, `src1_scale` gets `J_best` more entries in both branches: the stream-k fixup
+- **Retracted, see the top of this file** -- kept so the patch can be read: for NVFP4 with native FP4,
+  `src1_scale` gets `J_best` more entries in both branches: the stream-k fixup
   write-back is called as `write_back(..., y_scale, I, I, J)`, so it reads `y_scale[j]` for all `J` columns of
   its tile. Found by reading the code; the GPU runs here use the q8_1 path, since an sm_120 build without `120a`
   has no native FP4.
@@ -336,7 +392,8 @@ padding. That is a refactor of the amendment, not a different fix: `mmq_get_nbyt
 `nbs_ids + nbs_x + GGML_PAD(nbs_y, config.nthreads*sizeof(int))` with `nbs_y = config.J*sizeof(block_q8_1_mmq)`,
 so the helper returns the identical value. **It is not what the measurements here were taken with** -- they used
 the 8-line form, which keeps `mmq.cuh` untouched -- so it is a suggestion for the maintainer to weigh, not a
-tested change.
+tested change. **Since built and verified against the head**, with a compile-time check added: see "A guard
+upstream could take".
 
 ## A second architecture, and the dense branch (sm_75)
 
@@ -411,20 +468,31 @@ this as an intermittent crash rather than a test failure.
 
 ## Reproduce
 
-- **CPU only, nvcc and no GPU**, from a llama.cpp checkout at `dd266785c` or later:
-  `tasks/build-check.sh <checkout> tasks/mmq-rules-check.cu ./mmq-rules-check && ./mmq-rules-check`.
-  The exit code is non-zero only if the amended rule leaves a shape short. `--design` lists shapes where only
+- **CPU only, nvcc and no GPU**, from a llama.cpp checkout at `dd266785c` or later, #29953's head included:
+  `tasks/build-check.sh <checkout> tasks/mmq-rules-check.cu mmq-rules-check && <checkout>/mmq-rules-check`.
+  The check carries its own copy of `ggml_cuda_mmq_get_J_max()`, which #29953 deletes; built against
+  `dd266785c` and against `3070d927f`, it prints output byte-identical to the build that used the library's
+  copy -- every rule's counts over all 9.4 million shapes.
+  `build-check.sh` compiles from inside the checkout, so a relative output path lands there. The exit code is
+  non-zero only if 903's rule or #29953's head leaves a shape short. `--design` lists shapes where only
   `ids_dst` is short.
 - **GPU:** `tasks/mmq-rules-gpu.sh <llama.cpp checkout at dd266785c> [out dir]`. It applies the test cases,
   builds each rule twice (with and without the debug allocator) via
   [tasks/mmq-variant.py](tasks/mmq-variant.py), and runs every case stock, under `guard:ids_dst`, under
   `guard:src1`, and under memcheck with exact-size allocations. It is resumable.
-- **The shortest path for a reviewer**, on any NVIDIA GPU:
-  1. apply `tasks/mmq-tests.patch` to #29953's branch,
-  2. apply `tasks/mmq-variant.py`'s debug allocator (`mmq-variant.py <checkout> p29953 --debug`),
-  3. `MMQ_DEBUG_ALLOC=guard:src1 test-backend-ops test -o MUL_MAT_ID -b CUDA0 -p 'type_a=q4_0,type_b=f32,n_mats=256,n_used=16,b=0,m=640,n=16,k=2560,'`
-     -> illegal memory access,
-  4. apply `tasks/mmq-amend-29953.patch` and run it again -> passes. Then repeat with `guard:ids_dst`.
+- **The shortest path for a reviewer**, on any NVIDIA GPU, from a `dd266785c` checkout with
+  `tasks/mmq-tests.patch` applied:
+  1. `mmq-variant.py <checkout> p29953 --debug` -- #29953 as published, plus the debug allocator -- and build
+     `test-backend-ops`;
+  2. `MMQ_DEBUG_ALLOC=guard:src1 test-backend-ops test -o MUL_MAT_ID -b CUDA0 -p 'type_a=q4_0,type_b=f32,n_mats=256,n_used=16,b=0,m=640,n=16,k=2560,'`
+     -> illegal memory access;
+  3. `mmq-variant.py <checkout> head29953 --debug` -- #29953's head, byte-identical to `3070d927f` plus the
+     allocator -- rebuild, and run it again -> passes. Then repeat both with `guard:ids_dst`.
+
+  [tasks/mmq-rules-head29953.sh](tasks/mmq-rules-head29953.sh) runs this and the rest of "The exact head,
+  measured".
+- **The guard**, no GPU: [tasks/mmq-29953-guard-verify.sh](tasks/mmq-29953-guard-verify.sh), next to a
+  llama.cpp clone that has `3070d927f`.
 
 ## Measured on gfx1151 (MaxusAI/ollama#449)
 
@@ -460,18 +528,151 @@ Full reply on the issue; build was a Ryzen AI Max+ 395 / Radeon 8060S, ROCm 7.2.
   main's rule. Its MoE experts are q4_K/q6_K/q8_0 (`qwen3.6:35b-a3b`, `qwen3-vl:30b-a3b`), q4_K/q5_0/q8_0
   (`gemma4:26b-a4b`) and q5_0/q8_0 (`nemotron3:33b`), read from that host's GGUF headers.
 
+## The exact head, measured
+
+`3070d927f` itself, not an equivalent. On a `dd266785c` checkout, `mmq-variant.py`'s `head29953` swaps in the
+head's `mmq.cu` and `mmq.cuh` -- the only two files the head commit touches -- after checking that the checkout is
+the head's parent and that the commit touches nothing else; without `--debug` the result is byte-identical to
+`3070d927f`. Same harness as "Results", run by
+[tasks/mmq-rules-head29953.sh](tasks/mmq-rules-head29953.sh), with the published #29953 alongside as the positive
+control: a clean head means something only where the same harness shows the published version failing. The
+controls for the combined guard, `guard:ids_dst` and memcheck were a second pass,
+[tasks/mmq-rules-head29953-controls.sh](tasks/mmq-rules-head29953-controls.sh).
+
+Totals, from `mmq-rules-table.py --summary` ("aborted" is a process that died; memcheck errors are summed over the
+runs made under the sanitizer):
+
+**sm_120**, RTX PRO 6000 Blackwell, CUDA 12.8, every case in [tasks/mmq-tests.patch](tasks/mmq-tests.patch):
+
+| rule | mode | cases | runs | passed | aborted | memcheck errors |
+|---|---|---|---|---|---|---|
+| #29953 as published | exact | 2 | 2 | 0 | 0 | 13183 |
+| #29953 as published | guard | 2 | 6 | 0 | 6 | - |
+| #29953 as published | guard_ids_dst | 2 | 6 | 0 | 6 | - |
+| #29953 as published | guard_src1 | 3 | 30 | 12 | 18 | - |
+| #29953 head (`3070d927f`) | stock | 13 | 39 | 39 | 0 | - |
+| #29953 head (`3070d927f`) | exact | 13 | 13 | 13 | 0 | 0 |
+| #29953 head (`3070d927f`) | guard | 13 | 39 | 39 | 0 | - |
+| #29953 head (`3070d927f`) | guard_ids_dst | 13 | 39 | 39 | 0 | - |
+| #29953 head (`3070d927f`) | guard_src1 | 13 | 60 | 60 | 0 | - |
+
+**sm_75**, RTX 2080 Ti, CUDA 12.8: `ids16`, and `dense321`, the plain `MUL_MAT` that the published version
+regressed against master:
+
+| rule | mode | cases | runs | passed | aborted | memcheck errors |
+|---|---|---|---|---|---|---|
+| #29953 as published | guard_ids_dst | 2 | 6 | 3 | 3 | - |
+| #29953 as published | guard_src1 | 2 | 6 | 0 | 6 | - |
+| #29953 head (`3070d927f`) | stock | 2 | 6 | 6 | 0 | - |
+| #29953 head (`3070d927f`) | guard_ids_dst | 2 | 6 | 6 | 0 | - |
+| #29953 head (`3070d927f`) | guard_src1 | 2 | 6 | 6 | 0 | - |
+
+- **The head passes everything**: 190 runs on sm_120 across 13 cases and five modes, and 18 on sm_75, with no
+  abort and no memcheck error.
+- **The control fails in every mode, on both cards**: `ids16` aborts in every guarded run of the published
+  version; under memcheck `one113` reports 12,885 errors and `ids16` 298; `dense321` aborts 3/3 under
+  `guard:src1` on sm_75. The three sm_75 control runs that pass are `dense321` under `guard:ids_dst`, which has
+  no `ids_dst` to guard.
+- **The one-block cases are a rate, as before**: under `guard:src1` the published version aborted `j100_b1` 3
+  times in 10 and `j100_b0` 5 times in 10 (4 and 1 in the earlier run), the head 0 in 10 on both.
+- Per-case tables: [matrix-head29953-sm120.md](tasks/mmq-successor-results/matrix-head29953-sm120.md) and
+  [matrix-head29953-sm75.md](tasks/mmq-successor-results/matrix-head29953-sm75.md), raw rows beside them.
+
+## A guard upstream could take
+
+The head's fix is right, but it leaves the y tile computed twice: `mmq_get_nbytes_shared()` sizes the
+shared-memory tile with `GGML_PAD(nbs_y, config.nthreads*sizeof(int))`, and `ggml_cuda_mul_mat_q` sizes the global
+src1 padding with its own ceiling-division expression. Those two drifting apart is what this bug was, and nothing
+would see it happen again: on the stock pool the over-read lands in mapped memory and `test-backend-ops` passes.
+
+[tasks/mmq-29953-y-tile-guard.patch](tasks/mmq-29953-y-tile-guard.patch) applies to `3070d927f` (70 lines added,
+9 removed, `mmq.cu` and `mmq.cuh`):
+
+- **One helper.** `ggml_cuda_mmq_get_nbytes_y_tile(config)` in `mmq.cuh` returns
+  `GGML_PAD(config.J*sizeof(block_q8_1_mmq), config.nthreads*sizeof(int))`. `mmq_get_nbytes_shared()` calls it
+  for its y term, and the selection loop carries `nbytes_y_tile_best` out instead of `nthreads_best`, so the
+  padding *is* the tile. **No runtime value changes:** every config table uses 128, 256 or 512 threads, where
+  `GGML_PAD` and the head's division agree.
+- **A compile-time check.** For every config table and type, the helper must cover the load loop's extent,
+  written from the loop with exact ceiling division -- `ceil(J*MMQ_TILE_Y_K / nthreads) * nthreads` ints -- not
+  from the helper or from `GGML_PAD`. One `static_assert` per (table, type), host pass only: the structure of
+  the fork's compat 903 guard (#451), which the ROCm host verified under hipcc.
+
+Verified with nvcc 12.8 for sm_120 by [tasks/mmq-29953-guard-verify.sh](tasks/mmq-29953-guard-verify.sh), which
+scores each build by why it failed -- every error must be the guard's and the tables it names must be the expected
+ones -- not just by its exit code
+([tasks/mmq-successor-results/verify-29953-guard.txt](tasks/mmq-successor-results/verify-29953-guard.txt)):
+
+```
+host load at start (1/5/15 min): 7.73 12.80 17.77, 32 cores
+1-head-unpatched                   rc=0   ok   want=ok   as-expected     8.1s
+patch applies to a pristine 3070d927f6c172242d56a2c1fb897e838e7b2632
+2-head-with-guard                  rc=0   ok   want=ok   as-expected    11.2s
+3-helper-shrunk-to-J-blocks        rc=2   fail want=fail as-expected     3.0s
+    errors=220, from the guard=220, tables named: ampere blackwell cdna gcn pascal_dp4a pascal_older rdna2 rdna3 rdna3_5 rdna4
+4-helper-rounded-to-one-warp       rc=2   fail want=fail as-expected     3.0s
+    errors=220, from the guard=220, tables named: ampere blackwell cdna gcn pascal_dp4a pascal_older rdna2 rdna3 rdna3_5 rdna4
+mutation 5 rewrote:     CASE(GGML_TYPE_MXFP4, 256, 1, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true); ->     CASE(GGML_TYPE_MXFP4, 192, 1, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true);
+5-blackwell-config-at-192-threads  rc=2   fail want=fail as-expected     4.0s
+    errors=1, from the guard=1, tables named: blackwell
+6-guard-compiled-out               rc=0   ok   want=ok   as-expected     7.9s
+VERIFY29953GUARD COMPLETE 2026-10-05T02:33:16Z all_as_expected=1
+```
+
+- **3** puts the published #29953's padding, `J` blocks, back into the helper: every table fails, 220 (table,
+  type) pairs, and nothing else does.
+- **4** is a plausible wrong simplification, rounding to one warp's ints: the same.
+- **5** is what the exact arithmetic is for. `GGML_PAD` is a bit mask that is only right for power-of-two
+  granularities, the config tables only assert `nthreads % 32 == 0`, and the kernel places `tile_x` with the
+  same `GGML_PAD`. A 192-thread config would compile today and mis-size the tile; with the guard, the build fails
+  and names that table only.
+- **Cost:** about 3 s per compile of `mmq.cu`, paid once per build because the guard runs in the host pass only,
+  however many architectures are targeted: a median of 11.1 s with it against 8.2 s compiled out, over five
+  interleaved repeats at a load of 7-8 on 32 cores
+  ([tasks/mmq-29953-guard-timing.sh](tasks/mmq-29953-guard-timing.sh),
+  [output](tasks/mmq-successor-results/guard-timing.txt)). A single build is too noisy to quote: two one-off
+  runs gave 9.8 against 8.9 s and 11.2 against 7.9 s.
+- nvcc's front end stops at 100 errors by default, which on the first run cut mutation 3's list to the first five
+  tables; the script passes `-Xcudafe --error_limit=100000` so every failing table is named.
+
+**Limits.** Only nvcc was run here: the HIP and MUSA branches are untested on this host. HIP keys off
+`__HIP_DEVICE_COMPILE__` because `vendors/hip.h` defines `__CUDA_ARCH__` in every HIP pass -- the trap the ROCm
+host caught in #451's first draft -- and MUSA is excluded, as in 903, because nothing here builds it. The guard
+cannot see call sites: a later change that pads src1 by something other than the helper would pass it. `ids_dst`
+needs no guard of this kind; its `J_best - 1` depends on nothing but `J`.
+
+**This is material, not a submission.** Whether to offer it upstream is the maintainer's call, and anything posted
+to ggml-org is written by hand.
+
 ## For the fork
 
-- **Compat 903 needs the same amendment.** As merged in #448 it pads by the widest tile, which covers every
-  NVIDIA shape checked but is short by up to 5 blocks in 25,384 gfx1151 shapes (q2_K and q3_K, where `J = 80`
-  with `nthreads = 256` needs 85 blocks and the widest tile gives 80). The ROCm host serves gfx1151, so this is
-  not hypothetical for the fork even though it is not a shape production's models reach. Amended on 2026-10-05.
-  - **The gfx1151 column is modelled, not measured**, and that is asked of the ROCm host in
-    MaxusAI/ollama#449: those counts come from `nvcc` compiling the host-side AMD branch of
-    `ggml_cuda_mmq_get_config()`, so if hipcc resolves the config tables differently the shortfall could be
-    larger, smaller or absent. Until that comes back, treat the gfx1151 numbers here as a prediction.
+- **Compat 903 stays as merged in #451** -- the widest padded tile, one helper for the y tile, and the
+  compile-time guard -- until the pin includes #29953. It is not replaced by a backport of #29953:
+  - the pin, b11081, predates the `prec_src1` refactor (`mmq.cuh` differs from #29953's base `dd266785c` by 230
+    changed lines), so a backport would be an adaptation, not upstream's code, and would still be dropped or
+    re-cut at the pin move;
+  - #29953 is not merged (open, review required, last changed 2026-10-04 18:55Z) and can still change;
+  - 903 is measured on sm_120 and gfx1151 and guarded at compile time, and the head has no guard;
+  - nothing is waiting on it: even v0.35.1's pin, b11232, predates #29941.
+- **The two fixes agree on what matters.** Both pad src1 by a padded y tile and both pad `ids_dst`. 903 pads for
+  the widest tile any launch could pick, the head for the tile it launches, so 903 reserves up to 16 KiB more src1
+  per call on sm_120 (18,432 bytes against the head's 2,048 at `J = 8`) and a few hundred bytes more `ids_dst`.
+- **Retiring 903**, at the first pin move that includes #29953's merge commit:
+  1. the patch series, without 903, applies to the new pin with plain `git apply`;
+  2. the head's rule is covered in `mmq-rules-check.cu` built against the new pin with `build-check.sh` (it
+     already builds against #29953's head unchanged: it carries its own copy of the `get_J_max()` that #29953
+     deletes);
+  3. `ids16` under `guard:src1` and `guard:ids_dst` on sm_120, and `dense321` under `guard:src1` on sm_75, pass
+     on the new pin, with the padding set back to `J_best` blocks as the positive control in the same harness
+     (`mmq-variant.py`'s anchors will need moving with the pin);
+
+  then 903's padding goes. **The guard stays**, re-cut guard-only, unless upstream has taken one -- that was
+  already the compat README's plan, and [tasks/mmq-29953-y-tile-guard.patch](tasks/mmq-29953-y-tile-guard.patch)
+  is its form on the head. If #29953 changes before it merges, re-run "The exact head, measured" against what
+  merged.
+- **gfx1151 is measured, not modelled** (#449): the widest-tile rule's gap reproduced, aborting 6/6, and the
+  amended rule fixed it, passing 6/6, every passing run matching the CPU backend.
 - **Nothing production serves is affected today.** Its MoE GGUFs use 8 of 128 experts (gemma4:26b-a4b), 8 of 256
   (qwen3.6:35b-a3b) and 6 of 128 plus a shared one (nemotron3:33b), all with `J = 128` at the image ubatch sizes
   it runs, where `T == J*B` and #29941's padding already covers src1. The `ids_dst` read happens on every MoE
   call but lands in `expert_bounds`.
-- **Re-cutting 903 against the amendment is the maintainer's call.**
diff --git a/llama/compat/README.md b/llama/compat/README.md
index 75fd31db6ed..3989d88a8c8 100644
--- a/llama/compat/README.md
+++ b/llama/compat/README.md
@@ -196,19 +196,29 @@ intentionally skipped so a developer can iterate on a local llama.cpp tree.
       compiles the guard out silently. The first draft did exactly that and
       passed both mutations.
     - **Limits:** it cannot see the call sites, so the allocations must keep
-      using `J_pad`. nvcc has not built it yet.
+      using `J_pad`.
+    - **nvcc agrees with hipcc.** Built with nvcc 12.8 for sm_120 before the
+      merge: it compiles, and both mutations fail by `static_assert` in the
+      same four tables.
     - **Behaviour is unchanged:** on gfx1151 the device reports the same `J`,
       `nthreads`, `need` and `pad` as the amended 903 for every hand-routed
       shape, and the twelve cases pass.
   - **At the next pin move.** From the first llama.cpp pin that contains
-    `dd266785c`, the ids branch's context line reads `ne12`, and from the first
-    that contains #29953 there is no `get_J_max()` call left to replace - pad
-    by the padded tile of the `J_best` that PR already computes
-    (`docs/maxusai/tasks/mmq-amend-29953.patch` is that form). Re-cut the hunk
-    against the new context; the rule stays the same. Keep the guard: point
-    `ggml_cuda_mmq_get_J_pad()`, or a helper the new allocation calls, at the
-    rule the pin ships, so that a later upstream re-tightening fails the build
-    here. Check the series on a checkout before building.
+    `dd266785c`, the ids branch's context line reads `ne12`: re-cut the hunk
+    against the new context; the rule stays the same.
+    - **From the first pin that contains llama.cpp#29953**, 903's padding is
+      redundant if #29953 merged as its head `3070d927f` did: that pads src1 by
+      the launched config's padded tile and `ids_dst` by `J_best - 1`, which is
+      measured clean on sm_120, sm_75, GB10 and gfx1151. Retire the padding by
+      the gate under "For the fork" in
+      `docs/maxusai/upstream-mmq-29953-material.md`. If it merged in another
+      form, re-run that file's exact-head passes against what merged first.
+    - **Keep the guard** unless upstream has taken one: re-cut 903 as
+      guard-only, one helper for the y tile and a `static_assert` against the
+      load loop, so that a later upstream change that shrinks the padding fails
+      the build here. `docs/maxusai/tasks/mmq-29953-y-tile-guard.patch` is that
+      form on #29953's head, verified with nvcc.
+    - Check the series on a checkout before building.
   - **References.**
     - `docs/maxusai/qwen35moe-mmq-investigation.md`: the original diagnosis.
     - `docs/maxusai/mmq-padding-regression-window.md`: the affected build range.