Pull requests / #2392

#2392 hip: hipBLASLt tuning table for gfx1101 (RX 7800 XT, hipBLASLt 1.2.0)

open · @skuridin · 0 comentários · No GitHub

BenchmarksSetup & installAMD / HIPModels & quantsDocumentation

Descrição

## Title

Issue: none (follows up #254, which added gfx1101 support but no table)

## Summary

There was no hipBLASLt table for gfx1101, so the RX 7700 XT / 7800 XT ran the prompt's dense GEMMs on plain hipBLAS. This adds one for hipBLASLt 1.2.0 (100200), the version in the ROCm 7.10.0a20251120 wheels that setup installs for gfx110X. `hipblaslt_table()` finds it by arch + version, so setup picks it up with no code change. On my 7800 XT, prompts read about 70% faster and decode doesn't change.

## What changed

- `tools/hip/gfx1101-hipblaslt-100200.txt`: new, 43 rows. Same geometries as `gfx1100-hipblaslt-100401.txt` (the small-T expert GEMM included), calibrated with `tune_hipblaslt --workspace-mib 32`.
- `docs/AMD_HIP.md`: a bullet for it under "Shipped tables", and the "Community-validated cards" intro now mentions it (it keeps the "make your own" hint for gfx1101 on other hipBLASLt versions).

Nothing else changes. Other cards and other hipBLASLt versions don't load this file.

## Extra Notes

Setup: RX 7800 XT 16 GB, Ryzen 7 5800X3D, 60 GiB RAM, Ubuntu 26.04 (kernel 7.0.0), engine 0.1.44.1 (296b59d2) in Docker, Qwen3.8-Flash-Next GSQ-RCO IQ3_S, 131072 ctx, `--kv q4_0`.

Calibration: two `tune_hipblaslt` passes back to back picked the same solution id in all 43 rows. Every row beats plain hipBLAS, 2.98x to 16.6x per GEMM (median 4.9x / 6.0x in the two passes), so no rows were dropped.

Tests, with `STRATA_HIPBLASLT_TUNING` set and `STRATA_HIPBLASLT_VERBOSE=1`:
- `hip_prefill_hipblaslt_gemm` passes, `launches=4 fallbacks=0`
- `hip_prefill_hcd_exact_parity` skips (exit 77): the table sends no tested T to solution 1176 / 1177

End to end, three interleaved pairs (baseline, table, baseline, ...), fresh server each run, cold prompts:

| | plain hipBLAS (median, range) | with this table (median, range) |
|---|---|---|
| 35,537 tok prompt | 447.0 (446.7-449.1) tok/s | 761.9 (758.0-762.1) tok/s |
| 91,320 tok prompt | 443.6 (443.4-444.5) tok/s | 751.9 (751.6-752.5) tok/s |
| decode after prompt | 42.0 / 48.5 tok/s | 41.5 / 49.3 tok/s |

Each table run logged `hipBLASLt tuning enabled (43 rows, gfx1101, version 100200)` and `launches=13328 fallbacks=0 unique_fallback_shapes=0`.

Not tested: the full HIP ctest, the 7700 XT, and system-ROCm installs (they have a different hipBLASLt, so this table won't load there anyway). The output isn't bit-identical to plain hipBLAS, since the GEMM kernels differ.

No site

Links install, modelos, releases.