#981 hip: a rocBLAS solution table for the FP16 prompt GEMMs on gfx103x (the GDN projection 7x below 1,152 tokens; short prompts -10 to -14%)

Description

On gfx103x the prompt path's dense GEMMs run through rocBLAS in FP16 (#835). rocBLAS picks the kernel for each product from its own table by shape, and on gfx1030 that pick is far from the best kernel the same library holds for some shapes. The GDN input projection (N = 10240, K = 2560) below T = 1152 tokens runs at 5.2 TFLOPS (8.1 ms at n = 799 on an RX 6900 XT); at T = 1152 rocBLAS's choice flips to a 38-TFLOPS kernel, and the best of the 237 solutions `rocblas_gemm_ex_get_solutions` lists for n = 799 runs it in 1.36 ms (6x). It is not an alignment effect: n = 800 is as slow as 799. In the engine's phase timing this is why a 799-token prompt's "gdn" phase (445 ms) was a third of an 8,874-token prompt's (1,345 ms).

This PR does for rocBLAS what `tools/hip/tune_hipblaslt` and `STRATA_HIPBLASLT_TUNING` do for hipBLASLt (which has no gfx1030 kernels):

- **`tools/hip/tune_rocblas.cpp`**: for each of the engine's 14 FP16 dense shapes at T = 256 ... 8192 (`--case`, `--tokens` change them), times rocBLAS's default kernel and every solution the library lists, checks the fastest against the default kernel's output (relative L2 within 2e-3, no element beyond 4 FP16 ulps of the largest output, nothing written outside the N x T result) and that they also run at T-1, T-37, T-63 and T/2+1 (a kernel can refuse a token count it was not measured at), and writes one row per shape and bucket: the winning solution index where it beats the default by `--min-gain` (5%), `default` otherwise. The `default` rows keep a neighbouring bucket's solution from reaching token counts where rocBLAS's own choice measured best.
- **`src/prefill/rocblas_tuning.hpp`**, **`Gemm::f16_inplace`** (`gemm.cu`): `STRATA_ROCBLAS_TUNING` names a table file, or a directory from which the engine takes `<arch>-rocblas-<version>.txt` for its own rocBLAS build. The table header carries the architecture and the full `rocblas_get_version_string` (tweak hash included) and any other pair is refused. The shape's nearest bucket decides; a tuned row's solution runs through a rocBLAS handle on the prompt stream, a `default` row or no row runs the kernel as before. One tiny default product at startup opens rocBLAS's kernel library (a solution-index call before that fails with `internal_error` although the index is valid - measured). A solution rocBLAS refuses at run time writes nothing and the default kernel runs that call; three refusals drop it for that shape. `STRATA_ROCBLAS_VERBOSE=1` prints the pick per shape and a `tuned_launches=... fallbacks=...` summary.
- **`tools/hip/gfx1030-rocblas-5.6.0.8d1ae90e.txt`**: the table for ROCm 10.0.0's rocBLAS (`/opt/rocm/lib`), 112 rows, 76 with a solution. Note that a HIP build on Ubuntu 26.04 links Ubuntu's own `librocblas5` (a 5.1 build) unless `lib_dirs` puts `/opt/rocm/lib` first, as setup's configuration does; 5.1's default kernel for the shape above is as slow, and its indices differ, so it would need its own table (`tune_rocblas`, about five minutes).
- **setup.py**: points the engine at `tools/hip` when a `<arch>-rocblas-*.txt` exists for the card; the engine's log then says `rocBLAS tuning enabled (112 rows, 76 with a solution, gfx1030, rocBLAS 5.6.0.8d1ae90e)` or why not.
- **CMake**: `find_package(rocblas)` (quiet) sets `STRATA_ROCBLAS_AVAILABLE`; `tune_rocblas` and the test build with it.
- **`tests/hip/prefill_rocblas_gemm.cpp`** (`hip_prefill_rocblas_gemm`, skips without `STRATA_ROCBLAS_TUNING`): refuses a table for another architecture or build, and checks Gemm's FP16 route against hipBLASEx (FP32 out) on the table's first two shapes at their bucket and at a smaller odd T. With `STRATA_ROCBLAS_VERBOSE=1` it reported `tuned_launches=4 fallbacks=0`.
- **docs/AMD_HIP.md**: the RDNA2 section's new entry, a "rocBLAS table (gfx103x)" subsection under "Tuning table", a line in the setup section.

**Measured** (2x RX 6900 XT 16 GB, one card, Ryzen 5 5600X, 128 GB, Ubuntu 26.04, ROCm 10.0.0; `6f32ec0` + #835 + this, `-DSTRATA_PREFILL_MMQ=ON`; GSQ-RCO IQ3_S, `--kv int8 --kv-resident 32768 --max-context 131072 --adapt-every 100000`; `STRATA_PREFILL_TIMING=1`, temperature 0, `max_tokens` 8; one server per arm, four prompts from four offsets of one text; details and per-run lines in `bench/results/2026-10-05-rdna2-rocblas-table/`):

| prompt tokens read | GPU timeline, default kernels | with the table | gdn phase | hc read phase |
| ---: | ---: | ---: | ---: | ---: |
| 799 | 3,329 ms | **2,985 ms (-10.3%)** | 443 -> 177 | 270 -> 77 |
| 495 | 2,477 ms | **2,139 ms (-13.6%)** | 365 -> 146 | 81 -> 52 |
| 1,108 | 3,391 ms | 3,395 ms (0%) | 598 -> 201 | 158 -> 99 |
| 8,874 | 11,978 ms | **11,459 ms (-4.3%)** | 1,346 -> 1,128 | 1,200 -> 898 |

The answers are the same in all four pairs. The 1,108-token prompt did not change because its saving reappeared as `wait copy`: that prompt is bound by the experts' PCIe transfer, and a faster GEMM only brings the GPU to the wait sooner. Prompts above ~1,150 tokens gain through their tail chunk and `hc read` only. Decode does not use these GEMMs.

**ctest** (`-DSTRATA_BUILD_TESTS=ON`, gfx1030, ROCm 10.0, `STRATA_ROCBLAS_TUNING=tools/hip`): 66 of 70 pass. `hip_prefill_rocblas_gemm` passes. Failed, for reasons outside this change: `ple_parity` (a Q2_0 PLE file that is not on this machine), `expert_multi_test` (this Zen 3 CPU has no AVX-512), `platform_memory_test` (`mlock` at the shell's default `ulimit -l`), `expert_cache_segmented_test` ("--vram-elastic (a segmented expert cache) is CUDA-only for now").

Builds on #835 (its commits are in this branch until #835 lands): the table serves the FP16-out route that #835 adds. One machine, one card, one model, one run per arm; another rocBLAS build needs its own table; gfx1031 / gfx1032 were not checked and the engine's architecture check refuses this table on them.

Developed with an AI coding assistant; every number above was measured on 2x RX 6900 XT (gfx1030, PCIe 4.0 x8 each) / Ryzen 5 5600X, ROCm 10.0.