#2374 gfx1011 (Navi 12) on real hardware: 96/105 ctest pass; `hip_q2_zero` fails, fixed by adding `__gfx1011__` to the Q2_0 signed-zero workaround

Descripción

`docs/OLDER_GPUS.md` lists gfx1011 as "untested on hardware". Here is a hardware run.

**Hardware / software**
- 3× AMD BC-160 (Navi 12, gfx1011, 8 GB HBM2 each); host i5-6500 (4 cores, no AVX-512), 16 GB RAM
- Ubuntu, kernel 7.0.0-38-generic (in-tree amdgpu)
- Strata v0.1.43.1, commit `a0116b2`
- ROCm from TheRock wheels: `rocm[libraries,devel]==7.14.0a20260608`, index `https://rocm.nightlies.amd.com/v2/gfx101X-dgpu/`
- configure: `-DSTRATA_ENABLE_HIP=ON -DSTRATA_ENABLE_CUDA=OFF -DSTRATA_BUILD_TESTS=ON -DSTRATA_PREFILL_MMQ=ON -DCMAKE_HIP_ARCHITECTURES=gfx1011`

**Results**
- Build: 0 errors.
- `strata-device`: `HIP arch gfx1011 wave32 (compiled for gfx1011)`, 18 CUs, 7.984 GiB VRAM.
- `ctest`, run separately on each card (`HIP_VISIBLE_DEVICES=0/1/2`), with `ple_parity` and `platform_memory_test` excluded: identical on all three cards — **96 passed, 8 skipped, 1 failed** out of 105, ~170–180 s per card. No amdgpu errors in the kernel log.
- Skipped (expected for this hardware): `hip_prefill_wmma_gemm_parity`, `hip_prefill_hcd_exact_parity`, `hip_prompt_attn_wmma` (no WMMA on RDNA1), `hip_prefill_hipblaslt_gemm`, `hip_prefill_fused_moe`, `hip_prefill_fused_iq`, `expert_cache_segmented_test`, `expert_multi_test`.
- Failed: `hip_q2_zero`:
  ```
  index=65 got=0x0000 expected=0x8000
  ...
  hip_q2_zero: 1024 values, 128 failures
  ```

**Cause and fix**

This is the same signed-zero folding the guard in `dq_q2_0` (`src/kernels/cuda/iq_kernels.cu`) already works around for gfx1012 on HIP < 7, gfx1151 and gfx1200/gfx1201. gfx1011 is not in that list, and with ROCm 7.14 it shows the same behaviour. Adding it fixes the test on all three cards (1024 values, 0 failures):

```diff
-#if defined(__HIPCC__) && ((defined(__gfx1012__) && HIP_VERSION_MAJOR < 7) || defined(__gfx1151__) || \
+#if defined(__HIPCC__) && ((defined(__gfx1012__) && HIP_VERSION_MAJOR < 7) || defined(__gfx1011__) || defined(__gfx1151__) || \
     defined(__gfx1200__) || defined(__gfx1201__))
```

Since gfx1011 still folds on HIP 7.14, gfx1012 on HIP ≥ 7 may need the guard too. I have no gfx1012 card to check.

**Not tested:** running a model. This box has 16 GB of RAM, below the documented minimum. The memory planner (`--max-context 20480`) reports a 4105-slot expert cache (5.68 GB) plus 0.27 GB KV, which fits a single 8 GB card.

**Small doc note:** `OLDER_GPUS.md` says gfx1011 is "the same RDNA1 ISA as gfx1010 (Navi 12)". Navi 12 is gfx1011 itself; gfx1010 is Navi 10.