Change sets 6-7: CDNA MMQ tile retune + SSD chunk size (pp16384 1775 → 2590, +45.9%)

modified-files/ggml/src/ggml-cuda/ssm-scan.cu

@@ -22,7 +22,7 @@ using namespace cub;
 #define SSM_SSD_MAX_TOKENS (SSM_SSD_DT_BLOCK * SSM_SSD_DT_MAX_ITEMS)
 
 // Chunk size for chunked SSD. Caps matmul cost at O(chunk^2) per chunk.
-#define SSM_SSD_CHUNK_SIZE 256
+#define SSM_SSD_CHUNK_SIZE 128
 
 // We would like to keep pragma unroll for cases where L_template is not 0,
 // so we suppress the clang transformation warning.

patches/07-ssd-chunk-size-cdna.patch

@@ -0,0 +1,13 @@
+diff --git a/ggml/src/ggml-cuda/ssm-scan.cu b/ggml/src/ggml-cuda/ssm-scan.cu
+index 45f172c..147a924 100644
+--- a/ggml/src/ggml-cuda/ssm-scan.cu
++++ b/ggml/src/ggml-cuda/ssm-scan.cu
+@@ -22,7 +22,7 @@ using namespace cub;
+ #define SSM_SSD_MAX_TOKENS (SSM_SSD_DT_BLOCK * SSM_SSD_DT_MAX_ITEMS)
+
+ // Chunk size for chunked SSD. Caps matmul cost at O(chunk^2) per chunk.
+-#define SSM_SSD_CHUNK_SIZE 256
++#define SSM_SSD_CHUNK_SIZE 128
+
+ // We would like to keep pragma unroll for cases where L_template is not 0,
+ // so we suppress the clang transformation warning.

tools/patch_mmq_cdna_add_j128.py

@@ -0,0 +1,129 @@
+#!/usr/bin/env python3
+"""Add J=128 MMQ tile entries to the CDNA config table.
+
+WHY. `mul_mat_q_switch_J` (mmq.cuh) picks the tile width J at runtime:
+
+    for (int J = 8; J <= 128 && ntiles_J_best > 1; J += 8) {
+        config = ggml_cuda_mmq_get_config(type, J, fallback, cc);
+        if (config.type == GGML_TYPE_COUNT) continue;      // no entry -> skip
+        if (mmq_get_nbytes_shared(config, cc) > smpbo) continue;
+        ...keep the largest J that strictly reduces ceil(ncols_max/J)
+    }
+
+so J is capped by whatever the arch table happens to contain. Per-arch maxima:
+
+    ampere    128        blackwell 128        rdna4  128
+    cdna       64        pascal     64        rdna2   64
+
+CDNA stops at 64 not because 128 was measured and rejected, but because the
+pre-refactor code returned 64 for any HIP target lacking Turing-style MMA
+(`get_mmq_x_max_device`), and PR #24127 transcribed that constant into the new
+table. CDNA was excluded from the 128 branch for being MFMA rather than WMMA --
+an instruction-family test, not a capacity one.
+
+At prefill with -ub 2048 the J switch saturates at the table maximum, so on CDNA
+every large-batch matmul runs J=64 tiles when it could run J=128, halving the
+tokens covered per tile and doubling the tile count.
+
+LDS BUDGET. mmq_get_nbytes_shared is
+    J*4 + I*sram_stride*4 + pad(J*sizeof(block_q8_1_mmq), nthreads*4)
+with sizeof(block_q8_1_mmq)=144 and sram_stride 76 (Q8_0/Q8_1/Q6_K), 84 (Q3_K),
+100 (Q2_K). Against gfx90a's 64 KiB limit, at the current I=64 / nthreads=256:
+
+    J=64,  Q8_1  28928 B (28.2 KiB)      J=128, Q8_1  38400 B (37.5 KiB)
+    J=64,  Q2_K  35072 B (34.2 KiB)      J=128, Q2_K  44544 B (43.5 KiB)
+
+so J=128 fits with room to spare. Note this would NOT have fit as comfortably at
+the old I=128/nthreads=512 (Q2_K would have been 70912 B, over the limit), which
+is a reason the two changes belong together.
+
+CAVEAT, STATED UP FRONT. aviallon measured mmq_x=128 as catastrophic on MI210
+pre-refactor (-60% at n=128 for IQ3_XXS) due to 412 B of scratch spill at
+min_blocks=2. Whether that survives the refactor's SRAM-layout changes is
+unknown. This is an experiment; the switch will simply ignore any entry that
+does not fit, but it will happily select one that fits and is slower. A/B it.
+
+Entries are cloned from each type's existing J=64 line so sram_layout, K_vram,
+stream_k and fallback stay whatever that type already uses.
+
+    python patch_mmq_cdna_add_j128.py [--check] [--revert] [--j N]
+"""
+import re
+import sys
+
+TARGET = "ggml/src/ggml-cuda/mmq-config-cdna.cuh"
+MARKER = "MI210_ADD_J"
+
+CASE_RE = re.compile(
+    r"^(?P<indent>\s*)CASE\("
+    r"(?P<type>GGML_TYPE_[A-Z0-9_]+),\s*"
+    r"(?P<nthreads>\d+),\s*(?P<occupancy>\d+),\s*(?P<I>\d+),\s*(?P<J>\d+),\s*"
+    r"(?P<sram>[A-Z0-9_]+),\s*(?P<kvram>[A-Z0-9_]+),\s*"
+    r"(?P<streamk>true|false),\s*(?P<fallback>true|false)\s*\);\s*(?P<trail>//.*)?$"
+)
+
+
+def main() -> int:
+    check = "--check" in sys.argv
+    revert = "--revert" in sys.argv
+    new_j = 128
+    if "--j" in sys.argv:
+        new_j = int(sys.argv[sys.argv.index("--j") + 1])
+
+    with open(TARGET) as f:
+        lines = f.readlines()
+    patched = any(MARKER in ln for ln in lines)
+
+    if check:
+        n = sum(1 for ln in lines if MARKER in ln)
+        print(f"{'PATCHED' if patched else 'not patched'}  {TARGET}  ({n} added entries)")
+        return 0 if patched else 1
+
+    if revert:
+        if not patched:
+            print("not patched; nothing to revert")
+            return 0
+        out = [ln for ln in lines if MARKER not in ln]
+        with open(TARGET, "w") as f:
+            f.writelines(out)
+        print(f"reverted ({len(lines) - len(out)} entries removed)")
+        return 0
+
+    if patched:
+        print("already patched")
+        return 0
+
+    if new_j % 8:
+        print("ERROR: J must be a multiple of 8", file=sys.stderr)
+        return 1
+
+    # Clone each J=64 line into a new J=<new_j> line placed directly after it,
+    # so each type keeps its own layout/K_vram/stream_k/fallback settings and the
+    # table stays grouped by type.
+    out, n = [], 0
+    for ln in lines:
+        out.append(ln)
+        m = CASE_RE.match(ln.rstrip("\n"))
+        if not m or int(m["J"]) != 64:
+            continue
+        g = m.groupdict()
+        out.append(
+            f"{g['indent']}CASE({g['type']}, {g['nthreads']}, {g['occupancy']}, "
+            f"{g['I']}, {new_j}, {g['sram']}, {g['kvram']}, "
+            f"{g['streamk']}, {g['fallback']});  // {MARKER}\n"
+        )
+        n += 1
+
+    if n == 0:
+        print("ERROR: found no J=64 entries to clone -- table layout changed; "
+              "re-derive rather than forcing.", file=sys.stderr)
+        return 1
+
+    with open(TARGET, "w") as f:
+        f.writelines(out)
+    print(f"added {n} J={new_j} entries (cloned from each type's J=64 config)")
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())

tools/patch_mmq_cdna_no_streamk.py

@@ -25,25 +25,34 @@
 turned out 6.5% slower than the assumption behind it. A/B against the same
 binary, and read generated tokens both ways.
 
-Scoped to the K-quants (Q2_K..Q6_K) because that is what an i1-Q4_K_M model
-actually dispatches; leaving the other types alone keeps the experiment narrow
-and the result attributable.
+SCOPE. Initially this covered only the K-quants, on the assumption that an
+i1-Q4_K_M model dispatches nothing else. The profile disproved that: the three
+hottest MMQ kernels were
+
+    mul_mat_q<(ggml_type)12, 64, false>   23.5%   Q4_K
+    mul_mat_q<(ggml_type)6,  64, false>   14.8%   Q5_0
+    mul_mat_q<(ggml_type)8,  64, false>   12.6%   Q8_0
+
+so Q5_0 and Q8_0 -- 27.4% of all prefill GPU time -- were still on stream-k.
+The hypothesis is about the workload shape (MoE, many tiles per expert), not
+about any property of a particular quant, so this now covers every quantised
+type in the table. GGML_TYPE_COUNT is the unreachable sentinel and is already
+false, so it simply does not match.
 
     python patch_mmq_cdna_no_streamk.py [--check] [--revert]
 """
 import re
 import sys
 
 TARGET = "ggml/src/ggml-cuda/mmq-config-cdna.cuh"
-TYPES = ("Q2_K", "Q3_K", "Q4_K", "Q5_K", "Q6_K")
 MARKER = "MI210_NO_STREAMK"
 
 # CASE(type, nthreads, occupancy, I, J, sram_layout, K_vram, stream_k, fallback)
-# Rewrite only the 8th argument, and only on lines for the listed types. Anchor
-# on the two trailing args so a config whose shape changes upstream fails to
-# match instead of silently corrupting a different field.
+# Rewrite only the 8th argument. Anchor on the two trailing args so a config
+# whose shape changes upstream fails to match instead of silently corrupting a
+# different field.
 LINE_RE = re.compile(
-    r"^(\s*CASE\(GGML_TYPE_(?:" + "|".join(TYPES) + r"),[^;]*?,\s*)true(\s*,\s*(?:true|false)\s*\);)$"
+    r"^(\s*CASE\(GGML_TYPE_[A-Z0-9_]+,[^;]*?,\s*)true(\s*,\s*(?:true|false)\s*\);)$"
 )
 
 

tools/tune_mmq_cdna.py

@@ -0,0 +1,147 @@
+#!/usr/bin/env python3
+"""Rewrite the CDNA MMQ config fields so they can be swept empirically.
+
+The table in ggml-cuda/mmq-config-cdna.cuh is a list of
+
+    CASE(type, nthreads, occupancy, I, J, sram_layout, K_vram, stream_k, fallback)
+
+and mmq.cuh states outright that these "should not affect results, only
+speed/register pressure/shared memory use". Every CDNA entry currently uses
+nthreads=512, occupancy=1, I=128, and one file covers CDNA1/2/3, so nothing here
+was necessarily tuned on gfx90a.
+
+`occupancy` is the second argument of __launch_bounds__. That parameter means
+different things on the two platforms: CUDA reads it as minBlocksPerMultiprocessor,
+HIP reads it as MIN_WARPS_PER_EXECUTION_UNIT. So on gfx90a `occupancy=1` asks the
+compiler for as little as one wave per SIMD, which lets it spend registers freely
+and hide very little latency. Raising it forces the register budget down in
+exchange for more resident waves -- worth measuring rather than assuming.
+
+Usage:
+    python tune_mmq_cdna.py --nthreads 256 --occupancy 2 --tile-i 64
+    python tune_mmq_cdna.py --revert
+
+Only the fields named on the command line are touched; the rest keep their
+upstream values. Restricted to entries whose J matches --only-j when given, so a
+sweep can target just the hot configs instead of the whole table.
+"""
+import argparse
+import re
+import sys
+
+TARGET = "ggml/src/ggml-cuda/mmq-config-cdna.cuh"
+BACKUP = TARGET + ".tune-orig"
+
+# CASE(type_, nthreads_, occupancy_, I_, J_, sram_layout_, K_vram_, stream_k_, fallback_)
+CASE_RE = re.compile(
+    r"^(?P<indent>\s*)CASE\("
+    r"(?P<type>GGML_TYPE_[A-Z0-9_]+),\s*"
+    r"(?P<nthreads>\d+),\s*"
+    r"(?P<occupancy>\d+),\s*"
+    r"(?P<I>\d+),\s*"
+    r"(?P<J>\d+),\s*"
+    r"(?P<sram>[A-Z0-9_]+),\s*"
+    r"(?P<kvram>[A-Z0-9_]+),\s*"
+    r"(?P<streamk>true|false),\s*"
+    r"(?P<fallback>true|false)\s*\);\s*(?P<trail>//.*)?$"
+)
+
+
+def main() -> int:
+    ap = argparse.ArgumentParser()
+    ap.add_argument("--nthreads", type=int)
+    ap.add_argument("--occupancy", type=int)
+    ap.add_argument("--tile-i", type=int, dest="tile_i")
+    ap.add_argument("--stream-k", choices=("true", "false"))
+    ap.add_argument("--only-j", type=int,
+                    help="restrict to entries with this J (e.g. 64 for the hot configs)")
+    ap.add_argument("--revert", action="store_true")
+    ap.add_argument("--show", action="store_true")
+    a = ap.parse_args()
+
+    if a.revert:
+        import os
+        if not os.path.exists(BACKUP):
+            print("no backup; nothing to revert")
+            return 0
+        with open(BACKUP) as f:
+            src = f.read()
+        with open(TARGET, "w") as f:
+            f.write(src)
+        os.remove(BACKUP)
+        print("reverted to pre-tuning config")
+        return 0
+
+    with open(TARGET) as f:
+        lines = f.readlines()
+
+    # Keep one pristine copy so a sweep always starts from the same baseline
+    # rather than compounding edits from the previous iteration.
+    import os
+    if not os.path.exists(BACKUP):
+        with open(BACKUP, "w") as f:
+            f.writelines(lines)
+
+    if a.show:
+        seen = {}
+        for ln in lines:
+            m = CASE_RE.match(ln.rstrip("\n"))
+            if m:
+                k = (m["nthreads"], m["occupancy"], m["I"], m["J"], m["streamk"])
+                seen[k] = seen.get(k, 0) + 1
+        print(f"{'nthr':>5} {'occ':>4} {'I':>4} {'J':>4} {'stream_k':>9}  count")
+        for k, n in sorted(seen.items(), key=lambda kv: -kv[1]):
+            print(f"{k[0]:>5} {k[1]:>4} {k[2]:>4} {k[3]:>4} {k[4]:>9}  {n}")
+        return 0
+
+    out, n = [], 0
+    for ln in lines:
+        m = CASE_RE.match(ln.rstrip("\n"))
+        if not m:
+            out.append(ln)
+            continue
+        if a.only_j is not None and int(m["J"]) != a.only_j:
+            out.append(ln)
+            continue
+        g = m.groupdict()
+        if a.nthreads is not None:
+            g["nthreads"] = str(a.nthreads)
+        if a.occupancy is not None:
+            g["occupancy"] = str(a.occupancy)
+        if a.tile_i is not None:
+            g["I"] = str(a.tile_i)
+        if a.stream_k is not None:
+            g["streamk"] = a.stream_k
+        trail = ("  " + g["trail"]) if g.get("trail") else ""
+        out.append(
+            f"{g['indent']}CASE({g['type']}, {g['nthreads']}, {g['occupancy']}, "
+            f"{g['I']}, {g['J']}, {g['sram']}, {g['kvram']}, "
+            f"{g['streamk']}, {g['fallback']});{trail}\n"
+        )
+        n += 1
+
+    if n == 0:
+        print("ERROR: no CASE lines matched -- layout changed upstream, re-derive.",
+              file=sys.stderr)
+        return 1
+
+    # The static_asserts in the CASE macro are the real guard rails; mirror the
+    # cheap ones here so a bad sweep value fails before a 10-minute build.
+    if a.nthreads is not None and (a.nthreads % 32 or a.nthreads > 512):
+        print("ERROR: nthreads must be a multiple of 32 and <= 512", file=sys.stderr)
+        return 1
+    if a.occupancy is not None and a.occupancy > 8:
+        print("ERROR: occupancy must be <= 8", file=sys.stderr)
+        return 1
+    if a.tile_i is not None and a.tile_i % 32:
+        print("ERROR: I must be a multiple of 32", file=sys.stderr)
+        return 1
+
+    with open(TARGET, "w") as f:
+        f.writelines(out)
+    print(f"rewrote {n} CASE entries")
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())