#322 hip: fast packed-byte intrinsics for RDNA3/RDNA4 (v_perm_b32 + SWAR)

本文

## What

Speeds up the four emulated CUDA packed-byte intrinsics in `include/strata/hip_compat/intrinsics.hpp` on the validated wave32 architectures (gfx1100, gfx1201), and adds an exhaustive parity test.

`__byte_perm`, `__vsub4`, `__vsubss4` and `__vcmpne4` are used in the i-quant dequant/sign paths (`iq_kernels.cu`: 28 uses, `native_mmvq.cu`: 14 uses). Today:

- HIP's own `__byte_perm` lowers to a byte-array access in private memory — scratch traffic in the table lookups.
- The packed-byte ops are per-lane scalar loops (8–12 instructions each).

This patch:

- `__byte_perm` → **`v_perm_b32`** in one instruction (selector nibbles spread to bytes, masked to 0–7).
- `__vsub4` / `__vsubss4` / `__vcmpne4` → branchless **SWAR** (Hacker's Delight 2-18 style). I probed the gfx1201 ISA directly: the RDNA1/2-era packed ops (`v_sub_u8`, `v_add_u8`, `v_cmp_ne_u8`, `v_sub_u16`, …) are **not in the ISA**, so lane isolation is done with the top-bit trick, not ISA lanes. `v_perm_b32` is present and is the only 1:1 mapping.
- `tests/hip/fast_intrinsics_parity.cpp` (registered as `hip_fast_intrinsics_parity`, HIP-only like `hip_intrinsics`): 262,144 cases against an independent host reference — **all 65,536 byte pairs × 4,096 selectors exhaustively**, plus randomized mixed-lane words, plus output guards.

## Evidence

Radeon AI PRO R9700 (gfx1201), engine 0.1.30, Swift 1.5 IQ3_XXS, 131072 ctx, INT8 KV, MTP, same machine/config, three fresh 512-token trials per arm:

| arm | sustained decode |
|---|---:|
| stock 0.1.30 | 55.0 tok/s |
| + this patch | **67.5 tok/s** |

- `hip_fast_intrinsics_parity`: PASS (262144 cases / 1048576 operation results, all 65536 byte pairs, 4096 selectors, 16 guards).
- HIP suite with the patch: **43/43** (`-E '^(ple_parity|platform_memory_test|expert_parity|pool_test)$'`, the same exclusions as the release docs; `hip_prefill_hipblaslt_gemm` skips, no gfx1201 table shipped).
- Real-model checks on the patched engine: arithmetic, system-marker recall, multi-turn recall, streaming, and native expert parity on both IQ3_XXS shards (layer 0 / layer 47) all passed.

## Notes

- CUDA path untouched; portable fallbacks kept for non-wave32 HIP targets.
- This is the intrinsics portion of #262, rebased onto 0.1.30's native RDNA4 support and re-validated — 0.1.30 shipped the hardware support without this block.
- The SWAR forms are branchless and overflow-safe; the exhaustive test covers every byte-pair for every selector, saturating edges included.