HIP: on-device autotuning for gfx1100 decode (kernel shapes, draft and CPU settings)

CMakeLists.txt

@@ -250,6 +250,11 @@ if(STRATA_BUILD_TESTS)
   add_executable(tuning_header_arch_family_test tests/prefill/tuning_header_arch_family_test.cpp)
   target_include_directories(tuning_header_arch_family_test PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/src/prefill")
   add_test(NAME tuning_header_arch_family_test COMMAND tuning_header_arch_family_test)
+
+  # the decode tuning table's parser and its refusals (plain C++: no GPU)
+  add_executable(decode_tuning_test tests/core/decode_tuning_test.cpp src/kernels/decode_tuning.cpp)
+  target_include_directories(decode_tuning_test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include)
+  add_test(NAME decode_tuning_test COMMAND decode_tuning_test)
 endif()
 
 # ---- strata_plan ----------------------------------------------------------
@@ -349,6 +354,7 @@ if(STRATA_ENABLE_CUDA OR STRATA_ENABLE_HIP)
                           src/kernels/cuda/qsa_select.cu
                           src/kernels/cuda/verify_kernels.cu
                           src/kernels/cuda/iq_kernels.cu
+                          src/kernels/decode_tuning.cpp
                           src/kernels/ngram.cpp
                           src/ngram/ple_reader.cpp
                           src/platform/direct_file.cpp)
@@ -528,6 +534,11 @@ if(STRATA_ENABLE_CUDA OR STRATA_ENABLE_HIP)
     target_link_libraries(tune_hipblaslt PRIVATE hip::host roc::hipblas roc::hipblaslt)
   endif()
 
+  if(STRATA_ENABLE_HIP)
+    # on-device tuning of the decode kernels' block shapes (docs/AMD_HIP_AUTOTUNE.md; tools/hip/autotune.py runs it)
+    add_executable(tune_decode tools/hip/tune_decode.cpp)
+    target_link_libraries(tune_decode PRIVATE strata_kernels strata_hip_runtime)
+  endif()
   add_executable(strata src/program/generate.cpp)
   target_include_directories(strata PRIVATE ${_strata_gpu_include_directories})
   target_link_libraries(strata PRIVATE strata_engine strata_prefill strata_spec ${_strata_gpu_runtime_target})

docs/AMD_HIP.md

@@ -50,7 +50,8 @@ the kernel's amdgpu driver (no ROCm install needed):
 - **Limits for now:** images only through the CPU encoder (`--vision cpu`, 0.1.32). Setup does not offer the tuning
   (calibration) on AMD yet: its controls are being checked on HIP one at a time (#566). Since 0.1.39 a tuning run by
   hand (`./setup.sh --calibrate`) is saved for the AMD card it ran on and reused when setup runs again. The Monitor
-  shows the card's load, VRAM, temperature and power from Linux sysfs (0.1.32).
+  shows the card's load, VRAM, temperature and power from Linux sysfs (0.1.32). On-device tuning of the
+decode kernels (block shapes, draft and CPU settings) is available by hand: [AMD_HIP_AUTOTUNE.md](AMD_HIP_AUTOTUNE.md).
 
 The rest of setup is the same as on NVIDIA: the model download, the start script, the server.
 

docs/AMD_HIP_AUTOTUNE.md

@@ -0,0 +1,117 @@
+# On-device autotuning for the RX 7900 XTX (gfx1100)
+
+Strata's AMD backend runs CUDA-shaped kernels whose launch shapes were chosen on an RTX 5070. This adds
+two measured layers of tuning for the 7900 XTX, in the spirit of engines that tune kernels on the user's
+device instead of shipping one shape for a whole class of hardware. Both layers keep answers exactly the
+same.
+
+| Layer | What is measured | Changes answers? | Cost |
+|---|---|---|---|
+| Kernels | Rows per thread block of the decode pass's GPU expert kernels (gate/up, down) and `iq_mmvq` | No: bitwise-equal variants only, checked on the machine and again end to end | ~1–3 min + 4 engine starts |
+| Engine settings | Draft floor (`--spec-min-p`), CPU expert workers (`--pool-workers`) | No (greedy verification keeps the tokens) | 5–10 min |
+| Drafting (opt-in) | MTP window (`--spec`), prompt-lookup match (`--suffix-draft`) | No | +6 engine starts |
+
+## Run it
+
+Stop the server first: the timings need the GPU idle and about 1 GB of free VRAM. Then:
+
+```
+python tools/hip/autotune.py strata-iq3_xxs.json                 # kernels only
+python tools/hip/autotune.py strata-iq3_xxs.json --settings      # + draft floor and CPU workers
+python tools/hip/autotune.py strata-iq3_xxs.json --draft --prompts my-agent-prompts.json
+python tools/hip/autotune.py strata-iq3_xxs.json --off           # back to the default kernel shapes
+```
+
+`--prompts` takes a JSON list of strings, or plain text with prompts separated by a line `---`. Use requests your
+agent really sends. Prompt lookup only pays where the answer repeats the context (edits, refactors, quoting),
+so the built-in prompts are a weak proxy for an agent workload.
+
+The script needs the `tune_decode` tool. It builds it in `build-hip` when that build directory exists, or you
+can build it yourself:
+
+```
+cmake --build build-hip --target tune_decode -j2
+```
+
+When it finishes, the config's `env` carries `STRATA_DECODE_TUNING` (next to `STRATA_HIPBLASLT_TUNING`), and the
+engine log says `decode tuning: N shapes from ...` at start. A report goes to `strata-<model>.autotune.json`, and the
+original config is kept once as `strata-<model>.json.bak-autotune`.
+
+The settings (`--settings`, `--draft`) also work through setup: `./setup.sh --calibrate`, `--calibrate-draft`,
+`--calibrate-prompts FILE`. On AMD the PCIe share is left exactly as the config has it, because the HIP expert
+paths (mmap, resident CPU experts, the staging buffer) were measured with that share. `tools/calibrate.py
+--sweep-pcie` sweeps it anyway.
+
+## How the kernel layer stays exact
+
+`native_expert_grouped` (the GPU experts in every verify window) and `iq_mmvq` give one warp to each weight row.
+How many rows share a block changes occupancy and scheduling but never the arithmetic: each row is still reduced
+by one warp, over the same lanes, in the same order. `tune_decode` checks this anyway:
+
+1. Every candidate's output is compared **byte for byte** with the engine default's - the real `native_expert_grouped`
+   / `iq_mmvq` call, no table, with its own launch heuristics - before it may be timed. A difference excludes the
+   candidate, fails the run (exit 4), and nothing is written.
+2. The workload is the verify pass's: windows of 1, 3 and 5 tokens (weighted 0.2/0.3/0.5), `--hit-frac` of each
+   token's 10 experts on the GPU, and experts drawn from a 768 MB pool. The pool is far larger than the 96 MB
+   Infinity Cache, because decode streams around a gigabyte of experts per token and a cache-warm benchmark would lie.
+3. Candidates are timed interleaved over several rounds, in a rotating order. A winner must beat the default by more
+   than 3% on the weighted median, then again in a second interleaved confirmation.
+4. `autotune.py` then starts the engine without and with the table, twice each, interleaved. It keeps the table
+   only if the generated tokens are identical and decode is not slower. If two runs *without* the table already
+   differ (an adaptive expert tier, for example), the token check is reported as not checkable and speed decides.
+
+## What a table row means
+
+A row names the rows per block of the **plain one-warp-per-row kernels** (`native_gu_kernel`, `native_down_kernel`,
+`mmvq_kernel`). The no-table default for split formats instead runs the multi decode-once kernels, whose rows per
+block is their own - so `16` in a table and `16` in the default's launch geometry are different mechanisms, and the
+tuner compares whole launches against the engine's default entry, not numbers against numbers. Two consequences:
+
+- A table row replaces the whole call. Kernels the table leaves at their default (no row, or the default value) run
+  the default dispatch chain, at the grid the parameter sizes - not necessarily the grid the no-table path would
+  have chosen. Outputs stay bitwise equal; only the block count can differ.
+- The `mmvq` key is `(type, n_in, n_out)` with no token-count axis: the winner is a blend over the 1/3/5-token
+  verify windows, which is the decode workload.
+
+The table is refused at engine start, with a message, unless its header matches the running engine: the GPU
+architecture, the HIP runtime version, a hash of the compiler that built the kernels, and the kernel selection
+space (a version that bumps when the kernel set or the default dispatch changes - a table from an older engine is
+stale and the tuner is run again). It is also refused when a dispatch rerouting variable is set
+(`STRATA_OLD_IQ_MMVQ`, `STRATA_NO_SUB16_GU`, `STRATA_IQ_STAGE_GRID=0`, `STRATA_IQ_STAGE_GRID_MMVQ=1`,
+`STRATA_EXPERT_V2`, `STRATA_EXPERT_V2K`, `STRATA_TSUM`, `STRATA_GROUPED_V1`): the table was measured against the
+default dispatch, and a knob that reroutes it makes those measurements meaningless. CUDA builds compile only the
+default shapes and ignore the variable.
+
+A note on the staging knobs, measured on the RX 7900 XTX: the defaults stage the i-quant codebook grids into
+shared memory, which upstream chose on an RTX 5070. On this card the codebook already sits in the 96 MB Infinity
+Cache, so staging is a wash for the IQ3-class formats (within 0.2% on every shape measured) and a small win when
+OFF for IQ2_S gate/up (~2-5% with `STRATA_IQ_STAGE_GRID=0`); mmvq staging is a wash. Measured by the same tool,
+byte-checked; if you run an IQ2_S-class model it is worth one paired run on your machine.
+
+## What to expect
+
+Honestly: unknown until it runs on your card. The tunable kernels are a slice of each token's time. In the
+hybrid setup, CPU-side expert work and the GPU-reach wait are large too, and kernel shapes cannot touch those.
+Measure where your time goes before and after:
+
+```
+"env": {"STRATA_DECODE_TIMING": "1", ...}
+```
+
+With this set, each request logs one line: windows, average window size, tokens per window, and ms per window
+split into verify (GPU-reach wait, per-layer host and CPU expert time), commit and draft. It also logs VRAM hits,
+CPU expert entries and PCIe experts per layer-window, followed by the GPU stage table. If "verify" is dominated by
+the CPU expert part, the drafting and worker settings matter more than kernel shapes.
+
+## Drafting notes for Qwen3.8-Flash-Next
+
+- **MTP** is the model's own draft layer. Its window grows while the draft's confidence is at least
+  `--spec-min-p`. Every extra row in a window brings more experts the VRAM cache misses onto the CPU. On a
+  CPU-bound machine a higher floor or a shorter `--spec` often wins, so it is measured, not assumed.
+- **Prompt lookup** (`--suffix-draft`) drafts the continuation of an earlier repeat of the context. A built-in
+  policy learns the measured round cost per window size and lookup acceptance per match length, and takes a lookup
+  window only when its expected tokens per millisecond beat the MTP's. It fires only when its first token agrees
+  with the MTP's first guess.
+- **DFlash2** (block-diffusion drafters) exists for Qwen3.8-27B, not for Flash-Next. A drafter for Flash-Next would
+  have to be trained, and it would take VRAM from the expert cache (about 700 experts per GB). This is not part of
+  this tuning.

include/strata/kernels/decode_tuning.hpp

@@ -0,0 +1,111 @@
+// include/strata/kernels/decode_tuning.hpp - on-device tuning of the decode kernels' launch shapes (gfx1100).
+//
+// The decode hot path's GPU experts (`native_expert_grouped`: gate/up, then down) and the i-quant matrix-vector
+// kernel (`iq_mmvq`) give one warp to each weight row; how many rows share a thread block is a free choice.  It
+// changes occupancy and scheduling but not one floating-point operation: every row is still reduced by the same
+// warp, over the same lanes, in the same order, so every variant writes bitwise the same output as the default.
+// Which one is fastest depends on the GPU, the driver and the compiler, so it is measured on the machine
+// (tools/hip/tune_decode.cpp writes the table, tools/hip/autotune.py runs the whole procedure) instead of guessed.
+//
+// The table (STRATA_DECODE_TUNING=<path>) is plain text:
+//
+//   # comments
+//   STRATA_DECODE_TUNING_V1 <gpu arch> <HIP runtime version> <toolchain id> <selection space>
+//   gu   <ggml type> <n_in> <n_out> <rows per block>      gate/up of the grouped experts (n_in = n_embd, n_out = n_ff)
+//   down <ggml type> <n_in> <n_out> <rows per block>      down of the grouped experts   (n_in = n_ff,   n_out = n_embd)
+//   mmvq <ggml type> <n_in> <n_out> <rows per block>      iq_mmvq
+//
+// The header must match the running engine exactly - the architecture (e.g. gfx1100), the HIP runtime, the
+// compiler that built the kernels (the tuner checked bitwise equality with THAT build), and the kernel selection
+// space (below) - or the table is refused with a message and the defaults stay.  A shape the table does not list
+// keeps its default.  CUDA builds compile only the default shapes and ignore the variable (with a message), so
+// their behaviour is unchanged.
+//
+// What a row names: the rows per block of the PLAIN one-warp-per-row kernels (mmvq_kernel / native_gu_kernel /
+// native_down_kernel).  The no-table default for split types instead runs the multi decode-once kernels, whose
+// rows per block is their own - so the number 16 in a table and the number 16 in the default's launch geometry
+// are different mechanisms.  The tuner compares whole configurations against the engine's true default entry
+// (native_expert_grouped / iq_mmvq with no table), byte-checks every candidate against it first, and a table row
+// replaces the whole call's launch: the kernels the table leaves at their default run the default chain.
+#pragma once
+
+#include <cstdint>
+#include <istream>
+#include <string>
+#include <vector>
+
+namespace strata::kernels {
+
+enum class DecodeKernel : int { GateUp = 0, Down = 1, Mmvq = 2 };
+
+/// Defaults (the layouts before tuning existed): 8 rows per block for the grouped experts, 4 for iq_mmvq.
+inline constexpr int kDefaultExpertRows = 8;
+inline constexpr int kDefaultMmvqRows = 4;
+
+/// Which kernel selection space this build offers the tuner.  The value rides in the table header; a build whose
+/// space differs refuses the table ("run the tuner again").  Bump it when the plain kernels' arithmetic or the
+/// default dispatch chain changes enough that an older table no longer describes this build - a compiler or
+/// runtime bump alone does not change it, a kernel rework does (the sub-warp/EXACT_N rework is the case that
+/// motivated the field: same compiler, same runtime, different kernels).
+inline constexpr int kDecodeTuningSpace = 2;
+
+/// The rows-per-block variants HIP builds compile besides the default (one list: the parser's validity set, the
+/// launchers' compiled-in set and the tuner's candidate list all derive from these two macros).  CUDA builds
+/// compile only the defaults.
+#define STRATA_EXPERT_ROWS_VARIANTS(X) X(4) X(16) X(2)
+#define STRATA_MMVQ_ROWS_VARIANTS(X)  X(2) X(8) X(1)
+
+/// The rows-per-block values a kernel is compiled for (one warp of 32 lanes per row, at most 512 threads).
+bool decode_rows_valid(DecodeKernel k, int rows) noexcept;
+/// The candidate list the tuner measures for a kernel, default first.
+std::vector<int> decode_rows_candidates(DecodeKernel k);
+const char* decode_kernel_name(DecodeKernel k) noexcept;
+int decode_default_rows(DecodeKernel k) noexcept;
+
+struct DecodeTuningRow {
+    DecodeKernel kernel = DecodeKernel::GateUp;
+    int type = 0;
+    int64_t n_in = 0, n_out = 0;
+    int rows = 0;
+};
+
+/// Who the kernels were built and run by.  A table is only valid for exactly this identity.
+struct DecodeIdentity {
+    std::string arch;          ///< e.g. "gfx1100" (feature suffixes such as ":xnack-" stripped); "cuda" on CUDA
+    long long runtime = 0;     ///< hipRuntimeGetVersion / cudaRuntimeGetVersion
+    std::string toolchain;     ///< 16 hex digits: a hash of the device compiler's version string
+    int space = 0;             ///< kDecodeTuningSpace of the build the table was measured on
+};
+
+class DecodeTuningTable {
+public:
+    /// Parse a table; refuses (returns false, `err` says why) a malformed line, an unknown kernel, a rows value
+    /// the kernel is not compiled for, a duplicate shape, or a header that does not match `want`.
+    bool parse(std::istream& in, const DecodeIdentity& want, std::string& err);
+    bool load(const std::string& path, const DecodeIdentity& want, std::string& err);
+    /// Rows per block for this shape, or `fallback` when the table does not list it.
+    int rows(DecodeKernel k, int type, int64_t n_in, int64_t n_out, int fallback) const noexcept;
+    const std::vector<DecodeTuningRow>& entries() const noexcept { return rows_; }
+    bool empty() const noexcept { return rows_.empty(); }
+
+    static std::string header(const DecodeIdentity& id);
+    static std::string line(const DecodeTuningRow& r);
+
+private:
+    std::vector<DecodeTuningRow> rows_;
+};
+
+/// A stable 16-hex-digit FNV-1a hash of `s` (the toolchain id).
+std::string decode_toolchain_hash(const std::string& s);
+
+// ---- the running engine (implemented in the GPU translation unit) ----
+
+/// This process's identity: the current device's architecture, the runtime version and the kernels' compiler.
+DecodeIdentity decode_identity();
+/// Load STRATA_DECODE_TUNING once (later calls are no-ops) and say on stderr what was used.  Call it before any
+/// graph capture; captured graphs keep the kernels they captured.  Safe to skip: the first launch loads it lazily.
+void decode_tuning_init();
+/// The active table, or nullptr (none set, refused, or a CUDA build).
+const DecodeTuningTable* decode_tuning_active();
+
+}  // namespace strata::kernels

include/strata/kernels/iq_kernels.hpp

@@ -23,7 +23,12 @@ size_t iq_row_bytes(int ggml_type, int64_t n) noexcept;
 void quantize_q8_1_rows(const float* x, int64_t n_rows, int64_t n_cols, void* y, void* stream);
 
 /// y[c][r] = W[r] . x[c] for `ncols` columns of q8_1 activations (x stride n_in/32 blocks per column).
+/// Rows per thread block: the decode tuning table's choice for this shape (decode_tuning.hpp), else 4.
 void iq_mmvq(int ggml_type, const void* w, const void* x_q8_1, float* y, int n_in, int n_out, int ncols, void* stream);
+/// The same with an explicit rows-per-block (the tuner's entry; bitwise equal to iq_mmvq for every value
+/// `decode_rows_valid(DecodeKernel::Mmvq, rows)` accepts; CUDA builds compile only 4).
+void iq_mmvq_rows(int ggml_type, const void* w, const void* x_q8_1, float* y, int n_in, int n_out, int ncols, int rows,
+                  void* stream);
 
 /// Dequantize `n` contiguous values (n a multiple of 256) to fp16 / fp32.
 void iq_dequant_f16(int ggml_type, const void* src, int64_t n, uint16_t* dst, void* stream);
@@ -76,5 +81,12 @@ void native_expert_set_mode(int mode, int phase);
 /// bitwise the same results.  Set before graph capture; captured graphs keep the kernels they captured.
 void iq_set_old_kernels(bool old);
 bool iq_old_kernels();
+/// The same with explicit rows per block for the gate/up and the down kernel (the tuner's entry; bitwise equal to
+/// native_expert_grouped for every value decode_rows_valid accepts; CUDA builds compile only 8).  The default entry
+/// above reads the decode tuning table (decode_tuning.hpp) and routes here when it names a non-default rows value.
+void native_expert_grouped_rows(const NativeExpertLayout& L, const unsigned long long* grp_ptr,
+                                const int32_t* grp_start, const int32_t* n_groups, const int32_t* ent_dst,
+                                const int32_t* ent_tok, int64_t cap_groups, int64_t cap_entries, const void* x_q8_1,
+                                void* scratch, float* out, int gu_rows, int down_rows, void* stream);
 
 }  // namespace strata::kernels