CUDA/HIP: use hardware v_perm_b32 for Q1_0 vec_dot on AMD (+110% decode on gfx1201)

[The-Monk]

Overview

Follows the same approach as #26753, which did this for Q2_0.

HIP has no hardware __byte_perm; it lowers to a software routine. The current Q1_0 unpack calls it 20 times per 32-element chunk -- 2 for the nibble indices, 4 for the byte values, 4 to unshuffle, twice per block. That makes Q1_0 decode VALU-bound on AMD rather than bandwidth-bound.

This spreads the sign bits arithmetically and selects the byte constants with one hardware v_perm_b32 per nibble:

• n * 0x00204081 lands bit i at position 8i, so & 0x01010101 leaves one sign bit per byte
• 0x0D - spread selects the v_perm_b32 byte constant 0xFF (sel 0x0D) or 0x00 (sel 0x0C)
• OR the spread bit back in to get -1 / +1 as int8

One helper plus an #if defined(GGML_USE_HIP) / #else. The non-HIP branch is unchanged, so CUDA behaviour is identical.

Measured

1x Radeon AI PRO R9700 (gfx1201), ROCm/TheRock 7.13, Bonsai-27B Q1_0 (3.53 GiB), -p 0 -n 64 -r 5, interleaved against a clean build of the same base commit:

build	tg64
master	29.69 / 29.11 t/s
master + this patch	61.97 / 61.34 t/s

mean 29.402 -> 61.657 t/s (+109.7%)

Correctness

• Exhaustive: all 65536 possible q values, both paths compiled for gfx1201, all 4 nibbles each -- bit-identical to the __byte_perm reference. The change cannot alter numerics.
• test-backend-ops -o MUL_MAT: all 11 q1_0 cases pass (n=1..9, k=256 and 4096).

Caveats

• Tested on one model and one GPU (gfx1201). No other RDNA/CDNA parts available to me.
• No NVIDIA hardware, so the CUDA path is unverified by measurement. It is unchanged code inside the #else, but it has not been run.
• Q1_0 still will not beat Q2_0 on this hardware even with this fix -- that appears to be a property of sub-2-bpw formats rather than of any implementation (cf. ik_llama.cpp#45 on RTX-4080, and ggml-vulkan: Add support for TQ1_0 and TQ2_0 types for MUL_MAT #19743's RTX 4090 Vulkan numbers). This removes an AMD-specific software-lowering penalty; it does not change the format's ceiling.

Additional information

Context and further measurements in #27127, including a note on why the figure there was originally reported as +16% (that was measured against a fork where other changes already masked most of the loss, not against master).

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: YES. The unpack algorithm itself is my own prior work from my fork, not AI-generated. An AI agent (Claude) was used to: port that helper onto current master and adapt it to the surrounding upstream code; write the exhaustive 65536-value equivalence harness; build both the baseline and patched trees and run the interleaved benchmarks and test-backend-ops; and draft this description. I have reviewed the diff and am responsible for it.

[ggml-gh-bot]

Hi @The-Monk, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

• PR Template not respected: Please respect the template when creating a new pull request. Make sure to fill out all required sections.

---

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

[dfriehs]

Please test against master with #25628 reverted, for nvidia the previous seemingly naive unpacking resulted in extremely efficient code. You might also want to test applying the same changes to mmq_load_tiles.cuh and see if it improves prefill as well.

[dfriehs]

Can these two lines be a single __builtin_amdgcn_perm? I think it should be possible to directly select from 0x00/0x01 to 0xFF/0x01 via __builtin_amdgcn_perm(0x01FF, 0x01FF, spread).