mozempk/ggml-org-llama-cpp-mirror
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-10-10 23:10:48 +02:00
Code Issues Packages Projects Releases Wiki Activity

ggml-cuda: HIP: optimize packed byte subtraction (__vsubss4 -> __vsub4) (#29478)

Browse Source 
* ggml-cuda: HIP: optimize non-saturating packed byte subtraction (`__vsubss4`)

* CI: ignore 1 spilled vgpr in fattn_vec

---------

Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz>
This commit is contained in:
thelittlefireman
2026-09-29 19:59:19 +02:00
committed by GitHub
co-authored by Carl Philipp Klemm
parent cee37ffea0
commit 748d4225b9
4 changed files with 20 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/fattn-common.cuh
+7 -2
View File
@@ -416,9 +416,14 @@ static __device__ __forceinline__ void dequantize_V_q4_0(const void * __restrict
int q;
static_assert(ne == 2 || ne == 4, "bad ne");
ggml_cuda_memcpy_1<ne, 2>(&q, x[ib].qs + iqs);
#if defined(GGML_USE_HIP)
// Keep this VMEM read close to its packed-byte dequantization. Hoisting it too far
// increases VGPR pressure substantially in some FlashAttention vector kernels.
__builtin_amdgcn_sched_group_barrier(0x20, 1, 0);
#endif // defined(GGML_USE_HIP)
q >>= 4*shift;
q &= 0x0F0F0F0F;
q = __vsubss4(q, 0x08080808);
q = __vsub4(q, 0x08080808);
const int8_t * q8 = (const int8_t *) &q;
@@ -508,7 +513,7 @@ static __device__ __forceinline__ void dequantize_V_q5_0(const void * __restrict
}
}
q = __vsubss4(q, 0x10101010);
q = __vsub4(q, 0x10101010);
const int8_t * q8 = (const int8_t *) &q;
ggml/src/ggml-cuda/mmq-load-tiles.cuh
+10 -10
View File
@@ -218,8 +218,8 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
const int qs0 = get_int_b2(bxi->qs, kqsx);
#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
x_qs[i*sram_stride + kbx*(2*QI4_0) + kqsx + 0] = __vsubss4((qs0 >> 0) & 0x0F0F0F0F, 0x08080808);
x_qs[i*sram_stride + kbx*(2*QI4_0) + kqsx + QI4_0] = __vsubss4((qs0 >> 4) & 0x0F0F0F0F, 0x08080808);
x_qs[i*sram_stride + kbx*(2*QI4_0) + kqsx + 0] = __vsub4((qs0 >> 0) & 0x0F0F0F0F, 0x08080808);
x_qs[i*sram_stride + kbx*(2*QI4_0) + kqsx + QI4_0] = __vsub4((qs0 >> 4) & 0x0F0F0F0F, 0x08080808);
#else
x_qs[i*(MMQ_TILE_NE_K + 1) + txi] = qs0;
#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE)
@@ -350,14 +350,14 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
qs0 |= (qh << 11) & 0x00001000; // 1 -> 12
qs0 |= (qh << 18) & 0x00100000; // 2 -> 20
qs0 |= (qh << 25) & 0x10000000; // 3 -> 28
qs0 = __vsubss4(qs0, 0x10101010); // subtract 16
qs0 = __vsub4(qs0, 0x10101010); // subtract 16
int qs1 = (ql >> 4) & 0x0F0F0F0F;
qs1 |= (qh >> 12) & 0x00000010; // 16 -> 4
qs1 |= (qh >> 5) & 0x00001000; // 17 -> 12
qs1 |= (qh << 2) & 0x00100000; // 18 -> 20
qs1 |= (qh << 9) & 0x10000000; // 19 -> 28
qs1 = __vsubss4(qs1, 0x10101010); // subtract 16
qs1 = __vsub4(qs1, 0x10101010); // subtract 16
#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
x_qs[i*sram_stride + kbx*(2*QI5_0) + kqsx + 0] = qs0;
@@ -636,7 +636,7 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
const int x_ql_k = (x_ql_0 >> (2*l)) & 0x03030303;
const int x_qh_k = ((x_qh_0 >> l) << 2) & 0x04040404;
const int x_qs_k = __vsubss4(x_ql_k | x_qh_k, 0x04040404);
const int x_qs_k = __vsub4(x_ql_k | x_qh_k, 0x04040404);
#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
x_qs[i*sram_stride + k] = x_qs_k;
@@ -667,7 +667,7 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
const int shift_high = 2 * ksc;
const int sc_high = ((get_int_b2(bxi->scales, ksc_high) >> shift_high) << 4) & 0x30303030;
const int sc = __vsubss4(sc_low | sc_high, 0x20202020);
const int sc = __vsub4(sc_low | sc_high, 0x20202020);
#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
const int8_t * sc8 = (const int8_t *) &sc;
@@ -987,11 +987,11 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
const int kq1 = 2*txi - txi % (QI6_K/2) + QI6_K/2;
#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
x_qs[i*sram_stride + kq0] = __vsubss4(ql0 | qh0, 0x20202020);
x_qs[i*sram_stride + kq1] = __vsubss4(ql1 | qh1, 0x20202020);
x_qs[i*sram_stride + kq0] = __vsub4(ql0 | qh0, 0x20202020);
x_qs[i*sram_stride + kq1] = __vsub4(ql1 | qh1, 0x20202020);
#else
x_qs[i*(2*MMQ_TILE_NE_K + 1) + kq0] = __vsubss4(ql0 | qh0, 0x20202020);
x_qs[i*(2*MMQ_TILE_NE_K + 1) + kq1] = __vsubss4(ql1 | qh1, 0x20202020);
x_qs[i*(2*MMQ_TILE_NE_K + 1) + kq0] = __vsub4(ql0 | qh0, 0x20202020);
x_qs[i*(2*MMQ_TILE_NE_K + 1) + kq1] = __vsub4(ql1 | qh1, 0x20202020);
#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
}
ggml/src/ggml-cuda/vecdotq.cuh
+2 -2
View File
@@ -471,7 +471,7 @@ static __device__ __forceinline__ float vec_dot_q3_K_q8_1_impl_mmvq(
const int vih = ((vh >> i) << 2) & 0x04040404;
const int vi = __vsubss4(vil, vih);
const int vi = __vsub4(vil, vih);
sumf += d8[i] * (ggml_cuda_dp4a(vi, u[i], 0) * sc); // SIMD dot product
}
@@ -638,7 +638,7 @@ static __device__ __forceinline__ float vec_dot_q6_K_q8_1_impl_mmvq(
const int vih = ((vh >> (4*i)) << 4) & 0x30303030;
const int vi = __vsubss4((vil | vih), 0x20202020); // vi = (vil | vih) - 32
const int vi = __vsub4((vil | vih), 0x20202020); // vi = (vil | vih) - 32
sumf += d8[i] * (ggml_cuda_dp4a(vi, u[i], 0) * sc); // SIMD dot product
}
scripts/hip/gcn-cdna-vgpr-check.py
+1
View File
@@ -82,6 +82,7 @@ def main():
'_ZL18flash_attn_ext_f16ILi576ELi512ELi1ELi32ELb0ELb1ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil',
'_ZL18flash_attn_ext_f16ILi576ELi512ELi2ELi16ELb0ELb1ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil',
'_ZL18flash_attn_ext_f16ILi576ELi512ELi8ELi4ELb0ELb1ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil',
'_ZL18flash_attn_ext_vecILi256ELi1EL9ggml_type2ELS0_2ELb0EEvPKcS2_S2_S2_S2_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS6_IjLj3EEiiiiiiiiiiiliiliiiiil',
}
functions = parse_log_file(log_file)