diff --git a/examples/llama-bench/llama-bench.cpp b/examples/llama-bench/llama-bench.cpp index 89b3450..34b0233 100644 --- a/examples/llama-bench/llama-bench.cpp +++ b/examples/llama-bench/llama-bench.cpp @@ -265,6 +265,7 @@ struct cmd_params { bool fmoe = true; bool ger = false; // ger = Grouped Expert Routing bool no_fug = false; + bool no_mmad = false; bool use_thp = false; bool no_ooae = false; bool mqkv = false; @@ -314,6 +315,7 @@ static const cmd_params cmd_params_defaults = { /* fmoe */ true, /* ger */ false, /* no_fug */ false, + /* no_mmad */ false, /* use_thp */ false, /* no_ooae */ false, /* mqkv */ false, @@ -376,6 +378,7 @@ static void print_usage(int /* argc */, char ** argv) { printf(" -fmoe, --fused-moe <0|1> (default: %s)\n", cmd_params_defaults.fmoe? "1" : "0"); printf(" -ger, --grouped-expert-routing <0|1>(default: %s)\n", cmd_params_defaults.ger ? "1" : "0"); printf(" -no-fug, --no-fused-up-gate <0|1> (default: %s)\n", cmd_params_defaults.no_fug? "1" : "0"); + printf(" -no-mmad, --no-fused-mul-multiadd <0|1> (default: %s)\n", cmd_params_defaults.no_mmad? "1" : "0"); printf(" -no-ooae, --no-offload-only-active-experts <0|1> (default: %s)\n", cmd_params_defaults.no_ooae? "1" : "0"); printf(" -sas, --scheduler-async <0|1> (default: %s)\n", cmd_params_defaults.sas ? "1" : "0"); printf(" --fit <0|1> (default: %s)\n", cmd_params_defaults.fit ? "1" : "0"); @@ -872,6 +875,12 @@ static cmd_params parse_cmd_params(int argc, char ** argv) { break; } params.no_fug = std::stoi(argv[i]); + } else if (arg == "-no-mmad" || arg == "--no-fused-mul-multiadd") { + if (++i >= argc) { + invalid_param = true; + break; + } + params.no_mmad = std::stoi(argv[i]); } else if (arg == "-no-ooae" || arg == "--no-offload-only-active-experts") { if (++i >= argc) { invalid_param = true; @@ -982,6 +991,7 @@ struct cmd_params_instance { bool fmoe = true; bool ger = false; bool no_fug = false; + bool no_mmad = false; bool use_thp = false; bool no_ooae = false; bool mqkv = false; @@ -1057,6 +1067,7 @@ struct cmd_params_instance { cparams.grouped_expert_routing = ger; cparams.rope_cache = rcache; cparams.fused_up_gate = !no_fug; + cparams.fused_mmad = !no_mmad; cparams.only_active_experts = !no_ooae; cparams.min_experts = ser.first; cparams.thresh_experts = ser.second; @@ -1123,6 +1134,7 @@ static std::vector get_cmd_params_instances(const cmd_param /* .fmoe = */ params.fmoe, /* .ger = */ params.ger, /* .no_fug = */ params.no_fug, + /* .no_mmad = */ params.no_mmad, /* .use_thp = */ params.use_thp, /* .no_ooae = */ params.no_ooae, /* .mqkv = */ params.mqkv, @@ -1170,6 +1182,7 @@ static std::vector get_cmd_params_instances(const cmd_param /* .fmoe = */ params.fmoe, /* .ger = */ params.ger, /* .no_fug = */ params.no_fug, + /* .no_mmad = */ params.no_mmad, /* .use_thp = */ params.use_thp, /* .no_ooae = */ params.no_ooae, /* .mqkv = */ params.mqkv, @@ -1217,6 +1230,7 @@ static std::vector get_cmd_params_instances(const cmd_param /* .fmoe = */ params.fmoe, /* .ger = */ params.ger, /* .no_fug = */ params.no_fug, + /* .no_mmad = */ params.no_mmad, /* .use_thp = */ params.use_thp, /* .no_ooae = */ params.no_ooae, /* .mqkv = */ params.mqkv, @@ -1264,6 +1278,7 @@ static std::vector get_cmd_params_instances(const cmd_param /* .fmoe = */ params.fmoe, /* .ger = */ params.ger, /* .no_fug = */ params.no_fug, + /* .no_mmad = */ params.no_mmad, /* .use_thp = */ params.use_thp, /* .no_ooae = */ params.no_ooae, /* .mqkv = */ params.mqkv, @@ -1322,6 +1337,7 @@ struct test { bool fmoe = false; bool ger = false; bool no_fug = false; + bool no_mmad = false; bool use_thp = false; bool no_ooae = false; bool mqkv = false; @@ -1372,6 +1388,7 @@ struct test { defer_experts = inst.defer_experts; fmoe = inst.fmoe; ger = inst.ger; + no_mmad = inst.no_mmad; rcache = inst.rcache; sas = inst.sas; max_gpu = inst.max_gpu; @@ -1487,7 +1504,7 @@ struct test { if (field == "cuda" || field == "vulkan" || field == "kompute" || field == "metal" || field == "gpu_blas" || field == "blas" || field == "sycl" || field == "no_kv_offload" || field == "flash_attn" || field == "use_mmap" || field == "embeddings" || field == "repack" || field == "use_thp" || - field == "fused_moe" || field == "grouped_er" || field == "no_fused_up_gate" || field == "no_ooae" || field == "mqkv" || + field == "fused_moe" || field == "grouped_er" || field == "no_fused_up_gate" || field == "no_fused_mmad" || field == "no_ooae" || field == "mqkv" || field == "rcache" || field == "reuse" || field == "muge" || field == "defer_experts" || field == "sas") { return BOOL; } @@ -1532,7 +1549,7 @@ struct test { std::to_string(mla_attn), std::to_string(attn_max_batch), ser_to_string(ser), std::to_string(reuse), tensor_split_str, std::to_string(use_mmap), std::to_string(embeddings), std::to_string(repack), std::to_string(mqkv), std::to_string(muge), std::to_string(defer_experts), std::to_string(fmoe), std::to_string(ger), - std::to_string(no_fug), std::to_string(use_thp), std::to_string(no_ooae), std::to_string(rcache), std::to_string(sas), + std::to_string(no_fug), std::to_string(no_mmad), std::to_string(use_thp), std::to_string(no_ooae), std::to_string(rcache), std::to_string(sas), std::to_string(max_gpu), cuda_params, override_tensor, std::to_string(n_prompt), std::to_string(n_gen), test_time, @@ -1554,7 +1571,7 @@ struct test { "n_gpu_layers", "split_mode", "main_gpu", "no_kv_offload", "flash_attn", "mla_attn", "attn_max_batch", "ser", "reuse", "tensor_split", "use_mmap", "embeddings", "repack", "mqkv", "muge", "defer_experts", "fused_moe", "grouped_er", - "no_fused_up_gate", "use_thp", "no_ooae", "rcache", "sas", "max_gpu", "cuda_params", "override_tensor", + "no_fused_up_gate", "no_fused_mmad", "use_thp", "no_ooae", "rcache", "sas", "max_gpu", "cuda_params", "override_tensor", "n_prompt", "n_gen", "test_time", "avg_ns", "stddev_ns", "avg_ts", "stddev_ts", "test", @@ -1765,6 +1782,9 @@ struct markdown_printer : public printer { if (field == "no_fused_up_gate") { return 6; } + if (field == "no_fused_mmad") { + return 7; + } if (field == "no_ooae") { return 7; } @@ -1844,6 +1864,9 @@ struct markdown_printer : public printer { if (field == "no_fused_up_gate") { return "no-fug"; } + if (field == "no_fused_mmad") { + return "no-mmad"; + } if (field == "no_ooae") { return "no-ooae"; } @@ -1963,6 +1986,9 @@ struct markdown_printer : public printer { if (params.no_fug != cmd_params_defaults.no_fug) { fields.emplace_back("no_fused_up_gate"); } + if (params.no_mmad != cmd_params_defaults.no_mmad) { + fields.emplace_back("no_fused_mmad"); + } if (params.no_ooae != cmd_params_defaults.no_ooae) { fields.emplace_back("no_ooae"); } diff --git a/ggml/src/ggml-sycl.cpp b/ggml/src/ggml-sycl.cpp index e0f8536..fce7c66 100644 --- a/ggml/src/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl.cpp @@ -39,6 +39,7 @@ #include "ggml-sycl/backend.hpp" #include "ggml-sycl/common.hpp" #include "ggml-sycl/presets.hpp" +#include "ggml-sycl/vecdotq.hpp" bool ggml_sycl_loaded(void); void ggml_sycl_free_data(struct ggml_tensor * tensor); @@ -84,64 +85,12 @@ static __dpct_inline__ float op_add(const float a, const float b) { return a + b; } -static void ggml_backend_sycl_buffer_memset_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor, uint8_t value, - size_t offset, size_t size) { - GGML_SYCL_DEBUG(" [SYCL] call %s\n", __func__); - ggml_backend_sycl_buffer_context * ctx = (ggml_backend_sycl_buffer_context *) buffer->context; - SYCL_CHECK(ggml_sycl_set_device(ctx->device)); - auto stream = &(dpct::dev_mgr::instance().get_device(ctx->device).default_queue()); - if (size == 0) { - return; // Nothing to do - } - if (tensor->data == nullptr) { - GGML_ABORT("Error: Tensor data pointer is null.\n"); - } - void * target_ptr = static_cast(tensor->data) + offset; - SYCL_CHECK(CHECK_TRY_ERROR((*stream).memset(target_ptr, value, size))); - SYCL_CHECK(CHECK_TRY_ERROR((*stream).wait())); +static __dpct_inline__ float op_mul(const float a, const float b) { + return a * b; } -static void ggml_backend_sycl_buffer_reset(ggml_backend_buffer_t buffer) { - GGML_SYCL_DEBUG("[SYCL] call %s\n", __func__); - if (buffer == nullptr) { - return; - } - - ggml_backend_sycl_buffer_context * ctx = (ggml_backend_sycl_buffer_context *) buffer->context; - - if (ctx != nullptr) { - for (ggml_tensor_extra_gpu * extra : ctx->tensor_extras) { - release_extra_gpu(extra); - } - ctx->tensor_extras.clear(); // reset the tensor_extras vector - } -} - -static const ggml_backend_buffer_i ggml_backend_sycl_buffer_interface = { - /* .free_buffer = */ ggml_backend_sycl_buffer_free_buffer, - /* .get_base = */ ggml_backend_sycl_buffer_get_base, - /* .init_tensor = */ ggml_backend_sycl_buffer_init_tensor, - /* .memset_tensor = */ ggml_backend_sycl_buffer_memset_tensor, - /* .set_tensor = */ ggml_backend_sycl_buffer_set_tensor, - /* .get_tensor = */ ggml_backend_sycl_buffer_get_tensor, - /* .cpy_tensor = */ ggml_backend_sycl_buffer_cpy_tensor, - /* .clear = */ ggml_backend_sycl_buffer_clear, - /* .reset = */ ggml_backend_sycl_buffer_reset, -}; - -// sycl buffer type -struct ggml_backend_sycl_buffer_type_context { - int device; - std::string name; - - // each buffer type has its own stream - queue_ptr stream = nullptr; -}; - -static const char * ggml_backend_sycl_buffer_type_get_name(ggml_backend_buffer_type_t buft) { - ggml_backend_sycl_buffer_type_context * ctx = (ggml_backend_sycl_buffer_type_context *)buft->context; - - return ctx->name.c_str(); +static __dpct_inline__ float op_div(const float a, const float b) { + return a / b; } template @@ -311,6 +260,17 @@ static void hardsigmoid_f32(const float * x, float * dst, const int k, dst[i] = sycl::fmin(1.0f, sycl::fmax(0.0f, (x[i] + 3.0f) / 6.0f)); } +static void sigmoid_f32(const float * x, float * dst, const int k, + const sycl::nd_item<3> &item_ct1) { + const int i = item_ct1.get_local_range(2) * item_ct1.get_group(2) + + item_ct1.get_local_id(2); + + if (i >= k) { + return; + } + dst[i] = 1.0f / (1.0f + sycl::native::exp(-x[i])); +} + static void hardswish_f32(const float * x, float * dst, const int k, const sycl::nd_item<3> &item_ct1) { const int i = item_ct1.get_local_range(2) * item_ct1.get_group(2) + @@ -1340,6 +1300,18 @@ static void hardsigmoid_f32_sycl(const float *x, float *dst, const int k, }); } +static void sigmoid_f32_sycl(const float *x, float *dst, const int k, + queue_ptr stream) { + const int num_blocks = (k + SYCL_HARDSIGMOID_BLOCK_SIZE - 1) / SYCL_HARDSIGMOID_BLOCK_SIZE; + stream->parallel_for( + sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * + sycl::range<3>(1, 1, SYCL_HARDSIGMOID_BLOCK_SIZE), + sycl::range<3>(1, 1, SYCL_HARDSIGMOID_BLOCK_SIZE)), + [=](sycl::nd_item<3> item_ct1) { + sigmoid_f32(x, dst, k, item_ct1); + }); +} + static void hardswish_f32_sycl(const float *x, float *dst, const int k, queue_ptr stream) { const int num_blocks = (k + SYCL_HARDSWISH_BLOCK_SIZE - 1) / SYCL_HARDSWISH_BLOCK_SIZE; @@ -1729,6 +1701,148 @@ static void sum_rows_f32_sycl(const float *x, float *dst, const int ncols, }); } +static void multi_add_f32(const float *src0, float *dst, const int n_add, + const int64_t ne0, const int64_t ne1, + const int64_t nb1, const int64_t nb01, + const sycl::nd_item<3> &item_ct1) { + const int64_t i = item_ct1.get_local_range(2) * item_ct1.get_group(2) + + item_ct1.get_local_id(2); + const int64_t n = ne0 * ne1; + if (i >= n) { + return; + } + + const int64_t i1 = i / ne0; + const int64_t i0 = i % ne0; + const int64_t dst_stride = nb1 / sizeof(float); + const int64_t src_stride = nb01 / sizeof(float); + + const float *src_ptr = src0 + i1 * src_stride + i0; + float sum = src_ptr[0]; + for (int j = 1; j < n_add; ++j) { + sum += src_ptr[j * ne0]; + } + + dst[i1 * dst_stride + i0] = sum; +} + +static void multi_add_f32_sycl(const float *src0, float *dst, const int n_add, + const int64_t ne0, const int64_t ne1, + const int64_t nb1, const int64_t nb01, + queue_ptr stream) { + const int block_size = 256; + const int64_t n = ne0 * ne1; + const int64_t num_blocks = (n + block_size - 1) / block_size; + stream->parallel_for( + sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * + sycl::range<3>(1, 1, block_size), + sycl::range<3>(1, 1, block_size)), + [=](sycl::nd_item<3> item_ct1) { + multi_add_f32(src0, dst, n_add, ne0, ne1, nb1, nb01, item_ct1); + }); +} + +static void mul_multi_add_f32(const int n_add, const int64_t ne0, const int64_t ne1, + const int64_t nb1, const int64_t nb01, const int64_t nb02, + const int64_t nb11, const int64_t nb12, + const char *src0, const char *src1, char *dst, + const sycl::nd_item<3> &item_ct1) { + const int64_t i = item_ct1.get_local_range(2) * item_ct1.get_group(2) + + item_ct1.get_local_id(2); + const int64_t n = ne0 * ne1; + if (i >= n) { + return; + } + + const int64_t i1 = i / ne0; + const int64_t i0 = i % ne0; + float *result = (float *)(dst + i1 * nb1); + + const char *c0 = src0 + i1 * nb02; + const char *c1 = src1 + i1 * nb12; + + float sum = 0.0f; + for (int j = 0; j < n_add; ++j) { + const float *x0 = (const float *)c0; + const float *x1 = (const float *)c1; + sum += x0[i0] * x1[0]; + c0 += nb01; + c1 += nb11; + } + + result[i0] = sum; +} + +static void mul_multi_add_scaled_f32(const int n_add, const int64_t ne0, const int64_t ne1, + const int64_t nb1, const int64_t nb01, const int64_t nb02, + const int64_t nb11, const int64_t nb12, const int64_t nb31, + const char *src0, const char *src1, char *dst, + const float *scales, const char *cids, + const sycl::nd_item<3> &item_ct1) { + const int64_t i = item_ct1.get_local_range(2) * item_ct1.get_group(2) + + item_ct1.get_local_id(2); + const int64_t n = ne0 * ne1; + if (i >= n) { + return; + } + + const int64_t i1 = i / ne0; + const int64_t i0 = i % ne0; + float *result = (float *)(dst + i1 * nb1); + const int *ids = (const int *)(cids + i1 * nb31); + + const char *c0 = src0 + i1 * nb02; + const char *c1 = src1 + i1 * nb12; + + float sum = 0.0f; + for (int j = 0; j < n_add; ++j) { + const float *x0 = (const float *)c0; + const float *x1 = (const float *)c1; + sum += x0[i0] * x1[0] * scales[ids[j]]; + c0 += nb01; + c1 += nb11; + } + + result[i0] = sum; +} + +static void mul_multi_add_f32_sycl(const int n_add, const int64_t ne0, const int64_t ne1, + const int64_t nb1, const int64_t nb01, const int64_t nb02, + const int64_t nb11, const int64_t nb12, + const char *src0, const char *src1, char *dst, + queue_ptr stream) { + const int block_size = 256; + const int64_t n = ne0 * ne1; + const int64_t num_blocks = (n + block_size - 1) / block_size; + stream->parallel_for( + sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * + sycl::range<3>(1, 1, block_size), + sycl::range<3>(1, 1, block_size)), + [=](sycl::nd_item<3> item_ct1) { + mul_multi_add_f32(n_add, ne0, ne1, nb1, nb01, nb02, nb11, nb12, + src0, src1, dst, item_ct1); + }); +} + +static void mul_multi_add_scaled_f32_sycl(const int n_add, const int64_t ne0, const int64_t ne1, + const int64_t nb1, const int64_t nb01, const int64_t nb02, + const int64_t nb11, const int64_t nb12, const int64_t nb31, + const char *src0, const char *src1, char *dst, + const float *scales, const char *ids, + queue_ptr stream) { + const int block_size = 256; + const int64_t n = ne0 * ne1; + const int64_t num_blocks = (n + block_size - 1) / block_size; + stream->parallel_for( + sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * + sycl::range<3>(1, 1, block_size), + sycl::range<3>(1, 1, block_size)), + [=](sycl::nd_item<3> item_ct1) { + mul_multi_add_scaled_f32(n_add, ne0, ne1, nb1, nb01, nb02, nb11, nb12, nb31, + src0, src1, dst, scales, ids, item_ct1); + }); +} + static int next_power_of_2(int x) { int n = 1; while (n < x) { @@ -2114,12 +2228,12 @@ static dpct::err0 ggml_sycl_cpy_tensor_2d(void *dst, dpct::memcpy_direction kind; char * src_ptr; - if (src->backend == GGML_BACKEND_TYPE_CPU) { + if (src->__backend == GGML_BACKEND_TYPE_CPU) { kind = dpct::host_to_device; src_ptr = (char *) src->data; // GGML_SYCL_DEBUG("ggml_sycl_cpy_tensor_2d GGML_BACKEND_TYPE_CPU src_ptr %p\n", src_ptr); - } else if (src->backend == GGML_BACKEND_TYPE_GPU || src->backend == GGML_BACKEND_TYPE_GPU_SPLIT) { - GGML_ASSERT(src->backend != GGML_BACKEND_TYPE_GPU_SPLIT || (i1_low == 0 && i1_high == src->ne[1])); + } else if (src->__backend == GGML_BACKEND_TYPE_GPU || src->__backend == GGML_BACKEND_TYPE_GPU_SPLIT) { + GGML_ASSERT(src->__backend != GGML_BACKEND_TYPE_GPU_SPLIT || (i1_low == 0 && i1_high == src->ne[1])); kind = dpct::device_to_device; ggml_tensor_extra_gpu * extra = (ggml_tensor_extra_gpu *) src->extra; int id; @@ -2411,6 +2525,22 @@ static void ggml_sycl_op_hardswish(ggml_backend_sycl_context & ctx, const ggml_t (void) src1_dd; } +static void ggml_sycl_op_sigmoid(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, + const ggml_tensor *src1, ggml_tensor *dst, + const float *src0_dd, const float *src1_dd, + float *dst_dd, const queue_ptr &main_stream) { + + GGML_ASSERT(src0->type == GGML_TYPE_F32); + GGML_ASSERT( dst->type == GGML_TYPE_F32); + + sigmoid_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream); + + (void) ctx; + (void) src1; + (void) dst; + (void) src1_dd; +} + inline void ggml_sycl_op_leaky_relu(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1, ggml_tensor *dst, const float *src0_dd, const float *src1_dd, @@ -2749,6 +2879,72 @@ inline void ggml_sycl_op_sum_rows(ggml_backend_sycl_context & ctx, const ggml_te (void) src1_dd; } +inline void ggml_sycl_op_multi_add(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, + const ggml_tensor *src1, ggml_tensor *dst, + const float *src0_dd, const float *src1_dd, + float *dst_dd, + const queue_ptr &main_stream) { + + GGML_ASSERT(src0->type == GGML_TYPE_F32); + GGML_ASSERT( dst->type == GGML_TYPE_F32); + GGML_ASSERT(src0->nb[0] == sizeof(float)); + GGML_ASSERT( dst->nb[0] == sizeof(float)); + GGML_ASSERT(dst->ne[2] == 1 && dst->ne[3] == 1); + + const int n_add = dst->op_params[0]; + GGML_ASSERT(n_add > 0); + + multi_add_f32_sycl(src0_dd, dst_dd, n_add, dst->ne[0], dst->ne[1], + dst->nb[1], src0->nb[1], main_stream); + + (void) ctx; + (void) src1; + (void) src1_dd; +} + +inline void ggml_sycl_op_mul_multi_add(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, + const ggml_tensor *src1, ggml_tensor *dst, + const float *src0_dd, const float *src1_dd, + float *dst_dd, + const queue_ptr &main_stream) { + + GGML_ASSERT(src0->type == GGML_TYPE_F32); + GGML_ASSERT(src1->type == GGML_TYPE_F32); + GGML_ASSERT( dst->type == GGML_TYPE_F32); + GGML_ASSERT(src0->ne[0] == dst->ne[0]); + GGML_ASSERT(src0->ne[2] == dst->ne[1]); + GGML_ASSERT(src0->ne[1] == src1->ne[1]); + GGML_ASSERT(src0->ne[2] == src1->ne[2]); + GGML_ASSERT(src0->ne[3] == src1->ne[3]); + GGML_ASSERT(src0->ne[3] == 1); + GGML_ASSERT(src1->ne[0] == 1); + + const int n_add = src0->ne[1]; + const ggml_tensor *src2 = dst->src[2]; + const ggml_tensor *src3 = dst->src[3]; + + if (src2 && src3) { + GGML_ASSERT(src3->ne[0] == src0->ne[1]); + GGML_ASSERT(src3->type == GGML_TYPE_I32); + GGML_ASSERT(src2->type == GGML_TYPE_F32); + + mul_multi_add_scaled_f32_sycl(n_add, dst->ne[0], dst->ne[1], + dst->nb[1], src0->nb[1], src0->nb[2], + src1->nb[1], src1->nb[2], src3->nb[1], + (const char *)src0_dd, (const char *)src1_dd, (char *)dst_dd, + (const float *)src2->data, (const char *)src3->data, + main_stream); + } else { + mul_multi_add_f32_sycl(n_add, dst->ne[0], dst->ne[1], + dst->nb[1], src0->nb[1], src0->nb[2], + src1->nb[1], src1->nb[2], + (const char *)src0_dd, (const char *)src1_dd, (char *)dst_dd, + main_stream); + } + + (void) ctx; +} + inline void ggml_sycl_op_argsort(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1, ggml_tensor *dst, const float *src0_dd, const float *src1_dd, @@ -2848,8 +3044,8 @@ static void ggml_sycl_op_flatten(ggml_backend_sycl_context & ctx, const ggml_ten const bool use_src1 = src1 != nullptr; const int64_t nrows1 = use_src1 ? ggml_nrows(src1) : 1; - GGML_ASSERT(!use_src1 || src1->backend != GGML_BACKEND_TYPE_GPU_SPLIT); - GGML_ASSERT( dst->backend != GGML_BACKEND_TYPE_GPU_SPLIT); + GGML_ASSERT(!use_src1 || src1->__backend != GGML_BACKEND_TYPE_GPU_SPLIT); + GGML_ASSERT( dst->__backend != GGML_BACKEND_TYPE_GPU_SPLIT); ggml_tensor_extra_gpu * src0_extra = (ggml_tensor_extra_gpu *) src0->extra; ggml_tensor_extra_gpu * src1_extra = use_src1 ? (ggml_tensor_extra_gpu *) src1->extra : nullptr; @@ -2943,8 +3139,8 @@ static void ggml_sycl_op_mul_mat(ggml_backend_sycl_context & ctx, const ggml_ten const int nb2 = dst->nb[2]; const int nb3 = dst->nb[3]; - GGML_ASSERT(dst->backend != GGML_BACKEND_TYPE_GPU_SPLIT); - GGML_ASSERT(src1->backend != GGML_BACKEND_TYPE_GPU_SPLIT); + GGML_ASSERT(dst->__backend != GGML_BACKEND_TYPE_GPU_SPLIT); + GGML_ASSERT(src1->__backend != GGML_BACKEND_TYPE_GPU_SPLIT); GGML_ASSERT(src1->type == GGML_TYPE_F32 || (src1->ne[2] == 1 && src1->ne[3] == 1)); GGML_ASSERT(ne12 >= ne02 && ne12 % ne02 == 0); @@ -2965,7 +3161,7 @@ static void ggml_sycl_op_mul_mat(ggml_backend_sycl_context & ctx, const ggml_ten int64_t src1_padded_col_size = GGML_PAD(ne10, MATRIX_ROW_PADDING); - const bool split = src0->backend == GGML_BACKEND_TYPE_GPU_SPLIT; + const bool split = src0->__backend == GGML_BACKEND_TYPE_GPU_SPLIT; GGML_ASSERT(!(split && ne02 > 1)); GGML_ASSERT(!(split && ne03 > 1)); GGML_ASSERT(!(split && ne02 < ne12)); @@ -3329,6 +3525,12 @@ static void ggml_sycl_hardswish(ggml_backend_sycl_context & ctx, const ggml_tens GGML_SYCL_DEBUG("call %s done\n", __func__); } +static void ggml_sycl_sigmoid(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + GGML_SYCL_DEBUG("call %s\n", __func__); + ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_sigmoid); + GGML_SYCL_DEBUG("call %s done\n", __func__); +} + static void ggml_sycl_leaky_relu(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { GGML_SYCL_DEBUG("call %s\n", __func__); ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_leaky_relu); @@ -3376,7 +3578,7 @@ static void ggml_sycl_mul_mat_vec_p021(ggml_backend_sycl_context & ctx, const gg const ggml_tensor *src1, ggml_tensor *dst) try { GGML_ASSERT(ggml_is_permuted(src0) && ggml_is_permuted(src1)); - GGML_ASSERT(src0->backend != GGML_BACKEND_TYPE_GPU_SPLIT); + GGML_ASSERT(src0->__backend != GGML_BACKEND_TYPE_GPU_SPLIT); GGML_ASSERT(src0->nb[0] <= src0->nb[1] && src0->nb[2] <= src0->nb[3]); // 0213 permutation GGML_ASSERT(src1->nb[0] <= src1->nb[1] && src1->nb[2] <= src1->nb[3]); // 0213 permutation GGML_ASSERT(src0->type == GGML_TYPE_F16); @@ -3409,7 +3611,7 @@ static void ggml_sycl_mul_mat_vec_nc(ggml_backend_sycl_context & ctx, const ggml GGML_ASSERT(!ggml_is_transposed(src0)); GGML_ASSERT(!ggml_is_transposed(src1)); GGML_ASSERT(!ggml_is_permuted(src0)); - GGML_ASSERT(src0->backend != GGML_BACKEND_TYPE_GPU_SPLIT); + GGML_ASSERT(src0->__backend != GGML_BACKEND_TYPE_GPU_SPLIT); GGML_ASSERT(src0->type == GGML_TYPE_F16); GGML_ASSERT(src1->type == GGML_TYPE_F32); @@ -3471,7 +3673,7 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, ggml_tensor *dst) try { GGML_ASSERT(!ggml_is_transposed(src0)); GGML_ASSERT(!ggml_is_transposed(src1)); - GGML_ASSERT(src0->backend != GGML_BACKEND_TYPE_GPU_SPLIT); + GGML_ASSERT(src0->__backend != GGML_BACKEND_TYPE_GPU_SPLIT); GGML_ASSERT(src0->type == GGML_TYPE_F16); GGML_TENSOR_BINARY_OP_LOCALS @@ -3736,6 +3938,208 @@ __dpct_inline__ static void k_copy_dst_from_contiguous( } } +static __dpct_inline__ float ggml_sycl_apply_moe_gate_unary(const float x, const int unary_op, const float limit) { + float y; + switch ((ggml_unary_op) unary_op) { + case GGML_UNARY_OP_SILU: + y = x / (1.0f + sycl::native::exp(-x)); + break; + case GGML_UNARY_OP_GELU: + { + const float GELU_COEF_A = 0.044715f; + const float SQRT_2_OVER_PI = 0.79788456080286535587989211986876f; + y = 0.5f * x * (1.0f + sycl::tanh(SQRT_2_OVER_PI * x * (1.0f + GELU_COEF_A * x * x))); + } + break; + case GGML_UNARY_OP_RELU: + y = sycl::fmax(0.0f, x); + break; + default: + y = x; + break; + } + + return limit > 1e-6f ? sycl::fmin(y, limit) : y; +} + +static void moe_up_gate_combine_f32(const float * up, const float * gate, float * dst, + const int64_t n, const int unary_op, const float limit, + const sycl::nd_item<3> & item_ct1) { + const int64_t i = item_ct1.get_local_range(2) * item_ct1.get_group(2) + + item_ct1.get_local_id(2); + if (i >= n) { + return; + } + + dst[i] = up[i] * ggml_sycl_apply_moe_gate_unary(gate[i], unary_op, limit); +} + +static void moe_up_gate_combine_merged_f32(const float * up_gate, float * dst, + const int64_t ne0, const int64_t n, + const int unary_op, const float limit, + const sycl::nd_item<3> & item_ct1) { + const int64_t i = item_ct1.get_local_range(2) * item_ct1.get_group(2) + + item_ct1.get_local_id(2); + if (i >= n) { + return; + } + + const int64_t i0 = i % ne0; + const int64_t row = i / ne0; + const float * row_ptr = up_gate + row * 2 * ne0; + const float gate = row_ptr[i0]; + const float up = row_ptr[ne0 + i0]; + dst[i] = up * ggml_sycl_apply_moe_gate_unary(gate, unary_op, limit); +} + +static void moe_up_gate_combine_f32_sycl(const float * up, const float * gate, float * dst, + const int64_t n, const int unary_op, const float limit, + queue_ptr stream) { + const int block_size = 256; + const int64_t num_blocks = (n + block_size - 1) / block_size; + stream->parallel_for( + sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * + sycl::range<3>(1, 1, block_size), + sycl::range<3>(1, 1, block_size)), + [=](sycl::nd_item<3> item_ct1) { + moe_up_gate_combine_f32(up, gate, dst, n, unary_op, limit, item_ct1); + }); +} + +static void moe_up_gate_combine_merged_f32_sycl(const float * up_gate, float * dst, + const int64_t ne0, const int64_t n, + const int unary_op, const float limit, + queue_ptr stream) { + const int block_size = 256; + const int64_t num_blocks = (n + block_size - 1) / block_size; + stream->parallel_for( + sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * + sycl::range<3>(1, 1, block_size), + sycl::range<3>(1, 1, block_size)), + [=](sycl::nd_item<3> item_ct1) { + moe_up_gate_combine_merged_f32(up_gate, dst, ne0, n, unary_op, limit, item_ct1); + }); +} + +static __dpct_inline__ float moe_iq4_xs_q8_1_dot(const char * row, const block_q8_1 * y, + const int ncols, + const sycl::nd_item<3> & item_ct1) { + const int blocks_per_row = ncols / QK_K; + constexpr int qi = QI4_XS / 4; + constexpr int vdr = 1; + const int blocks_per_warp = vdr * WARP_SIZE / qi; + + float tmp = 0.0f; + const block_iq4_xs * x = (const block_iq4_xs *) row; + + for (int i = item_ct1.get_local_id(2) / (qi / vdr); i < blocks_per_row; + i += blocks_per_warp) { + const int iby = i * (QK_K / QK8_1); + const int iqs = vdr * (item_ct1.get_local_id(2) % (qi / vdr)); + tmp += vec_dot_iq4_xs_q8_1(&x[i], &y[iby], iqs); + } + +#pragma unroll + for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) { + tmp += dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask); + } + + return tmp; +} + +static void moe_up_gate_merged_iq4_xs_q8_1_sycl(const char * up_gate, const char * inp_q8_1, + const char * ids, float * dst, + const int64_t ncols, const int64_t n_ff, + const int64_t n_ids, const int64_t n_tokens, + const size_t nb01, const size_t nb02, + const size_t ids_nb0, const size_t ids_nb1, + const size_t dst_nb1, const size_t dst_nb2, + const int64_t inp_padded_col_size, + const int unary_op, const float limit, + queue_ptr stream) { + GGML_ASSERT(ncols % QK_K == 0); + const size_t q8_1_ts = sizeof(block_q8_1); + const size_t q8_1_bs = QK8_1; + const int block_rows = GGML_SYCL_MMV_Y; + const int64_t row_blocks = (n_ff + block_rows - 1) / block_rows; + + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for( + sycl::nd_range<3>( + sycl::range<3>(n_tokens, row_blocks, n_ids) * + sycl::range<3>(1, block_rows, WARP_SIZE), + sycl::range<3>(1, block_rows, WARP_SIZE)), + [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(WARP_SIZE)]] { + const int64_t token = item_ct1.get_group(0); + const int64_t row = item_ct1.get_group(1) * item_ct1.get_local_range(1) + + item_ct1.get_local_id(1); + const int64_t id = item_ct1.get_group(2); + if (row >= n_ff) { + return; + } + + const int32_t expert = *(const int32_t *)(ids + token * ids_nb1 + id * ids_nb0); + const char * expert_base = up_gate + expert * nb02; + const char * gate_row = expert_base + row * nb01; + const char * up_row = expert_base + (row + n_ff) * nb01; + const block_q8_1 * y = (const block_q8_1 *) + (inp_q8_1 + token * inp_padded_col_size * q8_1_ts / q8_1_bs); + + const float gate = moe_iq4_xs_q8_1_dot(gate_row, y, ncols, item_ct1); + const float up = moe_iq4_xs_q8_1_dot(up_row, y, ncols, item_ct1); + + if (item_ct1.get_local_id(2) == 0) { + float * dst_row = (float *)((char *)dst + token * dst_nb2 + id * dst_nb1); + dst_row[row] = up * ggml_sycl_apply_moe_gate_unary(gate, unary_op, limit); + } + }); + }); +} + +static void mul_mat_id_iq4_xs_q8_1_sycl(const char * weights, const char * src1_q8_1, + const char * ids, float * dst, + const int64_t ncols, const int64_t nrows, + const int64_t n_ids, const int64_t n_tokens, + const size_t weights_nb1, const size_t weights_nb2, + const size_t ids_nb0, const size_t ids_nb1, + const size_t dst_nb1, const size_t dst_nb2, + const int64_t src1_padded_col_size, + queue_ptr stream) { + GGML_ASSERT(ncols % QK_K == 0); + const size_t q8_1_ts = sizeof(block_q8_1); + const size_t q8_1_bs = QK8_1; + const int block_rows = GGML_SYCL_MMV_Y; + const int64_t row_blocks = (nrows + block_rows - 1) / block_rows; + + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for( + sycl::nd_range<3>( + sycl::range<3>(n_tokens, row_blocks, n_ids) * + sycl::range<3>(1, block_rows, WARP_SIZE), + sycl::range<3>(1, block_rows, WARP_SIZE)), + [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(WARP_SIZE)]] { + const int64_t token = item_ct1.get_group(0); + const int64_t row = item_ct1.get_group(1) * item_ct1.get_local_range(1) + + item_ct1.get_local_id(1); + const int64_t id = item_ct1.get_group(2); + if (row >= nrows) { + return; + } + + const int32_t expert = *(const int32_t *)(ids + token * ids_nb1 + id * ids_nb0); + const char * row_ptr = weights + expert * weights_nb2 + row * weights_nb1; + const block_q8_1 * y = (const block_q8_1 *) + (src1_q8_1 + (token * n_ids + id) * src1_padded_col_size * q8_1_ts / q8_1_bs); + const float value = moe_iq4_xs_q8_1_dot(row_ptr, y, ncols, item_ct1); + + if (item_ct1.get_local_id(2) == 0) { + float * dst_row = (float *)((char *)dst + token * dst_nb2 + id * dst_nb1); + dst_row[row] = value; + } + }); + }); +} + static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1, ggml_tensor *dst) try { @@ -3752,6 +4156,37 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx, const ggml_ten std::vector ids_host(ggml_nbytes(ids)); const char * ids_dev = (const char *) ids->data; + static const bool use_fast_iq4_xs_mmid = + get_sycl_env("GGML_SYCL_FAST_MUL_MAT_ID_IQ4_XS", 0) != 0; + + if (use_fast_iq4_xs_mmid && + src0->type == GGML_TYPE_IQ4_XS && + src1->type == GGML_TYPE_F32 && + dst->type == GGML_TYPE_F32 && + ids->type == GGML_TYPE_I32 && + ggml_is_contiguous(src0) && + ggml_is_contiguous(src1) && + ggml_is_contiguous(dst) && + src1->ne[0] == src0->ne[0] && + dst->ne[0] == src0->ne[1] && + src1->ne[1] == n_ids && + ids->ne[1] == dst->ne[2] && + ggml_nrows(src1) == n_ids * ids->ne[1]) { + const int64_t n_tokens = ids->ne[1]; + const int64_t src1_padded_col_size = GGML_PAD(src1->ne[0], MATRIX_ROW_PADDING); + const size_t q8_1_bytes = n_tokens * n_ids * src1_padded_col_size * sizeof(block_q8_1) / QK8_1; + ggml_sycl_pool_alloc src1_q8_1(ctx.pool(), q8_1_bytes); + + quantize_row_q8_1_sycl((const float *) src1->data, src1_q8_1.get(), src1->ne[0], + n_tokens * n_ids, src1_padded_col_size, stream); + mul_mat_id_iq4_xs_q8_1_sycl((const char *) src0->data, src1_q8_1.get(), ids_dev, + (float *) dst->data, src1->ne[0], dst->ne[0], + n_ids, n_tokens, src0->nb[1], src0->nb[2], + ids->nb[0], ids->nb[1], dst->nb[1], dst->nb[2], + src1_padded_col_size, stream); + return; + } + SYCL_CHECK(CHECK_TRY_ERROR( stream->memcpy(ids_host.data(), ids_dev, ggml_nbytes(ids)))); SYCL_CHECK(CHECK_TRY_ERROR(stream->wait())); @@ -3904,6 +4339,107 @@ catch (sycl::exception const &exc) { std::exit(1); } +static void ggml_sycl_moe_fused_up_gate(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, + const ggml_tensor * src1, ggml_tensor * dst) try { + GGML_UNUSED(src0); + GGML_UNUSED(src1); + + const ggml_tensor * up_exps = dst->src[0]; + const ggml_tensor * gate_exps = dst->src[1]; + const ggml_tensor * inp = dst->src[2]; + const ggml_tensor * ids = dst->src[3]; + + GGML_ASSERT(dst->type == GGML_TYPE_F32); + GGML_ASSERT(inp->type == GGML_TYPE_F32); + GGML_ASSERT(ids->type == GGML_TYPE_I32); + GGML_ASSERT(!dst->src[4] && !dst->src[5] && "SYCL MOE_FUSED_UP_GATE bias path is not implemented"); + GGML_ASSERT(ggml_is_contiguous(dst)); + + const int unary_op = dst->op_params[0]; + GGML_ASSERT(unary_op == GGML_UNARY_OP_SILU || unary_op == GGML_UNARY_OP_GELU || unary_op == GGML_UNARY_OP_RELU); + const float limit = *(const float *)(dst->op_params + 1); + + const int64_t n = ggml_nelements(dst); + const size_t dst_bytes = n * sizeof(float); + + queue_ptr stream = ctx.stream(); + + if (gate_exps) { + ggml_sycl_pool_alloc up_tmp(ctx.pool(), dst_bytes); + ggml_sycl_pool_alloc gate_tmp(ctx.pool(), dst_bytes); + + ggml_tensor up_node = *dst; + up_node.op = GGML_OP_MUL_MAT_ID; + up_node.data = up_tmp.get(); + up_node.src[0] = const_cast(up_exps); + up_node.src[1] = const_cast(inp); + up_node.src[2] = const_cast(ids); + up_node.src[3] = nullptr; + up_node.src[4] = nullptr; + up_node.src[5] = nullptr; + + ggml_tensor gate_node = up_node; + gate_node.data = gate_tmp.get(); + gate_node.src[0] = const_cast(gate_exps); + + ggml_sycl_mul_mat_id(ctx, up_exps, inp, &up_node); + ggml_sycl_mul_mat_id(ctx, gate_exps, inp, &gate_node); + + moe_up_gate_combine_f32_sycl((const float *) up_tmp.get(), (const float *) gate_tmp.get(), + (float *) dst->data, n, unary_op, limit, stream); + } else { + const int64_t ne0 = dst->ne[0]; + const int64_t n_tokens = dst->ne[2] * dst->ne[3]; + if (up_exps->type == GGML_TYPE_IQ4_XS && + ggml_is_contiguous(up_exps) && + ggml_is_contiguous(inp) && + inp->ne[0] == up_exps->ne[0] && + ids->ne[0] == dst->ne[1] && + n_tokens == ggml_nrows(inp)) { + const int64_t inp_padded_col_size = GGML_PAD(inp->ne[0], MATRIX_ROW_PADDING); + const size_t q8_1_bytes = n_tokens * inp_padded_col_size * sizeof(block_q8_1) / QK8_1; + ggml_sycl_pool_alloc inp_q8_1(ctx.pool(), q8_1_bytes); + + quantize_row_q8_1_sycl((const float *) inp->data, inp_q8_1.get(), inp->ne[0], + n_tokens, inp_padded_col_size, stream); + moe_up_gate_merged_iq4_xs_q8_1_sycl((const char *) up_exps->data, inp_q8_1.get(), + (const char *) ids->data, (float *) dst->data, + inp->ne[0], ne0, ids->ne[0], n_tokens, + up_exps->nb[1], up_exps->nb[2], + ids->nb[0], ids->nb[1], dst->nb[1], dst->nb[2], + inp_padded_col_size, unary_op, limit, stream); + return; + } + + ggml_sycl_pool_alloc up_gate_tmp(ctx.pool(), 2 * dst_bytes); + + ggml_tensor up_gate_node = *dst; + up_gate_node.op = GGML_OP_MUL_MAT_ID; + up_gate_node.ne[0] = 2 * ne0; + up_gate_node.nb[0] = sizeof(float); + up_gate_node.nb[1] = up_gate_node.ne[0] * up_gate_node.nb[0]; + up_gate_node.nb[2] = up_gate_node.ne[1] * up_gate_node.nb[1]; + up_gate_node.nb[3] = up_gate_node.ne[2] * up_gate_node.nb[2]; + up_gate_node.data = up_gate_tmp.get(); + up_gate_node.src[0] = const_cast(up_exps); + up_gate_node.src[1] = const_cast(inp); + up_gate_node.src[2] = const_cast(ids); + up_gate_node.src[3] = nullptr; + up_gate_node.src[4] = nullptr; + up_gate_node.src[5] = nullptr; + + ggml_sycl_mul_mat_id(ctx, up_exps, inp, &up_gate_node); + + moe_up_gate_combine_merged_f32_sycl((const float *) up_gate_tmp.get(), (float *) dst->data, + ne0, n, unary_op, limit, stream); + } +} +catch (sycl::exception const &exc) { + std::cerr << exc.what() << "Exception caught at file:" << __FILE__ + << ", line:" << __LINE__ << std::endl; + std::exit(1); +} + static void ggml_sycl_scale(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_scale); } @@ -3992,6 +4528,14 @@ static void ggml_sycl_sum_rows(ggml_backend_sycl_context & ctx, const ggml_tenso ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_sum_rows); } +static void ggml_sycl_multi_add(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_multi_add); +} + +static void ggml_sycl_mul_multi_add(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_mul_multi_add); +} + static void ggml_sycl_argsort(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { GGML_ASSERT(ggml_is_contiguous(src0)); ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_argsort); @@ -4032,6 +4576,10 @@ bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct ggml_tens if (!g_sycl_loaded) return false; ggml_sycl_func_t func; + static const bool op_timing = []() { + const char * env = getenv("GGML_SYCL_OP_TIMING"); + return env && strcmp(env, "0") != 0; + }(); switch (tensor->op) { case GGML_OP_CONV_TRANSPOSE_1D: @@ -4081,6 +4629,9 @@ bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct ggml_tens case GGML_UNARY_OP_HARDSWISH: func = ggml_sycl_hardswish; break; + case GGML_UNARY_OP_SIGMOID: + func = ggml_sycl_sigmoid; + break; default: return false; } @@ -4118,6 +4669,12 @@ bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct ggml_tens } func = ggml_sycl_mul_mat_id; break; + case GGML_OP_MOE_FUSED_UP_GATE: + if (tensor->src[0]->ne[3] != tensor->src[2]->ne[3]) { + return false; + } + func = ggml_sycl_moe_fused_up_gate; + break; case GGML_OP_SCALE: func = ggml_sycl_scale; break; @@ -4158,6 +4715,12 @@ bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct ggml_tens case GGML_OP_SUM_ROWS: func = ggml_sycl_sum_rows; break; + case GGML_OP_MULTI_ADD: + func = ggml_sycl_multi_add; + break; + case GGML_OP_MUL_MULTI_ADD: + func = ggml_sycl_mul_multi_add; + break; case GGML_OP_ARGSORT: func = ggml_sycl_argsort; break; @@ -4172,7 +4735,16 @@ bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct ggml_tens ggml_sycl_set_peer_access(tensor->src[1]->ne[1], ctx.device); } + const int64_t timing_start_us = op_timing ? ggml_time_us() : 0; func(ctx, tensor->src[0], tensor->src[1], tensor); + if (op_timing) { + SYCL_CHECK(CHECK_TRY_ERROR(ctx.stream()->wait())); + const int64_t timing_us = ggml_time_us() - timing_start_us; + fprintf(stderr, + "[SYCL_OP_TIMING] dev=%d op=%s name=%s ne=%" PRId64 "x%" PRId64 "x%" PRId64 "x%" PRId64 " us=%" PRId64 "\n", + ctx.device, ggml_op_name(tensor->op), tensor->name, + tensor->ne[0], tensor->ne[1], tensor->ne[2], tensor->ne[3], timing_us); + } return true; } @@ -4282,6 +4854,23 @@ GGML_CALL static bool ggml_backend_buffer_is_sycl(ggml_backend_buffer_t buffer) return buffer->iface.get_name == ggml_backend_sycl_buffer_get_name; } +static void ggml_backend_sycl_buffer_memset_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor, uint8_t value, + size_t offset, size_t size) { + GGML_SYCL_DEBUG(" [SYCL] call %s\n", __func__); + ggml_backend_sycl_buffer_context * ctx = (ggml_backend_sycl_buffer_context *) buffer->context; + SYCL_CHECK(ggml_sycl_set_device(ctx->device)); + auto stream = &(dpct::dev_mgr::instance().get_device(ctx->device).default_queue()); + if (size == 0) { + return; // Nothing to do + } + if (tensor->data == nullptr) { + GGML_ABORT("Error: Tensor data pointer is null.\n"); + } + void * target_ptr = static_cast(tensor->data) + offset; + SYCL_CHECK(CHECK_TRY_ERROR((*stream).memset(target_ptr, value, size))); + SYCL_CHECK(CHECK_TRY_ERROR((*stream).wait())); +} + static void ggml_backend_sycl_buffer_free_buffer(ggml_backend_buffer_t buffer) try { ggml_backend_sycl_buffer_context * ctx = ( ggml_backend_sycl_buffer_context *)buffer->context; @@ -4307,7 +4896,7 @@ ggml_backend_sycl_buffer_init_tensor(ggml_backend_buffer_t buffer, if (tensor->view_src != NULL && tensor->view_offs == 0) { assert(tensor->view_src->buffer->buft == buffer->buft); - tensor->backend = tensor->view_src->backend; + tensor->__backend = tensor->view_src->__backend; tensor->extra = tensor->view_src->extra; return; } @@ -4462,7 +5051,6 @@ static struct ggml_backend_buffer_i ggml_backend_sycl_buffer_interface = { /* .init_tensor = */ ggml_backend_sycl_buffer_init_tensor, /* .memset_tensor = */ ggml_backend_sycl_buffer_memset_tensor, /* .set_tensor = */ ggml_backend_sycl_buffer_set_tensor, - /* .memset_tensor =*/ NULL, /* .get_tensor = */ ggml_backend_sycl_buffer_get_tensor, /* .cpy_tensor = */ ggml_backend_sycl_buffer_cpy_tensor, /* .clear = */ ggml_backend_sycl_buffer_clear, @@ -4494,6 +5082,14 @@ ggml_backend_sycl_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft, void * dev_ptr; SYCL_CHECK(CHECK_TRY_ERROR(dev_ptr = (void *)sycl::malloc_device( size, *stream))); + if (dev_ptr == nullptr) { + const size_t max_alloc = dpct::get_current_device().get_max_mem_alloc_size(); + size_t free_mem = 0; + size_t total_mem = 0; + dpct::dev_mgr::instance().get_device(buft_ctx->device).get_memory_info(free_mem, total_mem); + GGML_ABORT("%s: sycl::malloc_device returned null: device=%d size=%zu max_alloc=%zu free=%zu total=%zu\n", + __func__, buft_ctx->device, size, max_alloc, free_mem, total_mem); + } ggml_backend_sycl_buffer_context * ctx = new ggml_backend_sycl_buffer_context(buft_ctx->device, dev_ptr, buft_ctx->stream); return ggml_backend_buffer_init(buft, ggml_backend_sycl_buffer_interface, ctx, size); } @@ -4608,6 +5204,14 @@ static void get_row_split(int64_t * row_low, int64_t * row_high, const ggml_tens } struct ggml_backend_sycl_split_buffer_context { + ggml_backend_sycl_split_buffer_context() { + streams.resize(ggml_sycl_info().device_count); + pools.resize(ggml_sycl_info().device_count); + for (int i = 0; i < ggml_sycl_info().device_count; ++i) { + streams[i] = &(dpct::dev_mgr::instance().get_device(i).default_queue()); + } + } + ~ggml_backend_sycl_split_buffer_context() try { for (ggml_tensor_extra_gpu * extra : tensor_extras) { for (int i = 0; i < ggml_sycl_info().device_count; ++i) { @@ -4623,20 +5227,17 @@ struct ggml_backend_sycl_split_buffer_context { dpct::destroy_event(extra->events[i][is]))); } } - if (extra->data_device[i] != nullptr) { - /* - DPCT1009:207: SYCL uses exceptions to report errors and does - not use the error codes. The original code was commented out - and a warning string was inserted. You need to rewrite this - code. - */ - ggml_sycl_set_device(i); - SYCL_CHECK(CHECK_TRY_ERROR(sycl::free( - extra->data_device[i], *(streams[i])))); - } } delete extra; } + for (int i = 0; i < int(pools.size()); ++i) { + ggml_sycl_set_device(i); + for (split_pool_chunk & chunk : pools[i]) { + if (chunk.data != nullptr) { + SYCL_CHECK(CHECK_TRY_ERROR(sycl::free(chunk.data, *(streams[i])))); + } + } + } } catch (sycl::exception const &exc) { std::cerr << exc.what() << "Exception caught at file:" << __FILE__ @@ -4644,8 +5245,149 @@ struct ggml_backend_sycl_split_buffer_context { std::exit(1); } + struct split_pool_chunk { + char * data = nullptr; + size_t size = 0; + size_t offset = 0; + }; + + static size_t align_up(size_t value, size_t alignment) { + return (value + alignment - 1) / alignment * alignment; + } + + size_t pool_size(int device) const { + size_t total = 0; + for (const split_pool_chunk & chunk : pools[device]) { + total += chunk.size; + } + return total; + } + + size_t pool_used(int device) const { + size_t total = 0; + for (const split_pool_chunk & chunk : pools[device]) { + total += chunk.offset; + } + return total; + } + + void log_pool_state(const char * prefix, int device) const { + fprintf(stderr, "%s: device=%d pool_used=%zu pool_size=%zu chunks=%zu\n", + prefix, device, pool_used(device), pool_size(device), pools[device].size()); + } + + char * alloc(int device, size_t size, const ggml_tensor * tensor = nullptr) { + constexpr size_t alignment = 256; + constexpr size_t preferred_chunk_size = size_t(1024) * 1024 * 1024; + const size_t aligned_size = align_up(size, alignment); + const int debug = get_sycl_env("GGML_SYCL_SPLIT_POOL_DEBUG", 0); + const bool exact_fallback = get_sycl_env("GGML_SYCL_SPLIT_ALLOC_EXACT", 0) != 0; + + for (split_pool_chunk & chunk : pools[device]) { + const size_t offset = align_up(chunk.offset, alignment); + if (offset + aligned_size <= chunk.size) { + chunk.offset = offset + aligned_size; + if (debug >= 2) { + fprintf(stderr, "split_pool_alloc: device=%d tensor=%s request=%zu aligned=%zu offset=%zu used=%zu pool=%zu chunks=%zu\n", + device, tensor ? tensor->name : "(null)", size, aligned_size, offset, + pool_used(device), pool_size(device), pools[device].size()); + } + return chunk.data + offset; + } + } + + if (debug) { + fprintf(stderr, "split_pool_spill: device=%d tensor=%s request=%zu aligned=%zu pool_used=%zu pool_size=%zu chunks=%zu\n", + device, tensor ? tensor->name : "(null)", size, aligned_size, + pool_used(device), pool_size(device), pools[device].size()); + } + + size_t alloc_size = exact_fallback ? aligned_size : std::max(aligned_size, preferred_chunk_size); + ggml_sycl_set_device(device); + char * data = nullptr; + SYCL_CHECK(CHECK_TRY_ERROR(data = (char *) sycl::malloc_device(alloc_size, *(streams[device])))); + if (data == nullptr && alloc_size != aligned_size) { + alloc_size = aligned_size; + SYCL_CHECK(CHECK_TRY_ERROR(data = (char *) sycl::malloc_device(alloc_size, *(streams[device])))); + } + if (data == nullptr) { + const size_t max_alloc = dpct::get_current_device().get_max_mem_alloc_size(); + size_t free_mem = 0; + size_t total_mem = 0; + dpct::dev_mgr::instance().get_device(device).get_memory_info(free_mem, total_mem); + log_pool_state(__func__, device); + GGML_ABORT("%s: sycl::malloc_device returned null: device=%d tensor=%s size=%zu chunk_size=%zu max_alloc=%zu free=%zu total=%zu\n", + __func__, device, tensor ? tensor->name : "(null)", aligned_size, alloc_size, max_alloc, free_mem, total_mem); + } + + pools[device].push_back({ data, alloc_size, aligned_size }); + if (debug) { + fprintf(stderr, "split_pool_spill_allocated: device=%d tensor=%s request=%zu chunk_size=%zu pool_used=%zu pool_size=%zu chunks=%zu\n", + device, tensor ? tensor->name : "(null)", size, alloc_size, + pool_used(device), pool_size(device), pools[device].size()); + } + return data; + } + + void reserve(size_t total_size) { + if (total_size == 0 || pools.empty()) { + return; + } + + constexpr size_t preferred_chunk_size = size_t(1024) * 1024 * 1024; + const size_t per_device_size = align_up(total_size / pools.size() + preferred_chunk_size, preferred_chunk_size); + + for (int device = 0; device < int(pools.size()); ++device) { + ggml_sycl_set_device(device); + char * data = nullptr; + SYCL_CHECK(CHECK_TRY_ERROR(data = (char *) sycl::malloc_device(per_device_size, *(streams[device])))); + if (data == nullptr) { + const size_t max_alloc = dpct::get_current_device().get_max_mem_alloc_size(); + size_t free_mem = 0; + size_t total_mem = 0; + dpct::dev_mgr::instance().get_device(device).get_memory_info(free_mem, total_mem); + GGML_ABORT("%s: sycl::malloc_device returned null: device=%d pool_size=%zu total_request=%zu max_alloc=%zu free=%zu total=%zu\n", + __func__, device, per_device_size, total_size, max_alloc, free_mem, total_mem); + } + pools[device].push_back({ data, per_device_size, 0 }); + } + } + + void reserve_per_device(size_t requested_size, size_t requested_chunk_size = 0) { + if (requested_size == 0 || pools.empty()) { + return; + } + + constexpr size_t preferred_chunk_size = size_t(1024) * 1024 * 1024; + const size_t pool_size = align_up(requested_size, preferred_chunk_size); + const size_t chunk_size = requested_chunk_size == 0 + ? pool_size + : align_up(requested_chunk_size, preferred_chunk_size); + + for (int device = 0; device < int(pools.size()); ++device) { + size_t remaining = pool_size; + while (remaining > 0) { + const size_t this_chunk = std::min(chunk_size, remaining); + ggml_sycl_set_device(device); + char * data = nullptr; + SYCL_CHECK(CHECK_TRY_ERROR(data = (char *) sycl::malloc_device(this_chunk, *(streams[device])))); + if (data == nullptr) { + const size_t max_alloc = dpct::get_current_device().get_max_mem_alloc_size(); + size_t free_mem = 0; + size_t total_mem = 0; + dpct::dev_mgr::instance().get_device(device).get_memory_info(free_mem, total_mem); + GGML_ABORT("%s: sycl::malloc_device returned null: device=%d chunk_size=%zu pool_size=%zu max_alloc=%zu free=%zu total=%zu\n", + __func__, device, this_chunk, pool_size, max_alloc, free_mem, total_mem); + } + pools[device].push_back({ data, this_chunk, 0 }); + remaining -= this_chunk; + } + } + } + std::vector tensor_extras; std::vector streams; + std::vector> pools; }; GGML_CALL static const char * ggml_backend_sycl_split_buffer_get_name(ggml_backend_buffer_t buffer) { @@ -4683,7 +5425,6 @@ ggml_backend_sycl_split_buffer_init_tensor(ggml_backend_buffer_t buffer, ggml_tensor_extra_gpu * extra = new ggml_tensor_extra_gpu{}; ctx->tensor_extras.push_back(extra); - ctx->streams.push_back(&(dpct::get_current_device().default_queue())); for (int i = 0; i < ggml_sycl_info().device_count; ++i) { int64_t row_low, row_high; @@ -4706,14 +5447,7 @@ ggml_backend_sycl_split_buffer_init_tensor(ggml_backend_buffer_t buffer, // currently, init_tensor cannot fail, it needs to be fixed in ggml-backend first ggml_sycl_set_device(i); const queue_ptr stream = ctx->streams[i]; - char * buf; - /* - DPCT1009:208: SYCL uses exceptions to report errors and does not use the - error codes. The original code was commented out and a warning string - was inserted. You need to rewrite this code. - */ - SYCL_CHECK(CHECK_TRY_ERROR(buf = (char *)sycl::malloc_device( - size, *stream))); + char * buf = ctx->alloc(i, size, tensor); // set padding to 0 to avoid possible NaN values if (size > original_size) { @@ -4740,7 +5474,7 @@ ggml_backend_sycl_split_buffer_init_tensor(ggml_backend_buffer_t buffer, CHECK_TRY_ERROR(extra->events[i][is] = new sycl::event())); } } - tensor->backend = GGML_BACKEND_TYPE_GPU_SPLIT; + tensor->__backend = GGML_BACKEND_TYPE_GPU_SPLIT; tensor->extra = extra; } catch (sycl::exception const &exc) { @@ -4790,6 +5524,11 @@ ggml_backend_sycl_split_buffer_set_tensor(ggml_backend_buffer_t buffer, */ ggml_sycl_set_device(i); const queue_ptr stream = ctx->streams[i]; + if (get_sycl_env("GGML_SYCL_SPLIT_COPY_DEBUG", 0) != 0) { + fprintf(stderr, + "split_copy_set_tensor: device=%d tensor=%s rows=[%" PRId64 ",%" PRId64 ") bytes=%zu padded_bytes=%zu offset=%zu\n", + i, tensor->name, row_low, row_high, original_size, size, offset_split); + } SYCL_CHECK(CHECK_TRY_ERROR( (*stream) .memcpy(extra->data_device[i], buf_host, original_size) @@ -4880,11 +5619,27 @@ GGML_CALL static const char * ggml_backend_sycl_split_buffer_type_name(ggml_back } GGML_CALL static ggml_backend_buffer_t ggml_backend_sycl_split_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft, size_t size) { - // since we don't know the exact split after rounding, we cannot allocate the device buffers at this point - // instead, we allocate them for each tensor separately in init_tensor - // however, the size still represents the maximum cumulative size of all the device buffers after the tensors are allocated, - // as returned by get_alloc_size. this limit is enforced during tensor allocation by ggml-alloc, so it must be correct. + // B70/Level Zero can fail after many USM allocations even when most VRAM is + // free. Pre-reserving coarse pools helps some models, but MiniMax graph + // split can overestimate the split buffer enough to exceed per-card VRAM. ggml_backend_sycl_split_buffer_context * ctx = new ggml_backend_sycl_split_buffer_context(); + if (const char * env = getenv("GGML_SYCL_SPLIT_POOL_MIB")) { + const int mib = std::atoi(env); + if (mib > 0) { + size_t chunk_size = 0; + if (const char * chunk_env = getenv("GGML_SYCL_SPLIT_POOL_CHUNK_MIB")) { + const int chunk_mib = std::atoi(chunk_env); + if (chunk_mib > 0) { + chunk_size = size_t(chunk_mib) * 1024 * 1024; + } + } + ctx->reserve_per_device(size_t(mib) * 1024 * 1024, chunk_size); + } + } else if (const char * env = getenv("GGML_SYCL_SPLIT_PRE_RESERVE")) { + if (std::atoi(env) != 0) { + ctx->reserve(size); + } + } return ggml_backend_buffer_init(buft, ggml_backend_sycl_split_buffer_interface, ctx, size); } @@ -5170,6 +5925,7 @@ GGML_CALL static bool ggml_backend_sycl_supports_op(ggml_backend_t backend, cons case GGML_UNARY_OP_HARDSWISH: case GGML_UNARY_OP_GELU_QUICK: case GGML_UNARY_OP_TANH: + case GGML_UNARY_OP_SIGMOID: return ggml_is_contiguous(op->src[0]); default: return false; @@ -5206,6 +5962,36 @@ GGML_CALL static bool ggml_backend_sycl_supports_op(ggml_backend_t backend, cons } return true; } break; + case GGML_OP_MOE_FUSED_UP_GATE: + { + struct ggml_tensor * up_exps = op->src[0]; + struct ggml_tensor * gate_exps = op->src[1]; + struct ggml_tensor * inp = op->src[2]; + struct ggml_tensor * ids = op->src[3]; + if (!up_exps || !inp || !ids || op->type != GGML_TYPE_F32 || inp->type != GGML_TYPE_F32 || ids->type != GGML_TYPE_I32) { + return false; + } + if (op->src[4] || op->src[5]) { + return false; + } + if (up_exps->ne[3] != inp->ne[3]) { + return false; + } + if (gate_exps && (up_exps->type != gate_exps->type || !ggml_are_same_shape(up_exps, gate_exps))) { + return false; + } + if (!ggml_is_contiguous(op)) { + return false; + } + const int unary_op = op->op_params[0]; + if (unary_op != GGML_UNARY_OP_SILU && unary_op != GGML_UNARY_OP_GELU && unary_op != GGML_UNARY_OP_RELU) { + return false; + } + if (up_exps->type == GGML_TYPE_BF16) { + return false; + } + return true; + } break; case GGML_OP_GET_ROWS: { switch (op->src[0]->type) { @@ -5278,6 +6064,8 @@ GGML_CALL static bool ggml_backend_sycl_supports_op(ggml_backend_t backend, cons case GGML_OP_IM2COL: case GGML_OP_POOL_2D: case GGML_OP_SUM_ROWS: + case GGML_OP_MULTI_ADD: + case GGML_OP_MUL_MULTI_ADD: case GGML_OP_ARGSORT: case GGML_OP_ACC: case GGML_OP_GROUP_NORM: diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index b7db11d..5062830 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -6524,6 +6524,12 @@ static struct ggml_tensor * ggml_fused_mul_unary_impl( enum ggml_unary_op op, bool inplace) { + const char * disable_fused_mul_unary = getenv("GGML_DISABLE_FUSED_MUL_UNARY"); + if (disable_fused_mul_unary && strcmp(disable_fused_mul_unary, "0") != 0) { + struct ggml_tensor * result = ggml_unary(ctx, a, op); + return ggml_mul_impl(ctx, result, b, inplace); + } + GGML_ASSERT(ggml_is_contiguous(a)); if (!ggml_are_same_shape(b, a)) { GGML_ASSERT(a->ne[0] == 1 && a->ne[1] == b->ne[1] && a->ne[2] == b->ne[2] && a->ne[3] == b->ne[3]); @@ -7518,6 +7524,13 @@ static struct ggml_tensor * ggml_fused_rms_norm_impl( return ggml_rms_norm_impl(ctx, a, eps, inplace); } + const char * disable_fused_rms = getenv("GGML_DISABLE_FUSED_RMS_NORM"); + if (disable_fused_rms && strcmp(disable_fused_rms, "0") != 0) { + struct ggml_tensor * result = ggml_rms_norm_impl(ctx, a, eps, inplace); + result = ggml_mul_impl(ctx, result, b, inplace); + return result; + } + if (ggml_nrows(b) > 1 || a->ne[0] != b->ne[0]) { struct ggml_tensor * result = ggml_rms_norm_impl(ctx, a, eps, inplace); result = ggml_mul_impl(ctx, result, b, inplace); diff --git a/ggml/src/iqk/iqk_flash_attn.cpp b/ggml/src/iqk/iqk_flash_attn.cpp index c0ece49..2325e30 100644 --- a/ggml/src/iqk/iqk_flash_attn.cpp +++ b/ggml/src/iqk/iqk_flash_attn.cpp @@ -493,6 +493,10 @@ extern "C" IQK_API bool iqk_flash_attn_noalibi(int type_q, int type_mask, float #else +size_t iqk_fa_work_buffer_size([[maybe_unused]] const struct ggml_tensor * dst, [[maybe_unused]] int nth) { + return 0; +} + bool iqk_flash_attn_noalibi([[maybe_unused]] int type_q, [[maybe_unused]] int type_mask, [[maybe_unused]] float max_bias, [[maybe_unused]] int neq3, [[maybe_unused]] int neq2, [[maybe_unused]] long nbq3, [[maybe_unused]] long nbq2, [[maybe_unused]] int nek3, [[maybe_unused]] int nek2, [[maybe_unused]] long nbk3, [[maybe_unused]] long nbk2, @@ -512,6 +516,7 @@ bool iqk_flash_attn_noalibi([[maybe_unused]] int type_q, [[maybe_unused]] int ty [[maybe_unused]] const void * k, // k matrix. Assumed to be fp16, nq x nk elements [[maybe_unused]] const void * v, // v matrix. Assumed to be fp16, nq x nk elements [[maybe_unused]] const void * mask, // mask. If not null, assumed to be fp16. nq x nk elements + [[maybe_unused]] const void * sinks, // sinks. If not null, assumed to be fp32. [[maybe_unused]] float scale, // scale applied before softmax [[maybe_unused]] float softcap, // if > 0, a "soft-cap" operation is applied before softmax [[maybe_unused]] float * qkv, // v*softmax(scale*(k*q)) @@ -521,4 +526,3 @@ bool iqk_flash_attn_noalibi([[maybe_unused]] int type_q, [[maybe_unused]] int ty } #endif -