perf(hip): keep prefill GEMMs spill-free on clang 23

CMakeLists.txt

@@ -406,6 +406,28 @@ if (BUILD_TESTING)
       $<TARGET_FILE:gufo>)
   set_tests_properties(serve_cli_test PROPERTIES LABELS "cpu;server;cli" TIMEOUT 30)
 
+  add_test(NAME kernel_resources_parser_test
+    COMMAND ${Python3_EXECUTABLE}
+      ${CMAKE_CURRENT_SOURCE_DIR}/tests/tools/kernel_resources_test.py)
+  set_tests_properties(kernel_resources_parser_test PROPERTIES
+    LABELS "cpu;kernels" TIMEOUT 30)
+
+  if(ENGINE_ENABLE_HIP)
+    # A compiler update can make a kernel spill with no source change. Clang 23
+    # increased the scratch of the Q8 prefill GEMM from 0 to 480 bytes. This
+    # test fails if a kernel uses more scratch than the baseline allows. It
+    # reads the code objects in the binary and does not use a GPU.
+    get_filename_component(GUFO_HIP_BIN_DIR ${CMAKE_HIP_COMPILER} DIRECTORY)
+    add_test(NAME kernel_resources_test
+      COMMAND ${Python3_EXECUTABLE}
+        ${CMAKE_CURRENT_SOURCE_DIR}/tools/ci/check-kernel-resources.py
+        $<TARGET_FILE:gufo>
+        --baseline ${CMAKE_CURRENT_SOURCE_DIR}/tools/ci/kernel-resources.json
+        --bundler ${GUFO_HIP_BIN_DIR}/clang-offload-bundler)
+    set_tests_properties(kernel_resources_test PROPERTIES
+      LABELS "cpu;kernels" TIMEOUT 120)
+  endif()
+
   add_subdirectory(tests/models/deepseek_v4_flash)
   add_subdirectory(tests/models/qwen38_flash_next)
 

cmake/Checks.cmake

@@ -22,7 +22,8 @@ list(REMOVE_ITEM gufo_pr_tests
 list(APPEND gufo_pr_tests
   "qwen38_flash_next\\.config" "qwen38_flash_next\\.mtp_sampling"
   "ds4\\.sampling" "ds4\\.template" "ds4\\.cli"
-  gufo_version gufo_help serve_cli_test eval_http_test functional_runner_test)
+  gufo_version gufo_help serve_cli_test eval_http_test functional_runner_test
+  kernel_resources_parser_test)
 list(JOIN gufo_pr_tests "|" gufo_pr_pattern)
 add_custom_target(check-pr
   COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure --no-tests=error

docs/PERFORMANCE.md

@@ -88,6 +88,8 @@ cost in the result.
 - Compile production kernels for `gfx1151` and treat Wave32 as explicit.
 - Prove alignment before vector loads and provide bounded tail paths.
 - Track VGPR, LDS, occupancy, and scratch for retained kernels.
+  `kernel_resources_test` fails if a kernel uses more scratch than
+  `tools/ci/kernel-resources.json` allows (see "Kernel resource check").
 - Keep distinct decode and prefill routes where their reuse differs.
 - Read only live KV spans and reuse GQA/MQA K/V across query heads.
 - Retain hipBLASLt or rocBLAS as the baseline for supported matrix shapes.
@@ -228,6 +230,27 @@ written to the artifact. `gufo bench` remains the direct model-path
 microbenchmark; use this serving harness for TTFT, ITL, queueing, and
 concurrency decisions.
 
+### Kernel resource check
+
+The check reads the code objects in a binary. It does not use a GPU. Run it
+after each compiler or ROCm update, and after each kernel change:
+
+```sh
+python3 tools/ci/check-kernel-resources.py result/bin/gufo \
+  --baseline tools/ci/kernel-resources.json
+```
+
+To compare two builds, for example two compilers, use `--compare`. It lists
+each kernel that has more scratch, more spilled VGPRs or less VGPR occupancy
+than in the base build:
+
+```sh
+python3 tools/ci/check-kernel-resources.py new/bin/gufo --compare old/bin/gufo
+```
+
+Change the baseline only after you decide that a new allowance is correct.
+Use `--write-baseline` and record the reason with `--note`.
+
 ### HIP allocation diagnostic
 
 Compare allocation, mapped-registration, first-touch, warm-access, copy,

docs/models/qwen3.8-27b/EXPERIMENTS.md

@@ -100,6 +100,7 @@
 | Concurrent generated-reply forks | Retained: freeze the live generated checkpoint before branching, within the snapshot budget. C1 avoids the copy; Q4/Q8 greedy/seeded forks, disk restoration and cancellation pass. |
 | Visible causal attention tails | Retained: FP32 FMA over visible tail keys removes chunk-boundary drift while complete tiles keep WMMA. Q8 full logits/features match at 8K across scheduling budgets; four 32K DFlash2 continuations match isolated AR. Later packed V loads and row-first scale reuse recover PP speed without spills; large shallow chunks use the same packed route. |
 | Separate causal-tail passes, helper functions and deeper V prefetch | Rejected: exact outputs, but extra work or register pressure makes attention slower. |
+| Clang 23 prefill GEMM spills | Retained: a `sched_barrier` fence per K block in BK=2 W8A8 tiles, no K-block unroll in BK=4 tiles, SLP off for the two K-quant units. Clang 23 (ROCm 10.0) Q8 pp2048 291 → 464, pp64 112 → 308, Q4 pp64 230 → 335 (published 0.8.1 image: Q8 454/322, Q4 327). Greedy output and prefill validation are byte-identical. Clang 22: pp2048 and decode do not change, 16–64 token prefill is 1.2–1.7% slower. |
 
 Q4 C1 AR leads the pinned d0/d32K controls; Q8 C1 AR matches them.
 DFlash2 retains AR output in the focused shallow/deep and concurrent checks.

docs/models/qwen3.8-flash-next/EXPERIMENTS.md

@@ -74,6 +74,7 @@
 | GPU greedy penalties and linear CPU anchor selection | Retained; exact FP64 penalties, unchanged proposals and snapshots. Short heat-pump tg400: 32.12 → 33.80 tok/s; C2/4/6/8 improve 7.6/8.5/16.3/14.2%. Unpenalized control unchanged. [Evidence](artifacts/penalty-verification.json). |
 | Prefill projection, attention and indexer kernels | Retained; bit-identical logit dumps and greedy hashes, plus a bitwise GEMM sweep against hipBLASLt for every routed n ≤ 4096. pp4096 +1.4% at d0 and +2.5–2.8% at 32K/115K (interleaved ABBA, `-r 6`, two machines); repository prompts 6.5K–102K +3.2% median; tg unchanged. The indexer projections use the own kernel only when hipBLASLt selects algorithm 4438 and otherwise stay on the library. |
 | 128 n-gram readers | Retained with 4096-token chunks, which double each gather; readers block in `pread`, so the pool sets the queue depth. Readers are woken per read, so small gathers do not wake the pool; 96 more threads. |
+| Paired routed GEMM code-cache blend | Retained: blend the four cached chunks with masks. A conditional over `uint4` lvalues becomes an address select, and clang 23 then keeps the cache in scratch (80 B, 0 spills). Clang 23 pp2048 +0.7%, pp8–pp64 +1.6–2.7%, MTP pp2048 +0.9%. Output is byte-identical. Clang 22 pp2048 is 0.7% slower. |
 
 Separate d32K pp2048 profiling attributes 29.1% of kernel time to MoE, 34.9%
 to dense projections and 12.6% to attention/indexing. Final-tile catch-up

src/models/qwen/CMakeLists.txt

@@ -97,4 +97,13 @@ if (ENGINE_ENABLE_HIP)
     TARGET_DIRECTORY gufo_core
     PROPERTIES COMPILE_OPTIONS
     "-mwavefrontsize64;-gline-tables-only;-Xarch_device;-mllvm=-amdgpu-sched-strategy=iterative-ilp")
+  # SLP vectorization puts the K-quant scales and offsets of the blocked
+  # prefill GEMMs into wide register tuples. With clang 23, the Q3_K and Q6_K
+  # 128x128 tiles then spill. Disable SLP for the two units on all compilers.
+  set_property(
+    SOURCE
+      ${PROJECT_SOURCE_DIR}/src/models/qwen/hip/kernels/prefill_quant_gemm.hip
+      ${PROJECT_SOURCE_DIR}/src/models/qwen/hip/kernels/prefill_quant_wave64.hip
+    TARGET_DIRECTORY gufo_core
+    APPEND PROPERTY COMPILE_OPTIONS -fno-slp-vectorize)
 endif()

src/models/qwen/hip/kernels/prefill_quant_gemm.hip

@@ -558,7 +558,11 @@ __launch_bounds__(256) __global__ void W8A8BlockedWmmaGEMMKernel(
       fetch_stage(kb0 + BK);
     }
 
-#pragma unroll
+    // Unroll the K blocks of a stage only if there are two. In a fully
+    // unrolled BK=4 stage, the scheduler can move the LDS operand reads of all
+    // blocks above the WMMAs of the first block. Clang 23 does this. Then the
+    // kernel needs more than 256 VGPRs and spills.
+#pragma unroll(BK <= 2 ? BK : 1)
     for (int kb = 0; kb < BK; ++kb) {
       int32x4_t a0[kWaveRowTiles];
       int32x4_t a1[kWaveRowTiles];
@@ -603,6 +607,10 @@ __launch_bounds__(256) __global__ void W8A8BlockedWmmaGEMMKernel(
           }
         }
       }
+      // Keep the LDS operand reads of the next K block below the WMMAs of this
+      // K block. This prevents the same spill in an unrolled BK=2 stage. The
+      // arithmetic order does not change.
+      __builtin_amdgcn_sched_barrier(0);
     }
     __syncthreads();
   }

src/models/qwen/hip/kernels/prefill_quant_gemm.hpp

@@ -385,7 +385,9 @@ __launch_bounds__(WM * WN * WaveSize, 1) __global__
       fetch_stage(kb0 + BK);
     }
 
-#pragma unroll
+    // Unroll the K blocks of a stage only if there are two. A fully unrolled
+    // BK=4 stage spills with clang 23 (see W8A8BlockedWmmaGEMMKernel).
+#pragma unroll(BK <= 2 ? BK : 1)
     for (int kb = 0; kb < BK; ++kb) {
       int32x4_t a0[kWaveRowTiles];
       int32x4_t a1[kWaveRowTiles];

src/models/qwen38_flash_next/kernels/rocm/kernels.hip.cpp

@@ -4075,10 +4075,23 @@ __launch_bounds__(256) __global__
         const int sb32 = (kb0 % 8) + f_c;
         if constexpr (kPair) {
           const int group = sb32 / 2;
-          f_codes[u] = group == 0   ? code_cache[u][0]
-                       : group == 1 ? code_cache[u][1]
-                       : group == 2 ? code_cache[u][2]
-                                    : code_cache[u][3];
+          // Blend the four cached chunks with masks. Do not select one slot
+          // with the runtime group: the compiler can change any such selection
+          // into an address select. Clang 23 does this, and then the cache is
+          // in scratch memory, not in VGPRs.
+          const std::uint32_t m0 = 0u - static_cast<std::uint32_t>(group == 0);
+          const std::uint32_t m1 = 0u - static_cast<std::uint32_t>(group == 1);
+          const std::uint32_t m2 = 0u - static_cast<std::uint32_t>(group == 2);
+          const std::uint32_t m3 = 0u - static_cast<std::uint32_t>(group == 3);
+          const uint4& c0 = code_cache[u][0];
+          const uint4& c1 = code_cache[u][1];
+          const uint4& c2 = code_cache[u][2];
+          const uint4& c3 = code_cache[u][3];
+          f_codes[u] =
+              make_uint4((c0.x & m0) | (c1.x & m1) | (c2.x & m2) | (c3.x & m3),
+                         (c0.y & m0) | (c1.y & m1) | (c2.y & m2) | (c3.y & m3),
+                         (c0.z & m0) | (c1.z & m1) | (c2.z & m2) | (c3.z & m3),
+                         (c0.w & m0) | (c1.w & m1) | (c2.w & m2) | (c3.w & m3));
         } else {
           f_codes[u] = blk[kCodeChunk + (sb32 / 2) * 2 + f_c];
         }

tests/tools/kernel_resources_test.py

@@ -0,0 +1,57 @@
+#!/usr/bin/env python3
+"""CPU-only regressions for compressed HIP bundle extraction."""
+
+import importlib.util
+import os
+from pathlib import Path
+import struct
+import subprocess
+import unittest
+from unittest import mock
+
+ROOT = Path(__file__).resolve().parents[2]
+SPEC = importlib.util.spec_from_file_location(
+    "kernel_resources", ROOT / "tools/ci/check-kernel-resources.py")
+RESOURCES = importlib.util.module_from_spec(SPEC)
+SPEC.loader.exec_module(RESOURCES)
+
+
+class KernelResourcesTest(unittest.TestCase):
+    def test_bundler_gets_one_output_per_target(self):
+        targets = ["host-x86_64-unknown-linux-gnu-", "hipv4-amdgcn-amd-amdhsa--gfx1151"]
+        payloads = [b"", b"device code"]
+        # An unsupported compression method requires the external fallback,
+        # independently of the Python version or installed zstd packages.
+        blob = b"CCOB" + struct.pack("<HH", 3, 65535) + bytes(24)
+
+        def bundler(command, **kwargs):
+            if "--list" in command:
+                return subprocess.CompletedProcess(command, 0, "\n".join(targets))
+            outputs = [arg.removeprefix("--output=") for arg in command
+                       if arg.startswith("--output=")]
+            self.assertEqual(len(outputs), len(targets))
+            self.assertIn("--targets=" + ",".join(targets), command)
+            for output, payload in zip(outputs, payloads):
+                Path(output).write_bytes(payload)
+            return subprocess.CompletedProcess(command, 0)
+
+        with mock.patch.object(RESOURCES.subprocess, "run", side_effect=bundler):
+            result = RESOURCES.inflate_ccob(blob, "/hip/clang-offload-bundler")
+        self.assertEqual(list(RESOURCES.bundle_entries(result)), list(zip(targets, payloads)))
+
+    def test_hip_bundler_precedes_unrelated_path_compiler(self):
+        with mock.patch.dict(os.environ, {}, clear=True), \
+             mock.patch.object(RESOURCES.shutil, "which", side_effect=lambda name: "/path/" + name), \
+             mock.patch.object(RESOURCES.os.path, "isfile", return_value=True), \
+             mock.patch.object(RESOURCES.subprocess, "run", return_value=
+                               subprocess.CompletedProcess([], 0, "/hip/bin\n")):
+            self.assertEqual(RESOURCES.find_bundler(None), "/hip/bin/clang-offload-bundler")
+
+    def test_missing_explicit_bundler_fails(self):
+        with mock.patch.object(RESOURCES.os.path, "isfile", return_value=False):
+            with self.assertRaisesRegex(SystemExit, "offload bundler not found"):
+                RESOURCES.find_bundler("/missing/bundler")
+
+
+if __name__ == "__main__":
+    unittest.main()