#1395 hip: a hipBLASLt tuning table for gfx1150 (Radeon 890M), and gfx1151's exact speed switches on gfx1150

Descrição

## Title
hip: a hipBLASLt tuning table for gfx1150 (Radeon 890M), and gfx1151's exact speed switches on gfx1150

<!-- If Applicable, reference the GitHub issue -->
Issue: Follow-up to #1217

## Summary

0.1.40.2 builds and runs on gfx1150 (Strix Point, Radeon 890M / 880M) since #1217, but there was no hipBLASLt table for
it, so the prompt's dense GEMMs ran on plain hipBLAS. This PR adds:

1. **`tools/hip/gfx1150-hipblaslt-100401.txt`**, calibrated on a Radeon 890M with ROCm 7.14.1's hipBLASLt. Prompts read
   **1.5-1.7x faster** (3,562 tokens: 124 -> 216 tok/s; 7,010 tokens: 131 -> 226 tok/s; medians of 3 interleaved
   rounds). Decode is unchanged.
2. **gfx1150 in `arch_default_env()`**: the gfx1151 table of exact speed switches less `STRATA_HCD_EXACT` (17 switches).
   With and without them, greedy answers were the same text. Decode is +2% and prompts +2-6%.
   `STRATA_GFX1150_DEFAULTS=0` turns them off.

The two commits stand alone: the second can be dropped without touching the first.

## What changed

**Commit 1: `hip: hipBLASLt tuning table for gfx1150 (Radeon 890M, ROCm 7.14.1 / hipBLASLt 100401)`**

- `tools/hip/gfx1150-hipblaslt-100401.txt` (new, 144 rows)
  - Covers the 16 dense GEMM geometries a Qwen3.8-Flash-Next prompt logs on the 890M, at T = 64, 128, ..., 16384
    - I collected them with a header-only table and `STRATA_HIPBLASLT_VERBOSE=1`
    - They are the same 16 geometries as in `gfx1201-hipblaslt-100500.txt`
  - Calibrated with `tools/hip/tune_hipblaslt.cpp` against AMD's TheRock ROCm 7.14.1 for gfx1150 (hipBLASLt 1.4.1
    `cd957402`, version number 100401), with the engine's 32 MiB workspace
  - The tool ran twice back to back. For each row, I kept the solution with the lowest **median of all six
    repetitions**
    - Now and then a single repetition stalled (a 100 ms call read up to 1.6 s)
    - Because of that, the tool's own mean-of-three picks differed in 26 rows between the two passes
  - Every row beats plain hipBLAS: 1.55-22.6x per GEMM, geometric mean 4.5x. For example, f16 N=12288 K=2560 at
    T=8192 takes 90 ms instead of 308 ms
  - No candidate failed the tool's accuracy gate
- `bench/results/2026-10-07-community-gfx1150/` (new)
  - `README.md`: the rig, the calibration, the measurements and the method
  - `runs.csv`: every run of the A/B
  - `tune-summary.csv`: per-row medians, the kept solution and the candidate count
  - `strata-gfx1150.json`: the server config used
- `bench/results/COMMUNITY.md`: one row for the new folder (the PR column still reads `-`)
- `docs/AMD_HIP.md`
  - Adds a "Shipped tables" entry for the gfx1150 table
  - The gfx1150 bullet now gives the measured speed and says to point `STRATA_HIPBLASLT_TUNING` at the table by hand,
    because setup does not install for gfx1150

**Commit 2: `core: gfx1151's exact speed switches on gfx1150 too, less STRATA_HCD_EXACT`**

- `include/strata/kernels/gfx_arch.hpp`: adds `gfx_arch_is_gfx1150()`
- `src/core/arch_defaults.cpp`
  - `arch_default_env()` returns the gfx1151 table for gfx1150 too, without `STRATA_HCD_EXACT`
    - That kernel reproduces gfx1151's hipBLASLt solutions 1176 / 1177
    - The gfx1150 table picks other solutions (539-555 for the bf16 rows), so on gfx1150 it would only report that it
      cannot run
  - Each architecture has its own opt-out: `STRATA_GFX1150_DEFAULTS=0` or `STRATA_GFX1151_DEFAULTS=0`
  - The start-up line names the architecture, e.g. `strata: gfx1150 (Strix Point): 17 exact speed switches on by
    default (STRATA_GFX1150_DEFAULTS=0 turns them off; ...)`
- `include/strata/core/arch_defaults.hpp`: the header comments cover gfx1150
- `src/core/arch_defaults_test.cpp`, new cases:
  - gfx1150, plain and with a feature suffix, gets the gfx1151 table less `STRATA_HCD_EXACT`
  - Each opt-out acts only on its own architecture
  - gfx1152, gfx1103 and gfx11500 stay empty
- `docs/STRIX_HALO.md`: section 4 notes that gfx1150 takes the same table less `STRATA_HCD_EXACT`
- `docs/AMD_HIP.md` and the bench README: a sentence each on the new default

**Measurements** (all from `bench/results/2026-10-07-community-gfx1150`)

Setup for the A/B:
- Qwen3.8-Flash-Next GSQ-RCO IQ3_XXS, all 24,576 experts in the GTT pool, `--prefill auto`
- Three rounds with the configurations interleaved, a fresh server for every run, one request at a time
- The engine's own `strata serve: prompt` timings

| Prompt | plain hipBLAS (+ the switches set by hand) | the table | the table + the switches |
|---|---:|---:|---:|
| 987 tokens | 99.1 tok/s | 152.4 | 153.1 |
| 3,562 tokens | 123.9 | 205.6 | 216.0 |
| 7,010 tokens | 130.5 | 219.4 | 225.8 |
| decode, 256-token chat | 15.3-15.4 tok/s | 14.9-15.0 | 15.2-15.3 |

- A 13,969-token prompt read in 107.7 s without the table and 61.5 s with it. A 24,383-token prompt read at 223 tok/s
  with it
- Answers (greedy, no thinking, four fixed prompts of 28 to 13,969 tokens)
  - The table: the first three answers were the same text with and without it. The 13,969-token answer diverged
    after its first sentence, because the summation order of the dense GEMMs changes
  - The switches: all four answers were the same text with and without them
- Tests with the table
  - `hip_prefill_hipblaslt_gemm` passes (`launches=4 fallbacks=0`)
  - A server run logged no fallbacks over prompts of 28-7,880 tokens
  - `hip_prefill_hcd_exact_parity` skips at every T (64-16384), because the table never sends the HC down projection
    to 1176 / 1177
- The engine built from commit 2, started with only `STRATA_HIPBLASLT_TUNING` in the config's env:
  - Printed the start-up line above
  - Gave the same four greedy answers as the run with the 17 switches set by hand
  - Read 151 / 220 / 231 tok/s at 1K / 3.6K / 7K, decode 15.4 tok/s
- Other tests
  - `arch_defaults_test` passes with Apple clang (C++17) and GCC 14.2
  - ctest on gfx1150 passes 58 of 61. The other three are the two hipBLASLt tests above, which skip without a table,
    and `ple_parity`, which needs `bench/micro` data. `iq_parity` by hand: all 10 formats ok

## Extra Notes

- **Machine**: MINISFORUM N5 PRO
  - Ryzen AI 9 HX PRO 370, Radeon 890M (gfx1150, 16 CUs)
  - 96 GB DDR5-5600, GTT raised to 70 GiB (`ttm.pages_limit=18350080`)
  - Proxmox VE 9.2 with kernel 7.0.14-20-pve; the engine runs in an unprivileged LXC container
  - Engine 0.1.40.2 built by hand with `-DCMAKE_HIP_ARCHITECTURES=gfx1150 -DSTRATA_PREFILL_MMQ=ON`
- **The table is for hipBLASLt 100401 (ROCm 7.14.1) only.** The engine refuses it with another version and falls back
  to hipBLAS, as with the other tables
- **setup.py is unchanged.** It still names gfx1150 as not supported, so the table and the build are set up by hand
- **Rounding-level switches** were measured but are not proposed as defaults. Each was added alone on top of the
  table + switches, one run each (base: 153 / 216 / 226 tok/s at 1K / 3.6K / 7K):
  - `STRATA_HIP_WMMA=1`: 155 / 252 / 277
  - `STRATA_PF_FUSED=1`: 155 / 257 / 255
  - `STRATA_PF_GEMM=1`: 166 / 243 / 262
  - The rest were flat or slower; the README has them
  - On this chip, `hip_prompt_attn_wmma` measured the WMMA attention 4.4-5.7x faster than the FP32 kernel, with an
    error against FP64 of about 3-4.5e-6 (the FP32 kernel: about 2e-6; output scale about 3.6)
- **Calibration note**: on this APU, `tune_hipblaslt`'s mean of three repetitions was sensitive to single stalled
  repetitions. A median over more repetitions (or the selection used here) gave steadier picks
- Not tested
  - Windows, setup's build path, `--prefill 16384` (the table has T=16384 rows, but the A/B used `--prefill auto`)
  - Models other than GSQ-RCO IQ3_XXS, and contexts above 24K tokens