cuda: align snake fusion matcher with the other backends (#25460 )

* cuda: fix snake fusion type predicate, a and inv_b are F32 The matcher required a->type == x->type while launch_snake reads both as const float *, matching the CPU and Metal contract where a and inv_b stay F32. F16/BF16 chains never fused and fell back to the naive path, and a hypothetical all F16 chain would have read F16 bits as float. Aligns the predicate and the comment with ggml-cpu.c * cuda: reject snake fusion on non-contiguous operands The kernel reads x[idx] and a[c] / inv_b[c] linearly, so a non-contiguous view passing the matcher would silently read wrong data. Mirror the contiguity guard already present in the CPU, Vulkan and Metal matchers.
server : respect min-step when splitting prompt batches (#25420 )
2026-07-10 06:25:55 +02:00 · 2026-07-09 11:00:06 +03:00 · 2026-07-09 01:23:30 -05:00 · 2026-07-08 21:55:00 -07:00 · 2026-07-09 08:34:19 +09:00 · 2026-07-08 15:52:21 -07:00
17 changed files with 757 additions and 100 deletions
@@ -3165,18 +3165,21 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
                              (a->ne[2]   == 1 && a->ne[3]   == 1);
        const bool shape_ok = ggml_are_same_shape(a, inv_b) && a->ne[0] == 1 && a->ne[1] == x->ne[1];

-        // x must be in the supported whitelist and every operand / intermediate
-        // result must share x's type, since launch_snake casts a / inv_b as
-        // float and templates the kernel on a single T. Mixed precision chains
-        // fall back to the naive path.
+        // x is in the supported whitelist and every chain intermediate shares
+        // x's type. launch_snake reads a and inv_b as const float *, so they
+        // stay F32.
        const ggml_tensor * sin1 = cgraph->nodes[i + 1];
        const bool types_ok = (x->type == GGML_TYPE_F32 || x->type == GGML_TYPE_F16 || x->type == GGML_TYPE_BF16) &&
-                              (a->type    == x->type) && (inv_b->type == x->type) &&
+                              (a->type    == GGML_TYPE_F32) && (inv_b->type == GGML_TYPE_F32) &&
                              (mul0->type == x->type) && (sin1->type  == x->type) &&
                              (sqr->type  == x->type) && (mul1->type  == x->type) &&
                              (add->type  == x->type);

-        if (types_ok && shape_ok && dim_ok && x_in_add == x) {
+        // kernel reads x[idx] and a[c] / inv_b[c] linearly, so every operand is contiguous
+        const bool contig_ok = ggml_is_contiguous(x) && ggml_is_contiguous(add) &&
+                               ggml_is_contiguous(a) && ggml_is_contiguous(inv_b);
+
+        if (types_ok && shape_ok && dim_ok && contig_ok && x_in_add == x) {
            ggml_cuda_op_snake_fused(*cuda_ctx, x, a, inv_b, add);
            return 4;
        }
@@ -3105,8 +3105,12 @@ static bool ggml_hexagon_supported_rope(const struct ggml_hexagon_session * sess

    int mode = op_params[2];

+    // n_dims == ne0/2, so the rotation spans the full row
    if (mode == GGML_ROPE_TYPE_VISION) {
-        return false;
+        const int n_dims = op_params[1];
+        if (n_dims != (int) (op->src[0]->ne[0] / 2)) {
+            return false;
+        }
    }
    if (mode & 1) {
        return false;
@@ -3137,16 +3141,23 @@ static bool ggml_hexagon_supported_rope(const struct ggml_hexagon_session * sess
    }

    if (src2) {
-        if (!ggml_is_contiguous(src0) || !ggml_is_contiguous(src1) || !ggml_is_contiguous(src2) ||
-            !ggml_is_contiguous(dst)) {
+        if (!ggml_is_contiguous(src1) || !ggml_is_contiguous(src2)) {
            return false;
        }
    } else {
-        if (!ggml_is_contiguous(src0) || !ggml_is_contiguous(src1) || !ggml_is_contiguous(dst)) {
+        if (!ggml_is_contiguous(src1)) {
            return false;
        }
    }

+    // src0/dst elements within a row must be contiguous (nb[0] == sizeof(float)).
+    // nb[1] may exceed ne[0]*sizeof(float) when the tensor is a strided view of a larger one
+    if (src0->nb[0] != sizeof(float) || dst->nb[0] != sizeof(float)) {
+        return false;
+    }
+    if (src0->nb[1] < src0->ne[0] * sizeof(float) || dst->nb[1] < dst->ne[0] * sizeof(float)) {
+        return false;
+    }
    return true;
 }

@@ -23,6 +23,7 @@
 #define HTP_ROPE_TYPE_NORMAL 0
 #define HTP_ROPE_TYPE_NEOX   2
 #define HTP_ROPE_TYPE_MROPE  8
+#define HTP_ROPE_TYPE_VISION 24
 #define HTP_ROPE_TYPE_IMROPE 40

 #define HTP_ROPE_SPAD_NROWS  16
@@ -70,7 +71,9 @@ struct htp_rope_context {
    struct htp_ops_context * octx;

    size_t src0_row_size;
+    size_t src0_row_stride;
    size_t dst_row_size;
+    size_t dst_row_stride;
    size_t src0_row_size_aligned;
    size_t dst_row_size_aligned;
    size_t theta_cache_offset;
@@ -210,6 +213,7 @@ static __attribute__((noinline)) void mrope_cache_init(const float    pos_t,
                             const float    pos_e,
                             const int32_t  sections[4],
                             const bool     is_imrope,
+                             const bool     indep_sects,
                             const float    freq_scale,
                             const float *  freq_factors,
                             float *        corr_dims,
@@ -231,6 +235,14 @@ static __attribute__((noinline)) void mrope_cache_init(const float    pos_t,
        const float ff     = freq_factors ? freq_factors[i0 / 2] : 1.0f;
        const int   sector = (i0 / 2) % sect_dims;

+        if (indep_sects) {
+            // Reset theta when crossing into a new section.
+            if      (sector == 0)           { theta_t = pos_t; }
+            else if (sector == sections[0]) { theta_h = pos_h; }
+            else if (sector == sec_w)       { theta_w = pos_w; }
+            else if (sector == sec_e)       { theta_e = pos_e; }
+        }
+
        float theta;
        if (is_imrope) {
            // Interleaved: sector mod 3 selects component
@@ -422,6 +434,17 @@ static void inline rope_neox_f32(struct htp_rope_context * rctx, uint8_t * restr
    }
 }

+static void inline rope_vision_f32(struct htp_rope_context * rctx, uint8_t * restrict dst, uint8_t * restrict src,
+                   uint32_t nr, uint32_t ne0, const float * restrict theta_cache) {
+    #pragma unroll(4)
+    for (uint32_t i = 0; i < nr; i++) {
+        float * d = (float *) (dst + i * rctx->dst_row_size_aligned);
+        float * s = (float *) (src + i * rctx->src0_row_size_aligned);
+
+        hvx_rope_neox_f32_aa(d, s, ne0, theta_cache);
+    }
+}
+
 static void rope_job_f32(unsigned int nth, unsigned int ith, void * data) {
    struct htp_rope_context * rctx = (struct htp_rope_context *) data;
    struct htp_ops_context * octx = rctx->octx;
@@ -447,8 +470,9 @@ static void rope_job_f32(unsigned int nth, unsigned int ith, void * data) {
    uint64_t tt = HAP_perf_get_qtimer_count();

    const int32_t mode    = rctx->mode;
-    // MROPE and IMROPE use NEOX-style pairing for the rotation
+    // MROPE, IMROPE and VISION use NEOX-style pairing for the rotation
    const bool    is_neox = (mode & HTP_ROPE_TYPE_NEOX) || (mode & HTP_ROPE_TYPE_MROPE);
+    const bool    is_vision = (mode == HTP_ROPE_TYPE_VISION);

    // VTCM setup
    uint8_t * src0_spad_base = octx->src0_spad.data + (ith * octx->src0_spad.size_per_thread);
@@ -496,8 +520,10 @@ static void rope_job_f32(unsigned int nth, unsigned int ith, void * data) {

                    const uint8_t * src_addr = (const uint8_t *) src0->data + i3 * nb03 + i2 * nb02 + pi1 * nb01;
                          uint8_t * src_spad = src0_spad_base + pr * rctx->src0_row_size_aligned;
-                    dma_queue_push_ddr_to_vtcm(dma_queue, dma_make_ptr(src_spad, src_addr),
-                        rctx->src0_row_size_aligned, rctx->src0_row_size, pnr);
+
+                    // Copy only the row payload while striding the DDR source
+                    dma_queue_push(dma_queue, dma_make_ptr(src_spad, src_addr),
+                        rctx->src0_row_size_aligned, rctx->src0_row_stride, rctx->src0_row_size, pnr);

                    // FARF(HIGH, "rope-prefetch %u: pr %u i1 %u i2 %u i3 %u src-spad %p src-addr %p pnr %u", ith, pir, pi1, i2, i3, src_spad, src_addr, pnr);
                }
@@ -516,7 +542,7 @@ static void rope_job_f32(unsigned int nth, unsigned int ith, void * data) {
                            (float) pos[i2 + ne2],
                            (float) pos[i2 + ne2 * 2],
                            (float) pos[i2 + ne2 * 3],
-                            rctx->sections, is_imrope,
+                            rctx->sections, is_imrope, is_vision,
                            rctx->freq_scale, freq_factors, rctx->corr_dims,
                            ne0, rctx->ext_factor, rctx->attn_factor,
                            theta_cache, rctx->theta_scale);
@@ -542,14 +568,19 @@ static void rope_job_f32(unsigned int nth, unsigned int ith, void * data) {
                    // FARF(HIGH, "rope-compute %u: ir %u i1 %u i2 %u i3 %u src-spad %p cnr %u : usec %u", ith, ir, i1, i2, i3, src_spad, cnr,
                    //         (unsigned) HAP_perf_qtimer_count_to_us(HAP_perf_get_qtimer_count() - rctx->t_start));

-                    if (is_neox) {
+                    if (is_vision) {
+                        rope_vision_f32(rctx, dst_spad, src_spad, cnr, ne0, theta_cache);
+                    } else if (is_neox) {
                        rope_neox_f32(rctx, dst_spad, src_spad, cnr, ne0, theta_cache);
                    } else {
                        rope_basic_f32(rctx, dst_spad, src_spad, cnr, ne0, theta_cache);
                    }

                    uint8_t * dst_addr = (uint8_t *) dst->data + i3 * nb3 + i2 * nb2 + i1 * nb1;
-                    dma_queue_push_vtcm_to_ddr(dma_queue, dma_make_ptr(dst_addr, dst_spad), rctx->dst_row_size, rctx->dst_row_size_aligned, cnr);
+
+                    // Write only the row payload while striding the DDR dst
+                    dma_queue_push(dma_queue, dma_make_ptr(dst_addr, dst_spad),
+                        rctx->dst_row_stride, rctx->dst_row_size_aligned, rctx->dst_row_size, cnr);

                    // Prefetch more rows (if any)
                    if ((cr + HTP_ROPE_SPAD_NROWS) < nrows) {
@@ -558,8 +589,8 @@ static void rope_job_f32(unsigned int nth, unsigned int ith, void * data) {
                        uint32_t pir = ir + HTP_ROPE_SPAD_NROWS;

                        const uint8_t * src_addr = (const uint8_t *) src0->data + i3 * nb03 + i2 * nb02 + pi1 * nb01;
-                        dma_queue_push_ddr_to_vtcm(dma_queue, dma_make_ptr(src_spad, src_addr),
-                            rctx->src0_row_size_aligned, rctx->src0_row_size, pnr);
+                        dma_queue_push(dma_queue, dma_make_ptr(src_spad, src_addr),
+                            rctx->src0_row_size_aligned, rctx->src0_row_stride, rctx->src0_row_size, pnr);

                        // FARF(HIGH, "rope-prefetch %u: pr %u i1 %u i2 %u i3 %u src-spad %p src-addr %p pnr %u", ith, pir, pi1, i2, i3, src_spad, src_addr, pnr);
                    }
@@ -598,12 +629,14 @@ static int execute_op_rope_f32(struct htp_ops_context * octx) {
    const uint32_t src0_nrows = src0->ne[1] * src0->ne[2] * src0->ne[3];
    const uint32_t n_threads = MIN(octx->n_threads, src0_nrows);

-    const size_t src0_row_size = src0->nb[1];
-    const size_t dst_row_size  = dst->nb[1];
+    const size_t src0_row_size   = src0->ne[0] * sizeof(float);
+    const size_t src0_row_stride = src0->nb[1];
+    const size_t dst_row_size    = dst->ne[0] * sizeof(float);
+    const size_t dst_row_stride  = dst->nb[1];

    // Aligned row sizes for VTCM
    const size_t src0_row_size_aligned    = hex_round_up(src0_row_size, VLEN);
-    const size_t dst_row_size_aligned     = hex_round_up(dst_row_size, VLEN);
+    const size_t dst_row_size_aligned     = hex_round_up(dst_row_stride, VLEN);
    const size_t theta_cache_size_aligned = hex_round_up(src0->ne[0] * sizeof(float), 256);

    // Calculate spad sizes per thread
@@ -652,8 +685,10 @@ static int execute_op_rope_f32(struct htp_ops_context * octx) {

    rope_corr_dims(rctx.n_dims, rctx.n_ctx_orig, rctx.freq_base, rctx.beta_fast, rctx.beta_slow, rctx.corr_dims);

-    rctx.src0_row_size = src0_row_size;
-    rctx.dst_row_size  = dst_row_size;
+    rctx.src0_row_size   = src0_row_size;
+    rctx.src0_row_stride = src0_row_stride;
+    rctx.dst_row_size    = dst_row_size;
+    rctx.dst_row_stride  = dst_row_stride;
    rctx.src0_row_size_aligned = src0_row_size_aligned;
    rctx.dst_row_size_aligned  = dst_row_size_aligned;
    rctx.theta_cache_offset    = theta_cache_size_aligned;
@@ -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
@@ -6229,8 +6242,14 @@ inline bool use_adreno_kernels(const ggml_backend_opencl_context *backend_ctx, c
        threshold_ne0 = 128;
        threshold_ne1 = 128;
    }
-    return tensor->ne[0] >= threshold_ne0 && tensor->ne[1] >= threshold_ne1 &&
+    bool threashold_ok = tensor->ne[0] >= threshold_ne0 && tensor->ne[1] >= threshold_ne1 &&
            tensor->ne[2] == 1 && tensor->ne[3] == 1;
+
+    // q6_K adreno kernels requires ne1 is multiple of 128
+    if (tensor->type == GGML_TYPE_Q6_K) {
+        return threashold_ok && tensor->ne[1] % 128 == 0;
+    }
+    return threashold_ok;
 }

 inline bool use_adreno_moe_kernels(const ggml_backend_opencl_context *backend_ctx, const ggml_tensor *tensor) {
@@ -6260,6 +6279,19 @@ static inline bool use_flat_gemv_for_large_m_q6_K(const ggml_tensor *tensor) {
    // threshold is well above typical hidden/FFN dims, but below typical vocab sizes.
    // q6_K flat gemv is worse for smaller K; 2048 seems to be a reasonable threshold.
    // note that this forces large M weights to use LM GEMM.
+    // The noshuffle (transposed-weight) layout packs 2 rows per 32-bit texel and the
+    // gemv reads it with a ne01/2 texel stride and an exact-cover dispatch of
+    // ceil(ne01/2 / 64)*64 work-items with no store guard; the gemm uses 4-row tiles.
+    // It is therefore only correct for ne01 % 128 == 0: an odd ne01 (e.g. granitemoe
+    // lm_head [1536, 49155] -- odd vocab) truncates the texel stride, misaligning every
+    // odd column of the transposed layout (gross garbage) and dropping the last row;
+    // other non-multiples over-dispatch and write past the end of dst. Route such
+    // tensors to the flat GEMV + regular convert; the matching GEMM (ne1>1) falls back
+    // to CPU (see supports_op). All standard even-vocab/hidden dims are multiples of
+    // 128 and keep the noshuffle path.
+    if ((tensor->ne[1] % 128 != 0) && tensor->ne[2] == 1 && tensor->ne[3] == 1) {
+        return true;
+    }
    return tensor->ne[1] >= 32768 && tensor->ne[0] >= 2048 && tensor->ne[2] == 1 && tensor->ne[3] == 1;
 }

@@ -9763,12 +9795,30 @@ static bool ggml_backend_opencl_buffer_type_supports_backend(ggml_backend_buffer
    UNUSED(buft);
 }

+static size_t ggml_backend_opencl_buffer_type_get_alloc_size(ggml_backend_buffer_type_t buft, const ggml_tensor * tensor) {
+    size_t size = ggml_nbytes(tensor);
+#ifdef GGML_OPENCL_SOA_Q
+    // set_tensor carves quantized weights into per-component subbuffers (d/q,
+    // ql/qh/s/d, ...) whose origins are each rounded up to the device base
+    // alignment. When a component's size is not a multiple of the alignment
+    // (e.g. q6_K [1536,49155]: size_s = 49155*96 leaves a 96-byte gap at 128-byte
+    // alignment), the aligned carve extends past ggml_nbytes and the last
+    // subbuffer would overlap the next tensor in the pool. Reserve the worst-case
+    // carve slack: at most 5 components (q5_K), i.e. 4 aligned gaps.
+    if (ggml_is_quantized(tensor->type)) {
+        ggml_backend_opencl_device_context * dev_ctx = (ggml_backend_opencl_device_context *) buft->device->context;
+        size += 4 * dev_ctx->backend_ctx->alignment;
+    }
+#endif // GGML_OPENCL_SOA_Q
+    return size;
+}
+
 static ggml_backend_buffer_type_i ggml_backend_opencl_buffer_type_interface = {
    /* .get_name         = */ ggml_backend_opencl_buffer_type_get_name,
    /* .alloc_buffer     = */ ggml_backend_opencl_buffer_type_alloc_buffer,
    /* .get_alignment    = */ ggml_backend_opencl_buffer_type_get_alignment,
    /* .get_max_size     = */ ggml_backend_opencl_buffer_type_get_max_size,
-    /* .get_alloc_size   = */ NULL,
+    /* .get_alloc_size   = */ ggml_backend_opencl_buffer_type_get_alloc_size,
    /* .is_host          = */ NULL,
 };

@@ -19338,6 +19388,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 +19616,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 +19794,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 +19973,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 +20232,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 +20412,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 +20589,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 +20774,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);
@@ -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) {
@@ -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) {
@@ -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) {
@@ -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) {
@@ -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) {
@@ -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) {
@@ -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) {
@@ -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) {
@@ -10310,7 +10310,8 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx
    }

    // Only use mask opt when the mask is fairly large. This hasn't been tuned extensively.
-    bool use_mask_opt = mask && nem1 >= 32 && nem0 * nem1 > 32768 && nem0 >= tuning_params.block_cols * 16;
+    bool use_mask_opt = mask && nem1 >= 32 && nem0 * nem1 > 32768 && nem0 >= tuning_params.block_cols * 16
+                        && (ctx->device->architecture != vk_device_architecture::AMD_GCN || HSK > 256 || HSV > 256);
    vk_fa_pipeline_state fa_pipeline_state = get_fa_pipeline_state(ctx->device, tuning_params, HSK, HSV, aligned, f32acc,
                                                                   mask != nullptr, use_mask_opt, logit_softcap != 0, k->type, v->type);

@@ -2821,23 +2821,16 @@ class ggml_webgpu_shader_lib {
            variant.resize(variant.size() - (sizeof("_mask") - 1));
            variant += "_mask_blk";
        }
-        uint32_t vec_ne = 1u;
-        if (key.common.k_type == GGML_TYPE_F16 && key.common.v_type == GGML_TYPE_F16 &&
-            key.common.head_dim_qk == key.common.head_dim_v) {
-            switch (key.common.head_dim_qk) {
-                case 64:
-                case 192:
-                case 576:
-                    vec_ne = 2u;
-                    break;
-                case 96:
-                    vec_ne = 4u;
-                    break;
-                default:
-                    break;
-            }
+
+        uint32_t d_split = context.min_subgroup_size;
+        if (key.common.k_type == GGML_TYPE_F16 && key.common.v_type == GGML_TYPE_F16) {
+            const uint32_t D     = key.common.head_dim_qk | key.common.head_dim_v;
+            const uint32_t D_lsb = D & (~(D - 1u));
+            d_split              = std::min(std::min(context.min_subgroup_size, 4u), std::max(D_lsb / 4u, 1u));
        }
-        defines.push_back(std::string("VEC_NE=") + std::to_string(vec_ne) + "u");
+
+        defines.push_back(std::string("D_SPLIT=") + std::to_string(d_split));
+        variant += "_dsplit" + std::to_string(d_split);

        auto            pipeline_decisions = std::make_shared<ggml_webgpu_flash_attn_vec_decisions>(decisions);
        webgpu_pipeline pipeline =
@@ -39,9 +39,6 @@ enable subgroups;
 #define KV_GRANULARITY 8
 #define KV_TILE 16
 #define WG_SIZE 64
-#ifndef VEC_NE
-#define VEC_NE 4u
-#endif

 #define KV_BLOCKS (KV_TILE / KV_GRANULARITY)

@@ -367,11 +364,11 @@ fn main(@builtin(workgroup_id) wg_id: vec3<u32>,

      // accumulate q block * k block into registers across the entire KV tile
      if (!skip_tile) {
-        let num_of_threads = subgroup_size / VEC_NE;
+        let num_of_threads:u32 = D_SPLIT;
        let tx = sg_inv_id % num_of_threads;
        let ty = sg_inv_id / num_of_threads;
          if (subgroup_id == 0u && q_row_start < params.seq_len_q) {
-              for (var kv_base : u32 = 0u; kv_base < KV_TILE; kv_base += VEC_NE) {
+              for (var kv_base : u32 = 0u; kv_base < KV_TILE; kv_base += subgroup_size / D_SPLIT) {
                  let kv_idx = kv_base + ty;
                  var partial_sum: f32 = 0.0;
                  let kv_valid = kv_idx < KV_TILE && (kv_tile + kv_idx) < params.seq_len_kv;
@@ -486,15 +483,18 @@ fn main(@builtin(workgroup_id) wg_id: vec3<u32>,
      if (!skip_tile) {
          // we have P (KV_TILE) in inter_shmem and V (KV_TILE x head_dim_v) in kv_shmem
          // we want to compute O += P * V across the full KV tile
-          let ne_threads : u32 = VEC_NE;
+          let ne_threads : u32 = subgroup_size / D_SPLIT;
          let nl_threads = max(1u, subgroup_size / ne_threads);
          let tx_pv = sg_inv_id % nl_threads;
          let ty_pv = sg_inv_id / nl_threads;
          if (subgroup_id == 0u && q_row_start < params.seq_len_q) {
              for (var vec_col = tx_pv; vec_col < (HEAD_DIM_V / 4u); vec_col += nl_threads) {
                  var lo = vec4<f32>(0.0, 0.0, 0.0, 0.0);
-                  for (var cc = 0u; cc < KV_TILE / ne_threads; cc += 1u) {
+                  for (var cc = 0u; cc * ne_threads < KV_TILE; cc += 1u) {
                      let kv_idx = cc * ne_threads + ty_pv;
+                      if (kv_idx >= KV_TILE) {
+                          continue;
+                      }
                      let v_row = kv_tile + kv_idx;
                      if (v_row >= params.seq_len_kv) {
                          continue;
@@ -379,6 +379,8 @@ bool llama_batch_allocr::init(
                    LLAMA_LOG_ERROR("%s: sequence %d positions are decreasing (not allowed)\n", __func__, seq_id);
                    return false;
                }
+
+                cur_seq_pos[seq_id] = pos;
            }
        }
    }
@@ -3435,9 +3435,14 @@ private:

                        slot.n_prompt_tokens_processed++;

-                        // stop the prompt batch exactly before a user message
-                        if (spans.is_user_start(slot.prompt.n_tokens())) {
-                            break;
+                        // break at the last user message, or at user messages at least min step past the last checkpoint
+                        if (do_checkpoint && spans.is_user_start(slot.prompt.n_tokens())) {
+                            const auto pos = slot.prompt.n_tokens();
+                            const auto & checkpoints = slot.prompt.checkpoints;
+
+                            if (pos == last_user_pos || checkpoints.empty() || pos > checkpoints.back().n_tokens + params_base.checkpoint_min_step) {
+                                break;
+                            }
                        }

                        // process the last few tokens of the prompt separately in order to allow for a checkpoint to be created.
Author	SHA1	Message	Date
Pascal	2021515a1a	cuda: align snake fusion matcher with the other backends (#25460 ) * cuda: fix snake fusion type predicate, a and inv_b are F32 The matcher required a->type == x->type while launch_snake reads both as const float , matching the CPU and Metal contract where a and inv_b stay F32. F16/BF16 chains never fused and fell back to the naive path, and a hypothetical all F16 chain would have read F16 bits as float. Aligns the predicate and the comment with ggml-cpu.c cuda: reject snake fusion on non-contiguous operands The kernel reads x[idx] and a[c] / inv_b[c] linearly, so a non-contiguous view passing the matcher would silently read wrong data. Mirror the contiguity guard already present in the CPU, Vulkan and Metal matchers.	2026-07-09 11:00:06 +03:00
Aldehir Rojas	64c8b7db72	server : respect min-step when splitting prompt batches (#25420 )	2026-07-09 01:23:30 -05:00
Aparna M P	f2d1c2f398	hexagon: add VISION RoPE support (#25216 ) * hexagon: add VISION RoPE support * hexagon: support RoPE on strided half-dim views for all modes * hex-rope: decouple src0 DMA copy size from row stride * hex-rope: support non-contiguous dst for RoPE * hex-rope: fix dst spad pitch for non-contiguous dst	2026-07-08 21:55:00 -07:00
Masashi Yoshimura	32e41fa5b4	ggml-webgpu: tune subgroup split (d_split) in flash_attn_vec (#25418 )	2026-07-09 08:34:19 +09:00
Hongqiang Wang	92366df30d	opencl: Q6_K GEMM/GEMV fix for ne01 of weights that are not multiples of 128. (#25464 ) * opencl: fix garbled output for Q6_K weights with ne01 % 128 != 0 on Adreno Observed with granite-3.1-3b-a800m-instruct, whose vocab is an odd number. Route Q6_K dense mul_mat with ne01 % 128 != 0 off the noshuffle path: decode (ne1==1) uses the correct flat GEMV and the matching GEMM (ne1>1) falls back to CPU (the flat convert has no verified small-batch GEMM kernel for these shapes). All standard hidden/FFN/vocab dims are multiples of 128 and keep the noshuffle path. * opencl: reserve alignment slack for the SOA subbuffer carve in alloc size set_tensor carves quantized weights into per-component subbuffers (d/q, ql/qh/s/d, ...) whose origins are each rounded up to the device base address alignment. When a component's size is not a multiple of the alignment, the carve extends past ggml_nbytes(tensor) and the last subbuffer overlaps the next tensor in the pool -- e.g. q6_K [1536, 49155]: size_s = 4915596 ends 32 bytes past a 128-byte boundary, so the d subbuffer ends 96 bytes past the tensor's allocation, and whichever of the two neighboring tensors is uploaded last silently corrupts the other (here: the last vocab rows' block scales). This affects any quant type whose component sizes can be misaligned, on any shape with ne01 not a multiple of the alignment granularity; standard power-of-two dims are unaffected. Implement get_alloc_size for the OpenCL buffer type and reserve the worst-case carve slack (4 aligned gaps; 5 components max, q5_K) for quantized tensors. Costs at most 512 bytes per quantized tensor at the observed 128-byte alignment. opencl: use lm based q6_k mm when ne1 is not multiple of 128 --------- Co-authored-by: Li He <lih@qti.qualcomm.com>	2026-07-08 15:52:21 -07:00
Ruben Ortlam	a646006f09	vulkan: disable FA mask_opt on GCN to improve performance (#24362 ) * vulkan: disable FA mask_opt on GCN to improve performance * reenable mask opt over attention head size 256	2026-07-08 19:01:25 +02:00
Hongqiang Wang	167d057604	opencl: ragged-tile MoE prefill FP16 GEMM optimization (skip padded expert tiles) (#25433 ) * opencl: ragged-tile MoE prefill GEMM (skip padded expert tiles) The MoE prefill GEMM groups tokens into TILESIZE_N=32 per-expert tiles; at low tokens-per-expert most tiles are mostly padding. When a tile's upper 16 slots are all padding (router index 0xFFFFFFFF), skip the second dotx16_reduce8 half. Numerically identical (skipped lanes are padding). Applied to all eight _f32_ns MoE GEMMs; default on, opt out with GGML_OPENCL_MOE_RAGGED_FP16=0. opencl: quarter-granularity ragged MoE tile-skip (8-col skip-groups) Replace the two half-tile dotx16_reduce8 calls in the 8 _f32_ns MoE GEMMs with four dotx8_reduce4 (8-column) calls, skipping each empty trailing skip-group independently. Padding is always trailing, so the kernel rounds the valid count up to the skip granularity and skips fully-padding groups. Byte-identical to the non-skipped path. New env GGML_OPENCL_MOE_RAGGED_GRAN={8,16,32} (quarter/half/ off); default quarter. opencl: move ragged moe env var in cl_init --------- Co-authored-by: Li He <lih@qti.qualcomm.com>	2026-07-08 09:44:55 -07:00
Aman Gupta	1ee093937f	llama-batch: fix allowed decreasing pos in a seq (#25449 )	2026-07-08 19:24:34 +03:00