From 6b261d771ee8d5b810749c54d411c482251c18ef Mon Sep 17 00:00:00 2001 From: Yap Sok Ann Date: Thu, 27 Aug 2026 13:40:14 +0700 Subject: [PATCH] cuda: fix iq3_k value table selection in MMVQ (#2364) The per-32-block extra bits select the iq3nl_values half for each 16-element half: sub-block 4g+m, half h uses bit 8g+2m+h. The four streams need local bits 0, 2, 4, 6, not 0, 1, 2, 3. --- ggml/src/ggml-cuda/template-instances/mmvq-instance-iq3_k.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/template-instances/mmvq-instance-iq3_k.cu b/ggml/src/ggml-cuda/template-instances/mmvq-instance-iq3_k.cu index e01df5f9..43516ed4 100644 --- a/ggml/src/ggml-cuda/template-instances/mmvq-instance-iq3_k.cu +++ b/ggml/src/ggml-cuda/template-instances/mmvq-instance-iq3_k.cu @@ -21,8 +21,8 @@ __device__ __forceinline__ void vec_dot_iq3_k_q8_1( const uint8_t extra = bq3->extra >> (8*ib128 + il8/2); 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); + uint32_t extra32_1 = (extra32 << 3) & 0x88888888; + uint32_t extra32_2 = (extra32 << 1) & 0x88888888; const int * q8; int sumi[4] = {0, 0, 0, 0};