[Cherry-pick][OpenCL][kernel] Support pad_left != pad_right or pad_up != pad_down (#5763)

Author: zhaoyang-star

lite/backends/opencl/cl_kernel/image/pool_kernel.cl

@@ -24,19 +24,18 @@ __kernel void pool_max(__read_only image2d_t input,
                        __private const int ksize_w,
                        __private const int stride_h,
                        __private const int stride_w,
-                       __private const int pad_top,
-                       __private const int pad_left) {
+                       __private const int4 pad) {
   const int out_c = get_global_id(0);
   const int out_w = get_global_id(1);
   const int out_nh = get_global_id(2);
   const int out_n = out_nh / out_height;
   const int out_h = out_nh % out_height;
 
-  int start_h = out_h * stride_h - pad_top;
+  int start_h = out_h * stride_h - (pad.x - pad.y);
   int end_h = min(start_h + ksize_h, in_height);
   start_h = max(start_h, 0);
 
-  int start_w = out_w * stride_w - pad_left;
+  int start_w = out_w * stride_w - (pad.w - pad.z);
   int end_w = min(start_w + ksize_w, in_width);
   start_w = max(start_w, 0);
 
@@ -65,19 +64,18 @@ __kernel void pool_avg(__read_only image2d_t input,
                        __private const int ksize_w,
                        __private const int stride_h,
                        __private const int stride_w,
-                       __private const int pad_top,
-                       __private const int pad_left) {
+                       __private const int4 pad) {
   const int out_c = get_global_id(0);
   const int out_w = get_global_id(1);
   const int out_nh = get_global_id(2);
   const int out_n = out_nh / out_height;
   const int out_h = out_nh % out_height;
 
-  int start_h = out_h * stride_h - pad_top;
+  int start_h = out_h * stride_h - pad.x;
   int end_h = min(start_h + ksize_h, in_height);
   start_h = max(start_h, 0);
 
-  int start_w = out_w * stride_w - pad_left;
+  int start_w = out_w * stride_w - pad.z;
   int end_w = min(start_w + ksize_w, in_width);
   start_w = max(start_w, 0);
 
@@ -96,7 +94,7 @@ __kernel void pool_avg(__read_only image2d_t input,
   div = (CL_DTYPE)((end_h - start_h)*(end_w - start_w));
 #else
   div = (CL_DTYPE)(ksize_w * ksize_h);
-#endif 
+#endif
   CL_DTYPE4 avg = sum / div;
   const int pos_out_x = mad24(out_c, out_width, out_w);
   WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(pos_out_x, out_nh), avg);
@@ -112,8 +110,7 @@ __kernel void pool_avg_global(__read_only image2d_t input,
                               __private const int ksize_w,
                               __private const int stride_h,
                               __private const int stride_w,
-                              __private const int pad_top,
-                              __private const int pad_left) {
+                              __private const int4 pad) {
   const int out_c = get_global_id(0);
   const int out_w = get_global_id(1);   // =1
   const int out_nh = get_global_id(2);  // = n*1
@@ -182,8 +179,7 @@ __kernel void pool_max_global(__read_only image2d_t input,
                               __private const int ksize_w,
                               __private const int stride_h,
                               __private const int stride_w,
-                              __private const int pad_top,
-                              __private const int pad_left) {
+                              __private const int4 pad) {
   const int out_c = get_global_id(0);
   const int out_w = get_global_id(1);   // =1
   const int out_nh = get_global_id(2);  // = n*1

lite/kernels/opencl/pool_image_compute.cc

@@ -27,8 +27,6 @@
 #endif
 #include "lite/backends/opencl/cl_utility.h"
 
-#undef LITE_WITH_LOG
-
 namespace paddle {
 namespace lite {
 namespace kernels {
@@ -44,7 +42,6 @@ class PoolComputeImage2D : public KernelLite<TARGET(kOpenCL),
 
   void PrepareForRun() override {
     const auto& param = *param_.get_mutable<param_t>();
-
     kernel_func_name_ += param.pooling_type;
     const bool global_pooling = param.global_pooling;
     const bool exclusive = param.exclusive;
@@ -93,6 +90,8 @@ class PoolComputeImage2D : public KernelLite<TARGET(kOpenCL),
       }
     }
 
+    cl_int4 pad = {paddings[0], paddings[1], paddings[2], paddings[3]};
+
 #ifdef LITE_WITH_LOG
     VLOG(4) << "in_dims : [" << in_dims.size() << "]" << in_dims[0] << "  "
             << in_dims[1] << "  " << in_dims[2] << "  " << in_dims[3];
@@ -108,12 +107,6 @@ class PoolComputeImage2D : public KernelLite<TARGET(kOpenCL),
             << paddings[1] << "  " << paddings[2] << "  " << paddings[3];
 #endif
 
-    bool pads_equal =
-        (paddings[0] == paddings[1]) && (paddings[2] == paddings[3]);
-    if (!pads_equal) {
-      LOG(FATAL)
-          << "padding requires pad_left == pad_right, pad_top == pad_bottom";
-    }
     auto& context = ctx_->As<OpenCLContext>();
     CHECK(context.cl_context() != nullptr);
 
@@ -162,9 +155,7 @@ class PoolComputeImage2D : public KernelLite<TARGET(kOpenCL),
     CL_CHECK_FATAL(status);
     status = kernel.setArg(++arg_idx, static_cast<const int>(strides[1]));
     CL_CHECK_FATAL(status);
-    status = kernel.setArg(++arg_idx, static_cast<const int>(paddings[2]));
-    CL_CHECK_FATAL(status);
-    status = kernel.setArg(++arg_idx, static_cast<const int>(paddings[0]));
+    status = kernel.setArg(++arg_idx, pad);
     CL_CHECK_FATAL(status);
 
     status = EnqueueNDRangeKernel(context,
@@ -203,4 +194,3 @@ REGISTER_LITE_KERNEL(pool2d,
                                        PRECISION(kFP16),
                                        DATALAYOUT(kImageDefault))})
     .Finalize();
-#define LITE_WITH_LOG