diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index 9d1d106784..10e183530e 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -517,6 +517,10 @@ struct ggml_backend_opencl_context {
     bool has_qcom_subgroup_shuffle = false;  // specifically cl_qcom_subgroup_shuffle
     bool disable_fusion;
 
+    // ragged moe, use int to directly pass to kernel
+    cl_uint  adreno_use_moe_ragged;
+    cl_uint  adreno_moe_ragged_skip_gran;
+
     bool adreno_has_large_buffer;
     bool adreno_use_large_buffer;
     bool adreno_use_bin_kernels;
@@ -5342,6 +5346,15 @@ static ggml_backend_opencl_context * ggml_cl_init(ggml_backend_dev_t dev) {
     backend_ctx->adreno_use_large_buffer = getenv("GGML_OPENCL_ADRENO_USE_LARGE_BUFFER") != nullptr &&
                                            backend_ctx->gpu_family == GPU_FAMILY::ADRENO;
 
+    // ragged moe, unspecified or non-zero means enabled, set to 0 to disable
+    static const char * ragged_fp16_env = getenv("GGML_OPENCL_MOE_RAGGED_FP16");
+    backend_ctx->adreno_use_moe_ragged = (ragged_fp16_env == NULL) ? 1 : (atoi(ragged_fp16_env) != 0);
+
+    // ragged moe, tile-skip granularity (columns per skip-group): 8 = quarter (default),
+    // 16 = half (legacy), 32 = disabled. Override with GGML_OPENCL_MOE_RAGGED_GRAN={8,16,32}
+    static const char * ragged_gran_env = getenv("GGML_OPENCL_MOE_RAGGED_GRAN");
+    backend_ctx->adreno_moe_ragged_skip_gran = (ragged_gran_env != NULL) ? atoi(ragged_gran_env) : 8;
+
 #ifdef GGML_OPENCL_USE_ADRENO_BIN_KERNELS
     // try loading adreno binary kernels if enabled
     // if fails to load, builtin kernels will be used
@@ -19338,6 +19351,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem),    &(backend_ctx->prealloc_total_tiles.buffer)));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne00));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne01));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),   &backend_ctx->adreno_use_moe_ragged));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),   &backend_ctx->adreno_moe_ragged_skip_gran));
 
                     // set thread grid
                     global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
@@ -19564,6 +19579,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem),    &(backend_ctx->prealloc_total_tiles.buffer)));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne00));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne01));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),   &backend_ctx->adreno_use_moe_ragged));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),   &backend_ctx->adreno_moe_ragged_skip_gran));
 
                     // set thread grid
                     global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
@@ -19740,6 +19757,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem),    &(backend_ctx->prealloc_total_tiles.buffer)));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne00));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne01));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),   &backend_ctx->adreno_use_moe_ragged));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),   &backend_ctx->adreno_moe_ragged_skip_gran));
 
                     // set thread grid
                     global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
@@ -19917,6 +19936,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem),    &(backend_ctx->prealloc_total_tiles.buffer)));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne00));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne01));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),       &backend_ctx->adreno_use_moe_ragged));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),       &backend_ctx->adreno_moe_ragged_skip_gran));
 
                     // set thread grid
                     global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
@@ -20174,6 +20195,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem),    &(backend_ctx->prealloc_total_tiles.buffer)));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne00));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne01));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),   &backend_ctx->adreno_use_moe_ragged));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),   &backend_ctx->adreno_moe_ragged_skip_gran));
 
                     // set thread grid
                     global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
@@ -20352,6 +20375,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem),    &(backend_ctx->prealloc_total_tiles.buffer)));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne00));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne01));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),   &backend_ctx->adreno_use_moe_ragged));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),   &backend_ctx->adreno_moe_ragged_skip_gran));
 
                     // set thread grid
                     global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
@@ -20527,6 +20552,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem),    &(backend_ctx->prealloc_total_tiles.buffer)));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne00));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne01));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),   &backend_ctx->adreno_use_moe_ragged));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),   &backend_ctx->adreno_moe_ragged_skip_gran));
 
                     // set thread grid
                     global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
@@ -20710,6 +20737,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem),    &(backend_ctx->prealloc_total_tiles.buffer)));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne00));
                     CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int),       &ne01));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),   &backend_ctx->adreno_use_moe_ragged));
+                    CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint),   &backend_ctx->adreno_moe_ragged_skip_gran));
 
                     // set thread grid
                     global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
diff --git a/ggml/src/ggml-opencl/kernels/gemm_moe_mxfp4_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemm_moe_mxfp4_f32_ns.cl
index 02cdbdd9fb..834050a4f9 100644
--- a/ggml/src/ggml-opencl/kernels/gemm_moe_mxfp4_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemm_moe_mxfp4_f32_ns.cl
@@ -132,6 +132,46 @@ static inline half8 mxfp4_to_fp16_packed8(ushort2 fp4x8) {
     c_reg.lo += convert_float8(acc.lo); \
     c_reg.hi += convert_float8(acc.hi); \
 
+// Quarter-tile variant: computes 8 output columns (one skip-group) into a float8
+// accumulator. Same reduction order / flush cadence as dotx16_reduce8, so the
+// non-skipped path is byte-identical; it just lets the caller skip empty
+// 8-column groups at finer granularity. Uses a private half8 `acc8`.
+#define dotx8_reduce4(a_reg, b_lm, c_reg, lm_offset) \
+    acc8.s0 = dot(a_reg.s0123, b_lm[lm_offset + 0]); \
+    acc8.s1 = dot(a_reg.s0123, b_lm[lm_offset + 1]); \
+    acc8.s2 = dot(a_reg.s0123, b_lm[lm_offset + 2]); \
+    acc8.s3 = dot(a_reg.s0123, b_lm[lm_offset + 3]); \
+    acc8.s4 = dot(a_reg.s0123, b_lm[lm_offset + 4]); \
+    acc8.s5 = dot(a_reg.s0123, b_lm[lm_offset + 5]); \
+    acc8.s6 = dot(a_reg.s0123, b_lm[lm_offset + 6]); \
+    acc8.s7 = dot(a_reg.s0123, b_lm[lm_offset + 7]); \
+    acc8.s0 += dot(a_reg.s4567, b_lm[lm_offset + 32]); \
+    acc8.s1 += dot(a_reg.s4567, b_lm[lm_offset + 33]); \
+    acc8.s2 += dot(a_reg.s4567, b_lm[lm_offset + 34]); \
+    acc8.s3 += dot(a_reg.s4567, b_lm[lm_offset + 35]); \
+    acc8.s4 += dot(a_reg.s4567, b_lm[lm_offset + 36]); \
+    acc8.s5 += dot(a_reg.s4567, b_lm[lm_offset + 37]); \
+    acc8.s6 += dot(a_reg.s4567, b_lm[lm_offset + 38]); \
+    acc8.s7 += dot(a_reg.s4567, b_lm[lm_offset + 39]); \
+    c_reg += convert_float8(acc8); \
+    acc8.s0 = dot(a_reg.s89ab, b_lm[lm_offset + 64]); \
+    acc8.s1 = dot(a_reg.s89ab, b_lm[lm_offset + 65]); \
+    acc8.s2 = dot(a_reg.s89ab, b_lm[lm_offset + 66]); \
+    acc8.s3 = dot(a_reg.s89ab, b_lm[lm_offset + 67]); \
+    acc8.s4 = dot(a_reg.s89ab, b_lm[lm_offset + 68]); \
+    acc8.s5 = dot(a_reg.s89ab, b_lm[lm_offset + 69]); \
+    acc8.s6 = dot(a_reg.s89ab, b_lm[lm_offset + 70]); \
+    acc8.s7 = dot(a_reg.s89ab, b_lm[lm_offset + 71]); \
+    acc8.s0 += dot(a_reg.scdef, b_lm[lm_offset + 96]); \
+    acc8.s1 += dot(a_reg.scdef, b_lm[lm_offset + 97]); \
+    acc8.s2 += dot(a_reg.scdef, b_lm[lm_offset + 98]); \
+    acc8.s3 += dot(a_reg.scdef, b_lm[lm_offset + 99]); \
+    acc8.s4 += dot(a_reg.scdef, b_lm[lm_offset + 100]); \
+    acc8.s5 += dot(a_reg.scdef, b_lm[lm_offset + 101]); \
+    acc8.s6 += dot(a_reg.scdef, b_lm[lm_offset + 102]); \
+    acc8.s7 += dot(a_reg.scdef, b_lm[lm_offset + 103]); \
+    c_reg += convert_float8(acc8); \
+
 
 static inline half e8m0_to_fp16(uchar x) {
     ushort bits;
@@ -157,7 +197,9 @@ kernel void kernel_gemm_moe_mxfp4_f32_ns(
         __write_only image1d_buffer_t dst,
         __global     int *            total_tiles,
         uint ne00,
-        uint ne01
+        uint ne01,
+        uint is_ragged,
+        uint skip_gran
 ) {
     uint block_id_m = get_global_id(1); // m_tile
     uint block_id_n = get_global_id(2); // n_tile
@@ -167,6 +209,28 @@ kernel void kernel_gemm_moe_mxfp4_f32_ns(
         return;
     }
 
+    // Ragged tile-skip: when is_ragged and the upper 16 token-slots of this tile are all
+    // padding (router 0xFFFFFFFF), skip the second (reg_c.hi) dotx16_reduce8 half -> ~half
+    // the GEMM dot for sparse tiles. Numerically identical (the skipped lanes are padding).
+    // Ragged tile-skip: tokens are packed contiguously per expert (moe_scatter fills
+    // lanes 0..V-1, moe_fill pre-pads the rest), so router padding (0xFFFFFFFF) is always
+    // trailing. Find the valid-token count V and round it UP to the skip granularity
+    // skip_gran (columns per skip-group: 8 = quarter, 16 = half/legacy, 32 = disabled).
+    // A 8-column group g is all-padding iff its first column (8*g) >= n_active, so its
+    // dotx8_reduce4 is skipped. Numerically identical (skipped lanes are padding).
+    uint n_active = TILESIZE_N;
+    if (is_ragged && skip_gran < TILESIZE_N) {
+        uint n_valid = TILESIZE_N;
+        for (uint _t = 0; _t < TILESIZE_N; ++_t) {
+            if (src2[block_id_n * TILESIZE_N + _t] == 0xFFFFFFFFu) { n_valid = _t; break; }
+        }
+        n_active = min((uint)TILESIZE_N, ((n_valid + skip_gran - 1) / skip_gran) * skip_gran);
+    }
+    // Group 0 (cols 0-7) always runs; groups 1-3 skip when fully padding.
+    bool skip_g1 = (8u  >= n_active);
+    bool skip_g2 = (16u >= n_active);
+    bool skip_g3 = (24u >= n_active);
+
     __private half16 reg_a;
     __private float32 reg_c = (float32)(0);
     __local half4 shared_b[128];
@@ -216,9 +280,11 @@ kernel void kernel_gemm_moe_mxfp4_f32_ns(
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
         // 32 16x16 fp16 dot product with 8 elements reduction for better precision
-        half16 acc;
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        half8 acc8;
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
 
         // Repeat for second sub-block
         uint half_step = step + TILESIZE_K;
@@ -244,8 +310,10 @@ kernel void kernel_gemm_moe_mxfp4_f32_ns(
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
         // 32 16x16 fp16 dot product with 3-levels reduction for better precision
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
     }
 
     if ((get_global_id(0) + block_id_m * TILESIZE_M) >= ne01) {
diff --git a/ggml/src/ggml-opencl/kernels/gemm_moe_q4_0_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemm_moe_q4_0_f32_ns.cl
index d403ed0cab..58f7932361 100644
--- a/ggml/src/ggml-opencl/kernels/gemm_moe_q4_0_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemm_moe_q4_0_f32_ns.cl
@@ -98,6 +98,46 @@
     c_reg.lo += convert_float8(acc.lo); \
     c_reg.hi += convert_float8(acc.hi); \
 
+// Quarter-tile variant: computes 8 output columns (one skip-group) into a float8
+// accumulator. Same reduction order / flush cadence as dotx16_reduce8, so the
+// non-skipped path is byte-identical; it just lets the caller skip empty
+// 8-column groups at finer granularity. Uses a private half8 `acc8`.
+#define dotx8_reduce4(a_reg, b_lm, c_reg, lm_offset) \
+    acc8.s0 = dot(a_reg.s0123, b_lm[lm_offset + 0]); \
+    acc8.s1 = dot(a_reg.s0123, b_lm[lm_offset + 1]); \
+    acc8.s2 = dot(a_reg.s0123, b_lm[lm_offset + 2]); \
+    acc8.s3 = dot(a_reg.s0123, b_lm[lm_offset + 3]); \
+    acc8.s4 = dot(a_reg.s0123, b_lm[lm_offset + 4]); \
+    acc8.s5 = dot(a_reg.s0123, b_lm[lm_offset + 5]); \
+    acc8.s6 = dot(a_reg.s0123, b_lm[lm_offset + 6]); \
+    acc8.s7 = dot(a_reg.s0123, b_lm[lm_offset + 7]); \
+    acc8.s0 += dot(a_reg.s4567, b_lm[lm_offset + 32]); \
+    acc8.s1 += dot(a_reg.s4567, b_lm[lm_offset + 33]); \
+    acc8.s2 += dot(a_reg.s4567, b_lm[lm_offset + 34]); \
+    acc8.s3 += dot(a_reg.s4567, b_lm[lm_offset + 35]); \
+    acc8.s4 += dot(a_reg.s4567, b_lm[lm_offset + 36]); \
+    acc8.s5 += dot(a_reg.s4567, b_lm[lm_offset + 37]); \
+    acc8.s6 += dot(a_reg.s4567, b_lm[lm_offset + 38]); \
+    acc8.s7 += dot(a_reg.s4567, b_lm[lm_offset + 39]); \
+    c_reg += convert_float8(acc8); \
+    acc8.s0 = dot(a_reg.s89ab, b_lm[lm_offset + 64]); \
+    acc8.s1 = dot(a_reg.s89ab, b_lm[lm_offset + 65]); \
+    acc8.s2 = dot(a_reg.s89ab, b_lm[lm_offset + 66]); \
+    acc8.s3 = dot(a_reg.s89ab, b_lm[lm_offset + 67]); \
+    acc8.s4 = dot(a_reg.s89ab, b_lm[lm_offset + 68]); \
+    acc8.s5 = dot(a_reg.s89ab, b_lm[lm_offset + 69]); \
+    acc8.s6 = dot(a_reg.s89ab, b_lm[lm_offset + 70]); \
+    acc8.s7 = dot(a_reg.s89ab, b_lm[lm_offset + 71]); \
+    acc8.s0 += dot(a_reg.scdef, b_lm[lm_offset + 96]); \
+    acc8.s1 += dot(a_reg.scdef, b_lm[lm_offset + 97]); \
+    acc8.s2 += dot(a_reg.scdef, b_lm[lm_offset + 98]); \
+    acc8.s3 += dot(a_reg.scdef, b_lm[lm_offset + 99]); \
+    acc8.s4 += dot(a_reg.scdef, b_lm[lm_offset + 100]); \
+    acc8.s5 += dot(a_reg.scdef, b_lm[lm_offset + 101]); \
+    acc8.s6 += dot(a_reg.scdef, b_lm[lm_offset + 102]); \
+    acc8.s7 += dot(a_reg.scdef, b_lm[lm_offset + 103]); \
+    c_reg += convert_float8(acc8); \
+
 
 __attribute__((qcom_wave_pair_mode(1))) // 1=force single 2=force pair
 kernel void kernel_gemm_moe_q4_0_f32_ns(
@@ -109,7 +149,9 @@ kernel void kernel_gemm_moe_q4_0_f32_ns(
         __write_only image1d_buffer_t dst,
         __global     int *            total_tiles,
         uint ne00,
-        uint ne01
+        uint ne01,
+        uint is_ragged,
+        uint skip_gran
 ) {
     uint block_id_m = get_global_id(1); // m_tile
     uint block_id_n = get_global_id(2); // n_tile
@@ -119,6 +161,28 @@ kernel void kernel_gemm_moe_q4_0_f32_ns(
         return;
     }
 
+    // Ragged tile-skip: when is_ragged and the upper 16 token-slots of this tile are all
+    // padding (router 0xFFFFFFFF), skip the second (reg_c.hi) dotx16_reduce8 half -> ~half
+    // the GEMM dot for sparse tiles. Numerically identical (the skipped lanes are padding).
+    // Ragged tile-skip: tokens are packed contiguously per expert (moe_scatter fills
+    // lanes 0..V-1, moe_fill pre-pads the rest), so router padding (0xFFFFFFFF) is always
+    // trailing. Find the valid-token count V and round it UP to the skip granularity
+    // skip_gran (columns per skip-group: 8 = quarter, 16 = half/legacy, 32 = disabled).
+    // A 8-column group g is all-padding iff its first column (8*g) >= n_active, so its
+    // dotx8_reduce4 is skipped. Numerically identical (skipped lanes are padding).
+    uint n_active = TILESIZE_N;
+    if (is_ragged && skip_gran < TILESIZE_N) {
+        uint n_valid = TILESIZE_N;
+        for (uint _t = 0; _t < TILESIZE_N; ++_t) {
+            if (src2[block_id_n * TILESIZE_N + _t] == 0xFFFFFFFFu) { n_valid = _t; break; }
+        }
+        n_active = min((uint)TILESIZE_N, ((n_valid + skip_gran - 1) / skip_gran) * skip_gran);
+    }
+    // Group 0 (cols 0-7) always runs; groups 1-3 skip when fully padding.
+    bool skip_g1 = (8u  >= n_active);
+    bool skip_g2 = (16u >= n_active);
+    bool skip_g3 = (24u >= n_active);
+
     __private half16 reg_a;
     __private float32 reg_c = (float32)(0);
     __local half4 shared_b[128];
@@ -167,9 +231,11 @@ kernel void kernel_gemm_moe_q4_0_f32_ns(
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
         // 32 16x16 fp16 dot product with 8 elements reduction for better precision
-        half16 acc;
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        half8 acc8;
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
 
         // Repeat for second sub-block
         uint half_step = step + TILESIZE_K;
@@ -194,8 +260,10 @@ kernel void kernel_gemm_moe_q4_0_f32_ns(
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
         // 32 16x16 fp16 dot product with 3-levels reduction for better precision
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
     }
 
     if ((get_global_id(0) + block_id_m * TILESIZE_M) >= ne01) {
diff --git a/ggml/src/ggml-opencl/kernels/gemm_moe_q4_1_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemm_moe_q4_1_f32_ns.cl
index b2bddf3f73..6593124463 100644
--- a/ggml/src/ggml-opencl/kernels/gemm_moe_q4_1_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemm_moe_q4_1_f32_ns.cl
@@ -98,6 +98,46 @@
     c_reg.lo += convert_float8(acc.lo); \
     c_reg.hi += convert_float8(acc.hi); \
 
+// Quarter-tile variant: computes 8 output columns (one skip-group) into a float8
+// accumulator. Same reduction order / flush cadence as dotx16_reduce8, so the
+// non-skipped path is byte-identical; it just lets the caller skip empty
+// 8-column groups at finer granularity. Uses a private half8 `acc8`.
+#define dotx8_reduce4(a_reg, b_lm, c_reg, lm_offset) \
+    acc8.s0 = dot(a_reg.s0123, b_lm[lm_offset + 0]); \
+    acc8.s1 = dot(a_reg.s0123, b_lm[lm_offset + 1]); \
+    acc8.s2 = dot(a_reg.s0123, b_lm[lm_offset + 2]); \
+    acc8.s3 = dot(a_reg.s0123, b_lm[lm_offset + 3]); \
+    acc8.s4 = dot(a_reg.s0123, b_lm[lm_offset + 4]); \
+    acc8.s5 = dot(a_reg.s0123, b_lm[lm_offset + 5]); \
+    acc8.s6 = dot(a_reg.s0123, b_lm[lm_offset + 6]); \
+    acc8.s7 = dot(a_reg.s0123, b_lm[lm_offset + 7]); \
+    acc8.s0 += dot(a_reg.s4567, b_lm[lm_offset + 32]); \
+    acc8.s1 += dot(a_reg.s4567, b_lm[lm_offset + 33]); \
+    acc8.s2 += dot(a_reg.s4567, b_lm[lm_offset + 34]); \
+    acc8.s3 += dot(a_reg.s4567, b_lm[lm_offset + 35]); \
+    acc8.s4 += dot(a_reg.s4567, b_lm[lm_offset + 36]); \
+    acc8.s5 += dot(a_reg.s4567, b_lm[lm_offset + 37]); \
+    acc8.s6 += dot(a_reg.s4567, b_lm[lm_offset + 38]); \
+    acc8.s7 += dot(a_reg.s4567, b_lm[lm_offset + 39]); \
+    c_reg += convert_float8(acc8); \
+    acc8.s0 = dot(a_reg.s89ab, b_lm[lm_offset + 64]); \
+    acc8.s1 = dot(a_reg.s89ab, b_lm[lm_offset + 65]); \
+    acc8.s2 = dot(a_reg.s89ab, b_lm[lm_offset + 66]); \
+    acc8.s3 = dot(a_reg.s89ab, b_lm[lm_offset + 67]); \
+    acc8.s4 = dot(a_reg.s89ab, b_lm[lm_offset + 68]); \
+    acc8.s5 = dot(a_reg.s89ab, b_lm[lm_offset + 69]); \
+    acc8.s6 = dot(a_reg.s89ab, b_lm[lm_offset + 70]); \
+    acc8.s7 = dot(a_reg.s89ab, b_lm[lm_offset + 71]); \
+    acc8.s0 += dot(a_reg.scdef, b_lm[lm_offset + 96]); \
+    acc8.s1 += dot(a_reg.scdef, b_lm[lm_offset + 97]); \
+    acc8.s2 += dot(a_reg.scdef, b_lm[lm_offset + 98]); \
+    acc8.s3 += dot(a_reg.scdef, b_lm[lm_offset + 99]); \
+    acc8.s4 += dot(a_reg.scdef, b_lm[lm_offset + 100]); \
+    acc8.s5 += dot(a_reg.scdef, b_lm[lm_offset + 101]); \
+    acc8.s6 += dot(a_reg.scdef, b_lm[lm_offset + 102]); \
+    acc8.s7 += dot(a_reg.scdef, b_lm[lm_offset + 103]); \
+    c_reg += convert_float8(acc8); \
+
 
 __attribute__((qcom_wave_pair_mode(1))) // 1=force single 2=force pair
 kernel void kernel_gemm_moe_q4_1_f32_ns(
@@ -110,7 +150,9 @@ kernel void kernel_gemm_moe_q4_1_f32_ns(
         __write_only image1d_buffer_t dst,
         __global     int *            total_tiles,
         uint ne00,
-        uint ne01
+        uint ne01,
+        uint is_ragged,
+        uint skip_gran
 ) {
     uint block_id_m = get_global_id(1); // m_tile
     uint block_id_n = get_global_id(2); // n_tile
@@ -120,6 +162,28 @@ kernel void kernel_gemm_moe_q4_1_f32_ns(
         return;
     }
 
+    // Ragged tile-skip: when is_ragged and the upper 16 token-slots of this tile are all
+    // padding (router 0xFFFFFFFF), skip the second (reg_c.hi) dotx16_reduce8 half -> ~half
+    // the GEMM dot for sparse tiles. Numerically identical (the skipped lanes are padding).
+    // Ragged tile-skip: tokens are packed contiguously per expert (moe_scatter fills
+    // lanes 0..V-1, moe_fill pre-pads the rest), so router padding (0xFFFFFFFF) is always
+    // trailing. Find the valid-token count V and round it UP to the skip granularity
+    // skip_gran (columns per skip-group: 8 = quarter, 16 = half/legacy, 32 = disabled).
+    // A 8-column group g is all-padding iff its first column (8*g) >= n_active, so its
+    // dotx8_reduce4 is skipped. Numerically identical (skipped lanes are padding).
+    uint n_active = TILESIZE_N;
+    if (is_ragged && skip_gran < TILESIZE_N) {
+        uint n_valid = TILESIZE_N;
+        for (uint _t = 0; _t < TILESIZE_N; ++_t) {
+            if (src2[block_id_n * TILESIZE_N + _t] == 0xFFFFFFFFu) { n_valid = _t; break; }
+        }
+        n_active = min((uint)TILESIZE_N, ((n_valid + skip_gran - 1) / skip_gran) * skip_gran);
+    }
+    // Group 0 (cols 0-7) always runs; groups 1-3 skip when fully padding.
+    bool skip_g1 = (8u  >= n_active);
+    bool skip_g2 = (16u >= n_active);
+    bool skip_g3 = (24u >= n_active);
+
     __private half16 reg_a;
     __private float32 reg_c = (float32)(0);
     __local half4 shared_b[128];
@@ -169,9 +233,11 @@ kernel void kernel_gemm_moe_q4_1_f32_ns(
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
         // 32 16x16 fp16 dot product with 8 elements reduction for better precision
-        half16 acc;
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        half8 acc8;
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
 
         // Repeat for second sub-block
         uint half_step = step + TILESIZE_K;
@@ -196,8 +262,10 @@ kernel void kernel_gemm_moe_q4_1_f32_ns(
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
         // 32 16x16 fp16 dot product with 3-levels reduction for better precision
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
     }
 
     if ((get_global_id(0) + block_id_m * TILESIZE_M) >= ne01) {
diff --git a/ggml/src/ggml-opencl/kernels/gemm_moe_q4_k_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemm_moe_q4_k_f32_ns.cl
index ab8228d18c..f96189692e 100644
--- a/ggml/src/ggml-opencl/kernels/gemm_moe_q4_k_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemm_moe_q4_k_f32_ns.cl
@@ -114,6 +114,46 @@ inline void get_scale_min_k4(
     c_reg.lo += convert_float8(acc.lo); \
     c_reg.hi += convert_float8(acc.hi); \
 
+// Quarter-tile variant: computes 8 output columns (one skip-group) into a float8
+// accumulator. Same reduction order / flush cadence as dotx16_reduce8, so the
+// non-skipped path is byte-identical; it just lets the caller skip empty
+// 8-column groups at finer granularity. Uses a private half8 `acc8`.
+#define dotx8_reduce4(a_reg, b_lm, c_reg, lm_offset) \
+    acc8.s0 = dot(a_reg.s0123, b_lm[lm_offset + 0]); \
+    acc8.s1 = dot(a_reg.s0123, b_lm[lm_offset + 1]); \
+    acc8.s2 = dot(a_reg.s0123, b_lm[lm_offset + 2]); \
+    acc8.s3 = dot(a_reg.s0123, b_lm[lm_offset + 3]); \
+    acc8.s4 = dot(a_reg.s0123, b_lm[lm_offset + 4]); \
+    acc8.s5 = dot(a_reg.s0123, b_lm[lm_offset + 5]); \
+    acc8.s6 = dot(a_reg.s0123, b_lm[lm_offset + 6]); \
+    acc8.s7 = dot(a_reg.s0123, b_lm[lm_offset + 7]); \
+    acc8.s0 += dot(a_reg.s4567, b_lm[lm_offset + 32]); \
+    acc8.s1 += dot(a_reg.s4567, b_lm[lm_offset + 33]); \
+    acc8.s2 += dot(a_reg.s4567, b_lm[lm_offset + 34]); \
+    acc8.s3 += dot(a_reg.s4567, b_lm[lm_offset + 35]); \
+    acc8.s4 += dot(a_reg.s4567, b_lm[lm_offset + 36]); \
+    acc8.s5 += dot(a_reg.s4567, b_lm[lm_offset + 37]); \
+    acc8.s6 += dot(a_reg.s4567, b_lm[lm_offset + 38]); \
+    acc8.s7 += dot(a_reg.s4567, b_lm[lm_offset + 39]); \
+    c_reg += convert_float8(acc8); \
+    acc8.s0 = dot(a_reg.s89ab, b_lm[lm_offset + 64]); \
+    acc8.s1 = dot(a_reg.s89ab, b_lm[lm_offset + 65]); \
+    acc8.s2 = dot(a_reg.s89ab, b_lm[lm_offset + 66]); \
+    acc8.s3 = dot(a_reg.s89ab, b_lm[lm_offset + 67]); \
+    acc8.s4 = dot(a_reg.s89ab, b_lm[lm_offset + 68]); \
+    acc8.s5 = dot(a_reg.s89ab, b_lm[lm_offset + 69]); \
+    acc8.s6 = dot(a_reg.s89ab, b_lm[lm_offset + 70]); \
+    acc8.s7 = dot(a_reg.s89ab, b_lm[lm_offset + 71]); \
+    acc8.s0 += dot(a_reg.scdef, b_lm[lm_offset + 96]); \
+    acc8.s1 += dot(a_reg.scdef, b_lm[lm_offset + 97]); \
+    acc8.s2 += dot(a_reg.scdef, b_lm[lm_offset + 98]); \
+    acc8.s3 += dot(a_reg.scdef, b_lm[lm_offset + 99]); \
+    acc8.s4 += dot(a_reg.scdef, b_lm[lm_offset + 100]); \
+    acc8.s5 += dot(a_reg.scdef, b_lm[lm_offset + 101]); \
+    acc8.s6 += dot(a_reg.scdef, b_lm[lm_offset + 102]); \
+    acc8.s7 += dot(a_reg.scdef, b_lm[lm_offset + 103]); \
+    c_reg += convert_float8(acc8); \
+
 
 __attribute__((qcom_wave_pair_mode(1)))
 kernel void kernel_gemm_moe_q4_k_f32_ns(
@@ -127,7 +167,9 @@ kernel void kernel_gemm_moe_q4_k_f32_ns(
         __write_only image1d_buffer_t dst,
         __global     int *            total_tiles,
         uint ne00,
-        uint ne01
+        uint ne01,
+        uint is_ragged,
+        uint skip_gran
 ) {
     uint block_id_m = get_global_id(1); // m_tile
     uint block_id_n = get_global_id(2); // n_tile
@@ -137,6 +179,25 @@ kernel void kernel_gemm_moe_q4_k_f32_ns(
         return;
     }
 
+    // Ragged tile-skip: tokens are packed contiguously per expert (moe_scatter fills
+    // lanes 0..V-1, moe_fill pre-pads the rest), so router padding (0xFFFFFFFF) is always
+    // trailing. Find the valid-token count V and round it UP to the skip granularity
+    // skip_gran (columns per skip-group: 8 = quarter, 16 = half/legacy, 32 = disabled).
+    // A 8-column group g is all-padding iff its first column (8*g) >= n_active, so its
+    // dotx8_reduce4 is skipped. Numerically identical (skipped lanes are padding).
+    uint n_active = TILESIZE_N;
+    if (is_ragged && skip_gran < TILESIZE_N) {
+        uint n_valid = TILESIZE_N;
+        for (uint _t = 0; _t < TILESIZE_N; ++_t) {
+            if (src2[block_id_n * TILESIZE_N + _t] == 0xFFFFFFFFu) { n_valid = _t; break; }
+        }
+        n_active = min((uint)TILESIZE_N, ((n_valid + skip_gran - 1) / skip_gran) * skip_gran);
+    }
+    // Group 0 (cols 0-7) always runs; groups 1-3 skip when fully padding.
+    bool skip_g1 = (8u  >= n_active);
+    bool skip_g2 = (16u >= n_active);
+    bool skip_g3 = (24u >= n_active);
+
     __private half16 reg_a;
     __private float32 reg_c = (float32)(0);
     __local half4 shared_b[128];
@@ -199,9 +260,11 @@ kernel void kernel_gemm_moe_q4_k_f32_ns(
 
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
-        half16 acc;
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        half8 acc8;
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
 
         // Second half (next 16 elements, same sub-block scale)
         uint half_step = step + TILESIZE_K;
@@ -221,8 +284,10 @@ kernel void kernel_gemm_moe_q4_k_f32_ns(
 
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
     }
 
     if ((get_global_id(0) + block_id_m * TILESIZE_M) >= ne01) {
diff --git a/ggml/src/ggml-opencl/kernels/gemm_moe_q5_0_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemm_moe_q5_0_f32_ns.cl
index d1a35d58bb..4b8dc52c29 100644
--- a/ggml/src/ggml-opencl/kernels/gemm_moe_q5_0_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemm_moe_q5_0_f32_ns.cl
@@ -98,6 +98,46 @@
     c_reg.lo += convert_float8(acc.lo); \
     c_reg.hi += convert_float8(acc.hi); \
 
+// Quarter-tile variant: computes 8 output columns (one skip-group) into a float8
+// accumulator. Same reduction order / flush cadence as dotx16_reduce8, so the
+// non-skipped path is byte-identical; it just lets the caller skip empty
+// 8-column groups at finer granularity. Uses a private half8 `acc8`.
+#define dotx8_reduce4(a_reg, b_lm, c_reg, lm_offset) \
+    acc8.s0 = dot(a_reg.s0123, b_lm[lm_offset + 0]); \
+    acc8.s1 = dot(a_reg.s0123, b_lm[lm_offset + 1]); \
+    acc8.s2 = dot(a_reg.s0123, b_lm[lm_offset + 2]); \
+    acc8.s3 = dot(a_reg.s0123, b_lm[lm_offset + 3]); \
+    acc8.s4 = dot(a_reg.s0123, b_lm[lm_offset + 4]); \
+    acc8.s5 = dot(a_reg.s0123, b_lm[lm_offset + 5]); \
+    acc8.s6 = dot(a_reg.s0123, b_lm[lm_offset + 6]); \
+    acc8.s7 = dot(a_reg.s0123, b_lm[lm_offset + 7]); \
+    acc8.s0 += dot(a_reg.s4567, b_lm[lm_offset + 32]); \
+    acc8.s1 += dot(a_reg.s4567, b_lm[lm_offset + 33]); \
+    acc8.s2 += dot(a_reg.s4567, b_lm[lm_offset + 34]); \
+    acc8.s3 += dot(a_reg.s4567, b_lm[lm_offset + 35]); \
+    acc8.s4 += dot(a_reg.s4567, b_lm[lm_offset + 36]); \
+    acc8.s5 += dot(a_reg.s4567, b_lm[lm_offset + 37]); \
+    acc8.s6 += dot(a_reg.s4567, b_lm[lm_offset + 38]); \
+    acc8.s7 += dot(a_reg.s4567, b_lm[lm_offset + 39]); \
+    c_reg += convert_float8(acc8); \
+    acc8.s0 = dot(a_reg.s89ab, b_lm[lm_offset + 64]); \
+    acc8.s1 = dot(a_reg.s89ab, b_lm[lm_offset + 65]); \
+    acc8.s2 = dot(a_reg.s89ab, b_lm[lm_offset + 66]); \
+    acc8.s3 = dot(a_reg.s89ab, b_lm[lm_offset + 67]); \
+    acc8.s4 = dot(a_reg.s89ab, b_lm[lm_offset + 68]); \
+    acc8.s5 = dot(a_reg.s89ab, b_lm[lm_offset + 69]); \
+    acc8.s6 = dot(a_reg.s89ab, b_lm[lm_offset + 70]); \
+    acc8.s7 = dot(a_reg.s89ab, b_lm[lm_offset + 71]); \
+    acc8.s0 += dot(a_reg.scdef, b_lm[lm_offset + 96]); \
+    acc8.s1 += dot(a_reg.scdef, b_lm[lm_offset + 97]); \
+    acc8.s2 += dot(a_reg.scdef, b_lm[lm_offset + 98]); \
+    acc8.s3 += dot(a_reg.scdef, b_lm[lm_offset + 99]); \
+    acc8.s4 += dot(a_reg.scdef, b_lm[lm_offset + 100]); \
+    acc8.s5 += dot(a_reg.scdef, b_lm[lm_offset + 101]); \
+    acc8.s6 += dot(a_reg.scdef, b_lm[lm_offset + 102]); \
+    acc8.s7 += dot(a_reg.scdef, b_lm[lm_offset + 103]); \
+    c_reg += convert_float8(acc8); \
+
 
 __attribute__((qcom_wave_pair_mode(1))) // 1=force single 2=force pair
 kernel void kernel_gemm_moe_q5_0_f32_ns(
@@ -110,7 +150,9 @@ kernel void kernel_gemm_moe_q5_0_f32_ns(
         __write_only image1d_buffer_t dst,
         __global     int *            total_tiles,
         uint ne00,
-        uint ne01
+        uint ne01,
+        uint is_ragged,
+        uint skip_gran
 ) {
     uint block_id_m = get_global_id(1); // m_tile
     uint block_id_n = get_global_id(2); // n_tile
@@ -120,6 +162,28 @@ kernel void kernel_gemm_moe_q5_0_f32_ns(
         return;
     }
 
+    // Ragged tile-skip: when is_ragged and the upper 16 token-slots of this tile are all
+    // padding (router 0xFFFFFFFF), skip the second (reg_c.hi) dotx16_reduce8 half -> ~half
+    // the GEMM dot for sparse tiles. Numerically identical (the skipped lanes are padding).
+    // Ragged tile-skip: tokens are packed contiguously per expert (moe_scatter fills
+    // lanes 0..V-1, moe_fill pre-pads the rest), so router padding (0xFFFFFFFF) is always
+    // trailing. Find the valid-token count V and round it UP to the skip granularity
+    // skip_gran (columns per skip-group: 8 = quarter, 16 = half/legacy, 32 = disabled).
+    // A 8-column group g is all-padding iff its first column (8*g) >= n_active, so its
+    // dotx8_reduce4 is skipped. Numerically identical (skipped lanes are padding).
+    uint n_active = TILESIZE_N;
+    if (is_ragged && skip_gran < TILESIZE_N) {
+        uint n_valid = TILESIZE_N;
+        for (uint _t = 0; _t < TILESIZE_N; ++_t) {
+            if (src2[block_id_n * TILESIZE_N + _t] == 0xFFFFFFFFu) { n_valid = _t; break; }
+        }
+        n_active = min((uint)TILESIZE_N, ((n_valid + skip_gran - 1) / skip_gran) * skip_gran);
+    }
+    // Group 0 (cols 0-7) always runs; groups 1-3 skip when fully padding.
+    bool skip_g1 = (8u  >= n_active);
+    bool skip_g2 = (16u >= n_active);
+    bool skip_g3 = (24u >= n_active);
+
     __private half16 reg_a;
     __private float32 reg_c = (float32)(0);
     __local half4 shared_b[128];
@@ -171,9 +235,11 @@ kernel void kernel_gemm_moe_q5_0_f32_ns(
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
         // 32 16x16 fp16 dot product with 8 elements reduction for better precision
-        half16 acc;
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        half8 acc8;
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
 
         // Repeat for second sub-block
         uint half_step = step + TILESIZE_K;
@@ -198,8 +264,10 @@ kernel void kernel_gemm_moe_q5_0_f32_ns(
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
         // 32 16x16 fp16 dot product with 3-levels reduction for better precision
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
     }
 
     if ((get_global_id(0) + block_id_m * TILESIZE_M) >= ne01) {
diff --git a/ggml/src/ggml-opencl/kernels/gemm_moe_q5_1_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemm_moe_q5_1_f32_ns.cl
index 90d345ecf5..ceb977096e 100644
--- a/ggml/src/ggml-opencl/kernels/gemm_moe_q5_1_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemm_moe_q5_1_f32_ns.cl
@@ -98,6 +98,46 @@
     c_reg.lo += convert_float8(acc.lo); \
     c_reg.hi += convert_float8(acc.hi); \
 
+// Quarter-tile variant: computes 8 output columns (one skip-group) into a float8
+// accumulator. Same reduction order / flush cadence as dotx16_reduce8, so the
+// non-skipped path is byte-identical; it just lets the caller skip empty
+// 8-column groups at finer granularity. Uses a private half8 `acc8`.
+#define dotx8_reduce4(a_reg, b_lm, c_reg, lm_offset) \
+    acc8.s0 = dot(a_reg.s0123, b_lm[lm_offset + 0]); \
+    acc8.s1 = dot(a_reg.s0123, b_lm[lm_offset + 1]); \
+    acc8.s2 = dot(a_reg.s0123, b_lm[lm_offset + 2]); \
+    acc8.s3 = dot(a_reg.s0123, b_lm[lm_offset + 3]); \
+    acc8.s4 = dot(a_reg.s0123, b_lm[lm_offset + 4]); \
+    acc8.s5 = dot(a_reg.s0123, b_lm[lm_offset + 5]); \
+    acc8.s6 = dot(a_reg.s0123, b_lm[lm_offset + 6]); \
+    acc8.s7 = dot(a_reg.s0123, b_lm[lm_offset + 7]); \
+    acc8.s0 += dot(a_reg.s4567, b_lm[lm_offset + 32]); \
+    acc8.s1 += dot(a_reg.s4567, b_lm[lm_offset + 33]); \
+    acc8.s2 += dot(a_reg.s4567, b_lm[lm_offset + 34]); \
+    acc8.s3 += dot(a_reg.s4567, b_lm[lm_offset + 35]); \
+    acc8.s4 += dot(a_reg.s4567, b_lm[lm_offset + 36]); \
+    acc8.s5 += dot(a_reg.s4567, b_lm[lm_offset + 37]); \
+    acc8.s6 += dot(a_reg.s4567, b_lm[lm_offset + 38]); \
+    acc8.s7 += dot(a_reg.s4567, b_lm[lm_offset + 39]); \
+    c_reg += convert_float8(acc8); \
+    acc8.s0 = dot(a_reg.s89ab, b_lm[lm_offset + 64]); \
+    acc8.s1 = dot(a_reg.s89ab, b_lm[lm_offset + 65]); \
+    acc8.s2 = dot(a_reg.s89ab, b_lm[lm_offset + 66]); \
+    acc8.s3 = dot(a_reg.s89ab, b_lm[lm_offset + 67]); \
+    acc8.s4 = dot(a_reg.s89ab, b_lm[lm_offset + 68]); \
+    acc8.s5 = dot(a_reg.s89ab, b_lm[lm_offset + 69]); \
+    acc8.s6 = dot(a_reg.s89ab, b_lm[lm_offset + 70]); \
+    acc8.s7 = dot(a_reg.s89ab, b_lm[lm_offset + 71]); \
+    acc8.s0 += dot(a_reg.scdef, b_lm[lm_offset + 96]); \
+    acc8.s1 += dot(a_reg.scdef, b_lm[lm_offset + 97]); \
+    acc8.s2 += dot(a_reg.scdef, b_lm[lm_offset + 98]); \
+    acc8.s3 += dot(a_reg.scdef, b_lm[lm_offset + 99]); \
+    acc8.s4 += dot(a_reg.scdef, b_lm[lm_offset + 100]); \
+    acc8.s5 += dot(a_reg.scdef, b_lm[lm_offset + 101]); \
+    acc8.s6 += dot(a_reg.scdef, b_lm[lm_offset + 102]); \
+    acc8.s7 += dot(a_reg.scdef, b_lm[lm_offset + 103]); \
+    c_reg += convert_float8(acc8); \
+
 
 __attribute__((qcom_wave_pair_mode(1))) // 1=force single 2=force pair
 kernel void kernel_gemm_moe_q5_1_f32_ns(
@@ -111,7 +151,9 @@ kernel void kernel_gemm_moe_q5_1_f32_ns(
         __write_only image1d_buffer_t dst,
         __global     int *            total_tiles,
         uint ne00,
-        uint ne01
+        uint ne01,
+        uint is_ragged,
+        uint skip_gran
 ) {
     uint block_id_m = get_global_id(1); // m_tile
     uint block_id_n = get_global_id(2); // n_tile
@@ -121,6 +163,28 @@ kernel void kernel_gemm_moe_q5_1_f32_ns(
         return;
     }
 
+    // Ragged tile-skip: when is_ragged and the upper 16 token-slots of this tile are all
+    // padding (router 0xFFFFFFFF), skip the second (reg_c.hi) dotx16_reduce8 half -> ~half
+    // the GEMM dot for sparse tiles. Numerically identical (the skipped lanes are padding).
+    // Ragged tile-skip: tokens are packed contiguously per expert (moe_scatter fills
+    // lanes 0..V-1, moe_fill pre-pads the rest), so router padding (0xFFFFFFFF) is always
+    // trailing. Find the valid-token count V and round it UP to the skip granularity
+    // skip_gran (columns per skip-group: 8 = quarter, 16 = half/legacy, 32 = disabled).
+    // A 8-column group g is all-padding iff its first column (8*g) >= n_active, so its
+    // dotx8_reduce4 is skipped. Numerically identical (skipped lanes are padding).
+    uint n_active = TILESIZE_N;
+    if (is_ragged && skip_gran < TILESIZE_N) {
+        uint n_valid = TILESIZE_N;
+        for (uint _t = 0; _t < TILESIZE_N; ++_t) {
+            if (src2[block_id_n * TILESIZE_N + _t] == 0xFFFFFFFFu) { n_valid = _t; break; }
+        }
+        n_active = min((uint)TILESIZE_N, ((n_valid + skip_gran - 1) / skip_gran) * skip_gran);
+    }
+    // Group 0 (cols 0-7) always runs; groups 1-3 skip when fully padding.
+    bool skip_g1 = (8u  >= n_active);
+    bool skip_g2 = (16u >= n_active);
+    bool skip_g3 = (24u >= n_active);
+
     __private half16 reg_a;
     __private float32 reg_c = (float32)(0);
     __local half4 shared_b[128];
@@ -173,9 +237,11 @@ kernel void kernel_gemm_moe_q5_1_f32_ns(
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
         // 32 16x16 fp16 dot product with 8 elements reduction for better precision
-        half16 acc;
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        half8 acc8;
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
 
         // Repeat for second sub-block
         uint half_step = step + TILESIZE_K;
@@ -200,8 +266,10 @@ kernel void kernel_gemm_moe_q5_1_f32_ns(
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
         // 32 16x16 fp16 dot product with 3-levels reduction for better precision
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
     }
 
     if ((get_global_id(0) + block_id_m * TILESIZE_M) >= ne01) {
diff --git a/ggml/src/ggml-opencl/kernels/gemm_moe_q5_k_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemm_moe_q5_k_f32_ns.cl
index 13c26f6f3b..0cbea13f08 100644
--- a/ggml/src/ggml-opencl/kernels/gemm_moe_q5_k_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemm_moe_q5_k_f32_ns.cl
@@ -114,6 +114,46 @@ inline void get_scale_min_k4(
     c_reg.lo += convert_float8(acc.lo); \
     c_reg.hi += convert_float8(acc.hi); \
 
+// Quarter-tile variant: computes 8 output columns (one skip-group) into a float8
+// accumulator. Same reduction order / flush cadence as dotx16_reduce8, so the
+// non-skipped path is byte-identical; it just lets the caller skip empty
+// 8-column groups at finer granularity. Uses a private half8 `acc8`.
+#define dotx8_reduce4(a_reg, b_lm, c_reg, lm_offset) \
+    acc8.s0 = dot(a_reg.s0123, b_lm[lm_offset + 0]); \
+    acc8.s1 = dot(a_reg.s0123, b_lm[lm_offset + 1]); \
+    acc8.s2 = dot(a_reg.s0123, b_lm[lm_offset + 2]); \
+    acc8.s3 = dot(a_reg.s0123, b_lm[lm_offset + 3]); \
+    acc8.s4 = dot(a_reg.s0123, b_lm[lm_offset + 4]); \
+    acc8.s5 = dot(a_reg.s0123, b_lm[lm_offset + 5]); \
+    acc8.s6 = dot(a_reg.s0123, b_lm[lm_offset + 6]); \
+    acc8.s7 = dot(a_reg.s0123, b_lm[lm_offset + 7]); \
+    acc8.s0 += dot(a_reg.s4567, b_lm[lm_offset + 32]); \
+    acc8.s1 += dot(a_reg.s4567, b_lm[lm_offset + 33]); \
+    acc8.s2 += dot(a_reg.s4567, b_lm[lm_offset + 34]); \
+    acc8.s3 += dot(a_reg.s4567, b_lm[lm_offset + 35]); \
+    acc8.s4 += dot(a_reg.s4567, b_lm[lm_offset + 36]); \
+    acc8.s5 += dot(a_reg.s4567, b_lm[lm_offset + 37]); \
+    acc8.s6 += dot(a_reg.s4567, b_lm[lm_offset + 38]); \
+    acc8.s7 += dot(a_reg.s4567, b_lm[lm_offset + 39]); \
+    c_reg += convert_float8(acc8); \
+    acc8.s0 = dot(a_reg.s89ab, b_lm[lm_offset + 64]); \
+    acc8.s1 = dot(a_reg.s89ab, b_lm[lm_offset + 65]); \
+    acc8.s2 = dot(a_reg.s89ab, b_lm[lm_offset + 66]); \
+    acc8.s3 = dot(a_reg.s89ab, b_lm[lm_offset + 67]); \
+    acc8.s4 = dot(a_reg.s89ab, b_lm[lm_offset + 68]); \
+    acc8.s5 = dot(a_reg.s89ab, b_lm[lm_offset + 69]); \
+    acc8.s6 = dot(a_reg.s89ab, b_lm[lm_offset + 70]); \
+    acc8.s7 = dot(a_reg.s89ab, b_lm[lm_offset + 71]); \
+    acc8.s0 += dot(a_reg.scdef, b_lm[lm_offset + 96]); \
+    acc8.s1 += dot(a_reg.scdef, b_lm[lm_offset + 97]); \
+    acc8.s2 += dot(a_reg.scdef, b_lm[lm_offset + 98]); \
+    acc8.s3 += dot(a_reg.scdef, b_lm[lm_offset + 99]); \
+    acc8.s4 += dot(a_reg.scdef, b_lm[lm_offset + 100]); \
+    acc8.s5 += dot(a_reg.scdef, b_lm[lm_offset + 101]); \
+    acc8.s6 += dot(a_reg.scdef, b_lm[lm_offset + 102]); \
+    acc8.s7 += dot(a_reg.scdef, b_lm[lm_offset + 103]); \
+    c_reg += convert_float8(acc8); \
+
 
 __attribute__((qcom_wave_pair_mode(1)))
 kernel void kernel_gemm_moe_q5_k_f32_ns(
@@ -128,7 +168,9 @@ kernel void kernel_gemm_moe_q5_k_f32_ns(
         __write_only image1d_buffer_t dst,
         __global     int *            total_tiles,
         uint ne00,
-        uint ne01
+        uint ne01,
+        uint is_ragged,
+        uint skip_gran
 ) {
     uint block_id_m = get_global_id(1); // m_tile
     uint block_id_n = get_global_id(2); // n_tile
@@ -138,6 +180,28 @@ kernel void kernel_gemm_moe_q5_k_f32_ns(
         return;
     }
 
+    // Ragged tile-skip: when is_ragged and the upper 16 token-slots of this tile are all
+    // padding (router 0xFFFFFFFF), skip the second (reg_c.hi) dotx16_reduce8 half -> ~half
+    // the GEMM dot for sparse tiles. Numerically identical (the skipped lanes are padding).
+    // Ragged tile-skip: tokens are packed contiguously per expert (moe_scatter fills
+    // lanes 0..V-1, moe_fill pre-pads the rest), so router padding (0xFFFFFFFF) is always
+    // trailing. Find the valid-token count V and round it UP to the skip granularity
+    // skip_gran (columns per skip-group: 8 = quarter, 16 = half/legacy, 32 = disabled).
+    // A 8-column group g is all-padding iff its first column (8*g) >= n_active, so its
+    // dotx8_reduce4 is skipped. Numerically identical (skipped lanes are padding).
+    uint n_active = TILESIZE_N;
+    if (is_ragged && skip_gran < TILESIZE_N) {
+        uint n_valid = TILESIZE_N;
+        for (uint _t = 0; _t < TILESIZE_N; ++_t) {
+            if (src2[block_id_n * TILESIZE_N + _t] == 0xFFFFFFFFu) { n_valid = _t; break; }
+        }
+        n_active = min((uint)TILESIZE_N, ((n_valid + skip_gran - 1) / skip_gran) * skip_gran);
+    }
+    // Group 0 (cols 0-7) always runs; groups 1-3 skip when fully padding.
+    bool skip_g1 = (8u  >= n_active);
+    bool skip_g2 = (16u >= n_active);
+    bool skip_g3 = (24u >= n_active);
+
     __private half16 reg_a;
     __private float32 reg_c = (float32)(0);
     __local half4 shared_b[128];
@@ -204,9 +268,11 @@ kernel void kernel_gemm_moe_q5_k_f32_ns(
 
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
-        half16 acc;
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        half8 acc8;
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
 
         // Second half
         uint half_step = step + TILESIZE_K;
@@ -226,8 +292,10 @@ kernel void kernel_gemm_moe_q5_k_f32_ns(
 
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
     }
 
     if ((get_global_id(0) + block_id_m * TILESIZE_M) >= ne01) {
diff --git a/ggml/src/ggml-opencl/kernels/gemm_moe_q6_k_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemm_moe_q6_k_f32_ns.cl
index 85ccebec78..f32114fcfb 100644
--- a/ggml/src/ggml-opencl/kernels/gemm_moe_q6_k_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemm_moe_q6_k_f32_ns.cl
@@ -98,6 +98,46 @@
     c_reg.lo += convert_float8(acc.lo); \
     c_reg.hi += convert_float8(acc.hi); \
 
+// Quarter-tile variant: computes 8 output columns (one skip-group) into a float8
+// accumulator. Same reduction order / flush cadence as dotx16_reduce8, so the
+// non-skipped path is byte-identical; it just lets the caller skip empty
+// 8-column groups at finer granularity. Uses a private half8 `acc8`.
+#define dotx8_reduce4(a_reg, b_lm, c_reg, lm_offset) \
+    acc8.s0 = dot(a_reg.s0123, b_lm[lm_offset + 0]); \
+    acc8.s1 = dot(a_reg.s0123, b_lm[lm_offset + 1]); \
+    acc8.s2 = dot(a_reg.s0123, b_lm[lm_offset + 2]); \
+    acc8.s3 = dot(a_reg.s0123, b_lm[lm_offset + 3]); \
+    acc8.s4 = dot(a_reg.s0123, b_lm[lm_offset + 4]); \
+    acc8.s5 = dot(a_reg.s0123, b_lm[lm_offset + 5]); \
+    acc8.s6 = dot(a_reg.s0123, b_lm[lm_offset + 6]); \
+    acc8.s7 = dot(a_reg.s0123, b_lm[lm_offset + 7]); \
+    acc8.s0 += dot(a_reg.s4567, b_lm[lm_offset + 32]); \
+    acc8.s1 += dot(a_reg.s4567, b_lm[lm_offset + 33]); \
+    acc8.s2 += dot(a_reg.s4567, b_lm[lm_offset + 34]); \
+    acc8.s3 += dot(a_reg.s4567, b_lm[lm_offset + 35]); \
+    acc8.s4 += dot(a_reg.s4567, b_lm[lm_offset + 36]); \
+    acc8.s5 += dot(a_reg.s4567, b_lm[lm_offset + 37]); \
+    acc8.s6 += dot(a_reg.s4567, b_lm[lm_offset + 38]); \
+    acc8.s7 += dot(a_reg.s4567, b_lm[lm_offset + 39]); \
+    c_reg += convert_float8(acc8); \
+    acc8.s0 = dot(a_reg.s89ab, b_lm[lm_offset + 64]); \
+    acc8.s1 = dot(a_reg.s89ab, b_lm[lm_offset + 65]); \
+    acc8.s2 = dot(a_reg.s89ab, b_lm[lm_offset + 66]); \
+    acc8.s3 = dot(a_reg.s89ab, b_lm[lm_offset + 67]); \
+    acc8.s4 = dot(a_reg.s89ab, b_lm[lm_offset + 68]); \
+    acc8.s5 = dot(a_reg.s89ab, b_lm[lm_offset + 69]); \
+    acc8.s6 = dot(a_reg.s89ab, b_lm[lm_offset + 70]); \
+    acc8.s7 = dot(a_reg.s89ab, b_lm[lm_offset + 71]); \
+    acc8.s0 += dot(a_reg.scdef, b_lm[lm_offset + 96]); \
+    acc8.s1 += dot(a_reg.scdef, b_lm[lm_offset + 97]); \
+    acc8.s2 += dot(a_reg.scdef, b_lm[lm_offset + 98]); \
+    acc8.s3 += dot(a_reg.scdef, b_lm[lm_offset + 99]); \
+    acc8.s4 += dot(a_reg.scdef, b_lm[lm_offset + 100]); \
+    acc8.s5 += dot(a_reg.scdef, b_lm[lm_offset + 101]); \
+    acc8.s6 += dot(a_reg.scdef, b_lm[lm_offset + 102]); \
+    acc8.s7 += dot(a_reg.scdef, b_lm[lm_offset + 103]); \
+    c_reg += convert_float8(acc8); \
+
 
 __attribute__((qcom_wave_pair_mode(1)))
 kernel void kernel_gemm_moe_q6_k_f32_ns(
@@ -111,7 +151,9 @@ kernel void kernel_gemm_moe_q6_k_f32_ns(
         __write_only image1d_buffer_t dst,
         __global     int *            total_tiles,
         uint ne00,
-        uint ne01
+        uint ne01,
+        uint is_ragged,
+        uint skip_gran
 ) {
     uint block_id_m = get_global_id(1); // m_tile
     uint block_id_n = get_global_id(2); // n_tile
@@ -121,6 +163,28 @@ kernel void kernel_gemm_moe_q6_k_f32_ns(
         return;
     }
 
+    // Ragged tile-skip: when is_ragged and the upper 16 token-slots of this tile are all
+    // padding (router 0xFFFFFFFF), skip the second (reg_c.hi) dotx16_reduce8 half -> ~half
+    // the GEMM dot for sparse tiles. Numerically identical (the skipped lanes are padding).
+    // Ragged tile-skip: tokens are packed contiguously per expert (moe_scatter fills
+    // lanes 0..V-1, moe_fill pre-pads the rest), so router padding (0xFFFFFFFF) is always
+    // trailing. Find the valid-token count V and round it UP to the skip granularity
+    // skip_gran (columns per skip-group: 8 = quarter, 16 = half/legacy, 32 = disabled).
+    // A 8-column group g is all-padding iff its first column (8*g) >= n_active, so its
+    // dotx8_reduce4 is skipped. Numerically identical (skipped lanes are padding).
+    uint n_active = TILESIZE_N;
+    if (is_ragged && skip_gran < TILESIZE_N) {
+        uint n_valid = TILESIZE_N;
+        for (uint _t = 0; _t < TILESIZE_N; ++_t) {
+            if (src2[block_id_n * TILESIZE_N + _t] == 0xFFFFFFFFu) { n_valid = _t; break; }
+        }
+        n_active = min((uint)TILESIZE_N, ((n_valid + skip_gran - 1) / skip_gran) * skip_gran);
+    }
+    // Group 0 (cols 0-7) always runs; groups 1-3 skip when fully padding.
+    bool skip_g1 = (8u  >= n_active);
+    bool skip_g2 = (16u >= n_active);
+    bool skip_g3 = (24u >= n_active);
+
     __private half16 reg_a;
     __private float32 reg_c = (float32)(0);
     __local half4 shared_b[128];
@@ -183,9 +247,11 @@ kernel void kernel_gemm_moe_q6_k_f32_ns(
 
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
-        half16 acc;
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        half8 acc8;
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
 
         // Second half
         uint half_step = step + TILESIZE_K;
@@ -205,8 +271,10 @@ kernel void kernel_gemm_moe_q6_k_f32_ns(
 
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
-        dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
-        dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
+        dotx8_reduce4(reg_a, shared_b, reg_c.lo.lo, 0);
+        if (!skip_g1) { dotx8_reduce4(reg_a, shared_b, reg_c.lo.hi, 8); }
+        if (!skip_g2) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.lo, 16); }
+        if (!skip_g3) { dotx8_reduce4(reg_a, shared_b, reg_c.hi.hi, 24); }
     }
 
     if ((get_global_id(0) + block_id_m * TILESIZE_M) >= ne01) {