Pull requests / #3058

#3058 Add a gfx1101 hipBLASLt table (RX 7800 XT, hipBLASLt 1.5.0): +14/+52/+72% prompt at 1K/8K/24K

open · @durikim1-cpu · 0 Kommentare · Auf GitHub

BenchmarksSetup & installAMD / HIPModels & quantsDocumentationWindowsLinux

Beschreibung

`docs/AMD_HIP.md` says gfx1101 has no hipBLASLt table ("make one with Tuning table and compare the prompt speed with
and without it"). This adds one, calibrated on an RX 7800 XT with the hipBLASLt shipped in the 0.1.44.1 Windows HIP
package (version 100500), plus the measurements.

## What changes

- **Adds** `tools/hip/gfx1101-hipblaslt-100500.txt` (70 rows). Setup picks it only for a gfx1101 card whose installed
  hipBLASLt reports 100500, by the existing `hipblaslt_table()` file-name rule; no code changes.
- **Adds** `bench/results/2026-10-16-community-rx-7800-xt-hipblaslt-100500/`: calibration method and script, per-row
  timings, launch configs and per-run prompt-speed data.
- **Edits** `docs/AMD_HIP.md`: the "gfx1101 has none" sentence and one entry in "Shipped tables".
- **Leaves alone**: the engine, setup, every other table. Other cards and other hipBLASLt versions are unaffected.

## Measured

Prebuilt Windows HIP engine 0.1.44.1, same `strata.exe` in both arms, only `STRATA_HIPBLASLT_TUNING` differs.
Swift 1.5 IQ3_XXS, 32K context, fresh prompts, medians of 3:

| Prompt tokens | Prompt tok/s without → with | Decode tok/s | TTFT s |
| ---: | --- | --- | --- |
| 1,060 | 161.7 → 184.5 (+14%) | 28.2 → 29.7 | 6.60 → 5.79 |
| 8,041 | 335.1 → 510.8 (+52%) | 29.6 → 29.0 | 24.07 → 15.80 |
| 24,054 | 404.6 → 695.2 (+72%) | 29.9 → 29.2 | 59.51 → 34.67 |

Per GEMM against plain hipBLAS: median 6.45x, geometric mean 5.80x, 0.81x–15.9x; all 70 rows pass the upstream
accuracy and padding gates. Correctness set 11/12 in both arms (the same arithmetic question fails in both).

## How it was calibrated

The Windows package has no `tune_hipblaslt`, so the folder includes a ctypes port of `tools/hip/tune_hipblaslt.cpp`
that calls the packaged DLLs unchanged and applies the same gates. Geometries: the 16 of `gfx1100-hipblaslt-100500.txt`
at T = 8192 / 6144 / 4096 / 1024, plus `f16 1280 2560 1280` at T = 64…8192. No solution id is copied from another GPU.

## Not checked

- `fallbacks=0` with `STRATA_HIPBLASLT_VERBOSE=1` (the summary line was not captured) and the HIP ctest
  (`hip_prefill_hipblaslt_gemm`); the prebuilt Windows package has no test binaries.
- Linux, other 1.5.0 builds (solution ids may be numbered differently), Flash-Next with only the table changed.
- One row (`bf16 1 2560 1 T=1024`) is 0.81x of hipBLAS (hipBLASLt 0.0327 ms vs hipBLAS 0.0266 ms). I kept it because the end-to-end runs
  used the full table; happy to drop that geometry's rows if you prefer.

I can run any extra check you want on this card.

Mehr auf der Site

Links zu Install, Modellen, Releases.