From 5caf6244a38bd3ea771f6ce181ff28dcc8bd28d8 Mon Sep 17 00:00:00 2001
From: "Li, Peng" <peng.li@intel.com>
Date: Fri, 23 Feb 2018 02:01:12 +0800
Subject: [PATCH] Merge pull request #10922 from pengli:dnn

* 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
---
 modules/dnn/src/layers/pooling_layer.cpp     |  4 +---
 modules/dnn/src/ocl4dnn/include/ocl4dnn.hpp  |  5 ++++-
 modules/dnn/src/ocl4dnn/src/ocl4dnn_pool.cpp |  6 ++++--
 modules/dnn/src/opencl/ocl4dnn_pooling.cl    | 11 ++++++++++-
 4 files changed, 19 insertions(+), 7 deletions(-)

diff --git a/modules/dnn/src/layers/pooling_layer.cpp b/modules/dnn/src/layers/pooling_layer.cpp
index ee552b5e70..3cdc786517 100644
--- a/modules/dnn/src/layers/pooling_layer.cpp
+++ b/modules/dnn/src/layers/pooling_layer.cpp
@@ -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));
         }
 
diff --git a/modules/dnn/src/ocl4dnn/include/ocl4dnn.hpp b/modules/dnn/src/ocl4dnn/include/ocl4dnn.hpp
index 93ac1a2bc4..70ced11276 100644
--- a/modules/dnn/src/ocl4dnn/include/ocl4dnn.hpp
+++ b/modules/dnn/src/ocl4dnn/include/ocl4dnn.hpp
@@ -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
diff --git a/modules/dnn/src/ocl4dnn/src/ocl4dnn_pool.cpp b/modules/dnn/src/ocl4dnn/src/ocl4dnn_pool.cpp
index 759408c755..4d4ea9e2a9 100644
--- a/modules/dnn/src/ocl4dnn/src/ocl4dnn_pool.cpp
+++ b/modules/dnn/src/ocl4dnn/src/ocl4dnn_pool.cpp
@@ -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())
diff --git a/modules/dnn/src/opencl/ocl4dnn_pooling.cl b/modules/dnn/src/opencl/ocl4dnn_pooling.cl
index bb3ba8a82d..13e4319172 100644
--- a/modules/dnn/src/opencl/ocl4dnn_pooling.cl
+++ b/modules/dnn/src/opencl/ocl4dnn_pooling.cl
@@ -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;