perf(qwen): tune Qwen3.8-27B verification GEMM geometries at 2-8 rows (#874)

Co-authored-by: D3rPapts <[email protected]>
Co-authored-by: Claude Opus 5.5 <[email protected]>

src/models/qwen/CMakeLists.txt

@@ -62,6 +62,7 @@ if (ENGINE_ENABLE_HIP)
     ${PROJECT_SOURCE_DIR}/src/models/qwen/hip/kernels/ssm_row_split.hip
     ${PROJECT_SOURCE_DIR}/src/models/qwen/hip/kernels/swiglu.hip
     ${PROJECT_SOURCE_DIR}/src/models/qwen/hip/kernels/attention_batched.hip
+    ${PROJECT_SOURCE_DIR}/src/models/qwen/hip/kernels/verify_gemm_27b.hip
     ${PROJECT_SOURCE_DIR}/src/models/qwen/hip/kernels/attention_decode.hip
     ${PROJECT_SOURCE_DIR}/src/models/qwen/hip/kernels/attention_decode_verify_wave64.hip
     ${PROJECT_SOURCE_DIR}/src/models/qwen/hip/kernels/attention_decode_graph.hip
@@ -119,4 +120,7 @@ if (ENGINE_ENABLE_HIP AND GUFO_BUILD_TOOLS)
   add_executable(qwen27b_continuation_probe EXCLUDE_FROM_ALL
     ${PROJECT_SOURCE_DIR}/tests/models/qwen27b/continuation_probe.cpp)
   target_link_libraries(qwen27b_continuation_probe PRIVATE gufo_core)
+  add_executable(qwen27b_verify_sweep EXCLUDE_FROM_ALL
+    ${PROJECT_SOURCE_DIR}/tools/qwen27b/verify_sweep/sweep.hip)
+  target_link_libraries(qwen27b_verify_sweep PRIVATE gufo_core hip::device)
 endif()

src/models/qwen/hip/batched_decode.cpp

@@ -11,6 +11,7 @@
 #include "src/models/qwen/hip/detail/attention_policy.hpp"
 #include "src/models/qwen/hip/detail/decode_step.hpp"
 #include "src/models/qwen/hip/executor.hpp"
+#include "src/models/qwen/hip/kernels/verify_gemm_27b.hpp"
 #include "src/models/qwen/hip/ops.hpp"
 
 namespace gufo::hip {
@@ -29,6 +30,9 @@ void LaunchProjection(const models::QwenTensorRef& weight, const float* input,
                       float* output, std::size_t batch_size,
                       std::size_t output_size, std::size_t input_size,
                       hipStream_t stream) {
+  if (TryLaunchVerifyGemm27(weight.type, weight.data, input, output, batch_size,
+                            output_size, input_size, stream))
+    return;
   if (SupportsExactSharedProjection(weight.type)) {
     LaunchBatchedQuantGEMMFp32(weight.type, weight.data, input, output,
                                batch_size, output_size, input_size, stream);

src/models/qwen/hip/dflash/executor.cpp

@@ -17,6 +17,7 @@
 #include "src/models/qwen/dflash_weights.hpp"
 #include "src/models/qwen/hip/dflash.hpp"
 #include "src/models/qwen/hip/kernels/dflash_kernels.hpp"
+#include "src/models/qwen/hip/kernels/verify_gemm_27b.hpp"
 #include "src/models/qwen/hip/mtp/detail/allocation.hpp"
 #include "src/models/qwen/hip/ops.hpp"
 
@@ -407,6 +408,9 @@ void QwenDFlashGpuExecutor::RunBlockGemm(const models::QwenTensorRef& weight,
     // matrix-core W8A8 route once the shared kernel read activations as float4
     // (11.3 vs 12.2 ms per draft block), so there is one route here.
     if (batch_size <= kDFlashMaxSharedBatch) {
+      if (TryLaunchVerifyGemm27(weight.type, weight.data, input, output,
+                                batch_size, output_size, input_size, stream_))
+        return;
       LaunchBatchedQuantGEMMFp32(weight.type, weight.data, input, output,
                                  batch_size, output_size, input_size, stream_);
       return;
@@ -471,7 +475,9 @@ void QwenDFlashGpuExecutor::RunInjectGemm(const models::QwenTensorRef& weight,
       LaunchExactBf16GEMMFp32SmallBatch(weight.data, chunk_input, chunk_output,
                                         chunk, output_size, input_size,
                                         stream_);
-    } else {
+    } else if (!TryLaunchVerifyGemm27(weight.type, weight.data, chunk_input,
+                                      chunk_output, chunk, output_size,
+                                      input_size, stream_)) {
       LaunchBatchedQuantGEMMFp32(weight.type, weight.data, chunk_input,
                                  chunk_output, chunk, output_size, input_size,
                                  stream_);

src/models/qwen/hip/kernels/verify_gemm_27b.hip

@@ -0,0 +1,322 @@
+#include "src/models/qwen/hip/kernels/small_batch_gemm.hpp"
+#include "src/models/qwen/hip/kernels/verify_gemm_27b.hpp"
+
+namespace gufo::hip {
+namespace {
+
+using core::GgmlType;
+
+// One lane walks its sub-blocks in the same order and the lanes reduce in the
+// same butterfly for every block shape, so these launches only change how
+// output rows are spread over waves and how much activation staging a block
+// shares. Lanes = 16 runs two sixteen-lane groups per wave; each keeps both
+// original lane partials and combines them in the same butterfly, so it is
+// exact too. MinWaves is only the occupancy hint of the launch bounds.
+template<GgmlType T, std::size_t Batch, std::uint32_t Waves, std::size_t Rows,
+         std::size_t Tiles, std::size_t TokensPerStep = Batch,
+         std::uint32_t MinWaves = 4, std::uint32_t Lanes = 32>
+bool Launch(const void* w, const float* x, float* y, std::size_t m,
+            std::size_t k, hipStream_t stream) {
+  constexpr std::size_t kRowsPerBlock = Waves * (32 / Lanes) * Rows;
+  SmallBatchKQuantExactFp32GEMMKernel<Waves, Batch, Rows, T, Tiles, MinWaves, 1,
+                                      true, TokensPerStep, 32, Lanes>
+      <<<static_cast<unsigned>((m + kRowsPerBlock - 1) / kRowsPerBlock),
+         Waves * 32, 0, stream>>>(w, x, y, m, k);
+  return true;
+}
+
+bool Is(std::size_t m, std::size_t k, std::size_t want_m, std::size_t want_k) {
+  return m == want_m && k == want_k;
+}
+
+// Measured on gfx1151 under Linux with ROCm 10.1 (tools/qwen27b/verify_sweep):
+// real UD-Q4_K_XL target and DFlash2 Q4_K_M draft weights rotated through
+// >= 512 MiB of copies, alternating rounds, median, every shape
+// bit-identical to the generic dispatch. A shape is listed only where it beat
+// LaunchBatchedQuantGEMMFp32 by 5% or more; the comment gives that gain.
+bool LaunchQ4K(const void* w, const float* x, float* y, std::size_t batch,
+               std::size_t m, std::size_t k, hipStream_t stream) {
+  constexpr auto T = GgmlType::kQ4_K;
+  if (Is(m, k, 1024, 5120)) {  // attn_k/attn_v
+    switch (batch) {
+      case 2:  // +15%
+        return Launch<T, 2, 8, 4, 2, 2, 4, 32>(w, x, y, m, k, stream);
+      case 3:  // +19%
+        return Launch<T, 3, 8, 4, 2, 3, 4, 32>(w, x, y, m, k, stream);
+      case 4:  // +14%
+        return Launch<T, 4, 8, 4, 2, 4, 4, 32>(w, x, y, m, k, stream);
+      case 5:  // +55%
+        return Launch<T, 5, 16, 4, 2, 5, 4, 32>(w, x, y, m, k, stream);
+      case 6:  // +50%
+        return Launch<T, 6, 16, 4, 2, 6, 4, 32>(w, x, y, m, k, stream);
+      case 7:  // +41%
+        return Launch<T, 7, 16, 4, 2, 7, 4, 32>(w, x, y, m, k, stream);
+      case 8:  // +65%
+        return Launch<T, 8, 16, 4, 2, 8, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  if (Is(m, k, 1280, 5120)) {  // draft conv
+    switch (batch) {
+      case 2:  // +10%
+        return Launch<T, 2, 16, 4, 2, 2, 4, 32>(w, x, y, m, k, stream);
+      case 3:  // +16%
+        return Launch<T, 3, 16, 4, 2, 3, 4, 32>(w, x, y, m, k, stream);
+      case 4:  // +9%
+        return Launch<T, 4, 8, 4, 2, 4, 4, 32>(w, x, y, m, k, stream);
+      case 5:  // +49%
+        return Launch<T, 5, 16, 4, 2, 5, 4, 32>(w, x, y, m, k, stream);
+      case 6:  // +40%
+        return Launch<T, 6, 8, 4, 1, 1, 4, 32>(w, x, y, m, k, stream);
+      case 7:  // +38%
+        return Launch<T, 7, 16, 4, 2, 7, 4, 32>(w, x, y, m, k, stream);
+      case 8:  // +10%
+        return Launch<T, 8, 16, 4, 2, 8, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  if (Is(m, k, 4096, 5120)) {  // draft attn_q
+    switch (batch) {
+      case 2:  // +12%
+        return Launch<T, 2, 16, 4, 1, 2, 8, 16>(w, x, y, m, k, stream);
+      case 5:  // +7%
+        return Launch<T, 5, 4, 4, 1, 5, 4, 16>(w, x, y, m, k, stream);
+      case 6:  // +11%
+        return Launch<T, 6, 4, 4, 1, 6, 8, 16>(w, x, y, m, k, stream);
+      case 8:  // +28%
+        return Launch<T, 8, 4, 4, 1, 8, 8, 16>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  if (Is(m, k, 5120, 4096)) {  // draft attn_output
+    switch (batch) {
+      case 2:  // +13%
+        return Launch<T, 2, 8, 4, 2, 2, 4, 32>(w, x, y, m, k, stream);
+      case 5:  // +13%
+        return Launch<T, 5, 8, 4, 1, 1, 4, 32>(w, x, y, m, k, stream);
+      case 6:  // +12%
+        return Launch<T, 6, 8, 4, 1, 1, 4, 32>(w, x, y, m, k, stream);
+      case 8:  // +36%
+        return Launch<T, 8, 8, 4, 1, 8, 8, 16>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  if (Is(m, k, 5120, 17408)) {  // ffn_down
+    switch (batch) {
+      case 4:  // +8%
+        return Launch<T, 4, 16, 4, 2, 4, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  if (Is(m, k, 6144, 5120)) {  // attn_gate
+    switch (batch) {
+      case 3:  // +9%
+        return Launch<T, 3, 8, 4, 2, 3, 4, 32>(w, x, y, m, k, stream);
+      case 4:  // +11%
+        return Launch<T, 4, 8, 4, 1, 1, 4, 32>(w, x, y, m, k, stream);
+      case 5:  // +9%
+        return Launch<T, 5, 8, 4, 2, 5, 4, 32>(w, x, y, m, k, stream);
+      case 6:  // +12%
+        return Launch<T, 6, 8, 4, 1, 1, 4, 32>(w, x, y, m, k, stream);
+      case 8:  // +5%
+        return Launch<T, 8, 8, 4, 2, 8, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  if (Is(m, k, 10240, 5120)) {  // attn_qkv
+    switch (batch) {
+      case 8:  // +8%
+        return Launch<T, 8, 16, 4, 2, 8, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  if (Is(m, k, 12288, 5120)) {  // attn_q
+    switch (batch) {
+      case 5:  // +9%
+        return Launch<T, 5, 16, 4, 2, 5, 4, 32>(w, x, y, m, k, stream);
+      case 6:  // +7%
+        return Launch<T, 6, 8, 4, 1, 1, 4, 32>(w, x, y, m, k, stream);
+      case 8:  // +15%
+        return Launch<T, 8, 8, 4, 1, 1, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  return false;
+}
+
+bool LaunchQ5K(const void* w, const float* x, float* y, std::size_t batch,
+               std::size_t m, std::size_t k, hipStream_t stream) {
+  constexpr auto T = GgmlType::kQ5_K;
+  if (Is(m, k, 1024, 5120)) {  // attn_k/attn_v
+    switch (batch) {
+      case 2:  // +10%
+        return Launch<T, 2, 16, 4, 2, 2, 4, 32>(w, x, y, m, k, stream);
+      case 3:  // +20%
+        return Launch<T, 3, 8, 4, 2, 3, 4, 32>(w, x, y, m, k, stream);
+      case 4:  // +19%
+        return Launch<T, 4, 8, 4, 2, 4, 4, 32>(w, x, y, m, k, stream);
+      case 5:  // +48%
+        return Launch<T, 5, 8, 4, 1, 1, 4, 32>(w, x, y, m, k, stream);
+      case 6:  // +42%
+        return Launch<T, 6, 16, 4, 2, 6, 4, 32>(w, x, y, m, k, stream);
+      case 7:  // +22%
+        return Launch<T, 7, 16, 4, 2, 7, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  return false;
+}
+
+bool LaunchQ6K(const void* w, const float* x, float* y, std::size_t batch,
+               std::size_t m, std::size_t k, hipStream_t stream) {
+  constexpr auto T = GgmlType::kQ6_K;
+  if (Is(m, k, 1024, 5120)) {  // attn_k/attn_v
+    switch (batch) {
+      case 3:  // +24%
+        return Launch<T, 3, 8, 4, 2, 3, 4, 32>(w, x, y, m, k, stream);
+      case 4:  // +24%
+        return Launch<T, 4, 8, 4, 1, 1, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  if (Is(m, k, 5120, 6144)) {  // ssm_out/attn_output
+    switch (batch) {
+      case 2:  // +9%
+        return Launch<T, 2, 8, 4, 2, 2, 4, 32>(w, x, y, m, k, stream);
+      case 3:  // +10%
+        return Launch<T, 3, 8, 2, 1, 3, 8, 16>(w, x, y, m, k, stream);
+      case 6:  // +6%
+        return Launch<T, 6, 8, 4, 1, 1, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  if (Is(m, k, 5120, 17408)) {  // ffn_down
+    switch (batch) {
+      case 2:  // +9%
+        return Launch<T, 2, 16, 4, 2, 2, 4, 32>(w, x, y, m, k, stream);
+      case 3:  // +12%
+        return Launch<T, 3, 16, 4, 1, 1, 4, 32>(w, x, y, m, k, stream);
+      case 7:  // +6%
+        return Launch<T, 7, 8, 4, 1, 1, 4, 32>(w, x, y, m, k, stream);
+      case 8:  // +13%
+        return Launch<T, 8, 16, 4, 2, 8, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  return false;
+}
+
+bool LaunchIQ4XS(const void* w, const float* x, float* y, std::size_t batch,
+                 std::size_t m, std::size_t k, hipStream_t stream) {
+  constexpr auto T = GgmlType::kIQ4_XS;
+  if (Is(m, k, 5120, 17408)) {  // ffn_down
+    switch (batch) {
+      case 5:  // +40%
+        return Launch<T, 5, 16, 4, 2, 5, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  if (Is(m, k, 6144, 5120)) {  // attn_gate
+    switch (batch) {
+      case 2:  // +15%
+        return Launch<T, 2, 4, 3, 1, 2, 8, 16>(w, x, y, m, k, stream);
+      case 3:  // +27%
+        return Launch<T, 3, 4, 3, 1, 3, 8, 16>(w, x, y, m, k, stream);
+      case 4:  // +20%
+        return Launch<T, 4, 8, 4, 2, 4, 4, 32>(w, x, y, m, k, stream);
+      case 5:  // +18%
+        return Launch<T, 5, 8, 4, 2, 5, 4, 32>(w, x, y, m, k, stream);
+      case 6:  // +8%
+        return Launch<T, 6, 8, 4, 2, 6, 4, 32>(w, x, y, m, k, stream);
+      case 7:  // +8%
+        return Launch<T, 7, 8, 4, 2, 7, 4, 32>(w, x, y, m, k, stream);
+      case 8:  // +10%
+        return Launch<T, 8, 4, 4, 1, 8, 4, 16>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  if (Is(m, k, 10240, 5120)) {  // attn_qkv
+    switch (batch) {
+      case 2:  // +9%
+        return Launch<T, 2, 16, 4, 1, 2, 4, 16>(w, x, y, m, k, stream);
+      case 3:  // +6%
+        return Launch<T, 3, 4, 2, 1, 3, 8, 16>(w, x, y, m, k, stream);
+      case 4:  // +6%
+        return Launch<T, 4, 4, 3, 1, 4, 8, 16>(w, x, y, m, k, stream);
+      case 5:  // +6%
+        return Launch<T, 5, 8, 4, 2, 5, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  if (Is(m, k, 12288, 5120)) {  // attn_q
+    switch (batch) {
+      case 3:  // +9%
+        return Launch<T, 3, 8, 4, 2, 3, 4, 32>(w, x, y, m, k, stream);
+      case 4:  // +8%
+        return Launch<T, 4, 8, 4, 2, 4, 4, 32>(w, x, y, m, k, stream);
+      case 5:  // +5%
+        return Launch<T, 5, 8, 4, 2, 5, 4, 32>(w, x, y, m, k, stream);
+      case 8:  // +5%
+        return Launch<T, 8, 8, 4, 1, 1, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  if (Is(m, k, 17408, 5120)) {  // ffn_gate/ffn_up
+    switch (batch) {
+      case 5:  // +45%
+        return Launch<T, 5, 8, 4, 2, 5, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  if (Is(m, k, 34816, 5120)) {  // joined ffn_gate+ffn_up
+    switch (batch) {
+      case 3:  // +6%
+        return Launch<T, 3, 16, 4, 1, 3, 4, 16>(w, x, y, m, k, stream);
+      case 5:  // +50%
+        return Launch<T, 5, 8, 4, 2, 5, 4, 32>(w, x, y, m, k, stream);
+      default:
+        return false;
+    }
+  }
+  return false;
+}
+
+}  // namespace
+
+bool TryLaunchVerifyGemm27(GgmlType type, const void* w, const float* x,
+                           float* y, std::size_t batch, std::size_t m,
+                           std::size_t k, hipStream_t stream) {
+  if (batch < 2 || batch > 8)
+    return false;
+  switch (type) {
+    case GgmlType::kQ4_K:
+      return LaunchQ4K(w, x, y, batch, m, k, stream);
+    case GgmlType::kQ5_K:
+      return LaunchQ5K(w, x, y, batch, m, k, stream);
+    case GgmlType::kQ6_K:
+      return LaunchQ6K(w, x, y, batch, m, k, stream);
+    case GgmlType::kIQ4_XS:
+      return LaunchIQ4XS(w, x, y, batch, m, k, stream);
+    default:
+      return false;
+  }
+}
+
+}  // namespace gufo::hip

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

@@ -0,0 +1,25 @@
+#ifndef GUFO_MODELS_QWEN_HIP_KERNELS_VERIFY_GEMM_27B_HPP_
+#define GUFO_MODELS_QWEN_HIP_KERNELS_VERIFY_GEMM_27B_HPP_
+
+#include <hip/hip_runtime.h>
+
+#include <cstddef>
+
+#include "src/core/gguf_reader.hpp"
+
+namespace gufo::hip {
+
+/// Qwen3.8-27B verification and DFlash2 draft projections at 2..8 rows:
+/// launch geometries measured faster on gfx1151 for the UD-Q4_K_XL target and
+/// the Q4_K_M DFlash2 draft shapes (tools/qwen27b/verify_sweep). Every
+/// geometry is the production exact small-batch kernel with another block
+/// shape; outputs are bit-identical to LaunchBatchedQuantGEMMFp32. Returns
+/// false when no tuned geometry applies.
+[[nodiscard]] bool TryLaunchVerifyGemm27(core::GgmlType type, const void* w,
+                                         const float* x, float* y,
+                                         std::size_t batch, std::size_t m,
+                                         std::size_t k, hipStream_t stream);
+
+}  // namespace gufo::hip
+
+#endif  // GUFO_MODELS_QWEN_HIP_KERNELS_VERIFY_GEMM_27B_HPP_

tools/README.md

@@ -75,6 +75,10 @@ compares separate and fused draft normalization/RoPE.
 `qwen27b/prefill_gemm_bench.hip` compares production Q8_0 gate/up GEMM
 with separate or fused SwiGLU, including every activation byte and scale.
 It replaces the old standalone W8A8 kernel copies.
+`qwen27b/verify_sweep/` measures the 2–8 row launch shapes of
+`verify_gemm_27b.hip` on a GGUF's real projections against the production
+dispatch (exact outputs, weights rotated beyond the MALL cache); build the
+`qwen27b_verify_sweep` target and run `sweep.py`, see its README.
 `qwen27b/dflash_reference.py` checks a GPU trace against
 pinned upstream PyTorch operators using the same GGUF weights. See
 `docs/models/qwen3.8-27b/BENCHMARKS.md` and its quality report for commands and evidence.

tools/qwen27b/verify_sweep/README.md

@@ -0,0 +1,78 @@
+# verify_sweep: tune verification launch shapes for a Qwen GGUF
+
+During speculative decoding the target model verifies 2-8 draft rows at a
+time, and the DFlash2 draft runs its block at the same widths. Each K-quant
+projection then runs `SmallBatchKQuantExactFp32GEMMKernel`, and how its work
+is split over waves (waves per block, rows per wave, 16- or 32-lane groups,
+staging tiles) decides how close the GPU gets to its memory bandwidth.
+`src/models/qwen/hip/kernels/verify_gemm_27b.hip` holds the shapes measured
+faster on gfx1151 for Qwen3.8-27B UD-Q4_K_XL and its Q4_K_M DFlash2 draft; any
+projection it does not list keeps the generic `LaunchBatchedQuantGEMMFp32`
+dispatch.
+
+A different GGUF, quantization mix, draft setting, GPU or toolchain changes
+which projections and row counts matter, so the shapes have to be measured
+again. This tool does that in a few minutes.
+
+## Use
+
+```sh
+# production optimization, as the timings should be
+nix develop -c cmake --preset release -DGUFO_BUILD_TOOLS=ON
+nix develop -c cmake --build --preset release --target qwen27b_verify_sweep
+# stop every gufo or llama.cpp server: the GPU has to be idle
+nix develop -c python3 tools/qwen27b/verify_sweep/sweep.py \
+  --binary build/release/qwen27b_verify_sweep /path/to/Qwen3.8-27B-UD-Q4_K_XL.gguf
+```
+
+`sweep.py` lists the model's projections (one example tensor per type and
+shape, with the joined gate/up matrix where both halves are stored back to
+back), sweeps every projection used by at least `--min-layers` layers at
+`--batches` rows, and prints the best exact candidate per projection plus the
+winning shapes as `Launch<...>` lines. Options: `--batches 3,4,5`,
+`--rounds 9`, `--min-gain 5` (percent), `--out sweep.log` for the raw lines.
+For a DFlash2 draft GGUF pass `--min-layers 1`: the draft has only five
+layers, so some projections occur once. Its shapes go into `LaunchQ4K`, the
+draft block of `verify_gemm_27b.hip`.
+
+The binary also runs on its own:
+`qwen27b_verify_sweep <gguf> <offset> <type> <m> <k> <batch> [rounds]`, with
+the absolute file offset of the tensor and the GGUF type id (12 Q4_K, 13 Q5_K,
+14 Q6_K, 23 IQ4_XS).
+
+## Method
+
+- Weights are the real tensors from the GGUF, read at their file offset.
+- Production is what `LaunchProjection` runs: `TryLaunchVerifyGemm27`, then
+  `LaunchBatchedQuantGEMMFp32`. The `generic` candidate is
+  `LaunchBatchedQuantGEMMFp32` alone, the dispatch without any listed shape;
+  where it beats production, the listed shape is a regression and is printed
+  for removal. Every candidate's output is compared bit for bit with
+  production; inexact candidates are reported and never timed.
+- Each launch uses another copy of the weights, enough copies for at least
+  512 MiB, so no run reads from the 32 MiB MALL cache. Without that, matrices
+  under ~35 MB look faster than the memory allows.
+- Production and candidates alternate within each round; the result is the
+  median over the rounds, 20 launches per sample after 3 warm-up launches.
+
+## Taking a result over
+
+Add the shape to `verify_gemm_27b.hip` under its format and projection,
+keyed by type, `m`, `k` and row count, only where it beats `generic` by the
+`--min-gain` margin. The target and the DFlash2 draft share some projection
+shapes (the Q4_K FFN and attention matrices), and one entry serves both, so
+sweep both GGUFs and judge a shared shape on both results. Then rebuild,
+rerun the sweep to confirm no listed shape loses to `generic`, confirm greedy
+output is unchanged end to end and measure a real workload: a per-launch gain
+only matters in proportion to the projection's share of the cycle.
+
+## Limits
+
+- Only the Qwen GPU path (`src/models/qwen/hip`: Qwen3.8-27B and the DFlash2
+  drafts) and the formats Q4_K, Q5_K, Q6_K and IQ4_XS.
+- The candidate set is the neighbourhood of shapes picked so far, not an
+  exhaustive search.
+- The current shapes were measured on gfx1151 under Linux with ROCm 10.1.
+  Shapes picked under Windows with TheRock ROCm 10.0.0 were up to 29% slower
+  than `generic` there (16-lane groups at six to eight rows), so a new
+  toolchain needs a new sweep.

tools/qwen27b/verify_sweep/sweep.hip

@@ -0,0 +1,275 @@
+// Launch-geometry sweep for the Qwen GPU verification projections.
+//
+// Loads one weight matrix from a GGUF, runs the production path
+// (TryLaunchVerifyGemm27, then LaunchBatchedQuantGEMMFp32), the generic
+// LaunchBatchedQuantGEMMFp32 dispatch alone and a set of
+// SmallBatchKQuantExactFp32GEMMKernel launch shapes on the same input, checks
+// every candidate bit for bit against production and times them with HIP
+// events: alternating rounds, median, weights rotated through >= 512 MiB of
+// copies so no run reads from the 32 MB Infinity Cache.
+//
+//   qwen27b_verify_sweep <gguf> <offset> <type> <m> <k> <batch> [rounds]
+//
+// <type> is the GGUF type id: 12 Q4_K, 13 Q5_K, 14 Q6_K, 23 IQ4_XS. Output:
+// one "base" line, a "listed" line (1 when verify_gemm_27b.hip lists the
+// shape) and one line per candidate ("generic" or Launch<> parameters, time,
+// GB/s, gain); see README.md.
+#include <hip/hip_runtime.h>
+
+#include <algorithm>
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <fstream>
+#include <functional>
+#include <random>
+#include <string>
+#include <vector>
+
+#include "src/models/qwen/hip/kernels/small_batch_gemm.hpp"
+#include "src/models/qwen/hip/kernels/verify_gemm_27b.hpp"
+#include "src/models/qwen/hip/ops/gemm.hpp"
+
+using gufo::core::GgmlType;
+
+#define CHECK(x)                                                               \
+  do {                                                                         \
+    hipError_t e_ = (x);                                                       \
+    if (e_ != hipSuccess) {                                                    \
+      std::fprintf(stderr, "HIP %s at %d\n", hipGetErrorString(e_), __LINE__); \
+      std::exit(2);                                                            \
+    }                                                                          \
+  } while (0)
+
+namespace {
+
+using Run = std::function<void(const void*, const float*, float*, std::size_t,
+                               std::size_t, hipStream_t)>;
+
+struct Candidate {
+  std::string
+      name;  // Launch<T, Batch, Waves, Rows, Tiles, Tokens, MinWaves, Lanes>
+  Run run;
+};
+
+template<GgmlType T, std::size_t Batch, std::uint32_t Waves, std::size_t Rows,
+         std::size_t Tiles, std::size_t TokensPerStep, std::uint32_t MinWaves,
+         std::uint32_t Lanes>
+Candidate Make() {
+  char name[96];
+  std::snprintf(name, sizeof(name), "Launch<T, %zu, %u, %zu, %zu, %zu, %u, %u>",
+                Batch, Waves, Rows, Tiles, TokensPerStep, MinWaves, Lanes);
+  return {
+      name, [](const void* w, const float* x, float* y, std::size_t m,
+               std::size_t k, hipStream_t s) {
+        constexpr std::size_t kRowsPerBlock = Waves * (32 / Lanes) * Rows;
+        gufo::hip::SmallBatchKQuantExactFp32GEMMKernel<Waves, Batch, Rows, T,
+                                                       Tiles, MinWaves, 1, true,
+                                                       TokensPerStep, 32, Lanes>
+            <<<static_cast<unsigned>((m + kRowsPerBlock - 1) / kRowsPerBlock),
+               Waves * 32, 0, s>>>(w, x, y, m, k);
+      }};
+}
+
+// The shapes picked in verify_gemm_27b.hip so far, plus their neighbours.
+template<GgmlType T, std::size_t B>
+std::vector<Candidate> Candidates() {
+  return {Make<T, B, 4, 4, 1, B, 8, 16>(),  Make<T, B, 4, 4, 1, B, 4, 16>(),
+          Make<T, B, 8, 4, 1, B, 8, 16>(),  Make<T, B, 8, 4, 1, B, 4, 16>(),
+          Make<T, B, 16, 4, 1, B, 4, 16>(), Make<T, B, 16, 4, 1, B, 8, 16>(),
+          Make<T, B, 4, 2, 1, B, 8, 16>(),  Make<T, B, 8, 2, 1, B, 8, 16>(),
+          Make<T, B, 4, 3, 1, B, 8, 16>(),  Make<T, B, 8, 3, 1, B, 4, 16>(),
+          Make<T, B, 8, 4, 1, 1, 4, 32>(),  Make<T, B, 16, 4, 1, 1, 4, 32>(),
+          Make<T, B, 8, 4, 2, B, 4, 32>(),  Make<T, B, 16, 4, 2, B, 4, 32>()};
+}
+
+template<GgmlType T>
+std::vector<Candidate> ForBatch(std::size_t batch) {
+  switch (batch) {
+    case 2:
+      return Candidates<T, 2>();
+    case 3:
+      return Candidates<T, 3>();
+    case 4:
+      return Candidates<T, 4>();
+    case 5:
+      return Candidates<T, 5>();
+    case 6:
+      return Candidates<T, 6>();
+    case 7:
+      return Candidates<T, 7>();
+    case 8:
+      return Candidates<T, 8>();
+    default:
+      return {};
+  }
+}
+
+}  // namespace
+
+int main(int argc, char** argv) {
+  if (argc < 7) {
+    std::fprintf(stderr,
+                 "usage: qwen27b_verify_sweep <gguf> <offset> <type> <m> <k> "
+                 "<batch> [rounds]\n");
+    return 1;
+  }
+  const std::string path = argv[1];
+  const std::uint64_t offset = std::strtoull(argv[2], nullptr, 10);
+  const int type_id = std::atoi(argv[3]);
+  const std::size_t m = std::strtoull(argv[4], nullptr, 10);
+  const std::size_t k = std::strtoull(argv[5], nullptr, 10);
+  const std::size_t batch = std::strtoull(argv[6], nullptr, 10);
+  const int rounds = argc > 7 ? std::atoi(argv[7]) : 9;
+  GgmlType type;
+  std::size_t block_bytes;  // per 256 weights
+  switch (type_id) {
+    case 12:
+      type = GgmlType::kQ4_K;
+      block_bytes = 144;
+      break;
+    case 13:
+      type = GgmlType::kQ5_K;
+      block_bytes = 176;
+      break;
+    case 14:
+      type = GgmlType::kQ6_K;
+      block_bytes = 210;
+      break;
+    case 23:
+      type = GgmlType::kIQ4_XS;
+      block_bytes = 136;
+      break;
+    default:
+      std::fprintf(stderr, "unsupported type %d\n", type_id);
+      return 1;
+  }
+  if (batch < 2 || batch > 8 || k % 256 != 0) {
+    std::fprintf(stderr, "batch must be 2-8 and k a multiple of 256\n");
+    return 1;
+  }
+  const std::size_t bytes = m * (k / 256) * block_bytes;
+
+  std::vector<char> host_w(bytes);
+  std::ifstream file(path, std::ios::binary);
+  file.seekg(static_cast<std::streamoff>(offset));
+  file.read(host_w.data(), static_cast<std::streamsize>(bytes));
+  if (!file) {
+    std::fprintf(stderr, "could not read %zu bytes at %llu\n", bytes,
+                 static_cast<unsigned long long>(offset));
+    return 1;
+  }
+  std::vector<float> host_x(batch * k);
+  std::mt19937 rng(42);
+  std::normal_distribution<float> dist(0.0F, 1.0F);
+  for (auto& v : host_x)
+    v = dist(rng);
+
+  const std::size_t copies =
+      std::max<std::size_t>(2, (512ull << 20) / bytes + 1);
+  std::vector<void*> weights(copies);
+  for (auto& p : weights) {
+    CHECK(hipMalloc(&p, bytes));
+    CHECK(hipMemcpy(p, host_w.data(), bytes, hipMemcpyHostToDevice));
+  }
+  std::size_t next = 0;
+  const auto rotate = [&] {
+    next = (next + 1) % copies;
+    return weights[next];
+  };
+  float *x = nullptr, *y = nullptr;
+  CHECK(hipMalloc(&x, host_x.size() * sizeof(float)));
+  CHECK(hipMalloc(&y, batch * m * sizeof(float)));
+  CHECK(hipMemcpy(x, host_x.data(), host_x.size() * sizeof(float),
+                  hipMemcpyHostToDevice));
+  hipStream_t s;
+  CHECK(hipStreamCreate(&s));
+
+  // The production path, as LaunchProjection takes it.
+  const auto production = [&](const void* w, float* out) {
+    if (!gufo::hip::TryLaunchVerifyGemm27(type, w, x, out, batch, m, k, s))
+      gufo::hip::LaunchBatchedQuantGEMMFp32(type, w, x, out, batch, m, k, s);
+  };
+  auto candidates = type == GgmlType::kQ4_K   ? ForBatch<GgmlType::kQ4_K>(batch)
+                    : type == GgmlType::kQ5_K ? ForBatch<GgmlType::kQ5_K>(batch)
+                    : type == GgmlType::kQ6_K
+                        ? ForBatch<GgmlType::kQ6_K>(batch)
+                        : ForBatch<GgmlType::kIQ4_XS>(batch);
+  // The dispatch without any listed shape; a listed shape slower than it is a
+  // regression.
+  candidates.insert(
+      candidates.begin(),
+      {"generic",
+       [type, batch](const void* w, const float* in, float* out,
+                     std::size_t rows, std::size_t cols, hipStream_t st) {
+         gufo::hip::LaunchBatchedQuantGEMMFp32(type, w, in, out, batch, rows,
+                                               cols, st);
+       }});
+
+  std::vector<float> reference(batch * m), got(batch * m);
+  const bool listed =
+      gufo::hip::TryLaunchVerifyGemm27(type, weights[0], x, y, batch, m, k, s);
+  if (!listed)
+    gufo::hip::LaunchBatchedQuantGEMMFp32(type, weights[0], x, y, batch, m, k,
+                                          s);
+  CHECK(hipStreamSynchronize(s));
+  CHECK(hipMemcpy(reference.data(), y, reference.size() * sizeof(float),
+                  hipMemcpyDeviceToHost));
+  std::vector<bool> exact(candidates.size());
+  for (std::size_t i = 0; i < candidates.size(); ++i) {
+    CHECK(hipMemset(y, 0, got.size() * sizeof(float)));
+    candidates[i].run(weights[0], x, y, m, k, s);
+    CHECK(hipStreamSynchronize(s));
+    CHECK(hipMemcpy(got.data(), y, got.size() * sizeof(float),
+                    hipMemcpyDeviceToHost));
+    exact[i] = std::memcmp(got.data(), reference.data(),
+                           got.size() * sizeof(float)) == 0;
+  }
+
+  hipEvent_t start, stop;
+  CHECK(hipEventCreate(&start));
+  CHECK(hipEventCreate(&stop));
+  constexpr int kWarmup = 3, kRepetitions = 20;
+  const auto time_us = [&](const std::function<void()>& launch) {
+    for (int i = 0; i < kWarmup; ++i)
+      launch();
+    CHECK(hipEventRecord(start, s));
+    for (int i = 0; i < kRepetitions; ++i)
+      launch();
+    CHECK(hipEventRecord(stop, s));
+    CHECK(hipEventSynchronize(stop));
+    float ms = 0;
+    CHECK(hipEventElapsedTime(&ms, start, stop));
+    return ms * 1000.0F / kRepetitions;
+  };
+  std::vector<std::vector<float>> samples(candidates.size() + 1);
+  for (int r = 0; r < rounds; ++r) {
+    samples[0].push_back(time_us([&] { production(rotate(), y); }));
+    for (std::size_t i = 0; i < candidates.size(); ++i)
+      if (exact[i])
+        samples[i + 1].push_back(
+            time_us([&] { candidates[i].run(rotate(), x, y, m, k, s); }));
+  }
+  const auto median = [](std::vector<float> v) {
+    std::sort(v.begin(), v.end());
+    return v.empty() ? 0.0F : v[v.size() / 2];
+  };
+  const float base = median(samples[0]);
+  std::printf(
+      "base type=%d m=%zu k=%zu batch=%zu MB=%.1f copies=%zu %.1f us %.0f "
+      "GB/s\n",
+      type_id, m, k, batch, bytes / 1e6, copies, base, bytes / (base * 1e3));
+  std::printf("listed %d\n", listed ? 1 : 0);
+  for (std::size_t i = 0; i < candidates.size(); ++i) {
+    if (!exact[i]) {
+      std::printf("cand %s NOT-EXACT\n", candidates[i].name.c_str());
+      continue;
+    }
+    const float t = median(samples[i + 1]);
+    std::printf("cand %s %.1f us %.0f GB/s %+.1f%%\n",
+                candidates[i].name.c_str(), t, bytes / (t * 1e3),
+                (base - t) / base * 100.0F);
+  }
+  return 0;
+}

tools/qwen27b/verify_sweep/sweep.py

@@ -0,0 +1,172 @@
+#!/usr/bin/env python3
+"""Sweep the Qwen verification launch geometries for one GGUF on this GPU.
+
+    python3 tools/qwen27b/verify_sweep/sweep.py --binary <qwen27b_verify_sweep> <model.gguf>
+
+Lists every distinct (type, m, k) projection of the transformer layers, runs
+qwen27b_verify_sweep on one example tensor of each at every requested row
+count, and prints the fastest exact candidate per projection and row count.
+Gains of at least --min-gain percent are also printed in the Launch<> syntax
+of src/models/qwen/hip/kernels/verify_gemm_27b.hip. Gate/up pairs of one
+format stored back to back are measured as the joined (2 * m) x k matrix the
+decode multiplies. The GPU must be idle for the timings to mean anything.
+"""
+from __future__ import annotations
+
+import argparse
+import collections
+from pathlib import Path
+import re
+import struct
+import subprocess
+import sys
+
+# GGUF type id -> bytes per 256 weights; the formats the sweep binary measures.
+BLOCK_BYTES = {12: 144, 13: 176, 14: 210, 23: 136}
+TYPE_NAMES = {12: "Q4_K", 13: "Q5_K", 14: "Q6_K", 23: "IQ4_XS"}
+CANDIDATE = re.compile(r"^cand (.+) [\d.]+ us (\d+) GB/s ([+-][\d.]+)%$")
+BASE = re.compile(r"^base .* [\d.]+ us (\d+) GB/s$")
+
+
+def read(f, fmt):
+    return struct.unpack("<" + fmt, f.read(struct.calcsize("<" + fmt)))[0]
+
+
+def read_string(f):
+    return f.read(read(f, "Q")).decode("utf-8", "replace")
+
+
+def read_value(f, kind):
+    scalars = {0: "B", 1: "b", 2: "H", 3: "h", 4: "I", 5: "i", 6: "f", 7: "?",
+               10: "Q", 11: "q", 12: "d"}
+    if kind in scalars:
+        return read(f, scalars[kind])
+    if kind == 8:
+        return read_string(f)
+    if kind == 9:
+        inner = read(f, "I")
+        for _ in range(read(f, "Q")):
+            read_value(f, inner)
+        return None
+    raise ValueError(f"unknown GGUF value type {kind}")
+
+
+def tensors(path):
+    """Return [(name, type, ne, absolute offset)] without reading weights."""
+    with open(path, "rb") as f:
+        if f.read(4) != b"GGUF":
+            sys.exit(f"{path}: not a GGUF file")
+        read(f, "I")
+        n_tensors, n_kv = read(f, "Q"), read(f, "Q")
+        alignment = 32
+        for _ in range(n_kv):
+            key = read_string(f)
+            value = read_value(f, read(f, "I"))
+            if key == "general.alignment":
+                alignment = value
+        infos = []
+        for _ in range(n_tensors):
+            name = read_string(f)
+            ne = [read(f, "Q") for _ in range(read(f, "I"))]
+            kind = read(f, "I")
+            infos.append((name, kind, ne, read(f, "Q")))
+        start = (f.tell() + alignment - 1) // alignment * alignment
+    return [(name, kind, ne, start + offset) for name, kind, ne, offset in infos]
+
+
+def projections(path):
+    """Return [(type, m, k, offset, layer_count, example_name)]."""
+    layers = collections.defaultdict(dict)
+    for name, kind, ne, offset in tensors(path):
+        parts = name.split(".")
+        if parts[0] != "blk" or len(ne) != 2 or kind not in BLOCK_BYTES:
+            continue
+        k, m = ne
+        layers[parts[1]][parts[2]] = (kind, m, k, offset, name)
+    found = collections.OrderedDict()
+    for layer in layers.values():
+        gate, up = layer.get("ffn_gate"), layer.get("ffn_up")
+        packed = False
+        if gate and up and gate[0] == up[0] and gate[1:3] == up[1:3]:
+            size = gate[1] * (gate[2] // 256) * BLOCK_BYTES[gate[0]]
+            packed = gate[3] + size == up[3]
+        for role, (kind, m, k, offset, name) in layer.items():
+            if packed and role in ("ffn_gate", "ffn_up"):
+                continue  # multiplied as the joined matrix below
+            found.setdefault((kind, m, k), [offset, 0, name])[1] += 1
+        if packed:
+            key = (gate[0], 2 * gate[1], gate[2])
+            found.setdefault(key, [gate[3], 0, "packed ffn_gate+ffn_up"])[1] += 1
+    return [(kind, m, k, offset, count, name)
+            for (kind, m, k), (offset, count, name) in found.items()]
+
+
+def main() -> int:
+    parser = argparse.ArgumentParser(description=__doc__.splitlines()[0])
+    parser.add_argument("model", type=Path)
+    parser.add_argument("--binary", type=Path, required=True,
+                        help="qwen27b_verify_sweep executable")
+    parser.add_argument("--batches", default="2,3,4,5,6,7,8")
+    parser.add_argument("--rounds", type=int, default=7)
+    parser.add_argument("--min-gain", type=float, default=5.0,
+                        help="percent over production to suggest a shape")
+    parser.add_argument("--min-layers", type=int, default=2,
+                        help="skip projections used by fewer layers; "
+                        "pass 1 for a DFlash2 draft")
+    parser.add_argument("--out", type=Path, help="write the raw sweep lines")
+    args = parser.parse_args()
+    batches = [int(b) for b in args.batches.split(",")]
+
+    log = []
+    rows = []
+    for kind, m, k, offset, count, name in projections(args.model):
+        if count < args.min_layers:
+            continue
+        for batch in batches:
+            lines = subprocess.run(
+                [str(args.binary), str(args.model), str(offset), str(kind),
+                 str(m), str(k), str(batch), str(args.rounds)],
+                check=True, capture_output=True, text=True).stdout.splitlines()
+            log += lines
+            base = next((int(match.group(1)) for line in lines
+                         if (match := BASE.match(line))), 0)
+            best = max(((match.group(1), int(match.group(2)),
+                         float(match.group(3))) for line in lines
+                        if (match := CANDIDATE.match(line))),
+                       key=lambda c: c[2], default=("-", 0, 0.0))
+            generic = next((float(match.group(3)) for line in lines
+                            if (match := CANDIDATE.match(line))
+                            and match.group(1) == "generic"), 0.0)
+            if "listed 1" not in lines:
+                generic = 0.0  # production is the generic dispatch
+            inexact = sum("NOT-EXACT" in line for line in lines)
+            rows.append((TYPE_NAMES[kind], m, k, count, name, batch, base,
+                         best[1], best[2], best[0], inexact, generic))
+            print(f"{TYPE_NAMES[kind]:6} {m:>6} {k:>6} layers={count:<3} "
+                  f"rows={batch} {base:>4} -> {best[1]:>4} GB/s "
+                  f"{best[2]:+6.1f}%  {best[0]}"
+                  + (f"  ({inexact} not exact)" if inexact else ""),
+                  flush=True)
+    if args.out:
+        args.out.write_text("\n".join(log) + "\n", encoding="utf-8")
+
+    print(f"\nGains of {args.min_gain:g}% or more, as verify_gemm_27b.hip "
+          "entries (check the shape is not also a draft shape there):")
+    for (type_name, m, k, count, name, batch, base, best_gbs, gain, launch,
+         _, _) in rows:
+        if gain >= args.min_gain and launch != "generic":
+            print(f"  // {type_name} {m}x{k} ({count} layers, {name}), "
+                  f"{batch} rows: {base} -> {best_gbs} GB/s, +{gain:g}%")
+            print(f"  {launch}(w, x, y, m, k, stream);")
+    slower = [row for row in rows if row[11] > 0]
+    if slower:
+        print("\nListed shapes slower than the generic dispatch; remove them "
+              "or replace them with a shape that wins by --min-gain:")
+        for (type_name, m, k, count, name, batch, *_, generic) in slower:
+            print(f"  {type_name} {m}x{k} ({count} layers, {name}), "
+                  f"{batch} rows: generic {generic:+g}%")
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())