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

opencl: add fwht (#31434)

Browse Source 
Co-authored-by: Li He <lih@qti.qualcomm.com>
This commit is contained in:
shaofeiqiandLi He authored and GitHub committed 2026-10-30 22:13:48 -07:00
1 parent 1052f7e1d6
commit 0ec4027a2a
5 files changed
+153 -19
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

docs/ops.md
+1 -1
View File
@@ -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 | ❌ | ❌ | ✅ | ❌ | ❌ | 🟡 | ❌ | 🟡 | ✅ | ❌ | ❌ | ❌ | ❌ |
docs/ops/OpenCL.csv
+18 -18
View File
@@ -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"
Can't render this file because it is too large.
ggml/src/ggml-opencl/CMakeLists.txt
+1
View File
@@ -70,6 +70,7 @@ set(GGML_OPENCL_KERNELS
lightning_indexer
tri
fill
fwht
clamp
cpy
cvt
ggml/src/ggml-opencl/ggml-opencl.cpp
+89
View File
@@ -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)
ggml/src/ggml-opencl/kernels/fwht.cl
+44
View File
@@ -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];
}
}