Skip to content

Commit

Permalink
add [ImageDefault] -> [ImageFolder] test=develop
Browse files Browse the repository at this point in the history
  • Loading branch information
zhaoyang-star committed Oct 9, 2021
1 parent 43407f8 commit 71ec3d8
Show file tree
Hide file tree
Showing 2 changed files with 150 additions and 0 deletions.
36 changes: 36 additions & 0 deletions lite/backends/opencl/cl_kernel/image/layout_kernel.cl
Original file line number Diff line number Diff line change
Expand Up @@ -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
////////////////////////////////////////////////////////
Expand Down
114 changes: 114 additions & 0 deletions lite/kernels/opencl/layout_image_compute.cc
Original file line number Diff line number Diff line change
Expand Up @@ -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),
Expand Down Expand Up @@ -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,
Expand Down

0 comments on commit 71ec3d8

Please sign in to comment.