char-fork/llama.cpp

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

Neo Zhangsycl : support kernel type fp16 for conv2d_dw (#25653)ae9291e16

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
7.3 KiB170 linesraw
1#include "conv2d-dw.hpp"
2
3struct conv2d_dw_params {
4    int in_w, in_h;
5    int out_w, out_h;
6    int kernel_w, kernel_h;
7    int stride_x, stride_y;
8    int padding_x, padding_y;
9    int dilation_x, dilation_y;
10    int channels, batches;
11};
12
13struct conv2d_dw_kernel_bounds {
14    int y_min, y_max;
15    int x_min, x_max;
16};
17
18static inline conv2d_dw_kernel_bounds dw_calculate_kernel_bounds(int out_x, int out_y,
19                                                                  const conv2d_dw_params & p) {
20    conv2d_dw_kernel_bounds bounds;
21    bounds.y_min = sycl::max(0, (p.padding_y - out_y * p.stride_y + p.dilation_y - 1) / p.dilation_y);
22    bounds.y_max = sycl::min(p.kernel_h,
23                             (p.in_h + p.padding_y - out_y * p.stride_y + p.dilation_y - 1) / p.dilation_y);
24    bounds.x_min = sycl::max(0, (p.padding_x - out_x * p.stride_x + p.dilation_x - 1) / p.dilation_x);
25    bounds.x_max = sycl::min(p.kernel_w,
26                             (p.in_w + p.padding_x - out_x * p.stride_x + p.dilation_x - 1) / p.dilation_x);
27    return bounds;
28}
29
30static inline int dw_calculate_input_coord(int out_coord, int kern_coord, int stride, int dilation, int padding) {
31    return out_coord * stride + kern_coord * dilation - padding;
32}
33
34// whcn layout: input/output stored as [N, C, H, W]
35struct dw_whcn_layout {
36    static int input_index(int n, int c, int y, int x, const conv2d_dw_params & p) {
37        return n * (p.channels * p.in_w * p.in_h) + c * p.in_w * p.in_h + y * p.in_w + x;
38    }
39    static int kernel_index(int c, int ky, int kx, const conv2d_dw_params & p) {
40        return c * p.kernel_h * p.kernel_w + ky * p.kernel_w + kx;
41    }
42    static int output_index(int n, int c, int y, int x, const conv2d_dw_params & p) {
43        return n * (p.channels * p.out_w * p.out_h) + c * p.out_w * p.out_h + y * p.out_w + x;
44    }
45    static void unpack_indices(int global_idx, const conv2d_dw_params & p,
46                               int & n, int & c, int & out_y, int & out_x) {
47        out_x  = global_idx % p.out_w;
48        out_y  = (global_idx / p.out_w) % p.out_h;
49        c      = (global_idx / (p.out_w * p.out_h)) % p.channels;
50        n      = global_idx / (p.out_w * p.out_h * p.channels);
51    }
52};
53
54// cwhn layout: input/output stored as [N, H, W, C]
55struct dw_cwhn_layout {
56    static int input_index(int n, int c, int y, int x, const conv2d_dw_params & p) {
57        return n * (p.channels * p.in_w * p.in_h) + (y * p.in_w + x) * p.channels + c;
58    }
59    static int kernel_index(int c, int ky, int kx, const conv2d_dw_params & p) {
60        return (ky * p.kernel_w + kx) * p.channels + c;
61    }
62    static int output_index(int n, int c, int y, int x, const conv2d_dw_params & p) {
63        return n * (p.channels * p.out_w * p.out_h) + y * (p.out_w * p.channels) + x * p.channels + c;
64    }
65    static void unpack_indices(int global_idx, const conv2d_dw_params & p,
66                               int & n, int & c, int & out_y, int & out_x) {
67        c      = global_idx % p.channels;
68        out_x  = (global_idx / p.channels) % p.out_w;
69        out_y  = (global_idx / (p.channels * p.out_w)) % p.out_h;
70        n      = global_idx / (p.channels * p.out_w * p.out_h);
71    }
72};
73
74template <typename KernelT, typename Layout>
75static void conv2d_dw_kernel(const float * input, const KernelT * kernel, float * output,
76                             const conv2d_dw_params p, const sycl::nd_item<3> & item_ct1) {
77    const int global_idx     = item_ct1.get_local_id(2) +
78                               item_ct1.get_group(2) * item_ct1.get_local_range(2);
79    const int total_elements = p.batches * p.channels * p.out_h * p.out_w;
80
81    if (global_idx >= total_elements) {
82        return;
83    }
84
85    int n, c, out_y, out_x;
86    Layout::unpack_indices(global_idx, p, n, c, out_y, out_x);
87
88    float acc = 0.0f;
89    const conv2d_dw_kernel_bounds bounds = dw_calculate_kernel_bounds(out_x, out_y, p);
90
91    for (int ky = bounds.y_min; ky < bounds.y_max; ++ky) {
92        const int in_y = dw_calculate_input_coord(out_y, ky, p.stride_y, p.dilation_y, p.padding_y);
93        for (int kx = bounds.x_min; kx < bounds.x_max; ++kx) {
94            const int in_x = dw_calculate_input_coord(out_x, kx, p.stride_x, p.dilation_x, p.padding_x);
95            acc += input[Layout::input_index(n, c, in_y, in_x, p)] *
96                   static_cast<float>(kernel[Layout::kernel_index(c, ky, kx, p)]);
97        }
98    }
99
100    output[Layout::output_index(n, c, out_y, out_x, p)] = acc;
101}
102
103template <typename KernelT, typename Layout>
104static void conv2d_dw_sycl(const float * x_d, const KernelT * w_d, float * y_d,
105                            const conv2d_dw_params p, const queue_ptr & stream) {
106    const int total      = p.batches * p.channels * p.out_h * p.out_w;
107    const int num_blocks = (total + SYCL_CONV2D_DW_BLOCK_SIZE - 1) / SYCL_CONV2D_DW_BLOCK_SIZE;
108    const sycl::range<3> block_dims(1, 1, SYCL_CONV2D_DW_BLOCK_SIZE);
109    const sycl::range<3> block_nums(1, 1, num_blocks);
110    stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
111        [=](sycl::nd_item<3> item_ct1) {
112            conv2d_dw_kernel<KernelT, Layout>(x_d, w_d, y_d, p, item_ct1);
113        });
114}
115
116void ggml_sycl_op_conv2d_dw(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
117    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
118
119    const ggml_tensor * kernel = dst->src[0];
120    const ggml_tensor * input  = dst->src[1];
121
122    GGML_ASSERT((kernel->type == GGML_TYPE_F32 || kernel->type == GGML_TYPE_F16) &&
123                input->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32);
124
125    const float * x_d = (const float *) input->data;
126    float *       y_d = (float *) dst->data;
127
128    const int32_t * p          = (const int32_t *) dst->op_params;
129    const int       stride_x   = p[0];
130    const int       stride_y   = p[1];
131    const int       padding_x  = p[2];
132    const int       padding_y  = p[3];
133    const int       dilation_x = p[4];
134    const int       dilation_y = p[5];
135
136    const int in_w     = input->ne[0];
137    const int in_h     = input->ne[1];
138    const int kernel_w = kernel->ne[0];
139    const int kernel_h = kernel->ne[1];
140    const int out_w    = dst->ne[0];
141    const int out_h    = dst->ne[1];
142    const int channels = dst->ne[2];
143    const int batches  = dst->ne[3];
144
145    const conv2d_dw_params params = { in_w, in_h, out_w, out_h, kernel_w, kernel_h,
146                                      stride_x, stride_y, padding_x, padding_y,
147                                      dilation_x, dilation_y, channels, batches };
148
149    const queue_ptr stream = ctx.stream();
150
151    if (kernel->type == GGML_TYPE_F16) {
152        const sycl::half * w_d = (const sycl::half *) kernel->data;
153        if (ggml_is_contiguous(input)) {
154            conv2d_dw_sycl<sycl::half, dw_whcn_layout>(x_d, w_d, y_d, params, stream);
155        } else if (ggml_is_contiguous_channels(input)) {
156            conv2d_dw_sycl<sycl::half, dw_cwhn_layout>(x_d, w_d, y_d, params, stream);
157        } else {
158            GGML_ABORT("Unsupported memory layout for conv2d_dw");
159        }
160    } else {
161        const float * w_d = (const float *) kernel->data;
162        if (ggml_is_contiguous(input)) {
163            conv2d_dw_sycl<float, dw_whcn_layout>(x_d, w_d, y_d, params, stream);
164        } else if (ggml_is_contiguous_channels(input)) {
165            conv2d_dw_sycl<float, dw_cwhn_layout>(x_d, w_d, y_d, params, stream);
166        } else {
167            GGML_ABORT("Unsupported memory layout for conv2d_dw");
168        }
169    }
170}