Fix __byte_perm sign-replicate selector modes on the device path

[The-Monk]

Fix __byte_perm sign-replicate selector modes on the AMD device path

Problem

The device implementation of __byte_perm(x, y, s) in
hipamd/include/hip/amd_detail/amd_device_functions.h ignores the
sign-replicate selector mode. In CUDA (and the PTX prmt.b32 instruction it
mirrors), each output byte is chosen by a 4-bit nibble of the selector s:
the low 3 bits index one of the eight source bytes {x, y}, and bit 3
requests "replicate" mode, in which the most-significant bit of the selected
byte is broadcast across all 8 bits (producing 0x00 or 0xFF). This is the
standard way to do a branchless per-byte sign extension.

The current implementation masks the index with 0x07/0x70 and never looks
at bit 3, so any selector nibble in [0x8, 0xF] silently returns the plain
byte instead of its sign broadcast. Selectors in [0x0, 0x7] (the common
case) are unaffected and remain correct.

This is silent: code ported from CUDA that relies on replicate mode compiles
and runs, but computes wrong results only for those selector values, which is
painful to track down. It is present on clr master (verified at a223bb1)
and in shipping ROCm toolchains.

Semantics target (please confirm)

This PR targets bit-compatibility with CUDA __byte_perm / PTX prmt.b32
(default mode), since that is the portability contract most code assumes.
If AMD prefers to define its own __byte_perm semantics (e.g. document that
HIP only supports plain selection and leave replicate undefined), that is a
reasonable alternative — but it should then be documented explicitly,
because the current behavior looks like an implementation bug rather than a
deliberate semantic choice. Flagging this for maintainer direction.

Evidence

Host reference implementing the CUDA/PTX semantics, compared against the
shipped emulation (test_byte_perm_host.c, run on the selector byte layout,
x=0x33221100, y=0x77665544):

selector s	CUDA/PTX ref	current emulation	note
0x00003210	0x33221100	0x33221100	plain select — OK
0x00001234	0x11223344	0x11223344	plain select — OK
0x0000ba98	0x00000000	0x33221100	replicate — WRONG

Aggregate over 200,000 random (x, y, s):

• selectors masked to plain mode (s & 0x77777777): 0 mismatches
• unmasked (replicate modes included): 186,995 / 200,000 mismatches

Every mismatch involves at least one replicate-mode nibble; no plain-mode
selector is affected.

Fix

Rewrite the byte assembly as a 4-iteration loop over the selector nibbles,
handling the replicate bit. Verified bit-exact against the CUDA/PTX reference
over 2,000,000 random inputs and the exhaustive 16-bit selector space
(0 mismatches).

Test coverage

• test_byte_perm_host.c — host reference vs the old and new emulations
  (proves the old bug and the new fix).
• test_byte_perm_device.hip — device-side check of the built-in against the
  host reference over the full 16-bit selector space; exits non-zero on any
  mismatch. Ready to adapt into the ROCm/hip-tests catch2 layout
  (catch/unit/deviceLib/), see checklist.

Related

• [Documentation]: ROCm Documentation Gap: WMMA Output Lane Mapping for gfx12 (RDNA4) legacy-rocm-build#6025 (comment) documented the HIP-vs-CUDA __byte_perm selector
  divergence for the >= 8 range while debugging RDNA4 kernels.
• ROCm/ROCm discussion #6618 lists this among RDNA4 LLM-inference findings.

[doplxyz]

Not a maintainer of clr or of llama.cpp — an outside contributor with gfx1201 hardware who came at
this from the llama.cpp side. Measurements below, plus one documentation finding that I think
changes the framing. My view, not a verdict.

Short version: the behaviour gap you describe is real, and I reproduced it. But NVIDIA documents
__byte_perm as masking each selector nibble with & 0x7, so I think this change would move HIP
away from the documented CUDA contract rather than toward it — and it measurably changes the
generated code of the Q1_0 and Q2_0 MMQ paths that are in llama.cpp master today.

Patch state reviewed: head 67646d07824e95835dad70b3f7127f2b5e150574 onto
ccbcc2fab6eacec571808c97a45e20f997704204.

1. The behaviour gap is real

RX 9070 XT (gfx1201), AMD clang 23.0.0git (46fcb339 +PATCHED:440716f8), ROCm 7.14: the shipped
device __byte_perm masks the index to 3 bits and ignores nibble bit 3. Over all 65,536 16-bit
selectors with x = 0x03020100, y = 0xF3F2F1F0, shipped and patched differ on 58,975 of them
(90.0% — for that operand pair; the rate depends on the operands).

2. What NVIDIA documents for __byte_perm

This is the part that I think decides the PR, and it is checkable without hardware. The CUDA Math
API reference specifies __byte_perm(x, y, s) as:

Create 8-byte source: uint64_t tmp64 = ((uint64_t)y << 32) | x;
Extract selector bits: selector0 = (s >> 0) & 0x7; selector1 = (s >> 4) & 0x7; selector2 = (s >> 8) & 0x7; selector3 = (s >> 12) & 0x7;
Return 4 selected bytes from 8-byte source: res[07:00] = tmp64[selector0]; ...

— https://docs.nvidia.com/cuda/cuda-math-api/cuda_math_api/group__CUDA__MATH__INTRINSIC__INT.html

Each nibble is masked with & 0x7 in the specification itself. Bit 3 is not part of the contract,
and sign replication is not mentioned anywhere in that entry.

PTX prmt.b32 in its default mode does have the replicate behaviour — but the CUDA C++ intrinsic
and the PTX instruction are not documented as sharing a contract, and here they visibly do not.
The PR body targets "bit-compatibility with CUDA __byte_perm / PTX prmt.b32 (default mode)",
which treats them as one thing.

If that reading is right, then HIP's current implementation already matches the documented CUDA
__byte_perm contract, and this change would introduce a divergence rather than remove one. What
it would genuinely fix is the absence of a PTX-prmt-equivalent in HIP — which seems to me like an
argument for adding an intrinsic that offers those semantics, rather than redefining __byte_perm.

I have no NVIDIA hardware, so I cannot tell you what a given CUDA toolchain actually emits, and a
runtime probe would only establish implementation behaviour rather than the contract either way.
But the documented contract is unambiguous, and it is the thing downstream code is entitled to rely
on.

3. It changes the generated code of llama.cpp's Q1_0 and Q2_0 MMQ paths

At llama.cpp master 95c409c13625a23da2aa37270339ce9179215a18 there are two __byte_perm users:

ggml/src/ggml-cuda/vecdotq.cuh — not affected. get_int_from_table_16 takes a
#if defined(GGML_USE_HIP) branch using __builtin_amdgcn_perm with explicit & 0x07070707
masks; the __byte_perm code is in the #elif branch. AMD never reaches it. (I assumed this one
was at risk before checking. It is not.)

ggml/src/ggml-cuda/mmq-load-tiles.cuh — affected. load_tiles_q1_0 and load_tiles_q2_0
pass runtime quant data as the selector, so nibbles cover 0x0–0xF:

const int n0 = __byte_perm(0x11100100, 0x11100100, q >> 0);   // q1_0
const int qe = __byte_perm(0x020100FF, 0x020100FF, q >> 0);   // q2_0

Rather than argue from #ifdefs, I compiled the MMQ instantiation TUs for gfx1201 twice — once
against the stock SDK header, once with your patch body applied to the real
amd_device_functions.h inside a throwaway container (stock sha256 4c49dd5f…, patched
4095e47c…, one textual replacement, asserted) — and compared the emitted device assembly:

TU	instruction lines, stock → patched	changed lines (of which instructions)
mmq-instance-q1_0.cu	82,692 → 93,916 (+11,224)	47,912 (47,590)
mmq-instance-q2_0.cu	79,373 → 87,032 (+7,659)	34,673 (34,377)
mmq-instance-q4_0.cu	unchanged	0
mmq-instance-q8_0.cu	unchanged	0

q4_0 and q8_0 are negative controls — they show the diff is not uniform build noise. (That the
patch was applied at all is established separately, by the header hash pair and the
one-replacement assertion in the script.)

Per kernel symbol, slicing each symbol from its label to its .Lfunc_end: in q1_0, 12 of 64
symbols changed size, the largest being mul_mat_q<GGML_TYPE_Q1_0, 64, true> at 2,352 → 3,423
instructions (+1,071, +46%). In q2_0, 12 of 64 changed, the largest being
mul_mat_q<GGML_TYPE_Q2_0, 128, true> at 3,080 → 3,778 (+698). (ggml_type 41 and 42 are
GGML_TYPE_Q1_0 and GGML_TYPE_Q2_0, per ggml.h.)

Largest opcode count deltas in the q1_0 TU (patched minus stock):

scratch_load_i8    +2304      s_wait_alu         +1659
scratch_load_u8    -2304      v_cndmask_b32_e64   +999
v_and_b32_e32      +2242      v_cndmask_b32_e32   +991
v_lshrrev_b16      +1728      v_lshlrev_b32_e32   -774

I would not read too much into which opcodes those are. A representative hunk inside the largest
mover is mostly spill/reload churn (scratch_store_b32 ... Folded Spill, s_clause,
scratch_load_b32), so a good part of the delta is register-pressure fallout from inlining a
larger function body rather than the replicate logic itself. The load of the claim here is just
the size and extent of the change, not its internal attribution.

Separately, sweeping all 65,536 possible int16_t quant words through those unpack sequences in
isolation, the two semantics disagree on 36,864 words (56.25%) for q1_0 and 44,800 (68.36%) for
q2_0.

What I have not shown: that a real workload dispatches to these kernels. GGML_TYPE_Q1_0 and
GGML_TYPE_Q2_0 are in the unconditional supported list in ggml_cuda_should_use_mmq, but I have
no Q1_0/Q2_0 GGUF to run. So: instantiated and code-changed, yes; observed at runtime, no. I am
also not claiming today's output is the correct one — that is what §2 is about.

Easy to miss from inside a fork: a tree that has already replaced these __byte_perm calls with
its own unpacks would see none of this.

4. Static codegen difference

gfx1201, -O3, tiny synthetic kernels, counting indented non-directive non-label opcode lines
between each kernel's .globl and its .Lfunc_end:

	current	this PR
compile-time constant selector	18	18
runtime selector	39	62 (4 of them s_cbranch/v_cndmask)

A static count in a synthetic kernel — not a benchmark, and not a claim about dynamic cost,
occupancy or scheduling. It only suggests constant selectors fold either way while runtime ones do
not, which lines up with the +11,224 instructions in §3.

5. v_perm_b32 selector map on gfx1201, if a branchless form is wanted

Measured because I could not find it spelled out, and because my first attempt at it was wrong.
One-hot MSB per source byte, plus all-MSB-clear and all-MSB-set controls, read across all four
destination byte lanes (all four agree; table below is any one of them):

selector value:              0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15
all bytes 0x01 (msb=0)      01 01 01 01 01 01 01 01 00 00 00 00 00 FF FF FF
all bytes 0x81 (msb=1)      81 81 81 81 81 81 81 81 FF FF FF FF 00 FF FF FF
only byte 0 has msb set     80 01 01 01 01 01 01 01 00 00 00 00 00 FF FF FF
only byte 1 has msb set     01 80 01 01 01 01 01 01 FF 00 00 00 00 FF FF FF
only byte 2 has msb set     01 01 80 01 01 01 01 01 00 00 00 00 00 FF FF FF
only byte 3 has msb set     01 01 01 80 01 01 01 01 00 FF 00 00 00 FF FF FF
only byte 4 has msb set     01 01 01 01 80 01 01 01 00 00 00 00 00 FF FF FF
only byte 5 has msb set     01 01 01 01 01 80 01 01 00 00 FF 00 00 FF FF FF
only byte 6 has msb set     01 01 01 01 01 01 80 01 00 00 00 00 00 FF FF FF
only byte 7 has msb set     01 01 01 01 01 01 01 80 00 00 00 FF 00 FF FF FF

Reading it: 0–7 select bytes 0–7 (0–3 from the second operand, 4–7 from the first); 8, 9, 10, 11
replicate the sign of bytes 1, 3, 5, 7 — the odd bytes only; 12 yields constant 0x00 and 13–15
constant 0xFF, both holding under the all-set and all-clear controls. Scope: gfx1201, these
operand patterns, all four destination lanes.

So the hardware has sign-replicate but only for odd byte lanes, and cannot directly express
"replicate the sign of whichever byte this nibble selected" for an arbitrary index. A branchless
implementation still looks possible by computing the sign mask separately; it is just not a
one-instruction substitution. Neither the current nor the patched implementation emits
v_perm_b32 at all in my synthetic kernels, which may be worth something separately.

What I would want before this landed, if it were mine

Reconcile §2 first — whether the target is the documented CUDA __byte_perm contract (which masks
to 3 bits, and which HIP already matches) or PTX prmt.b32 semantics (which HIP lacks). If it is
the latter, a separate intrinsic would give people those semantics without changing what existing
HIP code compiles to. If __byte_perm itself is changed, the mmq-load-tiles.cuh call sites will
need attention upstream too.

Reproducers, with full command lines, container digests, header hashes and the counting method:
https://gist.github.com/doplxyz/eb142625fbc0a1bae1bb4a312e8c85cb/90c10dc1286bd16022b37f437493e0168e2f141d

Happy to run further variants on gfx1201.

[The-Monk]

Thanks for this — reproduction, blast-radius measurement, and the v_perm_b32 lane map in one pass. This is the most thorough look the PR has gotten.

You're right about the docs, and I'll concede the framing point directly: the PR body treats "CUDA __byte_perm" and "PTX prmt.b32 default mode" as one contract, and they aren't. The CUDA Math API entry does specify & 0x7 per nibble, so against the documented contract, HIP's current behavior is already conforming. That part of the PR text is wrong as written.

What motivated the PR is the de facto contract: nvcc lowers __byte_perm to prmt.b32 in default mode, and on NVIDIA hardware selector nibbles in [0x8, 0xF] observably sign-replicate (that's the divergence documented in ROCm/legacy-rocm-build#6025 — found the hard way, debugging ported kernels on RDNA4). So NVIDIA's implementation diverges from NVIDIA's own docs, and real ported code relies on the implementation, not the spec. That's the porting hazard either way: code that computes correctly on NVIDIA silently computes something else here.

Given your doc finding, I agree redefining __byte_perm is probably the wrong lever, and I'd be happy to rework this PR along the lines you suggest:

1. Leave __byte_perm on the documented & 0x7 contract (unchanged behavior, no codegen impact on existing HIP binaries).
2. Add an explicit intrinsic exposing prmt.b32 default-mode semantics (__prmt? open to naming), implemented with the loop body from this PR — your odd-lane-only v_perm_b32 sign-replicate map suggests a branchless lowering is possible too, which we can pursue separately.
3. Document the divergence in the HIP porting guide (and ideally have hipify flag __byte_perm call sites for review).

On the MMQ instruction deltas: useful measurement, but note what it implies. load_tiles_q1_0/load_tiles_q2_0 feed runtime quant nibbles covering 0x0–0xF into the selector, and your own sweep shows the two semantics disagree on 56–68% of quant words — so those call sites cannot be correct under both contracts. Whichever semantics they were written against, one platform is currently miscomputing there, and that needs an upstream llama.cpp audit regardless of which way this PR goes. The +11k instructions is the price of some fix at those sites, not overhead this PR adds to otherwise-correct code.

Maintainers — the "Semantics target (please confirm)" section in the PR body was written for exactly this decision, and there are now measurements on both sides of it. Direction would be appreciated: documented contract + new intrinsic (I'll rework this PR), or bit-compatibility with observed CUDA behavior (PR stands as-is). Happy either way; the silent-wrong-results hazard is what I want closed.