Superminaren/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-09-27 21:46:57 +02:00
Code Issues Packages Projects Releases Wiki Activity

sycl : pinned memory use right device context instead of 0 (#28895)

Browse Source
This commit is contained in:
Łukasz Ślusarczyk
2026-09-21 13:58:59 +03:00
committed by GitHub
parent 1884824fda
commit af911149c5
2 changed files with 48 additions and 25 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/include/ggml-sycl.h
+2
View File
@@ -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);
ggml/src/ggml-sycl/ggml-sycl.cpp
+46 -25
View File
@@ -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);
return &ggml_backend_sycl_buffer_type_host;
// 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 &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) {