From 6623435813c1c049a7ccac158c50395c78bfbdbe Mon Sep 17 00:00:00 2001 From: Yee Man Chan Date: Tue, 19 May 2026 08:33:46 +0800 Subject: [PATCH 1/7] OP_GATED_DELTA_NET impl --- ggml/src/ggml-opencl/CMakeLists.txt | 1 + ggml/src/ggml-opencl/ggml-opencl.cpp | 300 +++++++++++++++++- .../ggml-opencl/kernels/gated_delta_net.cl | 209 ++++++++++++ 3 files changed, 509 insertions(+), 1 deletion(-) create mode 100644 ggml/src/ggml-opencl/kernels/gated_delta_net.cl diff --git a/ggml/src/ggml-opencl/CMakeLists.txt b/ggml/src/ggml-opencl/CMakeLists.txt index 7edb3eb4e9c1..1dd04b864e9d 100644 --- a/ggml/src/ggml-opencl/CMakeLists.txt +++ b/ggml/src/ggml-opencl/CMakeLists.txt @@ -154,6 +154,7 @@ set(GGML_OPENCL_KERNELS sqr sqrt ssm_conv + gated_delta_net sub sum_rows cumsum diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 73a58f74a94c..18a212fed17c 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -395,6 +395,10 @@ struct ggml_backend_opencl_context { size_t max_workgroup_size; bool fp16_support; bool has_vector_subgroup_broadcast; + bool has_subgroup_arithmetic = false; // cl_khr_subgroup_arithmetic + bool has_subgroup_clustered_reduce = false; // cl_khr_subgroup_clustered_reduce + bool has_qcom_subgroup_shuffle = false; // cl_qcom_subgroup_shuffle + int subgroup_size = 0; // wavefront / subgroup width bool disable_fusion; std::regex *opfilter = nullptr; // regex of ops to not claim @@ -601,6 +605,10 @@ struct ggml_backend_opencl_context { cl_kernel kernel_conv_2d_f32; cl_kernel kernel_conv_2d_f16_f32; cl_kernel kernel_ssm_conv_f32_f32, kernel_ssm_conv_f32_f32_4; + // [size_idx][kda][tile] where size_idx: 0=S_V=16, 1=32, 2=64, 3=128; kda: 0 or 1. + // tile 0 = TG variant (COLS_PER_LANE_GROUP=1), tile 1 = prefill variant (COLS_PER_LANE_GROUP=4). + cl_kernel kernel_gated_delta_net_f32[4][2][2] = {}; + cl_kernel kernel_timestep_embedding; cl_kernel kernel_gemv_moe_q4_0_f32_ns, kernel_gemm_moe_q4_0_f32_ns; cl_kernel kernel_gemv_moe_q4_1_f32_ns, kernel_gemm_moe_q4_1_f32_ns; @@ -2473,6 +2481,81 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve GGML_LOG_CONT("."); } + // gated_delta_net: one kernel per (S_V, KDA, tgpp) triple. + { + #ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "gated_delta_net.cl.h" + }; + #else + const std::string kernel_src = read_file("gated_delta_net.cl"); + #endif + + const int gdn_sizes[4] = { 16, 32, 64, 128 }; + const int sg_size = backend_ctx->subgroup_size > 0 ? backend_ctx->subgroup_size : 32; + + for (int si = 0; si < 4; si++) { + const int S_V = gdn_sizes[si]; + + // MUST match the dispatcher heuristic in ggml_cl_gated_delta_net exactly. + if (S_V >= 128) { + lanes_per_column = 8; + } else { + lanes_per_column = std::min(S_V, sg_size); + } + + // Round LANES_PER_COLUMN down until it is: + // * power-of-two + // * divides both S_V and sg_size + while (lanes_per_column > 1 && + (((lanes_per_column & (lanes_per_column - 1)) != 0) || + (S_V % lanes_per_column) != 0 || + (sg_size % lanes_per_column) != 0)) { + lanes_per_column >>= 1; + } + + GGML_ASSERT(lanes_per_column >= 1); + GGML_ASSERT(((lanes_per_column & (lanes_per_column - 1)) == 0)); + GGML_ASSERT((S_V % lanes_per_column) == 0); + GGML_ASSERT((sg_size % lanes_per_column) == 0); + + const bool is_partial_reduce = (lanes_per_column != 1) && (lanes_per_column < sg_size); + int use_qcom_shuffle = 0; + if (is_partial_reduce) { + if (backend_ctx->has_qcom_subgroup_shuffle) { + use_qcom_shuffle = 1; + } + } + for (int kda = 0; kda < 2; kda++) { + for (int tgpp = 0; tgpp < 2; tgpp++) { + const int cpl = (tgpp == 0) ? 1 : 4; + const int spw = (tgpp == 0) ? 1 : 1; + + std::string opts = compile_opts; + opts += " -DS_V=" + std::to_string(S_V); + opts += " -DKDA=" + std::to_string(kda); + opts += " -DSUBGROUP_SIZE=" + std::to_string(sg_size); + opts += " -DLANES_PER_COLUMN=" + std::to_string(lanes_per_column); + opts += " -DCOLS_PER_LANE_GROUP=" + std::to_string(cpl); + opts += " -DUSE_QCOM_SUBGROUP_SHUFFLE=" + std::to_string(use_qcom_shuffle); + + // Since spw=1 is found to be optimal, SUBGROUPS_PER_WG > 1 code in + // the kernel is removed. If you want to experiment with spw > 1, + // Please remember to implement code to handle it. + opts += " -DSUBGROUPS_PER_WG=" + std::to_string(spw); + + cl_program prog = build_program_from_source( + backend_ctx->context, backend_ctx->device, kernel_src.c_str(), opts); + + CL_CHECK((backend_ctx->kernel_gated_delta_net_f32[si][kda][tgpp] = + clCreateKernel(prog, "kernel_gated_delta_net", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + } + } + } + GGML_LOG_CONT("."); + } + // mul_mv_id_q4_0_f32_8x_flat { #ifdef GGML_OPENCL_EMBED_KERNELS @@ -3443,6 +3526,21 @@ static ggml_backend_opencl_context * ggml_cl2_init(ggml_backend_dev_t dev) { // check Adreno large buffer support backend_ctx->adreno_has_large_buffer = strstr(ext_buffer, "cl_qcom_large_buffer") != NULL; + // Detect subgroup arithmetic and clustered-reduce support. These are + // optional extensions; the gated_delta_net kernel uses them when + // available and falls back to a local-memory reduction otherwise. + backend_ctx->has_subgroup_arithmetic = strstr(ext_buffer, "cl_khr_subgroup_arithmetic") != NULL; + backend_ctx->has_subgroup_clustered_reduce = strstr(ext_buffer, "cl_khr_subgroup_clustered_reduce") != NULL; + // The kernel needs a subgroup-size hint at build time. Use the + // Adreno wave size when we already know it; otherwise default to 32 + // (a common subgroup size for Intel/Mali). The local-memory fallback + // is correct for any device whose actual subgroup_size divides this. + backend_ctx->subgroup_size = backend_ctx->gpu_family == ADRENO ? backend_ctx->adreno_wave_size : 32; + GGML_LOG_INFO("ggml_opencl: subgroup arithmetic support: %s\n", + backend_ctx->has_subgroup_arithmetic ? "true" : "false"); + GGML_LOG_INFO("ggml_opencl: subgroup clustered reduce support: %s\n", + backend_ctx->has_subgroup_clustered_reduce ? "true" : "false"); + // fp16 is required if (!backend_ctx->fp16_support) { GGML_LOG_ERROR("ggml_opencl: device does not support FP16\n"); @@ -3453,10 +3551,19 @@ static ggml_backend_opencl_context * ggml_cl2_init(ggml_backend_dev_t dev) { // optional in OpenCL 3.0 (cl_khr_subgroup is mandatory in OpenCL 2.x) if (opencl_c_version.major == 3 && strstr(ext_buffer, "cl_khr_subgroups") == NULL && strstr(ext_buffer, "cl_intel_subgroups") == NULL) { - GGML_LOG_ERROR("ggml_opencl: device does not support subgroups (cl_khr_subgroups or cl_intel_subgroups) " + GGML_LOG_INFO("ggml_opencl: device does not support subgroups (cl_khr_subgroups or cl_intel_subgroups) " "(note that subgroups is an optional feature in OpenCL 3.0)\n"); return nullptr; } + if (opencl_c_version.major == 3 && strstr(ext_buffer, "cl_khr_subgroups") != NULL) { + GGML_LOG_INFO("ggml_opencl: cl_khr_subgroups support: true\n"); + if (strstr(ext_buffer, "cl_qcom_subgroup_shuffle") != NULL) { + GGML_LOG_INFO("ggml_opencl: cl_qcom_subgroup_shuffle support: true\n"); + backend_ctx->has_qcom_subgroup_shuffle = true; + } + else + GGML_LOG_INFO("ggml_opencl: cl_qcom_subgroup_shuffle support: false\n"); + } cl_uint base_align_in_bits; CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_MEM_BASE_ADDR_ALIGN, sizeof(cl_uint), &base_align_in_bits, NULL)); @@ -4423,6 +4530,15 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32); case GGML_OP_SSM_CONV: return (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32); + case GGML_OP_GATED_DELTA_NET: + { + // Match the Vulkan backend: only F32 -> F32, S_v in {16, 32, 64, 128}. + if (op->src[0]->type != GGML_TYPE_F32 || op->type != GGML_TYPE_F32) { + return false; + } + const int64_t S_v = op->src[2]->ne[0]; + return S_v == 16 || S_v == 32 || S_v == 64 || S_v == 128; + } case GGML_OP_CONCAT: return op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32; case GGML_OP_TIMESTEP_EMBEDDING: @@ -14916,6 +15032,180 @@ static void ggml_cl_glu(ggml_backend_t backend, const ggml_tensor * src0, const backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); } +static void ggml_cl_gated_delta_net(ggml_backend_t backend, ggml_tensor * dst) { + GGML_ASSERT(dst); + GGML_ASSERT(dst->extra); + + const ggml_tensor * src_q = dst->src[0]; + const ggml_tensor * src_k = dst->src[1]; + const ggml_tensor * src_v = dst->src[2]; + const ggml_tensor * src_g = dst->src[3]; + const ggml_tensor * src_beta = dst->src[4]; + const ggml_tensor * src_state = dst->src[5]; + + GGML_ASSERT(src_q && src_q->extra); + GGML_ASSERT(src_k && src_k->extra); + GGML_ASSERT(src_v && src_v->extra); + GGML_ASSERT(src_g && src_g->extra); + GGML_ASSERT(src_beta && src_beta->extra); + GGML_ASSERT(src_state && src_state->extra); + + ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *) backend->context; + + const cl_uint S_v = (cl_uint) src_v->ne[0]; + const cl_uint H_v = (cl_uint) src_v->ne[1]; + const cl_uint n_tokens = (cl_uint) src_v->ne[2]; + const cl_uint n_seqs = (cl_uint) src_v->ne[3]; + + int si; + switch (S_v) { + case 16: si = 0; break; + case 32: si = 1; break; + case 64: si = 2; break; + case 128: si = 3; break; + default: + GGML_ASSERT(false && "ggml_cl_gated_delta_net: unsupported S_v"); + return; + } + + const int kda = (src_g->ne[0] == (int64_t) S_v) ? 1 : 0; + + // TODO: Optimize when S_v!=128. Not necessary for now as Qwen3.5/6 are all S_v=128 + // token generation mode (tgpp=0): + // process 1 token at a time, so columns per lane (cpl) == 1 + // prompt processing mode (tgpp=1): + // cpl=4 to process 4 tokens for single-token. 4 is chosen for Adreno 750 as per + // work-item/thread has at most 128 registers. + // All Qwen3.5/6 models are S_v == 128, so LANES_PER_COLUMN == 8 + // such that ROWS_PER_LANE = 128/8 = 16 + // Variables in the kernel: + // k_reg, q_reg, g_exp are all 16 floats + // s_shard has cpl*ROWS_PER_LANE = 4*16 = 64 floats + // Total 112 registers used. + // subgroups_per_workgroup (spw) can be set to 1,2,4,8,16 for tg and 1,2,4 for pp + // for S_v=128. + // Empirically found that when spw=1, we get the best performance for both tg and pp + const int tgpp = (n_tokens == 1) ? 0 : 1; + const int cpl = (tgpp == 0) ? 1 : 4; + // spw needs adjustment when S_v != 128 + const int spw = (tgpp == 0) ? 1 : 1; + + cl_kernel kernel = backend_ctx->kernel_gated_delta_net_f32[si][kda][tgpp]; + GGML_ASSERT(kernel != nullptr); + + const cl_uint s_off = S_v * H_v * n_tokens * n_seqs; + + const cl_uint sq1 = (cl_uint)(src_q->nb[1] / sizeof(float)); + const cl_uint sq2 = (cl_uint)(src_q->nb[2] / sizeof(float)); + const cl_uint sq3 = (cl_uint)(src_q->nb[3] / sizeof(float)); + const cl_uint sv1 = (cl_uint)(src_v->nb[1] / sizeof(float)); + const cl_uint sv2 = (cl_uint)(src_v->nb[2] / sizeof(float)); + const cl_uint sv3 = (cl_uint)(src_v->nb[3] / sizeof(float)); + const cl_uint sb1 = (cl_uint)(src_beta->nb[1] / sizeof(float)); + const cl_uint sb2 = (cl_uint)(src_beta->nb[2] / sizeof(float)); + const cl_uint sb3 = (cl_uint)(src_beta->nb[3] / sizeof(float)); + + const cl_uint H_k = (cl_uint) src_q->ne[1]; + const cl_uint rq3 = (cl_uint)(src_v->ne[3] / src_q->ne[3]); + + const float scale = 1.0f / sqrtf((float) S_v); + + ggml_tensor_extra_cl * extra_q = (ggml_tensor_extra_cl *) src_q->extra; + ggml_tensor_extra_cl * extra_k = (ggml_tensor_extra_cl *) src_k->extra; + ggml_tensor_extra_cl * extra_v = (ggml_tensor_extra_cl *) src_v->extra; + ggml_tensor_extra_cl * extra_g = (ggml_tensor_extra_cl *) src_g->extra; + ggml_tensor_extra_cl * extra_beta = (ggml_tensor_extra_cl *) src_beta->extra; + ggml_tensor_extra_cl * extra_state = (ggml_tensor_extra_cl *) src_state->extra; + ggml_tensor_extra_cl * extra_dst = (ggml_tensor_extra_cl *) dst->extra; + + const cl_ulong off_q = extra_q->offset + src_q->view_offs; + const cl_ulong off_k = extra_k->offset + src_k->view_offs; + const cl_ulong off_v = extra_v->offset + src_v->view_offs; + const cl_ulong off_g = extra_g->offset + src_g->view_offs; + const cl_ulong off_beta = extra_beta->offset + src_beta->view_offs; + const cl_ulong off_state = extra_state->offset + src_state->view_offs; + const cl_ulong off_dst = extra_dst->offset + dst->view_offs; + + int idx = 0; + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem), &extra_q->data_device)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_ulong), &off_q)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem), &extra_k->data_device)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_ulong), &off_k)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem), &extra_v->data_device)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_ulong), &off_v)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem), &extra_g->data_device)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_ulong), &off_g)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem), &extra_beta->data_device)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_ulong), &off_beta)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem), &extra_state->data_device)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_ulong), &off_state)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem), &extra_dst->data_device)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_ulong), &off_dst)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &H_v)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &n_tokens)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &n_seqs)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &s_off)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &sq1)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &sq2)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &sq3)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &sv1)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &sv2)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &sv3)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &sb1)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &sb2)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &sb3)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &H_k)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &rq3)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(float), &scale)); + + // Subgroup size is 64 for Adreno 750 + const int sg_size = backend_ctx->subgroup_size > 0 ? backend_ctx->subgroup_size : 32; + + // For the subgroup-shuffle kernel, we can safely prefer 8 lanes/column for S_v>=128 + // For the subgroup-shuffle kernel: + // S_v >= 128 -> prefer 8 lanes/column (good occupancy & register pressure tradeoff) + // else -> min(S_v, subgroup_size) + int lanes_per_column; + if ((int)S_v >= 128) { + lanes_per_column = 8; + } else { + lanes_per_column = std::min((int)S_v, sg_size); + } + + // Max workgroup size for Adreno 750 is 1024 + const int wg_size = sg_size * spw; + + // Ensure lanes_per_column is a power-of-two and divides both S_v and subgroup_size. + // (Required for lane-group shuffle-xor reduction correctness.) + while (lanes_per_column > 1 && + (((lanes_per_column & (lanes_per_column - 1)) != 0) || + (((int)S_v % lanes_per_column) != 0) || + (sg_size % lanes_per_column) != 0)) { + lanes_per_column >>= 1; + } + GGML_ASSERT(lanes_per_column >= 1); + GGML_ASSERT(((lanes_per_column & (lanes_per_column - 1)) == 0)); + GGML_ASSERT(((int)S_v % lanes_per_column) == 0); + GGML_ASSERT((sg_size % lanes_per_column) == 0); + + const int cols_per_wg = spw * (sg_size / lanes_per_column) * cpl; + GGML_ASSERT(cols_per_wg > 0); + GGML_ASSERT(((int)S_v % cols_per_wg) == 0); + + size_t global_work_size[3]; + size_t local_work_size[3]; + + global_work_size[0] = (size_t) H_v * (size_t) wg_size; + global_work_size[1] = (size_t) n_seqs; + global_work_size[2] = (size_t) S_v / (size_t) cols_per_wg; + + local_work_size[0] = (size_t) wg_size; + local_work_size[1] = 1; + local_work_size[2] = 1; + + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); +} + //------------------------------------------------------------------------------ // Op offloading //------------------------------------------------------------------------------ @@ -15161,6 +15451,14 @@ bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor } func = ggml_cl_ssm_conv; break; + case GGML_OP_GATED_DELTA_NET: + if (!any_on_device) { + return false; + } + // GDN has 6 source tensors, so it cannot use the standard + // (src0, src1, dst) func signature. Dispatch directly and return. + ggml_cl_gated_delta_net(backend, tensor); + return true; case GGML_OP_CONCAT: if (!any_on_device) { return false; diff --git a/ggml/src/ggml-opencl/kernels/gated_delta_net.cl b/ggml/src/ggml-opencl/kernels/gated_delta_net.cl new file mode 100644 index 000000000000..ceca98d13d20 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/gated_delta_net.cl @@ -0,0 +1,209 @@ +#pragma OPENCL EXTENSION cl_khr_subgroups : enable + +#ifndef S_V +#define S_V 128 +#endif +#ifndef KDA +#define KDA 0 +#endif +#ifndef SUBGROUP_SIZE +#define SUBGROUP_SIZE 64 +#endif +#ifndef LANES_PER_COLUMN +#define LANES_PER_COLUMN 8 +#endif +#ifndef COLS_PER_LANE_GROUP +#define COLS_PER_LANE_GROUP 1 +#endif +#ifndef SUBGROUPS_PER_WG +#define SUBGROUPS_PER_WG 1 +#endif +#ifndef USE_QCOM_SUBGROUP_SHUFFLE +#define USE_QCOM_SUBGROUP_SHUFFLE 0 +#endif + +#define WG_SIZE (SUBGROUP_SIZE * SUBGROUPS_PER_WG) +#define LANE_GROUPS_PER_SG (SUBGROUP_SIZE / LANES_PER_COLUMN) +#define COLS_PER_SG (LANE_GROUPS_PER_SG * COLS_PER_LANE_GROUP) +#define COLS_PER_WG (SUBGROUPS_PER_WG * COLS_PER_SG) +#define ROWS_PER_LANE (S_V / LANES_PER_COLUMN) + +#if USE_QCOM_SUBGROUP_SHUFFLE +#pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable +#endif + +// XOR-based parallel sum +// This does a reduction across groups of LANES_PER_COLUMN +static inline float reduce_add_shmem(float partial, __local float * temp, uint lane) { +#if USE_QCOM_SUBGROUP_SHUFFLE + #pragma unroll + for (uint s = LANES_PER_COLUMN / 2u; s > 0u; s >>= 1u) { + partial += qcom_sub_group_shuffle_xor(partial, s, CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM, partial); + } + return partial; +#else + temp[lane] = partial; + sub_group_barrier(CLK_LOCAL_MEM_FENCE); + #pragma unroll + for (uint s = LANES_PER_COLUMN / 2u; s > 0u; s >>= 1u) { + float other = temp[lane ^ s]; + sub_group_barrier(CLK_LOCAL_MEM_FENCE); + temp[lane] += other; + sub_group_barrier(CLK_LOCAL_MEM_FENCE); + } + const float result = temp[lane]; + sub_group_barrier(CLK_LOCAL_MEM_FENCE); + return result; +#endif +} + +#define REDUCE_PARTIAL(partial, temp_ptr, lid) \ + ((LANES_PER_COLUMN == 1u) ? (partial) : reduce_add_shmem((partial), (temp_ptr), (lid))) + +// force compiler to optimize kernel for a specific fixed work-group size +__attribute__((reqd_work_group_size(WG_SIZE, 1, 1))) +kernel void kernel_gated_delta_net( + global const char * q_buf, ulong off_q, + global const char * k_buf, ulong off_k, + global const char * v_buf, ulong off_v, + global const char * g_buf, ulong off_g, + global const char * beta_buf, ulong off_beta, + global const char * state_buf, ulong off_state, + global char * dst_buf, ulong off_dst, + uint H_v, + uint n_tokens, + uint n_seqs, + uint s_off, + uint sq1, uint sq2, uint sq3, + uint sv1, uint sv2, uint sv3, + uint sb1, uint sb2, uint sb3, + uint H_k, + uint rq3, + float scale) { + + global const float * data_q = (global const float *)(q_buf + off_q); + global const float * data_k = (global const float *)(k_buf + off_k); + global const float * data_v = (global const float *)(v_buf + off_v); + global const float * data_g = (global const float *)(g_buf + off_g); + global const float * data_beta = (global const float *)(beta_buf + off_beta); + global const float * data_state = (global const float *)(state_buf + off_state); + global float * data_dst = (global float *)(dst_buf + off_dst); + + const uint head_id = get_group_id(0); + const uint seq_id = get_group_id(1); + const uint tid = (uint)get_local_id(0); + + const uint sg_id = get_sub_group_id(); // subgroup id + const uint sg_lid = get_sub_group_local_id(); // subgroup lane id + + const uint lane = sg_lid % LANES_PER_COLUMN; + const uint lane_group = sg_lid / LANES_PER_COLUMN; + const uint wg_col_base = get_group_id(2) * COLS_PER_WG; + const uint sg_col_base = wg_col_base + sg_id * COLS_PER_SG; + + const uint iq1 = head_id % H_k; // head index for Q and K + const uint iq3 = seq_id / rq3; // seq index for Q and K + + const uint state_size = S_V * S_V; + const uint state_base = (seq_id * H_v + head_id) * state_size; + const uint q_off_base = iq3 * sq3 + iq1 * sq1; + const uint v_off_base = seq_id * sv3 + head_id * sv1; + const uint gb_off_base = seq_id * sb3 + head_id * sb1; + + __local float reduce_temp[WG_SIZE]; + __local float * temp_ptr = reduce_temp + sg_id * SUBGROUP_SIZE; + + float s_shard[COLS_PER_LANE_GROUP][ROWS_PER_LANE]; + #pragma unroll + for (uint cg = 0; cg < COLS_PER_LANE_GROUP; cg++) { + const uint col = sg_col_base + cg * LANE_GROUPS_PER_SG + lane_group; + #pragma unroll + for (uint r = 0; r < ROWS_PER_LANE; r++) { + s_shard[cg][r] = data_state[state_base + col * S_V + r * LANES_PER_COLUMN + lane]; + } + } + + uint attn_off = (seq_id * n_tokens * H_v + head_id) * S_V; + + for (uint t = 0; t < n_tokens; t++) { + const uint q_off = q_off_base + t * sq2; + const uint k_off = q_off; + const uint v_off = v_off_base + t * sv2; + const uint gb_off = gb_off_base + t * sb2; + const float beta_val = data_beta[gb_off]; + + float k_reg[ROWS_PER_LANE]; + float q_reg[ROWS_PER_LANE]; +#if KDA + float g_exp[ROWS_PER_LANE]; + #pragma unroll + for (uint r = 0; r < ROWS_PER_LANE; r++) { + const uint i = r * LANES_PER_COLUMN + lane; + k_reg[r] = data_k[k_off + i]; + q_reg[r] = data_q[q_off + i]; + g_exp[r] = exp(data_g[gb_off * S_V + i]); + } +#else + const float g_val = exp(data_g[gb_off]); + + #pragma unroll + for (uint r = 0; r < ROWS_PER_LANE; r++) { + const uint i = r * LANES_PER_COLUMN + lane; + k_reg[r] = data_k[k_off + i]; + q_reg[r] = data_q[q_off + i]; + } +#endif + + #pragma unroll + for (uint cg = 0; cg < COLS_PER_LANE_GROUP; cg++) { + const uint col = sg_col_base + cg * LANE_GROUPS_PER_SG + lane_group; + float v_val = data_v[v_off + col]; + + float kv_shard = 0.0f; + #pragma unroll + for (uint r = 0; r < ROWS_PER_LANE; r++) { +#if KDA + float gs = g_exp[r] * s_shard[cg][r]; + kv_shard += gs * k_reg[r]; +#else + kv_shard += s_shard[cg][r] * k_reg[r]; +#endif + } + +#if !KDA + kv_shard *= g_val; // Applied once instead of ROWS_PER_LANE times +#endif + + const float kv_col = REDUCE_PARTIAL(kv_shard, temp_ptr, sg_lid); + + const float delta_col = (v_val - kv_col) * beta_val; + + float attn_partial = 0.0f; + #pragma unroll + for (uint r = 0; r < ROWS_PER_LANE; r++) { +#if KDA + float gs = g_exp[r] * s_shard[cg][r]; +#else + float gs = g_val * s_shard[cg][r]; +#endif + s_shard[cg][r] = gs + k_reg[r] * delta_col; + attn_partial += s_shard[cg][r] * q_reg[r]; + } + const float attn_col = REDUCE_PARTIAL(attn_partial, temp_ptr, sg_lid); + + if (lane == 0) { + data_dst[attn_off + col] = attn_col * scale; + } + } + attn_off += S_V * H_v; + } + + #pragma unroll + for (uint cg = 0; cg < COLS_PER_LANE_GROUP; cg++) { + const uint col = sg_col_base + cg * LANE_GROUPS_PER_SG + lane_group; + #pragma unroll + for (uint r = 0; r < ROWS_PER_LANE; r++) { + data_dst[s_off + state_base + col * S_V + r * LANES_PER_COLUMN + lane] = s_shard[cg][r]; + } + } +} From 748c76ff7b39424ba485e7adfbb8ead2cb212c95 Mon Sep 17 00:00:00 2001 From: Yee Man Chan Date: Tue, 19 May 2026 08:50:57 +0800 Subject: [PATCH 2/7] add back lanes_per_column declaration --- ggml/src/ggml-opencl/ggml-opencl.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 18a212fed17c..a5bf212f0070 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -2498,6 +2498,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve const int S_V = gdn_sizes[si]; // MUST match the dispatcher heuristic in ggml_cl_gated_delta_net exactly. + int lanes_per_column; if (S_V >= 128) { lanes_per_column = 8; } else { From f12e625ff9b885dba491c7cc62ccb1024ef67e16 Mon Sep 17 00:00:00 2001 From: Yee Man Chan Date: Tue, 19 May 2026 11:33:57 +0800 Subject: [PATCH 3/7] removed has_subgroup_arithmetic and has_subgroup_clustered_reduce --- ggml/src/ggml-opencl/ggml-opencl.cpp | 23 +++-------------------- 1 file changed, 3 insertions(+), 20 deletions(-) diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index a5bf212f0070..bffb2bdb0e9e 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -395,8 +395,6 @@ struct ggml_backend_opencl_context { size_t max_workgroup_size; bool fp16_support; bool has_vector_subgroup_broadcast; - bool has_subgroup_arithmetic = false; // cl_khr_subgroup_arithmetic - bool has_subgroup_clustered_reduce = false; // cl_khr_subgroup_clustered_reduce bool has_qcom_subgroup_shuffle = false; // cl_qcom_subgroup_shuffle int subgroup_size = 0; // wavefront / subgroup width bool disable_fusion; @@ -605,8 +603,8 @@ struct ggml_backend_opencl_context { cl_kernel kernel_conv_2d_f32; cl_kernel kernel_conv_2d_f16_f32; cl_kernel kernel_ssm_conv_f32_f32, kernel_ssm_conv_f32_f32_4; - // [size_idx][kda][tile] where size_idx: 0=S_V=16, 1=32, 2=64, 3=128; kda: 0 or 1. - // tile 0 = TG variant (COLS_PER_LANE_GROUP=1), tile 1 = prefill variant (COLS_PER_LANE_GROUP=4). + // [size_idx][kda][tgpp] where size_idx: 0=S_V=16, 1=32, 2=64, 3=128; kda: 0 or 1. + // tgpp 0 = TG variant (COLS_PER_LANE_GROUP=1), tgpp 1 = prefill variant (COLS_PER_LANE_GROUP=4). cl_kernel kernel_gated_delta_net_f32[4][2][2] = {}; cl_kernel kernel_timestep_embedding; @@ -3527,21 +3525,6 @@ static ggml_backend_opencl_context * ggml_cl2_init(ggml_backend_dev_t dev) { // check Adreno large buffer support backend_ctx->adreno_has_large_buffer = strstr(ext_buffer, "cl_qcom_large_buffer") != NULL; - // Detect subgroup arithmetic and clustered-reduce support. These are - // optional extensions; the gated_delta_net kernel uses them when - // available and falls back to a local-memory reduction otherwise. - backend_ctx->has_subgroup_arithmetic = strstr(ext_buffer, "cl_khr_subgroup_arithmetic") != NULL; - backend_ctx->has_subgroup_clustered_reduce = strstr(ext_buffer, "cl_khr_subgroup_clustered_reduce") != NULL; - // The kernel needs a subgroup-size hint at build time. Use the - // Adreno wave size when we already know it; otherwise default to 32 - // (a common subgroup size for Intel/Mali). The local-memory fallback - // is correct for any device whose actual subgroup_size divides this. - backend_ctx->subgroup_size = backend_ctx->gpu_family == ADRENO ? backend_ctx->adreno_wave_size : 32; - GGML_LOG_INFO("ggml_opencl: subgroup arithmetic support: %s\n", - backend_ctx->has_subgroup_arithmetic ? "true" : "false"); - GGML_LOG_INFO("ggml_opencl: subgroup clustered reduce support: %s\n", - backend_ctx->has_subgroup_clustered_reduce ? "true" : "false"); - // fp16 is required if (!backend_ctx->fp16_support) { GGML_LOG_ERROR("ggml_opencl: device does not support FP16\n"); @@ -3552,7 +3535,7 @@ static ggml_backend_opencl_context * ggml_cl2_init(ggml_backend_dev_t dev) { // optional in OpenCL 3.0 (cl_khr_subgroup is mandatory in OpenCL 2.x) if (opencl_c_version.major == 3 && strstr(ext_buffer, "cl_khr_subgroups") == NULL && strstr(ext_buffer, "cl_intel_subgroups") == NULL) { - GGML_LOG_INFO("ggml_opencl: device does not support subgroups (cl_khr_subgroups or cl_intel_subgroups) " + GGML_LOG_ERROR("ggml_opencl: device does not support subgroups (cl_khr_subgroups or cl_intel_subgroups) " "(note that subgroups is an optional feature in OpenCL 3.0)\n"); return nullptr; } From ac971b6589d2c283647c82df6fb1aa1941109742 Mon Sep 17 00:00:00 2001 From: Yee Man Chan Date: Thu, 21 May 2026 14:45:25 +0800 Subject: [PATCH 4/7] removed trailing spaces and fixes indentation. Hard coded subgroup size for Adreno and Intel. Return not supported when K>1 state snapshot --- ggml/src/ggml-opencl/ggml-opencl.cpp | 172 ++++++++++++++------------- 1 file changed, 91 insertions(+), 81 deletions(-) diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 14d018c718ee..61b9e4c005fb 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -396,7 +396,6 @@ struct ggml_backend_opencl_context { bool fp16_support; bool has_vector_subgroup_broadcast; bool has_qcom_subgroup_shuffle = false; // cl_qcom_subgroup_shuffle - int subgroup_size = 0; // wavefront / subgroup width bool disable_fusion; std::regex *opfilter = nullptr; // regex of ops to not claim @@ -518,8 +517,7 @@ struct ggml_backend_opencl_context { cl_kernel kernel_tri; cl_kernel kernel_fill; cl_kernel kernel_clamp; - cl_kernel kernel_geglu, kernel_reglu, kernel_swiglu, kernel_swiglu_oai, kernel_geglu_erf, kernel_geglu_quick, - kernel_geglu_f16, kernel_reglu_f16, kernel_swiglu_f16, kernel_geglu_erf_f16, kernel_geglu_quick_f16; + cl_kernel kernel_geglu, kernel_reglu, kernel_swiglu, kernel_swiglu_oai, kernel_geglu_erf, kernel_geglu_quick, kernel_geglu_f16, kernel_reglu_f16, kernel_swiglu_f16, kernel_geglu_erf_f16, kernel_geglu_quick_f16; cl_kernel kernel_norm, kernel_norm_mul_add; cl_kernel kernel_rms_norm, kernel_rms_norm_mul; cl_kernel kernel_l2_norm_f32; @@ -795,16 +793,16 @@ struct ggml_backend_opencl_context { static std::vector g_ggml_backend_opencl_devices; inline std::string read_file(const std::string &path) { - std::ifstream ifs(path); - if (!ifs) { - return ""; - } - std::string text; - ifs.seekg(0, std::ios::end); - text.resize(ifs.tellg()); - ifs.seekg(0, std::ios::beg); - ifs.read(&text[0], text.size()); - return text; + std::ifstream ifs(path); + if (!ifs) { + return ""; + } + std::string text; + ifs.seekg(0, std::ios::end); + text.resize(ifs.tellg()); + ifs.seekg(0, std::ios::beg); + ifs.read(&text[0], text.size()); + return text; } static cl_program build_program_from_source(cl_context ctx, cl_device_id dev, const char* program_buffer, const std::string &compile_opts) { @@ -843,8 +841,8 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve auto opencl_c_std = std::string("CL") + std::to_string(opencl_c_version.major) + "." + std::to_string(opencl_c_version.minor); std::string compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable -cl-unsafe-math-optimizations" - " -cl-finite-math-only -cl-fast-relaxed-math"; + " -cl-mad-enable -cl-unsafe-math-optimizations" + " -cl-finite-math-only -cl-fast-relaxed-math"; if (backend_ctx->adreno_use_large_buffer) { compile_opts += " -qcom-enable-large-buffer "; @@ -2073,7 +2071,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve const std::string kernel_src = read_file("div.cl"); #endif std::string compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable -cl-finite-math-only "; + " -cl-mad-enable -cl-finite-math-only "; backend_ctx->program_div = build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts); @@ -2390,12 +2388,12 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts); CL_CHECK((backend_ctx->kernel_upscale = clCreateKernel(backend_ctx->program_upscale, "kernel_upscale", &err), err)); if (backend_ctx->program_upscale) { - cl_int err_bilinear; - backend_ctx->kernel_upscale_bilinear = clCreateKernel(backend_ctx->program_upscale, "kernel_upscale_bilinear", &err_bilinear); - if (err_bilinear != CL_SUCCESS) { + cl_int err_bilinear; + backend_ctx->kernel_upscale_bilinear = clCreateKernel(backend_ctx->program_upscale, "kernel_upscale_bilinear", &err_bilinear); + if (err_bilinear != CL_SUCCESS) { GGML_LOG_WARN("ggml_opencl: kernel_upscale_bilinear not found in upscale.cl. Bilinear upscale will not be available. Error: %d\n", err_bilinear); backend_ctx->kernel_upscale_bilinear = nullptr; - } + } } else { backend_ctx->kernel_upscale_bilinear = nullptr; } @@ -2465,8 +2463,8 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve GGML_LOG_CONT("."); } - // conv2d - { + // conv2d + { #ifdef GGML_OPENCL_EMBED_KERNELS const std::string kernel_src { #include "conv2d.cl.h" @@ -2535,7 +2533,11 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve #endif const int gdn_sizes[4] = { 16, 32, 64, 128 }; - const int sg_size = backend_ctx->subgroup_size > 0 ? backend_ctx->subgroup_size : 32; + const int sg_size = backend_ctx->gpu_family == GPU_FAMILY::ADRENO ? 64 : backend_ctx->gpu_family == GPU_FAMILY::INTEL ? 32 : -1; + if (sg_size < 0) { + GGML_LOG_ERROR("Unsupported GPU Family: only Adreno and Intel are supported.\n"); + exit(1); + } for (int si = 0; si < 4; si++) { const int S_V = gdn_sizes[si]; @@ -2552,7 +2554,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve // * power-of-two // * divides both S_V and sg_size while (lanes_per_column > 1 && - (((lanes_per_column & (lanes_per_column - 1)) != 0) || + (((lanes_per_column & (lanes_per_column - 1)) != 0) || (S_V % lanes_per_column) != 0 || (sg_size % lanes_per_column) != 0)) { lanes_per_column >>= 1; @@ -2592,7 +2594,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve backend_ctx->context, backend_ctx->device, kernel_src.c_str(), opts); CL_CHECK((backend_ctx->kernel_gated_delta_net_f32[si][kda][tgpp] = - clCreateKernel(prog, "kernel_gated_delta_net", &err), err)); + clCreateKernel(prog, "kernel_gated_delta_net", &err), err)); CL_CHECK(clReleaseProgram(prog)); } } @@ -2707,9 +2709,9 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve // gemv_noshuffle_general { std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable " - " -DSIMDGROUP_WIDTH=" + - std::to_string(backend_ctx->adreno_wave_size); + " -cl-mad-enable " + " -DSIMDGROUP_WIDTH=" + + std::to_string(backend_ctx->adreno_wave_size); if (backend_ctx->has_vector_subgroup_broadcast) { CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAST "; } @@ -2830,7 +2832,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve #ifdef GGML_OPENCL_EMBED_KERNELS const std::string kernel_src { #include "gemm_noshuffle_q4_1_f32.cl.h" - }; + }; #else const std::string kernel_src = read_file("gemm_noshuffle_q4_1_f32.cl"); #endif @@ -2843,7 +2845,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve // gemv_noshuffle_q4_1_f32 { std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable "; + " -cl-mad-enable "; if (backend_ctx->has_vector_subgroup_broadcast) { CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAT "; } @@ -2869,7 +2871,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve #ifdef GGML_OPENCL_EMBED_KERNELS const std::string kernel_src { #include "gemm_noshuffle_iq4_nl_f32.cl.h" - }; + }; #else const std::string kernel_src = read_file("gemm_noshuffle_iq4_nl_f32.cl"); #endif @@ -2882,7 +2884,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve // gemv_noshuffle_iq4_nl_f32 { std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable "; + " -cl-mad-enable "; if (backend_ctx->has_vector_subgroup_broadcast) { CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAST "; } @@ -2908,7 +2910,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve #ifdef GGML_OPENCL_EMBED_KERNELS const std::string kernel_src { #include "gemm_noshuffle_q8_0_f32.cl.h" - }; + }; #else const std::string kernel_src = read_file("gemm_noshuffle_q8_0_f32.cl"); #endif @@ -2921,9 +2923,9 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve // gemv_noshuffle_general_q8_0_f32 { std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable " - " -DSIMDGROUP_WIDTH=" + - std::to_string(backend_ctx->adreno_wave_size); + " -cl-mad-enable " + " -DSIMDGROUP_WIDTH=" + + std::to_string(backend_ctx->adreno_wave_size); if (backend_ctx->has_vector_subgroup_broadcast) { CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAT "; } @@ -2949,7 +2951,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve #ifdef GGML_OPENCL_EMBED_KERNELS const std::string kernel_src { #include "gemm_noshuffle_q4_k_f32.cl.h" - }; + }; #else const std::string kernel_src = read_file("gemm_noshuffle_q4_k_f32.cl"); #endif @@ -2962,7 +2964,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve // gemv_noshuffle_q4_k_f32 { std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable "; + " -cl-mad-enable "; if (backend_ctx->has_vector_subgroup_broadcast) { CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAST "; } @@ -3235,7 +3237,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve #endif std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable "; + " -cl-mad-enable "; if (backend_ctx->has_vector_subgroup_broadcast) { CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAT "; } @@ -3266,7 +3268,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve // gemv_noshuffle_q5_k_f32 { std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable "; + " -cl-mad-enable "; if (backend_ctx->has_vector_subgroup_broadcast) { CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAST "; } @@ -3534,11 +3536,11 @@ static std::vector ggml_opencl_probe_devices(ggml_backend_r if (found_devices.size()) { auto * dev_ctx = static_cast(found_devices.front().context); GGML_LOG_INFO("ggml_opencl: default device: '%s (%s)'\n", dev_ctx->device_name.c_str(), - dev_ctx->device_version.c_str()); + dev_ctx->device_version.c_str()); if (dev_ctx->device_type != CL_DEVICE_TYPE_GPU) { GGML_LOG_WARN("ggml_opencl: warning, the default device is not a GPU: '%s'.\n", - dev_ctx->device_name.c_str()); + dev_ctx->device_name.c_str()); } } @@ -3702,7 +3704,7 @@ static ggml_backend_opencl_context * ggml_cl2_init(ggml_backend_dev_t dev) { backend_ctx->non_uniform_workgroups = false; #if CL_TARGET_OPENCL_VERSION >= 300 CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT, sizeof(cl_bool), - &backend_ctx->non_uniform_workgroups, 0)); + &backend_ctx->non_uniform_workgroups, 0)); #endif } else { GGML_ASSERT(opencl_c_version.major == 2); @@ -3721,17 +3723,17 @@ static ggml_backend_opencl_context * ggml_cl2_init(ggml_backend_dev_t dev) { #ifdef GGML_OPENCL_USE_ADRENO_KERNELS backend_ctx->adreno_xmem_gemm_enabled = getenv("GGML_OPENCL_ADRENO_XMEM_GEMM") != nullptr && - backend_ctx->gpu_family == GPU_FAMILY::ADRENO; + backend_ctx->gpu_family == GPU_FAMILY::ADRENO; if (getenv("GGML_OPENCL_ADRENO_XMEM_GEMM") != nullptr) { GGML_LOG_INFO("ggml_opencl: Adreno xmem F16xF32 GEMM %s\n", - backend_ctx->adreno_xmem_gemm_enabled ? - "enabled (temporary weight prepack)" : "requested but unsupported by this driver"); + backend_ctx->adreno_xmem_gemm_enabled ? + "enabled (temporary weight prepack)" : "requested but unsupported by this driver"); } #endif // GGML_OPENCL_USE_ADRENO_KERNELS // determine whether to use large buffer for Adreno backend_ctx->adreno_use_large_buffer = getenv("GGML_OPENCL_ADRENO_USE_LARGE_BUFFER") != nullptr && - backend_ctx->gpu_family == GPU_FAMILY::ADRENO; + backend_ctx->gpu_family == GPU_FAMILY::ADRENO; if (backend_ctx->adreno_use_large_buffer) { if (!backend_ctx->adreno_has_large_buffer) { GGML_LOG_INFO("ggml_opencl: Adreno large buffer requested but not supported by driver, will use regular buffer\n"); @@ -3771,15 +3773,15 @@ static ggml_backend_opencl_context * ggml_cl2_init(ggml_backend_dev_t dev) { size_t max_B_d_bytes = MIN(required_B_d_bytes, backend_ctx->max_alloc_size); if (required_A_q_d_bytes > backend_ctx->max_alloc_size) { GGML_LOG_WARN("ggml_opencl: A_q_d buffer size reduced from %zu to %zu due to device limitations.\n", - required_A_q_d_bytes, max_A_q_d_bytes); + required_A_q_d_bytes, max_A_q_d_bytes); } if (required_A_s_d_bytes > backend_ctx->max_alloc_size) { GGML_LOG_WARN("ggml_opencl: A_s_d buffer size reduced from %zu to %zu due to device limitations.\n", - required_A_s_d_bytes, max_A_s_d_bytes); + required_A_s_d_bytes, max_A_s_d_bytes); } if (required_B_d_bytes > backend_ctx->max_alloc_size) { GGML_LOG_WARN("ggml_opencl: B_d buffer size reduced from %zu to %zu due to device limitations.\n", - required_B_d_bytes, max_B_d_bytes); + required_B_d_bytes, max_B_d_bytes); } backend_ctx->prealloc_quant_trans.allocate(context, max_A_q_d_bytes); @@ -4420,7 +4422,7 @@ static void ggml_backend_opencl_synchronize(ggml_backend_t backend) { // completed. static void sync_with_other_backends(ggml_backend_opencl_context * backend_ctx) { if (g_ggml_backend_opencl_devices.size() < 2) - return; // No other devices to synchronize with. + return; // No other devices to synchronize with. std::vector events; events.reserve(g_ggml_backend_opencl_devices.size()); @@ -4568,7 +4570,7 @@ inline bool use_adreno_kernels(const ggml_backend_opencl_context *backend_ctx, c int64_t threshold_ne0 = 512; int64_t threshold_ne1 = 512; if (!backend_ctx->adreno_cl_compiler_version.newer_than_or_same(E031, 38, 11, 0) && - backend_ctx->adreno_cl_compiler_version.type != DX) { + backend_ctx->adreno_cl_compiler_version.type != DX) { threshold_ne0 = 128; threshold_ne1 = 128; } @@ -4683,8 +4685,8 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te case GGML_OP_DIV: case GGML_OP_SUB: return (op->src[0]->type == op->src[1]->type) && - (op->src[0]->type == op->type) && - (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16); + (op->src[0]->type == op->type) && + (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16); case GGML_OP_ADD_ID: return op->src[0]->type == GGML_TYPE_F32; case GGML_OP_SQR: @@ -4698,17 +4700,17 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te case GGML_UNARY_OP_RELU: case GGML_UNARY_OP_GELU_ERF: case GGML_UNARY_OP_GELU_QUICK: - return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32; + return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32; case GGML_UNARY_OP_SIGMOID: return ggml_is_contiguous(op->src[0]); case GGML_UNARY_OP_TANH: case GGML_UNARY_OP_NEG: case GGML_UNARY_OP_EXP: - return op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16; + return op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16; case GGML_UNARY_OP_EXPM1: - return op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16; + return op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16; case GGML_UNARY_OP_SOFTPLUS: - return op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16; + return op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16; default: return false; } @@ -4749,22 +4751,26 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te ggml_scale_mode mode = (ggml_scale_mode)(ggml_get_op_params_i32(op, 0) & 0xFF); const bool antialias = (ggml_scale_mode)(ggml_get_op_params_i32(op, 0) & GGML_SCALE_FLAG_ANTIALIAS); return op->src[0]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 && - (mode == GGML_SCALE_MODE_NEAREST || mode == GGML_SCALE_MODE_BILINEAR) && !antialias; + (mode == GGML_SCALE_MODE_NEAREST || mode == GGML_SCALE_MODE_BILINEAR) && !antialias; } case GGML_OP_CONV_2D: return (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F16 && op->type == GGML_TYPE_F16) || - (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32) || - (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32); + (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32) || + (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32); case GGML_OP_SSM_CONV: return (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32); - case GGML_OP_GATED_DELTA_NET: + case GGML_OP_GATED_DELTA_NET: { // Match the Vulkan backend: only F32 -> F32, S_v in {16, 32, 64, 128}. if (op->src[0]->type != GGML_TYPE_F32 || op->type != GGML_TYPE_F32) { return false; } + // K>1 state snapshot is not supported + // This means state dimension must be S_v * S_v * H_v * n_seqs const int64_t S_v = op->src[2]->ne[0]; - return S_v == 16 || S_v == 32 || S_v == 64 || S_v == 128; + const int64_t H_v = op->src[2]->ne[1]; + const int64_t n_seqs = op->src[2]->ne[3]; + return (S_v == 16 || S_v == 32 || S_v == 64 || S_v == 128) && ggml_nelements(op->src[5]) == S_v * S_v * H_v * n_seqs; } case GGML_OP_CONCAT: return op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32; @@ -4778,11 +4784,11 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te } else if (op->src[0]->type == GGML_TYPE_F32) { return op->src[1]->type == GGML_TYPE_F32; } else if (op->src[0]->type == GGML_TYPE_Q4_0 || op->src[0]->type == GGML_TYPE_Q4_1 || - op->src[0]->type == GGML_TYPE_MXFP4 || - op->src[0]->type == GGML_TYPE_IQ4_NL || - op->src[0]->type == GGML_TYPE_Q4_K || - op->src[0]->type == GGML_TYPE_Q5_K || - op->src[0]->type == GGML_TYPE_Q6_K) { + op->src[0]->type == GGML_TYPE_MXFP4 || + op->src[0]->type == GGML_TYPE_IQ4_NL || + op->src[0]->type == GGML_TYPE_Q4_K || + op->src[0]->type == GGML_TYPE_Q5_K || + op->src[0]->type == GGML_TYPE_Q6_K) { 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) { return op->src[1]->type == GGML_TYPE_F32; @@ -7601,14 +7607,14 @@ static bool ggml_cl_can_mul_mat(const struct ggml_tensor * src0, const struct gg // TODO: find the optimal values for these return (src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) && src1->type == GGML_TYPE_F32 && - dst->type == GGML_TYPE_F32 && + dst->type == GGML_TYPE_F32 && (ne0 >= 32 && ne1 >= 32 && ne10 >= 32); } // Copy a noncontiguous tensor to contiguous tensor. ne[] remains the same but // nb[] is recalculated such that tensor is contiguous. static void ggml_cl_copy_to_contiguous(ggml_backend_t backend, const ggml_tensor * src, cl_mem dst, - cl_ulong &nb0, cl_ulong &nb1, cl_ulong &nb2, cl_ulong &nb3) { + cl_ulong &nb0, cl_ulong &nb1, cl_ulong &nb2, cl_ulong &nb3) { ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context; const int tensor_type_size = ggml_type_size(src->type); @@ -10156,7 +10162,7 @@ static void ggml_cl_pad(ggml_backend_t backend, const ggml_tensor * src0, ggml_t size_t local_work_size[] = { lws0, 1, 1 }; size_t * local_work_size_ptr = local_work_size; - if (d_ne0 % lws0 != 0 && !backend_ctx->non_uniform_workgroups) { + if (d_ne0 % lws0 != 0 && !backend_ctx->non_uniform_workgroups) { local_work_size_ptr = nullptr; } @@ -16160,19 +16166,19 @@ static void ggml_cl_gated_delta_net(ggml_backend_t backend, ggml_tensor * dst) { const int kda = (src_g->ne[0] == (int64_t) S_v) ? 1 : 0; // TODO: Optimize when S_v!=128. Not necessary for now as Qwen3.5/6 are all S_v=128 - // token generation mode (tgpp=0): + // token generation mode (tgpp=0): // process 1 token at a time, so columns per lane (cpl) == 1 - // prompt processing mode (tgpp=1): - // cpl=4 to process 4 tokens for single-token. 4 is chosen for Adreno 750 as per + // prompt processing mode (tgpp=1): + // cpl=4 to process 4 tokens for single-token. 4 is chosen for Adreno 750 as per // work-item/thread has at most 128 registers. // All Qwen3.5/6 models are S_v == 128, so LANES_PER_COLUMN == 8 // such that ROWS_PER_LANE = 128/8 = 16 // Variables in the kernel: // k_reg, q_reg, g_exp are all 16 floats // s_shard has cpl*ROWS_PER_LANE = 4*16 = 64 floats - // Total 112 registers used. + // Total 112 registers used. // subgroups_per_workgroup (spw) can be set to 1,2,4,8,16 for tg and 1,2,4 for pp - // for S_v=128. + // for S_v=128. // Empirically found that when spw=1, we get the best performance for both tg and pp const int tgpp = (n_tokens == 1) ? 0 : 1; const int cpl = (tgpp == 0) ? 1 : 4; @@ -16247,8 +16253,12 @@ static void ggml_cl_gated_delta_net(ggml_backend_t backend, ggml_tensor * dst) { CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &rq3)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(float), &scale)); - // Subgroup size is 64 for Adreno 750 - const int sg_size = backend_ctx->subgroup_size > 0 ? backend_ctx->subgroup_size : 32; + // Subgroup size is 64 for Adreno and 32 for Intel + const int sg_size = backend_ctx->gpu_family == GPU_FAMILY::ADRENO ? 64 : backend_ctx->gpu_family == GPU_FAMILY::INTEL ? 32 : -1; + if (sg_size < 0) { + GGML_LOG_ERROR("Unsupported GPU Family: only Adreno and Intel are supported.\n"); + exit(1); + } // For the subgroup-shuffle kernel, we can safely prefer 8 lanes/column for S_v>=128 // For the subgroup-shuffle kernel: @@ -16267,7 +16277,7 @@ static void ggml_cl_gated_delta_net(ggml_backend_t backend, ggml_tensor * dst) { // Ensure lanes_per_column is a power-of-two and divides both S_v and subgroup_size. // (Required for lane-group shuffle-xor reduction correctness.) while (lanes_per_column > 1 && - (((lanes_per_column & (lanes_per_column - 1)) != 0) || + (((lanes_per_column & (lanes_per_column - 1)) != 0) || (((int)S_v % lanes_per_column) != 0) || (sg_size % lanes_per_column) != 0)) { lanes_per_column >>= 1; @@ -16510,8 +16520,8 @@ bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor } func = ggml_cl_group_norm; break; - case GGML_OP_REPEAT: - if (!any_on_device) { + case GGML_OP_REPEAT: + if (!any_on_device) { return false; } func = ggml_cl_repeat; From 623fbf57728c15bf0e923be737fd876fc3a7267a Mon Sep 17 00:00:00 2001 From: Yee Man Chan Date: Thu, 21 May 2026 21:51:30 +0800 Subject: [PATCH 5/7] support for K>1 state snapshot --- ggml/src/ggml-opencl/ggml-opencl.cpp | 8 ++- .../ggml-opencl/kernels/gated_delta_net.cl | 52 ++++++++++++++++--- 2 files changed, 48 insertions(+), 12 deletions(-) diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 2c8a57311e11..7bed756a8d16 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -4971,12 +4971,8 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te if (op->src[0]->type != GGML_TYPE_F32 || op->type != GGML_TYPE_F32) { return false; } - // K>1 state snapshot is not supported - // This means state dimension must be S_v * S_v * H_v * n_seqs const int64_t S_v = op->src[2]->ne[0]; - const int64_t H_v = op->src[2]->ne[1]; - const int64_t n_seqs = op->src[2]->ne[3]; - return (S_v == 16 || S_v == 32 || S_v == 64 || S_v == 128) && ggml_nelements(op->src[5]) == S_v * S_v * H_v * n_seqs; + return S_v == 16 || S_v == 32 || S_v == 64 || S_v == 128; } case GGML_OP_CONCAT: return op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32; @@ -17163,6 +17159,7 @@ static void ggml_cl_gated_delta_net(ggml_backend_t backend, ggml_tensor * dst) { const cl_uint H_v = (cl_uint) src_v->ne[1]; const cl_uint n_tokens = (cl_uint) src_v->ne[2]; const cl_uint n_seqs = (cl_uint) src_v->ne[3]; + const cl_uint K = (cl_uint) src_state->ne[1]; int si; switch (S_v) { @@ -17264,6 +17261,7 @@ static void ggml_cl_gated_delta_net(ggml_backend_t backend, ggml_tensor * dst) { CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &H_k)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &rq3)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(float), &scale)); + CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &K)); // Subgroup size is 64 for Adreno and 32 for Intel const int sg_size = backend_ctx->gpu_family == GPU_FAMILY::ADRENO ? 64 : backend_ctx->gpu_family == GPU_FAMILY::INTEL ? 32 : -1; diff --git a/ggml/src/ggml-opencl/kernels/gated_delta_net.cl b/ggml/src/ggml-opencl/kernels/gated_delta_net.cl index ceca98d13d20..d11192f5802b 100644 --- a/ggml/src/ggml-opencl/kernels/gated_delta_net.cl +++ b/ggml/src/ggml-opencl/kernels/gated_delta_net.cl @@ -1,5 +1,17 @@ #pragma OPENCL EXTENSION cl_khr_subgroups : enable +#ifdef cl_intel_required_subgroup_size +#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable +#define INTEL_GPU 1 +#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16))) +#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32))) +#elif defined(cl_qcom_reqd_sub_group_size) +#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable +#define ADRENO_GPU 1 +#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half"))) +#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full"))) +#endif + #ifndef S_V #define S_V 128 #endif @@ -62,6 +74,11 @@ static inline float reduce_add_shmem(float partial, __local float * temp, uint l // force compiler to optimize kernel for a specific fixed work-group size __attribute__((reqd_work_group_size(WG_SIZE, 1, 1))) +#ifdef INTEL_GPU +REQD_SUBGROUP_SIZE_32 +#elif defined (ADRENO_GPU) +REQD_SUBGROUP_SIZE_64 +#endif kernel void kernel_gated_delta_net( global const char * q_buf, ulong off_q, global const char * k_buf, ulong off_k, @@ -79,7 +96,8 @@ kernel void kernel_gated_delta_net( uint sb1, uint sb2, uint sb3, uint H_k, uint rq3, - float scale) { + float scale, + uint K) { global const float * data_q = (global const float *)(q_buf + off_q); global const float * data_k = (global const float *)(k_buf + off_k); @@ -105,10 +123,12 @@ kernel void kernel_gated_delta_net( const uint iq3 = seq_id / rq3; // seq index for Q and K const uint state_size = S_V * S_V; - const uint state_base = (seq_id * H_v + head_id) * state_size; + const uint state_base = (seq_id * K * H_v + head_id) * state_size; const uint q_off_base = iq3 * sq3 + iq1 * sq1; const uint v_off_base = seq_id * sv3 + head_id * sv1; const uint gb_off_base = seq_id * sb3 + head_id * sb1; + const uint state_out_base = (seq_id * H_v + head_id) * state_size; + const uint state_size_per_snap = state_size * H_v * n_seqs; __local float reduce_temp[WG_SIZE]; __local float * temp_ptr = reduce_temp + sg_id * SUBGROUP_SIZE; @@ -123,6 +143,7 @@ kernel void kernel_gated_delta_net( } } + const int shift = (int)n_tokens - (int)K; uint attn_off = (seq_id * n_tokens * H_v + head_id) * S_V; for (uint t = 0; t < n_tokens; t++) { @@ -196,14 +217,31 @@ kernel void kernel_gated_delta_net( } } attn_off += S_V * H_v; + + if (K > 1u) { + const int target_slot = (int)t - shift; + if (target_slot >= 0 && target_slot < (int)K) { + #pragma unroll + for (uint cg = 0; cg < COLS_PER_LANE_GROUP; cg++) { + const uint col = sg_col_base + cg * LANE_GROUPS_PER_SG + lane_group; + const uint slot_base = s_off + (uint)target_slot * state_size_per_snap + state_out_base; + #pragma unroll + for (uint r = 0; r < ROWS_PER_LANE; r++) { + data_dst[slot_base + col * S_V + r * LANES_PER_COLUMN + lane] = s_shard[cg][r]; + } + } + } + } } - #pragma unroll - for (uint cg = 0; cg < COLS_PER_LANE_GROUP; cg++) { - const uint col = sg_col_base + cg * LANE_GROUPS_PER_SG + lane_group; + if (K == 1u) { #pragma unroll - for (uint r = 0; r < ROWS_PER_LANE; r++) { - data_dst[s_off + state_base + col * S_V + r * LANES_PER_COLUMN + lane] = s_shard[cg][r]; + for (uint cg = 0; cg < COLS_PER_LANE_GROUP; cg++) { + const uint col = sg_col_base + cg * LANE_GROUPS_PER_SG + lane_group; + #pragma unroll + for (uint r = 0; r < ROWS_PER_LANE; r++) { + data_dst[s_off + state_base + col * S_V + r * LANES_PER_COLUMN + lane] = s_shard[cg][r]; + } } } } From 32562126eb611c639aaeeebe29e8fc458239fa1e Mon Sep 17 00:00:00 2001 From: Yee Man Chan Date: Tue, 26 May 2026 17:09:18 +0800 Subject: [PATCH 6/7] removed picky indent multiple of 4 fixes --- ggml/src/ggml-opencl/ggml-opencl.cpp | 77 ++++++++++++++-------------- 1 file changed, 39 insertions(+), 38 deletions(-) diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 7bed756a8d16..8eea76c7cecc 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -534,7 +534,8 @@ struct ggml_backend_opencl_context { cl_kernel kernel_tri; cl_kernel kernel_fill; cl_kernel kernel_clamp; - cl_kernel kernel_geglu, kernel_reglu, kernel_swiglu, kernel_swiglu_oai, kernel_geglu_erf, kernel_geglu_quick, kernel_geglu_f16, kernel_reglu_f16, kernel_swiglu_f16, kernel_geglu_erf_f16, kernel_geglu_quick_f16; + cl_kernel kernel_geglu, kernel_reglu, kernel_swiglu, kernel_swiglu_oai, kernel_geglu_erf, kernel_geglu_quick, + kernel_geglu_f16, kernel_reglu_f16, kernel_swiglu_f16, kernel_geglu_erf_f16, kernel_geglu_quick_f16; cl_kernel kernel_norm, kernel_norm_mul_add; cl_kernel kernel_rms_norm, kernel_rms_norm_mul; cl_kernel kernel_l2_norm_f32; @@ -977,8 +978,8 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { auto opencl_c_std = std::string("CL") + std::to_string(backend_ctx->opencl_c_version.major) + "." + std::to_string(backend_ctx->opencl_c_version.minor); std::string compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable -cl-unsafe-math-optimizations" - " -cl-finite-math-only -cl-fast-relaxed-math"; + " -cl-mad-enable -cl-unsafe-math-optimizations" + " -cl-finite-math-only -cl-fast-relaxed-math"; if (backend_ctx->adreno_use_large_buffer) { compile_opts += " -qcom-enable-large-buffer "; @@ -2130,7 +2131,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { const std::string kernel_src = read_file("div.cl"); #endif std::string compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable -cl-finite-math-only "; + " -cl-mad-enable -cl-finite-math-only "; backend_ctx->program_div = build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts); @@ -2768,9 +2769,9 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { // gemv_noshuffle_general { std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable " - " -DSIMDGROUP_WIDTH=" + - std::to_string(backend_ctx->adreno_wave_size); + " -cl-mad-enable " + " -DSIMDGROUP_WIDTH=" + + std::to_string(backend_ctx->adreno_wave_size); if (backend_ctx->has_vector_subgroup_broadcast) { CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAST "; } @@ -2904,7 +2905,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { // gemv_noshuffle_q4_1_f32 { std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable "; + " -cl-mad-enable "; if (backend_ctx->has_vector_subgroup_broadcast) { CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAT "; } @@ -2943,7 +2944,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { // gemv_noshuffle_iq4_nl_f32 { std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable "; + " -cl-mad-enable "; if (backend_ctx->has_vector_subgroup_broadcast) { CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAST "; } @@ -2982,9 +2983,9 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { // gemv_noshuffle_general_q8_0_f32 { std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable " - " -DSIMDGROUP_WIDTH=" + - std::to_string(backend_ctx->adreno_wave_size); + " -cl-mad-enable " + " -DSIMDGROUP_WIDTH=" + + std::to_string(backend_ctx->adreno_wave_size); if (backend_ctx->has_vector_subgroup_broadcast) { CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAT "; } @@ -3023,7 +3024,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { // gemv_noshuffle_q4_k_f32 { std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable "; + " -cl-mad-enable "; if (backend_ctx->has_vector_subgroup_broadcast) { CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAST "; } @@ -3398,7 +3399,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { #endif std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable "; + " -cl-mad-enable "; if (backend_ctx->has_vector_subgroup_broadcast) { CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAT "; } @@ -3429,7 +3430,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { // gemv_noshuffle_q5_k_f32 { std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std + - " -cl-mad-enable "; + " -cl-mad-enable "; if (backend_ctx->has_vector_subgroup_broadcast) { CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAST "; } @@ -3699,11 +3700,11 @@ static std::vector ggml_opencl_probe_devices(ggml_backend_r if (found_devices.size()) { auto * dev_ctx = static_cast(found_devices.front().context); GGML_LOG_INFO("ggml_opencl: default device: '%s (%s)'\n", dev_ctx->device_name.c_str(), - dev_ctx->device_version.c_str()); + dev_ctx->device_version.c_str()); if (dev_ctx->device_type != CL_DEVICE_TYPE_GPU) { GGML_LOG_WARN("ggml_opencl: warning, the default device is not a GPU: '%s'.\n", - dev_ctx->device_name.c_str()); + dev_ctx->device_name.c_str()); } } @@ -3902,7 +3903,7 @@ static ggml_backend_opencl_context * ggml_cl_init(ggml_backend_dev_t dev) { backend_ctx->non_uniform_workgroups = false; #if CL_TARGET_OPENCL_VERSION >= 300 CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT, sizeof(cl_bool), - &backend_ctx->non_uniform_workgroups, 0)); + &backend_ctx->non_uniform_workgroups, 0)); #endif } else { GGML_ASSERT(opencl_c_version.major == 2); @@ -3921,17 +3922,17 @@ static ggml_backend_opencl_context * ggml_cl_init(ggml_backend_dev_t dev) { #ifdef GGML_OPENCL_USE_ADRENO_KERNELS backend_ctx->adreno_xmem_gemm_enabled = getenv("GGML_OPENCL_ADRENO_XMEM_GEMM") != nullptr && - backend_ctx->gpu_family == GPU_FAMILY::ADRENO; + backend_ctx->gpu_family == GPU_FAMILY::ADRENO; if (getenv("GGML_OPENCL_ADRENO_XMEM_GEMM") != nullptr) { GGML_LOG_INFO("ggml_opencl: Adreno xmem F16xF32 GEMM %s\n", - backend_ctx->adreno_xmem_gemm_enabled ? - "enabled (temporary weight prepack)" : "requested but unsupported by this driver"); + backend_ctx->adreno_xmem_gemm_enabled ? + "enabled (temporary weight prepack)" : "requested but unsupported by this driver"); } #endif // GGML_OPENCL_USE_ADRENO_KERNELS // determine whether to use large buffer for Adreno backend_ctx->adreno_use_large_buffer = getenv("GGML_OPENCL_ADRENO_USE_LARGE_BUFFER") != nullptr && - backend_ctx->gpu_family == GPU_FAMILY::ADRENO; + backend_ctx->gpu_family == GPU_FAMILY::ADRENO; if (backend_ctx->adreno_use_large_buffer) { if (!backend_ctx->adreno_has_large_buffer) { GGML_LOG_INFO("ggml_opencl: Adreno large buffer requested but not supported by driver, will use regular buffer\n"); @@ -3971,15 +3972,15 @@ static ggml_backend_opencl_context * ggml_cl_init(ggml_backend_dev_t dev) { size_t max_B_d_bytes = MIN(required_B_d_bytes, backend_ctx->max_alloc_size); if (required_A_q_d_bytes > backend_ctx->max_alloc_size) { GGML_LOG_WARN("ggml_opencl: A_q_d buffer size reduced from %zu to %zu due to device limitations.\n", - required_A_q_d_bytes, max_A_q_d_bytes); + required_A_q_d_bytes, max_A_q_d_bytes); } if (required_A_s_d_bytes > backend_ctx->max_alloc_size) { GGML_LOG_WARN("ggml_opencl: A_s_d buffer size reduced from %zu to %zu due to device limitations.\n", - required_A_s_d_bytes, max_A_s_d_bytes); + required_A_s_d_bytes, max_A_s_d_bytes); } if (required_B_d_bytes > backend_ctx->max_alloc_size) { GGML_LOG_WARN("ggml_opencl: B_d buffer size reduced from %zu to %zu due to device limitations.\n", - required_B_d_bytes, max_B_d_bytes); + required_B_d_bytes, max_B_d_bytes); } backend_ctx->prealloc_quant_trans.allocate(context, max_A_q_d_bytes); @@ -4776,7 +4777,7 @@ inline bool use_adreno_kernels(const ggml_backend_opencl_context *backend_ctx, c int64_t threshold_ne0 = 512; int64_t threshold_ne1 = 512; if (!backend_ctx->adreno_cl_compiler_version.newer_than_or_same(E031, 38, 11, 0) && - backend_ctx->adreno_cl_compiler_version.type != DX) { + backend_ctx->adreno_cl_compiler_version.type != DX) { threshold_ne0 = 128; threshold_ne1 = 128; } @@ -4891,8 +4892,8 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te case GGML_OP_DIV: case GGML_OP_SUB: return (op->src[0]->type == op->src[1]->type) && - (op->src[0]->type == op->type) && - (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16); + (op->src[0]->type == op->type) && + (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16); case GGML_OP_ADD_ID: return op->src[0]->type == GGML_TYPE_F32; case GGML_OP_SQR: @@ -4957,12 +4958,12 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te ggml_scale_mode mode = (ggml_scale_mode)(ggml_get_op_params_i32(op, 0) & 0xFF); const bool antialias = (ggml_scale_mode)(ggml_get_op_params_i32(op, 0) & GGML_SCALE_FLAG_ANTIALIAS); return op->src[0]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 && - (mode == GGML_SCALE_MODE_NEAREST || mode == GGML_SCALE_MODE_BILINEAR) && !antialias; + (mode == GGML_SCALE_MODE_NEAREST || mode == GGML_SCALE_MODE_BILINEAR) && !antialias; } case GGML_OP_CONV_2D: return (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F16 && op->type == GGML_TYPE_F16) || - (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32) || - (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32); + (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32) || + (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32); case GGML_OP_SSM_CONV: return (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32); case GGML_OP_GATED_DELTA_NET: @@ -4986,11 +4987,11 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te } else if (op->src[0]->type == GGML_TYPE_F32) { return op->src[1]->type == GGML_TYPE_F32; } else if (op->src[0]->type == GGML_TYPE_Q4_0 || op->src[0]->type == GGML_TYPE_Q4_1 || - op->src[0]->type == GGML_TYPE_MXFP4 || - op->src[0]->type == GGML_TYPE_IQ4_NL || - op->src[0]->type == GGML_TYPE_Q4_K || - op->src[0]->type == GGML_TYPE_Q5_K || - op->src[0]->type == GGML_TYPE_Q6_K) { + op->src[0]->type == GGML_TYPE_MXFP4 || + op->src[0]->type == GGML_TYPE_IQ4_NL || + op->src[0]->type == GGML_TYPE_Q4_K || + op->src[0]->type == GGML_TYPE_Q5_K || + op->src[0]->type == GGML_TYPE_Q6_K) { 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) { return op->src[1]->type == GGML_TYPE_F32; @@ -8083,14 +8084,14 @@ static bool ggml_cl_can_mul_mat(const struct ggml_tensor * src0, const struct gg // TODO: find the optimal values for these return (src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) && src1->type == GGML_TYPE_F32 && - dst->type == GGML_TYPE_F32 && + dst->type == GGML_TYPE_F32 && (ne0 >= 32 && ne1 >= 32 && ne10 >= 32); } // Copy a noncontiguous tensor to contiguous tensor. ne[] remains the same but // nb[] is recalculated such that tensor is contiguous. static void ggml_cl_copy_to_contiguous(ggml_backend_t backend, const ggml_tensor * src, cl_mem dst, - cl_ulong &nb0, cl_ulong &nb1, cl_ulong &nb2, cl_ulong &nb3) { + cl_ulong &nb0, cl_ulong &nb1, cl_ulong &nb2, cl_ulong &nb3) { ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context; const int tensor_type_size = ggml_type_size(src->type); From 2176f9575b7f6aef4d642869236b17baa87dc4f9 Mon Sep 17 00:00:00 2001 From: Yee Man Chan Date: Wed, 27 May 2026 07:40:27 +0800 Subject: [PATCH 7/7] removed return that won\'t be executed --- ggml/src/ggml-opencl/ggml-opencl.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 7a066d22966d..6d6c3e8973d1 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -17190,7 +17190,6 @@ static void ggml_cl_gated_delta_net(ggml_backend_t backend, ggml_tensor * dst) { case 128: si = 3; break; default: GGML_ASSERT(false && "ggml_cl_gated_delta_net: unsupported S_v"); - return; } const int kda = (src_g->ne[0] == (int64_t) S_v) ? 1 : 0;