mirror of
https://github.com/ggml-org/llama.cpp.git
synced 2026-07-09 05:55:54 +02:00
Compare commits
3 Commits
| Author | SHA1 | Date | |
|---|---|---|---|
 Ruben Ortlam
|
a646006f09 | ||
|
Hongqiang Wang
|
167d057604 | ||
|
Aman Gupta
|
1ee093937f |
@@ -517,6 +517,10 @@ struct ggml_backend_opencl_context {
|
||||
bool has_qcom_subgroup_shuffle = false; // specifically cl_qcom_subgroup_shuffle
|
||||
bool disable_fusion;
|
||||
|
||||
// ragged moe, use int to directly pass to kernel
|
||||
cl_uint adreno_use_moe_ragged;
|
||||
cl_uint adreno_moe_ragged_skip_gran;
|
||||
|
||||
bool adreno_has_large_buffer;
|
||||
bool adreno_use_large_buffer;
|
||||
bool adreno_use_bin_kernels;
|
||||
@@ -5342,6 +5346,15 @@ static ggml_backend_opencl_context * ggml_cl_init(ggml_backend_dev_t dev) {
|
||||
backend_ctx->adreno_use_large_buffer = getenv("GGML_OPENCL_ADRENO_USE_LARGE_BUFFER") != nullptr &&
|
||||
backend_ctx->gpu_family == GPU_FAMILY::ADRENO;
|
||||
|
||||
// ragged moe, unspecified or non-zero means enabled, set to 0 to disable
|
||||
static const char * ragged_fp16_env = getenv("GGML_OPENCL_MOE_RAGGED_FP16");
|
||||
backend_ctx->adreno_use_moe_ragged = (ragged_fp16_env == NULL) ? 1 : (atoi(ragged_fp16_env) != 0);
|
||||
|
||||
// ragged moe, tile-skip granularity (columns per skip-group): 8 = quarter (default),
|
||||
// 16 = half (legacy), 32 = disabled. Override with GGML_OPENCL_MOE_RAGGED_GRAN={8,16,32}
|
||||
static const char * ragged_gran_env = getenv("GGML_OPENCL_MOE_RAGGED_GRAN");
|
||||
backend_ctx->adreno_moe_ragged_skip_gran = (ragged_gran_env != NULL) ? atoi(ragged_gran_env) : 8;
|
||||
|
||||
#ifdef GGML_OPENCL_USE_ADRENO_BIN_KERNELS
|
||||
// try loading adreno binary kernels if enabled
|
||||
// if fails to load, builtin kernels will be used
|
||||
@@ -19338,6 +19351,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &(backend_ctx->prealloc_total_tiles.buffer)));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_use_moe_ragged));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_moe_ragged_skip_gran));
|
||||
|
||||
// set thread grid
|
||||
global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
|
||||
@@ -19564,6 +19579,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &(backend_ctx->prealloc_total_tiles.buffer)));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_use_moe_ragged));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_moe_ragged_skip_gran));
|
||||
|
||||
// set thread grid
|
||||
global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
|
||||
@@ -19740,6 +19757,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &(backend_ctx->prealloc_total_tiles.buffer)));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_use_moe_ragged));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_moe_ragged_skip_gran));
|
||||
|
||||
// set thread grid
|
||||
global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
|
||||
@@ -19917,6 +19936,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &(backend_ctx->prealloc_total_tiles.buffer)));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_use_moe_ragged));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_moe_ragged_skip_gran));
|
||||
|
||||
// set thread grid
|
||||
global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
|
||||
@@ -20174,6 +20195,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &(backend_ctx->prealloc_total_tiles.buffer)));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_use_moe_ragged));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_moe_ragged_skip_gran));
|
||||
|
||||
// set thread grid
|
||||
global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
|
||||
@@ -20352,6 +20375,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &(backend_ctx->prealloc_total_tiles.buffer)));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_use_moe_ragged));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_moe_ragged_skip_gran));
|
||||
|
||||
// set thread grid
|
||||
global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
|
||||
@@ -20527,6 +20552,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &(backend_ctx->prealloc_total_tiles.buffer)));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_use_moe_ragged));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_moe_ragged_skip_gran));
|
||||
|
||||
// set thread grid
|
||||
global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
|
||||
@@ -20710,6 +20737,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &(backend_ctx->prealloc_total_tiles.buffer)));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_use_moe_ragged));
|
||||
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_uint), &backend_ctx->adreno_moe_ragged_skip_gran));
|
||||
|
||||
// set thread grid
|
||||
global_size[1] = static_cast<size_t>((ne01 + 63) / 64);
|
||||
|
||||
@@ -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);
|
||||
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user