char-fork/llama.cpp

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

Intel AI Get-to Market Customer Success and Solutionssycl: add FILL, CUMSUM, DIAG, SOLVE_TRI, SSM_SCAN, GATED_DELTA_NET (#22149)ad0922465

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
1.8 KiB55 linesraw
1#include "fill.hpp"
2#include "common.hpp"
3
4#define SYCL_FILL_BLOCK_SIZE 256
5
6template <typename T>
7static void fill_kernel(T * dst, const int64_t k, const T value,
8                        const sycl::nd_item<1> & item) {
9    const int64_t i = (int64_t)item.get_global_id(0);
10    if (i >= k) {
11        return;
12    }
13    dst[i] = value;
14}
15
16inline void ggml_sycl_op_fill(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
17    GGML_ASSERT(ggml_is_contiguous(dst));
18
19    dpct::queue_ptr stream = ctx.stream();
20    SYCL_CHECK(ggml_sycl_set_device(ctx.device));
21
22    float value;
23    memcpy(&value, dst->op_params, sizeof(float));
24
25    const int64_t k = ggml_nelements(dst);
26    const int64_t num_blocks = (k + SYCL_FILL_BLOCK_SIZE - 1) / SYCL_FILL_BLOCK_SIZE;
27    void * dst_d = dst->data;
28
29    switch (dst->type) {
30        case GGML_TYPE_F32:
31            stream->parallel_for(
32                sycl::nd_range<1>(num_blocks * SYCL_FILL_BLOCK_SIZE, SYCL_FILL_BLOCK_SIZE),
33                [=](sycl::nd_item<1> item) {
34                    fill_kernel(static_cast<float *>(dst_d), k, value, item);
35                });
36            break;
37        case GGML_TYPE_F16:
38            {
39                sycl::half h_value = sycl::half(value);
40                stream->parallel_for(
41                    sycl::nd_range<1>(num_blocks * SYCL_FILL_BLOCK_SIZE, SYCL_FILL_BLOCK_SIZE),
42                    [=](sycl::nd_item<1> item) {
43                        fill_kernel(static_cast<sycl::half *>(dst_d), k, h_value, item);
44                    });
45            }
46            break;
47        default:
48            GGML_ABORT("unsupported type");
49    }
50}
51
52void ggml_sycl_fill(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
53    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/0);
54    ggml_sycl_op_fill(ctx, dst);
55}