char-fork/llama.cpp

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

Neo Zhangsycl : enhance concat to support Q4_0, Q4_1, Q5_0, Q5_1, Q8_0 (#26800)d415e65a5

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
20.9 KiB509 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.h"
14
15#include "concat.hpp"
16
17static inline size_t elem_size(ggml_type t) {
18    return ggml_type_size(t) / ggml_blck_size(t);
19}
20
21template <typename T>
22static void concat_T_dim0(const T *x, const T *y, T *dst,
23                            const int ne0, const int ne00,
24                            const sycl::nd_item<3> &item_ct1) {
25  int nidx = item_ct1.get_local_id(2) +
26             item_ct1.get_group(2) * item_ct1.get_local_range(2);
27  if (nidx >= ne0) {
28    return;
29  }
30  // operation
31  int offset_dst = nidx + item_ct1.get_group(1) * ne0 +
32                   item_ct1.get_group(0) * ne0 * item_ct1.get_group_range(1);
33  if (nidx < ne00) { // src0
34    int offset_src = nidx + item_ct1.get_group(1) * ne00 +
35                     item_ct1.get_group(0) * ne00 * item_ct1.get_group_range(1);
36    dst[offset_dst] = x[offset_src];
37  } else {
38    int offset_src =
39        nidx - ne00 + item_ct1.get_group(1) * (ne0 - ne00) +
40        item_ct1.get_group(0) * (ne0 - ne00) * item_ct1.get_group_range(1);
41    dst[offset_dst] = y[offset_src];
42  }
43}
44
45template <typename T>
46static void concat_T_dim1(const T *x, const T *y, T *dst,
47                            const int ne0, const int ne01,
48                            const sycl::nd_item<3> &item_ct1) {
49  int nidx = item_ct1.get_local_id(2) +
50             item_ct1.get_group(2) * item_ct1.get_local_range(2);
51  if (nidx >= ne0) {
52    return;
53  }
54  // operation
55  int offset_dst = nidx + item_ct1.get_group(1) * ne0 +
56                   item_ct1.get_group(0) * ne0 * item_ct1.get_group_range(1);
57  if (item_ct1.get_group(1) < (size_t) ne01) { // src0
58    int offset_src =
59        nidx + item_ct1.get_group(1) * ne0 + item_ct1.get_group(0) * ne0 * ne01;
60    dst[offset_dst] = x[offset_src];
61  } else {
62    int offset_src =
63        nidx + (item_ct1.get_group(1) - ne01) * ne0 +
64        item_ct1.get_group(0) * ne0 * (item_ct1.get_group_range(1) - ne01);
65    dst[offset_dst] = y[offset_src];
66  }
67}
68
69template <typename T>
70static void concat_T_dim2(const T *x, const T *y, T *dst,
71                            const int ne0, const int ne02,
72                            const sycl::nd_item<3> &item_ct1) {
73  int nidx = item_ct1.get_local_id(2) +
74             item_ct1.get_group(2) * item_ct1.get_local_range(2);
75  if (nidx >= ne0) {
76    return;
77  }
78  // operation
79  int offset_dst = nidx + item_ct1.get_group(1) * ne0 +
80                   item_ct1.get_group(0) * ne0 * item_ct1.get_group_range(1);
81  if (item_ct1.get_group(0) < (size_t) ne02) { // src0
82    int offset_src = nidx + item_ct1.get_group(1) * ne0 +
83                     item_ct1.get_group(0) * ne0 * item_ct1.get_group_range(1);
84    dst[offset_dst] = x[offset_src];
85  } else {
86    int offset_src =
87        nidx + item_ct1.get_group(1) * ne0 +
88        (item_ct1.get_group(0) - ne02) * ne0 * item_ct1.get_group_range(1);
89    dst[offset_dst] = y[offset_src];
90  }
91}
92
93template <typename T>
94static void concat_T_sycl(const T *x, const T *y, T *dst,
95                            int ne00, int ne01, int ne02, int ne0, int ne1,
96                            int ne2, int dim, queue_ptr stream) {
97  int num_blocks = (ne0 + SYCL_CONCAT_BLOCK_SIZE - 1) / SYCL_CONCAT_BLOCK_SIZE;
98  sycl::range<3> gridDim(ne2, ne1, num_blocks);
99  switch (dim) {
100  case 0:
101      stream->parallel_for(sycl::nd_range<3>(gridDim * sycl::range<3>(1, 1, SYCL_CONCAT_BLOCK_SIZE),
102                                          sycl::range<3>(1, 1, SYCL_CONCAT_BLOCK_SIZE)),
103                        [=](sycl::nd_item<3> item_ct1) { concat_T_dim0<T>(x, y, dst, ne0, ne00, item_ct1); });
104      break;
105  case 1:
106      stream->parallel_for(sycl::nd_range<3>(gridDim * sycl::range<3>(1, 1, SYCL_CONCAT_BLOCK_SIZE),
107                                          sycl::range<3>(1, 1, SYCL_CONCAT_BLOCK_SIZE)),
108                        [=](sycl::nd_item<3> item_ct1) { concat_T_dim1<T>(x, y, dst, ne0, ne01, item_ct1); });
109      break;
110  // dim >=2 will be dispatched to the default path
111  default:
112      stream->parallel_for(sycl::nd_range<3>(gridDim * sycl::range<3>(1, 1, SYCL_CONCAT_BLOCK_SIZE),
113                                          sycl::range<3>(1, 1, SYCL_CONCAT_BLOCK_SIZE)),
114                        [=](sycl::nd_item<3> item_ct1) { concat_T_dim2<T>(x, y, dst, ne0, ne02, item_ct1); });
115      break;
116  }
117}
118
119// non-contiguous kernel (slow)
120template<typename T>
121static void concat_T_sycl_non_cont(
122    queue_ptr stream, const char *src0, const char *src1, char *dst,
123    int64_t ne00, int64_t ne01, int64_t ne02, int64_t ne03, uint64_t nb00,
124    uint64_t nb01, uint64_t nb02, uint64_t nb03, int64_t /*ne10*/,
125    int64_t /*ne11*/, int64_t /*ne12*/, int64_t /*ne13*/, uint64_t nb10,
126    uint64_t nb11, uint64_t nb12, uint64_t nb13, int64_t ne0, int64_t ne1,
127    int64_t ne2, int64_t ne3, uint64_t nb0, uint64_t nb1, uint64_t nb2,
128    uint64_t nb3, int32_t dim) {
129  sycl::range<3> gridDim(ne3, ne2, ne1);
130
131  // Avoid oversubscribing device when there is not enough elements along the innermost dim to
132  // fill a full SYCL_CONCAT_BLOCK_SIZE. For larger # of elements, the full SYCL_CONCAT_BLOCK_SIZE
133  // is used.
134  const int64_t ne0_pad   = GGML_PAD(ne0, WARP_SIZE);
135  const int64_t block_ne0 = ne0_pad < SYCL_CONCAT_BLOCK_SIZE ? ne0_pad : (int64_t) SYCL_CONCAT_BLOCK_SIZE;
136  sycl::range<3> blockDim(1, 1, block_ne0);
137
138  stream->parallel_for(sycl::nd_range<3>(gridDim * blockDim, blockDim), [=](sycl::nd_item<3> item_ct1) {
139      int64_t i3 = item_ct1.get_group(0);
140      int64_t i2 = item_ct1.get_group(1);
141      int64_t i1 = item_ct1.get_group(2);
142
143      int64_t o[4] = { 0, 0, 0, 0 };
144      o[dim]       = dim == 0 ? ne00 : (dim == 1 ? ne01 : (dim == 2 ? ne02 : ne03));
145
146      const T * x;
147
148      for (int i0 = item_ct1.get_local_id(2); i0 < ne0; i0 += item_ct1.get_local_range(2)) {
149          if (i0 < ne00 && i1 < ne01 && i2 < ne02 && i3 < ne03) {
150              x = (const T *) (src0 + (i3) *nb03 + (i2) *nb02 + (i1) *nb01 + (i0) *nb00);
151          } else {
152              x = (const T *) (src1 + (i3 - o[3]) * nb13 + (i2 - o[2]) * nb12 + (i1 - o[1]) * nb11 +
153                                   (i0 - o[0]) * nb10);
154          }
155
156          T *y = (T *)(dst + i3 * nb3 + i2 * nb2 + i1 * nb1 + i0 * nb0);
157
158          *y = *x;
159      }
160  });
161}
162
163template <typename T>
164void concat_impl_sycl(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
165    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
166    const ggml_tensor *  src0   = dst->src[0];
167    const ggml_tensor *  src1   = dst->src[1];
168    queue_ptr            stream = ctx.stream();
169
170    const int32_t dim = ((int32_t *) dst->op_params)[0];
171
172    if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
173        const T * src0_d = (const T *) src0->data;
174        const T * src1_d = (const T *) src1->data;
175        T * dst_d = (T *) dst->data;
176        size_t type_size = elem_size(dst->type);
177        if (dim != 3) {
178            for (int i3 = 0; i3 < dst->ne[3]; i3++) {
179                concat_T_sycl<T>(src0_d + i3 * (src0->nb[3] / type_size), src1_d + i3 * (src1->nb[3] / type_size),
180                                dst_d + i3 * (dst->nb[3] / type_size), src0->ne[0], src0->ne[1], src0->ne[2], dst->ne[0],
181                                dst->ne[1], dst->ne[2], dim, stream);
182            }
183        } else {
184            const size_t size0 = ggml_nbytes(src0);
185            const size_t size1 = ggml_nbytes(src1);
186
187            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0)));
188            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d + size0 / type_size, src1_d, size1)));
189        }
190    } else {
191        concat_T_sycl_non_cont<T>(stream, (const char *) src0->data, (const char *) src1->data, (char *) dst->data,
192                                 src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3], src0->nb[0], src0->nb[1],
193                                 src0->nb[2], src0->nb[3], src1->ne[0], src1->ne[1], src1->ne[2], src1->ne[3],
194                                 src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3], dst->ne[0], dst->ne[1], dst->ne[2],
195                                 dst->ne[3], dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], dim);
196    }
197}
198
199static void concat_impl_q4_0_sycl(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
200    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
201    const ggml_tensor * src0   = dst->src[0];
202    const ggml_tensor * src1   = dst->src[1];
203    queue_ptr           stream = ctx.stream();
204
205    const int32_t dim = ((int32_t *) dst->op_params)[0];
206
207    GGML_ASSERT(src0->type == GGML_TYPE_Q4_0);
208    GGML_ASSERT(src1->type == GGML_TYPE_Q4_0);
209    GGML_ASSERT(dst->type == GGML_TYPE_Q4_0);
210    GGML_ASSERT(src0->ne[0] % QK4_0 == 0);
211    GGML_ASSERT(src1->ne[0] % QK4_0 == 0);
212    GGML_ASSERT(dst->ne[0] % QK4_0 == 0);
213
214    const int ne00_blk = src0->ne[0] / QK4_0;
215    const int ne0_blk  = dst->ne[0] / QK4_0;
216
217    if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
218        const block_q4_0 * src0_d = (const block_q4_0 *) src0->data;
219        const block_q4_0 * src1_d = (const block_q4_0 *) src1->data;
220        block_q4_0 * dst_d        = (block_q4_0 *) dst->data;
221        const size_t type_size    = sizeof(block_q4_0);
222
223        if (dim != 3) {
224            for (int i3 = 0; i3 < dst->ne[3]; i3++) {
225                concat_T_sycl<block_q4_0>(
226                    src0_d + i3 * (src0->nb[3] / type_size),
227                    src1_d + i3 * (src1->nb[3] / type_size),
228                    dst_d + i3 * (dst->nb[3] / type_size),
229                    ne00_blk, src0->ne[1], src0->ne[2], ne0_blk,
230                    dst->ne[1], dst->ne[2], dim, stream);
231            }
232        } else {
233            const size_t size0 = ggml_nbytes(src0);
234            const size_t size1 = ggml_nbytes(src1);
235
236            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0)));
237            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy((char *) dst_d + size0, src1_d, size1)));
238        }
239    } else {
240        concat_T_sycl_non_cont<block_q4_0>(
241            stream, (const char *) src0->data, (const char *) src1->data,
242            (char *) dst->data,
243            ne00_blk, src0->ne[1], src0->ne[2], src0->ne[3],
244            src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
245            src1->ne[0] / QK4_0, src1->ne[1], src1->ne[2], src1->ne[3],
246            src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
247            ne0_blk, dst->ne[1], dst->ne[2], dst->ne[3],
248            dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], dim);
249    }
250}
251
252static void concat_impl_q4_1_sycl(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
253    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
254    const ggml_tensor * src0   = dst->src[0];
255    const ggml_tensor * src1   = dst->src[1];
256    queue_ptr           stream = ctx.stream();
257
258    const int32_t dim = ((int32_t *) dst->op_params)[0];
259
260    GGML_ASSERT(src0->type == GGML_TYPE_Q4_1);
261    GGML_ASSERT(src1->type == GGML_TYPE_Q4_1);
262    GGML_ASSERT(dst->type == GGML_TYPE_Q4_1);
263    GGML_ASSERT(src0->ne[0] % QK4_1 == 0);
264    GGML_ASSERT(src1->ne[0] % QK4_1 == 0);
265    GGML_ASSERT(dst->ne[0] % QK4_1 == 0);
266
267    const int ne00_blk = src0->ne[0] / QK4_1;
268    const int ne0_blk  = dst->ne[0] / QK4_1;
269
270    if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
271        const block_q4_1 * src0_d = (const block_q4_1 *) src0->data;
272        const block_q4_1 * src1_d = (const block_q4_1 *) src1->data;
273        block_q4_1 * dst_d        = (block_q4_1 *) dst->data;
274        const size_t type_size    = sizeof(block_q4_1);
275
276        if (dim != 3) {
277            for (int i3 = 0; i3 < dst->ne[3]; i3++) {
278                concat_T_sycl<block_q4_1>(
279                    src0_d + i3 * (src0->nb[3] / type_size),
280                    src1_d + i3 * (src1->nb[3] / type_size),
281                    dst_d + i3 * (dst->nb[3] / type_size),
282                    ne00_blk, src0->ne[1], src0->ne[2], ne0_blk,
283                    dst->ne[1], dst->ne[2], dim, stream);
284            }
285        } else {
286            const size_t size0 = ggml_nbytes(src0);
287            const size_t size1 = ggml_nbytes(src1);
288
289            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0)));
290            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy((char *) dst_d + size0, src1_d, size1)));
291        }
292    } else {
293        concat_T_sycl_non_cont<block_q4_1>(
294            stream, (const char *) src0->data, (const char *) src1->data,
295            (char *) dst->data,
296            ne00_blk, src0->ne[1], src0->ne[2], src0->ne[3],
297            src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
298            src1->ne[0] / QK4_1, src1->ne[1], src1->ne[2], src1->ne[3],
299            src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
300            ne0_blk, dst->ne[1], dst->ne[2], dst->ne[3],
301            dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], dim);
302    }
303}
304
305static void concat_impl_q5_0_sycl(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
306    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
307    const ggml_tensor * src0   = dst->src[0];
308    const ggml_tensor * src1   = dst->src[1];
309    queue_ptr           stream = ctx.stream();
310
311    const int32_t dim = ((int32_t *) dst->op_params)[0];
312
313    GGML_ASSERT(src0->type == GGML_TYPE_Q5_0);
314    GGML_ASSERT(src1->type == GGML_TYPE_Q5_0);
315    GGML_ASSERT(dst->type == GGML_TYPE_Q5_0);
316    GGML_ASSERT(src0->ne[0] % QK5_0 == 0);
317    GGML_ASSERT(src1->ne[0] % QK5_0 == 0);
318    GGML_ASSERT(dst->ne[0] % QK5_0 == 0);
319
320    const int ne00_blk = src0->ne[0] / QK5_0;
321    const int ne0_blk  = dst->ne[0] / QK5_0;
322
323    if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
324        const block_q5_0 * src0_d = (const block_q5_0 *) src0->data;
325        const block_q5_0 * src1_d = (const block_q5_0 *) src1->data;
326        block_q5_0 * dst_d        = (block_q5_0 *) dst->data;
327        const size_t type_size    = sizeof(block_q5_0);
328
329        if (dim != 3) {
330            for (int i3 = 0; i3 < dst->ne[3]; i3++) {
331                concat_T_sycl<block_q5_0>(
332                    src0_d + i3 * (src0->nb[3] / type_size),
333                    src1_d + i3 * (src1->nb[3] / type_size),
334                    dst_d + i3 * (dst->nb[3] / type_size),
335                    ne00_blk, src0->ne[1], src0->ne[2], ne0_blk,
336                    dst->ne[1], dst->ne[2], dim, stream);
337            }
338        } else {
339            const size_t size0 = ggml_nbytes(src0);
340            const size_t size1 = ggml_nbytes(src1);
341
342            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0)));
343            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy((char *) dst_d + size0, src1_d, size1)));
344        }
345    } else {
346        concat_T_sycl_non_cont<block_q5_0>(
347            stream, (const char *) src0->data, (const char *) src1->data,
348            (char *) dst->data,
349            ne00_blk, src0->ne[1], src0->ne[2], src0->ne[3],
350            src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
351            src1->ne[0] / QK5_0, src1->ne[1], src1->ne[2], src1->ne[3],
352            src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
353            ne0_blk, dst->ne[1], dst->ne[2], dst->ne[3],
354            dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], dim);
355    }
356}
357
358static void concat_impl_q5_1_sycl(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
359    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
360    const ggml_tensor * src0   = dst->src[0];
361    const ggml_tensor * src1   = dst->src[1];
362    queue_ptr           stream = ctx.stream();
363
364    const int32_t dim = ((int32_t *) dst->op_params)[0];
365
366    GGML_ASSERT(src0->type == GGML_TYPE_Q5_1);
367    GGML_ASSERT(src1->type == GGML_TYPE_Q5_1);
368    GGML_ASSERT(dst->type == GGML_TYPE_Q5_1);
369    GGML_ASSERT(src0->ne[0] % QK5_1 == 0);
370    GGML_ASSERT(src1->ne[0] % QK5_1 == 0);
371    GGML_ASSERT(dst->ne[0] % QK5_1 == 0);
372
373    const int ne00_blk = src0->ne[0] / QK5_1;
374    const int ne0_blk  = dst->ne[0] / QK5_1;
375
376    if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
377        const block_q5_1 * src0_d = (const block_q5_1 *) src0->data;
378        const block_q5_1 * src1_d = (const block_q5_1 *) src1->data;
379        block_q5_1 * dst_d        = (block_q5_1 *) dst->data;
380        const size_t type_size    = sizeof(block_q5_1);
381
382        if (dim != 3) {
383            for (int i3 = 0; i3 < dst->ne[3]; i3++) {
384                concat_T_sycl<block_q5_1>(
385                    src0_d + i3 * (src0->nb[3] / type_size),
386                    src1_d + i3 * (src1->nb[3] / type_size),
387                    dst_d + i3 * (dst->nb[3] / type_size),
388                    ne00_blk, src0->ne[1], src0->ne[2], ne0_blk,
389                    dst->ne[1], dst->ne[2], dim, stream);
390            }
391        } else {
392            const size_t size0 = ggml_nbytes(src0);
393            const size_t size1 = ggml_nbytes(src1);
394
395            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0)));
396            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy((char *) dst_d + size0, src1_d, size1)));
397        }
398    } else {
399        concat_T_sycl_non_cont<block_q5_1>(
400            stream, (const char *) src0->data, (const char *) src1->data,
401            (char *) dst->data,
402            ne00_blk, src0->ne[1], src0->ne[2], src0->ne[3],
403            src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
404            src1->ne[0] / QK5_1, src1->ne[1], src1->ne[2], src1->ne[3],
405            src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
406            ne0_blk, dst->ne[1], dst->ne[2], dst->ne[3],
407            dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], dim);
408    }
409}
410
411static void concat_impl_q8_0_sycl(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
412    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
413    const ggml_tensor * src0   = dst->src[0];
414    const ggml_tensor * src1   = dst->src[1];
415    queue_ptr           stream = ctx.stream();
416
417    const int32_t dim = ((int32_t *) dst->op_params)[0];
418
419    GGML_ASSERT(src0->type == GGML_TYPE_Q8_0);
420    GGML_ASSERT(src1->type == GGML_TYPE_Q8_0);
421    GGML_ASSERT(dst->type == GGML_TYPE_Q8_0);
422    GGML_ASSERT(src0->ne[0] % QK8_0 == 0);
423    GGML_ASSERT(src1->ne[0] % QK8_0 == 0);
424    GGML_ASSERT(dst->ne[0] % QK8_0 == 0);
425
426    const int ne00_blk = src0->ne[0] / QK8_0;
427    const int ne0_blk  = dst->ne[0] / QK8_0;
428
429    if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
430        const block_q8_0 * src0_d = (const block_q8_0 *) src0->data;
431        const block_q8_0 * src1_d = (const block_q8_0 *) src1->data;
432        block_q8_0 * dst_d        = (block_q8_0 *) dst->data;
433        const size_t type_size    = sizeof(block_q8_0);
434
435        if (dim != 3) {
436            for (int i3 = 0; i3 < dst->ne[3]; i3++) {
437                concat_T_sycl<block_q8_0>(
438                    src0_d + i3 * (src0->nb[3] / type_size),
439                    src1_d + i3 * (src1->nb[3] / type_size),
440                    dst_d + i3 * (dst->nb[3] / type_size),
441                    ne00_blk, src0->ne[1], src0->ne[2], ne0_blk,
442                    dst->ne[1], dst->ne[2], dim, stream);
443            }
444        } else {
445            const size_t size0 = ggml_nbytes(src0);
446            const size_t size1 = ggml_nbytes(src1);
447            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0)));
448            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy((char *) dst_d + size0, src1_d, size1)));
449        }
450    } else {
451        concat_T_sycl_non_cont<block_q8_0>(
452            stream, (const char *) src0->data, (const char *) src1->data,
453            (char *) dst->data,
454            ne00_blk, src0->ne[1], src0->ne[2], src0->ne[3],
455            src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
456            src1->ne[0] / QK8_0, src1->ne[1], src1->ne[2], src1->ne[3],
457            src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
458            ne0_blk, dst->ne[1], dst->ne[2], dst->ne[3],
459            dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], dim);
460    }
461}
462
463void ggml_sycl_op_concat(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
464
465    switch (dst->type) {
466    case GGML_TYPE_F32:
467        concat_impl_sycl<float>(ctx, dst);
468        break;
469    case GGML_TYPE_F16:
470        concat_impl_sycl<sycl::half>(ctx, dst);
471        break;
472#ifdef GGML_SYCL_HAS_BF16
473    case GGML_TYPE_BF16:
474        concat_impl_sycl<sycl::ext::oneapi::bfloat16>(ctx, dst);
475        break;
476#endif
477    case GGML_TYPE_I32:
478        concat_impl_sycl<int32_t>(ctx, dst);
479        break;
480    case GGML_TYPE_I16:
481        concat_impl_sycl<int16_t>(ctx, dst);
482        break;
483    case GGML_TYPE_I64:
484        concat_impl_sycl<int64_t>(ctx, dst);
485        break;
486    case GGML_TYPE_I8:
487        concat_impl_sycl<int8_t>(ctx, dst);
488        break;
489    case GGML_TYPE_Q4_0:
490        concat_impl_q4_0_sycl(ctx, dst);
491        break;
492    case GGML_TYPE_Q4_1:
493        concat_impl_q4_1_sycl(ctx, dst);
494        break;
495    case GGML_TYPE_Q5_0:
496        concat_impl_q5_0_sycl(ctx, dst);
497        break;
498    case GGML_TYPE_Q5_1:
499        concat_impl_q5_1_sycl(ctx, dst);
500        break;
501    case GGML_TYPE_Q8_0:
502        concat_impl_q8_0_sycl(ctx, dst);
503        break;
504    default:
505        fprintf(stderr, "%s: unsupported types: dst: %s\n", __func__, ggml_type_name(dst->type));
506        GGML_ASSERT(false);
507    break;
508    }
509}