[bsorensen110]

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 HIP: gfx1201 (Radeon AI PRO R9700) support, faster packed byte intrinsics, and a gfx1201 hipBLASLt table #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.

[Copilot]

Copilot review overview

🟡 Changes recommended

The parity test omits selector encodings 8–F and the PR overstates its Cartesian coverage.

Review effort: Balanced
Findings: 1 · 1

Open (2)

• Parity test omits required selector encodings 8 through f · New
• Test coverage does not implement the claimed Cartesian product · New

What changed in this PR

Optimizes HIP packed-byte intrinsics for RDNA3/RDNA4 and adds parity testing.

Changes:

• Implements v_perm_b32 and SWAR packed-byte operations.
• Adds exhaustive/randomized HIP parity coverage.
• Registers the new test with CMake.

File	Description
include/​strata/​hip_compat/​intrinsics.hpp	Adds optimized packed-byte intrinsics.
tests/​hip/​fast_intrinsics_parity.cpp	Tests intrinsic parity and buffer guards.
CMakeLists.txt	Registers the parity test.
cmake/​hip_backend.cmake	Compiles the test as HIP.

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[bsorensen110]

@copilot Fix the code for this review comment.

When a review comment includes a suggested change, apply the suggestion exactly.

Do not make changes beyond what is described in the linked review comment.

[bsorensen110]

@copilot Fix the code for all comments in this review thread.

When a review comment includes a suggested change, apply the suggestion exactly.

Do not make changes beyond what is described in the linked review thread.

[clapbr]

Seems legit on my 7900 XTX. Good gain and correctness seems perfect after a few Astra reviews/parity tests/logit tests and benchs.

Three MTP runs are complete. Median decode improved **57\.4 → 68.7 tok/s (+19.8%)**; median prefill was essentially unchanged at **710 → 709 tok/s**. Every continuation matched exactly, and acceptance stayed **93/136**.

[Niko1221]

Thanks - the same intrinsics shipped in 0.1.31 from #262 (with a test that also covers selector bit 3). Closing as shipped.