Commit 4b77dee05 for llama.cpp
commit 4b77dee0524c64b8471869232e2958bd9589607d
Author: Shawn Gu <shawngu@qti.qualcomm.com>
Date: Thu Nov 5 21:11:07 2026 -0800
opencl: add q6_k multi column tiled GEMV for speculative verification and refactoring (#31758)
* opencl: use the multi-column q6_K vocab head GEMV for 2..4 tokens on X2E
Co-authored-by: Shangqing Gu <shawngu@qti.qualcomm.com>
Assisted-by: Claude Opus 5.5
* opencl: rename GGML_OPENCL_Q6K_MC3 to GGML_OPENCL_Q6_K_MC3
Co-authored-by: Shangqing Gu <shawngu@qti.qualcomm.com>
Assisted-by: Claude Opus 5.5
* opencl: rename GGML_OPENCL_Q40/Q41/Q4K_MC3 to GGML_OPENCL_Q4_0/Q4_1/Q4_K_MC3
GGML_OPENCL_Q4K_MC3_DQ and GGML_OPENCL_Q4K_MC3_LDS become GGML_OPENCL_Q4_K_MC3_DQ
and GGML_OPENCL_Q4_K_MC3_LDS.
Co-authored-by: Shangqing Gu <shawngu@qti.qualcomm.com>
Assisted-by: Claude Opus 5.5
* opencl: remove unnecessary debug probe
Co-authored-by: Shangqing Gu <shawngu@qti.qualcomm.com>
Assisted-by: Claude Opus 5.5
* opencl: clean up
---------
Co-authored-by: Li He <lih@qti.qualcomm.com>
diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index 983e61cca..33457f990 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -1423,7 +1423,7 @@ struct ggml_backend_opencl_context {
cl_kernel kernel_gemv_noshuffle_q6_K_f32_o4;
cl_kernel kernel_gemv_noshuffle_q6_K_f32_o4_global; // weights via __global (opt-in)
cl_kernel kernel_gemv_noshuffle_q6_K_f32_tiled; // tiled-wide layout (opt-in)
- cl_kernel kernel_gemv_noshuffle_q6_K_f32_tiled_mc3; // tiled multi-column (N=3) verify lm_head
+ cl_kernel kernel_gemv_noshuffle_q6_K_f32_tiled_mc3; // tiled multi-column (N 2..4) verify lm_head
cl_kernel kernel_gemm_noshuffle_q6_K_f32_tiled; // batched (N>1) over the tiled layout
cl_kernel kernel_convert_block_q6_k_tiled_ns; // tiled-wide convert (opt-in)
cl_kernel kernel_gemv_noshuffle_q6_K_f32_mc3; // multi-column (N=3) verify GEMV
@@ -6448,16 +6448,13 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
if (backend_ctx->adreno_gen == ADRENO_GPU_GEN::A6X) {
CL_gemv_compile_opts += " -DGGML_CL_A6X_CONSTFOLD_FIX";
}
- // Opt-in: dequant-once-per-block mc3 verify GEMV (factors q4_K dequant
- // out of the 3-column loop; byte-identical, lower spill). A/B vs the
- // shipped inline mc3 in the same binary.
- if (getenv("GGML_OPENCL_Q4K_MC3_DQ")) {
+ // dequant-once-per-block mc3 verify GEMV
+ if (getenv("GGML_OPENCL_Q4_K_MC3_DQ")) {
CL_gemv_compile_opts += " -DQ4K_MC3_DEQUANT_ONCE ";
}
- // Opt-in: LDS-staged dequant mc3 verify GEMV (stages the dequantized
- // q4_K weights in __local instead of private regs that spill to slow
- // global on Adreno; byte-identical). A/B vs inline + dequant-once.
- if (getenv("GGML_OPENCL_Q4K_MC3_LDS")) {
+ // LM-staged dequant mc3 verify GEMV, stages the dequantized
+ // q4_K weights in __local instead of regs
+ if (getenv("GGML_OPENCL_Q4_K_MC3_LDS")) {
CL_gemv_compile_opts += " -DQ4K_MC3_DEQUANT_LDS ";
}
@@ -7379,17 +7376,6 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q6_K_f32 = clCreateKernel(prog, "kernel_gemv_noshuffle_q6_K_f32", &err), err));
CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q6_K_f32_mc3 = clCreateKernel(prog, "kernel_gemv_noshuffle_q6_K_f32_mc3", &err), err));
- if (getenv("GGML_OPENCL_MC3_PROBE")) {
- cl_ulong pm6 = 0, pm4 = 0; size_t wg6 = 0, wg4 = 0, mult = 0;
- clGetKernelWorkGroupInfo(backend_ctx->kernel_gemv_noshuffle_q6_K_f32_mc3, backend_ctx->device, CL_KERNEL_PRIVATE_MEM_SIZE, sizeof(pm6), &pm6, NULL);
- clGetKernelWorkGroupInfo(backend_ctx->kernel_gemv_noshuffle_q6_K_f32_mc3, backend_ctx->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(wg6), &wg6, NULL);
- clGetKernelWorkGroupInfo(backend_ctx->kernel_gemv_noshuffle_q4_k_f32_mc3, backend_ctx->device, CL_KERNEL_PRIVATE_MEM_SIZE, sizeof(pm4), &pm4, NULL);
- clGetKernelWorkGroupInfo(backend_ctx->kernel_gemv_noshuffle_q4_k_f32_mc3, backend_ctx->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(wg4), &wg4, NULL);
- clGetKernelWorkGroupInfo(backend_ctx->kernel_gemv_noshuffle_q6_K_f32_mc3, backend_ctx->device, CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(mult), &mult, NULL);
- fprintf(stderr, "[MC3-PROBE] q4K_mc3 private=%llu wg_cap=%zu | q6K_mc3 private=%llu wg_cap=%zu | pref_mult=%zu\n",
- (unsigned long long)pm4, wg4, (unsigned long long)pm6, wg6, mult);
- fflush(stderr);
- }
GGML_LOG_CONT(".");
}
@@ -24715,14 +24701,8 @@ static void ggml_cl_mul_mat_q4_0_f32_adreno(ggml_backend_t backend, const ggml_t
int N = ne1;
int K = ne00;
- // Multi-column (N=3) verify GEMV for q4_0: route the spec/MTP verify batch
- // (ne1==3) onto the efficient GEMV path instead of the transposed-GEMM dead-
- // zone (gemm_noshuffle_q4_0 is ~50% of MTP decode on a Q4_0 model since q4_0
- // weights have no cok/mc3, unlike q4_K). Reuses the ne1==1 GEMV image setup
- // (activation image already sized by N=ne1). Byte-identical. Opt-in via
- // GGML_OPENCL_Q40_MC3=1. Per-layer only (ne01 < 32768); q4_0 lm_head doesn't
- // occur (token_embd/output stay Q6_K), guard kept for parity with q4_K mc3.
- static const bool q40_mc3 = (getenv("GGML_OPENCL_Q40_MC3") != nullptr);
+ // Multi-column (N=3) verify GEMV for q4_0
+ static const bool q40_mc3 = (getenv("GGML_OPENCL_Q4_0_MC3") != nullptr);
const bool use_q40_mc3 = q40_mc3 && (ne1 >= 2 && ne1 <= 4) && (ne01 < 32768);
const bool use_bin = use_q4_0_bin_kernels(backend_ctx, src0);
@@ -24731,7 +24711,7 @@ static void ggml_cl_mul_mat_q4_0_f32_adreno(ggml_backend_t backend, const ggml_t
if (use_q40_mc3) {
static bool warned = false;
if (!warned) {
- GGML_LOG_WARN("ggml_opencl: GGML_OPENCL_Q40_MC3 is bypassed by Q4_0 binary kernels\n");
+ GGML_LOG_WARN("ggml_opencl: GGML_OPENCL_Q4_0_MC3 is bypassed by Q4_0 binary kernels\n");
warned = true;
}
}
@@ -25109,11 +25089,8 @@ static void ggml_cl_mul_mat_q4_1_f32_adreno(ggml_backend_t backend, const ggml_t
int N = ne1;
int K = ne00;
- // Multi-column (N=3) verify GEMV for q4_1: route the spec/MTP verify batch
- // (ne1==3) onto the efficient GEMV path instead of the transposed-GEMM dead-
- // zone (gemm_noshuffle_q4_1). Reuses the ne1==1 GEMV image setup. Opt-in via
- // GGML_OPENCL_Q41_MC3=1. Per-layer only (ne01 < 32768).
- static const bool q41_mc3 = (getenv("GGML_OPENCL_Q41_MC3") != nullptr);
+ // Multi-column (N=3) verify GEMV for q4_1
+ static const bool q41_mc3 = (getenv("GGML_OPENCL_Q4_1_MC3") != nullptr);
const bool use_q41_mc3 = q41_mc3 && (ne1 >= 2 && ne1 <= 4) && (ne01 < 32768);
if (ne1 == 1 || use_q41_mc3) {
@@ -26930,15 +26907,8 @@ static void ggml_cl_mul_mat_q4_k_f32_adreno(ggml_backend_t backend, const ggml_t
cl_uchar mask_d4 = 0x0F;
cl_uchar mask_hi2 = 0xC0;
- // Multi-column verify GEMV: route the spec/MTP verify batch (ne1==3 = 2
- // drafts + 1 bonus) onto the efficient GEMV path (subgroup-broadcast, no
- // transpose) instead of the transposed-GEMM dead-zone. Reuses the ne1==1
- // GEMV setup (the activation image is already sized by N=ne1). Byte-
- // identical. Opt-in via GGML_OPENCL_Q4K_MC3=1 while validating.
- static const bool q4k_mc3 = (getenv("GGML_OPENCL_Q4K_MC3") != nullptr);
- // Per-layer only (ne01 < 32768): the batched large-vocab lm_head at ne1==3
- // is left to the existing routing (corrupts on the Adreno GEMV path; x2-
- // unified routes batched Q6_K lm_head to CPU). Per-layer mc3 is byte-identical.
+ // Multi-column verify GEMV
+ static const bool q4k_mc3 = (getenv("GGML_OPENCL_Q4_K_MC3") != nullptr);
const bool use_mc3 = q4k_mc3 && (ne1 == 3) && (ne01 < 32768);
const bool use_bin = use_q4_k_bin_kernels(backend_ctx, src0);
@@ -26947,7 +26917,7 @@ static void ggml_cl_mul_mat_q4_k_f32_adreno(ggml_backend_t backend, const ggml_t
if (use_mc3) {
static bool warned = false;
if (!warned) {
- GGML_LOG_WARN("ggml_opencl: GGML_OPENCL_Q4K_MC3 is bypassed by Q4_K binary kernels\n");
+ GGML_LOG_WARN("ggml_opencl: GGML_OPENCL_Q4_K_MC3 is bypassed by Q4_K binary kernels\n");
warned = true;
}
}
@@ -27772,20 +27742,15 @@ static void ggml_cl_mul_mat_q6_K_f32_adreno(ggml_backend_t backend, const ggml_t
cl_image_format img_fmt;
cl_image_desc img_desc;
- // subbuffer and image for activation
- // Multi-column verify GEMV: route the spec/MTP verify q6_K matmuls (ne1==3)
- // onto the efficient GEMV path instead of the transposed-GEMM dead-zone.
- // Reuses the ne1==1 image setup (activation image sized by N=ne1). Byte-
- // identical. Opt-in via GGML_OPENCL_Q6K_MC3=1 while validating.
- static const bool q6k_mc3 = (getenv("GGML_OPENCL_Q6K_MC3") != nullptr);
- // Per-layer only (ne01 < 32768): batched large-vocab lm_head stays on the
- // existing path (x2-unified routes batched Q6_K lm_head to CPU; the Adreno
- // GEMV corrupts it). Per-layer mc3 is byte-identical.
+ // Multi-column verify GEMV
+ static const bool q6k_mc3 = (getenv("GGML_OPENCL_Q6_K_MC3") != nullptr);
const bool use_q6k_mc3 = q6k_mc3 && (ne1 == 3) && (ne01 < 32768);
- // Batched verify lm_head/embed (ne1==3, tiled layout): multi-column tiled
- // GEMV — streams the large lm_head weight once across the 3 verify columns
- // (the #1 MTP bottleneck; mc3 above can't, it reads the noshuffle layout).
- const bool use_q6k_tiled_mc = q6k_mc3 && (ne1 == 3) && (ne01 >= 32768) && use_q6k_tiled(backend_ctx, src0);
+ // Multi-column tiled vocab GEMV for verify
+ // streams the large lm_head weight once across the verify columns
+ static const char * q6k_tiled_mc_env = getenv("GGML_OPENCL_Q6_K_TILED_MC");
+ const bool q6k_tiled_mc = q6k_tiled_mc_env ? atoi(q6k_tiled_mc_env) != 0
+ : (q6k_mc3 || backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E);
+ const bool use_q6k_tiled_mc = q6k_tiled_mc && (ne1 >= 2 && ne1 <= 4) && (ne01 >= 32768) && use_q6k_tiled(backend_ctx, src0);
const bool use_bin = use_q6_k_bin_kernels(backend_ctx, src0);
@@ -27793,7 +27758,7 @@ static void ggml_cl_mul_mat_q6_K_f32_adreno(ggml_backend_t backend, const ggml_t
if (use_q6k_mc3 || use_q6k_tiled_mc) {
static bool warned = false;
if (!warned) {
- GGML_LOG_WARN("ggml_opencl: GGML_OPENCL_Q6K_MC3 is bypassed by Q6_K binary kernels\n");
+ GGML_LOG_WARN("ggml_opencl: GGML_OPENCL_Q6_K_MC3 is bypassed by Q6_K binary kernels\n");
warned = true;
}
}
@@ -27881,6 +27846,9 @@ static void ggml_cl_mul_mat_q6_K_f32_adreno(ggml_backend_t backend, const ggml_t
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_ulong), &offsetd));
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_int), &ne00));
CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_int), &ne01));
+ if (use_q6k_tiled_mc) {
+ CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_int), &ne1));
+ }
const size_t gws_x = use_tiled
? (size_t) CEIL_DIV(ne01, 64) * 64
diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q6_k_f32_tiled.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q6_k_f32_tiled.cl
index c5049f396..a9ddc1130 100644
--- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q6_k_f32_tiled.cl
+++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q6_k_f32_tiled.cl
@@ -104,11 +104,11 @@ kernel void kernel_gemv_noshuffle_q6_K_f32_tiled(
}
}
-// Multi-column (N=3) variant of the tiled q6_K decode GEMV, for the speculative/
-// MTP VERIFY lm_head/embed (ne1=3 = 2 drafts + 1 bonus). Identical tiled weight
-// layout + unpack as the ne1=1 kernel above; each WI computes 3 output columns,
+// Multi-column (ne1 in [2..4]) variant of the tiled q6_K decode GEMV, for the
+// speculative/MTP VERIFY lm_head/embed (drafts + 1 bonus). Identical tiled weight
+// layout + unpack as the ne1=1 kernel above; each WI computes up to 4 output columns,
// streaming the (large) lm_head weight ONCE per superblock and reusing it across
-// the 3 verify activation columns (dequant once per code, MAC into 3 accs). This
+// the verify activation columns (dequant once per code, MAC into 4 accs). This
// is the lm_head analogue of the per-layer mc3 GEMV; the multiply order matches
// the ne1=1 kernel, so each column is byte-identical to a standalone tiled GEMV.
#if defined(ADRENO_GPU)
@@ -123,7 +123,8 @@ kernel void kernel_gemv_noshuffle_q6_K_f32_tiled_mc3(
global float * dst,
ulong offsetd,
int ne00,
- int ne01
+ int ne01,
+ int ne1
) {
int grp = get_local_id(1);
int row = get_global_id(0);
@@ -133,7 +134,12 @@ kernel void kernel_gemv_noshuffle_q6_K_f32_tiled_mc3(
int nb = ne00 / 256;
int col_stride = ne00 / 4; // activation float4 pixels per column
- float acc0 = 0.0f, acc1 = 0.0f, acc2 = 0.0f;
+ // columns past ne1 re-read the last column and are not stored
+ int col1 = min(1, ne1 - 1) * col_stride;
+ int col2 = min(2, ne1 - 1) * col_stride;
+ int col3 = min(3, ne1 - 1) * col_stride;
+
+ float acc0 = 0.0f, acc1 = 0.0f, acc2 = 0.0f, acc3 = 0.0f;
for (int sb = grp; sb < nb; sb += NSUBGROUPS) {
int tile_blk = rt * nb + sb;
@@ -157,9 +163,10 @@ kernel void kernel_gemv_noshuffle_q6_K_f32_tiled_mc3(
int act_base = sb * 64;
#pragma unroll
for (int e4 = 0; e4 < 64; ++e4) {
- float4 a0 = read_imagef(src1, 0*col_stride + act_base + e4);
- float4 a1 = read_imagef(src1, 1*col_stride + act_base + e4);
- float4 a2 = read_imagef(src1, 2*col_stride + act_base + e4);
+ float4 a0 = read_imagef(src1, act_base + e4);
+ float4 a1 = read_imagef(src1, col1 + act_base + e4);
+ float4 a2 = read_imagef(src1, col2 + act_base + e4);
+ float4 a3 = read_imagef(src1, col3 + act_base + e4);
#pragma unroll
for (int t = 0; t < 4; ++t) {
int e = e4 * 4 + t;
@@ -171,15 +178,17 @@ kernel void kernel_gemv_noshuffle_q6_K_f32_tiled_mc3(
float av0 = (t == 0) ? a0.x : (t == 1) ? a0.y : (t == 2) ? a0.z : a0.w;
float av1 = (t == 0) ? a1.x : (t == 1) ? a1.y : (t == 2) ? a1.z : a1.w;
float av2 = (t == 0) ? a2.x : (t == 1) ? a2.y : (t == 2) ? a2.z : a2.w;
+ float av3 = (t == 0) ? a3.x : (t == 1) ? a3.y : (t == 2) ? a3.z : a3.w;
acc0 += w * av0;
acc1 += w * av1;
acc2 += w * av2;
+ acc3 += w * av3;
}
}
}
local float4 reduce_lm[NSUBGROUPS * TILE_ROWS];
- reduce_lm[grp * TILE_ROWS + rit] = (float4)(acc0, acc1, acc2, 0.0f);
+ reduce_lm[grp * TILE_ROWS + rit] = (float4)(acc0, acc1, acc2, acc3);
barrier(CLK_LOCAL_MEM_FENCE);
if (grp == 0) {
@@ -188,9 +197,10 @@ kernel void kernel_gemv_noshuffle_q6_K_f32_tiled_mc3(
+ reduce_lm[2 * TILE_ROWS + rit]
+ reduce_lm[3 * TILE_ROWS + rit];
dst = (global float*)((global char*)dst + offsetd);
- // dst column-major [ne01 rows x 3 cols]: (row, col) at col*ne01 + row
+ // dst column-major [ne01 rows x ne1 cols]: (row, col) at col*ne01 + row
dst[0*ne01 + row] = total.x;
dst[1*ne01 + row] = total.y;
- dst[2*ne01 + row] = total.z;
+ if (ne1 > 2) dst[2*ne01 + row] = total.z;
+ if (ne1 > 3) dst[3*ne01 + row] = total.w;
}
}