#339 hip: hipBLASLt tuning table for gfx1201 (R9700) with ROCm 10.2.0a nightly (hipBLASLt 1.5.0)

本文

## What

Commits: (1) Adds `tools/hip/gfx1201-hipblaslt-100500.txt`, a hipBLASLt tuning table for the **Radeon AI PRO R9700 (gfx1201, 32 GB)** with **ROCm 10.2.0a20260914** (AMD's `gfx120X-all` nightly, hipBLASLt **1.5.0**, library build `d3164197`), plus a short "Shipped tables" list in `docs/AMD_HIP.md`. (2) A 4-line fix to how the tuner and the parity test encode BF16 fixtures (see Caveats; the table's BF16 rows depend on it). (3) Review follow-up: the stale "no gfx1201 table" note in `docs/AMD_HIP.md` is updated, and the test's comments and messages say it is a smoke test over two rows (behaviour and exit codes unchanged).

The table is for the engine's *prompt* GEMMs (dense bf16/f16). Setup chooses `tools/hip/<arch>-hipblaslt-<version>.txt` by the installed hipBLASLt version, and the engine refuses any other version, so it does not interfere with a table for another version. (#262's `gfx1201-hipblaslt-100401.txt`, hipBLASLt 1.4.1, was not merged when #262 closed; I use it below only as the example of a table that 1.5.0 rejects.)

## Why it is needed

With hipBLASLt 1.5.0, #262's table is rejected and the engine falls back to hipBLASEx (`prefill gemm: hipBLASLt version mismatch: file=100401 runtime=100500; using hipBLASEx` in the engine log; the parity test prints `tuning table rejected` and exits 1). Solution ids are not stable across versions: on the 24 shapes both tables cover, only 13 of the 24 chosen ids are the same.

## Contents

- 32 rows: 16 geometries x T=4096 / T=8192 (18 bf16 rows, 14 f16 rows). The 24 shapes of #262's table, plus 8 more, all bf16: N=1/K=2560 at T=8192, N=4/K=10240 at both T, N=512/K=2560 and N=2560/K=2560 at T=4096, and N=10240 with K=2560 and K=320 at both T.
- Calibrated with `tune_hipblaslt` as shipped (16 heuristic candidates per row, 32 MiB workspace, two warmups, three timing repetitions; a candidate must be finite, relative L2 <= 1e-4, max abs <= 1e-2, with no padding writes).

## Evidence

On a clean checkout of `main` (`30ec18e`) plus this PR, built for gfx1201 with ROCm 10.2.0a (re-run on the head `8a4d331`, which merges v0.1.31 `9259cad`: the same 4/4 PASS, `fallbacks=0`, and the same rejection of the 100401 table):

- `STRATA_HIPBLASLT_TUNING=tools/hip/gfx1201-hipblaslt-100500.txt hip_prefill_hipblaslt_gemm`: `tuning enabled (32 rows, gfx1201, version 100500)`, 4/4 cases PASS, exit 0 (relative L2 <= 1.3e-6 for f16, 3.6e-7 for bf16 at T=37; the T=4096 bf16 case is exactly 0).
- Same binary with #262's 100401 table: `tuning table rejected: hipBLASLt version mismatch: file=100401 runtime=100500`, exit 1.
- `setup.hipblaslt_table('gfx1201', ['/opt/rocm/lib'])` returns the new file.
- In full-model use (Swift 1.5 IQ3_XXS serving, `--prefill auto`): the engine log shows `tuning enabled (32 rows, gfx1201, version 100500)` on every start since, and a calibration run recorded 16,480 hipBLASLt launches with zero fallbacks across all 16 geometries.

## What it is worth

Small: on fresh 4,210-token prompts, 1,590 tok/s tuned vs 1,531 untuned (+3.9%), on one card. It does not change decode. I would not describe it as more than a modest prompt-GEMM gain, and I did not re-measure it on longer prompts, where other phases dominate.

## Caveats

- The version number (`100500`) is all the engine can check. A different 1.5.0 build (another nightly) could number its solutions differently. `hip_prefill_hipblaslt_gemm` is a smoke test, not a validator: it refuses a table for another architecture or version and compares two of the 32 rows with hipBLASEx. A solution id the library rejects makes the engine fall back to hipBLASEx and the test still passes (I checked: with those two rows set to an invalid id the test exits 0 and `STRATA_HIPBLASLT_VERBOSE=1` prints `launches=0 fallbacks=4`; with the shipped table it prints `launches=4 fallbacks=0`). The docs say to read that summary line and to recalibrate with `tune_hipblaslt` for another 1.5.0 build.
- One GPU (R9700), one ROCm build. Not tested on gfx1200 or on other gfx1201 cards.
- The table was calibrated with the corrected BF16 host encoding (second commit below). I did not inspect the stock encoder's bit patterns directly; what I observed is the parity test's output: on stock `main` the T=37 beta=1 BF16 case prints `relative_l2=0 max_abs=0`, with the fix it prints `relative_l2=3.6e-07 max_abs=2.6e-06` (the T=4096 BF16 case prints 0 either way, and the f16 cases are unchanged). So the stock BF16 inputs differ from the intended ones; the exact fault is inferred. The commit changes only the two helper files (`tests/hip/prefill_hipblaslt_gemm.cpp`, `tools/hip/tune_hipblaslt.cpp`, 4 lines), not the engine and not the f16 path; it is in this PR because anyone regenerating a table with the stock tuner would calibrate the BF16 rows on those inputs.