#638 hip: AMD Instinct MI50 / MI60 / Radeon VII (gfx906, wave64) as an opt-in build

Descrição

## What

The HIP backend (`STRATA_ENABLE_HIP`) is wave32 RDNA only, and the engine refuses a wave64 card. AMD Instinct MI50 / MI60 and Radeon VII (gfx906) are cheap 16/32 GB cards with ~1 TB/s HBM2. This PR adds an **opt-in** build for them: `-DSTRATA_HIP_GFX906=ON -DCMAKE_HIP_ARCHITECTURES=gfx906`.

It compiles the CUDA sources as HIP through a small compat layer (`include/strata/platform/hip_compat/`). A CUDA warp is a logical half of the 64-lane wavefront: 32-wide shuffles, and a ballot of its own half. **With the option off, nothing changes for CUDA or the RDNA backend.** Every new path is under `STRATA_HIP_GFX906`, a macro neither existing build defines.

## Why a separate option and not `STRATA_ENABLE_HIP` + gfx906

The RDNA backend's kernels and compat layer assume wave32 throughout. Treating the CUDA warp as half a wavefront let the CUDA path run unchanged on gfx906 first, and then the hot kernels got wave64 layouts of their own. If you'd rather fold gfx906 into `STRATA_ENABLE_HIP` / `cmake/hip_backend.cmake`, I'm happy to rework it that way. Tell me which you prefer.

## Correctness fixes found on the way (each changed results or hung)

- `__fmul_rn` / `__fadd_rn` / `__fsub_rn` contracted into FMA under HIP. They no longer do, so the elementwise and quantize_act parity tests pass bitwise.
- ggml MMQ got NVIDIA compute capability 9.0 for gfx906 and hung the prompt path. It now gets the AMD arch code.
- HIP's `__byte_perm` is a scratch-memory array. The IQ4 table lookup now uses llama.cpp's `v_perm_b32` sequence: agent-request decode on 2x MI50 went from 29.9 to 39.2 tok/s.
- gfx906's wall clock is 25 MHz, not 100: fixed in the stage profiler.
- The `CUDART_VERSION` toolkit check is skipped under HIP. `gdn_rec_parity` (cp.async, CUDA occupancy API) is not built for gfx906.

## wave64 kernels (gfx906 only; each has an env switch back to the CUDA layout for A/B)

Each was checked bitwise against the layout it replaces: parity tests, `native_expert_bench` on real GGUF rows, and the greedy text of a fixed request.

| kernel | change | effect | switch |
|---|---|---|---|
| grouped native experts | mode 7: signs as `dp4a(g ^ m, u) - dp4a(m, u)` (gfx906 has no byte SIMD; `__vcmpne4`/`__vsub4` were ~12 word ops each), grids and activations in LDS, one weight load per group | verify window 60.0 -> 51.3 ms | `STRATA_EXP_MODE=2` |
| MMVQ | one wavefront per row, 64-lane butterfly | +2.3% decode | `STRATA_MMVQ_WAVE=0` |
| router top-10 | one wavefront, register argmax | bit-identical | - |
| verify-window routing (256-expert router) | one BF16 projection + one top-k per window | -3 ms per window | `STRATA_ROUTE_PER_TOKEN=1` |
| `gr_down_multi` (opt-in) | split along K; sums in another order than the single-token read | +2.9% | `STRATA_GR_SPLIT=1` |
| hyper-connection read | loads in flight, 8 lanes per row | 51.3 -> 50.2 ms | `STRATA_GR_FAST=0` |

`docs/AMD_HIP.md` gets a gfx906 section: the build recipe on the community ROCm 7.14 image (AMD dropped gfx906 from current ROCm), the table above, and measurements.

## Measured (on the port this is cut from)

2x MI50 16 GB (85 W each, PCIe 3.0 x16), Xeon E5-2666 v3, 32 GB DDR4. Coder IQ1_M, 128K context, `--kv int8 --kv-resident 32768`, `--layer-split 27`, `--prefill 4096`, MTP `--spec 4`.

- All 12,288 experts are in VRAM.
- A 17K-token agent prompt is read in 43.5 s; decode on it runs at 57.7-59.3 tok/s.
- 6/6 long agent requests and a 900 s two-client load test x6 (407 requests) ran with no stalls or faults.
- For reference, llama.cpp on the same model and machine: 26.9 tok/s `tg128`.

These numbers also include #639 (layer-split weight trim) and #640 (mapped arena for small RAM), which are separate PRs.

## Verified for this branch

- Builds `all` with `STRATA_HIP_GFX906=ON` (gfx906, ROCm 7.14).
- Builds `all` with `STRATA_ENABLE_HIP=ON` for gfx1100, so the RDNA backend is intact.
- Builds `all` with CUDA 13.0 for sm_86, so the CUDA path is intact.

## Status

Re-run on 2x MI50 with this branch: ctest 35/35 on the cards and the greedy-text comparison with the port branch are in the comment below; builds checked for gfx906, gfx1100 and CUDA sm_86 at the branch head.

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