From f46bc30cb6a7f68a67e34a00061e20a4ad1eff43 Mon Sep 17 00:00:00 2001 From: Felix Ye Date: Wed, 23 Sep 2026 04:31:12 +0800 Subject: [PATCH] HIP : optimize IQ2/IQ3 (`__vsub4` `__vcmpne4`) using SWAR (#27962) * HIP : use bit manipulation for __vcmpne4 * HIP : use bit manipulation for __vsub4 --- ggml/src/ggml-cuda/vendors/hip.h | 28 ++++++++++++++++++---------- 1 file changed, 18 insertions(+), 10 deletions(-) diff --git a/ggml/src/ggml-cuda/vendors/hip.h b/ggml/src/ggml-cuda/vendors/hip.h index 48d4eb2ce3..0a2f2829e5 100644 --- a/ggml/src/ggml-cuda/vendors/hip.h +++ b/ggml/src/ggml-cuda/vendors/hip.h @@ -277,7 +277,15 @@ static __device__ __forceinline__ int __vsubss4(const int a, const int b) { } static __device__ __forceinline__ int __vsub4(const int a, const int b) { - return __vsubss4(a, b); + // do some small modifications to a and b to make the subtraction not underflow + const unsigned int a_large = a | 0x80808080; + const unsigned int b_small = b & 0x7f7f7f7f; + const unsigned int result_low_7bits = a_large - b_small; + + // if two ops share the same high bit, we should flip the high bit of the result + const unsigned int should_flip_high_1bit = (a ^ ~b) & 0x80808080; + + return result_low_7bits ^ should_flip_high_1bit; } static __device__ __forceinline__ unsigned int __vcmpeq4(unsigned int a, unsigned int b) { @@ -293,13 +301,13 @@ static __device__ __forceinline__ unsigned int __vcmpeq4(unsigned int a, unsigne } static __device__ __forceinline__ unsigned int __vcmpne4(unsigned int a, unsigned int b) { - const uint8x4_t& va = reinterpret_cast(a); - const uint8x4_t& vb = reinterpret_cast(b); - unsigned int c; - uint8x4_t& vc = reinterpret_cast(c); -#pragma unroll - for (int i = 0; i < 4; ++i) { - vc[i] = va[i] == vb[i] ? 0x00 : 0xff; - } - return c; + const unsigned int x = a ^ b; + + // any non-equal bit in a byte will set the high bit of that byte here + // the addition will not overflow in the byte as op1 and op2 are both less than 0x80 + const unsigned int ne_low_7bits = ((x & 0x7f7f7f7f) + 0x7f7f7f7f) & 0x80808080; + const unsigned int ne_high_1bit = x & 0x80808080; + const unsigned int ne_any_bit = ne_low_7bits | ne_high_1bit; + + return (ne_any_bit >> 7) * 0xff; }