wmantly/llama.cpp-xrip
1
0 Fork 0
Code Issues Pull Requests Actions 29 Packages Projects Releases Wiki Activity

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

Browse Source 
* 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>
This commit is contained in:
Shawn GuandLi He authored and GitHub committed 2026-10-16 15:15:04 -07:00
1 parent 01fc3aae84
commit 2f2763db6e
3 files changed
+115 -4
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

ggml/src/ggml-opencl/CMakeLists.txt
+1
View File
@@ -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
ggml/src/ggml-opencl/ggml-opencl.cpp
+79 -4
View File
@@ -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));
ggml/src/ggml-opencl/kernels/transpose_activation_q8.cl
+35
View File
@@ -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]);
}
}