#766 hip: a gfx1100 hipBLASLt tuning table for 100401 (packaged ROCm 10.0.0)

本文

## What

`tools/hip/gfx1100-hipblaslt-100401.txt`, plus the shipped-tables list in `docs/AMD_HIP.md`. No code change.

## Why

hipBLASLt solution ids are scoped to one version, and ROCm 10.0.0 ships **hipBLASLt 1.4.1** = version **100401**.
The repo ships gfx1100 tables for `100100` and `100200` (the 7.10 TheRock wheels) and a gfx1201 table for
`100500`, so on a packaged ROCm 10 an RX 7900 XTX gets this at setup:

```
[!] no hipBLASLt tuning table for gfx1100 with hipBLASLt 100401 (have: gfx1100-hipblaslt-100100.txt,
    gfx1100-hipblaslt-100200.txt): the prompt's dense matrix products use plain hipBLAS
```

PR #745 documents the same confound on ROCm 10.2 (their `100500`, +82% prefill at 128K once a table exists).
This is the `100401` / packaged-ROCm-10.0.0 case, for the card most people here own.

## How it was made

The documented recipe (`docs/AMD_HIP.md`, Tuning table) — the shipped gfx1100 table's 26 dense geometries,
re-tuned against the installed library:

```sh
CASES=$(awk 'NR>2 {printf " --case %s,%s,%s,%s,%s", $1, $5, $2, $3, $4}' tools/hip/gfx1100-hipblaslt-100200.txt)
./build-hip/tune_hipblaslt $CASES --tuning-out gfx1100-hipblaslt-100401.txt
```

Measured on an RX 7900 XTX (gfx1100), Ryzen 7 9800X3D, inside `rocm/dev-ubuntu-24.04:10.0.0-full`, engine
0.1.38 (`99f3dbd`), Swift 1.5 IQ3_XXS, 262144 ctx, `--kv int8`, greedy. Fresh prompts (a nonce per request —
the server checkpoints prompts, so the same text again reads nothing), medians of 3:

| fresh prompt | plain hipBLAS | with this table | |
|---|---:|---:|---:|
| 370 tok | 319 tok/s | 371 tok/s | +16% |
| 1,670 tok | 682 tok/s | **1,038 tok/s** | **+52%** |
| 6,471 tok | 857 tok/s | **1,524 tok/s** | **+78%** |

Decode is unchanged (79.0 vs 80.3 tok/s warm, inside this machine's measured 2.3% run-to-run drift) — the
table buys prompt speed, not generation speed, which matches #745.

Gate: `hip_prefill_hipblaslt_gemm` with `STRATA_HIPBLASLT_VERBOSE=1` reports

```
prefill gemm: hipBLASLt summary launches=4 fallbacks=0 unique_fallback_shapes=0
```

so the library honours the ids rather than the engine quietly falling back per-id (it still passes the smoke
test when it does, which is why the verbose run is the real check).

## Notes for review

- The second line of the file stamps `STRATA_HIPBLASLT_TUNING_V1 gfx1100 100401`, so setup picks it up only
  when the installed hipBLASLt reports 1.4.1, exactly as with the other tables.
- Per `docs/AMD_HIP.md`'s own caveat for `100500`: the version number is all the engine can check, so a
  *different* 1.4.1 build could number its solutions differently. This one was tuned against
  `libamdhip64.so.7.15.26333` / `libhipblaslt.so.1.4` as shipped in `rocm/dev-ubuntu-24.04:10.0.0-full`.
- Measured inside a container, but nothing about the measurement depends on that: the engine, the table and
  the tuner are the same files either way.

Not affiliated with the benchmark in #745; independent measurement on different hardware (9800X3D, IQ3_XXS,
packaged ROCm rather than a nightly SDK).