char-fork/llama.cpp

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

Neo Zhangsycl : support MUL_MAT and OUT_PROD with Q1_0 (#24721)dd69db292

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
3.4 KiB83 linesraw
1#include "outprod.hpp"
2#include "convert.hpp"
3
4void ggml_sycl_op_out_prod(ggml_backend_sycl_context& ctx, ggml_tensor* dst) {
5    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
6    const ggml_tensor *src0 = dst->src[0];
7    const ggml_tensor *src1 = dst->src[1];
8
9    GGML_ASSERT(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_Q1_0);
10    GGML_ASSERT(src1->type == GGML_TYPE_F32);
11    GGML_ASSERT(dst->type == GGML_TYPE_F32);
12    GGML_ASSERT(ggml_is_contiguous(src0));
13    GGML_ASSERT(ggml_is_contiguous(dst));
14
15    GGML_TENSOR_BINARY_OP_LOCALS
16
17    // Get SYCL queue
18    dpct::queue_ptr stream = ctx.stream();
19
20    // Dimension checks
21    GGML_ASSERT(ne01 == ne11);  // Inner dimensions must match
22    GGML_ASSERT(ne0 == ne00);   // Output rows match src0 rows
23    GGML_ASSERT(ne1 == ne10);   // Output cols match src1 cols
24    GGML_ASSERT(ne2 == ne12);
25    GGML_ASSERT(ne3 == ne13);
26    GGML_ASSERT(ne2 % ne02 == 0);
27    GGML_ASSERT(ne3 % ne03 == 0);
28
29    // Get data pointers
30    const float * src0_d = (const float *) src0->data;
31    const float * src1_d = (const float *) src1->data;
32    float * dst_d = (float *) dst->data;
33
34    ggml_sycl_pool_alloc<float> src0_as_f32(ctx.pool());
35    int64_t src0_nb02 = nb02;
36    int64_t src0_nb03 = nb03;
37    if (src0->type == GGML_TYPE_Q1_0) {
38        scope_op_debug_print scope_dbg_print(__func__, "/to_fp32_sycl", dst, /*num_src=*/2,
39                                             " : converting src0 Q1_0 to fp32");
40        src0_d = src0_as_f32.alloc(ne00 * ne01 * ne02 * ne03);
41        const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(src0->type, dst);
42        GGML_ASSERT(to_fp32_sycl != nullptr);
43        to_fp32_sycl(src0->data, const_cast<float *>(src0_d), ne00 * ne01 * ne02 * ne03, stream);
44
45        // Dequantized src0 buffer is contiguous fp32 [ne00, ne01, ne02, ne03].
46        src0_nb02 = ne00 * ne01 * (int64_t) sizeof(float);
47        src0_nb03 = ne00 * ne01 * ne02 * (int64_t) sizeof(float);
48    }
49
50    // GEMM parameters
51    const float alpha = 1.0f;
52    const float beta = 0.0f;
53
54    // Handle transposition of src1
55    const bool src1_T = ggml_is_transposed(src1);
56    const oneapi::mkl::transpose src1_op = src1_T ? oneapi::mkl::transpose::nontrans : oneapi::mkl::transpose::trans;
57    const int64_t ldb = (src1_T ? nb10 : nb11) / sizeof(float);
58
59    const int64_t r2 = ne2 / ne02;
60    const int64_t r3 = ne3 / ne03;
61
62    try {
63        // OUT_PROD applies independently to each (i2, i3) destination plane.
64        for (int64_t i3 = 0; i3 < ne3; ++i3) {
65            for (int64_t i2 = 0; i2 < ne2; ++i2) {
66                const int64_t i03 = i3 / r3;
67                const int64_t i02 = i2 / r2;
68
69                const float * src0_plane = (const float *) ((const char *) src0_d + i02 * src0_nb02 + i03 * src0_nb03);
70                const float * src1_plane = (const float *) ((const char *) src1_d + i2 * nb12 + i3 * nb13);
71                float * dst_plane = (float *) ((char *) dst_d + i2 * nb2 + i3 * nb3);
72
73                // Perform matrix multiplication using oneMKL GEMM
74                oneapi::mkl::blas::column_major::gemm(*stream, oneapi::mkl::transpose::nontrans, src1_op,
75                                                      ne0, ne1, ne01, alpha, src0_plane, ne00,
76                                                      src1_plane, ldb, beta, dst_plane, ne0);
77            }
78        }
79    } catch (sycl::exception const& exc) {
80        std::cerr << exc.what() << std::endl;
81        GGML_ASSERT(false);
82    }
83}