Public/koboldcpp
1
0
Fork 0
mirror of https://github.com/LostRuins/koboldcpp.git synced 2026-09-14 02:38:50 +02:00
Code Issues Packages Projects Releases Wiki Activity

opencl: fix several bugs where the backend aborts (#27630)

Browse Source
This commit is contained in:
Hongqiang Wang
2026-09-11 22:11:12 -07:00
committed by GitHub
parent 07fc97716f
commit 8a56aedd61
1 changed files with 226 additions and 130 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-opencl/ggml-opencl.cpp
+226 -130
View File
@@ -203,39 +203,67 @@ static ggml_cl_version get_opencl_platform_version(cl_platform_id platform) {
return parse_cl_version(param_value);
}
// Returns the DEVICE's OpenCL version. On an error returns ggml_cl_version with all zeroes.
static ggml_cl_version get_opencl_device_version(cl_device_id device) {
size_t param_size;
if (clGetDeviceInfo(device, CL_DEVICE_VERSION, 0, nullptr, &param_size) != CL_SUCCESS || !param_size) {
return {};
}
std::unique_ptr<char[]> param_storage(new char[param_size]);
if (clGetDeviceInfo(device, CL_DEVICE_VERSION, param_size, param_storage.get(), nullptr) != CL_SUCCESS) {
return {};
}
auto param_value = std::string_view(param_storage.get(), param_size);
const std::string version_prefix = "OpenCL "; // "OpenCL <major>.<minor> <device-specific-info>"
if (param_value.find(version_prefix) != 0) {
return {};
}
param_value.remove_prefix(version_prefix.length());
return parse_cl_version(param_value);
}
// Return a version to use in OpenCL C compilation. On an error returns ggml_cl_version with all zeroes.
static ggml_cl_version get_opencl_c_version(ggml_cl_version platform_version, cl_device_id device) {
size_t param_size;
#if CL_TARGET_OPENCL_VERSION >= 300
if (platform_version.major >= 3) {
CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, 0, nullptr, &param_size));
if (!param_size) {
return {};
// CL_DEVICE_OPENCL_C_ALL_VERSIONS is an OpenCL 3.0 *device* query, so gating it on the
// *platform* version is not enough: a 3.0 platform can expose 2.0 devices, where the
// query returns CL_INVALID_VALUE and the old CL_CHECK aborted during backend init.
// Gate on the device version, and treat a failure as "fall back to the legacy query"
// rather than fatal -- a device may advertise 3.0 and still refuse the property.
const ggml_cl_version device_version = get_opencl_device_version(device);
if (platform_version.major >= 3 && device_version.major >= 3) {
cl_int err = clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, 0, nullptr, &param_size);
if (err == CL_SUCCESS && param_size) {
std::unique_ptr<cl_name_version[]> versions(new cl_name_version[param_size]);
err = clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, param_size, versions.get(), nullptr);
if (err == CL_SUCCESS) {
unsigned versions_count = param_size / sizeof(cl_name_version);
cl_version version_max = 0;
for (unsigned i = 0; i < versions_count; i++) {
version_max = std::max<cl_version>(versions[i].version, version_max);
}
return { CL_VERSION_MAJOR(version_max), CL_VERSION_MINOR(version_max) };
}
}
std::unique_ptr<cl_name_version[]> versions(new cl_name_version[param_size]);
CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, param_size, versions.get(), nullptr));
unsigned versions_count = param_size / sizeof(cl_name_version);
cl_version version_max = 0;
for (unsigned i = 0; i < versions_count; i++) {
version_max = std::max<cl_version>(versions[i].version, version_max);
}
return { CL_VERSION_MAJOR(version_max), CL_VERSION_MINOR(version_max) };
// fall through to CL_DEVICE_OPENCL_C_VERSION below
}
#else
GGML_UNUSED(platform_version);
#endif // CL_TARGET_OPENCL_VERSION >= 300
CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_VERSION, 0, nullptr, &param_size));
if (!param_size) {
if (clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_VERSION, 0, nullptr, &param_size) != CL_SUCCESS || !param_size) {
return {};
}
std::unique_ptr<char[]> param_storage(new char[param_size]);
CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_VERSION, param_size, param_storage.get(), nullptr));
if (clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_VERSION, param_size, param_storage.get(), nullptr) != CL_SUCCESS) {
return {};
}
auto param_value = std::string_view(param_storage.get(), param_size);
const std::string version_prefix = "OpenCL C "; // Suffix: "XX.YY <platform-specific-info>"
@@ -1115,6 +1143,18 @@ struct ggml_backend_opencl_context {
}
void enqueue_ndrange_kernel(cl_kernel kernel, cl_uint work_dim, size_t *global_work_size, size_t *local_work_size, const ggml_tensor * tensor) {
// From the spec on clEnqueueNDRangeKernel:
// If the device associated with command_queue is an OpenCL 2.1 or newer device,
// and global_work_size is NULL or the value in any passed dimension is zero,
// then the kernel command will trivially succeed after its event dependencies
// are satisfied and subsequently update its completion event.
// So this ensures such cases always return trivially without causing errors in
// case of an older device.
for (cl_uint i = 0; i < work_dim; i++) {
if (global_work_size[i] == 0) {
return;
}
}
#ifdef GGML_OPENCL_PROFILING
cl_event evt;
CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, &evt));
@@ -9459,6 +9499,96 @@ static enum ggml_status ggml_backend_opencl_buffer_init_tensor(ggml_backend_buff
return GGML_STATUS_SUCCESS;
}
// Allocate a temporary upload buffer of `nbytes` and populate it with `data`
// from host. On Adreno X1-85 the device-only pool intermittently fails to
// allocate at hundreds of MB once model weights fragment the heap (observed
// on Qwen3.5-9B output.weight Q6_K at 834 MB). Three-step retry:
// 1. CL_MEM_READ_WRITE alloc + clEnqueueWriteBuffer (normal fast path).
// 2. clFinish + retry (drains in-flight allocs that may be holding heap;
// mirrors the proven pattern at the FD-split partial buffer alloc).
// 3. CL_MEM_ALLOC_HOST_PTR + map(WRITE_INVALIDATE) + memcpy + unmap —
// different memory pool (host-pinned); true zero-copy on Adreno per
// QCOM guidance. (CL_MEM_USE_HOST_PTR is NOT zero-copy on Adreno: the
// driver triggers an internal copy because arbitrary host pages aren't
// guaranteed mappable/coherent, AND it draws from the same exhausted
// device pool — so it doesn't solve the problem.)
// Returns the ready-to-read buffer (caller must clReleaseMemObject) or NULL
// if all three strategies fail. The buffer is opaque to the caller — it can
// be passed as a kernel argument like any normal cl_mem.
static cl_mem ggml_cl_create_temp_upload_buffer(
cl_context context, cl_command_queue queue,
size_t nbytes, const void * data,
const char * tensor_name_for_log)
{
cl_int err;
cl_mem buf = clCreateBuffer(context, CL_MEM_READ_WRITE, nbytes, NULL, &err);
if (err != CL_SUCCESS) {
clFinish(queue);
buf = clCreateBuffer(context, CL_MEM_READ_WRITE, nbytes, NULL, &err);
}
if (err == CL_SUCCESS) {
const cl_int werr = clEnqueueWriteBuffer(queue, buf, CL_TRUE, 0, nbytes, data, 0, NULL, NULL);
if (werr == CL_SUCCESS) {
return buf;
}
clReleaseMemObject(buf);
}
buf = clCreateBuffer(context,
CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR | CL_MEM_HOST_WRITE_ONLY,
nbytes, NULL, &err);
if (err != CL_SUCCESS) {
return NULL;
}
void * mapped = clEnqueueMapBuffer(queue, buf, CL_TRUE,
CL_MAP_WRITE_INVALIDATE_REGION, 0, nbytes, 0, NULL, NULL, &err);
if (err != CL_SUCCESS) {
clReleaseMemObject(buf);
return NULL;
}
memcpy(mapped, data, nbytes);
const cl_int uerr = clEnqueueUnmapMemObject(queue, buf, mapped, 0, NULL, NULL);
if (uerr != CL_SUCCESS) {
clReleaseMemObject(buf);
return NULL;
}
if (tensor_name_for_log) {
GGML_LOG_INFO("ggml_opencl: %s (%.1f MiB) — device alloc failed, using CL_MEM_ALLOC_HOST_PTR fallback\n",
tensor_name_for_log, nbytes / 1024.0 / 1024.0);
}
return buf;
}
// Allocate a temporary download buffer of `nbytes`. The caller runs a kernel
// that writes into it, then reads it back to host via clEnqueueReadBuffer (or
// equivalent). Mirrors ggml_cl_create_temp_upload_buffer; the host-pinned
// fallback flags are flipped (CL_MEM_WRITE_ONLY | HOST_READ_ONLY) and the
// helper doesn't populate the buffer.
static cl_mem ggml_cl_create_temp_download_buffer(
cl_context context, cl_command_queue queue,
size_t nbytes, const char * tensor_name_for_log)
{
cl_int err;
cl_mem buf = clCreateBuffer(context, CL_MEM_READ_WRITE, nbytes, NULL, &err);
if (err != CL_SUCCESS) {
clFinish(queue);
buf = clCreateBuffer(context, CL_MEM_READ_WRITE, nbytes, NULL, &err);
}
if (err == CL_SUCCESS) {
return buf;
}
buf = clCreateBuffer(context,
CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR | CL_MEM_HOST_READ_ONLY,
nbytes, NULL, &err);
if (err != CL_SUCCESS) {
return NULL;
}
if (tensor_name_for_log) {
GGML_LOG_INFO("ggml_opencl: %s download (%.1f MiB) — device alloc failed, using CL_MEM_ALLOC_HOST_PTR fallback\n",
tensor_name_for_log, nbytes / 1024.0 / 1024.0);
}
return buf;
}
static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor, const void * data, size_t offset, size_t size) {
ggml_backend_opencl_device_context * dev_ctx = (ggml_backend_opencl_device_context *) buffer->buft->device->context;
ggml_backend_opencl_context * backend_ctx = dev_ctx->backend_ctx;
@@ -9567,12 +9697,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
GGML_ASSERT(size_d + size_q == ggml_nbytes(tensor) && "Incorrect tensor size");
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
CL_CHECK(clEnqueueWriteBuffer(
queue, data_device, CL_TRUE, 0,
ggml_nbytes(tensor), data, 0, NULL, NULL));
cl_mem data_device = ggml_cl_create_temp_upload_buffer(context, queue, ggml_nbytes(tensor), data, tensor->name);
GGML_ASSERT(data_device != NULL && "set_tensor: temp upload buffer alloc failed");
// We consider the specified offset arg as always, although For weights
// the offset arg should be 0 (we do not assert this).
@@ -9730,12 +9856,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
GGML_ASSERT(size_d + size_m + size_q == ggml_nbytes(tensor) && "Incorrect tensor size");
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
CL_CHECK(clEnqueueWriteBuffer(
queue, data_device, CL_TRUE, 0,
ggml_nbytes(tensor), data, 0, NULL, NULL));
cl_mem data_device = ggml_cl_create_temp_upload_buffer(context, queue, ggml_nbytes(tensor), data, tensor->name);
GGML_ASSERT(data_device != NULL && "set_tensor: temp upload buffer alloc failed");
cl_buffer_region region;
@@ -9862,12 +9984,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
GGML_ASSERT(size_d + size_qs + size_qh == ggml_nbytes(tensor) && "Incorrect tensor size");
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
CL_CHECK(clEnqueueWriteBuffer(
queue, data_device, CL_TRUE, 0,
ggml_nbytes(tensor), data, 0, NULL, NULL));
cl_mem data_device = ggml_cl_create_temp_upload_buffer(context, queue, ggml_nbytes(tensor), data, tensor->name);
GGML_ASSERT(data_device != NULL && "set_tensor: temp upload buffer alloc failed");
cl_buffer_region region;
@@ -10026,12 +10144,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
GGML_ASSERT(size_d + size_m + size_qs + size_qh == ggml_nbytes(tensor) && "Incorrect tensor size");
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
CL_CHECK(clEnqueueWriteBuffer(
queue, data_device, CL_TRUE, 0,
ggml_nbytes(tensor), data, 0, NULL, NULL));
cl_mem data_device = ggml_cl_create_temp_upload_buffer(context, queue, ggml_nbytes(tensor), data, tensor->name);
GGML_ASSERT(data_device != NULL && "set_tensor: temp upload buffer alloc failed");
cl_buffer_region region;
@@ -10179,12 +10293,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
GGML_ASSERT(size_e + size_q == ggml_nbytes(tensor) && "Incorrect tensor size");
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
CL_CHECK(clEnqueueWriteBuffer(
queue, data_device, CL_TRUE, 0,
ggml_nbytes(tensor), data, 0, NULL, NULL));
cl_mem data_device = ggml_cl_create_temp_upload_buffer(context, queue, ggml_nbytes(tensor), data, tensor->name);
GGML_ASSERT(data_device != NULL && "set_tensor: temp upload buffer alloc failed");
// The original tensor memory is divided into scales and quants, i.e.,
// we first store scales, then quants.
@@ -10290,12 +10400,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
GGML_ASSERT(size_d + size_q == ggml_nbytes(tensor) && "Incorrect tensor size");
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
CL_CHECK(clEnqueueWriteBuffer(
queue, data_device, CL_TRUE, 0,
ggml_nbytes(tensor), data, 0, NULL, NULL));
cl_mem data_device = ggml_cl_create_temp_upload_buffer(context, queue, ggml_nbytes(tensor), data, tensor->name);
GGML_ASSERT(data_device != NULL && "set_tensor: temp upload buffer alloc failed");
// The original tensor memory is divided into scales and quants, i.e.,
// we first store scales, then quants.
@@ -10394,12 +10500,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
GGML_ASSERT(size_d + size_q == ggml_nbytes(tensor) && "Incorrect tensor size");
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
CL_CHECK(clEnqueueWriteBuffer(
queue, data_device, CL_TRUE, 0,
ggml_nbytes(tensor), data, 0, NULL, NULL));
cl_mem data_device = ggml_cl_create_temp_upload_buffer(context, queue, ggml_nbytes(tensor), data, tensor->name);
GGML_ASSERT(data_device != NULL && "set_tensor: temp upload buffer alloc failed");
cl_buffer_region region;
@@ -10478,12 +10580,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
GGML_ASSERT(size_d + size_dm + size_s + size_q == ggml_nbytes(tensor) && "Incorrect tensor size");
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
CL_CHECK(clEnqueueWriteBuffer(
queue, data_device, CL_TRUE, 0,
ggml_nbytes(tensor), data, 0, NULL, NULL));
cl_mem data_device = ggml_cl_create_temp_upload_buffer(context, queue, ggml_nbytes(tensor), data, tensor->name);
GGML_ASSERT(data_device != NULL && "q4_K set_tensor: temp upload buffer alloc failed");
cl_buffer_region region;
@@ -10680,9 +10778,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
"Incorrect tensor size");
cl_int err;
cl_mem data_device;
CL_CHECK((data_device = clCreateBuffer(context, CL_MEM_READ_WRITE, ggml_nbytes(tensor), NULL, &err), err));
CL_CHECK(clEnqueueWriteBuffer(queue, data_device, CL_TRUE, 0, ggml_nbytes(tensor), data, 0, NULL, NULL));
cl_mem data_device = ggml_cl_create_temp_upload_buffer(context, queue, ggml_nbytes(tensor), data, tensor->name);
GGML_ASSERT(data_device != NULL && "q5_K set_tensor: temp upload buffer alloc failed");
cl_buffer_region region;
@@ -10868,9 +10965,8 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
"Incorrect tensor size");
cl_int err;
cl_mem data_device;
CL_CHECK((data_device = clCreateBuffer(context, CL_MEM_READ_WRITE, ggml_nbytes(tensor), NULL, &err), err));
CL_CHECK(clEnqueueWriteBuffer(queue, data_device, CL_TRUE, 0, ggml_nbytes(tensor), data, 0, NULL, NULL));
cl_mem data_device = ggml_cl_create_temp_upload_buffer(context, queue, ggml_nbytes(tensor), data, tensor->name);
GGML_ASSERT(data_device != NULL && "q6_K set_tensor: temp upload buffer alloc failed");
cl_buffer_region region;
@@ -11211,9 +11307,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
cl_int err;
cl_kernel kernel = backend_ctx->kernel_restore_block_q4_0_trans4_ns;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
int ne00 = tensor->ne[0];
int ne01 = tensor->ne[1];
@@ -11282,10 +11377,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
}
#endif
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
cl_kernel kernel = backend_ctx->kernel_restore_block_q4_0;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra->q));
@@ -11310,10 +11403,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
if (use_adreno_moe_kernels(backend_ctx, tensor)) {
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
cl_kernel kernel = backend_ctx->kernel_restore_block_q4_1_trans4_ns;
int ne00 = tensor->ne[0];
@@ -11385,10 +11476,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
}
#endif
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
cl_kernel kernel = backend_ctx->kernel_restore_block_q4_1;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra->q));
@@ -11416,9 +11505,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
if (use_adreno_moe_kernels(backend_ctx, tensor)) {
cl_int err;
// TODO: use ggml_cl_buffer to manage this temporary buffer
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
cl_kernel kernel = backend_ctx->kernel_restore_block_q5_0_trans4_ns;
@@ -11520,9 +11608,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
if (use_adreno_moe_kernels(backend_ctx, tensor)) {
cl_int err;
// TODO: use ggml_cl_buffer to manage this temporary buffer
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
cl_kernel kernel = backend_ctx->kernel_restore_block_q5_1_trans4_ns;
@@ -11627,10 +11714,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
if (tensor->type == GGML_TYPE_MXFP4) {
ggml_tensor_extra_cl_mxfp4 * extra = (ggml_tensor_extra_cl_mxfp4 *)tensor->extra;
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
if (use_adreno_moe_kernels(backend_ctx, tensor)) {
@@ -11692,10 +11777,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
const ggml_tensor * extra_src = tensor->view_src != nullptr ? tensor->view_src : tensor;
ggml_tensor_extra_cl_q8_0 * extra = (ggml_tensor_extra_cl_q8_0 *)extra_src->extra;
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
if (enable_adreno_trans_weight(backend_ctx, tensor)) {
@@ -11748,10 +11831,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
if (tensor->type == GGML_TYPE_IQ4_NL) {
ggml_tensor_extra_cl_iq4_nl * extra = (ggml_tensor_extra_cl_iq4_nl *)tensor->extra;
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
if (use_adreno_kernels(backend_ctx, tensor)) {
@@ -11820,10 +11901,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
if (tensor->type == GGML_TYPE_Q4_K) {
ggml_tensor_extra_cl_q4_K * extra = (ggml_tensor_extra_cl_q4_K *)tensor->extra;
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
cl_uchar mask_0F = 0x0F;
cl_uchar mask_F0 = 0xF0;
@@ -11878,10 +11957,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
return;
}
if (use_adreno_moe_kernels(backend_ctx, tensor)) {
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
cl_kernel kernel = backend_ctx->kernel_restore_block_q4_k_trans4_ns;
@@ -11986,20 +12063,16 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
if (tensor->type == GGML_TYPE_Q5_K) {
ggml_tensor_extra_cl_q5_K * extra = (ggml_tensor_extra_cl_q5_K *)tensor->extra;
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
cl_uchar mask_0F = 0x0F;
cl_uchar mask_F0 = 0xF0;
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
if (use_adreno_moe_kernels(backend_ctx, tensor)) {
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
cl_kernel kernel = backend_ctx->kernel_restore_block_q5_k_trans4_ns;
int ne00 = tensor->ne[0];
@@ -12159,10 +12232,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
return;
}
if (use_adreno_moe_kernels(backend_ctx, tensor)) {
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
cl_kernel kernel = backend_ctx->kernel_restore_block_q6_k_trans4_ns;
@@ -12249,10 +12320,8 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
}
#endif // GGML_OPENCL_USE_ADRENO_KERNELS
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
cl_mem data_device = ggml_cl_create_temp_download_buffer(context, queue, ggml_nbytes(tensor), tensor->name);
GGML_ASSERT(data_device != NULL && "get_tensor: temp download buffer alloc failed");
cl_uchar mask = 0xFF;
cl_ulong n_blk = ggml_nelements(tensor)/ggml_blck_size(tensor->type);
@@ -12380,6 +12449,21 @@ static ggml_backend_buffer_t ggml_backend_opencl_buffer_type_alloc_buffer(ggml_b
cl_int err;
cl_mem mem = clCreateBuffer(backend_ctx->context, CL_MEM_READ_WRITE, size, NULL, &err);
// On Adreno X1-85 the device pool intermittently fails at hundreds of MB
// once the heap fragments (e.g. graph-allocator compute-buffer reserve
// after model load). Four-step retry:
// 1. normal alloc (fast path)
// 2. clFinish + retry (drains in-flight allocs)
// 3. cl_qcom_large_buffer (X2-class driver only, OpenCL 3.0 only)
// 4. ALLOC_HOST_PTR (host-pinned pool) — last-resort fallback. This
// buffer backs compute scratch read/written by every kernel in the
// graph, so kernel accesses fall to host memory and runtime perf
// degrades meaningfully. Better than failing to load, but the user
// should see the warning and consider -ngl reduction.
if (err != CL_SUCCESS) {
clFinish(backend_ctx->queue);
mem = clCreateBuffer(backend_ctx->context, CL_MEM_READ_WRITE, size, NULL, &err);
}
#if GGML_OPENCL_TARGET_VERSION >= 300
// clCreateBufferWithProperties and cl_mem_properties are OpenCL 3.0. Drivers older than
// that do not export the symbol, so a build targeting them fails to link. The large
@@ -12390,9 +12474,20 @@ static ggml_backend_buffer_t ggml_backend_opencl_buffer_type_alloc_buffer(ggml_b
mem = clCreateBufferWithProperties(backend_ctx->context, props, CL_MEM_READ_WRITE, size, NULL, &err);
}
#endif
if (err != CL_SUCCESS) {
mem = clCreateBuffer(backend_ctx->context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, size, NULL, &err);
if (err == CL_SUCCESS) {
GGML_LOG_WARN("%s: %.2f MiB allocated via CL_MEM_ALLOC_HOST_PTR fallback — "
"device pool exhausted; runtime perf will be degraded. "
"Consider lowering -ngl or context size.\n",
__func__, size / 1024.0 / 1024.0);
}
}
if (err != CL_SUCCESS) {
GGML_LOG_INFO("%s: failed to allocate %.2f MiB\n", __func__, size / 1024.0 / 1024.0);
GGML_LOG_ERROR("%s: failed to allocate %.2f MiB (err=%d). "
"Consider reducing -ngl, lowering -c / -ub, or using quantized KV cache.\n",
__func__, size / 1024.0 / 1024.0, err);
return nullptr;
}
@@ -13147,6 +13242,7 @@ static void ggml_cl_set_rows(ggml_backend_t backend, const ggml_tensor * src0, c
(size_t)ne03};
size_t local_work_size[] = {(size_t)nth0, (size_t)rows_per_workgroup, 1};
// ne01 == 0 makes global_work_size[0] zero here; enqueue_ndrange_kernel drops the empty range.
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
}