char-fork/llama.cpp

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

Ian Faustsycl: fix multiple warnings in compiling sycl backend (#26713)6602dd338

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
18.2 KiB400 linesraw
1//
2// MIT license
3// Copyright (C) 2026 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 "im2col.hpp"
14
15#define MAX_GRIDDIM_Z 65535
16
17template <typename T>
18static  void im2col_kernel(
19        const float * x, T * dst,
20        int64_t IC, int64_t IW, int64_t IH, int64_t OH, int64_t OW, int64_t KW, int64_t KH,
21        int64_t IC_IH_IW, int64_t IH_IW, int64_t N_OH, int64_t KH_KW, int64_t IC_KH_KW,
22        int s0, int s1, int p0, int p1, int d0, int d1) {
23    auto          item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
24    const int64_t i        = item_ct1.get_local_id(2) + item_ct1.get_group(2) * item_ct1.get_local_range(2);
25    if (i >= IC_KH_KW) {
26        return;
27    }
28
29    const int64_t iic = i / (KH_KW);
30    const int64_t rem = i - iic * KH_KW;
31    const int64_t ikh = rem / KW;
32    const int64_t ikw = rem - ikh * KW;
33
34    const int64_t iow = item_ct1.get_group(1);
35    for (int64_t iz = item_ct1.get_group(0); iz < N_OH; iz += MAX_GRIDDIM_Z) {
36        const int64_t  in = iz / OH;
37        const int64_t  ioh = iz - in * OH;
38
39        const int64_t iiw = iow * s0 + ikw * d0 - p0;
40        const int64_t iih = ioh * s1 + ikh * d1 - p1;
41
42        const int64_t offset_dst =
43            ((in * OH + ioh) * OW + iow) * IC_KH_KW + iic * KH_KW + ikh * KW + ikw;
44
45        if (iih < 0 || iih >= IH || iiw < 0 || iiw >= IW) {
46            dst[offset_dst] = 0.0f;
47        } else {
48            const int64_t offset_src = iic * IC_IH_IW + in * IH_IW;
49            dst[offset_dst] = x[offset_src + iih * IW + iiw];
50        }
51    }
52
53    GGML_UNUSED(IC);
54    GGML_UNUSED(KH);
55}
56
57// im2col: [N, IC, IH, IW] => [N, OH, OW, IC*KH*KW]
58template <typename T>
59static void im2col_sycl(const float *   x,
60                        T *             dst,
61                        int64_t         IW,
62                        int64_t         IH,
63                        int64_t         OW,
64                        int64_t         OH,
65                        int64_t         KW,
66                        int64_t         KH,
67                        int64_t         IC,
68                        int64_t         N,
69                        int64_t         IC_IH_IW,
70                        int64_t         IH_IW,
71                        int             s0,
72                        int             s1,
73                        int             p0,
74                        int             p1,
75                        int             d0,
76                        int             d1,
77                        dpct::queue_ptr stream) {
78    const int64_t IC_KH_KW = IC * KH * KW;
79    const int64_t num_blocks = (IC_KH_KW + SYCL_IM2COL_BLOCK_SIZE - 1) / SYCL_IM2COL_BLOCK_SIZE;
80    const int64_t N_OH = N * OH;
81    const int64_t KH_KW = KW*KH;
82    dpct::dim3    block_nums(num_blocks, OW, MIN(N_OH, MAX_GRIDDIM_Z));
83    /*
84    DPCT1049:73: The work-group size passed to the SYCL kernel may exceed the limit. To get the device limit, query info::device::max_work_group_size. Adjust the work-group size if needed.
85    */
86    stream->parallel_for(sycl::nd_range<3>(block_nums * sycl::range<3>(1, 1, MIN(IC_KH_KW, SYCL_IM2COL_BLOCK_SIZE)),
87                                           sycl::range<3>(1, 1, MIN(IC_KH_KW, SYCL_IM2COL_BLOCK_SIZE))),
88                         [=](sycl::nd_item<3>) {
89                             im2col_kernel(x, dst, IC, IW, IH, OH, OW, KW, KH, IC_IH_IW, IH_IW, N_OH, KH_KW, IC_KH_KW,
90                                           s0, s1, p0, p1, d0, d1);
91                         });
92}
93
94static void im2col_sycl_f16(const float *   x,
95                            sycl::half *    dst,
96                            int64_t         IW,
97                            int64_t         IH,
98                            int64_t         OW,
99                            int64_t         OH,
100                            int64_t         KW,
101                            int64_t         KH,
102                            int64_t         IC,
103                            int64_t         N,
104                            int64_t         IC_IH_IW,
105                            int64_t         IH_IW,
106                            int             s0,
107                            int             s1,
108                            int             p0,
109                            int             p1,
110                            int             d0,
111                            int             d1,
112                            dpct::queue_ptr stream) {
113    im2col_sycl<sycl::half>(x, dst, IW, IH, OW, OH, KW, KH, IC, N, IC_IH_IW, IH_IW, s0, s1, p0, p1, d0, d1, stream);
114}
115
116static void im2col_sycl_f32(const float *   x,
117                            float *         dst,
118                            int64_t         IW,
119                            int64_t         IH,
120                            int64_t         OW,
121                            int64_t         OH,
122                            int64_t         KW,
123                            int64_t         KH,
124                            int64_t         IC,
125                            int64_t         N,
126                            int64_t         IC_IH_IW,
127                            int64_t         IH_IW,
128                            int             s0,
129                            int             s1,
130                            int             p0,
131                            int             p1,
132                            int             d0,
133                            int             d1,
134                            dpct::queue_ptr stream) {
135    im2col_sycl<float>(x, dst, IW, IH, OW, OH, KW, KH, IC, N, IC_IH_IW, IH_IW, s0, s1, p0, p1, d0, d1, stream);
136}
137
138void ggml_sycl_op_im2col(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
139    const ggml_tensor * src0 = dst->src[0];
140    const ggml_tensor * src1 = dst->src[1];
141    const float * src1_d = (const float *)src1->data;
142    float * dst_d = (float *)dst->data;
143    dpct::queue_ptr     stream = ctx.stream();
144
145    GGML_ASSERT(src1->type == GGML_TYPE_F32);
146    GGML_ASSERT( dst->type == GGML_TYPE_F16 || dst->type == GGML_TYPE_F32);
147
148    const int32_t s0 = ((const int32_t*)(dst->op_params))[0];
149    const int32_t s1 = ((const int32_t*)(dst->op_params))[1];
150    const int32_t p0 = ((const int32_t*)(dst->op_params))[2];
151    const int32_t p1 = ((const int32_t*)(dst->op_params))[3];
152    const int32_t d0 = ((const int32_t*)(dst->op_params))[4];
153    const int32_t d1 = ((const int32_t*)(dst->op_params))[5];
154
155    const bool is_2D = ((const int32_t*)(dst->op_params))[6] == 1;
156
157    const int64_t IC = src1->ne[is_2D ? 2 : 1];
158    const int64_t IH = is_2D ? src1->ne[1] : 1;
159    const int64_t IW =         src1->ne[0];
160
161    const int64_t KH = is_2D ? src0->ne[1] : 1;
162    const int64_t KW =         src0->ne[0];
163
164    const int64_t OH = is_2D ? dst->ne[2] : 1;
165    const int64_t OW =         dst->ne[1];
166
167    const int64_t IC_IH_IW = src1->nb[is_2D ? 2 : 1] / 4; // nb is byte offset, src is type float32
168    const int64_t N        = src1->ne[is_2D ? 3 : 2];
169    const int64_t IH_IW    = src1->nb[is_2D ? 3 : 2] / 4; // nb is byte offset, src is type float32
170
171    if(dst->type == GGML_TYPE_F16) {
172        im2col_sycl_f16(src1_d, (sycl::half *) dst_d, IW, IH, OW, OH, KW, KH, IC, N, IC_IH_IW, IH_IW, s0, s1, p0, p1,
173                        d0, d1, stream);
174    } else {
175        im2col_sycl_f32(src1_d, (float *) dst_d, IW, IH, OW, OH, KW, KH, IC, N, IC_IH_IW, IH_IW, s0, s1, p0, p1, d0, d1, stream);
176    }
177}
178
179// [N*IC, ID, IH, IW] => [N*OD, OH, OW, IC * KD * KH * KW]
180template <typename T>
181static  void im2col_3d_kernel(
182        const float * src, T * dst,
183        int64_t N, int64_t IC, int64_t ID, int64_t IH, int64_t IW, int64_t OC,
184        int64_t KD, int64_t KH, int64_t KW, int64_t OD, int64_t OH, int64_t OW,
185        int64_t OH_OW, int64_t KD_KH_KW, int64_t ID_IH_IW, int64_t KH_KW, int64_t IH_IW, int64_t IC_ID_IH_IW,
186        int64_t IC_KD_KH_KW, int64_t OW_KD_KH_KW, int64_t OD_OH_OW_IC_KD_KH_KW, int64_t OH_OW_IC_KD_KH_KW,
187        int64_t OW_IC_KD_KH_KW, int64_t N_OD_OH, int64_t OD_OH,
188        int64_t stride_q, int64_t stride_z, int64_t stride_y, int64_t stride_x,
189        int s0, int s1, int s2, int p0, int p1, int p2, int d0, int d1, int d2) {
190    auto          item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
191    const int64_t i        = item_ct1.get_local_id(2) + item_ct1.get_group(2) * item_ct1.get_local_range(2);
192    if (i >= IC_KD_KH_KW) {
193        return;
194    }
195    GGML_UNUSED(N); GGML_UNUSED(OC); GGML_UNUSED(OH_OW); GGML_UNUSED(OD); GGML_UNUSED(OW); GGML_UNUSED(KD); GGML_UNUSED(KH);
196    GGML_UNUSED(ID_IH_IW); GGML_UNUSED(IH_IW); GGML_UNUSED(IC_ID_IH_IW); GGML_UNUSED(OW_KD_KH_KW);
197
198    const int64_t iic = i / KD_KH_KW;
199    const int64_t ikd = (i - iic * KD_KH_KW) / KH_KW;
200    const int64_t ikh = (i - iic * KD_KH_KW - ikd * KH_KW) / KW;
201    const int64_t ikw = i % KW;
202
203    const int64_t iow = item_ct1.get_group(1);
204    for (int64_t iz = item_ct1.get_group(0); iz < N_OD_OH; iz += MAX_GRIDDIM_Z) {
205        const int64_t in  = iz / OD_OH;
206        const int64_t iod = (iz - in*OD_OH) / OH;
207        const int64_t ioh = iz % OH;
208
209        const int64_t iiw = iow * s0 + ikw * d0 - p0;
210        const int64_t iih = ioh * s1 + ikh * d1 - p1;
211        const int64_t iid = iod * s2 + ikd * d2 - p2;
212
213        const int64_t offset_dst = in*OD_OH_OW_IC_KD_KH_KW + iod*OH_OW_IC_KD_KH_KW + ioh*OW_IC_KD_KH_KW + iow*IC_KD_KH_KW + iic*KD_KH_KW + ikd * KH_KW + ikh*KW + ikw;
214
215        if (iih < 0 || iih >= IH || iiw < 0 || iiw >= IW || iid < 0 || iid >= ID) {
216            dst[offset_dst] = 0.0f;
217        } else {
218            const int64_t offset_src = ((in * IC + iic) * stride_q) + (iid * stride_z) + (iih * stride_y) + (iiw * stride_x);
219            dst[offset_dst] = src[offset_src];
220        }
221    }
222}
223
224// [N*IC, ID, IH, IW] => [N*OD, OH, OW, IC * KD * KH * KW]
225template <typename T>
226static void im2col_3d_sycl(const float *   src,
227                           T *             dst,
228                           int64_t         N,
229                           int64_t         IC,
230                           int64_t         ID,
231                           int64_t         IH,
232                           int64_t         IW,
233                           int64_t         OC,
234                           int64_t         KD,
235                           int64_t         KH,
236                           int64_t         KW,
237                           int64_t         OD,
238                           int64_t         OH,
239                           int64_t         OW,
240                           int64_t         stride_q,
241                           int64_t         stride_z,
242                           int64_t         stride_y,
243                           int64_t         stride_x,
244                           int             s0,
245                           int             s1,
246                           int             s2,
247                           int             p0,
248                           int             p1,
249                           int             p2,
250                           int             d0,
251                           int             d1,
252                           int             d2,
253                           dpct::queue_ptr stream) {
254    const int64_t OH_OW = OH*OW;
255    const int64_t KD_KH_KW = KD*KH*KW;
256    const int64_t ID_IH_IW = ID*IH*IW;
257    const int64_t KH_KW = KH*KW;
258    const int64_t IH_IW = IH*IW;
259    const int64_t IC_KD_KH_KW = IC*KD*KH*KW;
260    const int64_t OW_KD_KH_KW = OW*KD*KH*KW;
261    const int64_t N_OD_OH = N*OD*OH;
262    const int64_t OD_OH = OD*OH;
263    const int64_t IC_ID_IH_IW = IC*ID*IH*IW;
264    const int64_t OD_OH_OW_IC_KD_KH_KW = OD*OH*OW*IC*KD*KH*KW;
265    const int64_t OH_OW_IC_KD_KH_KW = OH*OW*IC*KD*KH*KW;
266    const int64_t OW_IC_KD_KH_KW = OW*IC*KD*KH*KW;
267    const int64_t num_blocks = (IC_KD_KH_KW + SYCL_IM2COL_BLOCK_SIZE - 1) / SYCL_IM2COL_BLOCK_SIZE;
268    dpct::dim3    block_nums(num_blocks, OW, MIN(N_OD_OH, MAX_GRIDDIM_Z));
269    /*
270    DPCT1049:74: The work-group size passed to the SYCL kernel may exceed the limit. To get the device limit, query info::device::max_work_group_size. Adjust the work-group size if needed.
271    */
272    stream->parallel_for(sycl::nd_range<3>(block_nums * sycl::range<3>(1, 1, MIN(IC_KD_KH_KW, SYCL_IM2COL_BLOCK_SIZE)),
273                                           sycl::range<3>(1, 1, MIN(IC_KD_KH_KW, SYCL_IM2COL_BLOCK_SIZE))),
274                         [=](sycl::nd_item<3>) {
275                             im2col_3d_kernel(src, dst, N, IC, ID, IH, IW, OC, KD, KH, KW, OD, OH, OW, OH_OW, KD_KH_KW,
276                                              ID_IH_IW, KH_KW, IH_IW, IC_ID_IH_IW, IC_KD_KH_KW, OW_KD_KH_KW,
277                                              OD_OH_OW_IC_KD_KH_KW, OH_OW_IC_KD_KH_KW, OW_IC_KD_KH_KW, N_OD_OH, OD_OH,
278                                              stride_q, stride_z, stride_y, stride_x, s0, s1, s2, p0, p1, p2, d0, d1,
279                                              d2);
280                         });
281}
282
283static void im2col_3d_sycl_f16(const float *   src,
284                               sycl::half *    dst,
285                               int64_t         N,
286                               int64_t         IC,
287                               int64_t         ID,
288                               int64_t         IH,
289                               int64_t         IW,
290                               int64_t         OC,
291                               int64_t         KD,
292                               int64_t         KH,
293                               int64_t         KW,
294                               int64_t         OD,
295                               int64_t         OH,
296                               int64_t         OW,
297                               int64_t         stride_q,
298                               int64_t         stride_z,
299                               int64_t         stride_y,
300                               int64_t         stride_x,
301                               int             s0,
302                               int             s1,
303                               int             s2,
304                               int             p0,
305                               int             p1,
306                               int             p2,
307                               int             d0,
308                               int             d1,
309                               int             d2,
310                               dpct::queue_ptr stream) {
311    im2col_3d_sycl<sycl::half>(src, dst, N, IC, ID, IH, IW, OC, KD, KH, KW, OD, OH, OW, stride_q, stride_z, stride_y,
312                               stride_x, s0, s1, s2, p0, p1, p2, d0, d1, d2, stream);
313}
314
315static void im2col_3d_sycl_f32(const float *   src,
316                               float *         dst,
317                               int64_t         N,
318                               int64_t         IC,
319                               int64_t         ID,
320                               int64_t         IH,
321                               int64_t         IW,
322                               int64_t         OC,
323                               int64_t         KD,
324                               int64_t         KH,
325                               int64_t         KW,
326                               int64_t         OD,
327                               int64_t         OH,
328                               int64_t         OW,
329                               int64_t         stride_q,
330                               int64_t         stride_z,
331                               int64_t         stride_y,
332                               int64_t         stride_x,
333                               int             s0,
334                               int             s1,
335                               int             s2,
336                               int             p0,
337                               int             p1,
338                               int             p2,
339                               int             d0,
340                               int             d1,
341                               int             d2,
342                               dpct::queue_ptr stream) {
343    im2col_3d_sycl<float>(src, dst, N, IC, ID, IH, IW, OC, KD, KH, KW, OD, OH, OW,
344                          stride_q, stride_z, stride_y, stride_x,
345                          s0, s1, s2, p0, p1, p2, d0, d1, d2, stream);
346}
347
348void ggml_sycl_op_im2col_3d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
349    const ggml_tensor * src0 = dst->src[0];
350    const ggml_tensor * src1 = dst->src[1];
351    const float * src1_d = (const float *)src1->data;
352    float * dst_d = (float *)dst->data;
353    dpct::queue_ptr     stream = ctx.stream();
354
355    GGML_ASSERT(src1->type == GGML_TYPE_F32);
356    GGML_ASSERT( dst->type == GGML_TYPE_F16 || dst->type == GGML_TYPE_F32);
357
358    GGML_TENSOR_BINARY_OP_LOCALS
359
360    const int32_t s0 = ((const int32_t *)(dst->op_params))[0];
361    const int32_t s1 = ((const int32_t *)(dst->op_params))[1];
362    const int32_t s2 = ((const int32_t *)(dst->op_params))[2];
363    const int32_t p0 = ((const int32_t *)(dst->op_params))[3];
364    const int32_t p1 = ((const int32_t *)(dst->op_params))[4];
365    const int32_t p2 = ((const int32_t *)(dst->op_params))[5];
366    const int32_t d0 = ((const int32_t *)(dst->op_params))[6];
367    const int32_t d1 = ((const int32_t *)(dst->op_params))[7];
368    const int32_t d2 = ((const int32_t *)(dst->op_params))[8];
369    const int32_t IC = ((const int32_t *)(dst->op_params))[9];
370
371    const int64_t N  = ne13 / IC;
372    const int64_t ID = ne12;
373    const int64_t IH = ne11;
374    const int64_t IW = ne10;
375
376    const int64_t OC = ne03 / IC;
377    const int64_t KD = ne02;
378    const int64_t KH = ne01;
379    const int64_t KW = ne00;
380
381    const int64_t OD = ne3 / N;
382    const int64_t OH = ne2;
383    const int64_t OW = ne1;
384
385    const size_t  es       = ggml_element_size(src1);
386    const int64_t stride_x = src1->nb[0] / es;
387    const int64_t stride_y = src1->nb[1] / es;
388    const int64_t stride_z = src1->nb[2] / es;
389    const int64_t stride_q = src1->nb[3] / es;
390
391    if(dst->type == GGML_TYPE_F16) {
392        im2col_3d_sycl_f16(src1_d, (sycl::half *) dst_d, N, IC, ID, IH, IW, OC, KD, KH, KW, OD, OH, OW,
393                           stride_q, stride_z, stride_y, stride_x,
394                           s0, s1, s2, p0, p1, p2, d0, d1, d2, stream);
395    } else {
396        im2col_3d_sycl_f32(src1_d, (float *) dst_d, N, IC, ID, IH, IW, OC, KD, KH, KW, OD, OH, OW,
397                           stride_q, stride_z, stride_y, stride_x,
398                           s0, s1, s2, p0, p1, p2, d0, d1, d2, stream);
399    }
400}