diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index 0186b492..9f2b2300 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -1253,10 +1253,10 @@ static __device__ __forceinline__ float vec_dot_iq4_k_q8_1( const block_iq4_k * bq4 = (const block_iq4_k *) vbq + kbx; const uint8_t * all_values = (const uint8_t *)iq4k_values; - // iqs is 0...7 - const int ib32 = iqs; + // iqs is 0...28 + const int ib32 = iqs/4; // Why iqs/4 ? - const int32_t * q8 = (const int *)bq8_1[iqs/4].qs; + const int32_t * q8 = (const int *)bq8_1[ib32].qs; const uint16_t * q4 = (const uint16_t *)bq4->qs + 8*ib32; const uint16_t extra = bq4->extra >> 2*ib32; int v1, v2; @@ -1267,7 +1267,7 @@ static __device__ __forceinline__ float vec_dot_iq4_k_q8_1( sumi1 = ggml_cuda_dp4a(v1, q8[j+0], sumi1); sumi2 = ggml_cuda_dp4a(v2, q8[j+4], sumi2); } - const float d = __half2float(bq4->d) * __low2float(bq8_1[iqs/4].ds); + const float d = __half2float(bq4->d) * __low2float(bq8_1[ib32].ds); const uint8_t sh = bq4->scales_h[ib32/2] >> 4*(ib32%2); const int ls1 = ((bq4->scales_l[ib32] & 0xf) | ((sh << 4) & 0x30)) - 32; const int ls2 = ((bq4->scales_l[ib32] >> 4) | ((sh << 2) & 0x30)) - 32;