char-fork/llama.cpp

git clone https://git.t4t.associates/char-fork/llama.cpp

Hongqiang Wangopencl: flash attention improvement (#25069)ebd048fc5

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
20.3 KiB708 linesraw
1#pragma OPENCL EXTENSION cl_khr_fp16 : enable
2
3// v = { mp, L, d }
4inline uint fastdiv(uint n, uint4 v) {
5    uint msbs;
6    msbs = mul_hi(n, v.s0);
7    return (msbs + n) >> v.s1;
8}
9inline uint fastmod(uint n, uint4 v) {
10    uint q = fastdiv(n, v);
11    return n - q * v.s2;
12}
13
14kernel void kernel_set_rows_f32_i64(
15        global char * src0,
16        ulong         offset0,
17        global char * src1,
18        ulong         offset1,
19        global char * dst,
20        ulong         offsetd,
21        int           ne01,
22        ulong         nb01,
23        ulong         nb02,
24        ulong         nb03,
25        uint4         ne11,
26        uint4         ne12,
27        ulong         nb10,
28        ulong         nb11,
29        ulong         nb12,
30        int           nblk0,
31        ulong         nb1,
32        ulong         nb2,
33        ulong         nb3
34) {
35    src0 = src0 + offset0;
36    src1 = src1 + offset1;
37    dst  = dst  + offsetd;
38
39    int i03 = get_group_id(2);
40    int i02 = get_group_id(1);
41    int i01 = get_group_id(0)*get_local_size(1) + get_local_id(1);
42
43    if (i01 >= ne01) {
44        return;
45    }
46
47    //int i12 = i03%ne12;
48    //int i11 = i02%ne11;
49    int i12 = fastmod(i03, ne12);
50    int i11 = fastmod(i02, ne11);
51
52    int i10 = i01;
53    long i1 = ((global long *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
54
55    global float * dst_row = (global float *) (dst  +  i1*nb1  + i02*nb2  + i03*nb3);
56    global float * src_row = (global float *) (src0 + i01*nb01 + i02*nb02 + i03*nb03);
57
58    for (int ind = get_local_id(0); ind < nblk0; ind += get_local_size(0)) {
59        dst_row[ind] = (float)src_row[ind];
60    }
61}
62
63kernel void kernel_set_rows_f16_i64(
64        global char * src0,
65        ulong         offset0,
66        global char * src1,
67        ulong         offset1,
68        global char * dst,
69        ulong         offsetd,
70        int           ne01,
71        ulong         nb01,
72        ulong         nb02,
73        ulong         nb03,
74        uint4         ne11,
75        uint4         ne12,
76        ulong         nb10,
77        ulong         nb11,
78        ulong         nb12,
79        int           nblk0,
80        ulong         nb1,
81        ulong         nb2,
82        ulong         nb3
83) {
84    src0 = src0 + offset0;
85    src1 = src1 + offset1;
86    dst  = dst  + offsetd;
87
88    int i03 = get_group_id(2);
89    int i02 = get_group_id(1);
90    int i01 = get_group_id(0)*get_local_size(1) + get_local_id(1);
91
92    if (i01 >= ne01) {
93        return;
94    }
95
96    //int i12 = i03%ne12;
97    //int i11 = i02%ne11;
98    int i12 = fastmod(i03, ne12);
99    int i11 = fastmod(i02, ne11);
100
101    int i10 = i01;
102    long i1 = ((global long *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
103
104    global half  * dst_row = (global half  *) (dst  +  i1*nb1  + i02*nb2  + i03*nb3);
105    global float * src_row = (global float *) (src0 + i01*nb01 + i02*nb02 + i03*nb03);
106
107    for (int ind = get_local_id(0); ind < nblk0; ind += get_local_size(0)) {
108        dst_row[ind] = src_row[ind];
109    }
110}
111
112kernel void kernel_set_rows_f32_i32(
113        global char * src0,
114        ulong         offset0,
115        global char * src1,
116        ulong         offset1,
117        global char * dst,
118        ulong         offsetd,
119        int           ne01,
120        ulong         nb01,
121        ulong         nb02,
122        ulong         nb03,
123        uint4         ne11,
124        uint4         ne12,
125        ulong         nb10,
126        ulong         nb11,
127        ulong         nb12,
128        int           nblk0,
129        ulong         nb1,
130        ulong         nb2,
131        ulong         nb3
132) {
133    src0 = src0 + offset0;
134    src1 = src1 + offset1;
135    dst  = dst  + offsetd;
136
137    int i03 = get_group_id(2);
138    int i02 = get_group_id(1);
139    int i01 = get_group_id(0)*get_local_size(1) + get_local_id(1);
140
141    if (i01 >= ne01) {
142        return;
143    }
144
145    //int i12 = i03%ne12;
146    //int i11 = i02%ne11;
147    int i12 = fastmod(i03, ne12);
148    int i11 = fastmod(i02, ne11);
149
150    int i10 = i01;
151    int i1  = ((global int *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
152
153    global float * dst_row = (global float *) (dst  +  i1*nb1  + i02*nb2  + i03*nb3);
154    global float * src_row = (global float *) (src0 + i01*nb01 + i02*nb02 + i03*nb03);
155
156    for (int ind = get_local_id(0); ind < nblk0; ind += get_local_size(0)) {
157        dst_row[ind] = (float)src_row[ind];
158    }
159}
160
161// f32 -> q8_0 quantize set_rows. Block = half d + char qs[32].
162#define QK8_0 32
163
164inline void quantize_q8_0_block(global float * x, global char * qs, global half * d_out) {
165    float amax = 0.0f;
166    for (int j = 0; j < QK8_0; j++) {
167        amax = fmax(amax, fabs(x[j]));
168    }
169
170    float d  = amax / 127.0f;
171    float id = (d != 0.0f) ? 127.0f / amax : 0.0f;
172
173    vstore_half(d, 0, d_out);
174
175    for (int j = 0; j < QK8_0; j++) {
176        qs[j] = (char)((int)round(x[j] * id));
177    }
178}
179
180kernel void kernel_set_rows_q8_0_i64(
181        global char * src0,
182        ulong         offset0,
183        global char * src1,
184        ulong         offset1,
185        global char * dst,
186        ulong         offsetd,
187        int           ne01,
188        ulong         nb01,
189        ulong         nb02,
190        ulong         nb03,
191        uint4         ne11,
192        uint4         ne12,
193        ulong         nb10,
194        ulong         nb11,
195        ulong         nb12,
196        int           nblk0,
197        ulong         nb1,
198        ulong         nb2,
199        ulong         nb3
200) {
201    src0 = src0 + offset0;
202    src1 = src1 + offset1;
203    dst  = dst  + offsetd;
204
205    int i03 = get_group_id(2);
206    int i02 = get_group_id(1);
207    int i01 = get_group_id(0)*get_local_size(1) + get_local_id(1);
208
209    if (i01 >= ne01) {
210        return;
211    }
212
213    int i12 = fastmod(i03, ne12);
214    int i11 = fastmod(i02, ne11);
215
216    int i10 = i01;
217    long i1 = ((global long *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
218
219    global char  * dst_row = (global char  *) (dst  +  i1*nb1  + i02*nb2  + i03*nb3);
220    global float * src_row = (global float *) (src0 + i01*nb01 + i02*nb02 + i03*nb03);
221
222    for (int blk = get_local_id(0); blk < nblk0; blk += get_local_size(0)) {
223        global float * x = src_row + blk * QK8_0;
224        global char  * y = dst_row + blk * (2 + QK8_0);
225
226        quantize_q8_0_block(x, y + 2, (global half *)y);
227    }
228}
229
230kernel void kernel_set_rows_q8_0_i32(
231        global char * src0,
232        ulong         offset0,
233        global char * src1,
234        ulong         offset1,
235        global char * dst,
236        ulong         offsetd,
237        int           ne01,
238        ulong         nb01,
239        ulong         nb02,
240        ulong         nb03,
241        uint4         ne11,
242        uint4         ne12,
243        ulong         nb10,
244        ulong         nb11,
245        ulong         nb12,
246        int           nblk0,
247        ulong         nb1,
248        ulong         nb2,
249        ulong         nb3
250) {
251    src0 = src0 + offset0;
252    src1 = src1 + offset1;
253    dst  = dst  + offsetd;
254
255    int i03 = get_group_id(2);
256    int i02 = get_group_id(1);
257    int i01 = get_group_id(0)*get_local_size(1) + get_local_id(1);
258
259    if (i01 >= ne01) {
260        return;
261    }
262
263    int i12 = fastmod(i03, ne12);
264    int i11 = fastmod(i02, ne11);
265
266    int i10 = i01;
267    int i1  = ((global int *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
268
269    global char  * dst_row = (global char  *) (dst  +  i1*nb1  + i02*nb2  + i03*nb3);
270    global float * src_row = (global float *) (src0 + i01*nb01 + i02*nb02 + i03*nb03);
271
272    for (int blk = get_local_id(0); blk < nblk0; blk += get_local_size(0)) {
273        global float * x = src_row + blk * QK8_0;
274        global char  * y = dst_row + blk * (2 + QK8_0);
275
276        quantize_q8_0_block(x, y + 2, (global half *)y);
277    }
278}
279
280// SoA q8_0 variants. dst_q: int8[QK8_0] per block; dst_d: fp16 scale per block.
281// Layout matches kernel_convert_block_q8_0; block index follows dst element order.
282kernel void kernel_set_rows_q8_0_soa_i64(
283        global char * src0,
284        ulong         offset0,
285        global char * src1,
286        ulong         offset1,
287        global char * dst_q,
288        ulong         offset_q,
289        global char * dst_d,
290        ulong         offset_d,
291        int           ne01,
292        ulong         nb01,
293        ulong         nb02,
294        ulong         nb03,
295        uint4         ne11,
296        uint4         ne12,
297        ulong         nb10,
298        ulong         nb11,
299        ulong         nb12,
300        int           nblk0,
301        int           ne1_dst,
302        int           ne2_dst,
303        int           ne3_dst
304) {
305    src0  = src0  + offset0;
306    src1  = src1  + offset1;
307    dst_q = dst_q + offset_q;
308    dst_d = dst_d + offset_d;
309
310    int i03 = get_group_id(2);
311    int i02 = get_group_id(1);
312    int i01 = get_group_id(0)*get_local_size(1) + get_local_id(1);
313
314    if (i01 >= ne01) {
315        return;
316    }
317
318    int i12 = fastmod(i03, ne12);
319    int i11 = fastmod(i02, ne11);
320
321    int i10 = i01;
322    long i1 = ((global long *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
323
324    long row_blk_base = ((long)i03 * ne2_dst * ne1_dst + (long)i02 * ne1_dst + i1) * nblk0;
325
326    global half  * d_row = (global half  *)(dst_d) + row_blk_base;
327    global char  * q_row = (global char  *)(dst_q) + row_blk_base * QK8_0;
328    global float * src_row = (global float *)(src0 + i01*nb01 + i02*nb02 + i03*nb03);
329
330    for (int blk = get_local_id(0); blk < nblk0; blk += get_local_size(0)) {
331        global float * x = src_row + blk * QK8_0;
332        global char  * q = q_row + blk * QK8_0;
333
334        quantize_q8_0_block(x, q, d_row + blk);
335    }
336}
337
338kernel void kernel_set_rows_q8_0_soa_i32(
339        global char * src0,
340        ulong         offset0,
341        global char * src1,
342        ulong         offset1,
343        global char * dst_q,
344        ulong         offset_q,
345        global char * dst_d,
346        ulong         offset_d,
347        int           ne01,
348        ulong         nb01,
349        ulong         nb02,
350        ulong         nb03,
351        uint4         ne11,
352        uint4         ne12,
353        ulong         nb10,
354        ulong         nb11,
355        ulong         nb12,
356        int           nblk0,
357        int           ne1_dst,
358        int           ne2_dst,
359        int           ne3_dst
360) {
361    src0  = src0  + offset0;
362    src1  = src1  + offset1;
363    dst_q = dst_q + offset_q;
364    dst_d = dst_d + offset_d;
365
366    int i03 = get_group_id(2);
367    int i02 = get_group_id(1);
368    int i01 = get_group_id(0)*get_local_size(1) + get_local_id(1);
369
370    if (i01 >= ne01) {
371        return;
372    }
373
374    int i12 = fastmod(i03, ne12);
375    int i11 = fastmod(i02, ne11);
376
377    int i10 = i01;
378    int i1  = ((global int *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
379
380    long row_blk_base = ((long)i03 * ne2_dst * ne1_dst + (long)i02 * ne1_dst + i1) * nblk0;
381
382    global half  * d_row = (global half  *)(dst_d) + row_blk_base;
383    global char  * q_row = (global char  *)(dst_q) + row_blk_base * QK8_0;
384    global float * src_row = (global float *)(src0 + i01*nb01 + i02*nb02 + i03*nb03);
385
386    for (int blk = get_local_id(0); blk < nblk0; blk += get_local_size(0)) {
387        global float * x = src_row + blk * QK8_0;
388        global char  * q = q_row + blk * QK8_0;
389
390        quantize_q8_0_block(x, q, d_row + blk);
391    }
392}
393
394kernel void kernel_set_rows_f16_i32(
395        global char * src0,
396        ulong         offset0,
397        global char * src1,
398        ulong         offset1,
399        global char * dst,
400        ulong         offsetd,
401        int           ne01,
402        ulong         nb01,
403        ulong         nb02,
404        ulong         nb03,
405        uint4         ne11,
406        uint4         ne12,
407        ulong         nb10,
408        ulong         nb11,
409        ulong         nb12,
410        int           nblk0,
411        ulong         nb1,
412        ulong         nb2,
413        ulong         nb3
414) {
415    src0 = src0 + offset0;
416    src1 = src1 + offset1;
417    dst  = dst  + offsetd;
418
419    int i03 = get_group_id(2);
420    int i02 = get_group_id(1);
421    int i01 = get_group_id(0)*get_local_size(1) + get_local_id(1);
422
423    if (i01 >= ne01) {
424        return;
425    }
426
427    //int i12 = i03%ne12;
428    //int i11 = i02%ne11;
429    int i12 = fastmod(i03, ne12);
430    int i11 = fastmod(i02, ne11);
431
432    int i10 = i01;
433    int i1  = ((global int *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
434
435    global half  * dst_row = (global half  *) (dst  +  i1*nb1  + i02*nb2  + i03*nb3);
436    global float * src_row = (global float *) (src0 + i01*nb01 + i02*nb02 + i03*nb03);
437
438    for (int ind = get_local_id(0); ind < nblk0; ind += get_local_size(0)) {
439        dst_row[ind] = src_row[ind];
440    }
441}
442
443// f32 -> q4_0 quantize set_rows. Block = half d + uchar qs[16] (shuffled
444// nibbles: qs[j] low/high = elem j / j+16).
445// Dequant: val[i] = d * (nibble_i - 8)
446// nblk0 = number of q4_0 blocks per row = ne00 / 32.
447#define QK4_0 32
448#define Q4_0_BLOCK_SIZE 18
449
450inline void quantize_q4_0_block(global float * x, global uchar * qs, global half * d_out) {
451    // Find the signed value with the largest absolute magnitude (matches ggml ref).
452    float max  = 0.0f;
453    float amax = 0.0f;
454    for (int j = 0; j < QK4_0; j++) {
455        float v = x[j];
456        float a = fabs(v);
457        if (a > amax) {
458            amax = a;
459            max  = v;
460        }
461    }
462
463    float d  = max / -8.0f;
464    float id = (d != 0.0f) ? 1.0f / d : 0.0f;
465
466    vstore_half(d, 0, d_out);
467
468    for (int j = 0; j < QK4_0/2; j++) {
469        float x0 = x[j]           * id;
470        float x1 = x[j + QK4_0/2] * id;
471
472        int i0 = (int)(x0 + 8.5f);
473        int i1 = (int)(x1 + 8.5f);
474        if (i0 < 0)  i0 = 0;
475        if (i0 > 15) i0 = 15;
476        if (i1 < 0)  i1 = 0;
477        if (i1 > 15) i1 = 15;
478
479        qs[j] = (uchar)i0 | ((uchar)i1 << 4);
480    }
481}
482
483kernel void kernel_set_rows_q4_0_i64(
484        global char * src0,
485        ulong         offset0,
486        global char * src1,
487        ulong         offset1,
488        global char * dst,
489        ulong         offsetd,
490        int           ne01,
491        ulong         nb01,
492        ulong         nb02,
493        ulong         nb03,
494        uint4         ne11,
495        uint4         ne12,
496        ulong         nb10,
497        ulong         nb11,
498        ulong         nb12,
499        int           nblk0,
500        ulong         nb1,
501        ulong         nb2,
502        ulong         nb3
503) {
504    src0 = src0 + offset0;
505    src1 = src1 + offset1;
506    dst  = dst  + offsetd;
507
508    int i03 = get_group_id(2);
509    int i02 = get_group_id(1);
510    int i01 = get_group_id(0)*get_local_size(1) + get_local_id(1);
511
512    if (i01 >= ne01) {
513        return;
514    }
515
516    int i12 = fastmod(i03, ne12);
517    int i11 = fastmod(i02, ne11);
518
519    int i10 = i01;
520    long i1 = ((global long *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
521
522    global char  * dst_row = (global char  *) (dst  +  i1*nb1  + i02*nb2  + i03*nb3);
523    global float * src_row = (global float *) (src0 + i01*nb01 + i02*nb02 + i03*nb03);
524
525    for (int blk = get_local_id(0); blk < nblk0; blk += get_local_size(0)) {
526        global float * x    = src_row + blk * QK4_0;
527        global char  * y    = dst_row + blk * Q4_0_BLOCK_SIZE;
528        global half  * yd   = (global half  *)(y);
529        global uchar * yqs  = (global uchar *)(y + 2);
530
531        quantize_q4_0_block(x, yqs, yd);
532    }
533}
534
535kernel void kernel_set_rows_q4_0_i32(
536        global char * src0,
537        ulong         offset0,
538        global char * src1,
539        ulong         offset1,
540        global char * dst,
541        ulong         offsetd,
542        int           ne01,
543        ulong         nb01,
544        ulong         nb02,
545        ulong         nb03,
546        uint4         ne11,
547        uint4         ne12,
548        ulong         nb10,
549        ulong         nb11,
550        ulong         nb12,
551        int           nblk0,
552        ulong         nb1,
553        ulong         nb2,
554        ulong         nb3
555) {
556    src0 = src0 + offset0;
557    src1 = src1 + offset1;
558    dst  = dst  + offsetd;
559
560    int i03 = get_group_id(2);
561    int i02 = get_group_id(1);
562    int i01 = get_group_id(0)*get_local_size(1) + get_local_id(1);
563
564    if (i01 >= ne01) {
565        return;
566    }
567
568    int i12 = fastmod(i03, ne12);
569    int i11 = fastmod(i02, ne11);
570
571    int i10 = i01;
572    int i1  = ((global int *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
573
574    global char  * dst_row = (global char  *) (dst  +  i1*nb1  + i02*nb2  + i03*nb3);
575    global float * src_row = (global float *) (src0 + i01*nb01 + i02*nb02 + i03*nb03);
576
577    for (int blk = get_local_id(0); blk < nblk0; blk += get_local_size(0)) {
578        global float * x    = src_row + blk * QK4_0;
579        global char  * y    = dst_row + blk * Q4_0_BLOCK_SIZE;
580        global half  * yd   = (global half  *)(y);
581        global uchar * yqs  = (global uchar *)(y + 2);
582
583        quantize_q4_0_block(x, yqs, yd);
584    }
585}
586
587// SoA variants for q4_0 dst. Used when the backend has split block_q4_0 records
588// into separate quant (dst_q) and scale (dst_d) sub-buffers — same pattern as
589// the q8_0 SoA variants above.
590//
591// Layout (matches kernel_convert_block_q4_0, the "shuffled" variant):
592//   dst_q: contiguous 16 packed nibbles per block, block i at offset i * 16 bytes.
593//   dst_d: contiguous fp16 scales, block i at offset i * 2 bytes.
594// Nibble layout inside each byte is unchanged from AoS: qs[j] low nibble = element j,
595// qs[j] high nibble = element j+16. kernel_restore_block_q4_0 copies bytes as-is.
596kernel void kernel_set_rows_q4_0_soa_i64(
597        global char * src0,
598        ulong         offset0,
599        global char * src1,
600        ulong         offset1,
601        global char * dst_q,
602        ulong         offset_q,
603        global char * dst_d,
604        ulong         offset_d,
605        int           ne01,
606        ulong         nb01,
607        ulong         nb02,
608        ulong         nb03,
609        uint4         ne11,
610        uint4         ne12,
611        ulong         nb10,
612        ulong         nb11,
613        ulong         nb12,
614        int           nblk0,
615        int           ne1_dst,
616        int           ne2_dst,
617        int           ne3_dst
618) {
619    src0  = src0  + offset0;
620    src1  = src1  + offset1;
621    dst_q = dst_q + offset_q;
622    dst_d = dst_d + offset_d;
623
624    int i03 = get_group_id(2);
625    int i02 = get_group_id(1);
626    int i01 = get_group_id(0)*get_local_size(1) + get_local_id(1);
627
628    if (i01 >= ne01) {
629        return;
630    }
631
632    int i12 = fastmod(i03, ne12);
633    int i11 = fastmod(i02, ne11);
634
635    int i10 = i01;
636    long i1 = ((global long *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
637
638    long row_blk_base = ((long)i03 * ne2_dst * ne1_dst + (long)i02 * ne1_dst + i1) * nblk0;
639
640    global half  * d_row   = (global half  *)(dst_d) + row_blk_base;
641    global uchar * q_row   = (global uchar *)(dst_q) + row_blk_base * (QK4_0/2);
642    global float * src_row = (global float *)(src0 + i01*nb01 + i02*nb02 + i03*nb03);
643
644    for (int blk = get_local_id(0); blk < nblk0; blk += get_local_size(0)) {
645        global float * x    = src_row + blk * QK4_0;
646        global uchar * qs   = q_row   + blk * (QK4_0/2);
647        global half  * d_bk = d_row   + blk;
648
649        quantize_q4_0_block(x, qs, d_bk);
650    }
651}
652
653kernel void kernel_set_rows_q4_0_soa_i32(
654        global char * src0,
655        ulong         offset0,
656        global char * src1,
657        ulong         offset1,
658        global char * dst_q,
659        ulong         offset_q,
660        global char * dst_d,
661        ulong         offset_d,
662        int           ne01,
663        ulong         nb01,
664        ulong         nb02,
665        ulong         nb03,
666        uint4         ne11,
667        uint4         ne12,
668        ulong         nb10,
669        ulong         nb11,
670        ulong         nb12,
671        int           nblk0,
672        int           ne1_dst,
673        int           ne2_dst,
674        int           ne3_dst
675) {
676    src0  = src0  + offset0;
677    src1  = src1  + offset1;
678    dst_q = dst_q + offset_q;
679    dst_d = dst_d + offset_d;
680
681    int i03 = get_group_id(2);
682    int i02 = get_group_id(1);
683    int i01 = get_group_id(0)*get_local_size(1) + get_local_id(1);
684
685    if (i01 >= ne01) {
686        return;
687    }
688
689    int i12 = fastmod(i03, ne12);
690    int i11 = fastmod(i02, ne11);
691
692    int i10 = i01;
693    int i1  = ((global int *)(src1 + i10*nb10 + i11*nb11 + i12*nb12))[0];
694
695    long row_blk_base = ((long)i03 * ne2_dst * ne1_dst + (long)i02 * ne1_dst + i1) * nblk0;
696
697    global half  * d_row   = (global half  *)(dst_d) + row_blk_base;
698    global uchar * q_row   = (global uchar *)(dst_q) + row_blk_base * (QK4_0/2);
699    global float * src_row = (global float *)(src0 + i01*nb01 + i02*nb02 + i03*nb03);
700
701    for (int blk = get_local_id(0); blk < nblk0; blk += get_local_size(0)) {
702        global float * x    = src_row + blk * QK4_0;
703        global uchar * qs   = q_row   + blk * (QK4_0/2);
704        global half  * d_bk = d_row   + blk;
705
706        quantize_q4_0_block(x, qs, d_bk);
707    }
708}