Commit 1ddffe9f7 for llama.cpp
commit 1ddffe9f70a63e2d260d4bfad9ad200009007f62
Author: Arseniy Obolenskiy <good.doog.2012@gmail.com>
Date: Sat Oct 24 06:04:05 2026 +0200
HIP: add byte permute fallback when __builtin_amdgcn_perm is unavailable (#30915)
* HIP: add byte permute fallback when __builtin_amdgcn_perm is unavailable
amdgcnspirv rejects target-specific builtins unless they are guarded by `__builtin_amdgcn_is_invocable`, so 66 of 145 HIP files failed to build
Wrap the perm calls in a helper that keeps v_perm_b32 when the target has it and falls back to a portable byte select otherwise
* move perm helper to vendors/hip.h
diff --git a/ggml/src/ggml-cuda/mmq-load-tiles.cuh b/ggml/src/ggml-cuda/mmq-load-tiles.cuh
index b37d88c16..2c7f4cd0a 100644
--- a/ggml/src/ggml-cuda/mmq-load-tiles.cuh
+++ b/ggml/src/ggml-cuda/mmq-load-tiles.cuh
@@ -60,10 +60,10 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
const uint32_t q2_indices = (q2_bits & 0x01) | ((q2_bits & 0x02) << 7) | ((q2_bits & 0x04) << 14) | ((q2_bits & 0x08) << 21);
const uint32_t q3_indices = (q3_bits & 0x01) | ((q3_bits & 0x02) << 7) | ((q3_bits & 0x04) << 14) | ((q3_bits & 0x08) << 21);
- const int v0 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q0_indices);
- const int v1 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q1_indices);
- const int v2 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q2_indices);
- const int v3 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q3_indices);
+ const int v0 = ggml_hip_perm(0x000001FF, 0x000001FF, q0_indices);
+ const int v1 = ggml_hip_perm(0x000001FF, 0x000001FF, q1_indices);
+ const int v2 = ggml_hip_perm(0x000001FF, 0x000001FF, q2_indices);
+ const int v3 = ggml_hip_perm(0x000001FF, 0x000001FF, q3_indices);
#else
// unpack crumbs into nibble indices
const int n0 = __byte_perm(0x11100100, 0x11100100, q >> 0); // [0, 1, 4, 5] [ 8, 9, 12, 13]
@@ -162,8 +162,8 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
const uint32_t qx_indices = (q & 0x03) | ((q & 0x0C) << 6) | ((q & 0x30) << 12) | ((q & 0xC0) << 18);
const uint32_t qy_bits = q >> 8;
const uint32_t qy_indices = (qy_bits & 0x03) | ((qy_bits & 0x0C) << 6) | ((qy_bits & 0x30) << 12) | ((qy_bits & 0xC0) << 18);
- const int qx = __builtin_amdgcn_perm(0x020100FF, 0x020100FF, qx_indices);
- const int qy = __builtin_amdgcn_perm(0x020100FF, 0x020100FF, qy_indices);
+ const int qx = ggml_hip_perm(0x020100FF, 0x020100FF, qx_indices);
+ const int qy = ggml_hip_perm(0x020100FF, 0x020100FF, qy_indices);
#else
// unpack even and odd crumbs into byte values
const int qe = __byte_perm(0x020100FF, 0x020100FF, q >> 0);
diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh
index c7acdacaf..d5ff93b91 100644
--- a/ggml/src/ggml-cuda/vecdotq.cuh
+++ b/ggml/src/ggml-cuda/vecdotq.cuh
@@ -40,18 +40,18 @@ static __device__ __forceinline__ int2 get_int_from_table_16(const int & q4, con
const uint32_t q_odd = (q4 >> 4);
// Perform lookups in the lower half of the table (indices 0-7).
- uint32_t v_even_low = __builtin_amdgcn_perm(values[1], values[0], q_even & 0x07070707);
- uint32_t v_odd_low = __builtin_amdgcn_perm(values[1], values[0], q_odd & 0x07070707);
+ uint32_t v_even_low = ggml_hip_perm(values[1], values[0], q_even & 0x07070707);
+ uint32_t v_odd_low = ggml_hip_perm(values[1], values[0], q_odd & 0x07070707);
// Perform lookups in the upper half of the table (indices 8-15).
- uint32_t v_even_high = __builtin_amdgcn_perm(values[3], values[2], q_even & 0x07070707);
- uint32_t v_odd_high = __builtin_amdgcn_perm(values[3], values[2], q_odd & 0x07070707);
+ uint32_t v_even_high = ggml_hip_perm(values[3], values[2], q_even & 0x07070707);
+ uint32_t v_odd_high = ggml_hip_perm(values[3], values[2], q_odd & 0x07070707);
// Select between the low and high results based on the MSB of each index nibble.
uint32_t mask_even = 0x03020100 | ((q_even & 0x08080808) >> 1);
- uint32_t res_x = __builtin_amdgcn_perm(v_even_high, v_even_low, mask_even);
+ uint32_t res_x = ggml_hip_perm(v_even_high, v_even_low, mask_even);
uint32_t mask_odd = 0x03020100 | ((q_odd & 0x08080808) >> 1);
- uint32_t res_y = __builtin_amdgcn_perm(v_odd_high, v_odd_low, mask_odd);
+ uint32_t res_y = ggml_hip_perm(v_odd_high, v_odd_low, mask_odd);
return make_int2(res_x, res_y);
#elif !defined(GGML_USE_MUSA)
@@ -720,10 +720,10 @@ static __device__ __forceinline__ float vec_dot_q1_0_q8_1(
const uint32_t q2_indices = (q2_bits & 0x01) | ((q2_bits & 0x02) << 7) | ((q2_bits & 0x04) << 14) | ((q2_bits & 0x08) << 21);
const uint32_t q3_indices = (q3_bits & 0x01) | ((q3_bits & 0x02) << 7) | ((q3_bits & 0x04) << 14) | ((q3_bits & 0x08) << 21);
- const int v0 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q0_indices);
- const int v1 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q1_indices);
- const int v2 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q2_indices);
- const int v3 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q3_indices);
+ const int v0 = ggml_hip_perm(0x000001FF, 0x000001FF, q0_indices);
+ const int v1 = ggml_hip_perm(0x000001FF, 0x000001FF, q1_indices);
+ const int v2 = ggml_hip_perm(0x000001FF, 0x000001FF, q2_indices);
+ const int v3 = ggml_hip_perm(0x000001FF, 0x000001FF, q3_indices);
#else
// unpack crumbs into nibble indices
const int n0 = __byte_perm(0x11100100, 0x11100100, q >> 0); // [0, 1, 4, 5] [ 8, 9, 12, 13]
@@ -777,8 +777,8 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1(
const uint32_t qx_indices = (q & 0x03) | ((q & 0x0C) << 6) | ((q & 0x30) << 12) | ((q & 0xC0) << 18);
const uint32_t qy_bits = q >> 8;
const uint32_t qy_indices = (qy_bits & 0x03) | ((qy_bits & 0x0C) << 6) | ((qy_bits & 0x30) << 12) | ((qy_bits & 0xC0) << 18);
- const int qx = __builtin_amdgcn_perm(0x020100FF, 0x020100FF, qx_indices);
- const int qy = __builtin_amdgcn_perm(0x020100FF, 0x020100FF, qy_indices);
+ const int qx = ggml_hip_perm(0x020100FF, 0x020100FF, qx_indices);
+ const int qy = ggml_hip_perm(0x020100FF, 0x020100FF, qy_indices);
#else
// unpack even and odd crumbs into byte values
const int qe = __byte_perm(0x020100FF, 0x020100FF, q >> 0);
diff --git a/ggml/src/ggml-cuda/vendors/hip.h b/ggml/src/ggml-cuda/vendors/hip.h
index da7e3da0b..8e8f5a4ee 100644
--- a/ggml/src/ggml-cuda/vendors/hip.h
+++ b/ggml/src/ggml-cuda/vendors/hip.h
@@ -311,3 +311,22 @@ static __device__ __forceinline__ unsigned int __vcmpne4(unsigned int a, unsigne
return (ne_any_bit >> 7) * 0xff;
}
+
+// v_perm_b32 with selector bytes in [0, 7]: 0-3 pick bytes of lo, 4-7 pick bytes of hi.
+static __device__ __forceinline__ uint32_t ggml_hip_perm(const uint32_t hi, const uint32_t lo, const uint32_t sel) {
+#if __has_builtin(__builtin_amdgcn_is_invocable)
+ if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_perm)) {
+ return __builtin_amdgcn_perm(hi, lo, sel);
+ }
+ const uint64_t src = ((uint64_t) hi << 32) | lo;
+ uint32_t res = 0;
+#pragma unroll
+ for (int i = 0; i < 4; ++i) {
+ const uint32_t byte_sel = (sel >> (8 * i)) & 0x07;
+ res |= (uint32_t) ((src >> (8 * byte_sel)) & 0xFF) << (8 * i);
+ }
+ return res;
+#else
+ return __builtin_amdgcn_perm(hi, lo, sel);
+#endif // __has_builtin(__builtin_amdgcn_is_invocable)
+}