diff --git a/ggml/src/ggml-cuda/iqk_mmvq.cu b/ggml/src/ggml-cuda/iqk_mmvq.cu index 3620a2bc..4a115fe7 100644 --- a/ggml/src/ggml-cuda/iqk_mmvq.cu +++ b/ggml/src/ggml-cuda/iqk_mmvq.cu @@ -1014,41 +1014,32 @@ __device__ __forceinline__ void vec_dot_iq3_k_q8_1( const uint16_t sh = bq3->scales_h >> (8*ib128 + il8/2); const uint8_t extra = bq3->extra >> (8*ib128 + il8/2); - const uint16_t * values1 = iq3k_table + ((extra << 6) & 0x40); - const uint16_t * values2 = iq3k_table + ((extra << 5) & 0x40); - const uint16_t * values3 = iq3k_table + ((extra << 4) & 0x40); - const uint16_t * values4 = iq3k_table + ((extra << 3) & 0x40); + uint32_t extra32 = uint32_t(extra) * 0x01010101; + uint32_t extra32_1 = ((extra32 << 3) & 0x08080808) | ((extra32 << 5) & 0x80808080); + uint32_t extra32_2 = ((extra32 << 2) & 0x08080808) | ((extra32 << 4) & 0x80808080); const int * q8; int sumi[4] = {0, 0, 0, 0}; - int v; for (int i = 0; i < 2; ++i) { uint32_t vl = ql[2*i+0] | (ql[2*i+1] << 16); - uint32_t vh = ((qh[2*i+0] | (qh[2*i+1] << 16)) << hshift) >> 2; + uint32_t vh = ((qh[2*i+0] | (qh[2*i+1] << 16)) << hshift); + + uint32_t val1 = ((vl >> 0) & 0x33333333) | extra32_1 | ((vh >> 2) & 0x04040404) | ((vh >> 0) & 0x40404040); + uint32_t val2 = ((vl >> 2) & 0x33333333) | extra32_2 | ((vh >> 3) & 0x04040404) | ((vh >> 1) & 0x40404040); + int2 v1 = get_int_from_table_16(val1, iq3nl_values); + int2 v2 = get_int_from_table_16(val2, iq3nl_values); q8 = (const int *)bq8_1[4*ib128+0].qs + 2*il8; - aux32 = (vl & 0x03030303) | (vh & 0x04040404); - v = int_from_table_2(aux8, values1); - sumi[0] = ggml_cuda_dp4a(v, q8[i], sumi[0]); - vl >>= 2; vh >>= 1; + sumi[0] = ggml_cuda_dp4a(v1.x, q8[i], sumi[0]); q8 += sizeof(block_q8_1)/4; - aux32 = (vl & 0x03030303) | (vh & 0x04040404); - v = int_from_table_2(aux8, values2); - sumi[1] = ggml_cuda_dp4a(v, q8[i], sumi[1]); - vl >>= 2; vh >>= 1; + sumi[1] = ggml_cuda_dp4a(v2.x, q8[i], sumi[1]); q8 += sizeof(block_q8_1)/4; - aux32 = (vl & 0x03030303) | (vh & 0x04040404); - v = int_from_table_2(aux8, values3); - sumi[2] = ggml_cuda_dp4a(v, q8[i], sumi[2]); - vl >>= 2; vh >>= 1; + sumi[2] = ggml_cuda_dp4a(v1.y, q8[i], sumi[2]); q8 += sizeof(block_q8_1)/4; - aux32 = (vl & 0x03030303) | (vh & 0x04040404); - v = int_from_table_2(aux8, values4); - sumi[3] = ggml_cuda_dp4a(v, q8[i], sumi[3]); - + sumi[3] = ggml_cuda_dp4a(v2.y, q8[i], sumi[3]); } const float d = __half2float(bq3->d); const uint16_t * sl16 = (const uint16_t *)bq3->scales_l + 2*ib128;