Chengxiangming/add pixel unshuffle opencl

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

@@ -0,0 +1,108 @@
+/* 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_unshuffle(__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 downscale_factor) {
+ const int in_c4 = get_global_id(0);
+ const int in_w = get_global_id(1);
+ const int in_nh = get_global_id(2);
+
+ int in_h = in_nh % in_H;
+ int in_n = in_nh / in_H;
+
+ int out_h = in_h * downscale_factor;
+ int out_w = in_w * downscale_factor;
+ int out_nh = in_n * out_H + out_h;
+
+ CL_DTYPE4 res;
+ int in_c;
+ int out_c;
+ CL_DTYPE4 in;
+ int2 out_pos;
+
+ in_c = in_c4 * 4 + 0;
+ out_c = in_c / (downscale_factor * downscale_factor);
+ out_pos.x = (out_c / 4) * out_W + out_w + (in_c % downscale_factor);
+ out_pos.y = out_nh + (in_c / (out_C * 4)) * out_H;
+ in = READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, SAMPLER, out_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;
+ }
+
+ in_c = in_c4 * 4 + 1;
+ out_c = in_c / (downscale_factor * downscale_factor);
+ out_pos.x = (out_c / 4) * out_W + out_w + (in_c % downscale_factor);
+ out_pos.y = out_nh + (in_c / (out_C * 4)) * out_H;
+ in = READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, SAMPLER, out_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;
+ }
+
+ in_c = in_c4 * 4 + 2;
+ out_c = in_c / (downscale_factor * downscale_factor);
+ out_pos.x = (out_c / 4) * out_W + out_w + (in_c % downscale_factor);
+ out_pos.y = out_nh + (in_c / (out_C * 4)) * out_H;
+ in = READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, SAMPLER, out_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;
+ }
+
+ in_c = in_c4 * 4 + 3;
+ out_c = in_c / (downscale_factor * downscale_factor);
+ out_pos.x = (out_c / 4) * out_W + out_w + (in_c % downscale_factor);
+ out_pos.y = out_nh + (in_c / (out_C * 4)) * out_H;
+ in = READ_IMG_TYPE(CL_DTYPE_CHAR, input_image, SAMPLER, out_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 in_pos;
+ in_pos.x = in_c4 * in_W + in_w;
+ in_pos.y = in_nh;
+ WRITE_IMG_TYPE(CL_DTYPE_CHAR, output_image, in_pos, res);
+}

lite/kernels/opencl/CMakeLists.txt

@@ -42,6 +42,7 @@ add_kernel(dropout_opencl_image OPENCL basic SRCS dropout_image_compute.cc)
 add_kernel(pad2d_opencl_image OPENCL basic SRCS pad2d_image_compute.cc)
 add_kernel(box_coder_opencl_image OPENCL basic SRCS box_coder_image_compute.cc)
 add_kernel(pixel_shuffle_opencl_image OPENCL basic SRCS pixel_shuffle_image_compute.cc)
+add_kernel(pixel_unshuffle_opencl_image OPENCL basic SRCS pixel_unshuffle_image_compute.cc)
 add_kernel(expand_opencl_image OPENCL basic SRCS expand_image_compute.cc)
 add_kernel(shuffle_channel_opencl_image OPENCL basic SRCS shuffle_channel_image_compute.cc)
 add_kernel(trigonometric_opencl_image OPENCL basic SRCS trigonometric_image_compute.cc)

lite/kernels/opencl/pixel_unshuffle_image_compute.cc

@@ -0,0 +1,192 @@
+// 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 PixelUnShuffleComputeImage2D
+ : public KernelLite<TARGET(kOpenCL),
+ PRECISION(kFP16),
+ DATALAYOUT(kImageDefault)> {
+ public:
+ using param_t = operators::PixelUnShuffleParam;
+
+ std::string doc() const override {
+ return "PixelUnShuffle 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_unshuffle_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_unshuffle_param_ = param_.get_mutable<param_t>();
+ auto x_dims = pixel_unshuffle_param_->x->dims();
+ auto out_dims = pixel_unshuffle_param_->output->dims();
+ VLOG(1) << "x_dims: " << x_dims;
+ VLOG(1) << "out_dims: " << out_dims;
+ VLOG(1) << "downscale_factor: " << pixel_unshuffle_param_->downscale_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_unshuffle_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 = GET_DATA_GPU(pixel_unshuffle_param_->x);
+ auto* out_img = MUTABLE_DATA_GPU(pixel_unshuffle_param_->output,
+ out_img_shape_[0],
+ out_img_shape_[1],
+ nullptr);
+
+ auto x_dims = pixel_unshuffle_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_unshuffle_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 downscale_factor = pixel_unshuffle_param_->downscale_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, downscale_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_unshuffle"};
+ std::string build_options_{""};
+ std::string time_stamp_{GetTimeStamp()};
+
+ param_t* pixel_unshuffle_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_unshuffle,
+ kOpenCL,
+ kFP16,
+ kImageDefault,
+ paddle::lite::kernels::opencl::PixelUnShuffleComputeImage2D,
+ image2d)
+ .BindInput("X",
+ {LiteType::GetTensorTy(TARGET(kOpenCL),
+ PRECISION(kFP16),
+ DATALAYOUT(kImageDefault))})
+ .BindOutput("Out",
+ {LiteType::GetTensorTy(TARGET(kOpenCL),
+ PRECISION(kFP16),
+ DATALAYOUT(kImageDefault))})
+ .Finalize();