Skip to content

Commit b898521

Browse files
rmatifrmatif
committed
add conv2d kernel
1 parent 086cf81 commit b898521

File tree

3 files changed

+286
-0
lines changed

3 files changed

+286
-0
lines changed

‎ggml/src/ggml-opencl/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -105,6 +105,7 @@ set(GGML_OPENCL_KERNELS
105105
pad
106106
repeat
107107
mul_mat_f16_f32
108+
conv2d
108109
)
109110

110111
foreach (K ${GGML_OPENCL_KERNELS})

‎ggml/src/ggml-opencl/ggml-opencl.cpp‎

Lines changed: 92 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -390,6 +390,7 @@ struct ggml_backend_opencl_context {
390390
cl_program program_tanh;
391391
cl_program program_upscale;
392392
cl_program program_concat;
393+
cl_program program_conv_2d;
393394
cl_program program_tsembd;
394395
cl_program program_mul_mv_id_q4_0_f32_8x_flat;
395396

@@ -441,6 +442,7 @@ struct ggml_backend_opencl_context {
441442
cl_kernel kernel_upscale_bilinear;
442443
cl_kernel kernel_concat_f32_contiguous;
443444
cl_kernel kernel_concat_f32_non_contiguous;
445+
cl_kernel kernel_conv_2d;
444446
cl_kernel kernel_timestep_embedding;
445447
cl_kernel kernel_mul_mv_id_q4_0_f32_8x_flat;
446448

@@ -1478,6 +1480,27 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve
14781480
GGML_LOG_CONT(".");
14791481
}
14801482

1483+
// conv2d
1484+
{
1485+
#ifdef GGML_OPENCL_EMBED_KERNELS
1486+
const std::string kernel_src {
1487+
#include "conv2d.cl.h"
1488+
};
1489+
#else
1490+
const std::string kernel_src = read_file("conv2d.cl");
1491+
#endif
1492+
if (!kernel_src.empty()) {
1493+
backend_ctx->program_conv_2d =
1494+
build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
1495+
CL_CHECK((backend_ctx->kernel_conv_2d = clCreateKernel(backend_ctx->program_conv_2d, "kernel_conv_2d", &err), err));
1496+
GGML_LOG_CONT(".");
1497+
} else {
1498+
GGML_LOG_WARN("ggml_opencl: conv2d kernel source not found or empty. This op will not be available.\n");
1499+
backend_ctx->program_conv_2d = nullptr;
1500+
backend_ctx->kernel_conv_2d = nullptr;
1501+
}
1502+
}
1503+
14811504
// mul_mv_id_q4_0_f32_8x_flat
14821505
{
14831506
#ifdef GGML_OPENCL_EMBED_KERNELS
@@ -2361,6 +2384,8 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
23612384
op->src[0]->ne[3] == 1 && op->ne[3] == 1;
23622385
case GGML_OP_UPSCALE:
23632386
return op->src[0]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32;
2387+
case GGML_OP_CONV_2D:
2388+
return op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32;
23642389
case GGML_OP_CONCAT:
23652390
return op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32;
23662391
case GGML_OP_TIMESTEP_EMBEDDING:
@@ -4946,7 +4971,12 @@ static void ggml_cl_timestep_embedding(ggml_backend_t backend, const ggml_tensor
49464971
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, NULL, dst);
49474972
}
49484973

4974+
<<<<<<< HEAD
49494975
static void ggml_cl_mul_mat_f16_f32_tiled(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
4976+
=======
4977+
static void ggml_cl_conv_2d(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
4978+
GGML_TENSOR_BINARY_OP_LOCALS;
4979+
>>>>>>> 4d5d5a83 (add conv2d kernel)
49504980
ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context;
49514981

49524982
ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *)src0->extra;
@@ -4957,6 +4987,7 @@ static void ggml_cl_mul_mat_f16_f32_tiled(ggml_backend_t backend, const ggml_ten
49574987
cl_ulong offset1 = extra1->offset + src1->view_offs;
49584988
cl_ulong offsetd = extrad->offset + dst->view_offs;
49594989

4990+
<<<<<<< HEAD
49604991
const int M = src0->ne[1];
49614992
const int N = src1->ne[1];
49624993
const int K = src0->ne[0];
@@ -4996,6 +5027,61 @@ static void ggml_cl_mul_mat_f16_f32_tiled(ggml_backend_t backend, const ggml_ten
49965027
};
49975028

49985029
backend_ctx->enqueue_ndrange_kernel(kernel, 2, global_work_size, local_work_size, dst);
5030+
=======
5031+
const cl_uint Cout = ne03; const cl_uint Cin = ne02; const cl_uint N = ne13;
5032+
const cl_uint KW = ne00; const cl_uint KH = ne01; const cl_uint W = ne10; const cl_uint H = ne11; const cl_uint OW = ne0; const cl_uint OH = ne1;
5033+
5034+
const cl_uint s0 = dst->op_params[0]; const cl_uint s1 = dst->op_params[1];
5035+
const cl_uint p0 = dst->op_params[2]; const cl_uint p1 = dst->op_params[3];
5036+
const cl_uint d0 = dst->op_params[4]; const cl_uint d1 = dst->op_params[5];
5037+
5038+
const cl_uint cl_nb01 = nb01/nb00; const cl_uint cl_nb02 = nb02/nb00; const cl_uint cl_nb03 = nb03/nb00;
5039+
const cl_uint cl_nb11 = nb11/nb10; const cl_uint cl_nb12 = nb12/nb10; const cl_uint cl_nb13 = nb13/nb10;
5040+
const cl_uint cl_nb1 = nb1/nb0; const cl_uint cl_nb2 = nb2/nb0; const cl_uint cl_nb3 = nb3/nb0;
5041+
5042+
const int64_t NPQ = (int64_t)N * OW * OH;
5043+
5044+
const uint32_t WG_SIZE = 128;
5045+
const uint32_t BS_K = 128;
5046+
const uint32_t BS_CRS = 16;
5047+
const uint32_t BS_NPQ = 64;
5048+
const uint32_t VEC_SIZE = 4;
5049+
5050+
auto splitWork = [](uint32_t work_size, uint32_t block_size) { return (block_size + work_size - 1) / block_size; };
5051+
const uint32_t NB_K = splitWork(Cout, BS_K);
5052+
const uint32_t NB_NPQ = splitWork(NPQ, BS_NPQ);
5053+
5054+
const size_t shmem_size = (size_t)(BS_K * (BS_CRS + 1) * sizeof(cl_half) + BS_CRS * (BS_NPQ / VEC_SIZE + 1) * sizeof(cl_half4));
5055+
5056+
cl_kernel kernel = backend_ctx->kernel_conv_2d;
5057+
cl_uint idx = 0;
5058+
CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem), &extra0->data_device)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_ulong), &offset0));
5059+
CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem), &extra1->data_device)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_ulong), &offset1));
5060+
CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem), &extrad->data_device)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_ulong), &offsetd));
5061+
CL_CHECK(clSetKernelArg(kernel, idx++, shmem_size, NULL));
5062+
CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &Cout)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &Cin)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &N));
5063+
CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &KW)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &KH)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &W)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &H));
5064+
CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &OW)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &OH));
5065+
CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &s0)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &s1)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &p0)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &p1));
5066+
CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &d0)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &d1));
5067+
CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &cl_nb01)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &cl_nb02)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &cl_nb03));
5068+
CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &cl_nb11)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &cl_nb12)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &cl_nb13));
5069+
CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &cl_nb1)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &cl_nb2)); CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint), &cl_nb3));
5070+
5071+
size_t global_work_size[] = { (size_t)NB_K * WG_SIZE, (size_t)NB_NPQ, 1 };
5072+
size_t local_work_size[] = { (size_t)WG_SIZE, 1, 1 };
5073+
5074+
#ifdef GGML_OPENCL_PROFILING
5075+
cl_event evt;
5076+
CL_CHECK(clEnqueueNDRangeKernel(backend_ctx->queue, kernel, 3, NULL, global_work_size, local_work_size, 0, NULL, &evt));
5077+
5078+
backend_ctx->profiling_info.emplace_back();
5079+
populateProfilingInfo(backend_ctx->profiling_info.back(), evt, kernel, 3, global_work_size, local_work_size, dst);
5080+
#else
5081+
GGML_UNUSED(dst);
5082+
CL_CHECK(clEnqueueNDRangeKernel(backend_ctx->queue, kernel, 3, NULL, global_work_size, local_work_size, 0, NULL, NULL));
5083+
#endif
5084+
>>>>>>> 4d5d5a83 (add conv2d kernel)
49995085
}
50005086

50015087
static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
@@ -6752,6 +6838,12 @@ bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor
67526838
}
67536839
ggml_cl_upscale(backend, tensor->src[0], tensor);
67546840
return true;
6841+
case GGML_OP_CONV_2D:
6842+
if (!any_on_device) {
6843+
return false;
6844+
}
6845+
func = ggml_cl_conv_2d;
6846+
break;
67556847
case GGML_OP_CONCAT:
67566848
if (!any_on_device) {
67576849
return false;

‎ggml/src/ggml-opencl/kernels/conv2d.cl‎

Lines changed: 193 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,193 @@
1+
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
2+
3+
#ifdef cl_intel_required_subgroup_size
4+
#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable
5+
#define INTEL_GPU 1
6+
#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16)))
7+
#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32)))
8+
#elif defined(cl_qcom_reqd_sub_group_size)
9+
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
10+
#define ADRENO_GPU 1
11+
#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half")))
12+
#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
13+
#else
14+
#define REQD_SUBGROUP_SIZE_64
15+
#endif
16+
17+
#define T_FLOAT half
18+
#define T_FLOAT4 half4
19+
#define T_ACCUM float4
20+
#define VEC_SIZE 4
21+
22+
#define BS_K 128
23+
#define BS_CRS 16
24+
#define BS_NPQ 64
25+
#define TS_K 8
26+
#define TS_NPQ 8
27+
#define WG_SIZE 128
28+
29+
#define BS_NPQ_VEC (BS_NPQ / VEC_SIZE)
30+
#define TS_NPQ_VEC (TS_NPQ / VEC_SIZE)
31+
32+
#define NT_K (BS_K / TS_K)
33+
#define NT_NPQ (BS_NPQ / TS_NPQ)
34+
35+
static inline uint splitWork(uint work_size, uint block_size){
36+
return (work_size + block_size - 1) / block_size;
37+
}
38+
39+
REQD_SUBGROUP_SIZE_64
40+
kernel void kernel_conv_2d(
41+
global void* p_knl,
42+
ulong off_knl,
43+
global void* p_src,
44+
ulong off_src,
45+
global void* p_dst,
46+
ulong off_dst,
47+
local void* shared,
48+
uint Cout, uint Cin, uint N,
49+
uint KW, uint KH, uint W, uint H, uint OW, uint OH,
50+
uint s0, uint s1, uint p0, uint p1, uint d0, uint d1,
51+
uint nb01, uint nb02, uint nb03,
52+
uint nb11, uint nb12, uint nb13,
53+
uint nb1, uint nb2, uint nb3
54+
) {
55+
global float* knl_data = (global float*) ((global char*)p_knl + off_knl);
56+
global float* src_data = (global float*) ((global char*)p_src + off_src);
57+
global float* dst_data = (global float*) ((global char*)p_dst + off_dst);
58+
59+
const uint tid = get_local_id(0);
60+
61+
const uint K = Cout;
62+
const uint CRS = Cin*KH*KW;
63+
const uint NPQ = N*OH*OW;
64+
65+
const uint NB_CRS = splitWork(CRS, BS_CRS);
66+
67+
const uint Ash_stride = BS_CRS + 1;
68+
const uint Bsh_stride_vec = BS_NPQ_VEC + 1;
69+
70+
local T_FLOAT* Ash = (local T_FLOAT*)shared;
71+
local T_FLOAT4* Bsh = (local T_FLOAT4*) &Ash[BS_K * Ash_stride];
72+
73+
T_ACCUM regC[TS_K][TS_NPQ_VEC];
74+
for (int i = 0; i < TS_K; ++i) {
75+
for (int j = 0; j < TS_NPQ_VEC; ++j) {
76+
regC[i][j] = (T_ACCUM)(0.0f);
77+
}
78+
}
79+
80+
const uint B_idx_K = get_group_id(0);
81+
const uint B_idx_NPQ = get_group_id(1);
82+
83+
const uint T_y = tid / NT_NPQ;
84+
const uint T_x = tid % NT_NPQ;
85+
86+
for (uint B_idx_CRS = 0; B_idx_CRS < NB_CRS; ++B_idx_CRS) {
87+
for(uint i = tid; i < BS_K * BS_CRS; i += WG_SIZE){
88+
uint k_l = i / BS_CRS;
89+
uint crs_l = i % BS_CRS;
90+
uint k_g = B_idx_K*BS_K + k_l;
91+
uint crs_g = B_idx_CRS*BS_CRS + crs_l;
92+
if(k_g < K && crs_g < CRS){
93+
uint Cin_idx = crs_g / (KW*KH);
94+
uint KH_idx = (crs_g - Cin_idx*KW*KH) / KW;
95+
uint KW_idx = crs_g - Cin_idx*KW*KH - KH_idx*KW;
96+
uint knl_idx = KW_idx + KH_idx*nb01 + Cin_idx*nb02 + k_g*nb03;
97+
Ash[k_l * Ash_stride + crs_l] = (T_FLOAT)knl_data[knl_idx];
98+
} else {
99+
Ash[k_l * Ash_stride + crs_l] = (T_FLOAT)0.0h;
100+
}
101+
}
102+
103+
for (uint i = tid; i < BS_CRS * BS_NPQ_VEC; i += WG_SIZE) {
104+
uint crs_l = i / BS_NPQ_VEC;
105+
uint npq_l_vec = i % BS_NPQ_VEC;
106+
107+
float4 val_f = (float4)(0.0f);
108+
uint crs_g = B_idx_CRS * BS_CRS + crs_l;
109+
110+
if (crs_g < CRS) {
111+
uint Cin_idx = crs_g / (KW * KH);
112+
uint KH_idx = (crs_g - Cin_idx * KW * KH) / KW;
113+
uint KW_idx = crs_g - Cin_idx * KW * KH - KH_idx * KW;
114+
115+
for (int v = 0; v < VEC_SIZE; ++v) {
116+
uint npq_g = B_idx_NPQ * BS_NPQ + npq_l_vec * VEC_SIZE + v;
117+
if (npq_g < NPQ) {
118+
uint N_idx = npq_g / (OH * OW);
119+
uint pq_idx = npq_g % (OH * OW);
120+
uint OH_idx = pq_idx / OW;
121+
uint OW_idx = pq_idx % OW;
122+
int H_idx = (int)(OH_idx * s1 + KH_idx * d1 - p1);
123+
int W_idx = (int)(OW_idx * s0 + KW_idx * d0 - p0);
124+
125+
if (H_idx >= 0 && H_idx < H && W_idx >= 0 && W_idx < W) {
126+
uint src_idx = W_idx + H_idx * nb11 + Cin_idx * nb12 + N_idx * nb13;
127+
switch (v) {
128+
case 0: val_f.s0 = src_data[src_idx]; break;
129+
case 1: val_f.s1 = src_data[src_idx]; break;
130+
case 2: val_f.s2 = src_data[src_idx]; break;
131+
case 3: val_f.s3 = src_data[src_idx]; break;
132+
}
133+
}
134+
}
135+
}
136+
}
137+
Bsh[crs_l * Bsh_stride_vec + npq_l_vec] = convert_half4(val_f);
138+
}
139+
barrier(CLK_LOCAL_MEM_FENCE);
140+
141+
for(uint crs_l = 0; crs_l < BS_CRS; ++crs_l){
142+
T_FLOAT regA[TS_K];
143+
for(uint k_l_reg = 0; k_l_reg < TS_K; ++k_l_reg){
144+
regA[k_l_reg] = Ash[(T_y*TS_K + k_l_reg)*Ash_stride + crs_l];
145+
}
146+
for(uint npq_l_vec_reg = 0; npq_l_vec_reg < TS_NPQ_VEC; ++npq_l_vec_reg){
147+
T_FLOAT4 regB = Bsh[crs_l*Bsh_stride_vec + T_x*TS_NPQ_VEC + npq_l_vec_reg];
148+
for(uint k_l_reg = 0; k_l_reg < TS_K; ++k_l_reg){
149+
regC[k_l_reg][npq_l_vec_reg] = mad((float)regA[k_l_reg], convert_float4(regB), regC[k_l_reg][npq_l_vec_reg]);
150+
}
151+
}
152+
}
153+
barrier(CLK_LOCAL_MEM_FENCE);
154+
}
155+
156+
for(uint k_l_reg = 0; k_l_reg < TS_K; ++k_l_reg){
157+
uint k_g = B_idx_K * BS_K + T_y * TS_K + k_l_reg;
158+
if(k_g >= K) continue;
159+
160+
for(uint npq_l_vec_reg = 0; npq_l_vec_reg < TS_NPQ_VEC; ++npq_l_vec_reg){
161+
uint npq_g_base = B_idx_NPQ * BS_NPQ + (T_x * TS_NPQ_VEC + npq_l_vec_reg) * VEC_SIZE;
162+
uint N_idx = npq_g_base / (OH*OW);
163+
uint pq_idx = npq_g_base % (OH*OW);
164+
uint OH_idx = pq_idx / OW;
165+
uint OW_idx = pq_idx % OW;
166+
167+
if (nb1 == OW && OW_idx + VEC_SIZE <= OW && npq_g_base + VEC_SIZE <= NPQ) {
168+
uint dst_idx = OW_idx + OH_idx*nb1 + k_g*nb2 + N_idx*nb3;
169+
vstore4(regC[k_l_reg][npq_l_vec_reg], 0, &dst_data[dst_idx]);
170+
} else {
171+
T_ACCUM res = regC[k_l_reg][npq_l_vec_reg];
172+
for (int v = 0; v < VEC_SIZE; ++v) {
173+
uint npq_g = npq_g_base + v;
174+
if (npq_g < NPQ) {
175+
uint N_idx_s = npq_g / (OH*OW);
176+
uint pq_idx_s = npq_g % (OH*OW);
177+
uint OH_idx_s = pq_idx_s / OW;
178+
uint OW_idx_s = pq_idx_s % OW;
179+
uint dst_idx_s = OW_idx_s + OH_idx_s*nb1 + k_g*nb2 + N_idx_s*nb3;
180+
float val_f;
181+
switch(v) {
182+
case 0: val_f = res.s0; break;
183+
case 1: val_f = res.s1; break;
184+
case 2: val_f = res.s2; break;
185+
default:val_f = res.s3; break;
186+
}
187+
dst_data[dst_idx_s] = val_f;
188+
}
189+
}
190+
}
191+
}
192+
}
193+
}

0 commit comments

Comments
 (0)