#3343 [Engine] rocBLAS table for gfx1030 with rocBLAS 5.4.0 (TheRock gfx103X-all 7.13 wheels): short follow-up prompts 12% faster

Descrição

## Problem and result

The rocBLAS tables in `tools/hip` cover gfx1030 with rocBLAS 5.6.0.8d1ae90e (ROCm 10.0.0) only. An RX 6000 engine built
against TheRock's `gfx103X-all` 7.13.0a20260515 wheels - what setup installs into `.venv` when the PC has no system
ROCm 7 (`ROCM_FAMILY_PINS`), and what the `rx6900` package of #3329 carries - loads rocBLAS 5.4.0.ad84daa81c, so with
`STRATA_HIP_PROMPT_F16=1` the engine logs

```
prefill gemm: cannot open rocBLAS tuning file (.../tools/hip/gfx1030-rocblas-5.4.0.ad84daa81c.txt); rocBLAS's default kernels run
```

This adds that table, made with `tools/hip/tune_rocblas` on an RX 6900 XT against that rocBLAS, and names it in
docs/AMD_HIP.md. With it a 268-395-token follow-up on a cached conversation reads 11.7% faster (median), every pair
10-13%; longer fresh prompts change little (numbers below).

Related: #3329 (the package; independent of this PR). Its docs sentence "The package's rocBLAS (5.4.0) has no table in
`tools/hip` yet" becomes stale once both are merged; I will update whichever lands second.

## Approach

Data only: the engine already picks `<arch>-rocblas-<version>.txt` from `STRATA_ROCBLAS_TUNING` for its own rocBLAS
build and refuses any other, and setup already points `STRATA_ROCBLAS_TUNING` at `tools/hip` when a table for the
card's arch exists. Nothing changes for an engine on another rocBLAS build (it keeps refusing this file) or without
`STRATA_HIP_PROMPT_F16=1` (the table serves that route only). The packaged engines do not carry `tools/hip` tables
yet, so a Terminal user gets it only with `STRATA_ROCBLAS_TUNING` set by hand until packaging includes the tables.

`tune_rocblas` was built from this tree against the wheels' rocBLAS and run with the rx6900 package's libraries
(`LD_LIBRARY_PATH=<package>`; `ldd` shows its `librocblas.so.5` and `libamdhip64.so.7`). It ran one token bucket per
process (`--tokens 256`, ..., `--tokens 4096`) and the 8192 bucket one shape per process (`--case 8192,N,K,N` for the
14 default shapes), with pauses between runs; every row is one shape at one bucket, so the parts were merged in the
tuner's own order under one header. 112 rows (the same shape/bucket keys as the 5.6.0 table), 79 with a solution,
33 `default`.

## Validation

RX 6900 XT 16 GB (gfx1030), Ryzen 9 5950X, 128 GB DDR4, Arch Linux; engine 0.1.45.3 (e88ee39d), the rx6900 package of
#3329 (rocBLAS 5.4.0.ad84daa81c from `rocm[libraries,devel]==7.13.0a20260515`).

**Passed**

- `hip_prefill_rocblas_gemm` built from this tree against the same rocBLAS, `STRATA_HIP_PROMPT_F16=1`:
  `STRATA_ROCBLAS_TUNING=tools/hip/gfx1030-rocblas-5.4.0.ad84daa81c.txt` and `=tools/hip` (the engine's own pick): 4 of
  4 cases PASS, "outputs match hipBLASEx within FP16 rounding" (relative_l2 0.000208, max_abs 0.00196, 4 ulp = 0.0156),
  exit 0. With the 5.6.0 table: "tuning table rejected: rocBLAS version mismatch: file=5.6.0.8d1ae90e
  runtime=5.4.0.ad84daa81c", exit 0.
- The engine with the table: `prefill gemm: rocBLAS tuning enabled (112 rows, 79 with a solution, gfx1030, rocBLAS
  5.4.0.ad84daa81c)`.
- A/B on the server, UD-Q4_K_XL, the same settings in both arms except `STRATA_ROCBLAS_TUNING=tools/hip`
  (`STRATA_HIP_PROMPT_F16=1 STRATA_SH_STREAM=1 STRATA_ROUTE_TAIL_SKIP=7 STRATA_PREFILL_CPU_SHARE=auto
  STRATA_HIP_ADAPT_KERNEL_COPY=1`, `--resident-budget-gib 71 --pcie-frac 0.35 --mtp-draft-vocab data/draft_vocab_en.bin
  --kv int8 --kv-resident 32768 --max-context 131072 --spec 4`). Arms in the order without, with, with, without, each on
  a fresh server; per arm 3 rounds of a fresh prompt, a follow-up on it and a fresh longer prompt, different text per
  round, the same text in both arms. Server prompt times in ms:

  | prompt | without the table | with the table | median |
  |---|---|---|---|
  | follow-up, 268-395 new tokens on a cached conversation | 1585 1971 1873 / 1566 1930 1830 | 1427 1709 1636 / 1363 1715 1634 | 1,851 -> 1,635 (-11.7%) |
  | fresh, 745-1,569 tokens | 4121 3612 3459 / 2785 3631 3496 | 2630 5531 3437 / 2873 3554 3434 | 3,554 -> 3,435 (-3.3%) |
  | fresh, 4,859-6,690 tokens | 7906 8984 10417 / 8043 9156 9851 | 8049 8980 9660 / 7964 8972 9662 | 9,070 -> 8,976 (-1.0%) |

  The first value of each arm is the first request after the load (not a stable measurement); the 5531 is one slow
  request in the table arm.

**Failed**

- None observed.

**Not run**

- Another gfx1030 card (RX 6800 / 6800 XT / 6950 XT): the solution indices are valid there, the picks may differ.
- gfx1031: a table is per architecture; none was made.
- A setup install from scratch on the wheels (the same rocBLAS build as the package, but setup's path itself was not run).
- A second tuning run to see how stable the picks are; answer-quality benchmarks (the smoke test checks the first two
  shapes against hipBLASEx; every chosen solution passed the tuner's own check against the default kernel's output).

## Limits

- Only short follow-ups gain much; fresh prompts of 5K+ tokens are within noise.
- Packaged engines need `STRATA_ROCBLAS_TUNING` set by hand until the packages carry the `tools/hip` tables.

Measured and written with Claude Code.

🤖 Generated with [Claude Code](https://claude.com/claude-code)

https://claude.ai/code/session_012oqtnhtRwYGZ7yDfoVULrX