char-fork/llama.cpp

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

Neo Zhang[SYCL] support OPs: conv_2d, conv_2d_dw, conv2d_transpose (#24600)6f1034b32

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
5.3 KiB125 linesraw
1#include "conv2d-transpose.hpp"
2#include "convert.hpp"
3
4template <typename kernel_t>
5static void conv2d_transpose_kernel(const float * input, const kernel_t * kernel, float * output,
6                                    const int in_w, const int in_h,
7                                    const int out_w, const int out_h,
8                                    const int kernel_w, const int kernel_h,
9                                    const int stride,
10                                    const int c_in, const int c_out, const int batches,
11                                    const sycl::nd_item<3> & item_ct1) {
12    const int global_idx     = item_ct1.get_local_id(2) +
13                               item_ct1.get_group(2) * item_ct1.get_local_range(2);
14    const int total_elements = out_w * out_h * c_out * batches;
15
16    if (global_idx >= total_elements) {
17        return;
18    }
19
20    const int out_x = global_idx % out_w;
21    const int out_y = (global_idx / out_w) % out_h;
22    const int c_idx = (global_idx / (out_w * out_h)) % c_out;
23    const int n_idx = global_idx / (out_w * out_h * c_out);
24
25    float acc = 0.0f;
26
27    for (int c_in_idx = 0; c_in_idx < c_in; ++c_in_idx) {
28        for (int kh = 0; kh < kernel_h; ++kh) {
29            int in_y = out_y - kh;
30            if (in_y < 0 || in_y % stride) {
31                continue;
32            }
33            in_y /= stride;
34            if (in_y >= in_h) {
35                continue;
36            }
37
38            for (int kw = 0; kw < kernel_w; ++kw) {
39                int in_x = out_x - kw;
40                if (in_x < 0 || in_x % stride) {
41                    continue;
42                }
43                in_x /= stride;
44                if (in_x >= in_w) {
45                    continue;
46                }
47
48                const int input_idx  = (in_w * in_h * c_in) * n_idx + (in_w * in_h) * c_in_idx + in_w * in_y + in_x;
49                const int kernel_idx = (kernel_h * kernel_w * c_out) * c_in_idx + (kernel_h * kernel_w) * c_idx +
50                                       kernel_w * kh + kw;
51
52                acc += input[input_idx] * ggml_sycl_cast<float>(kernel[kernel_idx]);
53            }
54        }
55    }
56
57    output[(out_w * out_h * c_out) * n_idx + (out_w * out_h) * c_idx + out_w * out_y + out_x] = acc;
58}
59
60template <typename kernel_t>
61static void conv2d_transpose_sycl(const float * input_d, const kernel_t * kernel_d, float * output_d,
62                                   const int in_w, const int in_h,
63                                   const int out_w, const int out_h,
64                                   const int kernel_w, const int kernel_h,
65                                   const int stride,
66                                   const int c_in, const int c_out, const int batches,
67                                   const queue_ptr & stream) {
68    const int total      = out_w * out_h * c_out * batches;
69    const int num_blocks = (total + SYCL_CONV2D_TRANSPOSE_BLOCK_SIZE - 1) / SYCL_CONV2D_TRANSPOSE_BLOCK_SIZE;
70    const sycl::range<3> block_dims(1, 1, SYCL_CONV2D_TRANSPOSE_BLOCK_SIZE);
71    const sycl::range<3> block_nums(1, 1, num_blocks);
72    stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
73        [=](sycl::nd_item<3> item_ct1) {
74            conv2d_transpose_kernel<kernel_t>(input_d, kernel_d, output_d,
75                                             in_w, in_h, out_w, out_h, kernel_w, kernel_h,
76                                             stride, c_in, c_out, batches, item_ct1);
77        });
78}
79
80// input:  (W, H, C_in, N)
81// kernel: (W, H, C_out, C_in)
82// output: (W, H, C_out, N)
83void ggml_sycl_op_conv2d_transpose(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
84    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
85
86    const ggml_tensor * kernel = dst->src[0];
87    const ggml_tensor * input  = dst->src[1];
88
89    GGML_ASSERT(kernel->type == GGML_TYPE_F16 || kernel->type == GGML_TYPE_F32);
90    GGML_ASSERT(input->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32);
91
92    GGML_ASSERT(ggml_is_contiguous(input));
93    GGML_ASSERT(ggml_is_contiguous(kernel));
94    GGML_ASSERT(ggml_is_contiguous(dst));
95
96    const float * input_d  = (const float *) input->data;
97    float *       output_d = (float *) dst->data;
98    const void *  kernel_d = kernel->data;
99
100    const int input_w      = input->ne[0];
101    const int input_h      = input->ne[1];
102    const int channels_in  = input->ne[2];
103    const int batches      = input->ne[3];
104    const int output_w     = dst->ne[0];
105    const int output_h     = dst->ne[1];
106    const int channels_out = kernel->ne[2];
107    const int kernel_w     = kernel->ne[0];
108    const int kernel_h     = kernel->ne[1];
109    const int stride       = dst->op_params[0];
110
111    GGML_ASSERT(channels_in == kernel->ne[3]);
112    GGML_ASSERT(stride > 0);
113
114    const queue_ptr stream = ctx.stream();
115
116    if (kernel->type == GGML_TYPE_F16) {
117        conv2d_transpose_sycl<sycl::half>(input_d, (const sycl::half *) kernel_d, output_d,
118                                          input_w, input_h, output_w, output_h, kernel_w, kernel_h,
119                                          stride, channels_in, channels_out, batches, stream);
120    } else {
121        conv2d_transpose_sycl<float>(input_d, (const float *) kernel_d, output_d,
122                                     input_w, input_h, output_w, output_h, kernel_w, kernel_h,
123                                     stride, channels_in, channels_out, batches, stream);
124    }
125}