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 <lih@qti.qualcomm.com>
This commit is contained in:
kurquharandLi He authored and 
GitHub committed
1 parent
0a1ef309cd
commit
9b56ccf00d
4 files changed
+320
No files matched your search
@@ -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
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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]);
|
||||
}
|
||||
}
|
||||
@@ -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];
|
||||
}
|
||||
}
|
||||
Reference in new issue
Block a user