char-fork/llama.cpp

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

Neo Zhang[SYCL] support op col2im_1d (#25264)47e1de77a

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
3.2 KiB102 linesraw
1#include "col2im-1d.hpp"
2
3template <typename T>
4static void col2im_1d_sycl(
5        const T * col,
6        T * dst,
7        const int T_in,
8        const sycl::uint3 T_out_fd,
9        const int K,
10        const int K_OC,
11        const int32_t s0,
12        const int32_t p0,
13        const int total,
14        dpct::queue_ptr stream) {
15
16    const uint32_t block_size = SYCL_COL2IM_1D_BLOCK_SIZE;
17    const uint32_t num_blocks = (uint32_t) ((total + block_size - 1) / block_size);
18
19    stream->parallel_for(
20        sycl::nd_range<3>(
21            sycl::range<3>(1, 1, num_blocks * block_size),
22            sycl::range<3>(1, 1, block_size)),
23        [=](sycl::nd_item<3> item_ct1) {
24            const int idx = (int) item_ct1.get_global_id(2);
25            if (idx >= total) {
26                return;
27            }
28
29            const sycl::uint2 qr = fast_div_modulo((uint32_t) idx, T_out_fd);
30            const int oc    = (int) qr.x();
31            const int t_out = (int) qr.y();
32            const int t_abs = t_out + p0;
33
34            int t_in_min = (t_abs - K + s0) / s0;
35            if (t_in_min < 0) {
36                t_in_min = 0;
37            }
38            int t_in_max = t_abs / s0;
39            if (t_in_max >= T_in) {
40                t_in_max = T_in - 1;
41            }
42
43            float sum = 0.0f;
44            for (int t_in = t_in_min; t_in <= t_in_max; ++t_in) {
45                const int k = t_abs - t_in * s0;
46                sum += static_cast<float>(col[(oc * K + k) + t_in * K_OC]);
47            }
48
49            dst[idx] = static_cast<T>(sum);
50        });
51}
52
53void ggml_sycl_op_col2im_1d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
54    const ggml_tensor * src0 = dst->src[0];
55
56    GGML_ASSERT(src0 != nullptr);
57    GGML_ASSERT(ggml_is_contiguous(src0));
58    GGML_ASSERT(src0->type == dst->type);
59
60    const int32_t s0 = ((const int32_t *) dst->op_params)[0];
61    const int32_t OC = ((const int32_t *) dst->op_params)[1];
62    const int32_t p0 = ((const int32_t *) dst->op_params)[2];
63
64    const int K_OC  = (int) src0->ne[0];
65    const int T_in  = (int) src0->ne[1];
66    const int K     = K_OC / OC;
67    const int T_out = (int) dst->ne[0];
68
69    GGML_ASSERT(OC > 0);
70    GGML_ASSERT(K_OC % OC == 0);
71
72    const sycl::uint3 T_out_fd = init_fastdiv_values((uint32_t) T_out);
73
74    const int total = T_out * OC;
75
76    dpct::queue_ptr stream = ctx.stream();
77
78    switch (src0->type) {
79        case GGML_TYPE_F32:
80            col2im_1d_sycl<float>(
81                (const float *) src0->data,
82                (float *) dst->data,
83                T_in, T_out_fd, K, K_OC, s0, p0, total, stream);
84            break;
85        case GGML_TYPE_F16:
86            col2im_1d_sycl<sycl::half>(
87                (const sycl::half *) src0->data,
88                (sycl::half *) dst->data,
89                T_in, T_out_fd, K, K_OC, s0, p0, total, stream);
90            break;
91#ifdef GGML_SYCL_HAS_BF16
92        case GGML_TYPE_BF16:
93            col2im_1d_sycl<sycl::ext::oneapi::bfloat16>(
94                (const sycl::ext::oneapi::bfloat16 *) src0->data,
95                (sycl::ext::oneapi::bfloat16 *) dst->data,
96                T_in, T_out_fd, K, K_OC, s0, p0, total, stream);
97            break;
98#endif
99        default:
100            GGML_ABORT("col2im_1d: unsupported type %d", src0->type);
101    }
102}