From 9b56ccf00d755b0bf780140f6352a67db1ac0378 Mon Sep 17 00:00:00 2001 From: kurquhar Date: Wed, 21 Oct 2026 07:59:55 -0700 Subject: [PATCH] ggml-opencl : add IM2COL_3D support (#30895) * ggml-opencl : add IM2COL_3D support Assisted-by: OpenCode * opencl: use same type for kernel and host * opencl: fix for older adreno --------- Co-authored-by: Li He --- ggml/src/ggml-opencl/CMakeLists.txt | 2 + ggml/src/ggml-opencl/ggml-opencl.cpp | 138 ++++++++++++++++++ ggml/src/ggml-opencl/kernels/im2col_3d_f16.cl | 90 ++++++++++++ ggml/src/ggml-opencl/kernels/im2col_3d_f32.cl | 90 ++++++++++++ 4 files changed, 320 insertions(+) create mode 100644 ggml/src/ggml-opencl/kernels/im2col_3d_f16.cl create mode 100644 ggml/src/ggml-opencl/kernels/im2col_3d_f32.cl diff --git a/ggml/src/ggml-opencl/CMakeLists.txt b/ggml/src/ggml-opencl/CMakeLists.txt index 84087162d..917578c9f 100644 --- a/ggml/src/ggml-opencl/CMakeLists.txt +++ b/ggml/src/ggml-opencl/CMakeLists.txt @@ -80,6 +80,8 @@ set(GGML_OPENCL_KERNELS solve_tri im2col_f32 im2col_f16 + im2col_3d_f32 + im2col_3d_f16 mean mul_mv_f16_f16 mul_mv_f16_f32_1row diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 6071b6e58..2312c6c7c 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -971,6 +971,7 @@ struct ggml_backend_opencl_context { cl_kernel kernel_mul_mv_iq4_nl_f32_flat; cl_kernel kernel_solve_tri_f32; cl_kernel kernel_im2col_f32, kernel_im2col_f16; + cl_kernel kernel_im2col_3d_f32, kernel_im2col_3d_f16; cl_kernel kernel_argsort_f32_i32; cl_kernel kernel_sum_rows_f32, kernel_sum_rows_f32_4; cl_kernel kernel_cumsum_blk, kernel_cumsum_add; @@ -2101,6 +2102,40 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { GGML_LOG_CONT("."); } + // im2col_3d_f32 + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "im2col_3d_f32.cl.h" + }; +#else + const std::string kernel_src = read_file("im2col_3d_f32.cl"); +#endif + cl_program prog = + build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts); + + CL_CHECK((backend_ctx->kernel_im2col_3d_f32 = clCreateKernel(prog, "kernel_im2col_3d_f32", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + + // im2col_3d_f16 + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "im2col_3d_f16.cl.h" + }; +#else + const std::string kernel_src = read_file("im2col_3d_f16.cl"); +#endif + cl_program prog = + build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts); + + CL_CHECK((backend_ctx->kernel_im2col_3d_f16 = clCreateKernel(prog, "kernel_im2col_3d_f16", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + // mul_mv_q4_0_f32 { #ifdef GGML_OPENCL_EMBED_KERNELS @@ -9656,6 +9691,8 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te return op->src[0]->type == GGML_TYPE_F32 && ggml_is_contiguous(op->src[0]); case GGML_OP_IM2COL: return true; + case GGML_OP_IM2COL_3D: + return true; case GGML_OP_ARGSORT: { load_cl_kernels_argsort(backend_ctx); @@ -30032,6 +30069,101 @@ static void ggml_cl_im2col(ggml_backend_t backend, const ggml_tensor * src0, con backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); } +static void ggml_cl_im2col_3d(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + GGML_ASSERT(src0); + GGML_ASSERT(src1); + GGML_ASSERT(src1->extra); + GGML_ASSERT(dst); + GGML_ASSERT(dst->extra); + + GGML_ASSERT(src1->type == GGML_TYPE_F32); + GGML_ASSERT(dst->type == GGML_TYPE_F16 || dst->type == GGML_TYPE_F32); + + ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context; + + ggml_tensor_extra_cl * extra1 = (ggml_tensor_extra_cl *)src1->extra; + ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra; + + cl_ulong offset1 = extra1->offset + src1->view_offs; + cl_ulong offsetd = extrad->offset + dst->view_offs; + + const int32_t s0 = ((const int32_t*)(dst->op_params))[0]; + const int32_t s1 = ((const int32_t*)(dst->op_params))[1]; + const int32_t s2 = ((const int32_t*)(dst->op_params))[2]; + const int32_t p0 = ((const int32_t*)(dst->op_params))[3]; + const int32_t p1 = ((const int32_t*)(dst->op_params))[4]; + const int32_t p2 = ((const int32_t*)(dst->op_params))[5]; + const int32_t d0 = ((const int32_t*)(dst->op_params))[6]; + const int32_t d1 = ((const int32_t*)(dst->op_params))[7]; + const int32_t d2 = ((const int32_t*)(dst->op_params))[8]; + const cl_long IC = ((const int32_t*)(dst->op_params))[9]; + + const cl_long N = src1->ne[3] / IC; + const cl_long ID = src1->ne[2]; + const cl_long IH = src1->ne[1]; + const cl_long IW = src1->ne[0]; + + const cl_long KD = src0->ne[2]; + const cl_long KH = src0->ne[1]; + const cl_long KW = src0->ne[0]; + + const cl_long OD = dst->ne[3] / N; + const cl_long OH = dst->ne[2]; + const cl_long OW = dst->ne[1]; + + const size_t es = ggml_element_size(src1); + const cl_long stride_x = src1->nb[0] / es; + const cl_long stride_y = src1->nb[1] / es; + const cl_long stride_z = src1->nb[2] / es; + const cl_long stride_q = src1->nb[3] / es; + + const cl_long IC_KD_KH_KW = IC * KD * KH * KW; + + cl_kernel kernel; + + if (dst->type == GGML_TYPE_F16) { + kernel = backend_ctx->kernel_im2col_3d_f16; + } else { + kernel = backend_ctx->kernel_im2col_3d_f32; + } + + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_long), &IC)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_long), &ID)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_long), &IH)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_long), &IW)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_long), &KD)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_long), &KH)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_long), &KW)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_long), &OD)); + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_long), &OH)); + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_long), &OW)); + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_long), &N)); + CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_long), &stride_x)); + CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_long), &stride_y)); + CL_CHECK(clSetKernelArg(kernel, 17, sizeof(cl_long), &stride_z)); + CL_CHECK(clSetKernelArg(kernel, 18, sizeof(cl_long), &stride_q)); + CL_CHECK(clSetKernelArg(kernel, 19, sizeof(int), &s0)); + CL_CHECK(clSetKernelArg(kernel, 20, sizeof(int), &s1)); + CL_CHECK(clSetKernelArg(kernel, 21, sizeof(int), &s2)); + CL_CHECK(clSetKernelArg(kernel, 22, sizeof(int), &p0)); + CL_CHECK(clSetKernelArg(kernel, 23, sizeof(int), &p1)); + CL_CHECK(clSetKernelArg(kernel, 24, sizeof(int), &p2)); + CL_CHECK(clSetKernelArg(kernel, 25, sizeof(int), &d0)); + CL_CHECK(clSetKernelArg(kernel, 26, sizeof(int), &d1)); + CL_CHECK(clSetKernelArg(kernel, 27, sizeof(int), &d2)); + + const size_t local_col = 64; + const size_t num_blocks = (IC_KD_KH_KW + local_col - 1) / local_col; + size_t global_work_size[] = { num_blocks * local_col, (size_t)OW, (size_t)(N * OD * OH) }; + size_t local_work_size[] = { local_col, 1, 1 }; + + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); +} + static void ggml_cl_argsort(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { GGML_ASSERT(src0); GGML_ASSERT(src0->extra); @@ -30933,6 +31065,12 @@ bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor } func = ggml_cl_im2col; break; + case GGML_OP_IM2COL_3D: + if (!any_on_device) { + return false; + } + func = ggml_cl_im2col_3d; + break; case GGML_OP_ARGSORT: if (!any_on_device) { return false; diff --git a/ggml/src/ggml-opencl/kernels/im2col_3d_f16.cl b/ggml/src/ggml-opencl/kernels/im2col_3d_f16.cl new file mode 100644 index 000000000..a165e4460 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/im2col_3d_f16.cl @@ -0,0 +1,90 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +kernel void kernel_im2col_3d_f16( + global float * src1, + ulong offset1, + global half * dst, + ulong offsetd, + long IC, + long ID, + long IH, + long IW, + long KD, + long KH, + long KW, + long OD, + long OH, + long OW, + long N, + long stride_x, + long stride_y, + long stride_z, + long stride_q, + int s0, + int s1, + int s2, + int p0, + int p1, + int p2, + int d0, + int d1, + int d2 +) { + const int ic = (int)IC; + const int kd = (int)KD; + const int kh = (int)KH; + const int kw = (int)KW; + const int od = (int)OD; + const int oh = (int)OH; + const int ow = (int)OW; + const int n = (int)N; + + const int KD_KH_KW = kd * kh * kw; + const int KH_KW = kh * kw; + const int IC_KD_KH_KW = ic * KD_KH_KW; + + const int i = (int)get_global_id(0); + if (i >= IC_KD_KH_KW) { + return; + } + + const int iow = (int)get_global_id(1); + if (iow >= ow) { + return; + } + + const int OD_OH = od * oh; + const int iz = (int)get_global_id(2); + if (iz >= n * OD_OH) { + return; + } + + src1 = (global float*)((global char*)src1 + offset1); + dst = (global half*)((global char*)dst + offsetd); + + const int iic = i / KD_KH_KW; + const int ikd = (i - iic * KD_KH_KW) / KH_KW; + const int ikh = (i - iic * KD_KH_KW - ikd * KH_KW) / kw; + const int ikw = i % kw; + + const int in = iz / OD_OH; + const int iod = (iz - in * OD_OH) / oh; + const int ioh = iz % oh; + + const int iiw = iow * s0 + ikw * d0 - p0; + const int iih = ioh * s1 + ikh * d1 - p1; + const int iid = iod * s2 + ikd * d2 - p2; + + const long offset_dst = + (long)(iz * ow + iow) * (long)IC_KD_KH_KW + + (long)(iic * KD_KH_KW + ikd * KH_KW + ikh * kw + ikw); + + if (iih < 0 || iih >= IH || iiw < 0 || iiw >= IW || iid < 0 || iid >= ID) { + dst[offset_dst] = 0.0f; + } else { + const long offset_src = + (long)(in * ic + iic) * stride_q + (long)iid * stride_z + + (long)iih * stride_y + (long)iiw * stride_x; + dst[offset_dst] = convert_half_rte(src1[offset_src]); + } +} diff --git a/ggml/src/ggml-opencl/kernels/im2col_3d_f32.cl b/ggml/src/ggml-opencl/kernels/im2col_3d_f32.cl new file mode 100644 index 000000000..ea94d5a5d --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/im2col_3d_f32.cl @@ -0,0 +1,90 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +kernel void kernel_im2col_3d_f32( + global float * src1, + ulong offset1, + global float * dst, + ulong offsetd, + long IC, + long ID, + long IH, + long IW, + long KD, + long KH, + long KW, + long OD, + long OH, + long OW, + long N, + long stride_x, + long stride_y, + long stride_z, + long stride_q, + int s0, + int s1, + int s2, + int p0, + int p1, + int p2, + int d0, + int d1, + int d2 +) { + const int ic = (int)IC; + const int kd = (int)KD; + const int kh = (int)KH; + const int kw = (int)KW; + const int od = (int)OD; + const int oh = (int)OH; + const int ow = (int)OW; + const int n = (int)N; + + const int KD_KH_KW = kd * kh * kw; + const int KH_KW = kh * kw; + const int IC_KD_KH_KW = ic * KD_KH_KW; + + const int i = (int)get_global_id(0); + if (i >= IC_KD_KH_KW) { + return; + } + + const int iow = (int)get_global_id(1); + if (iow >= ow) { + return; + } + + const int OD_OH = od * oh; + const int iz = (int)get_global_id(2); + if (iz >= n * OD_OH) { + return; + } + + src1 = (global float*)((global char*)src1 + offset1); + dst = (global float*)((global char*)dst + offsetd); + + const int iic = i / KD_KH_KW; + const int ikd = (i - iic * KD_KH_KW) / KH_KW; + const int ikh = (i - iic * KD_KH_KW - ikd * KH_KW) / kw; + const int ikw = i % kw; + + const int in = iz / OD_OH; + const int iod = (iz - in * OD_OH) / oh; + const int ioh = iz % oh; + + const int iiw = iow * s0 + ikw * d0 - p0; + const int iih = ioh * s1 + ikh * d1 - p1; + const int iid = iod * s2 + ikd * d2 - p2; + + const long offset_dst = + (long)(iz * ow + iow) * (long)IC_KD_KH_KW + + (long)(iic * KD_KH_KW + ikd * KH_KW + ikh * kw + ikw); + + if (iih < 0 || iih >= IH || iiw < 0 || iiw >= IW || iid < 0 || iid >= ID) { + dst[offset_dst] = 0.0f; + } else { + const long offset_src = + (long)(in * ic + iic) * stride_q + (long)iid * stride_z + + (long)iih * stride_y + (long)iiw * stride_x; + dst[offset_dst] = src1[offset_src]; + } +}