nicolas
/
ik_llama_opt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

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.
This commit is contained in:
Yap Sok Ann 2026-08-27 13:40:14 +07:00 committed by GitHub
parent ef40550042
commit 6b261d771e
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
1 changed files with 2 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
ggml/src/ggml-cuda/template-instances/mmvq-instance-iq3_k.cu
View File

@ -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};