#695 hip: the tuning table's arch gate accepts the RDNA4 sibling (gfx1200 <-> gfx1201)

Descrição

## Summary

`TuningTable::load()` compares the tuning table's architecture to the runtime's with strict equality (`arch != expected_arch`). A card that reports its RDNA4 *sibling* — which happens routinely under `HSA_OVERRIDE_GFX_VERSION` — rejects its own calibration at startup and loses the tuned prefill route wholesale:

```
prefill gemm: tuning architecture mismatch: file=gfx1201 runtime=gfx1200; using hipBLASEx
```

This patch makes the header gate family-wide (`gfx1200` ↔ `gfx1201`, and symmetrically `gfx1100/1101/1102`), matching the scope of the gate that actually protects correctness — the per-solution check in `src/prefill/gemm.cu`. The hipBLASLt version gate is untouched; cross-family pairs stay rejected. Adds a host-only unit test (synthetic tables, no model, no GPU): RED on the strict gate, GREEN on the family gate.

## The symptom (live service)

Radeon AI PRO R9700 (gfx1201), Linux, ROCm nightly 10.2.0a (hipBLASLt 100500), engine 0.1.38, profile `tools/hip/gfx1201-hipblaslt-100500.txt` (32 rows). Engine startup log:

```
strata generate: GPU 0: AMD Radeon AI PRO R9700 (gfx1200)
prefill gemm: tuning architecture mismatch: file=gfx1201 runtime=gfx1200; using hipBLASEx
```

The card is a gfx1201 (`rocminfo`: `Name: gfx1201`, "AMD Radeon AI PRO R9700"), yet the engine sees `gfx1200`.

## Root cause

The engine runs with `HSA_OVERRIDE_GFX_VERSION=12.0.0` (the llama.cpp-documented override practice, set in this machine's systemd user environment). A minimal `hipGetDeviceProperties` probe on the live card:

| build `--offload-arch` | env | `gcnArchName` |
|---|---|---|
| gfx1201 | (none) | **gfx1201** |
| gfx1200 | (none) | **gfx1201** |
| gfx1201 gfx1200 | (none) | **gfx1201** |
| gfx1201 | `HSA_OVERRIDE_GFX_VERSION=12.0.0` | **gfx1200** |
| gfx1201 | `HSA_OVERRIDE_GFX_VERSION=12.0.1` | **gfx1201** |

So the runtime arch string is environment-dependent, not card-dependent. The strict gate then traps nightly users in a corner:

- the shipped `gfx1200-hipblaslt-100202.txt` table is rejected by the **version** gate on a 100500 runtime (measured: `tuning table rejected: hipBLASLt version mismatch: file=100202 runtime=100500`);
- the only table that passes the version gate is the gfx1201/100500 one — which the **arch** gate rejects, because the override renamed the runtime.

Net effect: no tuned route at all, with a startup line that reads like a config error.

## Why the header gate is the wrong place for strict equality

The header check is a coarse guard; the real safety net is per-solution and per-shape, in `resolve_hipblaslt_algo()` (`src/prefill/gemm.cu`):

1. `hipblaslt_ext::getAlgosFromIndex()` — the table's `solution_id` must exist in the running library's index;
2. `hipblaslt_ext::matmulIsAlgoSupported()` — the algo must support the actual geometry/beta;
3. on any failure the call falls back to hipBLASEx and is counted (`lt_launches` / `fallbacks`, reported by `STRATA_HIPBLASLT_VERBOSE=1`).

A family-wide header gate therefore cannot make a wrong kernel run: every row is validated against the live library before its first launch. The family scope is also what ROCm itself uses — `rocminfo` on this card lists `amdgcn-amd-amdhsa--gfx12-generic` alongside `--gfx1201`.

## The change

9 lines in `src/prefill/hipblaslt_tuning.hpp` (header gate compares `arch.substr(0, 5)`), plus a host-only test and its CTest registration. No kernel, tuner, table format, or version-gate changes.

## Evidence

**1. Header gate matrix** (`tests/prefill/tuning_header_arch_family_test.cpp`, 10 cases, synthetic tables):

| file arch/ver | runtime arch/ver | strict (base) | family (patched) |
|---|---|---|---|
| gfx1201/100500 | gfx1201/100500 | ACCEPT | ACCEPT |
| gfx1200/100500 | gfx1200/100500 | ACCEPT | ACCEPT |
| **gfx1201/100500** | **gfx1200/100500** | **REJECT** | **ACCEPT** |
| **gfx1200/100500** | **gfx1201/100500** | **REJECT** | **ACCEPT** |
| gfx1201/100202 | gfx1201/100500 | REJECT (version) | REJECT (version) |
| gfx1200/100202 | gfx1201/100500 | REJECT (arch) | REJECT (version) |
| gfx1100/100202 | gfx1201/100202 | REJECT | REJECT |
| gfx1201/100202 | gfx1100/100202 | REJECT | REJECT |
| **gfx1100/100202** | **gfx1101/100202** | **REJECT** | **ACCEPT** |
| gfx1030/100202 | gfx1201/100202 | REJECT | REJECT |

Base: `10 cases, 3 failures` (the sibling rows). Patched: `10 cases, 0 failures`. Version gate and cross-family rejections are unchanged.

**2. Production Gemm probe** (`build-hip/hip_prefill_hipblaslt_gemm`, `STRATA_HIPBLASLT_VERBOSE=1`), same binary, same table (gfx1201/100500):

| runtime | header result | launches | fallbacks |
|---|---|---|---|
| native (gfx1201) | `tuning enabled (32 rows, gfx1201, version 100500)` | **4** | **0** |
| override 12.0.0 (reports gfx1200) | `tuning enabled (32 rows, gfx1200, version 100500)` | 0 | 4 (`solution 110416/135699 unavailable …`) |

This is the honest nuance: the patch removes the wholesale startup rejection, and the per-solution gate then decides. Under the 12.0.0 override the gfx1201 solution IDs do **not** resolve in the library's gfx1200 view, so every call falls back to hipBLASEx — identical performance to today, no regression, and the diagnostics change from one misleading "architecture mismatch … using hipBLASEx" line to explicit per-shape fallback records.

**3. Live service** (engine `dafa4bf79a64272b`, source fingerprint `a2a2d19670350e0f`, v0.1.38): after the patch the engine starts with `prefill gemm: hipBLASLt tuning enabled (32 rows, gfx1200, version 100500)`, service healthy, arithmetic/recall/streaming checks pass. With verbose on, the running engine shows the per-shape `Lt fallback; solution … unavailable` records for every projection geometry — safe degradation, exactly as the probe predicts.

**4. What this does *not* fix** (measured, labelled as such): tuned performance under `HSA_OVERRIDE_GFX_VERSION=12.0.0` on a gfx1201 card. The solution IDs are scoped to the family *view* the library is given; under 12.0.0 the gfx1201 IDs don't resolve. For an R9700 the override is unnecessary (native gfx1201 works: 4 launches / 0 fallbacks); if an override is wanted, `12.0.1` keeps the card reporting gfx1201 (probe above). A docs note in `docs/AMD_HIP.md` recommending `12.0.1` (not `12.0.0`) for gfx1201 cards would complement this patch — I did not include it here to keep the diff minimal.

## Test plan

- [x] `tuning_header_arch_family_test`: RED on `origin/main` (3 failures), GREEN with the patch (0 failures)
- [x] Production probe, native runtime: 4 launches / 0 fallbacks (unchanged)
- [x] Production probe, override runtime: table loads, 4 per-shape fallbacks, outputs match hipBLASEx (4/4 parity cases PASS)
- [x] Live service on R9700: table accepted, health 200, arithmetic/recall/streaming pass
- [ ] CI (no GPU on the test target — the new test is host-only)

## Reproduction

```sh
# header matrix (host-only, no GPU)
g++ -std=c++20 -Isrc/prefill tests/prefill/tuning_header_arch_family_test.cpp -o t && ./t

# probe, native vs override
LD_LIBRARY_PATH=/opt/rocm/lib STRATA_HIPBLASLT_TUNING=tools/hip/gfx1201-hipblaslt-100500.txt \
  STRATA_HIPBLASLT_VERBOSE=1 build-hip/hip_prefill_hipblaslt_gemm
HSA_OVERRIDE_GFX_VERSION=12.0.0 LD_LIBRARY_PATH=/opt/rocm/lib \
  STRATA_HIPBLASLT_TUNING=tools/hip/gfx1201-hipblaslt-100500.txt \
  STRATA_HIPBLASLT_VERBOSE=1 build-hip/hip_prefill_hipblaslt_gemm
```

Environment: AMD Radeon AI PRO R9700 (gfx1201), CachyOS/Arch, ROCm nightly 10.2.0a20260914 (hipBLASLt 100500, HIP runtime 71726391), Strata 0.1.38.

🤖 Generated with Hermes Agent