From 125099f709351b42edf75877bdce839cd3d26e30 Mon Sep 17 00:00:00 2001 From: Akarshan Biswas Date: Wed, 26 Mar 2025 09:48:27 +0530 Subject: [PATCH 1/3] SYCL: implement memset ggml backend buffer interface --- ggml/src/ggml-sycl/ggml-sycl.cpp | 22 +++++++++++++++++++++- 1 file changed, 21 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 9fa24b98098f..26e1f445d427 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -37,6 +37,7 @@ #include "ggml-backend-impl.h" #include "ggml-sycl/backend.hpp" +#include "ggml-sycl/common.hpp" #include "ggml-sycl/presets.hpp" #include "ggml-sycl/gemm.hpp" #include "ggml-sycl/sycl_hw.hpp" @@ -490,6 +491,25 @@ catch (sycl::exception const &exc) { std::exit(1); } +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()); + SYCL_CHECK(CHECK_TRY_ERROR(dpct::dev_mgr::instance().get_device(ctx->device).queues_wait_and_throw())); + if (size == 0) { + return; // Nothing to do + } + if (tensor->data == nullptr) { + std::cerr << "Error: Tensor data pointer is null." << std::endl; + std::exit(1); + } + 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_reset(ggml_backend_buffer_t buffer) { GGML_SYCL_DEBUG("[SYCL] call %s\n", __func__); if (buffer == nullptr) { @@ -510,7 +530,7 @@ 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 = */ NULL, + /* .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, From 6d5b114c78d3e644fead5cfb79c7e9cde2d25efb Mon Sep 17 00:00:00 2001 From: Akarshan Biswas Date: Wed, 26 Mar 2025 10:16:29 +0530 Subject: [PATCH 2/3] use GGML_ABORT macro --- ggml/src/ggml-sycl/ggml-sycl.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 26e1f445d427..e66170fc5627 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -502,8 +502,7 @@ static void ggml_backend_sycl_buffer_memset_tensor(ggml_backend_buffer_t buffer, return; // Nothing to do } if (tensor->data == nullptr) { - std::cerr << "Error: Tensor data pointer is null." << std::endl; - std::exit(1); + 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))); From d5f513db3c9e8ee4b85909820d8d97d796e9bac1 Mon Sep 17 00:00:00 2001 From: Akarshan Biswas Date: Wed, 26 Mar 2025 18:01:14 +0530 Subject: [PATCH 3/3] Do not wait for all queues to finish for memset operation --- ggml/src/ggml-sycl/ggml-sycl.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index e66170fc5627..39d53da33035 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -497,7 +497,6 @@ static void ggml_backend_sycl_buffer_memset_tensor(ggml_backend_buffer_t buffer, 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()); - SYCL_CHECK(CHECK_TRY_ERROR(dpct::dev_mgr::instance().get_device(ctx->device).queues_wait_and_throw())); if (size == 0) { return; // Nothing to do }