char-fork/llama.cpp

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

Neo Zhang[SYCL] add to support pool_1d, move pool_1d/2d code to pool.cpp/hpp (#24584)987fbd821

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
6.5 KiB185 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 "pool.hpp"
14#include <float.h>
15
16template <typename Ti, typename To>
17static void pool2d_nchw_kernel(
18        const int ih, const int iw, const int oh, const int ow,
19        const int kh, const int kw, const int sh, const int sw,
20        const int ph, const int pw, const int parallel_elements,
21        const Ti* src, To* dst, const enum ggml_op_pool op,
22        const sycl::nd_item<3> &item_ct1) {
23        int idx = item_ct1.get_local_id(2) +
24                  item_ct1.get_group(2) * item_ct1.get_local_range(2);
25        if (idx >= parallel_elements) {
26            return;
27        }
28
29        const int I_HW = ih * iw;
30        const int O_HW = oh * ow;
31        const int nc = idx / O_HW;
32        const int cur_oh = idx % O_HW / ow;
33        const int cur_ow = idx % O_HW % ow;
34        const Ti* i_ptr = src + nc * I_HW;
35        To* o_ptr = dst + nc * O_HW;
36        const int start_h = cur_oh * sh - ph;
37        const int bh = sycl::max(0, start_h);
38        const int eh = sycl::min(ih, start_h + kh);
39        const int start_w = cur_ow * sw - pw;
40        const int bw = sycl::max(0, start_w);
41        const int ew = sycl::min(iw, start_w + kw);
42
43        To res = 0;
44
45        switch (op) {
46            case GGML_OP_POOL_AVG: res = 0; break;
47            case GGML_OP_POOL_MAX: res = -FLT_MAX; break;
48            default:
49                res      = (To) sycl::nan(uint32_t(0));
50                break;
51        }
52
53        for (int i = bh; i < eh; i += 1) {
54            for (int j = bw; j < ew; j += 1) {
55                Ti cur = i_ptr[i * iw + j];
56                switch (op) {
57                    case GGML_OP_POOL_AVG: res += (cur / (kh * kw)); break;
58                    case GGML_OP_POOL_MAX: res = sycl::max(res, (To)cur); break;
59                    default:
60                        res = (To) sycl::nan(uint32_t(0));
61                        break;
62                }
63            }
64        }
65        o_ptr[cur_oh * ow + cur_ow] = res;
66}
67
68template <typename Ti, typename To>
69static void pool1d_ncw_kernel(
70        const int iw, const int ow,
71        const int k, const int s,
72        const int p, const int parallel_elements,
73        const Ti * src, To * dst, const enum ggml_op_pool op,
74        const sycl::nd_item<3> & item_ct1) {
75    int idx = item_ct1.get_local_id(2) +
76              item_ct1.get_group(2) * item_ct1.get_local_range(2);
77    if (idx >= parallel_elements) {
78        return;
79    }
80
81    const int nc     = idx / ow;
82    const int cur_ow = idx % ow;
83    const Ti * i_ptr = src + nc * iw;
84    To *       o_ptr = dst + nc * ow;
85    const int  start = cur_ow * s - p;
86    const int  b     = sycl::max(0, start);
87    const int  e     = sycl::min(iw, start + k);
88
89    To res = 0;
90    switch (op) {
91        case GGML_OP_POOL_AVG: res = 0;        break;
92        case GGML_OP_POOL_MAX: res = -FLT_MAX; break;
93        default:
94            res = (To) sycl::nan(uint32_t(0));
95            break;
96    }
97
98    for (int j = b; j < e; j += 1) {
99        Ti cur = i_ptr[j];
100        switch (op) {
101            case GGML_OP_POOL_AVG: res += cur; break;
102            case GGML_OP_POOL_MAX: res = sycl::max(res, (To) cur); break;
103            default:
104                res = (To) sycl::nan(uint32_t(0));
105                break;
106        }
107    }
108
109    const int count = e - b;
110    if (op == GGML_OP_POOL_AVG) {
111        res = (count > 0) ? (res / count) : (To) 0;
112    }
113    o_ptr[cur_ow] = res;
114}
115
116void ggml_sycl_op_pool2d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
117    GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32);
118    GGML_ASSERT( dst->type == GGML_TYPE_F32);
119    dpct::queue_ptr main_stream = ctx.stream();
120    SYCL_CHECK(ggml_sycl_set_device(ctx.device));
121    const float * src0_dd = static_cast<const float *>(dst->src[0]->data);
122    float *       dst_dd  = static_cast<float *>(dst->data);
123
124    const int32_t * opts = (const int32_t *)dst->op_params;
125    enum ggml_op_pool op = static_cast<ggml_op_pool>(opts[0]);
126    const int k0 = opts[1];
127    const int k1 = opts[2];
128    const int s0 = opts[3];
129    const int s1 = opts[4];
130    const int p0 = opts[5];
131    const int p1 = opts[6];
132
133    const int64_t IH = dst->src[0]->ne[1];
134    const int64_t IW = dst->src[0]->ne[0];
135
136    const int64_t N = dst->ne[3];
137    const int64_t OC = dst->ne[2];
138    const int64_t OH = dst->ne[1];
139    const int64_t OW = dst->ne[0];
140
141    const int parallel_elements = N * OC * OH * OW;
142    const int num_blocks = (parallel_elements + SYCL_POOL2D_BLOCK_SIZE - 1) / SYCL_POOL2D_BLOCK_SIZE;
143    sycl::range<3> block_nums(1, 1, num_blocks);
144    main_stream->parallel_for(
145        sycl::nd_range<3>(block_nums *
146                              sycl::range<3>(1, 1, SYCL_IM2COL_BLOCK_SIZE),
147                          sycl::range<3>(1, 1, SYCL_IM2COL_BLOCK_SIZE)),
148        [=](sycl::nd_item<3> item_ct1) {
149            pool2d_nchw_kernel(IH, IW, OH, OW, k1, k0, s1, s0, p1, p0,
150                               parallel_elements, src0_dd, dst_dd, op,
151                               item_ct1);
152        });
153}
154
155void ggml_sycl_op_pool1d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
156    GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32);
157    GGML_ASSERT( dst->type == GGML_TYPE_F32);
158    dpct::queue_ptr main_stream = ctx.stream();
159    SYCL_CHECK(ggml_sycl_set_device(ctx.device));
160    const float * src0_dd = static_cast<const float *>(dst->src[0]->data);
161    float *       dst_dd  = static_cast<float *>(dst->data);
162
163    const int32_t * opts = (const int32_t *)dst->op_params;
164    enum ggml_op_pool op = static_cast<ggml_op_pool>(opts[0]);
165    const int k0 = opts[1];
166    const int s0 = opts[2];
167    const int p0 = opts[3];
168
169    const int64_t IW = dst->src[0]->ne[0];
170    const int64_t OW = dst->ne[0];
171    const int64_t NC = dst->ne[3] * dst->ne[2] * dst->ne[1];
172
173    const int parallel_elements = NC * OW;
174    const int num_blocks = (parallel_elements + SYCL_POOL1D_BLOCK_SIZE - 1) / SYCL_POOL1D_BLOCK_SIZE;
175    sycl::range<3> block_nums(1, 1, num_blocks);
176    main_stream->parallel_for(
177        sycl::nd_range<3>(block_nums *
178                              sycl::range<3>(1, 1, SYCL_POOL1D_BLOCK_SIZE),
179                          sycl::range<3>(1, 1, SYCL_POOL1D_BLOCK_SIZE)),
180        [=](sycl::nd_item<3> item_ct1) {
181            pool1d_ncw_kernel(IW, OW, k0, s0, p0,
182                              parallel_elements, src0_dd, dst_dd, op,
183                              item_ct1);
184        });
185}