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;
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
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);
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);
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);
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);
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);
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);
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);
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);
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;
__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
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];
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;
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) {
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(
__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
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];
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;
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) {
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(
__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
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];
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;
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) {
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(
__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
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];
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;
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) {
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(
__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
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];
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;
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) {
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(
__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
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];
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;
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) {
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(
__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
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];
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;
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) {
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(
__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
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];
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;
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) {
overview / pkg / ggml / sources / llama.cpp / commitdiff