diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index ad65de00a..70f199dba 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -831,6 +831,7 @@ struct ggml_backend_opencl_context { cl_program program_mul_mv_id_mxfp4_f32_flat; cl_program program_mul_mm_f32_f32_l4_lm; cl_program program_mul_mm_f16_f32_l4_lm; + cl_program program_mul_mm_f16_f32_l4_lm_n8; cl_program program_mul_mm_q8_0_f32_l4_lm; cl_kernel kernel_add, kernel_add_row, kernel_add_f16, kernel_add_row_f16; @@ -1052,6 +1053,7 @@ struct ggml_backend_opencl_context { cl_kernel kernel_mul_mm_f32_f32_l4_lm; cl_kernel kernel_gemv_f32_f32_mc; // multi-column (small-N) f32 GEMV for spec/MTP verify cl_kernel kernel_mul_mm_f16_f32_l4_lm; + cl_kernel kernel_mul_mm_f16_f32_l4_lm_n8; cl_kernel kernel_mul_mm_q1_0_f32_l4_lm; cl_kernel kernel_mul_mm_q4_0_f32_l4_lm; cl_kernel kernel_mul_mm_q4_1_f32_l4_lm; @@ -1247,6 +1249,8 @@ struct ggml_backend_opencl_context { cl_kernel kernel_gemv_noshuffle_q8_0_f32_splitk; // split-K across WGs (small-M decode) cl_kernel kernel_gemm_noshuffle_q1_0_f32; cl_kernel kernel_gemv_noshuffle_q1_0_f32; + cl_kernel kernel_gemv_noshuffle_q1_0_f32_nsg1 = nullptr; + int min_m_gemv_nsg1_q1_0 = 0; cl_kernel kernel_gemv_noshuffle_q4_k_f32; cl_kernel kernel_gemv_noshuffle_q4_k_f32_o4; // 4-output-per-WI, long-vocab lm_head cl_kernel kernel_gemv_noshuffle_q4_k_f32_tiled; // tiled-wide layout (opt-in) @@ -1261,6 +1265,11 @@ struct ggml_backend_opencl_context { cl_kernel kernel_gemv_noshuffle_q4_k_f32_32b_trans; cl_kernel kernel_gemm_noshuffle_q4_k_q8_1_dp4a = nullptr; // dp4a (int8) dense prefill GEMM cl_kernel kernel_gemm_noshuffle_q4_k_q8_1_dp4a_wimg = nullptr; // dp4a dense prefill GEMM, weights via texture (X1 opt-in) + cl_kernel kernel_gemm_noshuffle_q4_k_q8_1_dp4a_narrow = nullptr; + cl_kernel kernel_gemm_noshuffle_q4_k_q8_1_dp4a_narrow_wimg = nullptr; + int tile_n_dp4a_q4_k = 32; + int tile_n_dp4a_narrow_q4_k = 32; + int max_n_dp4a_narrow_q4_k = 16; cl_kernel kernel_gemm_noshuffle_q5_k_q8_1_dp4a = nullptr; // dp4a (int8) dense q5_K prefill GEMM cl_kernel kernel_gemm_noshuffle_q6_k_q8_1_dp4a = nullptr; // dp4a (int8) dense q6_K prefill GEMM cl_kernel kernel_quant_a_q8_1; // plain activation q8_1 pre-pass @@ -2781,6 +2790,11 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts); CL_CHECK((backend_ctx->kernel_mul_mm_f16_f32_l4_lm = clCreateKernel(backend_ctx->program_mul_mm_f16_f32_l4_lm, "kernel_mul_mm_f16_f32_l4_lm", &err), err)); + + const std::string narrow_opts = compile_opts + " -DBN=8 -DTN=1 -DKERNEL_NAME_LM=kernel_mul_mm_f16_f32_l4_lm_n8"; + backend_ctx->program_mul_mm_f16_f32_l4_lm_n8 = + build_program_from_source(backend_ctx, kernel_src.c_str(), narrow_opts); + CL_CHECK((backend_ctx->kernel_mul_mm_f16_f32_l4_lm_n8 = clCreateKernel(backend_ctx->program_mul_mm_f16_f32_l4_lm_n8, "kernel_mul_mm_f16_f32_l4_lm_n8", &err), err)); GGML_LOG_CONT("."); } @@ -3983,6 +3997,16 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q1_0_f32 = clCreateKernel(prog, "kernel_gemv_noshuffle_q1_0_f32", &err), err)); CL_CHECK(clReleaseProgram(prog)); + + const char * nsg1_min_m_env = getenv("GGML_OPENCL_Q1_0_GEMV_NSG1_MIN_M"); + backend_ctx->min_m_gemv_nsg1_q1_0 = nsg1_min_m_env ? atoi(nsg1_min_m_env) : + (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E ? 8192 : 0); + if (backend_ctx->min_m_gemv_nsg1_q1_0 > 0) { + const std::string nsg1_opts = CL_gemv_compile_opts + " -DN_SIMDGROUP=1"; + cl_program prog_nsg1 = build_program_from_source(backend_ctx, kernel_src_CL_gemv_general.c_str(), nsg1_opts); + CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q1_0_f32_nsg1 = clCreateKernel(prog_nsg1, "kernel_gemv_noshuffle_q1_0_f32", &err), err)); + CL_CHECK(clReleaseProgram(prog_nsg1)); + } GGML_LOG_CONT("."); } @@ -4453,16 +4477,32 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { #else const std::string kernel_src = read_file("gemm_noshuffle_q4_k_q8_1_dp4a.cl"); #endif - // Per-device dp4a dense tile. The X2-tuned TILESIZE_N=32 over-occupies LDS on - // X1 (1152 B/WG -> few resident WGs); TILESIZE_N=8 (288 B) lifts occupancy on - // X1, byte-identical. X2E keeps 32. Env override wins. - int q4k_dp4a_ts = (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X1E) ? 8 : 32; - if (const char * e = getenv("GGML_OPENCL_Q4K_DP4A_TS")) { q4k_dp4a_ts = atoi(e); } - std::string dp4a_opts = compile_opts + " -DTILESIZE_N=" + std::to_string(q4k_dp4a_ts); + // X1 uses a smaller N tile to reduce local memory and register use. + int tile_n = (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X1E) ? 8 : 32; + if (const char * e = getenv("GGML_OPENCL_Q4K_DP4A_TS")) { tile_n = atoi(e); } + GGML_ASSERT(tile_n >= 4 && tile_n <= 64 && tile_n % 4 == 0); + backend_ctx->tile_n_dp4a_q4_k = tile_n; + std::string dp4a_opts = compile_opts + " -DTILESIZE_N=" + std::to_string(tile_n); cl_program prog = build_program_from_source(backend_ctx, kernel_src.c_str(), dp4a_opts); CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q4_k_q8_1_dp4a = clCreateKernel(prog, "kernel_gemm_noshuffle_q4_k_q8_1_dp4a", &err), err)); CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q4_k_q8_1_dp4a_wimg = clCreateKernel(prog, "kernel_gemm_noshuffle_q4_k_q8_1_dp4a_wimg", &err), err)); CL_CHECK(clReleaseProgram(prog)); + + // Compile a smaller N tile to avoid unused columns in small batches. + int tile_n_narrow = MIN(tile_n, 16); + if (const char * e = getenv("GGML_OPENCL_Q4_K_DP4A_TS_NARROW")) { tile_n_narrow = atoi(e); } + GGML_ASSERT(tile_n_narrow >= 4 && tile_n_narrow <= 64 && tile_n_narrow % 4 == 0); + if (const char * e = getenv("GGML_OPENCL_Q4_K_DP4A_NARROW_MAX")) { + backend_ctx->max_n_dp4a_narrow_q4_k = atoi(e); + } + backend_ctx->tile_n_dp4a_narrow_q4_k = tile_n_narrow; + if (tile_n_narrow != tile_n) { + std::string narrow_opts = compile_opts + " -DTILESIZE_N=" + std::to_string(tile_n_narrow); + cl_program narrow_prog = build_program_from_source(backend_ctx, kernel_src.c_str(), narrow_opts); + CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q4_k_q8_1_dp4a_narrow = clCreateKernel(narrow_prog, "kernel_gemm_noshuffle_q4_k_q8_1_dp4a", &err), err)); + CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q4_k_q8_1_dp4a_narrow_wimg = clCreateKernel(narrow_prog, "kernel_gemm_noshuffle_q4_k_q8_1_dp4a_wimg", &err), err)); + CL_CHECK(clReleaseProgram(narrow_prog)); + } GGML_LOG_CONT("."); } @@ -9531,22 +9571,6 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te op->src[0]->ne[1] >= 32768) { // vocab-scale weight; no FFN/attn weight is this tall return false; } - // The generic mul_mv (GEMV) kernels are wrong for large-batch prefill on - // Adreno. A quant mul_mat only avoids the GEMV when it reaches the Adreno - // trans-weight GEMM, which needs both a GEMM kernel for the type and - // use_adreno_kernels(). Decline the large-N shapes that would otherwise - // fall through to the GEMV. - { - const ggml_type t = op->src[0]->type; - const bool type_has_gemm = (t == GGML_TYPE_Q4_0 || t == GGML_TYPE_Q4_1 || - t == GGML_TYPE_IQ4_NL || t == GGML_TYPE_Q8_0 || - t == GGML_TYPE_Q4_K || t == GGML_TYPE_Q5_K || - t == GGML_TYPE_Q6_K); - const bool uses_gemm = type_has_gemm && use_adreno_kernels(backend_ctx, op->src[0]); - if (!uses_gemm && op->src[1]->ne[1] >= 512) { - return false; - } - } return op->src[1]->type == GGML_TYPE_F32 && ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op->src[1]); } else if (op->src[0]->type == GGML_TYPE_Q8_0) { // ggml_cl_mul_mat_q8_0_f32_adreno now honors src1/dst view_offs (the @@ -19834,7 +19858,9 @@ static void ggml_cl_mul_mat_q1_0_f32_adreno(ggml_backend_t backend, const ggml_t 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_q1_0_f32; + // Tall M has enough row parallelism to skip the K split and its reduction. + const bool use_nsg1 = backend_ctx->min_m_gemv_nsg1_q1_0 > 0 && M >= backend_ctx->min_m_gemv_nsg1_q1_0; + kernel = use_nsg1 ? backend_ctx->kernel_gemv_noshuffle_q1_0_f32_nsg1 : backend_ctx->kernel_gemv_noshuffle_q1_0_f32; int r2 = 1; int r3 = 1; @@ -19856,8 +19882,9 @@ static void ggml_cl_mul_mat_q1_0_f32_adreno(ggml_backend_t backend, const ggml_t CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &r3)); size_t wavesize = backend_ctx->adreno_wave_size; - size_t local_work_size[] = { wavesize, 4, 1 }; - size_t global_work_size[] = { CEIL_DIV(M, wavesize)*wavesize, 4, 1 }; + const size_t nsg = use_nsg1 ? 1 : 4; + size_t local_work_size[] = { wavesize, nsg, 1 }; + size_t global_work_size[] = { CEIL_DIV(M, wavesize)*wavesize, nsg, 1 }; backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); @@ -20321,7 +20348,14 @@ static void ggml_cl_mul_mat_q4_0_f32_adreno(ggml_backend_t backend, const ggml_t // per-lane K-walk) for small M. Layout stride is fixed (4 uints/block), so only // the K-split count changes; the mc3 kernel reads it via get_local_size(1). The // ne1==1 base kernel hardcodes N_SIMDGROUP=4, so it always stays at 4. - const int mc3_nsg = (use_q40_mc3 && ne01 < 4096) ? 8 : 4; + int mc3_nsg = (use_q40_mc3 && ne01 < 4096) ? 8 : 4; + if (use_q40_mc3) { + // Register pressure can lower the kernel limit below the device limit. + const size_t max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); + while (mc3_nsg > 1 && (size_t)(64 * mc3_nsg) > max_workgroup_size) { + mc3_nsg /= 2; + } + } size_t local_work_size[3] = {64, (size_t)mc3_nsg, 1}; size_t global_work_size[3] = {(size_t)CEIL_DIV(ne01/2, 64)*64, (size_t)mc3_nsg, 1}; @@ -22667,8 +22701,13 @@ static void ggml_cl_mul_mat_q4_k_f32_adreno(ggml_backend_t backend, const ggml_t } } - cl_kernel dk = use_wimg ? backend_ctx->kernel_gemm_noshuffle_q4_k_q8_1_dp4a_wimg - : backend_ctx->kernel_gemm_noshuffle_q4_k_q8_1_dp4a; + const bool use_narrow = backend_ctx->kernel_gemm_noshuffle_q4_k_q8_1_dp4a_narrow != nullptr + && N <= backend_ctx->max_n_dp4a_narrow_q4_k; + cl_kernel dk = use_narrow + ? (use_wimg ? backend_ctx->kernel_gemm_noshuffle_q4_k_q8_1_dp4a_narrow_wimg + : backend_ctx->kernel_gemm_noshuffle_q4_k_q8_1_dp4a_narrow) + : (use_wimg ? backend_ctx->kernel_gemm_noshuffle_q4_k_q8_1_dp4a_wimg + : backend_ctx->kernel_gemm_noshuffle_q4_k_q8_1_dp4a); int ai = 0; if (use_wimg) { CL_CHECK(clSetKernelArg(dk, ai++, sizeof(cl_mem), &q4k_q_img)); @@ -22689,12 +22728,9 @@ static void ggml_cl_mul_mat_q4_k_f32_adreno(ggml_backend_t backend, const ggml_t CL_CHECK(clSetKernelArg(dk, ai++, sizeof(cl_uchar), &mask_d6)); CL_CHECK(clSetKernelArg(dk, ai++, sizeof(cl_uchar), &mask_d4)); CL_CHECK(clSetKernelArg(dk, ai++, sizeof(cl_uchar), &mask_hi2)); - // Must match the compile-time TILESIZE_N chosen at program build (per-device, - // X1E=8 else 32; env override). Same inputs -> same value. - int q4k_dp4a_ts = (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X1E) ? 8 : 32; - if (const char * e = getenv("GGML_OPENCL_Q4K_DP4A_TS")) q4k_dp4a_ts = atoi(e); + const int tile_n = use_narrow ? backend_ctx->tile_n_dp4a_narrow_q4_k : backend_ctx->tile_n_dp4a_q4_k; size_t d_local[3] = { 64, 1, 1 }; - size_t d_global[3] = { 64, (size_t)(M / 64), (size_t)CEIL_DIV(N, q4k_dp4a_ts) }; + size_t d_global[3] = { 64, (size_t)(M / 64), (size_t)CEIL_DIV(N, tile_n) }; backend_ctx->enqueue_ndrange_kernel(dk, 3, d_global, d_local, dst); if (q4k_q_img != nullptr) { @@ -24815,7 +24851,9 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co return; } #endif - kernel = backend_ctx->kernel_mul_mm_f16_f32_l4_lm; + static const char * mm_narrow_n_env = getenv("GGML_OPENCL_MM_NARROW_N"); + const bool use_narrow_n = (mm_narrow_n_env == nullptr || mm_narrow_n_env[0] != '0') && ne11 >= 2 && ne11 <= 8; + kernel = use_narrow_n ? backend_ctx->kernel_mul_mm_f16_f32_l4_lm_n8 : backend_ctx->kernel_mul_mm_f16_f32_l4_lm; nth0 = 128; // calculated as (BM*BN)/(TM*TN) int batch_stride_a = ne00*ne01; @@ -24874,8 +24912,9 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int), &r2)); CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int), &r3)); - // 64 is block tile size BM and BN - change here when BM and BN in the kernel are changed. - size_t global_work_size[] = {(size_t)(CEIL_DIV(ne01, 64)*nth0), (size_t)(CEIL_DIV(ne11, 64)), (size_t)ne12*ne13}; + // BN=8/TN=1 keeps 128 threads for small batches, including attention KQ/KQV. + const int bn = use_narrow_n ? 8 : 64; + size_t global_work_size[] = {(size_t)(CEIL_DIV(ne01, 64)*nth0), (size_t)(CEIL_DIV(ne11, bn)), (size_t)ne12*ne13}; size_t local_work_size[] = {(size_t)nth0, 1, 1}; backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q1_0_f32.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q1_0_f32.cl index d935a489c..572f11637 100644 --- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q1_0_f32.cl +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q1_0_f32.cl @@ -8,7 +8,9 @@ #endif #define QK1_0 128 +#ifndef N_SIMDGROUP #define N_SIMDGROUP 4 +#endif #define dequantizeBlockAccum_q1(total, bits, scale, regB, lb) \ total += (2.0f*(float)((bits >> 0) & 1u) - 1.0f) * scale * sub_group_broadcast(regB.s0, lb+0); \ @@ -106,15 +108,19 @@ __kernel void kernel_gemv_noshuffle_q1_0_f32( dequantizeBlockAccum_q1(totalSum, regA.s3, scale, regB, 12); } - // reduction in local memory, assumes #wave = N_SIMDGROUP = 4 - local float reduceLM[SIMDGROUP_WIDTH * 3]; - if (groupId == 1) reduceLM[SIMDGROUP_WIDTH * 0 + slid] = totalSum; - if (groupId == 2) reduceLM[SIMDGROUP_WIDTH * 1 + slid] = totalSum; - if (groupId == 3) reduceLM[SIMDGROUP_WIDTH * 2 + slid] = totalSum; +#if N_SIMDGROUP > 1 + // Sum the K slices from the other subgroups. + local float reduceLM[SIMDGROUP_WIDTH * (N_SIMDGROUP - 1)]; + if (groupId > 0) { + reduceLM[SIMDGROUP_WIDTH * (groupId - 1) + slid] = totalSum; + } barrier(CLK_LOCAL_MEM_FENCE); - if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 0 + slid]; - if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 1 + slid]; - if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 2 + slid]; + if (groupId == 0) { + for (uint w = 0; w < N_SIMDGROUP - 1; ++w) { + totalSum += reduceLM[SIMDGROUP_WIDTH * w + slid]; + } + } +#endif if (groupId == 0) { dst = (global float*)((global char*)dst + offsetd); @@ -122,6 +128,8 @@ __kernel void kernel_gemv_noshuffle_q1_0_f32( // so when ne01 is not a multiple of the wave size the tail work-items run past // row ne01 and would overrun dst into the adjacent tensor. No-op / byte-identical // when ne01 is wave-aligned (no padding). - if (gid < M) dst[gid] = totalSum; + if (gid < M) { + dst[gid] = totalSum; + } } } diff --git a/ggml/src/ggml-opencl/kernels/mul_mm_f16_f32_l4_lm.cl b/ggml/src/ggml-opencl/kernels/mul_mm_f16_f32_l4_lm.cl index 6982f8f51..eba1778af 100644 --- a/ggml/src/ggml-opencl/kernels/mul_mm_f16_f32_l4_lm.cl +++ b/ggml/src/ggml-opencl/kernels/mul_mm_f16_f32_l4_lm.cl @@ -4,12 +4,20 @@ #define LOAD_VEC_B 4 #define BM 64 +#ifndef BN #define BN 64 +#endif #define BK 16 #define TM 4 +#ifndef TN #define TN 8 +#endif -kernel void kernel_mul_mm_f16_f32_l4_lm( +#ifndef KERNEL_NAME_LM +#define KERNEL_NAME_LM kernel_mul_mm_f16_f32_l4_lm +#endif + +kernel void KERNEL_NAME_LM( global half4 * src0, ulong offset0, global float4 * src1, @@ -93,7 +101,8 @@ kernel void kernel_mul_mm_f16_f32_l4_lm( } } - for (int l = 0; l < BN; l += loadstride_b) { + // BN=8 uses only the threads that load a column inside the tile. + for (int l = 0; loadc_b + l < BN; l += loadstride_b) { if (ic*BN + loadc_b + l < ne11) { const int idx = pos_b + (loadc_b + l) * stride_b / LOAD_VEC_B + loadr_b; buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = src1[idx].s0; diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_mxfp4_f32_flat.cl b/ggml/src/ggml-opencl/kernels/mul_mv_mxfp4_f32_flat.cl index 3d5a923ee..2e31831ca 100644 --- a/ggml/src/ggml-opencl/kernels/mul_mv_mxfp4_f32_flat.cl +++ b/ggml/src/ggml-opencl/kernels/mul_mv_mxfp4_f32_flat.cl @@ -132,7 +132,7 @@ kernel void kernel_mul_mv_mxfp4_f32_flat( global float4 * y4 = (global float4 *)yb; #pragma unroll - for (short row = 0; row < N_R0_MXFP4; row++) { + for (short row = 0; row < N_R0_MXFP4 && first_row + row < ne0; row++) { uchar xb_e = x_e[row * nb + ib]; #ifdef SRC0Q_IMG ushort4 xb_q = as_ushort4(read_imageui(src0_q, (offset_q + row * nb + ib) * 2 + it).xy);