[OpenCL][kernel]add depthwise transpose conv opencl op and ut (#6816)

zhenlin-work · web-flow · commit 37cb085cb076 · 2021-09-07T13:53:19.000+08:00
* add depthwise transpose conv op test=develop

* add depthwise transpose conv op test=develop

* add depthwise transpose conv op test=develop
diff --git a/lite/backends/opencl/cl_kernel/image/depthwise_conv2d_transpose_kernel.cl b/lite/backends/opencl/cl_kernel/image/depthwise_conv2d_transpose_kernel.cl
@@ -0,0 +1,110 @@
+/* Copyright (c) 2021 PaddlePaddle Authors. All Rights Reserved.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "cl_common.h"
+
+__kernel void depthwise_conv2d_transpose(
+    __private const int global_size_dim0,  // (out_c + 3) / 4
+    __private const int global_size_dim1,  // out_w
+    __private const int global_size_dim2,  // out_n * out_h
+    __read_only image2d_t input,
+    __read_only image2d_t filter,
+    __read_only image2d_t bias,
+    __write_only image2d_t output,
+    __private const int2 input_shape,
+    __private const int2 output_shape,
+    __private const int2 stride_shape,
+    __private const int2 align_shape,
+    __private const int2 padding_shape,
+    __private const int2 kernel_shape,
+    __private const int kernel_size,
+    __private const int input_c_blks) {
+  const int out_c_blk_idx = get_global_id(0);
+  const int out_w_idx = get_global_id(1);
+  const int out_nh_idx = get_global_id(2);
+
+  if (out_c_blk_idx >= global_size_dim0 || out_w_idx >= global_size_dim1 ||
+      out_nh_idx >= global_size_dim2) {
+    return;
+  }
+
+  const int out_n_idx = out_nh_idx / output_shape.y;
+  const int out_h_idx = out_nh_idx % output_shape.y;
+
+  int kernel_start_x = max(0, (out_w_idx + align_shape.x) / stride_shape.x);
+  int kernel_start_y = max(0, (out_h_idx + align_shape.y) / stride_shape.y);
+  int valid_kernel_width =
+      kernel_shape.x - mad24(kernel_start_x, stride_shape.x, padding_shape.x) +
+      out_w_idx - 1;
+  int valid_kernel_height =
+      kernel_shape.y - mad24(kernel_start_y, stride_shape.y, padding_shape.y) +
+      out_h_idx - 1;
+
+  int kernel_x_0, kernel_x_1, kernel_x_2, kernel_x_3, kernel_y;
+  CL_DTYPE4 in0;
+  CL_DTYPE4 weights0, weights1, weights2, weights3;
+  int ic = out_c_blk_idx;
+
+#ifdef BIASE_CH
+  CL_DTYPE4 out0 =
+      READ_IMG_TYPE(CL_DTYPE_CHAR, bias, SAMPLER, (int2)(out_c_blk_idx, 0));
+#else
+  CL_DTYPE4 out0 = 0.f;
+#endif
+  kernel_x_0 = ic << 2;
+  kernel_x_1 = kernel_x_0 + 1;
+  kernel_x_2 = kernel_x_0 + 2;
+  kernel_x_3 = kernel_x_0 + 3;
+  int in_idx = mul24(ic, input_shape.x);
+  for (int k_y = valid_kernel_height, idx_h = kernel_start_y; k_y >= 0;
+       k_y -= stride_shape.y, idx_h++) {
+    int in_y_idx = mad24(out_n_idx, input_shape.y, idx_h);
+    int in_nh_value = select(in_y_idx, -1, idx_h < 0 || idx_h >= input_shape.y);
+    int in_width0 = kernel_start_x;
+
+    for (int k_x = valid_kernel_width; k_x >= 0; k_x -= stride_shape.x) {
+      kernel_y = mad24(k_y, kernel_shape.x, k_x);
+
+      weights0 = READ_IMG_TYPE(
+          CL_DTYPE_CHAR, filter, SAMPLER, (int2)(kernel_x_0, kernel_y));
+      weights1 = READ_IMG_TYPE(
+          CL_DTYPE_CHAR, filter, SAMPLER, (int2)(kernel_x_1, kernel_y));
+      weights2 = READ_IMG_TYPE(
+          CL_DTYPE_CHAR, filter, SAMPLER, (int2)(kernel_x_2, kernel_y));
+      weights3 = READ_IMG_TYPE(
+          CL_DTYPE_CHAR, filter, SAMPLER, (int2)(kernel_x_3, kernel_y));
+
+      int in_width_value0 = in_width0;
+      in_width_value0 =
+          select(in_idx + in_width_value0,
+                 -1,
+                 (in_width_value0 < 0 || in_width_value0 >= input_shape.x));
+      in0 = READ_IMG_TYPE(
+          CL_DTYPE_CHAR, input, SAMPLER, (int2)(in_width_value0, in_nh_value));
+
+      out0.x += in0.x * weights0.x;
+      out0.y += in0.y * weights1.x;
+      out0.z += in0.z * weights2.x;
+      out0.w += in0.w * weights3.x;
+      in_width0++;
+    }
+  }
+  int2 out_pos0 =
+      (int2)(out_c_blk_idx * output_shape.x + out_w_idx, out_nh_idx);
+  out0 = activation_type4(out0, 0.f);
+#ifdef SCALE_ACTIVATION
+  out0 = fuse_scale(out0, 1.f, 0.f, 0.f);
+#endif
+  WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, out_pos0, out0);
+}
diff --git a/lite/kernels/opencl/CMakeLists.txt b/lite/kernels/opencl/CMakeLists.txt
@@ -65,6 +65,9 @@ lite_cc_test(test_depthwise_conv2d_image_opencl SRCS depthwise_conv2d_image_comp
 lite_cc_test(test_conv_transpose_image_opencl SRCS conv_transpose_image_compute_test.cc
              DEPS conv_transpose_opencl_image core)
 
+lite_cc_test(test_depthwise_conv2d_transpose_image_opencl SRCS depthwise_conv2d_transpose_image_compute_test.cc
+             DEPS conv_transpose_opencl_image core)
+
 lite_cc_test(test_nearest_interp_image_opencl SRCS nearest_interp_image_compute_test.cc
              DEPS nearest_interp_opencl_image layout_opencl_image core)
 
diff --git a/lite/kernels/opencl/conv_transpose_image_compute.cc b/lite/kernels/opencl/conv_transpose_image_compute.cc
@@ -27,6 +27,18 @@ void ConvTransposeImageCompute::PrepareForRun() {
   const bool is_mali = context.cl_context()->IsArmMali();
 
   conv_param_ = param_.get_mutable<param_t>();
+  auto x_dims = conv_param_->x->dims();
+  input_tensor_n_ = x_dims[0];
+  input_tensor_c_ = x_dims[1];
+  input_tensor_h_ = x_dims[2];
+  input_tensor_w_ = x_dims[3];
+
+  auto output_dims = conv_param_->output->dims();
+  output_tensor_n_ = output_dims[0];
+  output_tensor_c_ = output_dims[1];
+  output_tensor_h_ = output_dims[2];
+  output_tensor_w_ = output_dims[3];
+
   auto filter_dims = conv_param_->filter->dims();
   filter_tensor_c_ = filter_dims[1];
   filter_tensor_h_ = filter_dims[2];
@@ -56,14 +68,14 @@ void ConvTransposeImageCompute::PrepareForRun() {
   std::vector<float> filter_cpu_trans(conv_param_->filter->numel());
   DDimLite filter_trans_dims{
       {filter_dims[1], filter_dims[0], filter_dims[2], filter_dims[3]}};
+
   // Convert filter layout from IOHW to OIHW
   IOHW2OIHW<float, int64_t>(filter_cpu,
                             filter_cpu_trans.data(),
                             filter_trans_dims[0],
                             filter_trans_dims[1],
                             filter_trans_dims[2],
                             filter_trans_dims[3]);
-
   filter_gpu_image_ = std::unique_ptr<Tensor>(new Tensor);
   tensor_hold_filter_image_ = std::unique_ptr<Tensor>(new Tensor);
   tensor_hold_bias_image_ = std::unique_ptr<Tensor>(new Tensor);
@@ -72,6 +84,23 @@ void ConvTransposeImageCompute::PrepareForRun() {
     std::string kernel_name = "conv2d_transpose";
     kernel_func_names_.push_back(kernel_name);
 
+    CLImageConverterNBlock converter;
+    const DDim& filter_image_dims =
+        converter.InitImageDimInfoWith(filter_trans_dims);
+    filter_image_w_ = filter_image_dims[0];  // ((C + 3) / 4) * 4;
+    filter_image_h_ = filter_image_dims[1];  // ((N + 3) / 4) * H * W;
+    tensor_hold_filter_image_->Resize({1, filter_image_w_, filter_image_h_, 4});
+    auto* filter_image_data = MUTABLE_DATA_CPU(tensor_hold_filter_image_);
+
+    converter.NCHWToImage(
+        filter_cpu_trans.data(), filter_image_data, filter_trans_dims);
+    MUTABLE_DATA_GPU(
+        filter_gpu_image_, filter_image_w_, filter_image_h_, filter_image_data);
+  } else if ((groups_ == input_tensor_c_) && (groups_ == output_tensor_c_)) {
+    // for depthwsie conv transpose
+    std::string kernel_name = "depthwise_conv2d_transpose";
+    kernel_func_names_.push_back(kernel_name);
+
     CLImageConverterNBlock converter;
     const DDim& filter_image_dims =
         converter.InitImageDimInfoWith(filter_trans_dims);
@@ -86,7 +115,8 @@ void ConvTransposeImageCompute::PrepareForRun() {
         filter_gpu_image_, filter_image_w_, filter_image_h_, filter_image_data);
   } else {
     LOG(FATAL)
-        << "conv2d_transpose image compute not support this condition yet!";
+        << "conv2d_transpose image compute not support this condition yet! "
+        << groups_ << " " << input_tensor_c_ << " " << output_tensor_c_;
   }
 
   // build options
@@ -180,7 +210,11 @@ void ConvTransposeImageCompute::PrepareForRun() {
                << conv_param_->scale_activation_type;
   }
 
-  kernel_func_paths_.push_back("image/conv2d_transpose_kernel.cl");
+  if (groups_ == 1) {
+    kernel_func_paths_.push_back("image/conv2d_transpose_kernel.cl");
+  } else if ((groups_ == input_tensor_c_) && (groups_ == output_tensor_c_)) {
+    kernel_func_paths_.push_back("image/depthwise_conv2d_transpose_kernel.cl");
+  }
   VLOG(1) << "kernel_func_names_[0]:" << kernel_func_names_[0]
           << " kernel_func_paths_[0]:" << kernel_func_paths_[0];
   build_options_.push_back(build_options_single);
@@ -209,19 +243,11 @@ void ConvTransposeImageCompute::ReInitWhenNeeded() {
     is_first_epoch_for_run_ = false;
     last_input_dims_ = x_dims;
 
-    input_tensor_n_ = x_dims[0];
-    input_tensor_c_ = x_dims[1];
-    input_tensor_h_ = x_dims[2];
-    input_tensor_w_ = x_dims[3];
     auto x_image_shape = InitImageDimInfoWith(x_dims);
     input_image_h_ = x_image_shape["height"];
     input_image_w_ = x_image_shape["width"];
 
     auto output_dims = conv_param_->output->dims();
-    output_tensor_n_ = output_dims[0];
-    output_tensor_c_ = output_dims[1];
-    output_tensor_h_ = output_dims[2];
-    output_tensor_w_ = output_dims[3];
     auto output_image_shape = InitImageDimInfoWith(output_dims);
     output_image_h_ = output_image_shape["height"];
     output_image_w_ = output_image_shape["width"];
@@ -247,13 +273,17 @@ void ConvTransposeImageCompute::SetGlobalWorkSize() {
   global_work_size_ = cl::NDRange{static_cast<size_t>(gws[0]),
                                   static_cast<size_t>(gws[1]),
                                   static_cast<size_t>(gws[2])};
+  LOG(INFO) << "global_work_size_: " << gws[0] << " " << gws[1] << " "
+            << gws[2];
 }
 
 void ConvTransposeImageCompute::SetArgs() {
   const int pad_w = filter_tensor_w_ - 1 - pad_left_;
   const int pad_h = filter_tensor_h_ - 1 - pad_up_;
   const int align_w = stride_w_ - 1 - pad_w;
   const int align_h = stride_h_ - 1 - pad_h;
+  LOG(INFO) << "pad_w, pad_h: " << pad_w << " " << pad_h;
+  LOG(INFO) << "align_w, align_h: " << align_w << " " << align_h;
   cl_int2 pad_wh = {pad_w, pad_h};
   cl_int2 align_wh = {align_w, align_h};
 
@@ -371,3 +401,22 @@ REGISTER_LITE_KERNEL(conv2d_transpose,
                                        DATALAYOUT(kImageDefault))})
     .BindPaddleOpVersion("conv2d_transpose", 1)
     .Finalize();
+
+REGISTER_LITE_KERNEL(depthwise_conv2d_transpose,
+                     kOpenCL,
+                     kFP16,
+                     kImageDefault,
+                     paddle::lite::kernels::opencl::ConvTransposeImageCompute,
+                     image2d)
+    .BindInput("Input",
+               {LiteType::GetTensorTy(TARGET(kOpenCL),
+                                      PRECISION(kFP16),
+                                      DATALAYOUT(kImageDefault))})
+    .BindInput("Bias", {LiteType::GetTensorTy(TARGET(kARM))})
+    .BindInput("Filter", {LiteType::GetTensorTy(TARGET(kARM))})
+    .BindOutput("Output",
+                {LiteType::GetTensorTy(TARGET(kOpenCL),
+                                       PRECISION(kFP16),
+                                       DATALAYOUT(kImageDefault))})
+    .BindPaddleOpVersion("depthwise_conv2d_transpose", 1)
+    .Finalize();
diff --git a/lite/kernels/opencl/depthwise_conv2d_transpose_image_compute_test.cc b/lite/kernels/opencl/depthwise_conv2d_transpose_image_compute_test.cc
