#2513 prefill: STRATA_HIPBLASLT_TUNING may list one table per architecture

Descrição

## Title
prefill: STRATA_HIPBLASLT_TUNING may list one table per architecture

## Summary
On a layer split over two AMD architectures, the second card rejects the one tuning table and runs its prompt GEMMs on untuned hipBLASEx, although `tools/hip` ships its table. This lets the variable list one table per architecture. Measured on an R9700 (gfx1201) + RX 7900 XTX (gfx1100): a fresh 32K prompt reads at 1,870 tok/s instead of 1,277 (+46%), a 157K prompt at 3,006 instead of 1,846-1,866 (+62%). Decode is unchanged.

## What changed
`src/prefill/gemm.cu`, `create_hipblaslt_state` only (inside `#if defined(__HIPCC__) && defined(STRATA_HIPBLASLT_AVAILABLE)`):
- `STRATA_HIPBLASLT_TUNING` may hold several paths, separated by `:` (`;` on Windows). Each card loads the first table whose architecture family and hipBLASLt version match it.
- If none matches, the per-file errors are printed joined with `; ` and the card falls back to hipBLASEx, as before.
- **A single path behaves exactly as before**, apart from the error text when that one file is rejected, which is the same message.

Left alone:
- the table format and `hipblaslt_tuning.hpp`
- the per-solution checks and the GEMM path
- setup (it still writes one table; a split with two architectures can now be given both by hand)
- the SYCL copy of the loader (`sycl/src/prefill/gemm.dp.cpp`)

Before, on this split, the log said `prefill gemm: tuning architecture mismatch: file=gfx1201 runtime=gfx1100; using hipBLASEx`. Now it loads both: `hipBLASLt tuning enabled (26 rows, gfx1100, version 100202)` and `(24 rows, gfx1201, version 100202)`.

## How I checked
- **Machine:** Radeon AI PRO R9700 32 GB + RX 7900 XTX 24 GB, layer split 26 (R9700 first), Core Ultra 7 270K Plus, 125 GiB RAM, Ubuntu 26.04.1, ROCm 7.2.4, hipBLASLt 1.2.2 (100202). Engine 0.1.44.1 at `296b59d` built with setup's HIP build for `gfx1100;gfx1201`.
- **Model and settings:** Flash-Next GSQ-RCO IQ3_S, `--max-context 262144`, int8 KV, `--prefill auto:32768`.
- **Method:** one engine restart per arm, and a fresh prompt with a random nonce first so nothing is reused.

| Arm | 32K prompt tok/s (per restart) | 157K prompt tok/s | Short decode tok/s |
|---|---|---|---|
| gfx1201 table only (the XTX falls back) | 1,277 / 1,278 / 1,276 | 1,846 / 1,866 | 94.4 / 95.8 |
| both tables | 1,874 / 1,873 / 1,868 | 3,006 | 95.3 / 94.3 / 95.2 |

- **Default path:** with one path the loop runs once with the same `load()` call, so the result is the same as before.
- **Builds:** HIP was built. CUDA and SYCL were not. The changed function sits inside the HIP + hipBLASLt block, so the CUDA build compiles the same source as before, and SYCL has its own file.
- **Answer quality:** the XTX's prompt GEMMs now use the tuned hipBLASLt solutions from the shipped gfx1100 table, which round differently from hipBLASEx. Answer quality was not measured separately.

## Extra Notes
- Raw logs, scripts and the full write-up of this machine's tuning are on my fork: [docs/XTX_R9700_SPLIT.md](https://github.com/rindi-ismail/Strata/blob/xtx-r9700-tuning/docs/XTX_R9700_SPLIT.md) and [bench/results/2026-10-14-community-r9700-7900xtx-split](https://github.com/rindi-ismail/Strata/tree/xtx-r9700-tuning/bench/results/2026-10-14-community-r9700-7900xtx-split).
- A possible follow-up: setup could write both tables when it finds two architectures. It is not in this PR.

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