Commit af911149c for llama.cpp
commit af911149c58a4a2c17470d7ccb2e91140bc73a9f
Author: Łukasz Ślusarczyk <lukasz.slusarczyk@intel.com>
Date: Mon Sep 21 12:58:59 2026 +0200
sycl : pinned memory use right device context instead of 0 (#28895)
diff --git a/ggml/include/ggml-sycl.h b/ggml/include/ggml-sycl.h
index 093fa4a7e..1c18f706a 100644
--- a/ggml/include/ggml-sycl.h
+++ b/ggml/include/ggml-sycl.h
@@ -36,6 +36,8 @@ GGML_BACKEND_API void ggml_backend_sycl_comm_free(void * comm_ctx);
GGML_BACKEND_API bool ggml_backend_sycl_comm_allreduce_tensor(void * comm_ctx, struct ggml_tensor ** tensors);
// pinned host buffer for use with the CPU backend for faster copies between CPU and GPU
+// pins on device 0 - a copy between another device and this memory can fail,
+// use ggml_backend_dev_host_buffer_type to pin on the device that does the copy
GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_sycl_host_buffer_type(void);
GGML_BACKEND_API void ggml_backend_sycl_print_sycl_devices(void);
diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index b24664a0b..a7fbd1644 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -1553,14 +1553,18 @@ static const char * ggml_backend_sycl_host_buffer_type_name(ggml_backend_buffer_
GGML_UNUSED(buft);
}
+static int ggml_backend_sycl_host_buffer_type_device(ggml_backend_buffer_type_t buft) {
+ return static_cast<const ggml_backend_sycl_device_context *>(buft->device->context)->device;
+}
+
//host pinned memory
-static void * ggml_backend_sycl_host_malloc(size_t size) {
- GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_host_malloc\n");
+static void * ggml_backend_sycl_host_malloc(int device, size_t size) {
+ GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_host_malloc of size %.2f MiB on device %d\n", size / 1024.0 / 1024.0, device);
void * ptr = nullptr;
try {
ggml_check_sycl();
// USM host memory is page-locked and device-accessible by construction
- auto & q = dpct::dev_mgr::instance().get_device(0).default_queue();
+ auto & q = dpct::dev_mgr::instance().get_device(device).default_queue();
ptr = sycl::malloc_host(size, q, sycl::property_list{});
} catch (...) {
ptr = nullptr;
@@ -1578,7 +1582,8 @@ static void ggml_backend_sycl_host_buffer_free_buffer(ggml_backend_buffer_t buff
return;
}
if (g_ggml_sycl_enable_host_pinned_mem) {
- auto & q = dpct::dev_mgr::instance().get_device(0).default_queue();
+ const int device = ggml_backend_sycl_host_buffer_type_device(buffer->buft);
+ auto & q = dpct::dev_mgr::instance().get_device(device).default_queue();
SYCL_CHECK(CHECK_TRY_ERROR(sycl::free(buffer->context, q)));
} else {
free_aligned_mem_host((void *) buffer->context);
@@ -1586,8 +1591,9 @@ static void ggml_backend_sycl_host_buffer_free_buffer(ggml_backend_buffer_t buff
}
static ggml_backend_buffer_t ggml_backend_sycl_host_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft, size_t size) {
- void * ptr = g_ggml_sycl_enable_host_pinned_mem ? ggml_backend_sycl_host_malloc(size) :
- aligned_malloc_host(TENSOR_ALIGNMENT, size);
+ void * ptr = g_ggml_sycl_enable_host_pinned_mem ?
+ ggml_backend_sycl_host_malloc(ggml_backend_sycl_host_buffer_type_device(buft), size) :
+ aligned_malloc_host(TENSOR_ALIGNMENT, size);
if (ptr == nullptr) {
// fallback to cpu buffer
return ggml_backend_buft_alloc_buffer(ggml_backend_cpu_buffer_type(), size);
@@ -1604,8 +1610,8 @@ static ggml_backend_buffer_t ggml_backend_sycl_host_buffer_type_alloc_buffer(ggm
static size_t ggml_backend_sycl_host_buffer_type_get_max_size(ggml_backend_buffer_type_t buft) {
if (g_ggml_sycl_enable_host_pinned_mem) {
- ggml_backend_sycl_device_context * dev_ctx = (ggml_backend_sycl_device_context *) buft->device->context;
- size_t max_alloc_size = dpct::dev_mgr::instance().get_device(dev_ctx->device).get_max_mem_alloc_size();
+ const int device = ggml_backend_sycl_host_buffer_type_device(buft);
+ size_t max_alloc_size = dpct::dev_mgr::instance().get_device(device).get_max_mem_alloc_size();
if (g_ggml_sycl_host_pinned_mem_2g) {
return std::min(max_alloc_size, (size_t) 2LL*1024*1024*1024);
} else {
@@ -1616,22 +1622,37 @@ static size_t ggml_backend_sycl_host_buffer_type_get_max_size(ggml_backend_buffe
}
}
-ggml_backend_buffer_type_t ggml_backend_sycl_host_buffer_type() {
- GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_host_buffer_type\n");
- static struct ggml_backend_buffer_type ggml_backend_sycl_buffer_type_host = {
- /* .iface = */ {
- /* .get_name = */ ggml_backend_sycl_host_buffer_type_name,
- /* .alloc_buffer = */ ggml_backend_sycl_host_buffer_type_alloc_buffer,
- /* .get_alignment = */ ggml_backend_cpu_buffer_type()->iface.get_alignment,
- /* .get_max_size = */ ggml_backend_sycl_host_buffer_type_get_max_size,
- /* .get_alloc_size = */ ggml_backend_cpu_buffer_type()->iface.get_alloc_size,
- /* .is_host = */ ggml_backend_cpu_buffer_type()->iface.is_host,
- },
- /* .device = */ ggml_backend_reg_dev_get(ggml_backend_sycl_reg(), 0),
- /* .context = */ nullptr,
- };
+static ggml_backend_buffer_type_t ggml_backend_sycl_host_buffer_type_for_device(int device) {
+ GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_host_buffer_type_for_device on device %d\n", device);
+
+ // the vector is never resized after this, so the returned pointers stay valid
+ static std::vector<ggml_backend_buffer_type> buffer_types_host = [] {
+ std::vector<ggml_backend_buffer_type> bufts(ggml_backend_sycl_get_device_count());
+ for (size_t i = 0; i < bufts.size(); i++) {
+ bufts[i] = {
+ /* .iface = */ {
+ /* .get_name = */ ggml_backend_sycl_host_buffer_type_name,
+ /* .alloc_buffer = */ ggml_backend_sycl_host_buffer_type_alloc_buffer,
+ /* .get_alignment = */ ggml_backend_cpu_buffer_type()->iface.get_alignment,
+ /* .get_max_size = */ ggml_backend_sycl_host_buffer_type_get_max_size,
+ /* .get_alloc_size = */ ggml_backend_cpu_buffer_type()->iface.get_alloc_size,
+ /* .is_host = */ ggml_backend_cpu_buffer_type()->iface.is_host,
+ },
+ /* .device = */ ggml_backend_reg_dev_get(ggml_backend_sycl_reg(), i),
+ /* .context = */ nullptr,
+ };
+ }
+ return bufts;
+ }();
+
+ GGML_ASSERT(device >= 0 && device < (int) buffer_types_host.size());
- return &ggml_backend_sycl_buffer_type_host;
+ return &buffer_types_host[device];
+}
+
+// TODO: this function is unused and is a temporary hack to avoid breaking changes
+ggml_backend_buffer_type_t ggml_backend_sycl_host_buffer_type() {
+ return ggml_backend_sycl_host_buffer_type_for_device(0);
}
// buffer pool for sycl (legacy)
@@ -6299,8 +6320,8 @@ static ggml_backend_buffer_type_t ggml_backend_sycl_device_get_buffer_type(ggml_
}
static ggml_backend_buffer_type_t ggml_backend_sycl_device_get_host_buffer_type(ggml_backend_dev_t dev) {
- GGML_UNUSED(dev);
- return ggml_backend_sycl_host_buffer_type();
+ ggml_backend_sycl_device_context * ctx = (ggml_backend_sycl_device_context *) dev->context;
+ return ggml_backend_sycl_host_buffer_type_for_device(ctx->device);
}
static ggml_backend_buffer_t ggml_backend_sycl_device_buffer_from_host_ptr(ggml_backend_dev_t dev, void * ptr, size_t size, size_t max_tensor_size) {