Commit 2f2763db6 for llama.cpp
commit 2f2763db6eaac591e78023d1e2c4a31607c64517
Author: Shawn Gu <shawngu@qti.qualcomm.com>
Date: Fri Oct 16 15:15:04 2026 -0700
opencl: Add MxFP4 MoE bin kernel for Adreno X1 devices (#30691)
* opencl: add bin kernel `gemm_moe_mxfp4_q8_1_dp4a_ila_ldck` for X1
* opencl: fix whitespace
* opencl: undo unnecessary change
* opencl: enable dp4a clc for X1E
---------
Co-authored-by: Li He <lih@qti.qualcomm.com>
diff --git a/ggml/src/ggml-opencl/CMakeLists.txt b/ggml/src/ggml-opencl/CMakeLists.txt
index 97b862aff..5ba63dade 100644
--- a/ggml/src/ggml-opencl/CMakeLists.txt
+++ b/ggml/src/ggml-opencl/CMakeLists.txt
@@ -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
diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index 17ceacc7a..de93b0134 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -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
@@ -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
@@ -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
@@ -4903,6 +4909,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
@@ -4984,6 +5008,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
@@ -28258,11 +28299,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;
}
@@ -28351,6 +28400,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
@@ -28375,13 +28430,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));
diff --git a/ggml/src/ggml-opencl/kernels/transpose_activation_q8.cl b/ggml/src/ggml-opencl/kernels/transpose_activation_q8.cl
new file mode 100644
index 000000000..34d54b57e
--- /dev/null
+++ b/ggml/src/ggml-opencl/kernels/transpose_activation_q8.cl
@@ -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]);
+ }
+
+}