workgroup tuning for Hamoa

Author: shawngu-quic

ggml/src/ggml-opencl/ggml-opencl.cpp

@@ -7754,8 +7754,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
             if (true) { // condition todo
                 cl_int status;
 
-                size_t local_size[3] = {64, 4, 1};
-                size_t global_size[3] = {64, 4, 1};
+                size_t local_size[3] = {64, 2, 1};
+                size_t global_size[3] = {64, 2, 1};
 
                 cl_mem src1_sub_buffer, buf_src1_image, buf_src2;
 
@@ -7771,7 +7771,9 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
 
                     // set thread grid
                     global_size[0] = static_cast<size_t>(ne01);
+                    global_size[1] = 4;
                     global_size[2] = static_cast<size_t>(ne20);
+                    local_size[1] = 4;
                 } else { // for gemm
                     kernel = backend_ctx->kernel_gemm_moe_mxfp4_f32;
 

ggml/src/ggml-opencl/kernels/gemm_moe_mxfp4_f32.cl

@@ -3,7 +3,7 @@
 #pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
 
 #define QK_MXFP4 32
-#define N_SIMDGROUP 4
+#define N_SIMDGROUP 2
 #define SIMDGROUP_WIDTH 64
 #define TILE_SIZE 320
 
@@ -142,12 +142,12 @@ __kernel void kernel_gemm_moe_mxfp4_f32(
     // reduction in local memory, assumes #subgroups=4
     __local float reduceLM[SIMDGROUP_WIDTH * (N_SIMDGROUP - 1)];
     if (sgid == 1) reduceLM[SIMDGROUP_WIDTH * 0 + slid] = sum;
-    if (sgid == 2) reduceLM[SIMDGROUP_WIDTH * 1 + slid] = sum;
-    if (sgid == 3) reduceLM[SIMDGROUP_WIDTH * 2 + slid] = sum;
+    // if (sgid == 2) reduceLM[SIMDGROUP_WIDTH * 1 + slid] = sum;
+    // if (sgid == 3) reduceLM[SIMDGROUP_WIDTH * 2 + slid] = sum;
     barrier(CLK_LOCAL_MEM_FENCE);
     if (sgid == 0) sum += reduceLM[SIMDGROUP_WIDTH * 0 + slid];
-    if (sgid == 0) sum += reduceLM[SIMDGROUP_WIDTH * 1 + slid];
-    if (sgid == 0) sum += reduceLM[SIMDGROUP_WIDTH * 2 + slid];
+    // if (sgid == 0) sum += reduceLM[SIMDGROUP_WIDTH * 1 + slid];
+    // if (sgid == 0) sum += reduceLM[SIMDGROUP_WIDTH * 2 + slid];
 
     // 1 outputs per thread in subgroup 0
     if (sgid == 0) {