From 7743b8f2be3f53ca50d327f9ddfab76aebb94604 Mon Sep 17 00:00:00 2001 From: BadStorm Date: Fri, 10 Jul 2026 14:28:11 +0200 Subject: [PATCH] Fix opencl support --- .../llamacpp/llamacpp-opencl.Containerfile | 5 + containers/llamacpp/opencl-mesa-rusticl.patch | 125 +++++++++++++++++- 2 files changed, 125 insertions(+), 5 deletions(-) diff --git a/containers/llamacpp/llamacpp-opencl.Containerfile b/containers/llamacpp/llamacpp-opencl.Containerfile index 607350b..6cf9baf 100644 --- a/containers/llamacpp/llamacpp-opencl.Containerfile +++ b/containers/llamacpp/llamacpp-opencl.Containerfile @@ -58,6 +58,11 @@ RUN git clone --depth 1 https://github.com/ggml-org/llama.cpp # 6. mul_mv_f16_f32_l4.cl: la macro sub_group_shuffle_xor->qcom_* scattava # con la sola macro del punto 3, ma richiede una diversa estensione # proprietaria (cl_qcom_subgroup_shuffle) che Mesa non implementa +# 7. clampa local_work_size anche contro CL_DEVICE_MAX_WORK_ITEM_SIZES +# (1024 su questo device, non solo il limite di prodotto CL_DEVICE_ +# MAX_WORK_GROUP_SIZE=2048) in 9 punti di dispatch - senza questo, +# QUALSIASI modello con n_embd (o dimensione analoga) sopra 1024 +# falliva con CL_INVALID_WORK_ITEM_SIZE (es. Qwen2.5-1.5B, n_embd=1536) COPY opencl-mesa-rusticl.patch /tmp/ RUN cd llama.cpp && git apply /tmp/opencl-mesa-rusticl.patch \ && rm -f /tmp/opencl-mesa-rusticl.patch diff --git a/containers/llamacpp/opencl-mesa-rusticl.patch b/containers/llamacpp/opencl-mesa-rusticl.patch index 3bb0801..29b4298 100644 --- a/containers/llamacpp/opencl-mesa-rusticl.patch +++ b/containers/llamacpp/opencl-mesa-rusticl.patch @@ -1,8 +1,16 @@ diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp -index 5c96b9a9f..90ce16da8 100644 +index 5c96b9a9f..6e1b19c89 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp -@@ -1160,7 +1160,8 @@ static void load_cl_kernels_argsort(ggml_backend_opencl_context *backend_ctx) { +@@ -529,6 +529,7 @@ struct ggml_backend_opencl_context { + size_t global_mem_size; + size_t max_alloc_size; + size_t max_workgroup_size; ++ size_t max_work_item_size0; + bool fp16_support; + bool has_vector_subgroup_broadcast; + bool has_subgroup_shuffle = false; // cl_khr_subgroup_shuffle or cl_qcom_subgroup_shuffle +@@ -1160,7 +1161,8 @@ static void load_cl_kernels_argsort(ggml_backend_opencl_context *backend_ctx) { std::string("CL") + std::to_string(backend_ctx->opencl_c_version.major) + "." + std::to_string(backend_ctx->opencl_c_version.minor); std::string compile_opts = std::string("-cl-std=") + opencl_c_std + " -cl-mad-enable -cl-unsafe-math-optimizations" @@ -12,7 +20,7 @@ index 5c96b9a9f..90ce16da8 100644 // argsort if (!backend_ctx->kernels_loaded_argsort) { -@@ -1203,7 +1204,8 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { +@@ -1203,7 +1205,8 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { std::string("CL") + std::to_string(backend_ctx->opencl_c_version.major) + "." + std::to_string(backend_ctx->opencl_c_version.minor); std::string compile_opts = std::string("-cl-std=") + opencl_c_std + " -cl-mad-enable -cl-unsafe-math-optimizations" @@ -22,7 +30,7 @@ index 5c96b9a9f..90ce16da8 100644 if (backend_ctx->adreno_use_large_buffer) { compile_opts += " -qcom-enable-large-buffer "; -@@ -5343,6 +5345,7 @@ static bool ggml_opencl_is_device_supported(ggml_backend_dev_t dev) { +@@ -5343,6 +5346,7 @@ static bool ggml_opencl_is_device_supported(ggml_backend_dev_t dev) { GGML_ASSERT(dev_ctx->device); if (strstr(dev_ctx->device_name.c_str(), "Adreno") || @@ -30,7 +38,7 @@ index 5c96b9a9f..90ce16da8 100644 strstr(dev_ctx->device_name.c_str(), "Qualcomm") || strstr(dev_ctx->device_version.c_str(), "Adreno")) { dev_ctx->gpu_family = GPU_FAMILY::ADRENO; -@@ -5394,7 +5397,8 @@ static bool ggml_opencl_is_device_supported(ggml_backend_dev_t dev) { +@@ -5394,7 +5398,8 @@ static bool ggml_opencl_is_device_supported(ggml_backend_dev_t dev) { // If OpenCL 3.0 is supported, then check for cl_khr_subgroups, which becomes // optional in OpenCL 3.0 (cl_khr_subgroup is mandatory in OpenCL 2.x) if (opencl_c_version.major == 3 && strstr(ext_buffer, "cl_khr_subgroups") == NULL && @@ -40,6 +48,113 @@ index 5c96b9a9f..90ce16da8 100644 GGML_LOG_WARN("ggml_opencl: device does not support subgroups (cl_khr_subgroups or cl_intel_subgroups) " "(note that subgroups is an optional feature in OpenCL 3.0)\n"); return false; +@@ -5502,6 +5507,15 @@ static ggml_backend_opencl_context * ggml_cl_init(ggml_backend_dev_t dev) { + CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof(size_t), &backend_ctx->image2d_max_width, NULL)); + CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof(size_t), &backend_ctx->image2d_max_height, NULL)); + CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), &backend_ctx->max_workgroup_size, NULL)); ++ { ++ // CL_DEVICE_MAX_WORK_ITEM_SIZES is a per-dimension cap that can be ++ // stricter than CL_DEVICE_MAX_WORK_GROUP_SIZE (the overall product ++ // cap) - e.g. this Mesa/rusticl Adreno device reports 1024 here but ++ // 2048 for MAX_WORK_GROUP_SIZE. A 1-D dispatch must respect both. ++ size_t max_work_item_sizes[3] = {0, 0, 0}; ++ CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(max_work_item_sizes), max_work_item_sizes, NULL)); ++ backend_ctx->max_work_item_size0 = max_work_item_sizes[0]; ++ } + CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_SVM_CAPABILITIES, sizeof(cl_device_svm_capabilities), &backend_ctx->svm_caps, 0)); + + if (opencl_c_version.major >= 3) { +@@ -6748,7 +6762,7 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te + load_cl_kernels_argsort(backend_ctx); + + cl_kernel kernel = backend_ctx->kernel_argsort_f32_i32; +- int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); ++ int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); + + int cols = 1; + while (cols < op->ne[0]) { +@@ -10491,7 +10505,7 @@ static void ggml_cl_get_rows(ggml_backend_t backend, const ggml_tensor * src0, c + CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_ulong), &nb2)); + CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_ulong), &nb3)); + +- int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); ++ int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); + int nth = 1; + while (nth < ne00 && 2*nth <= max_workgroup_size) { + nth *= 2; +@@ -10685,7 +10699,7 @@ static void ggml_cl_set_rows(ggml_backend_t backend, const ggml_tensor * src0, c + nth0 = 64; + } + +- int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); ++ int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); + while (nth0 < nblk0 && nth0 < max_workgroup_size) { + nth0 *= 2; + } +@@ -10945,7 +10959,7 @@ static void ggml_cl_add_id(ggml_backend_t backend, const ggml_tensor * src0, con + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne0)); + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &ne1)); + +- int nth = MIN(ne00, (int) backend_ctx->get_kernel_workgroup_size(kernel)); ++ int nth = MIN(ne00, MIN((int) backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0)); + size_t global_work_size[] = { (size_t)ne01*nth, (size_t)ne02, 1 }; + size_t local_work_size[] = { (size_t)nth, 1, 1 }; + +@@ -12095,7 +12109,7 @@ static void ggml_opencl_op_rms_norm_fused(ggml_backend_t backend, ggml_tensor * + cl_kernel kernel = backend_ctx->kernel_rms_norm_mul; + + int nth = sgs; +- int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); ++ int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); + while (nth < ne00 && nth < max_workgroup_size) { + nth *= 2; + } +@@ -12173,7 +12187,7 @@ static void ggml_opencl_op_norm_fused(ggml_backend_t backend, ggml_tensor * norm + cl_kernel kernel = backend_ctx->kernel_norm_mul_add; + + int nth = sgs; +- int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); ++ int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); + while (nth < ne00/4 && nth < max_workgroup_size) nth *= 2; + nth = MIN(nth, max_workgroup_size); + nth = MIN(nth, ne00/4); +@@ -12246,7 +12260,7 @@ static void ggml_opencl_op_group_norm_fused(ggml_backend_t backend, ggml_tensor + memcpy(&eps, (char *)gn_tensor->op_params + sizeof(int), sizeof(float)); + + cl_kernel kernel = backend_ctx->kernel_group_norm_mul_add; +- int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); ++ int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); + int ne = ggml_nelements(src0); + int group_size = ne / groups; + +@@ -12352,7 +12366,8 @@ static void ggml_cl_l2_norm(ggml_backend_t backend, const ggml_tensor * src0, co + cl_kernel kernel = backend_ctx->kernel_l2_norm_f32; + + int nth = sgs; +- while (nth < ne00 && nth < (int)backend_ctx->get_kernel_workgroup_size(kernel)) { ++ int max_workgroup_size_l2n = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); ++ while (nth < ne00 && nth < max_workgroup_size_l2n) { + nth *= 2; + } + +@@ -21401,7 +21416,7 @@ static void ggml_cl_set(ggml_backend_t backend, const ggml_tensor * src0, const + CL_CHECK(clSetKernelArg(kernel, 18, sizeof(cl_ulong), &pnb2)); + CL_CHECK(clSetKernelArg(kernel, 19, sizeof(cl_ulong), &pnb3)); + +- int max_local_size = backend_ctx->get_kernel_workgroup_size(kernel); ++ int max_local_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); + + const int nth = MIN(max_local_size, ne00); + +@@ -22112,7 +22127,7 @@ static void ggml_cl_cumsum(ggml_backend_t backend, const ggml_tensor * src0, con + + cl_kernel kernel = backend_ctx->kernel_cumsum_blk; + +- int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); ++ int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); + int nth = 1; + while (nth < ne00 && 2*nth <= max_workgroup_size) { + nth *= 2; diff --git a/ggml/src/ggml-opencl/kernels/mul_mm_q4_k_f32_l4_lm.cl b/ggml/src/ggml-opencl/kernels/mul_mm_q4_k_f32_l4_lm.cl index 2235b1ae8..dc191b1d2 100644 --- a/ggml/src/ggml-opencl/kernels/mul_mm_q4_k_f32_l4_lm.cl