Skip to content

Commit 9d1878a

Browse files
thelittlefiremanthelittlefireman
authored
ggml-cuda: HIP: use __builtin_amdgcn_perm for q1_0 unpacking (ggml-org#29927)
1 parent c0fd786 commit 9d1878a

2 files changed

Lines changed: 40 additions & 0 deletions

File tree

‎ggml/src/ggml-cuda/mmq-load-tiles.cuh‎

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -46,6 +46,25 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
4646
for (int j = 0; j < 2; ++j) {
4747
const int q = qxi[j];
4848

49+
#if defined(GGML_USE_HIP)
50+
const uint32_t q_bits = (uint16_t) q;
51+
52+
// Expand each group of four Q1_0 bits into byte selectors for v_perm_b32.
53+
const uint32_t q0_bits = q_bits >> 0;
54+
const uint32_t q1_bits = q_bits >> 4;
55+
const uint32_t q2_bits = q_bits >> 8;
56+
const uint32_t q3_bits = q_bits >> 12;
57+
58+
const uint32_t q0_indices = (q0_bits & 0x01) | ((q0_bits & 0x02) << 7) | ((q0_bits & 0x04) << 14) | ((q0_bits & 0x08) << 21);
59+
const uint32_t q1_indices = (q1_bits & 0x01) | ((q1_bits & 0x02) << 7) | ((q1_bits & 0x04) << 14) | ((q1_bits & 0x08) << 21);
60+
const uint32_t q2_indices = (q2_bits & 0x01) | ((q2_bits & 0x02) << 7) | ((q2_bits & 0x04) << 14) | ((q2_bits & 0x08) << 21);
61+
const uint32_t q3_indices = (q3_bits & 0x01) | ((q3_bits & 0x02) << 7) | ((q3_bits & 0x04) << 14) | ((q3_bits & 0x08) << 21);
62+
63+
const int v0 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q0_indices);
64+
const int v1 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q1_indices);
65+
const int v2 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q2_indices);
66+
const int v3 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q3_indices);
67+
#else
4968
// unpack crumbs into nibble indices
5069
const int n0 = __byte_perm(0x11100100, 0x11100100, q >> 0); // [0, 1, 4, 5] [ 8, 9, 12, 13]
5170
const int n1 = __byte_perm(0x11100100, 0x11100100, q >> 2); // [2, 3, 6, 7] [10, 11, 14, 15]
@@ -59,6 +78,7 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
5978
const int v1 = __byte_perm(s0, s1, 0x7632);
6079
const int v2 = __byte_perm(s2, s3, 0x5410);
6180
const int v3 = __byte_perm(s2, s3, 0x7632);
81+
#endif // defined(GGML_USE_HIP)
6282

6383
#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
6484
x_qs[i*sram_stride + dst_offset + j*4+0] = v0;

‎ggml/src/ggml-cuda/vecdotq.cuh‎

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -706,6 +706,25 @@ static __device__ __forceinline__ float vec_dot_q1_0_q8_1(
706706
const int u2 = get_int_b4(bq8_1_chunk->qs, j*4+2);
707707
const int u3 = get_int_b4(bq8_1_chunk->qs, j*4+3);
708708

709+
#if defined(GGML_USE_HIP)
710+
const uint32_t q_bits = (uint16_t) q;
711+
712+
// Expand each group of four Q1_0 bits into byte selectors for v_perm_b32.
713+
const uint32_t q0_bits = q_bits >> 0;
714+
const uint32_t q1_bits = q_bits >> 4;
715+
const uint32_t q2_bits = q_bits >> 8;
716+
const uint32_t q3_bits = q_bits >> 12;
717+
718+
const uint32_t q0_indices = (q0_bits & 0x01) | ((q0_bits & 0x02) << 7) | ((q0_bits & 0x04) << 14) | ((q0_bits & 0x08) << 21);
719+
const uint32_t q1_indices = (q1_bits & 0x01) | ((q1_bits & 0x02) << 7) | ((q1_bits & 0x04) << 14) | ((q1_bits & 0x08) << 21);
720+
const uint32_t q2_indices = (q2_bits & 0x01) | ((q2_bits & 0x02) << 7) | ((q2_bits & 0x04) << 14) | ((q2_bits & 0x08) << 21);
721+
const uint32_t q3_indices = (q3_bits & 0x01) | ((q3_bits & 0x02) << 7) | ((q3_bits & 0x04) << 14) | ((q3_bits & 0x08) << 21);
722+
723+
const int v0 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q0_indices);
724+
const int v1 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q1_indices);
725+
const int v2 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q2_indices);
726+
const int v3 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q3_indices);
727+
#else
709728
// unpack crumbs into nibble indices
710729
const int n0 = __byte_perm(0x11100100, 0x11100100, q >> 0); // [0, 1, 4, 5] [ 8, 9, 12, 13]
711730
const int n1 = __byte_perm(0x11100100, 0x11100100, q >> 2); // [2, 3, 6, 7] [10, 11, 14, 15]
@@ -719,6 +738,7 @@ static __device__ __forceinline__ float vec_dot_q1_0_q8_1(
719738
const int v1 = __byte_perm(s0, s1, 0x7632);
720739
const int v2 = __byte_perm(s2, s3, 0x5410);
721740
const int v3 = __byte_perm(s2, s3, 0x7632);
741+
#endif // defined(GGML_USE_HIP)
722742

723743
sumi = ggml_cuda_dp4a(v0, u0, sumi);
724744
sumi = ggml_cuda_dp4a(v1, u1, sumi);

0 commit comments

Comments
 (0)