#3230 [Engine] perf(hip): enable wave32 QSA decode attention and MMVQ row reuse on gfx12

Beschreibung

**Problem and result**
This PR combines two bitwise-exact RDNA4 (gfx12, wave32) decode/speculative-verification optimizations and enables their parity/benchmark harnesses under `STRATA_ENABLE_HIP`:
1. **D5 — Wave32 QSA decode attention**: Replaces cross-warp `__syncthreads()` with `qda_sync_lds()` to avoid stalling on global KV VMEM loads, enables wave32 `attn_chunk_kernel_pre75<M,8,true>` by default on gfx12, enables bounded `attn_merge_v2_kernel<16>` by default on gfx12.
2. **D6 — Native MMVQ cross-row fragment reuse & LDS barriers**: Replaces cross-warp partial LDS reduction barriers with `mmvq_sync_lds()` on gfx12/gfx103x, enables `v_perm_b32` IQ4 lookup path across HIP targets, extends `b6_mmvq_rows_kernel` to SmallTraits (Q8_0, IQ4_NL) with preload of activation fragments, adds `b6_default_rows_gfx12` for auto row count on gfx12.

**Related discussion**
none

**Approach**
On gfx12, the prior implementation suffered from shared-memory barriers draining in-flight global KV VMEM loads due to `__syncthreads()` emitting full system scope barriers. The workgroup-scoped `qda_sync_lds()` and `mmvq_sync_lds()` avoid L0 vector cache invalidation and global-load drain while preserving synchronization. For MMVQ, activation fragments and reusable weight-code transforms are preloaded once across output rows, reducing redundant global loads. The auto row count table (`b6_default_rows_gfx12`) selects measured optimal rows per shape on gfx12, falling back to off elsewhere.

**Validation**

**Passed:**
- `qsa_scores_parity`: revision e88ee39d, GPU AMD Radeon AI PRO R9700 (gfx1201), observed PASS (zero bitwise differences)
- `qsa_topk_active_parity`: revision e88ee39d, GPU AMD Radeon AI PRO R9700 (gfx1201), observed PASS (selected IDs bitwise identical)
- `mmvq_multi_parity`: revision e88ee39d, GPU AMD Radeon AI PRO R9700 (gfx1201), observed PASS (0 bit-different with multi_exact on)
- `b6_mmvq_bench 80 0 -1`: revision e88ee39d, GPU AMD Radeon AI PRO R9700 (gfx1201), observed PASS (0 failed comparisons)
- `qsa_merge_daily_parity`: revision e88ee39d, GPU AMD Radeon AI PRO R9700 (gfx1201), observed PASS (0 bitwise differences)
- `arch_defaults_test`: revision e88ee39d, GPU AMD Radeon AI PRO R9700 (gfx1201), observed PASS

**Failed:**
- none observed

**Not run:**
- CUDA builds (no CUDA toolchain available)
- SYCL builds (no SYCL toolchain available)
- Metal builds (no Metal support)

**Limits**
- One-column MMVQ (`ncols == 1`) is unchanged: `b6_launch` is only invoked from `launch_multi_n` when `ncols > 1` and `multi_exact()` is true. Evaluated `STRATA_Q6_PACKED=1` for one-column `Q6_K` heads on gfx1201, but at the default 640-block grid it regressed the 106,299-row MTP head (339.9 → 408.8 us) and 248,320-row output head (794.1 → 960.4 us) while requiring ~213–497 MiB of additional persistent VRAM; both `STRATA_Q8_PACKED` and `STRATA_Q6_PACKED` remain off by default.
- Target hardware validation: Runtime parity, microbenchmarks, and end-to-end 397B measurements were executed on gfx1201 (AMD Radeon AI PRO R9700). gfx906, gfx1030, gfx1100, gfx1200 (HIP) and sm_75, sm_89, sm_120 (CUDA nvcc) were verified by compilation only; non-gfx12 default dispatch paths are unchanged.