Skip to content

Commit e34b64c

Browse files
alalekalalek
committed
dnn(ocl4dnn): refactor pooling OpenCL calls
1 parent f071a48 commit e34b64c

File tree

3 files changed

+107
-111
lines changed

3 files changed

+107
-111
lines changed

‎modules/dnn/src/ocl4dnn/include/ocl4dnn.hpp

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -351,8 +351,6 @@ class OCL4DNNPool
351351
UMat& top_data,
352352
UMat& top_mask);
353353
private:
354-
UMat mask_idx_;
355-
356354
// Pooling parameters
357355
std::vector<int32_t> pad_;
358356
std::vector<int32_t> stride_;

‎modules/dnn/src/ocl4dnn/src/ocl4dnn_pool.cpp

Lines changed: 66 additions & 62 deletions
Original file line numberDiff line numberDiff line change
@@ -88,7 +88,7 @@ OCL4DNNPool<Dtype>::OCL4DNNPool(OCL4DNNPoolConfig config)
8888
template<typename Dtype>
8989
OCL4DNNPool<Dtype>::~OCL4DNNPool()
9090
{
91-
mask_idx_.release();
91+
// nothing
9292
}
9393

9494
template<typename Dtype>
@@ -99,99 +99,103 @@ bool OCL4DNNPool<Dtype>::Forward(const UMat& bottom,
9999
bool ret = true;
100100
size_t global[] = { 128 * 128 };
101101
size_t local[] = { 128 };
102-
cl_uint argIdx = 0;
103102

104103
// support 2D case
105104
switch (pool_method_)
106105
{
107106
case LIBDNN_POOLING_METHOD_MAX:
108107
{
109-
if (top_mask.empty() && mask_idx_.empty())
110-
{
111-
mask_idx_.create(1, count_, CV_32FC1);
112-
}
113-
ocl::Kernel oclk_max_pool_forward(CL_KERNEL_SELECT("max_pool_forward"),
114-
cv::ocl::dnn::ocl4dnn_pooling_oclsrc);
108+
bool haveMask = !top_mask.empty();
109+
ocl::Kernel oclk_max_pool_forward(
110+
haveMask ? CL_KERNEL_SELECT("max_pool_forward_mask") : CL_KERNEL_SELECT("max_pool_forward"),
111+
ocl::dnn::ocl4dnn_pooling_oclsrc,
112+
format("-D KERNEL_MAX_POOL=1 -D KERNEL_W=%d -D KERNEL_H=%d"
113+
" -D STRIDE_W=%d -D STRIDE_H=%d"
114+
" -D PAD_W=%d -D PAD_H=%d%s",
115+
kernel_w_, kernel_h_,
116+
stride_w_, stride_h_,
117+
pad_w_, pad_h_,
118+
haveMask ? " -D HAVE_MASK=1" : ""
119+
));
115120

116121
if (oclk_max_pool_forward.empty())
117122
return false;
118123

119-
argIdx = 0;
120-
oclk_max_pool_forward.set(argIdx++, count_);
121-
oclk_max_pool_forward.set(argIdx++, ocl::KernelArg::PtrReadOnly(bottom));
122-
oclk_max_pool_forward.set(argIdx++, batch_size_);
123-
oclk_max_pool_forward.set(argIdx++, channels_);
124-
oclk_max_pool_forward.set(argIdx++, height_);
125-
oclk_max_pool_forward.set(argIdx++, width_);
126-
oclk_max_pool_forward.set(argIdx++, pooled_height_);
127-
oclk_max_pool_forward.set(argIdx++, pooled_width_);
128-
oclk_max_pool_forward.set(argIdx++, kernel_h_);
129-
oclk_max_pool_forward.set(argIdx++, kernel_w_);
130-
oclk_max_pool_forward.set(argIdx++, stride_h_);
131-
oclk_max_pool_forward.set(argIdx++, stride_w_);
132-
oclk_max_pool_forward.set(argIdx++, pad_h_);
133-
oclk_max_pool_forward.set(argIdx++, pad_w_);
134-
oclk_max_pool_forward.set(argIdx++, ocl::KernelArg::PtrWriteOnly(top));
135-
oclk_max_pool_forward.set(argIdx++, mask_idx_.empty() ? 0 : 1);
136-
if (mask_idx_.empty())
137-
oclk_max_pool_forward.set(argIdx++, (void *)NULL);
138-
else
139-
oclk_max_pool_forward.set(argIdx++, ocl::KernelArg::PtrWriteOnly(mask_idx_));
140-
oclk_max_pool_forward.set(argIdx++, ocl::KernelArg::PtrWriteOnly(top_mask));
124+
oclk_max_pool_forward.args(
125+
count_,
126+
ocl::KernelArg::PtrReadOnly(bottom),
127+
batch_size_,
128+
channels_,
129+
height_,
130+
width_,
131+
pooled_height_,
132+
pooled_width_,
133+
ocl::KernelArg::PtrWriteOnly(top),
134+
ocl::KernelArg::PtrWriteOnly(top_mask)
135+
);
141136

142137
ret = oclk_max_pool_forward.run(1, global, local, false);
143138
}
144139
break;
145140
case LIBDNN_POOLING_METHOD_AVE:
146141
{
142+
CV_Assert(top_mask.empty());
143+
147144
ocl::Kernel oclk_ave_pool_forward(CL_KERNEL_SELECT("ave_pool_forward"),
148-
cv::ocl::dnn::ocl4dnn_pooling_oclsrc);
145+
ocl::dnn::ocl4dnn_pooling_oclsrc,
146+
format("-D KERNEL_AVE_POOL=1 -D KERNEL_W=%d -D KERNEL_H=%d"
147+
" -D STRIDE_W=%d -D STRIDE_H=%d"
148+
" -D PAD_W=%d -D PAD_H=%d",
149+
kernel_w_, kernel_h_,
150+
stride_w_, stride_h_,
151+
pad_w_, pad_h_
152+
));
149153

150154
if (oclk_ave_pool_forward.empty())
151155
return false;
152156

153-
argIdx = 0;
154-
oclk_ave_pool_forward.set(argIdx++, count_);
155-
oclk_ave_pool_forward.set(argIdx++, ocl::KernelArg::PtrReadOnly(bottom));
156-
oclk_ave_pool_forward.set(argIdx++, batch_size_);
157-
oclk_ave_pool_forward.set(argIdx++, channels_);
158-
oclk_ave_pool_forward.set(argIdx++, height_);
159-
oclk_ave_pool_forward.set(argIdx++, width_);
160-
oclk_ave_pool_forward.set(argIdx++, pooled_height_);
161-
oclk_ave_pool_forward.set(argIdx++, pooled_width_);
162-
oclk_ave_pool_forward.set(argIdx++, kernel_h_);
163-
oclk_ave_pool_forward.set(argIdx++, kernel_w_);
164-
oclk_ave_pool_forward.set(argIdx++, stride_h_);
165-
oclk_ave_pool_forward.set(argIdx++, stride_w_);
166-
oclk_ave_pool_forward.set(argIdx++, pad_h_);
167-
oclk_ave_pool_forward.set(argIdx++, pad_w_);
168-
oclk_ave_pool_forward.set(argIdx++, ocl::KernelArg::PtrWriteOnly(top));
157+
oclk_ave_pool_forward.args(
158+
count_,
159+
ocl::KernelArg::PtrReadOnly(bottom),
160+
batch_size_,
161+
channels_,
162+
height_,
163+
width_,
164+
pooled_height_,
165+
pooled_width_,
166+
ocl::KernelArg::PtrWriteOnly(top)
167+
);
169168

170169
ret = oclk_ave_pool_forward.run(1, global, local, false);
171170
}
172171
break;
173172
case LIBDNN_POOLING_METHOD_STO:
174173
{
174+
CV_Assert(top_mask.empty());
175+
175176
ocl::Kernel oclk_sto_pool_forward(CL_KERNEL_SELECT("sto_pool_forward_test"),
176-
cv::ocl::dnn::ocl4dnn_pooling_oclsrc);
177+
ocl::dnn::ocl4dnn_pooling_oclsrc,
178+
format("-D KERNEL_STO_POOL=1 -D KERNEL_W=%d -D KERNEL_H=%d"
179+
" -D STRIDE_W=%d -D STRIDE_H=%d",
180+
kernel_w_, kernel_h_,
181+
stride_w_, stride_h_
182+
));
183+
177184

178185
if (oclk_sto_pool_forward.empty())
179186
return false;
180187

181-
argIdx = 0;
182-
oclk_sto_pool_forward.set(argIdx++, count_);
183-
oclk_sto_pool_forward.set(argIdx++, ocl::KernelArg::PtrReadOnly(bottom));
184-
oclk_sto_pool_forward.set(argIdx++, batch_size_);
185-
oclk_sto_pool_forward.set(argIdx++, channels_);
186-
oclk_sto_pool_forward.set(argIdx++, height_);
187-
oclk_sto_pool_forward.set(argIdx++, width_);
188-
oclk_sto_pool_forward.set(argIdx++, pooled_height_);
189-
oclk_sto_pool_forward.set(argIdx++, pooled_width_);
190-
oclk_sto_pool_forward.set(argIdx++, kernel_h_);
191-
oclk_sto_pool_forward.set(argIdx++, kernel_w_);
192-
oclk_sto_pool_forward.set(argIdx++, stride_h_);
193-
oclk_sto_pool_forward.set(argIdx++, stride_w_);
194-
oclk_sto_pool_forward.set(argIdx++, ocl::KernelArg::PtrWriteOnly(top));
188+
oclk_sto_pool_forward.args(
189+
count_,
190+
ocl::KernelArg::PtrReadOnly(bottom),
191+
batch_size_,
192+
channels_,
193+
height_,
194+
width_,
195+
pooled_height_,
196+
pooled_width_,
197+
ocl::KernelArg::PtrWriteOnly(top)
198+
);
195199

196200
ret = oclk_sto_pool_forward.run(1, global, local, false);
197201
}

‎modules/dnn/src/opencl/ocl4dnn_pooling.cl

Lines changed: 41 additions & 47 deletions
Original file line numberDiff line numberDiff line change
@@ -44,14 +44,23 @@
4444
#define TEMPLATE(name,type) CONCAT(name,type)
4545
#define Dtype float
4646

47-
void TEMPLATE(max_pool_forward_impl, Dtype)(
47+
#if defined KERNEL_MAX_POOL
48+
49+
__kernel void
50+
#ifdef HAVE_MASK
51+
TEMPLATE(max_pool_forward_mask, Dtype)
52+
#else
53+
TEMPLATE(max_pool_forward, Dtype)
54+
#endif
55+
(
4856
const int nthreads, __global const Dtype* bottom_data, const int num,
4957
const int channels, const int height, const int width,
50-
const int pooled_height, const int pooled_width, const int kernel_h,
51-
const int kernel_w, const int stride_h, const int stride_w, const int pad_h,
52-
const int pad_w,
53-
__global Dtype* top_data,
54-
const int use_mask, __global int* mask, __global Dtype* top_mask, bool no_mask)
58+
const int pooled_height, const int pooled_width,
59+
__global Dtype* top_data
60+
#ifdef HAVE_MASK
61+
, __global Dtype* mask
62+
#endif
63+
)
5564
{
5665
for (int index = get_global_id(0); index < nthreads;
5766
index += get_global_size(0))
@@ -60,10 +69,10 @@ void TEMPLATE(max_pool_forward_impl, Dtype)(
6069
const int ph = (index / pooled_width) % pooled_height;
6170
const int c = (index / pooled_width / pooled_height) % channels;
6271
const int n = index / pooled_width / pooled_height / channels;
63-
int hstart = ph * stride_h - pad_h;
64-
int wstart = pw * stride_w - pad_w;
65-
const int hend = min(hstart + kernel_h, height);
66-
const int wend = min(wstart + kernel_w, width);
72+
int hstart = ph * STRIDE_H - PAD_H;
73+
int wstart = pw * STRIDE_W - PAD_W;
74+
const int hend = min(hstart + KERNEL_H, height);
75+
const int wend = min(wstart + KERNEL_W, width);
6776
hstart = max(hstart, (int)0);
6877
wstart = max(wstart, (int)0);
6978
Dtype maxval = -FLT_MAX;
@@ -79,38 +88,19 @@ void TEMPLATE(max_pool_forward_impl, Dtype)(
7988
}
8089
}
8190
top_data[index] = maxval;
82-
if (!no_mask) {
83-
if (use_mask == 1) {
84-
mask[index] = maxidx;
85-
} else {
86-
top_mask[index] = maxidx;
87-
}
88-
}
91+
#ifdef HAVE_MASK
92+
mask[index] = maxidx;
93+
#endif
8994
}
9095
}
9196

92-
__kernel void TEMPLATE(max_pool_forward, Dtype)(
93-
const int nthreads, __global const Dtype* bottom_data, const int num,
94-
const int channels, const int height, const int width,
95-
const int pooled_height, const int pooled_width, const int kernel_h,
96-
const int kernel_w, const int stride_h, const int stride_w, const int pad_h,
97-
const int pad_w,
98-
__global Dtype* top_data,
99-
const int use_mask, __global int* mask, __global Dtype* top_mask)
100-
{
101-
TEMPLATE(max_pool_forward_impl, Dtype)(
102-
nthreads, bottom_data, num, channels, height, width,
103-
pooled_height, pooled_width, kernel_h,
104-
kernel_w, stride_h, stride_w, pad_h, pad_w, top_data, use_mask, mask, top_mask, false
105-
);
106-
}
97+
#elif defined KERNEL_AVE_POOL
10798

10899
__kernel void TEMPLATE(ave_pool_forward, Dtype)(
109100
const int nthreads, __global const Dtype* const bottom_data, const int num,
110101
const int channels, const int height, const int width,
111-
const int pooled_height, const int pooled_width, const int kernel_h,
112-
const int kernel_w, const int stride_h, const int stride_w, const int pad_h,
113-
const int pad_w, __global Dtype* top_data)
102+
const int pooled_height, const int pooled_width,
103+
__global Dtype* top_data)
114104
{
115105
for (int index = get_global_id(0); index < nthreads;
116106
index += get_global_size(0))
@@ -120,10 +110,10 @@ __kernel void TEMPLATE(ave_pool_forward, Dtype)(
120110
const int ph = (index / pooled_width) % pooled_height;
121111
const int c = (index / pooled_width / pooled_height) % channels;
122112
const int n = index / pooled_width / pooled_height / channels;
123-
int hstart = ph * stride_h - pad_h;
124-
int wstart = pw * stride_w - pad_w;
125-
int hend = min(hstart + kernel_h, height + pad_h);
126-
int wend = min(wstart + kernel_w, width + pad_w);
113+
int hstart = ph * STRIDE_H - PAD_H;
114+
int wstart = pw * STRIDE_W - PAD_W;
115+
int hend = min(hstart + KERNEL_H, height + PAD_H);
116+
int wend = min(wstart + KERNEL_W, width + PAD_W);
127117
const int pool_size = (hend - hstart) * (wend - wstart);
128118
hstart = max(hstart, (int)0);
129119
wstart = max(wstart, (int)0);
@@ -142,11 +132,12 @@ __kernel void TEMPLATE(ave_pool_forward, Dtype)(
142132
}
143133
}
144134

135+
#elif defined KERNEL_STO_POOL
136+
145137
__kernel void TEMPLATE(sto_pool_forward_test,Dtype)(
146138
const int nthreads, __global const Dtype* const bottom_data, const int num,
147139
const int channels, const int height, const int width,
148-
const int pooled_height, const int pooled_width, const int kernel_h,
149-
const int kernel_w, const int stride_h, const int stride_w,
140+
const int pooled_height, const int pooled_width,
150141
__global Dtype* top_data)
151142
{
152143
for (int index = get_global_id(0); index < nthreads;
@@ -156,10 +147,10 @@ __kernel void TEMPLATE(sto_pool_forward_test,Dtype)(
156147
const int ph = (index / pooled_width) % pooled_height;
157148
const int c = (index / pooled_width / pooled_height) % channels;
158149
const int n = index / pooled_width / pooled_height / channels;
159-
const int hstart = ph * stride_h;
160-
const int hend = min(hstart + kernel_h, height);
161-
const int wstart = pw * stride_w;
162-
const int wend = min(wstart + kernel_w, width);
150+
const int hstart = ph * STRIDE_H;
151+
const int hend = min(hstart + KERNEL_H, height);
152+
const int wstart = pw * STRIDE_W;
153+
const int wend = min(wstart + KERNEL_W, width);
163154
// We set cumsum to be 0 to avoid divide-by-zero problems
164155
Dtype cumsum = FLT_MIN;
165156
Dtype cumvalues = 0.;
@@ -168,10 +159,13 @@ __kernel void TEMPLATE(sto_pool_forward_test,Dtype)(
168159
// First pass: get sum
169160
for (int h = hstart; h < hend; ++h) {
170161
for (int w = wstart; w < wend; ++w) {
171-
cumsum += bottom_slice[h * width + w];
172-
cumvalues += bottom_slice[h * width + w] * bottom_slice[h * width + w];
162+
Dtype v = bottom_slice[h * width + w];
163+
cumsum += v;
164+
cumvalues += v * v;
173165
}
174166
}
175167
top_data[index] = cumvalues / cumsum;
176168
}
177169
}
170+
171+
#endif // KERNEL_*

0 commit comments

Comments
 (0)