Skip to content

opencl: Add MxFP4 MoE bin kernel for Adreno X1 devices #30691

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
max-krasnyansky merged 4 commits into from
Oct 16, 2026
+115 −4
Merged

opencl: Add MxFP4 MoE bin kernel for Adreno X1 devices #30691

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
ad5ef51
opencl: add bin kernel `gemm_moe_mxfp4_q8_1_dp4a_ila_ldck` for X1
shawngu-quic Oct 15, 2026
042fc56
opencl: fix whitespace
lhez Oct 16, 2026
4f7aa34
opencl: undo unnecessary change
lhez Oct 16, 2026
a7d11d3
opencl: enable dp4a clc for X1E
shawngu-quic Oct 16, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions ggml/src/ggml-opencl/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -132,6 +132,7 @@ set(GGML_OPENCL_KERNELS
gemm_moe_q6_k_q8_1_dp4a
gemm_moe_q8_1_dp4a
moe_reorder_quant_a_q8_1
transpose_activation_q8
gemm_noshuffle_q4_k_q8_1_dp4a
gemm_noshuffle_q5_k_q8_1_dp4a
gemm_noshuffle_q6_k_q8_1_dp4a
Expand Down
83 changes: 79 additions & 4 deletions ggml/src/ggml-opencl/ggml-opencl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -720,6 +720,10 @@ struct ggml_backend_opencl_context {
ggml_cl_buffer prealloc_moe_qa; // int8 quants [tok_slots * ne00]
ggml_cl_buffer prealloc_moe_da; // per-block d [tok_slots * ne00/32] (half)
ggml_cl_buffer prealloc_moe_sa; // per-block s [tok_slots * ne00/32] (half)
// per-tile transposed qa/da and on-chip scratch for the ldck dp4a MoE GEMM
ggml_cl_buffer prealloc_moe_qa_trans; // int8 quants [tok_slots * ne00]
ggml_cl_buffer prealloc_moe_da_trans; // per-block d [tok_slots * ne00/32] (float)
ggml_cl_buffer prealloc_moe_qda_on_chip;
// scratch copy of the router weights to avoid dst aliasing
ggml_cl_buffer prealloc_moe_combine_w;
ggml_cl_buffer prealloc_splitk_partial; // [ksplit * M] partials for split-K GEMV
Expand Down Expand Up @@ -1016,6 +1020,7 @@ struct ggml_backend_opencl_context {
cl_kernel kernel_gemv_moe_q4_k_f32_ns_wimg = nullptr; // weight-as-texture MoE decode GEMV (opt-in)
cl_kernel kernel_gemm_moe_q4_k_q8_1_dp4a = nullptr; // dp4a (int8) prefill GEMM variant
cl_kernel kernel_moe_reorder_quant_a_q8_1; // fused reorder + q8_1 quant for the dp4a GEMM
cl_kernel kernel_transpose_activation_q8; // per-tile transpose of q8_1 activations for the ldck dp4a GEMM
cl_kernel kernel_gemm_moe_q8_1_dp4a_q80 = nullptr; // generic dp4a MoE GEMM (MOE_QT=80), opt-in
cl_kernel kernel_moe_expand_scale_q8_0 = nullptr; // q8_0 per-block d -> uniform scale[16]
cl_kernel kernel_gemm_moe_q8_1_dp4a_q50 = nullptr; // generic dp4a MoE GEMM (MOE_QT=50, q5_0), opt-in
Expand All @@ -1031,6 +1036,7 @@ struct ggml_backend_opencl_context {
cl_kernel kernel_gemm_moe_mxfp4_q8_1_dp4a = nullptr; // dp4a (int8) mxfp4 MoE prefill GEMM
cl_kernel kernel_gemm_moe_q4_0_q8_1_dp4a = nullptr; // dp4a (int8) q4_0 MoE prefill GEMM
cl_kernel kernel_gemm_moe_mxfp4_q8_1_dp4a_bin = nullptr; // binary dp4a (int8) mxfp4 MoE prefill GEMM
cl_kernel kernel_gemm_moe_mxfp4_q8_1_dp4a_ldck_bin = nullptr; // binary dp4a (int8) mxfp4 MoE prefill GEMM using ldck
cl_kernel kernel_gemm_moe_q4_0_q8_1_dp4a_bin = nullptr; // binary dp4a (int8) q4_0 MoE prefill GEMM
cl_kernel kernel_gemm_moe_q4_k_q8_1_dp4a_bin = nullptr; // binary dp4a (int8) q4_k MoE prefill GEMM
cl_kernel kernel_gemm_moe_q6_k_q8_1_dp4a_bin = nullptr; // binary dp4a (int8) q6_k MoE prefill GEMM
Expand Down Expand Up @@ -4902,6 +4908,24 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
}
}

// gemm_moe_mxfp4_q8_1_dp4a_ldck_bin (dp4a prefill GEMM)
if (backend_ctx->has_integer_dot) {
size_t bin_size = 0;
backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_ldck_bin = nullptr;

if (use_adreno_bin_kernels(backend_ctx)) {
const char * kernel_bin = (const char *)backend_ctx->get_adreno_bin_kernel("gemm_moe_mxfp4_q8_1_dp4a_ila_ldck", &bin_size);
if (kernel_bin && bin_size > 0) {
cl_program prog =
build_program_from_binary(backend_ctx->context, backend_ctx->device, kernel_bin, CL_moe_compile_opts, bin_size);

CL_CHECK((backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_ldck_bin = clCreateKernel(prog, "kernel_gemm_moe_mxfp4_q8_1_dp4a_ila_ldck", &err), err));
CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
}
}
}

// gemm_moe_q4_0_q8_1_dp4a (dp4a prefill GEMM)
if (backend_ctx->has_integer_dot) {
#ifdef GGML_OPENCL_EMBED_KERNELS
Expand Down Expand Up @@ -4983,6 +5007,23 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
GGML_LOG_CONT(".");
}

// transpose_activation_q8 (per-tile transpose of q8_1 activations for the ldck dp4a GEMM)
{
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "transpose_activation_q8.cl.h"
};
#else
const std::string kernel_src = read_file("transpose_activation_q8.cl");
#endif
cl_program prog =
build_program_from_source(backend_ctx, kernel_src.c_str(), CL_moe_compile_opts);

CL_CHECK((backend_ctx->kernel_transpose_activation_q8 = clCreateKernel(prog, "kernel_transpose_activation_q8", &err), err));
CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
}

// gemv_moe_q5_k_f32_ns
{
#ifdef GGML_OPENCL_EMBED_KERNELS
Expand Down Expand Up @@ -28217,11 +28258,19 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
static const char * mxfp4_moe_dp4a_env = getenv("GGML_OPENCL_MXFP4_MOE_DP4A");
bool use_moe_dp4a = mxfp4_moe_dp4a_env
? (atoi(mxfp4_moe_dp4a_env) != 0)
: (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E);
: (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E || backend_ctx->adreno_gen == ADRENO_GPU_GEN::X1E);
// dot prod has to be available
use_moe_dp4a = backend_ctx->has_integer_dot && use_moe_dp4a;
// bin kernel takes precedence, dp4a bin kernel has higher priority than normal bin kernel
if (backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_bin == nullptr) {
// ldck dp4a bin kernel (X1) has the same priority as normal bin kernel (X2/A8x), they are not expected to coexist.
const bool use_moe_dp4a_ldck = backend_ctx->has_integer_dot &&
(!mxfp4_moe_dp4a_env || atoi(mxfp4_moe_dp4a_env) != 0) &&
backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_bin == nullptr &&
backend_ctx->kernel_gemm_moe_mxfp4_f32_ns_bin == nullptr &&
backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_ldck_bin != nullptr;
if (use_moe_dp4a_ldck) {
use_moe_dp4a = true;
} else if (backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_bin == nullptr) {
use_moe_dp4a = use_moe_dp4a && backend_ctx->kernel_gemm_moe_mxfp4_f32_ns_bin == nullptr;
}

Expand Down Expand Up @@ -28310,6 +28359,12 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
backend_ctx->prealloc_moe_qa.allocate(backend_ctx->context, tok_slots * ne00 * sizeof(cl_char));
backend_ctx->prealloc_moe_da.allocate(backend_ctx->context, n_blocks * sizeof(cl_half));
backend_ctx->prealloc_moe_sa.allocate(backend_ctx->context, n_blocks * sizeof(cl_half));
if (use_moe_dp4a_ldck) {
backend_ctx->prealloc_moe_qa_trans.allocate(backend_ctx->context, tok_slots * ne00 * sizeof(cl_char));
backend_ctx->prealloc_moe_da_trans.allocate(backend_ctx->context, n_blocks * sizeof(cl_float));
// matches max_constant_size of src1_a_cm in the ldck kernel
backend_ctx->prealloc_moe_qda_on_chip.allocate(backend_ctx->context, 2304);
}

// fused reorder + q8_1 quant straight from the original
// activations (no intermediate f32 reorder buffer). mxfp4 has no
Expand All @@ -28334,13 +28389,33 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
cl_kernel dk = backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a;
if (backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_bin) {
dk = backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_bin;
} else if (use_moe_dp4a_ldck) {
dk = backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_ldck_bin;

// per-tile transpose of qa/da for the ldck kernel
cl_kernel tk = backend_ctx->kernel_transpose_activation_q8;
CL_CHECK(clSetKernelArg(tk, 0, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer));
CL_CHECK(clSetKernelArg(tk, 1, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa_trans.buffer));
CL_CHECK(clSetKernelArg(tk, 2, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer));
CL_CHECK(clSetKernelArg(tk, 3, sizeof(cl_mem), &backend_ctx->prealloc_moe_da_trans.buffer));
CL_CHECK(clSetKernelArg(tk, 4, sizeof(cl_mem), &(backend_ctx->prealloc_total_tiles.buffer)));
CL_CHECK(clSetKernelArg(tk, 5, sizeof(cl_uint), &ne00));
size_t tk_local[3] = { 32, 8, 1 };
size_t tk_global[3] = { 32, (size_t)(ne00 / 4), (size_t)max_post_router_tile };
backend_ctx->enqueue_ndrange_kernel(tk, 3, tk_global, tk_local, dst);
}

int aidx = 0;
CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &extra0_mxfp4->q_img));
CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &extra0_mxfp4->e));
CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer));
CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer));
if (use_moe_dp4a_ldck) {
CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa_trans.buffer));
CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &backend_ctx->prealloc_moe_da_trans.buffer));
CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &backend_ctx->prealloc_moe_qda_on_chip.buffer));
} else {
CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer));
CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer));
}
CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &buf_src2));
CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &buf_src2_emap));
CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &buf_dst_image));
Expand Down
35 changes: 35 additions & 0 deletions ggml/src/ggml-opencl/kernels/transpose_activation_q8.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,35 @@
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#pragma OPENCL EXTENSION cl_khr_subgroups : enable

#define tile_size 32

kernel void kernel_transpose_activation_q8(
__global uint * src1_qa,
__global uint * src1_qa_trans,
__global half * src1_sa,
__global float * src1_sa_trans,
__global const int * total_tiles,
uint K
) {
uint row_in_tile = get_global_id(0);
uint blk_k = get_group_id(1);
uint col_in_tile = get_local_id(1);
uint tile_n = get_global_id(2);

if (blk_k >= (K / tile_size) || tile_n >= (uint)total_tiles[0]) {
return;
}

uint src_idx = (tile_n * tile_size + row_in_tile) * K / 4 + blk_k * 8 + col_in_tile;
uint dst_idx = tile_n * tile_size * K / 4 + (blk_k * 8 + col_in_tile) * tile_size + row_in_tile;

src1_qa_trans[dst_idx] = src1_qa[src_idx];

if (col_in_tile == 0) {
src_idx = (tile_n * tile_size + row_in_tile) * K / 32 + blk_k;
dst_idx = tile_n * tile_size * K / 32 + blk_k * tile_size + row_in_tile;

src1_sa_trans[dst_idx] = (float)(src1_sa[src_idx]);
}

}
Loading