Skip to content

docs(mmq): conclude llama.cpp#29953 -- the head is correct, measured, and a guard for upstream #454

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
glennneuber merged 1 commit into from
Oct 5, 2026
Merged

docs(mmq): conclude llama.cpp#29953 -- the head is correct, measured, and a guard for upstream #454

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
34 changes: 34 additions & 0 deletions docs/maxusai/tasks/mmq-29953-guard-timing.sh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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)"
102 changes: 102 additions & 0 deletions docs/maxusai/tasks/mmq-29953-guard-verify.sh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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"
129 changes: 129 additions & 0 deletions docs/maxusai/tasks/mmq-29953-y-tile-guard.patch
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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>
Loading
Loading