Pull requests / #565

#565 tools/hip: add gfx1100 hipBLASLt tuning table for hipBLASLt 1.2.2 (100202)

closed · @lawsirlawsir-png · 0 comments · View on GitHub

BenchmarksSetup & installAMD / HIPDocumentation

Description

## What

Adds one file: `tools/hip/gfx1100-hipblaslt-100202.txt` — 26 rows, calibrated with `tools/hip/tune_hipblaslt.cpp` on a Radeon PRO **W7800 (gfx1100)** under **ROCm 7.2.4**.

## Why

`setup.py:hipblaslt_table()` resolves `tools/hip/<arch>-hipblaslt-<version>.txt`, where the version comes from `hipblaslt-version.h` (`MAJOR*100000 + MINOR*100 + PATCH`). ROCm 7.2.4 ships hipBLASLt **1.2.2 → 100202** (verified on this box: `HIPBLASLT_VERSION_MAJOR 1 / MINOR 2 / PATCH 2`), and `tools/hip/` on `main` carries only `100100` and `100200` for gfx1100. With no matching file, setup warns and the dense prefill GEMMs fall back to plain hipBLAS.

Measured on this machine on **engine 0.1.27** with `tools/hip/bench_prefill.py` (4 fresh prompts, 128 max_tokens, temp 0, seed 42, `--prefill auto --kv int8 --mtp mtp/rt --expert-cache auto`, thinking off):

| | prefill |
|---|---|
| no matching table (hipBLASEx fallback) | 622 tok/s |
| with this table | **1,088 tok/s** (+75%) |

## What the table actually is

- Header: `STRATA_HIPBLASLT_TUNING_V1 gfx1100 100202`.
- The 26 `(type, n, k, ldy, t_bucket)` shapes are **identical** to the shipped `gfx1100-hipblaslt-100200.txt` (mechanical diff of columns 1–5).
- 24 of 26 solution ids are the same as the 100200 table; **2 differ**:
  - `bf16 1 2560 1 4096`: 541 → **540**
  - `bf16 4 10240 4 8192`: 542 → **541**

So the gain comes from the version gate accepting a table at all, not from a wholesale re-tune. That is also why the file is worth shipping rather than leaving to each user: the fallback penalty is the whole 75%, not the two rows.

## How it was verified

1. Tuner run (verbatim):
   ```sh
   /opt/rocm-7.2.4/llvm/bin/clang++ -x hip -O2 -std=c++20 --offload-arch=gfx1100 \
     --rocm-path=/opt/rocm-7.2.4 -I/opt/rocm-7.2.4/include tools/hip/tune_hipblaslt.cpp \
     -L/opt/rocm-7.2.4/lib -lhipblaslt -lhipblas -lamdhip64 \
     -Wl,-rpath,/opt/rocm-7.2.4/lib -o /tmp/tune_hipblaslt
   ARGS=$(awk 'NR>2 {printf "--case %s,%s,%s,%s,%s ", $1,$5,$2,$3,$4}' tools/hip/gfx1100-hipblaslt-100200.txt)
   /tmp/tune_hipblaslt --workspace-mib 32 $ARGS --tuning-out tools/hip/gfx1100-hipblaslt-100202.txt
   ```
2. `TuningTable::load()` accepts it (arch and version match the runtime pair) — the engine logs `prefill gemm: hipBLASLt tuning enabled (26 rows, gfx1100, version 100202)`.
3. `setup.py` picks it up by name alone; no `STRATA_HIPBLASLT_TUNING` needed once the file is in `tools/hip/`.
4. End-to-end on the production server (engine 0.1.34): `/health` loaded, prefill 1,043–1,231 tok/s on fresh prompts, decode 62.6–67.0 tok/s.

File integrity: `sha256 a594869223200c62019d3c997345e89664d5126fcbd089a44f4f95a267ad2ff3`, 845 bytes.

## Scope and limits (so the table is not over-read)

- Calibrated on **one** gfx1100 board (W7800 48 GB). Solution ids are scoped to architecture + hipBLASLt version, not to board model; a 7900 XTX on hipBLASLt 100100 is unaffected by this file.
- Anyone on a different hipBLASLt build should re-run the tuner; the naming rule `tools/hip/<arch>-hipblaslt-<version>.txt` is what makes setup find it.
- The 622 → 1,088 figure is this machine's measurement under the `bench_prefill.py` conditions, not a cross-card claim.

## Two small notes

- Requested in #505 (the closing comment asked for exactly this table). PR #170 previously offered a 100202 table; `tools/hip/` on `main` still has only 100100/100200 as of today.
- The closing comment points to `tools/hip/README` for the tuner; there is no README in `tools/hip/` on `main` — the tuner documentation is in `docs/AMD_HIP.md` ("Tuning table").

Related on strata.com

Editorial links to help you install, pick models, or read release notes — not part of the upstream thread.