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 →
2.1 KiB67 linesraw
1#include "diag.hpp"
2#include "common.hpp"
3
4#define SYCL_DIAG_BLOCK_SIZE 256
5
6template <typename T>
7static void diag_kernel(T * __restrict__ dst, const T * __restrict__ src,
8                        const int64_t ne0, const int64_t ne1,
9                        const int64_t ne2, const int64_t ne3,
10                        const int64_t total_elements,
11                        const sycl::nd_item<1> & item) {
12    const int64_t i = item.get_global_id(0);
13    if (i >= total_elements) {
14        return;
15    }
16
17    const int64_t i0 = i % ne0;
18    const int64_t i1 = (i / ne0) % ne1;
19    const int64_t i2 = (i / (ne0 * ne1)) % ne2;
20    const int64_t i3 = i / (ne0 * ne1 * ne2);
21
22    const int64_t dst_idx = ((i3 * ne2 + i2) * ne1 + i1) * ne0 + i0;
23
24    if (i0 == i1) {
25        const int64_t batch_idx = i3 * ne2 + i2;
26        dst[dst_idx] = src[batch_idx * ne0 + i0];
27    } else {
28        dst[dst_idx] = T(0);
29    }
30
31    (void)ne3;
32}
33
34inline void ggml_sycl_op_diag(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
35    const ggml_tensor * src0 = dst->src[0];
36
37    GGML_ASSERT(ggml_is_contiguous(dst));
38    GGML_ASSERT(ggml_is_contiguous(src0));
39    GGML_ASSERT(src0->ne[1] == 1);
40
41    dpct::queue_ptr stream = ctx.stream();
42    SYCL_CHECK(ggml_sycl_set_device(ctx.device));
43
44    const void * src0_d = src0->data;
45    void * dst_d = dst->data;
46
47    const int64_t ne0 = dst->ne[0];
48    const int64_t ne1 = dst->ne[1];
49    const int64_t ne2 = dst->ne[2];
50    const int64_t ne3 = dst->ne[3];
51    const int64_t n_elems = ggml_nelements(dst);
52    const int64_t num_blocks = (n_elems + SYCL_DIAG_BLOCK_SIZE - 1) / SYCL_DIAG_BLOCK_SIZE;
53
54    GGML_ASSERT(dst->type == GGML_TYPE_F32);
55    stream->parallel_for(
56        sycl::nd_range<1>(num_blocks * SYCL_DIAG_BLOCK_SIZE, SYCL_DIAG_BLOCK_SIZE),
57        [=](sycl::nd_item<1> item) {
58            diag_kernel(static_cast<float *>(dst_d),
59                        static_cast<const float *>(src0_d),
60                        ne0, ne1, ne2, ne3, n_elems, item);
61        });
62}
63
64void ggml_sycl_diag(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
65    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
66    ggml_sycl_op_diag(ctx, dst);
67}