char-fork/llama.cpp

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

Intel AI Get-to Market Customer Success and Solutionssycl: support non-contiguous input in PAD op (#22148)c5703e03a

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
3.9 KiB97 linesraw
1//
2// MIT license
3// Copyright (C) 2025 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 "common.hpp"
14#include "pad.hpp"
15
16static void pad_f32(const float * src, size_t s00, size_t s01, size_t s02, size_t s03,
17                    float * dst,
18                    const int lp0, const int rp0, const int lp1, const int rp1,
19                    const int lp2, const int rp2, const int lp3, const int rp3,
20                    const int ne0, const int ne1, const int ne2, const int ne3,
21                    sycl::nd_item<3> item_ct1) {
22    int i0 = item_ct1.get_local_id(2) +
23             item_ct1.get_group(2) * item_ct1.get_local_range(2);
24    int i1 = item_ct1.get_group(1);
25    int i2 = item_ct1.get_group(0) % ne2;
26    int i3 = item_ct1.get_group(0) / ne2;
27    if (i0 >= ne0 || i1 >= ne1 || i2 >= ne2 || i3 >= ne3) {
28        return;
29    }
30
31    const int64_t dst_idx = i3*(ne0*ne1*ne2) + i2*(ne0*ne1) + i1*ne0 + i0;
32    if ((i0 >= lp0 && i0 < ne0 - rp0) &&
33        (i1 >= lp1 && i1 < ne1 - rp1) &&
34        (i2 >= lp2 && i2 < ne2 - rp2) &&
35        (i3 >= lp3 && i3 < ne3 - rp3)) {
36        const int64_t i00 = i0 - lp0;
37        const int64_t i01 = i1 - lp1;
38        const int64_t i02 = i2 - lp2;
39        const int64_t i03 = i3 - lp3;
40
41        const int64_t src_idx = i03 * s03 + i02 * s02 + i01 * s01 + i00 * s00;
42
43        dst[dst_idx] = src[src_idx];
44    } else {
45        dst[dst_idx] = 0.0f;
46    }
47}
48
49static void pad_f32_sycl(const float * src, size_t s00, size_t s01, size_t s02, size_t s03,
50                         float * dst, const int lp0, const int rp0, const int lp1, const int rp1,
51                         const int lp2, const int rp2, const int lp3, const int rp3,
52                         const int ne0, const int ne1, const int ne2, const int ne3,
53                         dpct::queue_ptr stream) {
54    int num_blocks = (ne0 + SYCL_PAD_BLOCK_SIZE - 1) / SYCL_PAD_BLOCK_SIZE;
55    sycl::range<3> grid(ne2 * ne3, ne1, num_blocks);
56    stream->parallel_for(
57        sycl::nd_range<3>(grid * sycl::range<3>(1, 1, SYCL_PAD_BLOCK_SIZE),
58                          sycl::range<3>(1, 1, SYCL_PAD_BLOCK_SIZE)),
59        [=](sycl::nd_item<3> item_ct1) {
60            pad_f32(src, s00, s01, s02, s03, dst, lp0, rp0, lp1, rp1, lp2, rp2, lp3, rp3,
61                    ne0, ne1, ne2, ne3, item_ct1);
62        });
63}
64
65void ggml_sycl_op_pad(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
66    const ggml_tensor * src0 = dst->src[0];
67    const float * src0_d = (const float *)src0->data;
68    float * dst_d = (float *)dst->data;
69    dpct::queue_ptr stream = ctx.stream();
70
71    GGML_ASSERT(src0->type == GGML_TYPE_F32);
72    GGML_ASSERT(dst->type == GGML_TYPE_F32);
73
74    const size_t ts = ggml_type_size(src0->type);
75    const size_t s00 = src0->nb[0] / ts;
76    const size_t s01 = src0->nb[1] / ts;
77    const size_t s02 = src0->nb[2] / ts;
78    const size_t s03 = src0->nb[3] / ts;
79
80    const int32_t lp0 = ((const int32_t *)(dst->op_params))[0];
81    const int32_t rp0 = ((const int32_t *)(dst->op_params))[1];
82    const int32_t lp1 = ((const int32_t *)(dst->op_params))[2];
83    const int32_t rp1 = ((const int32_t *)(dst->op_params))[3];
84    const int32_t lp2 = ((const int32_t *)(dst->op_params))[4];
85    const int32_t rp2 = ((const int32_t *)(dst->op_params))[5];
86    const int32_t lp3 = ((const int32_t *)(dst->op_params))[6];
87    const int32_t rp3 = ((const int32_t *)(dst->op_params))[7];
88
89    pad_f32_sycl(src0_d, s00, s01, s02, s03, dst_d,
90                 lp0, rp0, lp1, rp1, lp2, rp2, lp3, rp3,
91                 dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3], stream);
92}
93
94void ggml_sycl_pad(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
95    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
96    ggml_sycl_op_pad(ctx, dst);
97}