diff --git a/ggml/src/ggml-opencl/CMakeLists.txt b/ggml/src/ggml-opencl/CMakeLists.txt index 97b862afffd8..5ba63dade180 100644 --- a/ggml/src/ggml-opencl/CMakeLists.txt +++ b/ggml/src/ggml-opencl/CMakeLists.txt @@ -132,6 +132,7 @@ set(GGML_OPENCL_KERNELS gemm_moe_q6_k_q8_1_dp4a gemm_moe_q8_1_dp4a moe_reorder_quant_a_q8_1 + transpose_activation_q8 gemm_noshuffle_q4_k_q8_1_dp4a gemm_noshuffle_q5_k_q8_1_dp4a gemm_noshuffle_q6_k_q8_1_dp4a diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 52c38c36b017..2bf7e45f603f 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -720,6 +720,10 @@ struct ggml_backend_opencl_context { ggml_cl_buffer prealloc_moe_qa; // int8 quants [tok_slots * ne00] ggml_cl_buffer prealloc_moe_da; // per-block d [tok_slots * ne00/32] (half) ggml_cl_buffer prealloc_moe_sa; // per-block s [tok_slots * ne00/32] (half) + // per-tile transposed qa/da and on-chip scratch for the ldck dp4a MoE GEMM + ggml_cl_buffer prealloc_moe_qa_trans; // int8 quants [tok_slots * ne00] + ggml_cl_buffer prealloc_moe_da_trans; // per-block d [tok_slots * ne00/32] (float) + ggml_cl_buffer prealloc_moe_qda_on_chip; // scratch copy of the router weights to avoid dst aliasing ggml_cl_buffer prealloc_moe_combine_w; ggml_cl_buffer prealloc_splitk_partial; // [ksplit * M] partials for split-K GEMV @@ -1016,6 +1020,7 @@ struct ggml_backend_opencl_context { cl_kernel kernel_gemv_moe_q4_k_f32_ns_wimg = nullptr; // weight-as-texture MoE decode GEMV (opt-in) cl_kernel kernel_gemm_moe_q4_k_q8_1_dp4a = nullptr; // dp4a (int8) prefill GEMM variant cl_kernel kernel_moe_reorder_quant_a_q8_1; // fused reorder + q8_1 quant for the dp4a GEMM + cl_kernel kernel_transpose_activation_q8; // per-tile transpose of q8_1 activations for the ldck dp4a GEMM cl_kernel kernel_gemm_moe_q8_1_dp4a_q80 = nullptr; // generic dp4a MoE GEMM (MOE_QT=80), opt-in cl_kernel kernel_moe_expand_scale_q8_0 = nullptr; // q8_0 per-block d -> uniform scale[16] cl_kernel kernel_gemm_moe_q8_1_dp4a_q50 = nullptr; // generic dp4a MoE GEMM (MOE_QT=50, q5_0), opt-in @@ -1031,6 +1036,7 @@ struct ggml_backend_opencl_context { cl_kernel kernel_gemm_moe_mxfp4_q8_1_dp4a = nullptr; // dp4a (int8) mxfp4 MoE prefill GEMM cl_kernel kernel_gemm_moe_q4_0_q8_1_dp4a = nullptr; // dp4a (int8) q4_0 MoE prefill GEMM cl_kernel kernel_gemm_moe_mxfp4_q8_1_dp4a_bin = nullptr; // binary dp4a (int8) mxfp4 MoE prefill GEMM + cl_kernel kernel_gemm_moe_mxfp4_q8_1_dp4a_ldck_bin = nullptr; // binary dp4a (int8) mxfp4 MoE prefill GEMM using ldck cl_kernel kernel_gemm_moe_q4_0_q8_1_dp4a_bin = nullptr; // binary dp4a (int8) q4_0 MoE prefill GEMM cl_kernel kernel_gemm_moe_q4_k_q8_1_dp4a_bin = nullptr; // binary dp4a (int8) q4_k MoE prefill GEMM cl_kernel kernel_gemm_moe_q6_k_q8_1_dp4a_bin = nullptr; // binary dp4a (int8) q6_k MoE prefill GEMM @@ -4902,6 +4908,24 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { } } + // gemm_moe_mxfp4_q8_1_dp4a_ldck_bin (dp4a prefill GEMM) + if (backend_ctx->has_integer_dot) { + size_t bin_size = 0; + backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_ldck_bin = nullptr; + + if (use_adreno_bin_kernels(backend_ctx)) { + const char * kernel_bin = (const char *)backend_ctx->get_adreno_bin_kernel("gemm_moe_mxfp4_q8_1_dp4a_ila_ldck", &bin_size); + if (kernel_bin && bin_size > 0) { + cl_program prog = + build_program_from_binary(backend_ctx->context, backend_ctx->device, kernel_bin, CL_moe_compile_opts, bin_size); + + CL_CHECK((backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_ldck_bin = clCreateKernel(prog, "kernel_gemm_moe_mxfp4_q8_1_dp4a_ila_ldck", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + } + } + // gemm_moe_q4_0_q8_1_dp4a (dp4a prefill GEMM) if (backend_ctx->has_integer_dot) { #ifdef GGML_OPENCL_EMBED_KERNELS @@ -4983,6 +5007,23 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { GGML_LOG_CONT("."); } + // transpose_activation_q8 (per-tile transpose of q8_1 activations for the ldck dp4a GEMM) + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "transpose_activation_q8.cl.h" + }; +#else + const std::string kernel_src = read_file("transpose_activation_q8.cl"); +#endif + cl_program prog = + build_program_from_source(backend_ctx, kernel_src.c_str(), CL_moe_compile_opts); + + CL_CHECK((backend_ctx->kernel_transpose_activation_q8 = clCreateKernel(prog, "kernel_transpose_activation_q8", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + // gemv_moe_q5_k_f32_ns { #ifdef GGML_OPENCL_EMBED_KERNELS @@ -28217,11 +28258,19 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0, static const char * mxfp4_moe_dp4a_env = getenv("GGML_OPENCL_MXFP4_MOE_DP4A"); bool use_moe_dp4a = mxfp4_moe_dp4a_env ? (atoi(mxfp4_moe_dp4a_env) != 0) - : (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E); + : (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E || backend_ctx->adreno_gen == ADRENO_GPU_GEN::X1E); // dot prod has to be available use_moe_dp4a = backend_ctx->has_integer_dot && use_moe_dp4a; // bin kernel takes precedence, dp4a bin kernel has higher priority than normal bin kernel - if (backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_bin == nullptr) { + // ldck dp4a bin kernel (X1) has the same priority as normal bin kernel (X2/A8x), they are not expected to coexist. + const bool use_moe_dp4a_ldck = backend_ctx->has_integer_dot && + (!mxfp4_moe_dp4a_env || atoi(mxfp4_moe_dp4a_env) != 0) && + backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_bin == nullptr && + backend_ctx->kernel_gemm_moe_mxfp4_f32_ns_bin == nullptr && + backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_ldck_bin != nullptr; + if (use_moe_dp4a_ldck) { + use_moe_dp4a = true; + } else if (backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_bin == nullptr) { use_moe_dp4a = use_moe_dp4a && backend_ctx->kernel_gemm_moe_mxfp4_f32_ns_bin == nullptr; } @@ -28310,6 +28359,12 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0, backend_ctx->prealloc_moe_qa.allocate(backend_ctx->context, tok_slots * ne00 * sizeof(cl_char)); backend_ctx->prealloc_moe_da.allocate(backend_ctx->context, n_blocks * sizeof(cl_half)); backend_ctx->prealloc_moe_sa.allocate(backend_ctx->context, n_blocks * sizeof(cl_half)); + if (use_moe_dp4a_ldck) { + backend_ctx->prealloc_moe_qa_trans.allocate(backend_ctx->context, tok_slots * ne00 * sizeof(cl_char)); + backend_ctx->prealloc_moe_da_trans.allocate(backend_ctx->context, n_blocks * sizeof(cl_float)); + // matches max_constant_size of src1_a_cm in the ldck kernel + backend_ctx->prealloc_moe_qda_on_chip.allocate(backend_ctx->context, 2304); + } // fused reorder + q8_1 quant straight from the original // activations (no intermediate f32 reorder buffer). mxfp4 has no @@ -28334,13 +28389,33 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0, cl_kernel dk = backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a; if (backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_bin) { dk = backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_bin; + } else if (use_moe_dp4a_ldck) { + dk = backend_ctx->kernel_gemm_moe_mxfp4_q8_1_dp4a_ldck_bin; + + // per-tile transpose of qa/da for the ldck kernel + cl_kernel tk = backend_ctx->kernel_transpose_activation_q8; + CL_CHECK(clSetKernelArg(tk, 0, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer)); + CL_CHECK(clSetKernelArg(tk, 1, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa_trans.buffer)); + CL_CHECK(clSetKernelArg(tk, 2, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer)); + CL_CHECK(clSetKernelArg(tk, 3, sizeof(cl_mem), &backend_ctx->prealloc_moe_da_trans.buffer)); + CL_CHECK(clSetKernelArg(tk, 4, sizeof(cl_mem), &(backend_ctx->prealloc_total_tiles.buffer))); + CL_CHECK(clSetKernelArg(tk, 5, sizeof(cl_uint), &ne00)); + size_t tk_local[3] = { 32, 8, 1 }; + size_t tk_global[3] = { 32, (size_t)(ne00 / 4), (size_t)max_post_router_tile }; + backend_ctx->enqueue_ndrange_kernel(tk, 3, tk_global, tk_local, dst); } int aidx = 0; CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &extra0_mxfp4->q_img)); CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &extra0_mxfp4->e)); - CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer)); - CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer)); + if (use_moe_dp4a_ldck) { + CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa_trans.buffer)); + CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &backend_ctx->prealloc_moe_da_trans.buffer)); + CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &backend_ctx->prealloc_moe_qda_on_chip.buffer)); + } else { + CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer)); + CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer)); + } CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &buf_src2)); CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &buf_src2_emap)); CL_CHECK(clSetKernelArg(dk, aidx++, sizeof(cl_mem), &buf_dst_image)); diff --git a/ggml/src/ggml-opencl/kernels/transpose_activation_q8.cl b/ggml/src/ggml-opencl/kernels/transpose_activation_q8.cl new file mode 100644 index 000000000000..34d54b57e24a --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/transpose_activation_q8.cl @@ -0,0 +1,35 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +#pragma OPENCL EXTENSION cl_khr_subgroups : enable + +#define tile_size 32 + +kernel void kernel_transpose_activation_q8( + __global uint * src1_qa, + __global uint * src1_qa_trans, + __global half * src1_sa, + __global float * src1_sa_trans, + __global const int * total_tiles, + uint K +) { + uint row_in_tile = get_global_id(0); + uint blk_k = get_group_id(1); + uint col_in_tile = get_local_id(1); + uint tile_n = get_global_id(2); + + if (blk_k >= (K / tile_size) || tile_n >= (uint)total_tiles[0]) { + return; + } + + uint src_idx = (tile_n * tile_size + row_in_tile) * K / 4 + blk_k * 8 + col_in_tile; + uint dst_idx = tile_n * tile_size * K / 4 + (blk_k * 8 + col_in_tile) * tile_size + row_in_tile; + + src1_qa_trans[dst_idx] = src1_qa[src_idx]; + + if (col_in_tile == 0) { + src_idx = (tile_n * tile_size + row_in_tile) * K / 32 + blk_k; + dst_idx = tile_n * tile_size * K / 32 + blk_k * tile_size + row_in_tile; + + src1_sa_trans[dst_idx] = (float)(src1_sa[src_idx]); + } + +}