Commit 0ec4027a2 for llama.cpp
commit 0ec4027a2a0183de87c007069adddc06e665f7de
Author: shaofeiqi <shaoqi@qti.qualcomm.com>
Date: Fri Oct 30 22:13:48 2026 -0700
opencl: add fwht (#31434)
Co-authored-by: Li He <lih@qti.qualcomm.com>
diff --git a/docs/ops.md b/docs/ops.md
index 832a967bf..2885441c5 100644
--- a/docs/ops.md
+++ b/docs/ops.md
@@ -74,7 +74,7 @@ Legend:
| MEAN | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | 🟡 | ✅ | ❌ | ❌ | ❌ |
| MUL | ❌ | ✅ | ✅ | ✅ | 🟡 | 🟡 | 🟡 | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
| MUL_MAT | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 |
-| MUL_MAT_HADAMARD | ❌ | ❌ | ✅ | ❌ | ✅ | ✅ | ❌ | 🟡 | ✅ | ✅ | ✅ | ❌ | 🟡 |
+| MUL_MAT_HADAMARD | ❌ | ❌ | ✅ | ❌ | ✅ | ✅ | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | 🟡 |
| MUL_MAT_ID | ❌ | 🟡 | ✅ | ✅ | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | ✅ | 🟡 | 🟡 | ❌ |
| MUL_MAT_ID_W4A4 | ❌ | ❌ | ✅ | ❌ | ❌ | 🟡 | ❌ | 🟡 | ✅ | ❌ | ❌ | ❌ | ❌ |
| MUL_MAT_ID_W4A8 | ❌ | ❌ | ✅ | ❌ | ❌ | 🟡 | ❌ | 🟡 | ✅ | ❌ | ❌ | ❌ | ❌ |
diff --git a/docs/ops/OpenCL.csv b/docs/ops/OpenCL.csv
index ad786cfee..db66c5655 100644
--- a/docs/ops/OpenCL.csv
+++ b/docs/ops/OpenCL.csv
@@ -9242,31 +9242,31 @@
"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=128,n=32,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=128,n=4,k=128,bs=[2,3],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=256,n=512,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=64,n=1,k=64,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=1,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=512,n=1,k=512,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=32,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=512,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=64,n=1,k=64,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=1,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=512,n=1,k=512,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=32,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=512,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=32,n=1,k=32,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=16384,n=1,k=16384,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=64,n=1,k=64,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=1,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=512,n=1,k=512,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=32,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=4,k=128,bs=[2,3],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=512,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=64,n=1,k=64,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=1,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=512,n=1,k=512,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=32,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=128,n=4,k=128,bs=[2,3],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=256,n=512,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=1024,n=1,k=1024,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=2048,n=1,k=2048,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=4096,n=1,k=4096,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=8192,n=1,k=8192,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f32,m=1024,n=7,k=1024,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=1024,n=1,k=1024,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=2048,n=1,k=2048,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=4096,n=1,k=4096,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=8192,n=1,k=8192,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
-"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=1024,n=7,k=1024,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=1024,n=1,k=1024,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=2048,n=1,k=2048,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=4096,n=1,k=4096,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=8192,n=1,k=8192,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
+"OpenCL","MUL_MAT_HADAMARD","type_a=f32,type_b=f16,m=1024,n=7,k=1024,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","1","yes","OpenCL"
"OpenCL","MUL_MAT_W4A8","type_a=nvfp4,type_b=f32,m=32,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
"OpenCL","MUL_MAT_W4A8","type_a=nvfp4,type_b=f32,m=32,n=32,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
"OpenCL","MUL_MAT_W4A8","type_a=nvfp4,type_b=f32,m=64,n=16,k=512,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0,m_v=0,pad=0","support","0","no","OpenCL"
diff --git a/ggml/src/ggml-opencl/CMakeLists.txt b/ggml/src/ggml-opencl/CMakeLists.txt
index 3f4afac41..046e42acd 100644
--- a/ggml/src/ggml-opencl/CMakeLists.txt
+++ b/ggml/src/ggml-opencl/CMakeLists.txt
@@ -70,6 +70,7 @@ set(GGML_OPENCL_KERNELS
lightning_indexer
tri
fill
+ fwht
clamp
cpy
cvt
diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index 179f04260..5b05c9b36 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -661,6 +661,7 @@ struct ggml_backend_opencl_context {
size_t global_mem_size;
size_t max_alloc_size;
size_t max_workgroup_size;
+ cl_ulong local_mem_size;
bool fp16_support;
bool has_vector_subgroup_broadcast;
bool has_khr_subgroup_shuffle = false;
@@ -885,6 +886,8 @@ struct ggml_backend_opencl_context {
cl_kernel kernel_cpy_f32_f32_pack = nullptr;
cl_kernel kernel_cpy_f32_f32_flat = nullptr;
cl_kernel kernel_mul_mat_f32_f32;
+ cl_kernel kernel_fwht[2] = {};
+ size_t fwht_workgroup_size[2] = {};
cl_kernel kernel_mul_mat_f16_f16;
cl_kernel kernel_mul_mat_f16_f32_1row;
cl_program program_mul_mv_f16_f32_mrow;
@@ -2724,6 +2727,30 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
GGML_LOG_CONT(".");
}
+ // fwht
+ {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+ const std::string kernel_src {
+ #include "fwht.cl.h"
+ };
+#else
+ const std::string kernel_src = read_file("fwht.cl");
+#endif
+ for (int i = 0; i < 2; ++i) {
+ const std::string opts = compile_opts + (i == 0 ? " -DSRC_TYPE=float" : " -DSRC_TYPE=half");
+ cl_program prog = build_program_from_source(backend_ctx, kernel_src.c_str(), opts);
+ CL_CHECK((backend_ctx->kernel_fwht[i] = clCreateKernel(prog, "kernel_fwht", &err), err));
+ CL_CHECK(clReleaseProgram(prog));
+ size_t nth = 128;
+ const size_t max_workgroup = backend_ctx->get_kernel_workgroup_size(backend_ctx->kernel_fwht[i]);
+ while (nth > max_workgroup) {
+ nth /= 2;
+ }
+ backend_ctx->fwht_workgroup_size[i] = nth;
+ }
+ GGML_LOG_CONT(".");
+ }
+
// mul_mv_q2_0_f32_flat
{
#ifdef GGML_OPENCL_EMBED_KERNELS
@@ -8552,6 +8579,7 @@ static ggml_backend_opencl_context * ggml_cl_init(ggml_backend_dev_t dev) {
CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof(size_t), &backend_ctx->image2d_max_width, NULL));
CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof(size_t), &backend_ctx->image2d_max_height, NULL));
CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), &backend_ctx->max_workgroup_size, NULL));
+ CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &backend_ctx->local_mem_size, NULL));
CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_SVM_CAPABILITIES, sizeof(cl_device_svm_capabilities), &backend_ctx->svm_caps, 0));
if (opencl_c_version.major >= 3) {
@@ -10959,6 +10987,19 @@ inline bool use_q5_k_bin_kernels(const ggml_backend_opencl_context *backend_ctx,
#endif
}
+static bool ggml_cl_mul_mat_can_use_fwht(const ggml_backend_opencl_context * backend_ctx, const ggml_tensor * op) {
+ const ggml_tensor * a = op->src[0];
+ const ggml_tensor * b = op->src[1];
+ const int64_t n = a->ne[0];
+
+ return ggml_get_op_params_i32(op, 1) == GGML_HINT_SRC0_IS_HADAMARD &&
+ a->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 &&
+ (b->type == GGML_TYPE_F32 || b->type == GGML_TYPE_F16) &&
+ n >= 2 && n <= std::numeric_limits<int>::max() && (n & (n - 1)) == 0 && a->ne[1] == n &&
+ (uint64_t)n <= backend_ctx->local_mem_size/sizeof(float) &&
+ ggml_is_contiguous_rows(b) && ggml_is_contiguous(op) && ggml_are_same_shape(b, op);
+}
+
static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_tensor * op) {
ggml_backend_opencl_device_context * dev_ctx = (ggml_backend_opencl_device_context *)dev->context;
ggml_backend_opencl_context * backend_ctx = dev_ctx->backend_ctx;
@@ -11182,6 +11223,9 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
case GGML_OP_GROUP_NORM:
return ggml_is_contiguous(op->src[0]);
case GGML_OP_MUL_MAT:
+ if (ggml_cl_mul_mat_can_use_fwht(backend_ctx, op)) {
+ return true;
+ }
if (op->src[0]->type == GGML_TYPE_F16) {
return true;
} else if (op->src[0]->type == GGML_TYPE_BF16) {
@@ -28627,6 +28671,46 @@ static void ggml_cl_mul_mat_iq2_xxs_f32_adreno(ggml_backend_t backend, const ggm
}
#endif
+static void ggml_cl_fwht(ggml_backend_t backend, const ggml_tensor * src, ggml_tensor * dst) {
+ auto * backend_ctx = (ggml_backend_opencl_context *) backend->context;
+
+ auto * extra_src = (ggml_tensor_extra_cl *) src->extra;
+ auto * extra_dst = (ggml_tensor_extra_cl *) dst->extra;
+
+ const cl_ulong offset_src = extra_src->offset + src->view_offs;
+ const cl_ulong offset_dst = extra_dst->offset + dst->view_offs;
+
+ const cl_ulong ne1 = src->ne[1];
+ const cl_ulong ne2 = src->ne[2];
+ const cl_ulong nb1 = src->nb[1];
+ const cl_ulong nb2 = src->nb[2];
+ const cl_ulong nb3 = src->nb[3];
+
+ const int n = src->ne[0];
+ const float scale = 1.0f/sqrtf((float)n);
+ const int type_index = src->type == GGML_TYPE_F16 ? 1 : 0;
+
+ cl_kernel kernel = backend_ctx->kernel_fwht[type_index];
+ const size_t nth = std::min((size_t)n/2, backend_ctx->fwht_workgroup_size[type_index]);
+
+ CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra_src->data_device));
+ CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset_src));
+ CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra_dst->data_device));
+ CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offset_dst));
+ CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), &n));
+ CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &ne1));
+ CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_ulong), &ne2));
+ CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_ulong), &nb1));
+ CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_ulong), &nb2));
+ CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_ulong), &nb3));
+ CL_CHECK(clSetKernelArg(kernel, 10, sizeof(float), &scale));
+ CL_CHECK(clSetKernelArg(kernel, 11, n*sizeof(float), nullptr));
+
+ size_t global[] = { nth, (size_t)ggml_nrows(src) };
+ size_t local[] = { nth, 1 };
+ backend_ctx->enqueue_ndrange_kernel(kernel, 2, global, local, dst);
+}
+
static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
GGML_ASSERT(src0);
GGML_ASSERT(src0->extra);
@@ -28641,6 +28725,11 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context;
+ if (ggml_cl_mul_mat_can_use_fwht(backend_ctx, dst)) {
+ ggml_cl_fwht(backend, src1, dst);
+ return;
+ }
+
// quant kv without FA
// used for non-contiguous src0 (the usual head-major permuted K view when n_head_kv>1)
// AND for the contiguous case that occurs when n_head_kv==1 (e.g. Gemma-4 E2B)
diff --git a/ggml/src/ggml-opencl/kernels/fwht.cl b/ggml/src/ggml-opencl/kernels/fwht.cl
new file mode 100644
index 000000000..b1fafc34f
--- /dev/null
+++ b/ggml/src/ggml-opencl/kernels/fwht.cl
@@ -0,0 +1,44 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+kernel void kernel_fwht(
+ global const char * src,
+ ulong offset_src,
+ global char * dst,
+ ulong offset_dst,
+ int n,
+ ulong ne1,
+ ulong ne2,
+ ulong nb1,
+ ulong nb2,
+ ulong nb3,
+ float scale,
+ local float * work) {
+ const ulong row = get_group_id(1);
+ const ulong i1 = row % ne1;
+ const ulong i2 = row / ne1 % ne2;
+ const ulong i3 = row / ne1 / ne2;
+ global const SRC_TYPE * x = (global const SRC_TYPE *)(src + offset_src + i1*nb1 + i2*nb2 + i3*nb3);
+ global float * y = (global float *)(dst + offset_dst) + row*n;
+ const int tid = get_local_id(0);
+ const int nth = get_local_size(0);
+
+ for (int i = tid; i < n; i += nth) {
+ work[i] = (float)x[i] * scale;
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ for (int h = 1; h < n; h *= 2) {
+ for (int i = tid; i < n/2; i += nth) {
+ const int j = 2*i - (i & (h - 1));
+ const float a = work[j];
+ const float b = work[j + h];
+ work[j] = a + b;
+ work[j + h] = a - b;
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ }
+
+ for (int i = tid; i < n; i += nth) {
+ y[i] = work[i];
+ }
+}