add [ImageDefault] -> [ImageFolder] test=develop

Author: zhaoyang-star

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

@@ -312,6 +312,42 @@ __kernel void image2d_to_buffer_with_post255(__read_only image2d_t input,
   }
 }
 
+////////////////////////////////////////////////////////
+// image2d_default -> image2d_folder
+////////////////////////////////////////////////////////
+__kernel void image2d_default_to_image2d_folder(__read_only image2d_t input,
+                                                __write_only image2d_t output,
+                                                __private const int in_img_w,
+                                                __private const int in_img_h) {
+  const int pos_x = get_global_id(0);
+  const int pos_y = get_global_id(1);
+
+  CL_DTYPE4 in =
+      READ_IMG_TYPE(CL_DTYPE_CHAR, input, SAMPLER, (int2)(pos_x, pos_y));
+
+  CL_DTYPE4 in0 = 0.f;
+  CL_DTYPE4 in1 = 0.f;
+  CL_DTYPE4 in2 = 0.f;
+  CL_DTYPE4 in3 = 0.f;
+
+  in0 = READ_IMG_TYPE(CL_DTYPE_CHAR, input, SAMPLER, (int2)(pos_x * 4, pos_y));
+  if (pos_x * 4 + 1 < in_img_w) {
+    in1 = READ_IMG_TYPE(
+        CL_DTYPE_CHAR, input, SAMPLER, (int2)(pos_x * 4 + 1, pos_y));
+  }
+  if (pos_x * 4 + 2 < in_img_w) {
+    in2 = READ_IMG_TYPE(
+        CL_DTYPE_CHAR, input, SAMPLER, (int2)(pos_x * 4 + 2, pos_y));
+  }
+  if (pos_x * 4 + 3 < in_img_w) {
+    in3 = READ_IMG_TYPE(
+        CL_DTYPE_CHAR, input, SAMPLER, (int2)(pos_x * 4 + 3, pos_y));
+  }
+
+  CL_DTYPE4 out = (CL_DTYPE4)(in0.x, in1.x, in2.x, in3.x);
+  WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, (int2)(pos_x, pos_y), out);
+}
+
 ////////////////////////////////////////////////////////
 // image2d_folder -> image2d_default
 ////////////////////////////////////////////////////////

lite/kernels/opencl/layout_image_compute.cc

@@ -417,6 +417,102 @@ class LayoutComputeBufferChwToImage2DNw
   std::string build_options_{"-DCL_DTYPE_float "};
 };
 
+// [ImageDefault] -> [ImageFolder]
+class LayoutComputeImageDefaultToImageFolder
+    : public KernelLite<TARGET(kOpenCL),
+                        PRECISION(kAny),
+                        DATALAYOUT(kImageFolder)> {
+ public:
+  using param_t = operators::LayoutParam;
+
+  void PrepareForRun() override {
+    auto& param = Param<param_t>();
+    VLOG(1) << "kernel_func_name_:" << kernel_func_name_;
+    auto& context = ctx_->As<OpenCLContext>();
+    context.cl_context()->AddKernel(kernel_func_name_,
+                                    "image/layout_kernel.cl",
+                                    build_options_,
+                                    time_stamp_);
+  }
+
+#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
+
+  void Run() override {
+    auto& param = Param<param_t>();
+    auto x_dims = param.x->dims();
+    auto y_dims = param.y->dims();
+
+    CLImageConverterDefault default_converter;
+    CLImageConverterFolder folder_converter;
+    auto x_image_shape = default_converter.InitImageDimInfoWith(x_dims);
+    auto y_image_shape = folder_converter.InitImageDimInfoWith(y_dims);
+
+    const cl::Image2D* y_data =
+        MUTABLE_DATA_GPU(param.y, y_image_shape[0], y_image_shape[1], nullptr);
+    auto* x_data = GET_DATA_GPU(param.x);
+
+#ifdef LITE_WITH_LOG
+    VLOG(2) << "x_dims:" << x_dims;
+    VLOG(2) << "y_dims:" << y_dims;
+    VLOG(2) << "x_image_shape(w,h):" << x_image_shape[0] << " "
+            << x_image_shape[1];
+    VLOG(2) << "y_image_shape(w,h):" << y_image_shape[0] << " "
+            << y_image_shape[1];
+#endif
+
+    auto& context = ctx_->As<OpenCLContext>();
+    CHECK(context.cl_context() != nullptr);
+    STL::stringstream kernel_key;
+    kernel_key << kernel_func_name_ << build_options_ << time_stamp_;
+    auto kernel = context.cl_context()->GetKernel(kernel_key.str());
+
+    int arg_idx = 0;
+    cl_int status;
+    status = kernel.setArg(arg_idx, *x_data);
+    CL_CHECK_FATAL(status);
+    status = kernel.setArg(++arg_idx, *y_data);
+    CL_CHECK_FATAL(status);
+    status = kernel.setArg(++arg_idx, static_cast<const int>(x_image_shape[0]));
+    CL_CHECK_FATAL(status);
+    status = kernel.setArg(++arg_idx, static_cast<const int>(x_image_shape[1]));
+    CL_CHECK_FATAL(status);
+
+    auto global_work_size =
+        cl::NDRange{static_cast<cl::size_type>(y_image_shape[0]),
+                    static_cast<cl::size_type>(y_image_shape[1])};
+#ifdef LITE_WITH_LOG
+    for (auto i = 0; i < global_work_size.dimensions(); i++) {
+      VLOG(2) << "global_work_size[" << i << "]: " << global_work_size[i];
+    }
+#endif
+
+    status = EnqueueNDRangeKernel(context,
+                                  kernel,
+                                  cl::NullRange,
+                                  global_work_size,
+                                  cl::NullRange,
+                                  nullptr,
+                                  event_);
+    CL_CHECK_FATAL(status);
+  }
+
+  std::string doc() const override {
+    return "Trans Layout from cl::Image2D(ImageDefault/RGBA) to "
+           "cl::Image2D(ImageFolder)";
+  }
+
+ private:
+  std::string time_stamp_{GetTimeStamp()};
+  std::string kernel_func_name_{"image2d_default_to_image2d_folder"};
+  std::string build_options_{""};
+};
+
 // [ImageFolder] -> [ImageDefault]
 class LayoutComputeImageFolderToImageDefault
     : public KernelLite<TARGET(kOpenCL),
@@ -698,6 +794,24 @@ REGISTER_LITE_KERNEL(
                                        DATALAYOUT(kNCHW))})
     .Finalize();
 
+// [ImageDefault] -> [ImageFolder]
+REGISTER_LITE_KERNEL(
+    layout,
+    kOpenCL,
+    kAny,
+    kImageFolder,
+    paddle::lite::kernels::opencl::LayoutComputeImageDefaultToImageFolder,
+    ImageDefault_to_ImageFolder)
+    .BindInput("Input",
+               {LiteType::GetTensorTy(TARGET(kOpenCL),
+                                      PRECISION(kAny),
+                                      DATALAYOUT(kImageDefault))})
+    .BindOutput("Out",
+                {LiteType::GetTensorTy(TARGET(kOpenCL),
+                                       PRECISION(kAny),
+                                       DATALAYOUT(kImageFolder))})
+    .Finalize();
+
 // [ImageFolder] -> [ImageDefault]
 REGISTER_LITE_KERNEL(
     layout,