From 9cf9172afeeb05dc5dc1cba49fdd55e686ecdd27 Mon Sep 17 00:00:00 2001 From: Iwan Kawrakow Date: Thu, 21 Aug 2025 18:44:57 +0300 Subject: [PATCH] Use bperm trick for iq2_k gemv -> ~3% gain --- ggml/src/ggml-cuda/iqk_mmvq.cu | 36 +++++++++++++++++++++++++++------- 1 file changed, 29 insertions(+), 7 deletions(-) diff --git a/ggml/src/ggml-cuda/iqk_mmvq.cu b/ggml/src/ggml-cuda/iqk_mmvq.cu index 78533b07..689d2fe2 100644 --- a/ggml/src/ggml-cuda/iqk_mmvq.cu +++ b/ggml/src/ggml-cuda/iqk_mmvq.cu @@ -849,6 +849,34 @@ __device__ __forceinline__ void vec_dot_iq2_k_q8_1( const uint32_t * q2 = (const uint32_t *)bq2->qs + 8*(i4/4) + 2*(i4%4); const uint16_t extra = bq2->extra >> (8*(i4/4) + (i4%4)/2); + const uint32_t * scales = (const uint32_t *)bq2->scales; + uint32_t s32 = __vsub4((scales[i4/4] >> 4*(((i4%4)/2)%2)) & 0x0f0f0f0f, 0x08080808); + const int8_t * s8 = (const int8_t *)&s32; + + // Block of 16: (32*(4*(i4/4)+k)+8*(i4%4))/16 = 8*(i4/4) + 2*k + (i4%4)/2 + // -> scales_l[4*(i4/4) + k] >> 4*(((i4%4)/2)%2) + +#ifdef __CUDA_ARCH__ + uint32_t extra32 = uint32_t(extra & 0xff) * 0x01010101; + uint32_t extra32_1 = (extra32 << 2) & 0x44444444; + uint32_t extra32_2 = (extra32 << 0) & 0x44444444; + + uint32_t val1, val2; + + val1 = ((q2[0] >> 0) & 0x33333333) | extra32_1; val2 = ((q2[1] >> 0) & 0x33333333) | extra32_1; + int2 v1 = get_int_from_table_8(val1, iq2nl_values); + int2 v2 = get_int_from_table_8(val2, iq2nl_values); + int sumi1 = ggml_cuda_dp4a(v2.x, q8_1[1], ggml_cuda_dp4a(v1.x, q8_1[0], 0)) * s8[0]; + int sumi3 = ggml_cuda_dp4a(v2.y, q8_3[1], ggml_cuda_dp4a(v1.y, q8_3[0], 0)) * s8[2]; + + val1 = ((q2[0] >> 2) & 0x33333333) | extra32_2; val2 = ((q2[1] >> 2) & 0x33333333) | extra32_2; + v1 = get_int_from_table_8(val1, iq2nl_values); + v2 = get_int_from_table_8(val2, iq2nl_values); + int sumi2 = ggml_cuda_dp4a(v2.x, q8_2[1], ggml_cuda_dp4a(v1.x, q8_2[0], 0)) * s8[1]; + int sumi4 = ggml_cuda_dp4a(v2.y, q8_4[1], ggml_cuda_dp4a(v1.y, q8_4[0], 0)) * s8[3]; + +#else + const int * all_values = (const int *)iq2k_table; const int * values; @@ -857,13 +885,6 @@ __device__ __forceinline__ void vec_dot_iq2_k_q8_1( uint32_t aux32[2]; int v1, v2; - // Block of 16: (32*(4*(i4/4)+k)+8*(i4%4))/16 = 8*(i4/4) + 2*k + (i4%4)/2 - // -> scales_l[4*(i4/4) + k] >> 4*(((i4%4)/2)%2) - - const uint32_t * scales = (const uint32_t *)bq2->scales; - uint32_t s32 = __vsub4((scales[i4/4] >> 4*(((i4%4)/2)%2)) & 0x0f0f0f0f, 0x08080808); - const int8_t * s8 = (const int8_t *)&s32; - aux32[0] = ((val1 >> 0) & 0x03030303); aux32[1] = ((val2 >> 0) & 0x03030303); values = all_values + ((extra & 0x01) << 8); v1 = int_from_table_4(aux32[0], values); v2 = int_from_table_4(aux32[1], values); @@ -883,6 +904,7 @@ __device__ __forceinline__ void vec_dot_iq2_k_q8_1( v1 = int_from_table_4(aux32[0], values); v2 = int_from_table_4(aux32[1], values); int sumi4 = ggml_cuda_dp4a(v2, q8_4[1], ggml_cuda_dp4a(v1, q8_4[0], 0)) * s8[3]; +#endif *result += __half2float(bq2->d) * (__low2float(bq8_1[4*(i4/4)+0].ds) * sumi1 + __low2float(bq8_1[4*(i4/4)+1].ds) * sumi2