[OPENCL] develop pixel_shuffle opencl kernel & unit test ,test=develop (#3733)

* [OPENCL] develop pixel_shuffle opencl kernel & unit test ,test=develop

* [OPENCL] develop pixel_shuffle opencl kernel & unit test ,test=develop

* [OPENCL] develop pixel_shuffle opencl kernel & unit test ,test=develop

* [OPENCL] develop pixel_shuffle opencl kernel & unit test ,test=develop

* [OPENCL] develop pixel_shuffle opencl kernel & unit test ,test=develop

* [OPENCL] develop pixel_shuffle opencl kernel & unit test ,test=develop

* [OPENCL] develop pixel_shuffle opencl kernel & unit test ,test=develop

* [OPENCL] develop pixel_shuffle opencl kernel & unit test ,test=develop

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

+/* Copyright (c) 2018 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 pixel_shuffle(__read_only image2d_t input_image,
+                            __write_only image2d_t output_image,
+                            __private const int in_N,
+                            __private const int in_C,
+                            __private const int in_H,
+                            __private const int in_W,
+                            __private const int out_N,
+                            __private const int out_C,
+                            __private const int out_H,
+                            __private const int out_W,
+                            __private const int upscale_factor) {
+  const int out_c4 = get_global_id(0);
+  const int out_w = get_global_id(1);
+  const int out_nh = get_global_id(2);
+
+  int out_h = out_nh % out_H;
+  int out_n = out_nh / out_H;
+
+  const sampler_t sampler =
+      CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
+
+  int in_h = out_h / upscale_factor;
+  int in_w = out_w / upscale_factor;
+  int in_nh = out_n * in_H + in_h;
+
+  CL_DTYPE4 res;
+  int out_c;
+  int in_c;
+  CL_DTYPE4 in;
+  int2 in_pos;
+
+  out_c = out_c4 * 4 + 0;
+  in_c = out_c * upscale_factor * upscale_factor +
+         (out_h % upscale_factor) * upscale_factor + (out_w % upscale_factor);
+  in_pos.x = (in_c / 4) * in_W + in_w;
+  in_pos.y = in_nh;
+  in = READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, sampler, in_pos);
+  if (in_c % 4 == 0) {
+    res.x = in.x;
+  } else if (in_c % 4 == 1) {
+    res.x = in.y;
+  } else if (in_c % 4 == 2) {
+    res.x = in.z;
+  } else if (in_c % 4 == 3) {
+    res.x = in.w;
+  }
+
+  out_c = out_c4 * 4 + 1;
+  in_c = out_c * upscale_factor * upscale_factor +
+         (out_h % upscale_factor) * upscale_factor + (out_w % upscale_factor);
+  in_pos.x = (in_c / 4) * in_W + in_w;
+  in_pos.y = in_nh;
+  in = READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, sampler, in_pos);
+  if (in_c % 4 == 0) {
+    res.y = in.x;
+  } else if (in_c % 4 == 1) {
+    res.y = in.y;
+  } else if (in_c % 4 == 2) {
+    res.y = in.z;
+  } else if (in_c % 4 == 3) {
+    res.y = in.w;
+  }
+
+  out_c = out_c4 * 4 + 2;
+  in_c = out_c * upscale_factor * upscale_factor +
+         (out_h % upscale_factor) * upscale_factor + (out_w % upscale_factor);
+  in_pos.x = (in_c / 4) * in_W + in_w;
+  in_pos.y = in_nh;
+  in = READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, sampler, in_pos);
+  if (in_c % 4 == 0) {
+    res.z = in.x;
+  } else if (in_c % 4 == 1) {
+    res.z = in.y;
+  } else if (in_c % 4 == 2) {
+    res.z = in.z;
+  } else if (in_c % 4 == 3) {
+    res.z = in.w;
+  }
+
+  out_c = out_c4 * 4 + 3;
+  in_c = out_c * upscale_factor * upscale_factor +
+         (out_h % upscale_factor) * upscale_factor + (out_w % upscale_factor);
+  in_pos.x = (in_c / 4) * in_W + in_w;
+  in_pos.y = in_nh;
+  in = READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, sampler, in_pos);
+  if (in_c % 4 == 0) {
+    res.w = in.x;
+  } else if (in_c % 4 == 1) {
+    res.w = in.y;
+  } else if (in_c % 4 == 2) {
+    res.w = in.z;
+  } else if (in_c % 4 == 3) {
+    res.w = in.w;
+  }
+
+  int2 out_pos;
+  out_pos.x = out_c4 * out_W + out_w;
+  out_pos.y = out_nh;
+  WRITE_IMG_TYPE(CL_DTYPE_CHAR, output_image, out_pos, res);
+}

lite/kernels/opencl/CMakeLists.txt

@@ -34,6 +34,7 @@ add_kernel(instance_norm_opencl OPENCL basic SRCS instance_norm_image_compute.cc
 add_kernel(dropout_opencl OPENCL basic SRCS dropout_image_compute.cc DEPS ${cl_kernel_deps})
 add_kernel(pad2d_opencl OPENCL basic SRCS pad2d_image_compute.cc DEPS ${cl_kernel_deps})
 add_kernel(box_coder_opencl OPENCL basic SRCS box_coder_image_compute.cc DEPS ${cl_kernel_deps})
+add_kernel(pixel_shuffle_opencl OPENCL basic SRCS pixel_shuffle_image_compute.cc DEPS ${cl_kernel_deps})
 # extra
 # wait to add ...
 
@@ -73,6 +74,9 @@ lite_cc_test(test_concat_image_opencl SRCS concat_image_compute_test.cc
 lite_cc_test(test_layout_image_opencl SRCS layout_image_compute_test.cc
              DEPS layout_opencl op_registry program context)
 
+lite_cc_test(test_pixel_shuffle_image_opencl SRCS pixel_shuffle_image_compute_test.cc
+             DEPS pixel_shuffle_opencl op_registry program context)
+
 lite_cc_test(test_elementwise_add_image_opencl SRCS elementwise_add_image_compute_test.cc
              DEPS elementwise_add_opencl fusion_elementwise_add_activation_opencl op_registry program context)
 lite_cc_test(test_elementwise_sub_image_opencl SRCS elementwise_sub_image_compute_test.cc

lite/kernels/opencl/pixel_shuffle_image_compute.cc

+// Copyright (c) 2019 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 <vector>
+#include "lite/backends/opencl/cl_half.h"
+#include "lite/backends/opencl/cl_include.h"
+#include "lite/core/kernel.h"
+#include "lite/core/op_registry.h"
+#include "lite/kernels/opencl/image_helper.h"
+#include "lite/operators/op_params.h"
+#include "lite/utils/replace_stl/stream.h"
+#include "lite/utils/string.h"
+#ifdef LITE_WITH_PROFILE
+#include "lite/core/profile/profiler.h"
+#endif
+#include "lite/backends/opencl/cl_utility.h"
+namespace paddle {
+namespace lite {
+namespace kernels {
+namespace opencl {
+
+class PixelShuffleComputeImage2D
+    : public KernelLite<TARGET(kOpenCL),
+                        PRECISION(kFP16),
+                        DATALAYOUT(kImageDefault)> {
+ public:
+  using param_t = operators::PixelShuffleParam;
+
+  std::string doc() const override {
+    return "PixelShuffle using cl::Image2D, kFP16";
+  }
+
+  void PrepareForRun() override {
+    VLOG(1) << "kernel_func_name_:" << kernel_func_name_;
+
+    auto& context = ctx_->As<OpenCLContext>();
+    context.cl_context()->AddKernel(kernel_func_name_,
+                                    "image/pixel_shuffle_kernel.cl",
+                                    build_options_,
+                                    time_stamp_);
+
+    STL::stringstream kernel_key;
+    kernel_key << kernel_func_name_ << build_options_ << time_stamp_;
+    kernel_ = context.cl_context()->GetKernel(kernel_key.str());
+  }
+
+  void ReInitWhenNeeded() override {
+    VLOG(1) << "ReInitWhenNeeded:  " << kernel_func_name_;
+    pixel_shuffle_param_ = param_.get_mutable<param_t>();
+    auto x_dims = pixel_shuffle_param_->x->dims();
+    auto out_dims = pixel_shuffle_param_->output->dims();
+    VLOG(1) << "x_dims:  " << x_dims;
+    VLOG(1) << "out_dims:  " << out_dims;
+    VLOG(1) << "upscale_factor:  " << pixel_shuffle_param_->upscale_factor;
+
+    if ((!first_epoch_for_reinit_ && x_dims != last_x_dims_) ||
+        first_epoch_for_reinit_) {
+      last_x_dims_ = x_dims;
+      first_epoch_for_reinit_ = false;
+      // compute image shape
+      paddle::lite::CLImageConverterDefault default_convertor;
+      out_img_shape_ = default_convertor.InitImageDimInfoWith(
+          pixel_shuffle_param_->output->dims());
+      VLOG(1) << "out_img_shape_:  " << out_img_shape_[0] << "  "
+              << out_img_shape_[1];
+
+      // compute global work size
+      auto image_width = out_dims[3] * ((out_dims[1] + 3) / 4);
+      size_t work_size_0 = image_width / out_dims[3];
+      size_t work_size_1 = out_dims[3];
+      size_t work_size_2 = out_dims[0] * out_dims[2];
+      global_work_size_ = cl::NDRange{work_size_0, work_size_1, work_size_2};
+      VLOG(1) << "global_work_size_:  " << global_work_size_[0] << " "
+              << global_work_size_[1] << " " << global_work_size_[2];
+    }
+  }
+
+  void Run() override {
+    auto* x_img = pixel_shuffle_param_->x->data<half_t, cl::Image2D>();
+    auto* out_img =
+        pixel_shuffle_param_->output->mutable_data<half_t, cl::Image2D>(
+            out_img_shape_[0], out_img_shape_[1]);
+
+    auto x_dims = pixel_shuffle_param_->x->dims();
+
+    int in_n = x_dims[0];
+    int in_c = x_dims[1];
+    int in_h = x_dims[2];
+    int in_w = x_dims[3];
+
+    auto out_dims = pixel_shuffle_param_->output->dims();
+
+    int out_n = out_dims[0];
+    int out_c = out_dims[1];
+    int out_h = out_dims[2];
+    int out_w = out_dims[3];
+
+    const int upscale_factor = pixel_shuffle_param_->upscale_factor;
+
+    auto& context = ctx_->As<OpenCLContext>();
+    CHECK(context.cl_context() != nullptr);
+
+    auto kernel = kernel_;
+    cl_int status;
+    status = kernel.setArg(0, *x_img);
+    CL_CHECK_FATAL(status);
+    status = kernel.setArg(1, *out_img);
+    CL_CHECK_FATAL(status);
+    status = kernel.setArg(2, in_n);
+    CL_CHECK_FATAL(status);
+    status = kernel.setArg(3, in_c);
+    CL_CHECK_FATAL(status);
+    status = kernel.setArg(4, in_h);
+    CL_CHECK_FATAL(status);
+    status = kernel.setArg(5, in_w);
+    CL_CHECK_FATAL(status);
+    status = kernel.setArg(6, out_n);
+    CL_CHECK_FATAL(status);
+    status = kernel.setArg(7, out_c);
+    CL_CHECK_FATAL(status);
+    status = kernel.setArg(8, out_h);
+    CL_CHECK_FATAL(status);
+    status = kernel.setArg(9, out_w);
+    CL_CHECK_FATAL(status);
+    status = kernel.setArg(10, upscale_factor);
+    CL_CHECK_FATAL(status);
+
+    status = EnqueueNDRangeKernel(context,
+                                  kernel,
+                                  cl::NullRange,
+                                  global_work_size_,
+                                  cl::NullRange,
+                                  nullptr,
+                                  event_);
+    CL_CHECK_FATAL(status);
+  }
+
+#ifdef LITE_WITH_PROFILE
+  void SetProfileRuntimeKernelInfo(paddle::lite::profile::OpCharacter* ch) {
+    ch->kernel_func_name = kernel_func_name_;
+    ch->cl_event =
+        event_;  // `event_` defined in `kernel.h`, valid after kernel::Run
+  }
+#endif
+ private:
+  std::string kernel_func_name_{"pixel_shuffle"};
+  std::string build_options_{"-DCL_DTYPE_half"};
+  std::string time_stamp_{GetTimeStamp()};
+
+  param_t* pixel_shuffle_param_{nullptr};
+  cl::Kernel kernel_;
+  bool first_epoch_for_reinit_{true};
+  DDim last_x_dims_;
+  DDim out_img_shape_ = DDim(std::vector<DDim::value_type>(
+      {static_cast<DDim::value_type>(1), static_cast<DDim::value_type>(1)}));
+  cl::NDRange global_work_size_ = cl::NDRange{
+      static_cast<size_t>(1), static_cast<size_t>(1), static_cast<size_t>(1)};
+};
+
+}  // namespace opencl
+}  // namespace kernels
+}  // namespace lite
+}  // namespace paddle
+
+REGISTER_LITE_KERNEL(pixel_shuffle,
+                     kOpenCL,
+                     kFP16,
+                     kImageDefault,
+                     paddle::lite::kernels::opencl::PixelShuffleComputeImage2D,
+                     image2d)
+    .BindInput("X",
+               {LiteType::GetTensorTy(TARGET(kOpenCL),
+                                      PRECISION(kFP16),
+                                      DATALAYOUT(kImageDefault))})
+    .BindOutput("Out",
+                {LiteType::GetTensorTy(TARGET(kOpenCL),
+                                       PRECISION(kFP16),
+                                       DATALAYOUT(kImageDefault))})
+    .Finalize();

lite/kernels/opencl/pixel_shuffle_image_compute_test.cc

+// Copyright (c) 2019 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 <random>
+
+#include <gtest/gtest.h>
+#include "lite/backends/opencl/target_wrapper.h"
+#include "lite/core/op_registry.h"
+#include "lite/core/tensor.h"
+#include "lite/kernels/opencl/test_helper.h"
+
+#define FP16_MAX_DIFF (5e-1)
+
+namespace paddle {
+namespace lite {
+
+TEST(pixel_shuffle_image2d, compute) {
+  LOG(INFO) << "create kernel ...";
+  auto kernels = KernelRegistry::Global().Create("pixel_shuffle",
+                                                 TARGET(kOpenCL),
+                                                 PRECISION(kFP16),
+                                                 DATALAYOUT(kImageDefault));
+  ASSERT_FALSE(kernels.empty());
+
+  const int INPUT_N = 1;
+  const int INPUT_C = 4;
+  const int INPUT_H = 2;
+  const int INPUT_W = 2;
+  const int UPSCALE_FACTOR = 2;
+
+  auto kernel = std::move(kernels.front());
+
+  LOG(INFO) << "prepare to test kernel ====> " << kernel->doc();
+
+  lite::Tensor x, out;
+  operators::PixelShuffleParam param;
+  param.x = &x;
+  param.output = &out;
+  param.upscale_factor = UPSCALE_FACTOR;
+
+  std::unique_ptr<KernelContext> context(new KernelContext);
+  context->As<OpenCLContext>().InitOnce();
+
+  kernel->SetParam(param);
+  std::unique_ptr<KernelContext> pixel_shuffle_context(new KernelContext);
+  context->As<OpenCLContext>().CopySharedTo(
+      &(pixel_shuffle_context->As<OpenCLContext>()));
+
+  kernel->SetContext(std::move(pixel_shuffle_context));
+
+  const DDim in_dim =
+      DDim(std::vector<DDim::value_type>{INPUT_N, INPUT_C, INPUT_H, INPUT_W});
+  const DDim out_dim = DDim(
+      std::vector<DDim::value_type>{INPUT_N,
+                                    INPUT_C / UPSCALE_FACTOR / UPSCALE_FACTOR,
+                                    INPUT_H * UPSCALE_FACTOR,
+                                    INPUT_W * UPSCALE_FACTOR});
+  LOG(INFO) << "in_dim: " << in_dim;
+  LOG(INFO) << "UPSCALE_FACTOR: " << UPSCALE_FACTOR;
+  LOG(INFO) << "out_dim: " << out_dim;
+
+  x.Resize(in_dim);
+  out.Resize(out_dim);
+
+  std::default_random_engine engine;
+  std::uniform_real_distribution<float> dist(-2, 2);
+  std::vector<float> input_v(INPUT_N * INPUT_C * INPUT_H * INPUT_W);
+
+  int index = 0;
+  for (auto& i : input_v) {
+    i = index++;
+  }
+  VLOG(1) << "input_v ..... ";
+  for (size_t i = 0; i < input_v.size(); i++) {
+    VLOG(10) << input_v[i];
+  }
+
+  LOG(INFO) << "prepare input";
+  CLImageConverterDefault* default_converter = new CLImageConverterDefault();
+  DDim x_image_shape = default_converter->InitImageDimInfoWith(in_dim);
+  LOG(INFO) << "x_image_shape = " << x_image_shape[0] << " "
+            << x_image_shape[1];
+  std::vector<half_t> x_image_data(x_image_shape.production() * 4);  // 4 : RGBA
+  default_converter->NCHWToImage(input_v.data(), x_image_data.data(), in_dim);
+  auto* x_image = x.mutable_data<half_t, cl::Image2D>(
+      x_image_shape[0], x_image_shape[1], x_image_data.data());
+  VLOG(1) << "x_image_data ..... ";
+  for (size_t i = 0; i < x_image_data.size(); i++) {
+    VLOG(10) << Half2Float(x_image_data[i]);
+  }
+  DDim out_image_shape = default_converter->InitImageDimInfoWith(out_dim);
+  LOG(INFO) << "out_image_shape = " << out_image_shape[0] << " "
+            << out_image_shape[1];
+  auto* out_image = out.mutable_data<half_t, cl::Image2D>(out_image_shape[0],
+                                                          out_image_shape[1]);
+  kernel->Launch();
+  CLRuntime::Global()->command_queue().finish();
+  std::vector<float> out_data_v{
+      0, 4, 1, 5, 8, 12, 9, 13, 2, 6, 3, 7, 10, 14, 11, 15};
+
+  const size_t cl_image2d_row_pitch{0};
+  const size_t cl_image2d_slice_pitch{0};
+  half_t* out_image_data = new half_t[out_image_shape.production() * 4];
+  TargetWrapperCL::ImgcpySync(out_image_data,
+                              out_image,
+                              out_image_shape[0],
+                              out_image_shape[1],
+                              cl_image2d_row_pitch,
+                              cl_image2d_slice_pitch,
+                              IoDirection::DtoH);
+  VLOG(1) << "out_image_data ..... ";
+  for (size_t i = 0; i < out_image_shape.production() * 4; i++) {
+    VLOG(10) << Half2Float(out_image_data[i]);
+  }
+  float* out_data = new float[out_image_shape.production() * 4];
+  default_converter->ImageToNCHW(
+      out_image_data, out_data, out_image_shape, out_dim);
+
+  VLOG(1) << "out_data ..... ";
+  for (int i = 0; i < out_dim.production(); i++) {
+    VLOG(10) << out_data[i];
+  }
+
+  for (int i = 0; i < out_dim.production(); i++) {
+    auto abs_diff = abs(out_data[i] - out_data_v[i]);
+    auto relative_diff = COMPUTE_RELATIVE_DIFF(out_data[i], out_data_v[i]);
+    EXPECT_EQ((relative_diff <= FP16_MAX_DIFF) || (abs_diff <= FP16_MAX_DIFF),
+              true);
+    if ((relative_diff > FP16_MAX_DIFF) && (abs_diff > FP16_MAX_DIFF)) {
+      LOG(ERROR) << "error idx:" << i << " out_data[" << i
+                 << "]:" << out_data[i] << " "
+                                           "out_ref["
+                 << i << "]:" << out_data_v[i] << " abs_diff:" << abs_diff
+                 << " relative_diff:" << relative_diff
+                 << " FP16_MAX_DIFF:" << FP16_MAX_DIFF;
+    }
+  }
+}
+
+}  // namespace lite
+}  // namespace paddle
+
+USE_LITE_KERNEL(pixel_shuffle, kOpenCL, kFP16, kImageDefault, image2d);

lite/operators/op_params.h

@@ -377,17 +377,17 @@ struct ConvParam : ParamBase {
   lite::Tensor* output{};
   std::vector<int> strides{1, 1};
   /* paddings type change
-  * from std::vector<int> to std::shared_ptr<std::vector<int>>
-  * to support dynamically modify padding
-  * let kernel param and operator param Synchronous update
-  */
+   * from std::vector<int> to std::shared_ptr<std::vector<int>>
+   * to support dynamically modify padding
+   * let kernel param and operator param Synchronous update
+   */
   std::shared_ptr<std::vector<int>> paddings;
   int groups{1};
   /* dilations type change
-  * from std::vector<int> to std::shared_ptr<std::vector<int>>
-  * to support dynamically modify padding
-  * let kernel param and operator param Synchronous update
-  */
+   * from std::vector<int> to std::shared_ptr<std::vector<int>>
+   * to support dynamically modify padding
+   * let kernel param and operator param Synchronous update
+   */
   std::shared_ptr<std::vector<int>> dilations;
   bool fuse_relu_before_depthwise_conv{false};
   bool use_mkldnn{false};
@@ -471,10 +471,10 @@ struct PoolParam : ParamBase {
       false};  // if true, knernel size and paddings will be ignored
   std::vector<int> strides{1, 1};
   /* paddings type change
-  * from std::vector<int> to std::shared_ptr<std::vector<int>>
-  * to support dynamically modify padding
-  * let kernel param and operator param Synchronous update
-  */
+   * from std::vector<int> to std::shared_ptr<std::vector<int>>
+   * to support dynamically modify padding
+   * let kernel param and operator param Synchronous update
+   */
   std::shared_ptr<std::vector<int>> paddings;
   bool exclusive{true};
   bool adaptive{false};
@@ -1515,6 +1515,11 @@ struct XPUFcParam : ParamBase {
   std::string activation_type{""};
 };
 
+struct PixelShuffleParam : ParamBase {
+  lite::Tensor* x{nullptr};
+  lite::Tensor* output{nullptr};
+  int upscale_factor{1};
+};
 }  // namespace operators
 }  // namespace lite
 }  // namespace paddle