Pull requests / #766
#766 hip: a gfx1100 hipBLASLt tuning table for 100401 (packaged ROCm 10.0.0)
closed · @miskahm · 0 comments · View on GitHub
BenchmarksSetup & installAMD / HIPModels & quantsDocumentation
Description
## 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).Related on strata.com
Editorial links to help you install, pick models, or read release notes — not part of the upstream thread.