hey! this is a static git repo viewer. commit diffs, file history, older commits and other revisions won't load without javascript.

you are still able to view source files on branch/tag tips, but the experience is degraded.

char-fork/llama.cpp

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

Sigbjørn Skjæretsycl : add Q2_K reordered MMVQ and ESIMD kernels (again) (#27490)8672290032026-08-22

1#include "convert.hpp"
2#include "dequantize.hpp"
3#include "presets.hpp"
4
5template <int qk, int qr, dequantize_kernel_t dequantize_kernel, typename dst_t>
6static void dequantize_block(const void * __restrict__ vx, dst_t * __restrict__ y, const int64_t k,
7                             const sycl::nd_item<3> &item_ct1) {
8    const int64_t i = 2 * (item_ct1.get_local_range(2) * item_ct1.get_group(2) +
9                       item_ct1.get_local_id(2));
10
11    if (i >= k) {
12        return;
13    }
14
15    const int64_t ib = i/qk; // block index
16    const int64_t iqs = (i%qk)/qr; // quant index
17    const int64_t iybs = i - i%qk; // y block start index
18    const int64_t y_offset = qr == 1 ? 1 : qk/2;
19
20    // dequantize
21    dfloat2 v;
22    dequantize_kernel(vx, ib, iqs, v);
23
24    y[iybs + iqs + 0] = v.x();
25    y[iybs + iqs + y_offset] = v.y();
26}
27
28template <int qk, int qr, dequantize_kernel_t dequantize_kernel, typename dst_t>
29static void dequantize_block_sycl(const void *__restrict__ vx,
30                                  dst_t *__restrict__ y, const int64_t k,
31                                  dpct::queue_ptr stream) {
32    const int64_t num_blocks = (k + 2*SYCL_DEQUANTIZE_BLOCK_SIZE - 1) / (2*SYCL_DEQUANTIZE_BLOCK_SIZE);
33    {
34        dpct::has_capability_or_fail(stream->get_device(),
35                                     {sycl::aspect::fp16});
36        stream->parallel_for(
37            sycl::nd_range<3>(
38                sycl::range<3>(1, 1, num_blocks) *
39                    sycl::range<3>(1, 1, SYCL_DEQUANTIZE_BLOCK_SIZE),
40                sycl::range<3>(1, 1, SYCL_DEQUANTIZE_BLOCK_SIZE)),
41            [=](sycl::nd_item<3> item_ct1) {
42                dequantize_block<qk, qr, dequantize_kernel>(vx, y, k, item_ct1);
43            });
44    }
45}
46
47template <typename dst_t>
48static void dequantize_row_q2_K_sycl(const void *vx, dst_t *y, const int64_t k,
49                                     dpct::queue_ptr stream) {
50    const int64_t nb = k / QK_K;
51#if QK_K == 256
52    {
53        dpct::has_capability_or_fail(stream->get_device(),
54                                     {sycl::aspect::fp16});
55
56        stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
57                                                   sycl::range<3>(1, 1, 64),
58                                               sycl::range<3>(1, 1, 64)),
59                             [=](sycl::nd_item<3> item_ct1) {
60                                 dequantize_block_q2_K(vx, y, item_ct1);
61                             });
62    }
63#else
64    {
65        dpct::has_capability_or_fail(stream->get_device(),
66                                     {sycl::aspect::fp16});
67
68        stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
69                                                   sycl::range<3>(1, 1, 32),
70                                               sycl::range<3>(1, 1, 32)),
71                             [=](sycl::nd_item<3> item_ct1) {
72                                 dequantize_block_q2_K(vx, y, item_ct1);
73                             });
74    }
75
76#endif
77}
78
79template <typename dst_t>
80static void dequantize_row_q2_K_sycl_reorder(const void *vx, dst_t *y, const int64_t k,
81                                             dpct::queue_ptr stream) {
82    const int64_t nb = k / QK_K;
83
84    dpct::has_capability_or_fail(stream->get_device(), { sycl::aspect::fp16 });
85    stream->parallel_for(
86        sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
87        [=](sycl::nd_item<3> item_ct1) {
88            dequantize_block_q2_K_reorder(vx, y, item_ct1, nb);
89        });
90}
91
92template <typename dst_t>
93static void dequantize_row_q3_K_sycl(const void *vx, dst_t *y, const int64_t k,
94                                     dpct::queue_ptr stream) {
95    const int64_t nb = k / QK_K;
96#if QK_K == 256
97    {
98        dpct::has_capability_or_fail(stream->get_device(),
99                                     {sycl::aspect::fp16});
100
101        stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
102                                                   sycl::range<3>(1, 1, 64),
103                                               sycl::range<3>(1, 1, 64)),
104                             [=](sycl::nd_item<3> item_ct1) {
105                                 dequantize_block_q3_K(vx, y, item_ct1);
106                             });
107    }
108#else
109    {
110        dpct::has_capability_or_fail(stream->get_device(),
111                                     {sycl::aspect::fp16});
112
113        stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
114                                                   sycl::range<3>(1, 1, 32),
115                                               sycl::range<3>(1, 1, 32)),
116                             [=](sycl::nd_item<3> item_ct1) {
117                                 dequantize_block_q3_K(vx, y, item_ct1);
118                             });
119    }
120#endif
121}
122
123template <typename dst_t>
124static void dequantize_row_q3_K_sycl_reorder(const void *vx, dst_t *y, const int64_t k,
125                                             dpct::queue_ptr stream) {
126    const int64_t nb = k / QK_K;
127
128    dpct::has_capability_or_fail(stream->get_device(), { sycl::aspect::fp16 });
129    stream->parallel_for(
130        sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
131        [=](sycl::nd_item<3> item_ct1) {
132            dequantize_block_q3_K_reorder(vx, y, item_ct1, nb);
133        });
134}
135
136template <typename dst_t>
137static void dequantize_row_q4_0_sycl(const void *vx, dst_t *y, const int64_t k,
138                                     dpct::queue_ptr stream) {
139    const int64_t nb32 = k / 32;
140    const int64_t nb = (k + 255) / 256;
141    {
142        dpct::has_capability_or_fail(stream->get_device(),
143                                     {sycl::aspect::fp16});
144
145        stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
146                                                   sycl::range<3>(1, 1, 32),
147                                               sycl::range<3>(1, 1, 32)),
148                             [=](sycl::nd_item<3> item_ct1) {
149                                 dequantize_block_q4_0(vx, y, nb32, item_ct1);
150                             });
151    }
152}
153
154template <typename dst_t>
155static void dequantize_row_q4_0_sycl_reorder(const void *vx, dst_t *y, const int64_t k,
156                                     dpct::queue_ptr stream) {
157
158    dpct::has_capability_or_fail(stream->get_device(),
159                                    {sycl::aspect::fp16});
160
161    int constexpr WARP_K = WARP_SIZE * QK4_0;
162    const int n_warp = (k + WARP_K - 1) / WARP_K;
163    GGML_ASSERT(k % 2 == 0);
164    stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, n_warp) *
165        sycl::range<3>(1, 1, WARP_SIZE),
166        sycl::range<3>(1, 1, WARP_SIZE)),
167        [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
168            dequantize_block_q4_0_reorder(vx, y, k, item_ct1);
169        });
170
171}
172
173template <typename dst_t>
174static void dequantize_row_q8_0_sycl_reorder(const void *vx, dst_t *y, const int64_t k,
175                                     dpct::queue_ptr stream) {
176
177    dpct::has_capability_or_fail(stream->get_device(),
178                                    {sycl::aspect::fp16});
179
180    int constexpr WARP_K = WARP_SIZE * QK8_0;
181    const int n_warp = (k + WARP_K - 1) / WARP_K;
182    GGML_ASSERT(k % QK8_0 == 0);
183    stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, n_warp) *
184        sycl::range<3>(1, 1, WARP_SIZE),
185        sycl::range<3>(1, 1, WARP_SIZE)),
186        [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
187            dequantize_block_q8_0_reorder(vx, y, k, item_ct1);
188        });
189
190}
191
192template <typename dst_t>
193static void dequantize_row_q4_1_sycl(const void *vx, dst_t *y, const int64_t k,
194                                     dpct::queue_ptr stream) {
195    const int64_t nb32 = k / 32;
196    const int64_t nb = (k + 255) / 256;
197    {
198        dpct::has_capability_or_fail(stream->get_device(),
199                                     {sycl::aspect::fp16});
200
201        stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
202                                                   sycl::range<3>(1, 1, 32),
203                                               sycl::range<3>(1, 1, 32)),
204                             [=](sycl::nd_item<3> item_ct1) {
205                                 dequantize_block_q4_1(vx, y, nb32, item_ct1);
206                             });
207    }
208}
209
210
211template <typename dst_t>
212static void dequantize_row_q4_K_sycl(const void *vx, dst_t *y, const int64_t k,
213                                     dpct::queue_ptr stream) {
214    const int64_t nb = k / QK_K;
215    {
216        dpct::has_capability_or_fail(stream->get_device(),
217                                     {sycl::aspect::fp16});
218
219        stream->submit([&](sycl::handler &cgh) {
220            sycl::local_accessor<uint8_t, 1> scale_local_acc(sycl::range<1>(12), cgh);
221            cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
222                                                   sycl::range<3>(1, 1, 32),
223                                               sycl::range<3>(1, 1, 32)),
224                             [=](sycl::nd_item<3> item_ct1) {
225                                 dequantize_block_q4_K(vx, y, get_pointer(scale_local_acc), item_ct1);
226                             });
227        });
228    }
229}
230
231template <typename dst_t>
232static void dequantize_row_q4_K_sycl_reorder(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr stream) {
233    const int64_t nb = k / QK_K;
234    const size_t  local_size  = 32;
235    const size_t  global_size = nb * local_size;
236
237    dpct::has_capability_or_fail(stream->get_device(), { sycl::aspect::fp16 });
238
239    stream->submit([&](sycl::handler & cgh) {
240        sycl::local_accessor<uint8_t, 1> scale_local_acc(sycl::range<1>(12), cgh);
241
242        cgh.parallel_for(sycl::nd_range<1>(sycl::range<1>(global_size), sycl::range<1>(local_size)),
243                         [=](sycl::nd_item<1> item_ct1) {
244                             dequantize_block_q4_K_reorder(vx, y, get_pointer(scale_local_acc), item_ct1, nb);
245                         });
246    });
247}
248
249template <typename dst_t>
250static void dequantize_row_q5_K_sycl(const void *vx, dst_t *y, const int64_t k,
251                                     dpct::queue_ptr stream) {
252    const int64_t nb = k / QK_K;
253#if QK_K == 256
254    {
255        dpct::has_capability_or_fail(stream->get_device(),
256                                     {sycl::aspect::fp16});
257
258        stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
259                                                   sycl::range<3>(1, 1, 64),
260                                               sycl::range<3>(1, 1, 64)),
261                             [=](sycl::nd_item<3> item_ct1) {
262                                 dequantize_block_q5_K(vx, y, item_ct1);
263                             });
264    }
265#else
266    {
267        dpct::has_capability_or_fail(stream->get_device(),
268                                     {sycl::aspect::fp16});
269
270        stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
271                                                   sycl::range<3>(1, 1, 32),
272                                               sycl::range<3>(1, 1, 32)),
273                             [=](sycl::nd_item<3> item_ct1) {
274                                 dequantize_block_q5_K(vx, y, item_ct1);
275                             });
276    }
277
278#endif
279}
280
281template <typename dst_t>
282static void dequantize_row_q5_K_sycl_reorder(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr stream) {
283    const int64_t nb = k / QK_K;
284
285    dpct::has_capability_or_fail(stream->get_device(), { sycl::aspect::fp16 });
286
287    stream->submit([&](sycl::handler & cgh) {
288        sycl::local_accessor<uint8_t, 1> scale_local_acc(sycl::range<1>(K_SCALE_SIZE), cgh);
289
290        cgh.parallel_for(
291            sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
292            [=](sycl::nd_item<3> item_ct1) {
293                dequantize_block_q5_K_reorder(vx, y, get_pointer(scale_local_acc), item_ct1, nb);
294            });
295    });
296}
297
298template <typename dst_t>
299static void dequantize_row_q6_K_sycl(const void *vx, dst_t *y, const int64_t k,
300                                     dpct::queue_ptr stream) {
301    const int64_t nb = k / QK_K;
302#if QK_K == 256
303    {
304        dpct::has_capability_or_fail(stream->get_device(),
305                                     {sycl::aspect::fp16});
306
307        stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
308                                                   sycl::range<3>(1, 1, 64),
309                                               sycl::range<3>(1, 1, 64)),
310                             [=](sycl::nd_item<3> item_ct1) {
311                                 dequantize_block_q6_K(vx, y, item_ct1);
312                             });
313    }
314#else
315    {
316        dpct::has_capability_or_fail(stream->get_device(),
317                                     {sycl::aspect::fp16});
318
319        stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
320                                                   sycl::range<3>(1, 1, 32),
321                                               sycl::range<3>(1, 1, 32)),
322                             [=](sycl::nd_item<3> item_ct1) {
323                                 dequantize_block_q6_K(vx, y, item_ct1);
324                             });
325    }
326
327#endif
328}
329
330template <typename dst_t>
331static void dequantize_row_q6_K_sycl_reorder(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr stream) {
332    const int64_t nb = k / QK_K;
333
334    dpct::has_capability_or_fail(stream->get_device(), { sycl::aspect::fp16 });
335
336    stream->parallel_for(
337        sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
338        [=](sycl::nd_item<3> item_ct1) { dequantize_block_q6_K_reorder(vx, y, item_ct1, nb); });
339}
340
341template <typename dst_t>
342static void dequantize_row_iq1_s_sycl(const void *vx, dst_t *y, const int64_t k,
343                                        dpct::queue_ptr stream) {
344    const int64_t nb = k / QK_K;
345    {
346        dpct::has_capability_or_fail(stream->get_device(),
347                                     {sycl::aspect::fp16});
348
349        stream->submit([&](sycl::handler &cgh) {
350            cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
351                                                   sycl::range<3>(1, 1, 32),
352                                               sycl::range<3>(1, 1, 32)),
353                             [=](sycl::nd_item<3> item_ct1) {
354                                 dequantize_block_iq1_s(
355                                     vx, y, item_ct1, iq1s_grid_gpu
356                                     );
357                             });
358        });
359    }
360}
361
362template <typename dst_t>
363static void dequantize_row_iq1_m_sycl(const void *vx, dst_t *y, const int64_t k,
364                                        dpct::queue_ptr stream) {
365    const int64_t nb = k / QK_K;
366    {
367        dpct::has_capability_or_fail(stream->get_device(),
368                                     {sycl::aspect::fp16});
369
370        stream->submit([&](sycl::handler &cgh) {
371            cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
372                                                   sycl::range<3>(1, 1, 32),
373                                               sycl::range<3>(1, 1, 32)),
374                             [=](sycl::nd_item<3> item_ct1) {
375                                 dequantize_block_iq1_m(
376                                     vx, y, item_ct1, iq1s_grid_gpu
377                                     );
378                             });
379        });
380    }
381}
382
383template <typename dst_t>
384static void dequantize_row_iq2_xxs_sycl(const void *vx, dst_t *y, const int64_t k,
385                                        dpct::queue_ptr stream) {
386    const int64_t nb = k / QK_K;
387    {
388        dpct::has_capability_or_fail(stream->get_device(),
389                                     {sycl::aspect::fp16});
390
391        stream->submit([&](sycl::handler &cgh) {
392            cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
393                                                   sycl::range<3>(1, 1, 32),
394                                               sycl::range<3>(1, 1, 32)),
395                             [=](sycl::nd_item<3> item_ct1) {
396                                 dequantize_block_iq2_xxs(
397                                     vx, y, item_ct1, iq2xxs_grid,
398                                     ksigns_iq2xs, kmask_iq2xs);
399                             });
400        });
401    }
402}
403
404template <typename dst_t>
405static void dequantize_row_iq2_xs_sycl(const void *vx, dst_t *y, const int64_t k,
406                                       dpct::queue_ptr stream) {
407    const int64_t nb = k / QK_K;
408    {
409        dpct::has_capability_or_fail(stream->get_device(),
410                                     {sycl::aspect::fp16});
411
412        stream->submit([&](sycl::handler &cgh) {
413            cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
414                                                   sycl::range<3>(1, 1, 32),
415                                               sycl::range<3>(1, 1, 32)),
416                             [=](sycl::nd_item<3> item_ct1) {
417                                 dequantize_block_iq2_xs(
418                                     vx, y, item_ct1, iq2xs_grid,
419                                     ksigns_iq2xs, kmask_iq2xs);
420                             });
421        });
422    }
423}
424
425template <typename dst_t>
426static void dequantize_row_iq2_s_sycl(const void *vx, dst_t *y, const int64_t k,
427                                      dpct::queue_ptr stream) {
428    const int64_t nb = k / QK_K;
429    {
430        dpct::has_capability_or_fail(stream->get_device(),
431                                     {sycl::aspect::fp16});
432
433        stream->submit([&](sycl::handler &cgh) {
434            cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
435                                                   sycl::range<3>(1, 1, 32),
436                                               sycl::range<3>(1, 1, 32)),
437                             [=](sycl::nd_item<3> item_ct1) {
438                                 dequantize_block_iq2_s(vx, y, item_ct1);
439                             });
440        });
441    }
442}
443
444
445template <typename dst_t>
446static void dequantize_row_iq3_xxs_sycl(const void *vx, dst_t *y, const int64_t k,
447                                        dpct::queue_ptr stream) {
448    const int64_t nb = k / QK_K;
449    {
450        dpct::has_capability_or_fail(stream->get_device(),
451                                     {sycl::aspect::fp16});
452
453        stream->submit([&](sycl::handler &cgh) {
454            cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
455                                                   sycl::range<3>(1, 1, 32),
456                                               sycl::range<3>(1, 1, 32)),
457                             [=](sycl::nd_item<3> item_ct1) {
458                                 dequantize_block_iq3_xxs(
459                                     vx, y, item_ct1, iq3xxs_grid,
460                                     ksigns_iq2xs, kmask_iq2xs);
461                             });
462        });
463    }
464}
465
466template <typename dst_t>
467static void dequantize_row_iq3_s_sycl(const void *vx, dst_t *y, const int64_t k,
468                                        dpct::queue_ptr stream) {
469    const int64_t nb = k / QK_K;
470    {
471        dpct::has_capability_or_fail(stream->get_device(),
472                                     {sycl::aspect::fp16});
473
474        stream->submit([&](sycl::handler &cgh) {
475            cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
476                                                   sycl::range<3>(1, 1, 32),
477                                               sycl::range<3>(1, 1, 32)),
478                             [=](sycl::nd_item<3> item_ct1) {
479                                 dequantize_block_iq3_s(
480                                     vx, y, item_ct1, kmask_iq2xs, iq3s_grid);
481                             });
482        });
483    }
484}
485
486template <typename dst_t>
487static void dequantize_row_iq4_xs_sycl(const void *vx, dst_t *y, const int64_t k,
488                                       dpct::queue_ptr stream) {
489    const int64_t nb = (k + QK_K - 1) / QK_K;
490#if QK_K == 64
491    dequantize_row_iq4_nl_sycl(vx, y, k, stream);
492#else
493      {
494            dpct::has_capability_or_fail(stream->get_device(),
495                                         {sycl::aspect::fp16});
496
497            stream->submit([&](sycl::handler &cgh) {
498                  cgh.parallel_for(
499                      sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
500                                            sycl::range<3>(1, 1, 32),
501                                        sycl::range<3>(1, 1, 32)),
502                      [=](sycl::nd_item<3> item_ct1) {
503                            dequantize_block_iq4_xs(vx, y, item_ct1);
504                      });
505            });
506      }
507#endif
508}
509
510template <typename dst_t>
511static void dequantize_row_iq4_nl_sycl(const void *vx, dst_t *y, const int64_t k,
512                                       dpct::queue_ptr stream) {
513    const int64_t nb = (k + QK_K - 1) / QK_K;
514      {
515            dpct::has_capability_or_fail(stream->get_device(),
516                                         {sycl::aspect::fp16});
517
518            stream->submit([&](sycl::handler &cgh) {
519                  cgh.parallel_for(
520                      sycl::nd_range<3>(sycl::range<3>(1, 1, nb) *
521                                            sycl::range<3>(1, 1, 32),
522                                        sycl::range<3>(1, 1, 32)),
523                      [=](sycl::nd_item<3> item_ct1) {
524                            dequantize_block_iq4_nl(vx, y, item_ct1);
525                      });
526            });
527      }
528}
529
530template <typename dst_t>
531static void dequantize_row_mxfp4_sycl(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr stream) {
532    const int nb = (k + QK_K - 1) / QK_K;
533    stream->parallel_for(
534        sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)),
535        [=](sycl::nd_item<3> item_ct1) {
536            dequantize_block_mxfp4(vx, y, item_ct1);
537        });
538}
539
540template <typename dst_t>
541static void dequantize_row_nvfp4_sycl(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr stream) {
542    GGML_ASSERT(k % QK_NVFP4 == 0);
543    const int nb = k / QK_NVFP4;
544    stream->parallel_for(
545        sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)),
546        [=](sycl::nd_item<3> /*item_ct1*/) {
547            dequantize_block_nvfp4(vx, y, k);
548        });
549}
550
551
552template <int qk, int qr, dequantize_kernel_t dequantize_kernel, typename dst_t>
553static void dequantize_block_nc(const void * __restrict__ vx, dst_t * __restrict__ y,
554        const int64_t ne00, const int64_t ne01, const int64_t ne02,
555        const int64_t s01, const int64_t s02, const int64_t s03) {
556    auto          item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
557    const int64_t i00 = 2 * (int64_t(item_ct1.get_local_range(2)) * item_ct1.get_group(2) + item_ct1.get_local_id(2));
558
559    if (i00 >= ne00) {
560        return;
561    }
562
563    const int64_t i01 = item_ct1.get_group(1);
564    const int64_t i02 = item_ct1.get_group(0) % ne02;
565    const int64_t i03 = item_ct1.get_group(0) / ne02;
566
567    const int64_t ibx0 = i03*s03 + i02*s02 + i01*s01;
568
569    const int64_t ib = ibx0 + i00/qk; // block index
570    const int64_t iqs = (i00%qk)/qr; // quant index
571    const int64_t iybs = i00 - i00%qk; // y block start index
572    const int64_t y_offset = qr == 1 ? 1 : qk/2;
573
574    // dequantize
575    #ifdef GGML_SYCL_F16
576        sycl::half2 v;
577    #else
578        sycl::float2 v;
579    #endif
580
581    dequantize_kernel(vx, ib, iqs, v);
582
583    const int64_t iy0 = ((i03*ne02 + i02)*ne01 + i01)*ne00 + iybs + iqs;
584    y[iy0 + 0]        = ggml_sycl_cast<dst_t>(v.x());
585    y[iy0 + y_offset] = ggml_sycl_cast<dst_t>(v.y());
586}
587
588
589template <int qk, int qr, dequantize_kernel_t dequantize_kernel, typename dst_t>
590static void dequantize_block_nc_sycl(const void *    vx,
591                                  dst_t *         y,
592                                  const int64_t   ne00,
593                                  const int64_t   ne01,
594                                  const int64_t   ne02,
595                                  const int64_t   ne03,
596                                  const int64_t   s01,
597                                  const int64_t   s02,
598                                  const int64_t   s03,
599                                  dpct::queue_ptr stream) {
600    const dpct::dim3 num_blocks((ne00 + 2 * SYCL_DEQUANTIZE_BLOCK_SIZE - 1) / (2 * SYCL_DEQUANTIZE_BLOCK_SIZE), ne01,
601                                ne02 * ne03);
602    stream->parallel_for(sycl::nd_range<3>(num_blocks * sycl::range<3>(1, 1, SYCL_DEQUANTIZE_BLOCK_SIZE),
603                                           sycl::range<3>(1, 1, SYCL_DEQUANTIZE_BLOCK_SIZE)),
604                         [=](sycl::nd_item<3> item_ct1) {
605                             GGML_UNUSED(item_ct1);
606                             dequantize_block_nc<qk, qr, dequantize_kernel>(vx, y, ne00, ne01, ne02, s01, s02, s03);
607                         });
608}
609template <typename src_t, typename dst_t>
610static void convert_unary_nc(const void * __restrict__ vx, dst_t * __restrict__ y, const int64_t ne00, const int64_t ne01,
611                          const int64_t ne02, const int64_t s01, const int64_t s02, const int64_t s03,
612                          const sycl::nd_item<3> & item_ct1) {
613
614    const int64_t work_group_size = item_ct1.get_local_range(2);
615    const int64_t global_id       = item_ct1.get_local_id(2) + work_group_size * item_ct1.get_group(2);
616
617    const int64_t i01 = item_ct1.get_group(1);
618    const int64_t i02 = item_ct1.get_group(0) % ne02;
619    const int64_t i03 = item_ct1.get_group(0) / ne02;
620
621    // make each work-item deal with more elements since sycl global range can not exceed max int
622    const src_t * x = static_cast<const src_t *>(vx);
623    const int64_t ix = i03 * s03 + i02 * s02 + i01 * s01;
624    const int64_t iy = ((i03 * ne02 + i02) * ne01 + i01) * ne00;
625
626#pragma unroll
627    for (int64_t i00 = global_id; i00 < ne00; i00 += work_group_size * item_ct1.get_group_range(2)) {
628        y[iy + i00] = static_cast<dst_t>(x[ix + i00]);
629    }
630}
631
632template <typename src_t, typename dst_t>
633static void convert_unary_nc_sycl(const void * __restrict__ vx, dst_t * __restrict__ y,
634                                  const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03,
635                                  const int64_t s01, const int64_t s02, const int64_t s03, dpct::queue_ptr queue) {
636    dpct::has_capability_or_fail(queue->get_device(), { sycl::aspect::fp16 });
637
638    sycl::range<3> global_size(ne02 * ne03, ne01, ceil_div(ne00, SYCL_DEQUANTIZE_BLOCK_SIZE));
639
640    // decrease global range when it exceeds the max int
641    // TODO: Downsample logic is separated from the kernel, a rewrite is desirable
642    int64_t        downsized_workgroup = downsample_sycl_global_range(global_size[0], SYCL_DEQUANTIZE_BLOCK_SIZE);
643    sycl::range<3> workgroup_size(1, 1, downsized_workgroup);
644
645    queue->parallel_for(sycl::nd_range<3>(global_size * workgroup_size, workgroup_size), [=](sycl::nd_item<3> item_ct1) {
646        convert_unary_nc<src_t>(vx, y, ne00, ne01, ne02, s01, s02, s03, item_ct1);
647    });
648}
649
650template <typename src_t, typename dst_t>
651static void convert_unary_sycl(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr queue) {
652    convert_unary_nc_sycl<src_t>(vx, y, k, 1, 1, 1, k, k, k, queue);
653}
654
655
656to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) {
657    switch (type) {
658        case GGML_TYPE_Q1_0:
659            return dequantize_block_sycl<QK1_0, QR1_0, dequantize_q1_0>;
660        case GGML_TYPE_Q2_0:
661            return dequantize_block_sycl<QK2_0, QR2_0, dequantize_q2_0>;
662        case GGML_TYPE_Q4_0:
663            if (dst->src[0]->extra &&
664                ((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) {
665                return dequantize_row_q4_0_sycl_reorder;
666            } else {
667                return dequantize_block_sycl<QK4_0, QR4_0, dequantize_q4_0>;
668            }
669        case GGML_TYPE_Q4_1:
670            return dequantize_block_sycl<QK4_1, QR4_1, dequantize_q4_1>;
671        case GGML_TYPE_Q5_0:
672            return dequantize_block_sycl<QK5_0, QR5_0, dequantize_q5_0>;
673        case GGML_TYPE_Q5_1:
674            return dequantize_block_sycl<QK5_1, QR5_1, dequantize_q5_1>;
675        case GGML_TYPE_Q8_0:
676            if (dst->src[0]->extra &&
677                ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
678                return dequantize_row_q8_0_sycl_reorder;
679            } else {
680                return dequantize_block_sycl<QK8_0, QR8_0, dequantize_q8_0>;
681            }
682        case GGML_TYPE_Q2_K:
683            if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
684                return dequantize_row_q2_K_sycl_reorder;
685            } else {
686                return dequantize_row_q2_K_sycl;
687            }
688        case GGML_TYPE_Q3_K:
689            if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
690                return dequantize_row_q3_K_sycl_reorder;
691            } else {
692                return dequantize_row_q3_K_sycl;
693            }
694        case GGML_TYPE_Q4_K:
695            if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
696                return dequantize_row_q4_K_sycl_reorder;
697            } else {
698                return dequantize_row_q4_K_sycl;
699            }
700        case GGML_TYPE_Q5_K:
701            if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
702                return dequantize_row_q5_K_sycl_reorder;
703            } else {
704                return dequantize_row_q5_K_sycl;
705            }
706        case GGML_TYPE_Q6_K:
707            if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
708                return dequantize_row_q6_K_sycl_reorder;
709            } else {
710                return dequantize_row_q6_K_sycl;
711            }
712        case GGML_TYPE_IQ1_S:
713            return dequantize_row_iq1_s_sycl;
714        case GGML_TYPE_IQ1_M:
715            return dequantize_row_iq1_m_sycl;
716        case GGML_TYPE_IQ2_XXS:
717            return dequantize_row_iq2_xxs_sycl;
718        case GGML_TYPE_IQ2_XS:
719            return dequantize_row_iq2_xs_sycl;
720        case GGML_TYPE_IQ2_S:
721            return dequantize_row_iq2_s_sycl;
722        case GGML_TYPE_IQ3_XXS:
723            return dequantize_row_iq3_xxs_sycl;
724        case GGML_TYPE_IQ3_S:
725            return dequantize_row_iq3_s_sycl;
726        case GGML_TYPE_IQ4_XS:
727            return dequantize_row_iq4_xs_sycl;
728        case GGML_TYPE_IQ4_NL:
729            return dequantize_row_iq4_nl_sycl;
730        case GGML_TYPE_MXFP4:
731            return dequantize_row_mxfp4_sycl;
732        case GGML_TYPE_NVFP4:
733            return dequantize_row_nvfp4_sycl;
734        case GGML_TYPE_F32:
735            return convert_unary_sycl<float>;
736#ifdef GGML_SYCL_HAS_BF16
737        case GGML_TYPE_BF16:
738            return convert_unary_sycl<sycl::ext::oneapi::bfloat16>;
739#endif
740        default:
741            GGML_ABORT("fatal error: unsupport data type=%s\n", ggml_type_name(type));
742            return nullptr;
743    }
744}
745
746to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor *dst) {
747    switch (type) {
748        case GGML_TYPE_Q1_0:
749            return dequantize_block_sycl<QK1_0, QR1_0, dequantize_q1_0>;
750        case GGML_TYPE_Q2_0:
751            return dequantize_block_sycl<QK2_0, QR2_0, dequantize_q2_0>;
752        case GGML_TYPE_Q4_0:
753            if (dst->src[0]->extra &&
754                ((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) {
755                return dequantize_row_q4_0_sycl_reorder;
756            } else {
757                return dequantize_row_q4_0_sycl;
758            }
759        case GGML_TYPE_Q4_1:
760            return dequantize_row_q4_1_sycl;
761        case GGML_TYPE_Q5_0:
762            return dequantize_block_sycl<QK5_0, QR5_0, dequantize_q5_0>;
763        case GGML_TYPE_Q5_1:
764            return dequantize_block_sycl<QK5_1, QR5_1, dequantize_q5_1>;
765        case GGML_TYPE_Q8_0:
766            if (dst->src[0]->extra &&
767                ((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) {
768                return dequantize_row_q8_0_sycl_reorder;
769            } else {
770                return dequantize_block_sycl<QK8_0, QR8_0, dequantize_q8_0>;
771            }
772        case GGML_TYPE_Q2_K:
773            if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
774                return dequantize_row_q2_K_sycl_reorder;
775            } else {
776                return dequantize_row_q2_K_sycl;
777            }
778        case GGML_TYPE_Q3_K:
779            if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
780                return dequantize_row_q3_K_sycl_reorder;
781            } else {
782                return dequantize_row_q3_K_sycl;
783            }
784        case GGML_TYPE_Q4_K:
785            if (dst->src[0]->extra &&
786                ((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) {
787                return dequantize_row_q4_K_sycl_reorder;
788            } else {
789                return dequantize_row_q4_K_sycl;
790            }
791        case GGML_TYPE_Q5_K:
792            if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
793                return dequantize_row_q5_K_sycl_reorder;
794            } else {
795                return dequantize_row_q5_K_sycl;
796            }
797        case GGML_TYPE_Q6_K:
798            if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
799                return dequantize_row_q6_K_sycl_reorder;
800            } else {
801                return dequantize_row_q6_K_sycl;
802            }
803        case GGML_TYPE_IQ1_S:
804            return dequantize_row_iq1_s_sycl;
805        case GGML_TYPE_IQ1_M:
806            return dequantize_row_iq1_m_sycl;
807        case GGML_TYPE_IQ2_XXS:
808            return dequantize_row_iq2_xxs_sycl;
809        case GGML_TYPE_IQ2_XS:
810            return dequantize_row_iq2_xs_sycl;
811        case GGML_TYPE_IQ2_S:
812            return dequantize_row_iq2_s_sycl;
813        case GGML_TYPE_IQ3_XXS:
814            return dequantize_row_iq3_xxs_sycl;
815        case GGML_TYPE_IQ3_S:
816            return dequantize_row_iq3_s_sycl;
817        case GGML_TYPE_IQ4_XS:
818            return dequantize_row_iq4_xs_sycl;
819        case GGML_TYPE_IQ4_NL:
820            return dequantize_row_iq4_nl_sycl;
821        case GGML_TYPE_MXFP4:
822            return dequantize_row_mxfp4_sycl;
823        case GGML_TYPE_NVFP4:
824            return dequantize_row_nvfp4_sycl;
825        case GGML_TYPE_F16:
826            return convert_unary_sycl<sycl::half>;
827#ifdef GGML_SYCL_HAS_BF16
828        case GGML_TYPE_BF16:
829            return convert_unary_sycl<sycl::ext::oneapi::bfloat16>;
830#endif
831        default:
832            GGML_ABORT("fatal error: unsupport data type=%s\n", ggml_type_name(type));
833            return nullptr;
834    }
835}
836
837
838#ifdef GGML_SYCL_HAS_BF16
839to_bf16_sycl_t ggml_get_to_bf16_sycl(ggml_type type, ggml_tensor * /*dst*/) {
840    switch (type) {
841        case GGML_TYPE_F32:
842            return convert_unary_sycl<float>;
843        case GGML_TYPE_F16:
844            return convert_unary_sycl<sycl::half>;
845        case GGML_TYPE_BF16:
846            return convert_unary_sycl<sycl::ext::oneapi::bfloat16>;
847        default:
848            GGML_ABORT("fatal error: unsupport data type=%s\n", ggml_type_name(type));
849            return nullptr;
850    }
851}
852#endif
853
854to_fp16_nc_sycl_t ggml_get_to_fp16_nc_sycl(ggml_type type) {
855    switch (type) {
856        case GGML_TYPE_F32:
857            return convert_unary_nc_sycl<float>;
858#ifdef GGML_SYCL_HAS_BF16
859        case GGML_TYPE_BF16:
860            return convert_unary_nc_sycl<sycl::ext::oneapi::bfloat16>;
861#endif
862        case GGML_TYPE_Q1_0:
863            return dequantize_block_nc_sycl<QK1_0, QR1_0, dequantize_q1_0>;
864        case GGML_TYPE_Q4_0:
865            return dequantize_block_nc_sycl<QK4_0, QR4_0, dequantize_q4_0>;
866        case GGML_TYPE_Q4_1:
867            return dequantize_block_nc_sycl<QK4_1, QR4_1, dequantize_q4_1>;
868        case GGML_TYPE_Q5_0:
869            return dequantize_block_nc_sycl<QK5_0, QR5_0, dequantize_q5_0>;
870        case GGML_TYPE_Q5_1:
871            return dequantize_block_nc_sycl<QK5_1, QR5_1, dequantize_q5_1>;
872        case GGML_TYPE_Q8_0:
873            return dequantize_block_nc_sycl<QK8_0, QR8_0, dequantize_q8_0>;
874        default:
875            return nullptr;
876    }
877}