From 0ec4027a2a0183de87c007069adddc06e665f7de Mon Sep 17 00:00:00 2001 From: shaofeiqi Date: Fri, 30 Oct 2026 22:13:48 -0700 Subject: [PATCH] opencl: add fwht (#31434) Co-authored-by: Li He --- docs/ops.md | 2 +- docs/ops/OpenCL.csv | 36 +++++------ ggml/src/ggml-opencl/CMakeLists.txt | 1 + ggml/src/ggml-opencl/ggml-opencl.cpp | 89 ++++++++++++++++++++++++++++ ggml/src/ggml-opencl/kernels/fwht.cl | 44 ++++++++++++++ 5 files changed, 153 insertions(+), 19 deletions(-) create mode 100644 ggml/src/ggml-opencl/kernels/fwht.cl 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::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]; + } +}