bdi_podman_serverconf/containers/llamacpp/opencl-mesa-rusticl.patch

diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index 5c96b9a9f..6e1b19c89 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -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"
-                               " -cl-finite-math-only -cl-fast-relaxed-math";
+                               " -cl-finite-math-only -cl-fast-relaxed-math"
+                               " -D cl_qcom_reqd_sub_group_size";
 
     // argsort
     if (!backend_ctx->kernels_loaded_argsort) {
@@ -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"
-                               " -cl-finite-math-only -cl-fast-relaxed-math";
+                               " -cl-finite-math-only -cl-fast-relaxed-math"
+                               " -D cl_qcom_reqd_sub_group_size";
 
     if (backend_ctx->adreno_use_large_buffer) {
         compile_opts += " -qcom-enable-large-buffer ";
@@ -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") ||
+        strstr(dev_ctx->device_name.c_str(), "FD6") ||
         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 +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 &&
-        strstr(ext_buffer, "cl_intel_subgroups") == NULL) {
+        strstr(ext_buffer, "cl_intel_subgroups") == NULL &&
+        strstr(ext_buffer, "cl_khr_subgroup_ballot") == NULL) {
         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
+++ b/ggml/src/ggml-opencl/kernels/mul_mm_q4_k_f32_l4_lm.cl
@@ -90,7 +90,7 @@ kernel void kernel_mul_mm_q4_k_f32_l4_lm(
                 int is = 2 * n + b;
                 int qsi = n * 32 + (iqs % 16) * 2;
 
-                char * scales = src0_s + ib * 12;
+                global uchar * scales = src0_s + ib * 12;
 
                 int scidx0 = (is < 4) ? is : (is + 4);
                 int scidx1 = (is < 4) ? is : (is - 4);
diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl b/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl
index da2e14ae9..b293c73d6 100644
--- a/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl
+++ b/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl
@@ -180,10 +180,12 @@ kernel void kernel_mul_mat_f16_f32_l4_dr(
 // Kernels for decoding, Adreno only for now
 #define MUL_MAT_F16_F32_L4_DR_LS_R2_MAX 8
 
-#ifdef ADRENO_GPU
+#if defined(ADRENO_GPU) && defined(cl_qcom_subgroup_shuffle)
 #pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable
 #define sub_group_shuffle_xor(val, mask) qcom_sub_group_shuffle_xor((val), (mask), CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM, 0.0f)
+#endif
 
+#ifdef ADRENO_GPU
 REQD_SUBGROUP_SIZE_64
 kernel void kernel_mul_mat_f16_f32_l4_dr_ls(
         global char * src0,
diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32.cl b/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32.cl
index 71ab98982..64cff077f 100644
--- a/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32.cl
+++ b/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32.cl
@@ -151,8 +151,8 @@ kernel void kernel_mul_mv_q4_K_f32(
                 acc2.s3 += yh[i+9] * (q2[i/2] & 0xF000);
             }
 
-            float dall = dh[0];
-            float dmin = dh[1];
+            float dall = vload_half(0, dh);
+            float dmin = vload_half(1, dh);
             sumf[row] += dall * ((acc1.s0 + 1.f/256.f * acc1.s1) * sc8[0] +
                                  (acc1.s2 + 1.f/256.f * acc1.s3) * sc8[1] * 1.f/16.f +
                                  (acc2.s0 + 1.f/256.f * acc2.s1) * sc8[4] +