char-fork/llama.cpp

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

Neo Zhangsycl : fix get_rows Q2_K, Q4_K, Q5_K (#25656)d3fba0c79

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
15.9 KiB368 linesraw
1//
2// MIT license
3// Copyright (C) 2024 Intel Corporation
4// SPDX-License-Identifier: MIT
5//
6
7//
8// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
9// See https://llvm.org/LICENSE.txt for license information.
10// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
11//
12
13#include "ggml-impl.h"
14#include "common.hpp"
15#include "dequantize.hpp"
16#include "getrows.hpp"
17
18
19template<int qk, int qr, dequantize_kernel_t dequantize_kernel, typename dst_t>
20static void k_get_rows(
21            const void * src0, const int32_t * src1, dst_t * dst,
22            int64_t ne00, /*int64_t ne01, int64_t ne02, int64_t ne03,*/
23            /*int64_t ne10, int64_t ne11,*/ int64_t ne12, /*int64_t ne13,*/
24            /*size_t s0,*/ size_t s1, size_t s2, size_t s3,
25            /*size_t nb00,*/ size_t nb01, size_t nb02, size_t nb03,
26            size_t s10, size_t s11, size_t s12,
27            const sycl::nd_item<3> &item_ct1/*, size_t s13*/) {
28
29    const int i00 = (item_ct1.get_group(2) * item_ct1.get_local_range(2) +
30                     item_ct1.get_local_id(2)) *
31                    2;
32    const int i10 = item_ct1.get_local_range(1) * item_ct1.get_group(1) +
33                    item_ct1.get_local_id(1);
34    const int i11 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
35                     item_ct1.get_local_id(0)) /
36                    ne12;
37    const int i12 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
38                     item_ct1.get_local_id(0)) %
39                    ne12;
40
41    if (i00 >= ne00) {
42        return;
43    }
44
45    const int i01 = src1[i10*s10 + i11*s11 + i12*s12];
46
47    dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3;
48    const void * src0_row = (const char *)src0 + i01*nb01 + i11*nb02 + i12*nb03;
49
50    const int ib = i00/qk; // block index
51    const int iqs = (i00%qk)/qr; // quant index
52    const int iybs = i00 - i00%qk; // dst block start index
53    const int y_offset = qr == 1 ? 1 : qk/2;
54
55    // dequantize
56    dfloat2 v;
57    dequantize_kernel(src0_row, ib, iqs, v);
58
59    dst_row[iybs + iqs + 0] = v.x();
60    dst_row[iybs + iqs + y_offset] = v.y();
61}
62
63template<int qk, int qr, dequantize_kernel_f32_t dequantize_kernel, typename dst_t>
64static void k_get_rows_f32(
65            const void * src0, const int32_t * src1, dst_t * dst,
66            int64_t ne00,
67            int64_t ne12,
68            size_t s1, size_t s2, size_t s3,
69            size_t nb01, size_t nb02, size_t nb03,
70            size_t s10, size_t s11, size_t s12,
71            const sycl::nd_item<3> &item_ct1) {
72
73    const int i00 = (item_ct1.get_group(2) * item_ct1.get_local_range(2) +
74                     item_ct1.get_local_id(2)) *
75                    2;
76    const int i10 = item_ct1.get_local_range(1) * item_ct1.get_group(1) +
77                    item_ct1.get_local_id(1);
78    const int i11 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
79                     item_ct1.get_local_id(0)) /
80                    ne12;
81    const int i12 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
82                     item_ct1.get_local_id(0)) %
83                    ne12;
84
85    if (i00 >= ne00) {
86        return;
87    }
88
89    const int i01 = src1[i10*s10 + i11*s11 + i12*s12];
90
91    dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3;
92    const void * src0_row = (const char *)src0 + i01*nb01 + i11*nb02 + i12*nb03;
93
94    const int ib = i00/qk;
95    const int iqs = (i00%qk)/qr;
96    const int iybs = i00 - i00%qk;
97    const int y_offset = qr == 1 ? 1 : qk/2;
98
99    float v0;
100    float v1;
101    dequantize_kernel(src0_row, ib, iqs, v0, v1);
102
103    dst_row[iybs + iqs + 0] = (dst_t) v0;
104    dst_row[iybs + iqs + y_offset] = (dst_t) v1;
105}
106
107template<typename src0_t, typename dst_t>
108static void k_get_rows_float(
109            const src0_t * src0, const int32_t * src1, dst_t * dst,
110            int64_t ne00, /*int64_t ne01, int64_t ne02, int64_t ne03,*/
111            /*int64_t ne10, int64_t ne11,*/ int64_t ne12, /*int64_t ne13,*/
112            /*size_t s0,*/ size_t s1, size_t s2, size_t s3,
113            /*size_t nb00,*/ size_t nb01, size_t nb02, size_t nb03,
114            size_t s10, size_t s11, size_t s12,
115            const sycl::nd_item<3> &item_ct1/*, size_t s13*/) {
116
117    const int i00 = item_ct1.get_group(2) * item_ct1.get_local_range(2) +
118                    item_ct1.get_local_id(2);
119    const int i10 = item_ct1.get_local_range(1) * item_ct1.get_group(1) +
120                    item_ct1.get_local_id(1);
121    const int i11 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
122                     item_ct1.get_local_id(0)) /
123                    ne12;
124    const int i12 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
125                     item_ct1.get_local_id(0)) %
126                    ne12;
127
128    if (i00 >= ne00) {
129        return;
130    }
131
132    const int i01 = src1[i10*s10 + i11*s11 + i12*s12];
133
134    dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3;
135    const src0_t * src0_row = (const src0_t *)((const char *)src0 + i01*nb01 + i11*nb02 + i12*nb03);
136
137    dst_row[i00] = src0_row[i00];
138}
139
140template <int qk, int qr, dequantize_kernel_t dq>
141static void get_rows_sycl(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1,
142                          ggml_tensor *dst, const void *src0_dd,
143                          const int32_t *src1_dd, float *dst_dd,
144                          queue_ptr stream) {
145
146    GGML_TENSOR_BINARY_OP_LOCALS
147
148    const sycl::range<3> block_dims(1, 1, SYCL_GET_ROWS_BLOCK_SIZE);
149    const int block_num_x = (ne00 + 2*SYCL_GET_ROWS_BLOCK_SIZE - 1) / (2*SYCL_GET_ROWS_BLOCK_SIZE);
150    const sycl::range<3> block_nums(ne11 * ne12, ne10, block_num_x);
151
152    // strides in elements
153    //const size_t s0 = nb0 / ggml_element_size(dst);
154    const size_t s1 = nb1 / ggml_element_size(dst);
155    const size_t s2 = nb2 / ggml_element_size(dst);
156    const size_t s3 = nb3 / ggml_element_size(dst);
157
158    const size_t s10 = nb10 / ggml_element_size(src1);
159    const size_t s11 = nb11 / ggml_element_size(src1);
160    const size_t s12 = nb12 / ggml_element_size(src1);
161    //const size_t s13 = nb13 / ggml_element_size(src1);
162
163    GGML_ASSERT(ne00 % 2 == 0);
164
165    stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
166                         [=](sycl::nd_item<3> item_ct1) {
167                             k_get_rows<qk, qr, dq>(
168                                 src0_dd, src1_dd, dst_dd, ne00, ne12, s1, s2,
169                                 s3, nb01, nb02, nb03, s10, s11, s12, item_ct1);
170                         });
171
172    GGML_UNUSED(dst);
173    GGML_UNUSED(ctx);
174}
175
176template <int qk, int qr, dequantize_kernel_f32_t dq>
177static void get_rows_sycl_f32(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1,
178                              ggml_tensor *dst, const void *src0_dd,
179                              const int32_t *src1_dd, float *dst_dd,
180                              queue_ptr stream) {
181
182    GGML_TENSOR_BINARY_OP_LOCALS
183
184    const sycl::range<3> block_dims(1, 1, SYCL_GET_ROWS_BLOCK_SIZE);
185    const int block_num_x = (ne00 + 2*SYCL_GET_ROWS_BLOCK_SIZE - 1) / (2*SYCL_GET_ROWS_BLOCK_SIZE);
186    const sycl::range<3> block_nums(ne11 * ne12, ne10, block_num_x);
187
188    const size_t s1 = nb1 / ggml_element_size(dst);
189    const size_t s2 = nb2 / ggml_element_size(dst);
190    const size_t s3 = nb3 / ggml_element_size(dst);
191
192    const size_t s10 = nb10 / ggml_element_size(src1);
193    const size_t s11 = nb11 / ggml_element_size(src1);
194    const size_t s12 = nb12 / ggml_element_size(src1);
195
196    GGML_ASSERT(ne00 % 2 == 0);
197
198    stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
199                         [=](sycl::nd_item<3> item_ct1) {
200                             k_get_rows_f32<qk, qr, dq>(
201                                 src0_dd, src1_dd, dst_dd, ne00, ne12, s1, s2,
202                                 s3, nb01, nb02, nb03, s10, s11, s12, item_ct1);
203                         });
204
205    GGML_UNUSED(dst);
206    GGML_UNUSED(ctx);
207}
208
209template <typename src0_t, typename dst_t>
210static void get_rows_sycl_float(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
211                                const ggml_tensor *src1, ggml_tensor *dst,
212                                const src0_t *src0_dd, const int32_t *src1_dd,
213                                dst_t *dst_dd, queue_ptr stream) {
214
215    GGML_TENSOR_BINARY_OP_LOCALS
216
217    const sycl::range<3> block_dims(1, 1, SYCL_GET_ROWS_BLOCK_SIZE);
218    const int block_num_x = (ne00 + SYCL_GET_ROWS_BLOCK_SIZE - 1) / SYCL_GET_ROWS_BLOCK_SIZE;
219    const sycl::range<3> block_nums(ne11 * ne12, ne10, block_num_x);
220
221    // strides in elements
222    //const size_t s0 = nb0 / ggml_element_size(dst);
223    const size_t s1 = nb1 / ggml_element_size(dst);
224    const size_t s2 = nb2 / ggml_element_size(dst);
225    const size_t s3 = nb3 / ggml_element_size(dst);
226
227    const size_t s10 = nb10 / ggml_element_size(src1);
228    const size_t s11 = nb11 / ggml_element_size(src1);
229    const size_t s12 = nb12 / ggml_element_size(src1);
230    //const size_t s13 = nb13 / ggml_element_size(src1);
231
232    {
233        dpct::has_capability_or_fail(stream->get_device(),
234                                     {sycl::aspect::fp16});
235
236        stream->parallel_for(
237            sycl::nd_range<3>(block_nums * block_dims, block_dims),
238            [=](sycl::nd_item<3> item_ct1) {
239                k_get_rows_float(src0_dd, src1_dd, dst_dd, ne00, ne12, s1, s2,
240                                 s3, nb01, nb02, nb03, s10, s11, s12, item_ct1);
241            });
242    }
243
244    GGML_UNUSED(dst);
245    GGML_UNUSED(ctx);
246}
247
248void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
249    GGML_ASSERT(dst->src[1]->type == GGML_TYPE_I32);
250    GGML_ASSERT(dst->type == GGML_TYPE_F32 || dst->type == GGML_TYPE_I32 );
251
252    GGML_ASSERT(dst->src[0]->nb[0] == ggml_type_size(dst->src[0]->type));
253    GGML_ASSERT(dst->src[1]->nb[0] == ggml_type_size(dst->src[1]->type));
254    GGML_ASSERT(dst->nb[0] == ggml_type_size(dst->type));
255
256    const int32_t * src1_i32 = (const int32_t *) dst->src[1]->data;
257    /* TODO: Refactor and remove duplicates */
258    switch (dst->src[0]->type) {
259        case GGML_TYPE_F16:
260            get_rows_sycl_float(ctx, dst->src[0], dst->src[1], dst, (const sycl::half *)dst->src[0]->data,
261                                src1_i32, (float *)dst->data, ctx.stream());
262            break;
263        case GGML_TYPE_BF16:
264            get_rows_sycl_float(ctx, dst->src[0], dst->src[1], dst, (const sycl::ext::oneapi::bfloat16 *)dst->src[0]->data,
265                                src1_i32, (float *)dst->data, ctx.stream());
266            break;
267        case GGML_TYPE_F32:
268            get_rows_sycl_float(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
269            src1_i32, (float *)dst->data, ctx.stream());
270            break;
271        case GGML_TYPE_I32:
272            get_rows_sycl_float(ctx, dst->src[0], dst->src[1], dst, (const int32_t *)dst->src[0]->data,
273            src1_i32, (int32_t *)dst->data, ctx.stream());
274            break;
275        case GGML_TYPE_Q1_0:
276            get_rows_sycl<QK1_0, 1, dequantize_q1_0>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
277            src1_i32, (float *)dst->data, ctx.stream());
278            break;
279        case GGML_TYPE_MXFP4:
280            get_rows_sycl<QK_MXFP4, 2, dequantize_mxfp4>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
281            src1_i32, (float *)dst->data, ctx.stream());
282            break;
283        case GGML_TYPE_NVFP4:
284            get_rows_sycl<QK_NVFP4, 1, dequantize_nvfp4>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
285            src1_i32, (float *)dst->data, ctx.stream());
286            break;
287        case GGML_TYPE_IQ2_XXS:
288            get_rows_sycl<QK_K, 1, dequantize_iq2_xxs>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
289            src1_i32, (float *)dst->data, ctx.stream());
290            break;
291        case GGML_TYPE_IQ2_XS:
292            get_rows_sycl<QK_K, 1, dequantize_iq2_xs>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
293            src1_i32, (float *)dst->data, ctx.stream());
294            break;
295        case GGML_TYPE_IQ2_S:
296            get_rows_sycl<QK_K, 1, dequantize_iq2_s>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
297            src1_i32, (float *)dst->data, ctx.stream());
298            break;
299        case GGML_TYPE_IQ3_XXS:
300            get_rows_sycl<QK_K, 1, dequantize_iq3_xxs>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
301            src1_i32, (float *)dst->data, ctx.stream());
302            break;
303        case GGML_TYPE_IQ1_S:
304            get_rows_sycl<QK_K, 1, dequantize_iq1_s>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
305            src1_i32, (float *)dst->data, ctx.stream());
306            break;
307        case GGML_TYPE_IQ1_M:
308            get_rows_sycl<QK_K, 1, dequantize_iq1_m>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
309            src1_i32, (float *)dst->data, ctx.stream());
310            break;
311        case GGML_TYPE_IQ3_S:
312            get_rows_sycl<QK_K, 1, dequantize_iq3_s>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
313            src1_i32, (float *)dst->data, ctx.stream());
314            break;
315        case GGML_TYPE_IQ4_NL:
316            get_rows_sycl<QK4_NL, 1, dequantize_iq4_nl>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
317            src1_i32, (float *)dst->data, ctx.stream());
318            break;
319        case GGML_TYPE_IQ4_XS:
320            get_rows_sycl<QK_K, 1, dequantize_iq4_xs>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
321            src1_i32, (float *)dst->data, ctx.stream());
322            break;
323        case GGML_TYPE_Q2_K:
324            get_rows_sycl_f32<QK_K, 1, dequantize_q2_K_f32>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
325            src1_i32, (float *)dst->data, ctx.stream());
326            break;
327        case GGML_TYPE_Q3_K:
328            get_rows_sycl<QK_K, 1, dequantize_q3_K>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
329            src1_i32, (float *)dst->data, ctx.stream());
330            break;
331        case GGML_TYPE_Q4_0:
332            get_rows_sycl<QK4_0, QR4_0, dequantize_q4_0>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
333            src1_i32, (float *)dst->data, ctx.stream());
334            break;
335        case GGML_TYPE_Q4_1:
336            get_rows_sycl<QK4_1, QR4_1, dequantize_q4_1>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
337            src1_i32, (float *)dst->data, ctx.stream());
338            break;
339        case GGML_TYPE_Q4_K:
340            get_rows_sycl_f32<QK_K, 1, dequantize_q4_K_f32>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
341            src1_i32, (float *)dst->data, ctx.stream());
342            break;
343        case GGML_TYPE_Q5_0:
344            get_rows_sycl<QK5_0, QR5_0, dequantize_q5_0>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
345            src1_i32, (float *)dst->data, ctx.stream());
346            break;
347        case GGML_TYPE_Q5_1:
348            get_rows_sycl<QK5_1, QR5_1, dequantize_q5_1>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
349            src1_i32, (float *)dst->data, ctx.stream());
350            break;
351        case GGML_TYPE_Q5_K:
352            get_rows_sycl_f32<QK_K, 1, dequantize_q5_K_f32>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
353            src1_i32, (float *)dst->data, ctx.stream());
354            break;
355        case GGML_TYPE_Q6_K:
356            get_rows_sycl<QK_K, 1, dequantize_q6_K>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
357            src1_i32, (float *)dst->data, ctx.stream());
358            break;
359        case GGML_TYPE_Q8_0:
360            get_rows_sycl<QK8_0, QR8_0, dequantize_q8_0>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
361            src1_i32, (float *)dst->data, ctx.stream());
362            break;
363        default:
364            // TODO: k-quants
365            GGML_LOG_ERROR("%s: unsupported type: %s\n", __func__, ggml_type_name(dst->src[0]->type));
366            GGML_ABORT("fatal error");
367    }
368}