mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2025-06-11 11:45:30 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Merge pull request #10922 from pengli:dnn

Browse Source 
* ave pooling ocl fix

support the padded area control in ave pooling

Signed-off-by: Li Peng <peng.li@intel.com>

* warning fix: ununitialized field
This commit is contained in:
Li, Peng 2018-02-23 02:01:12 +08:00 committed by Alexander Alekhin
parent eaaba6462e
commit 5caf6244a3
4 changed files with 19 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
modules/dnn/src/layers/pooling_layer.cpp
View File

@ -145,9 +145,6 @@ public:
inps.getUMatVector(inputs);
outs.getUMatVector(outputs);
if (type == AVE && padMode == "SAME")
return false;
if (poolOp.empty())
{
OCL4DNNPoolConfig config;
@ -161,6 +158,7 @@ public:
config.pool_method = type == MAX ? LIBDNN_POOLING_METHOD_MAX :
(type == AVE ? LIBDNN_POOLING_METHOD_AVE :
LIBDNN_POOLING_METHOD_STO);
config.avePoolPaddedArea = avePoolPaddedArea;
poolOp = Ptr<OCL4DNNPool<float> >(new OCL4DNNPool<float>(config));
}

5
modules/dnn/src/ocl4dnn/include/ocl4dnn.hpp
View File

@ -344,7 +344,8 @@ struct OCL4DNNPoolConfig
dilation(1, 1),
channels(0),
pool_method(LIBDNN_POOLING_METHOD_MAX),
global_pooling(false)
global_pooling(false),
avePoolPaddedArea(false)
{}
MatShape in_shape;
MatShape out_shape;
@ -356,6 +357,7 @@ struct OCL4DNNPoolConfig
int channels;
ocl4dnnPoolingMethod_t pool_method; // = LIBDNN_POOLING_METHOD_MAX;
bool global_pooling; // = false;
bool avePoolPaddedArea;
};
template<typename Dtype>
@ -388,6 +390,7 @@ class OCL4DNNPool
int32_t width_;
int32_t pooled_height_;
int32_t pooled_width_;
bool avePoolPaddedArea;
};
struct OCL4DNNInnerProductConfig

6
modules/dnn/src/ocl4dnn/src/ocl4dnn_pool.cpp
View File

@ -56,6 +56,7 @@ OCL4DNNPool<Dtype>::OCL4DNNPool(OCL4DNNPoolConfig config)
channels_ = config.channels;
pool_method_ = config.pool_method;
avePoolPaddedArea = config.avePoolPaddedArea;
for (int i = 0; i < spatial_dims; ++i)
{
@ -143,10 +144,11 @@ bool OCL4DNNPool<Dtype>::Forward(const UMat& bottom,
ocl::dnn::ocl4dnn_pooling_oclsrc,
format("-D KERNEL_AVE_POOL=1 -D KERNEL_W=%d -D KERNEL_H=%d"
" -D STRIDE_W=%d -D STRIDE_H=%d"
" -D PAD_W=%d -D PAD_H=%d",
" -D PAD_W=%d -D PAD_H=%d%s",
kernel_w_, kernel_h_,
stride_w_, stride_h_,
pad_w_, pad_h_
pad_w_, pad_h_,
avePoolPaddedArea ? " -D AVE_POOL_PADDING_AREA" : ""
));
if (oclk_ave_pool_forward.empty())

11
modules/dnn/src/opencl/ocl4dnn_pooling.cl
View File

@ -114,11 +114,20 @@ __kernel void TEMPLATE(ave_pool_forward, Dtype)(
int wstart = pw * STRIDE_W - PAD_W;
int hend = min(hstart + KERNEL_H, height + PAD_H);
int wend = min(wstart + KERNEL_W, width + PAD_W);
const int pool_size = (hend - hstart) * (wend - wstart);
int pool_size;
#ifdef AVE_POOL_PADDING_AREA
pool_size = (hend - hstart) * (wend - wstart);
hstart = max(hstart, (int)0);
wstart = max(wstart, (int)0);
hend = min(hend, height);
wend = min(wend, width);
#else
hstart = max(hstart, (int)0);
wstart = max(wstart, (int)0);
hend = min(hend, height);
wend = min(wend, width);
pool_size = (hend - hstart) * (wend - wstart);
#endif
Dtype aveval = 0;
__global const Dtype* bottom_slice = bottom_data
+ (n * channels + c) * height * width;