From a25c9865fe03c954c93fd755b5d79ae86ba99750 Mon Sep 17 00:00:00 2001 From: shaofeiqi Date: Fri, 25 Sep 2026 07:35:18 -0700 Subject: [PATCH] opencl: add bin kernel `kernel_gemm_noshuffle_q5_k_f32_32b_trans_ila_a8_bin`, `kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8_bin` (#29401) * opencl: add A8 Q5_K non-MoE non dp4a + dp4a binary kernel * opencl: fix s transpose - s only transposed for bin kernels --------- Co-authored-by: Li He --- ggml/src/ggml-opencl/CMakeLists.txt | 1 + ggml/src/ggml-opencl/ggml-opencl.cpp | 329 +++++++++++++++++- .../gemv_noshuffle_q5_k_f32_32b_trans.cl | 141 ++++++++ 3 files changed, 466 insertions(+), 5 deletions(-) create mode 100644 ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_k_f32_32b_trans.cl diff --git a/ggml/src/ggml-opencl/CMakeLists.txt b/ggml/src/ggml-opencl/CMakeLists.txt index ff5e8ef46b..97b862afff 100644 --- a/ggml/src/ggml-opencl/CMakeLists.txt +++ b/ggml/src/ggml-opencl/CMakeLists.txt @@ -194,6 +194,7 @@ set(GGML_OPENCL_KERNELS gemv_noshuffle_q6_k_f32_32b_trans gemv_noshuffle_q5_k_f32 gemm_noshuffle_q5_k_f32 + gemv_noshuffle_q5_k_f32_32b_trans mul neg norm diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 2398218940..0b685e697f 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -1264,6 +1264,9 @@ struct ggml_backend_opencl_context { cl_kernel kernel_gemv_noshuffle_q5_k_f32; cl_kernel kernel_gemv_noshuffle_q5_k_f32_mc3; // multi-column (N=3) verify GEMV (spec/MTP) cl_kernel kernel_gemm_noshuffle_q5_k_f32; + cl_kernel kernel_gemm_noshuffle_q5_k_f32_32b_trans_ila_a8_bin; + cl_kernel kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8_bin; + cl_kernel kernel_gemv_noshuffle_q5_k_f32_32b_trans; cl_kernel kernel_gemv_noshuffle_q5_0_f32; cl_kernel kernel_gemm_noshuffle_q5_0_f32; cl_kernel kernel_gemm_noshuffle_q5_0_q8_1_dp4a = nullptr; // dp4a (int8) dense q5_0 prefill GEMM @@ -4455,6 +4458,55 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { } } + backend_ctx->kernel_gemv_noshuffle_q5_k_f32_32b_trans = nullptr; + backend_ctx->kernel_gemm_noshuffle_q5_k_f32_32b_trans_ila_a8_bin = nullptr; + backend_ctx->kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8_bin = nullptr; + if (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E) { + { + std::string opts = std::string("-cl-std=") + opencl_c_std + + " -cl-mad-enable " + " -DSIMDGROUP_WIDTH=" + + std::to_string(backend_ctx->adreno_wave_size); +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "gemv_noshuffle_q5_k_f32_32b_trans.cl.h" + }; +#else + const std::string kernel_src = read_file("gemv_noshuffle_q5_k_f32_32b_trans.cl"); +#endif + cl_program prog = build_program_from_source(backend_ctx, kernel_src.c_str(), opts); + CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q5_k_f32_32b_trans = + clCreateKernel(prog, "gemv_noshuffle_q5_k_f32_32b_trans", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + + if (use_adreno_bin_kernels(backend_ctx)) { + size_t bin_size = 0; + const char * kernel_bin = (const char *)backend_ctx->get_adreno_bin_kernel("gemm_noshuffle_q5_k_f32_32b_trans_ila_a8", &bin_size); + if (kernel_bin && bin_size > 0) { + cl_program bin_prog = + build_program_from_binary(backend_ctx->context, backend_ctx->device, kernel_bin, "", bin_size); + + CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q5_k_f32_32b_trans_ila_a8_bin = + clCreateKernel(bin_prog, "kernel_gemm_noshuffle_q5_k_f32_32b_trans_ila_a8", &err), err)); + CL_CHECK(clReleaseProgram(bin_prog)); + GGML_LOG_CONT("."); + } + + kernel_bin = (const char *)backend_ctx->get_adreno_bin_kernel("gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8", &bin_size); + if (kernel_bin && bin_size > 0) { + cl_program bin_prog = + build_program_from_binary(backend_ctx->context, backend_ctx->device, kernel_bin, "", bin_size); + + CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8_bin = + clCreateKernel(bin_prog, "kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8", &err), err)); + CL_CHECK(clReleaseProgram(bin_prog)); + GGML_LOG_CONT("."); + } + } + } + std::string CL_moe_compile_opts = std::string("-cl-std=") + opencl_c_std + " -cl-mad-enable " " -cl-fast-relaxed-math"; @@ -8749,6 +8801,20 @@ inline bool use_q4_k_bin_kernels(const ggml_backend_opencl_context *backend_ctx, #endif } +inline bool use_q5_k_bin_kernels(const ggml_backend_opencl_context *backend_ctx, const ggml_tensor *tensor) { +#ifdef GGML_OPENCL_USE_ADRENO_KERNELS + if (!backend_ctx->kernel_gemv_noshuffle_q5_k_f32_32b_trans || + !backend_ctx->kernel_gemm_noshuffle_q5_k_f32_32b_trans_ila_a8_bin) { + return false; + } + return (tensor->ne[0] % 256 == 0) && (tensor->ne[1] % 64 == 0); +#else + GGML_UNUSED(backend_ctx); + GGML_UNUSED(tensor); + return false; +#endif +} + 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; @@ -11147,8 +11213,29 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, GGML_ASSERT(K % 32 == 0); - // Transpose q, d, dm as ushort, qh as uchar - transpose_2d_as_16b(backend_ctx, extra->q, extra->q, size_q, K/4, M); + if (use_q5_k_bin_kernels(backend_ctx, tensor)) { + cl_int err; + cl_image_format wimg_fmt; + cl_image_desc wimg_desc; + + // transpose q as 32-bit words (M-first); qh/d/dm stay in their existing layout + // (both new ILA kernels read qh via the existing [K/8][M] uchar plane directly). + GGML_ASSERT(M % 64 == 0); + transpose_2d_as_32b(backend_ctx, extra->q, extra->q, size_q, K/8, M); + + wimg_fmt = { CL_R, CL_UNSIGNED_INT32 }; + memset(&wimg_desc, 0, sizeof(wimg_desc)); + wimg_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + wimg_desc.image_width = (size_t)M * K / 8; + wimg_desc.buffer = extra->q; + CL_CHECK((extra->q_img = clCreateImage(context, CL_MEM_READ_ONLY, &wimg_fmt, &wimg_desc, NULL, &err), err)); + + // Transpose s as uchar + transpose_2d_as_8b(backend_ctx, extra->s, extra->s, size_s, K/256*12, M, true, true); + } else { + // Transpose q as ushort + transpose_2d_as_16b(backend_ctx, extra->q, extra->q, size_q, K/4, M); + } transpose_2d_as_8b (backend_ctx, extra->qh, extra->qh, size_qh, K/8, M); transpose_2d_as_16b(backend_ctx, extra->d, extra->d, size_d, K/256, M); transpose_2d_as_16b(backend_ctx, extra->dm, extra->dm, size_dm, K/256, M); @@ -12333,21 +12420,28 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer, size_t size_q = extra->size_q; size_t size_qh = extra->size_qh; + size_t size_s = extra->size_s; size_t size_d = extra->size_d; size_t size_dm = extra->size_dm; static ggml_cl_buffer buf_trans_q; static ggml_cl_buffer buf_trans_qh; + static ggml_cl_buffer buf_trans_s; static ggml_cl_buffer buf_trans_d; static ggml_cl_buffer buf_trans_dm; buf_trans_q.allocate(backend_ctx->context, size_q); buf_trans_qh.allocate(backend_ctx->context, size_qh); + buf_trans_s.allocate(backend_ctx->context, size_s); buf_trans_d.allocate(backend_ctx->context, size_d); buf_trans_dm.allocate(backend_ctx->context, size_dm); - // Reverse transpose q, qh, d, dm - transpose_2d_as_16b(backend_ctx, extra->q, buf_trans_q.buffer, size_q, M, K/4); + if (use_q5_k_bin_kernels(backend_ctx, tensor)) { + transpose_2d_as_32b(backend_ctx, extra->q, buf_trans_q.buffer, size_q, M, K/8); + transpose_2d_as_8b (backend_ctx, extra->s, buf_trans_s.buffer, size_s, M, K/256*12, true, true); + } else { + transpose_2d_as_16b(backend_ctx, extra->q, buf_trans_q.buffer, size_q, M, K/4); + } transpose_2d_as_8b (backend_ctx, extra->qh, buf_trans_qh.buffer, size_qh, M, K/8); transpose_2d_as_16b(backend_ctx, extra->d, buf_trans_d.buffer, size_d, M, K/256); transpose_2d_as_16b(backend_ctx, extra->dm, buf_trans_dm.buffer, size_dm, M, K/256); @@ -12355,7 +12449,7 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer, cl_kernel kernel = backend_ctx->kernel_restore_block_q5_K_noshuffle; CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &buf_trans_q.buffer)); CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &buf_trans_qh.buffer)); - CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra->s)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &buf_trans_s.buffer)); CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &buf_trans_d.buffer)); CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &buf_trans_dm.buffer)); CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &data_device)); @@ -22443,6 +22537,217 @@ static void ggml_cl_mul_mat_q6_K_f32_adreno(ggml_backend_t backend, const ggml_t #endif } +#ifdef GGML_OPENCL_USE_ADRENO_KERNELS +static void ggml_cl_mul_mat_q5_K_f32_adreno_ila(ggml_backend_t backend, const ggml_tensor * src0, + const ggml_tensor * src1, ggml_tensor * dst) { + GGML_ASSERT(src0); + GGML_ASSERT(src0->extra); + GGML_ASSERT(src1); + GGML_ASSERT(src1->extra); + GGML_ASSERT(dst); + GGML_ASSERT(dst->extra); + + ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context; + + ggml_tensor_extra_cl_q5_K * extra0_q5_k = (ggml_tensor_extra_cl_q5_K *)src0->extra; + ggml_tensor_extra_cl * extra1 = (ggml_tensor_extra_cl *)src1->extra; + ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra; + + cl_ulong offset1 = extra1->offset + src1->view_offs; + cl_ulong offsetd = extrad->offset + dst->view_offs; + + const int ne00 = src0->ne[0]; + const int ne01 = src0->ne[1]; + + const int ne1 = dst->ne[1]; + + GGML_ASSERT(ne00 % ggml_blck_size(src0->type) == 0); + + cl_context context = backend_ctx->context; + cl_kernel kernel; + + cl_int err; + cl_buffer_region region; + cl_image_format img_fmt; + cl_image_desc img_desc; + + const int M = ne01; + const int N = ne1; + const int K = ne00; + + if (ne1 == 1) { + cl_mem b_sub_buf = nullptr; + cl_mem b_img = nullptr; + + region.origin = offset1; + region.size = (size_t)K * N * sizeof(float); + CL_CHECK((b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err)); + + img_fmt = { CL_RGBA, CL_FLOAT }; + memset(&img_desc, 0, sizeof(img_desc)); + img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + img_desc.image_width = (size_t)K * N / 4; + img_desc.buffer = b_sub_buf; + CL_CHECK((b_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, NULL, &err), err)); + + kernel = backend_ctx->kernel_gemv_noshuffle_q5_k_f32_32b_trans; + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q5_k->q_img)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q5_k->qh)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q5_k->d)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0_q5_k->dm)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra0_q5_k->s)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &b_img)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_int), &ne01)); + + size_t local_work_size[3] = { 64, 8, 1 }; + size_t global_work_size[3] = { (size_t)ne01, 8, 1 }; + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); + + CL_CHECK(clReleaseMemObject(b_img)); + CL_CHECK(clReleaseMemObject(b_sub_buf)); + } else { + static const char * q5_k_bin_dp4a_env = getenv("GGML_OPENCL_Q5_K_BIN_DP4A"); + bool q5_k_bin_dp4a_on = q5_k_bin_dp4a_env + ? (atoi(q5_k_bin_dp4a_env) != 0) + : true; + // dot prod has to be available + q5_k_bin_dp4a_on = backend_ctx->has_integer_dot && q5_k_bin_dp4a_on; + + if (q5_k_bin_dp4a_on && backend_ctx->kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8_bin) { + const int dp4a_N_pad = CEIL_DIV(N, 32) * 32; + const size_t n_blocks = (size_t)dp4a_N_pad * (K / 32); + + backend_ctx->prealloc_moe_qa.allocate(context, (size_t)dp4a_N_pad * K * sizeof(cl_char)); + backend_ctx->prealloc_moe_da.allocate(context, n_blocks * sizeof(cl_half)); + backend_ctx->prealloc_moe_sa.allocate(context, n_blocks * sizeof(cl_half)); + + cl_mem b_sub = nullptr; + region.origin = offset1; + region.size = (size_t)K * N * sizeof(float); + CL_CHECK((b_sub = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err)); + + cl_int tb = (cl_int)((size_t)N * (K / 32)); + cl_kernel qk = backend_ctx->kernel_quant_a_q8_1; + CL_CHECK(clSetKernelArg(qk, 0, sizeof(cl_mem), &b_sub)); + CL_CHECK(clSetKernelArg(qk, 1, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer)); + CL_CHECK(clSetKernelArg(qk, 2, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer)); + CL_CHECK(clSetKernelArg(qk, 3, sizeof(cl_mem), &backend_ctx->prealloc_moe_sa.buffer)); + CL_CHECK(clSetKernelArg(qk, 4, sizeof(cl_int), &tb)); + size_t q_local[1] = { 64 }; + size_t q_global[1] = { (size_t)CEIL_DIV(tb, 64) * 64 }; + backend_ctx->enqueue_ndrange_kernel(qk, 1, q_global, q_local, dst); + + cl_mem d_sub = nullptr; + cl_mem d_img = nullptr; + region.origin = offsetd; + region.size = (size_t)M * N * sizeof(float); + CL_CHECK((d_sub = clCreateSubBuffer(extrad->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err)); + + img_fmt = { CL_R, CL_FLOAT }; + memset(&img_desc, 0, sizeof(img_desc)); + img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + img_desc.image_width = (size_t)M * N; + img_desc.buffer = d_sub; + CL_CHECK((d_img = clCreateImage(context, CL_MEM_WRITE_ONLY, &img_fmt, &img_desc, NULL, &err), err)); + + kernel = backend_ctx->kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8_bin; + + cl_uint k_arg = 0; + CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &extra0_q5_k->q_img)); + CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &extra0_q5_k->qh)); + CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &extra0_q5_k->d)); + CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &extra0_q5_k->dm)); + CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &extra0_q5_k->s)); + CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer)); + CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer)); + CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &backend_ctx->prealloc_moe_sa.buffer)); + CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &d_img)); + CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_uint), &ne00)); + CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_uint), &ne01)); + CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_int), &N)); + + size_t local_work_size_dp4a[3] = { 64, 1, 1 }; + size_t global_work_size_dp4a[3] = { 64, (size_t)(M / 64), (size_t)(dp4a_N_pad / 32) }; + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size_dp4a, local_work_size_dp4a, dst); + + CL_CHECK(clReleaseMemObject(b_sub)); + CL_CHECK(clReleaseMemObject(d_img)); + CL_CHECK(clReleaseMemObject(d_sub)); + return; + } + + const int gemm_tile_n = 64; + int N_pad = CEIL_DIV(N, gemm_tile_n) * gemm_tile_n; + + cl_mem b_sub_buf = nullptr; + cl_mem b_padded = nullptr; + cl_mem b_buf = nullptr; + if (N_pad == N) { + region.origin = offset1; + region.size = (size_t)K * N * sizeof(float); + CL_CHECK((b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err)); + b_buf = b_sub_buf; + } else { + CL_CHECK((b_padded = clCreateBuffer(context, CL_MEM_READ_WRITE, (size_t)K * N_pad * sizeof(float), NULL, &err), err)); + const float zero = 0.0f; + CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, b_padded, &zero, sizeof(zero), 0, (size_t)K * N_pad * sizeof(float), 0, NULL, NULL)); + CL_CHECK(clEnqueueCopyBuffer(backend_ctx->queue, extra1->data_device, b_padded, offset1, 0, (size_t)K * N * sizeof(float), 0, NULL, NULL)); + b_buf = b_padded; + } + + img_fmt = { CL_R, CL_FLOAT }; + memset(&img_desc, 0, sizeof(img_desc)); + img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + img_desc.image_width = (size_t)K * N_pad; + img_desc.buffer = b_buf; + cl_mem b_img; + CL_CHECK((b_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, NULL, &err), err)); + + region.origin = offsetd; + region.size = (size_t)M * N * sizeof(float); + cl_mem d_sub_buf; + CL_CHECK((d_sub_buf = clCreateSubBuffer(extrad->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err)); + img_fmt = { CL_R, CL_FLOAT }; + memset(&img_desc, 0, sizeof(img_desc)); + img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + img_desc.image_width = (size_t)M * N; + img_desc.buffer = d_sub_buf; + cl_mem d_img; + CL_CHECK((d_img = clCreateImage(context, CL_MEM_WRITE_ONLY, &img_fmt, &img_desc, NULL, &err), err)); + + kernel = backend_ctx->kernel_gemm_noshuffle_q5_k_f32_32b_trans_ila_a8_bin; + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q5_k->q_img)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q5_k->qh)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q5_k->d)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0_q5_k->dm)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra0_q5_k->s)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &b_img)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), &d_img)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_uint), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_uint), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &N)); + + size_t local_work_size[3] = { 64, 2, 2 }; + size_t m_tiles = (size_t)CEIL_DIV(M, 64); + size_t global_work_size[3] = { 64, m_tiles, (size_t)CEIL_DIV(N_pad, gemm_tile_n) }; + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); + + CL_CHECK(clReleaseMemObject(b_img)); + if (b_sub_buf) { + CL_CHECK(clReleaseMemObject(b_sub_buf)); + } + if (b_padded) { + CL_CHECK(clReleaseMemObject(b_padded)); + } + CL_CHECK(clReleaseMemObject(d_img)); + CL_CHECK(clReleaseMemObject(d_sub_buf)); + } +} +#endif // GGML_OPENCL_USE_ADRENO_KERNELS + static void ggml_cl_mul_mat_q5_K_f32_adreno(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { #ifdef GGML_OPENCL_USE_ADRENO_KERNELS GGML_ASSERT(src0); @@ -22491,6 +22796,20 @@ static void ggml_cl_mul_mat_q5_K_f32_adreno(ggml_backend_t backend, const ggml_t static const bool q5k_mc3 = (getenv("GGML_OPENCL_Q5K_MC3") != nullptr); const bool use_q5k_mc3 = q5k_mc3 && (ne1 >= 2 && ne1 <= 4) && (ne01 < 32768); + const bool use_bin = use_q5_k_bin_kernels(backend_ctx, src0); + + if (use_bin) { + if (use_q5k_mc3) { + static bool warned = false; + if (!warned) { + GGML_LOG_WARN("ggml_opencl: GGML_OPENCL_Q5K_MC3 is bypassed by Q5_K binary kernels\n"); + warned = true; + } + } + ggml_cl_mul_mat_q5_K_f32_adreno_ila(backend, src0, src1, dst); + return; + } + if (ne1 == 1 || use_q5k_mc3) { cl_mem q_img = nullptr; cl_mem qh_img = nullptr; diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_k_f32_32b_trans.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_k_f32_32b_trans.cl new file mode 100644 index 0000000000..ecf15137b5 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_k_f32_32b_trans.cl @@ -0,0 +1,141 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +#pragma OPENCL EXTENSION cl_khr_subgroups : enable +#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable + +#define QK_K 256 +#define K_SCALE_SIZE 12 +#define N_SIMDGROUP 8 +#define SIMDGROUP_WIDTH 64 + +inline void get_scale_min_k4( + int j, + global const uchar * q, + uint stride, + uchar * d, + uchar * m +) { + if (j < 4) { + *d = q[j*stride] & 63; + *m = q[(j+4)*stride] & 63; + } else { + *d = (q[(j+4)*stride] & 0x0F) | ((q[(j-4)*stride] & 0xC0) >> 2); + *m = ((q[(j+4)*stride] >> 4) & 0x0F) | ((q[j*stride] & 0xC0) >> 2); + } +} + +static inline float8 q5_k_to_fp32_packed8(ushort2 q4x8, uint qh_byte, float scale, float minv) { + float8 fp32x8; + fp32x8.s0 = (float)(( q4x8.s0 & 0x000F) | (((qh_byte >> 0) & 1) << 4)) * scale - minv; + fp32x8.s1 = (float)(((q4x8.s0 >> 4) & 0x000F) | (((qh_byte >> 1) & 1) << 4)) * scale - minv; + fp32x8.s2 = (float)(((q4x8.s0 >> 8) & 0x000F) | (((qh_byte >> 2) & 1) << 4)) * scale - minv; + fp32x8.s3 = (float)(((q4x8.s0 >> 12) & 0x000F) | (((qh_byte >> 3) & 1) << 4)) * scale - minv; + fp32x8.s4 = (float)(( q4x8.s1 & 0x000F) | (((qh_byte >> 4) & 1) << 4)) * scale - minv; + fp32x8.s5 = (float)(((q4x8.s1 >> 4) & 0x000F) | (((qh_byte >> 5) & 1) << 4)) * scale - minv; + fp32x8.s6 = (float)(((q4x8.s1 >> 8) & 0x000F) | (((qh_byte >> 6) & 1) << 4)) * scale - minv; + fp32x8.s7 = (float)(((q4x8.s1 >> 12) & 0x000F) | (((qh_byte >> 7) & 1) << 4)) * scale - minv; + return fp32x8; +} + +__attribute__((qcom_reqd_sub_group_size("half"))) +__kernel void gemv_noshuffle_q5_k_f32_32b_trans( + read_only image1d_buffer_t src0_q, + __global uchar * src0_qh, + __global half * src0_d, + __global half * src0_dm, + __global uchar * src0_s, + __read_only image1d_buffer_t src1, + __global float * dst, + ulong offsetd, + int ne00, + int ne01 +) { + uint i01 = get_global_id(0); + uint sgid = get_local_id(1); + uint slid = get_sub_group_local_id(); + + int num_subblocks = ne00 / 32; + + __private float sum = 0.0f; + + // Loop over sub-blocks of 32 elements, N_SIMDGROUP sub-blocks per iter. + for (uint ib = sgid; ib < num_subblocks; ib += N_SIMDGROUP) { + uint sb = ib / 8; + uint j = ib % 8; + + // Load d and dmin for this super-block. + half d_val = src0_d[sb * ne01 + i01]; + half dm_val = src0_dm[sb * ne01 + i01]; + + // Load sub-block scale and min. s is transposed [nb][12][M]; stride ne01 per code. + global const uchar * sc = src0_s + sb * K_SCALE_SIZE * ne01 + i01; + uchar sv, mn; + get_scale_min_k4(j, sc, ne01, &sv, &mn); + + float scale = (float)d_val * (float)sv; + float minv = (float)dm_val * (float)mn; + + // Load 4 uints of quants (32 nibbles = 32 elements), column-major stride ne01. + uint q_base = ib * ne01 * 4 + i01; + + uint4 regQ; + regQ.s0 = read_imageui(src0_q, q_base).x; + regQ.s1 = read_imageui(src0_q, q_base + ne01).x; + regQ.s2 = read_imageui(src0_q, q_base + ne01 * 2).x; + regQ.s3 = read_imageui(src0_q, q_base + ne01 * 3).x; + + uint qh_grp = ib * 4; + uint qh_word = (uint)src0_qh[(qh_grp + 0) * ne01 + i01] + | ((uint)src0_qh[(qh_grp + 1) * ne01 + i01] << 8) + | ((uint)src0_qh[(qh_grp + 2) * ne01 + i01] << 16) + | ((uint)src0_qh[(qh_grp + 3) * ne01 + i01] << 24); + + // Load activations: 32 floats = 8 float4s. + uint y_offset = ib * 8; + + float4 y_local = (slid < 8) ? read_imagef(src1, (y_offset + slid)) : (float4)0.0f; + float4 y0 = sub_group_broadcast(y_local, 0); + float4 y1 = sub_group_broadcast(y_local, 1); + float4 y2 = sub_group_broadcast(y_local, 2); + float4 y3 = sub_group_broadcast(y_local, 3); + float4 y4 = sub_group_broadcast(y_local, 4); + float4 y5 = sub_group_broadcast(y_local, 5); + float4 y6 = sub_group_broadcast(y_local, 6); + float4 y7 = sub_group_broadcast(y_local, 7); + + float8 fp32x8 = q5_k_to_fp32_packed8(as_ushort2(regQ.s0), qh_word & 0xFF, scale, minv); + float4 acc = y0 * fp32x8.lo; + acc += y1 * fp32x8.hi; + + fp32x8 = q5_k_to_fp32_packed8(as_ushort2(regQ.s1), (qh_word >> 8) & 0xFF, scale, minv); + acc += y2 * fp32x8.lo; + acc += y3 * fp32x8.hi; + + fp32x8 = q5_k_to_fp32_packed8(as_ushort2(regQ.s2), (qh_word >> 16) & 0xFF, scale, minv); + acc += y4 * fp32x8.lo; + acc += y5 * fp32x8.hi; + + fp32x8 = q5_k_to_fp32_packed8(as_ushort2(regQ.s3), (qh_word >> 24) & 0xFF, scale, minv); + acc += y6 * fp32x8.lo; + acc += y7 * fp32x8.hi; + + sum += ((acc.s0 + acc.s1) + (acc.s2 + acc.s3)); + } + + // reduction in local memory over N_SIMDGROUP subgroups + __local float reduceLM[SIMDGROUP_WIDTH * (N_SIMDGROUP - 1)]; + if (sgid > 0) { + reduceLM[SIMDGROUP_WIDTH * (sgid - 1) + slid] = sum; + } + barrier(CLK_LOCAL_MEM_FENCE); + if (sgid == 0) { + for (uint i = 0; i < N_SIMDGROUP - 1; ++i) { + sum += reduceLM[SIMDGROUP_WIDTH * i + slid]; + } + } + + // 1 output per thread in subgroup 0 + if (sgid == 0) { + dst = dst + (offsetd >> 2); + dst[i01] = sum; + } +}