Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[OpTestPy] Fix opencl concat impl #8375

Merged
merged 5 commits into from
Jan 29, 2022
Merged

[OpTestPy] Fix opencl concat impl #8375

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
7fbc628
v1
zhaoyang-star Jan 26, 2022
66d0bb9
Merge branch 'develop' into concat
zhaoyang-star Jan 26, 2022
1206b66
done
zhaoyang-star Jan 28, 2022
2dbd5c4
Merge branch 'develop' into concat
zhaoyang-star Jan 28, 2022
b429c28
delete debug info
zhaoyang-star Jan 28, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
148 changes: 61 additions & 87 deletions lite/backends/opencl/cl_kernel/image/concat_default_kernel.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,8 +1,11 @@
/* 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.
Expand Down Expand Up @@ -58,96 +61,67 @@ limitations under the License. */
// axis = 1
CONCAT2(2Input, Axis1)

__kernel void concat2(__read_only image2d_t input0,
__read_only image2d_t input1,
__write_only image2d_t output,
int flag,
int C_0,
int out_C,
int out_W,
int width) {
const int out_w = get_global_id(0); // image_width cxw/4
const int out_c = get_global_id(1); // image_width cxw/4
const int out_nh = get_global_id(2); // image_height nxh
__kernel void Concat2InputAxis1Common(__read_only image2d_t input0,
__read_only image2d_t input1,
__write_only image2d_t output,
__private const int in0_dims_axis,
__private const int out_dims_last) {
const int width_idx = get_global_id(0);
const int channel_blk_idx = get_global_id(1);
const int hb_idx = get_global_id(2);
const int c = channel_blk_idx * 4;

if (flag == 1) { // by channel
int c_in = out_c;
int2 output_pos;
output_pos.x = out_c * out_W + out_w;
output_pos.y = out_nh;
CL_DTYPE4 output_data;
for (int i = 0; i < 4; i++) {
int c = out_c * 4 + i;
if (c >= out_C) {
break;
}
int c_in;
CL_DTYPE4 input_data;
if (c < C_0) {
c_in = c;
int2 input_pos;
input_pos.x = (c_in / 4) * out_W + out_w;
input_pos.y = out_nh;
input_data = READ_IMG_TYPE(CL_DTYPE_CHAR, input0, SAMPLER, input_pos);
const int2 pos = (int2)(channel_blk_idx * out_dims_last + width_idx, hb_idx);
CL_DTYPE4 in0_data = (CL_DTYPE4)0;

// write all input0 data to output directly
if (c < in0_dims_axis) {
in0_data = READ_IMG_TYPE(CL_DTYPE_CHAR, input0, SAMPLER, pos);
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, pos, in0_data);
}

// deal with output1
int channel_remain = in0_dims_axis % 4;
if (c + channel_remain >= in0_dims_axis) {
// only theads for output1 hit this
const int2 in1_pos = (int2)(
(channel_blk_idx - in0_dims_axis / 4) * out_dims_last + width_idx,
hb_idx);
CL_DTYPE4 in1_data = READ_IMG_TYPE(CL_DTYPE_CHAR, input1, SAMPLER, in1_pos);
if (channel_remain == 0) {
// write all input1 data to output directly
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, pos, in1_data);
} else {
CL_DTYPE4 remain, combined_val;
if (in1_pos.x - out_dims_last < 0) {
// combine input0 & input1
remain = in0_data;
if (channel_remain == 1) {
combined_val =
(CL_DTYPE4)(remain.x, in1_data.x, in1_data.y, in1_data.z);
} else if (channel_remain == 2) {
combined_val =
(CL_DTYPE4)(remain.x, remain.y, in1_data.x, in1_data.y);
} else if (channel_remain == 3) {
combined_val = (CL_DTYPE4)(remain.x, remain.y, remain.z, in1_data.x);
}
} else {
c_in = c - C_0;
int2 input_pos;
input_pos.x = (c_in / 4) * out_W + out_w;
input_pos.y = out_nh;
input_data = READ_IMG_TYPE(CL_DTYPE_CHAR, input1, SAMPLER, input_pos);
}
int value_offset = c_in % 4;
CL_DTYPE value;
if (value_offset == 0) {
value = input_data.x;
} else if (value_offset == 1) {
value = input_data.y;
} else if (value_offset == 2) {
value = input_data.z;
} else if (value_offset == 3) {
value = input_data.w;
// only deal with input1
remain = READ_IMG_TYPE(CL_DTYPE_CHAR,
input1,
SAMPLER,
(int2)(in1_pos.x - out_dims_last, in1_pos.y));
if (channel_remain == 1) {
combined_val =
(CL_DTYPE4)(remain.w, in1_data.x, in1_data.y, in1_data.z);
} else if (channel_remain == 2) {
combined_val =
(CL_DTYPE4)(remain.z, remain.w, in1_data.x, in1_data.y);
} else if (channel_remain == 3) {
combined_val = (CL_DTYPE4)(remain.y, remain.z, remain.w, in1_data.x);
}
}
if (i == 0) {
output_data.x = value;
} else if (i == 1) {
output_data.y = value;
} else if (i == 2) {
output_data.z = value;
} else if (i == 3) {
output_data.w = value;
}
}
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, output_pos, output_data);
} else if (flag == 2) { // by height, width == n
int2 input_pos;
input_pos.x = out_c * out_W + out_w;
int h = out_nh / width;
CL_DTYPE4 input;
if (h < C_0) {
input_pos.y = out_nh;
input = READ_IMG_TYPE(CL_DTYPE_CHAR, input0, SAMPLER, input_pos);
} else {
input_pos.y = (h - C_0) * width;
input = READ_IMG_TYPE(CL_DTYPE_CHAR, input1, SAMPLER, input_pos);
}
int2 output_pos;
output_pos.x = out_c * out_W + out_w;
output_pos.y = out_nh;
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, output_pos, input);
} else if (flag == 3) { // by width, width == C
int2 input_pos;
input_pos.y = out_nh;
CL_DTYPE4 input;
if (out_w < C_0) {
input_pos.x = out_c * out_W + out_w;
input = READ_IMG_TYPE(CL_DTYPE_CHAR, input0, SAMPLER, input_pos);
} else {
input_pos.x = out_c * out_W + (out_w - C_0);
input = READ_IMG_TYPE(CL_DTYPE_CHAR, input1, SAMPLER, input_pos);
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, pos, combined_val);
}
int2 output_pos;
output_pos.x = out_c * out_W + out_w;
output_pos.y = out_nh;
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output, output_pos, input);
}
}
3 changes: 3 additions & 0 deletions lite/backends/opencl/cl_kernel/image/concat_kernel.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,8 +1,11 @@
/* 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.
Expand Down
116 changes: 45 additions & 71 deletions lite/kernels/opencl/concat_image_compute.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,53 +42,39 @@ class ConcatComputeImage : public KernelLite<TARGET(kOpenCL),
axis_ = concat_param_->axis;
auto* axis_tensor = concat_param_->axis_tensor;
if (axis_tensor != nullptr) {
auto* d_image = DATA_GPU(axis_tensor);
cl::Image2D* cl_image = static_cast<cl::Image2D*>(d_image);
size_t cl_image2d_width, cl_image2d_height;
cl_image->getImageInfo(CL_IMAGE_WIDTH, &cl_image2d_width);
cl_image->getImageInfo(CL_IMAGE_HEIGHT, &cl_image2d_height);

const bool fp16_support =
CLRuntime::Global()->get_precision() == lite_api::CL_PRECISION_FP16;
if (fp16_support) {
auto* h_image = static_cast<half_t*>(TargetWrapperCL::MapImage(
d_image, cl_image2d_width, cl_image2d_height, 0, 0));
axis_ = Half2Float(h_image[0]);
TargetWrapperCL::Unmap(d_image, h_image);
} else {
auto* h_image = static_cast<float*>(TargetWrapperCL::MapImage(
d_image, cl_image2d_width, cl_image2d_height, 0, 0));
axis_ = h_image[0];
TargetWrapperCL::Unmap(d_image, h_image);
}
auto* axis_tensor_data = axis_tensor->data<int>();
axis_ = axis_tensor_data[0];
}
/*
if (axis_tensor != nullptr) {
auto* d_image = DATA_GPU(axis_tensor);
cl::Image2D* cl_image = static_cast<cl::Image2D*>(d_image);
size_t cl_image2d_width, cl_image2d_height;
cl_image->getImageInfo(CL_IMAGE_WIDTH, &cl_image2d_width);
cl_image->getImageInfo(CL_IMAGE_HEIGHT, &cl_image2d_height);

const bool fp16_support =
CLRuntime::Global()->get_precision() ==
lite_api::CL_PRECISION_FP16;
if (fp16_support) {
auto* h_image = static_cast<half_t*>(TargetWrapperCL::MapImage(
d_image, cl_image2d_width, cl_image2d_height, 0, 0));
axis_ = Half2Float(h_image[0]);
TargetWrapperCL::Unmap(d_image, h_image);
} else {
auto* h_image = static_cast<float*>(TargetWrapperCL::MapImage(
d_image, cl_image2d_width, cl_image2d_height, 0, 0));
axis_ = h_image[0];
TargetWrapperCL::Unmap(d_image, h_image);
}
}
*/
axis_ = axis_ >= 0 ? axis_ : axis_ + concat_param_->x[0]->dims().size();

auto inputs = concat_param_->x;
auto output_tensor_dims = concat_param_->output->dims();

std::vector<size_t> new_dims = {1, 1, 1, 1};
size_t offset = 4 - output_tensor_dims.size();
for (auto i = 0; i < output_tensor_dims.size(); ++i) {
new_dims[offset + i] = output_tensor_dims[i];
}
flag_ = offset + axis_;
switch (flag_) {
case 0:
width_ = new_dims[2];
break;
case 1:
width_ = new_dims[3];
break;
case 2:
width_ = new_dims[0];
break;
case 3:
width_ = new_dims[1];
break;
default:
LOG(FATAL) << "Unsupported axis:" << axis_;
}
auto new_dims = Broadcast2GpuShape(output_tensor_dims);
flag_ = 4 - output_tensor_dims.size() + axis_;

for (auto& input : inputs) {
cl_int4 in_shape = {1, 1, 1, 1};
Expand All @@ -103,15 +89,14 @@ class ConcatComputeImage : public KernelLite<TARGET(kOpenCL),
if (!(flag_ == 1 && !align_) && inputs.size() <= 6) {
kernel_func_name_ = "Concat" + std::to_string(inputs.size()) +
"InputAxis" + std::to_string(flag_);
} else if (inputs.size() == 2) {
kernel_func_name_ = "concat2";
} else if (inputs.size() == 2 && flag_ == 1) {
kernel_func_name_ = "Concat2InputAxis1Common";
} else if (inputs.size() == 3 && axis_ == 1 &&
output_tensor_dims.size() == 4) {
kernel_func_name_ = "concatByCWith3Inputs";
} else if (inputs.size() == 4 && axis_ == 1 &&
output_tensor_dims.size() == 4) {
kernel_func_name_ = "concatByCWith4Inputs";
"InputAxis" + std::to_string(flag_);
} else {
// note: do layout transform between image and buffer,
// before and after concat(buffer impl.)
Expand Down Expand Up @@ -151,27 +136,26 @@ class ConcatComputeImage : public KernelLite<TARGET(kOpenCL),
}

std::string kernel_file;
if (kernel_func_name_ == "concat2" ||
kernel_func_name_ == "Concat2InputAxis1") {
if (kernel_func_name_ == "concat_mul_buffer") {
kernel_file = "buffer/concat_kernel.cl";
} else if (kernel_func_name_ == "Concat2InputAxis1" ||
kernel_func_name_ == "Concat2InputAxis1Common") {
kernel_file = "image/concat_default_kernel.cl";
} else {
kernel_file = "image/concat_kernel.cl";
}

VLOG(1) << "kernel_func_name_:" << kernel_func_name_;

context.cl_context()->AddKernel(kernel_func_name_,
(kernel_func_name_ == "concat_mul_buffer")
? "buffer/concat_kernel.cl"
: kernel_file,
build_options_,
time_stamp_);
context.cl_context()->AddKernel(
kernel_func_name_, kernel_file, build_options_, time_stamp_);
}

void Run() override {
const auto& output_tensor_dims = concat_param_->output->dims();
int output_tensor_w = output_tensor_dims[output_tensor_dims.size() - 1];
int output_tensor_c = output_tensor_dims[1];
auto new_dim = Broadcast2GpuShape(output_tensor_dims);
int output_tensor_w = new_dim[3];
int output_tensor_c = new_dim[1];
auto output_image_shape = InitImageDimInfoWith(output_tensor_dims);
auto* output_image_p = MUTABLE_DATA_GPU(concat_param_->output,
output_image_shape["width"],
Expand Down Expand Up @@ -213,7 +197,6 @@ class ConcatComputeImage : public KernelLite<TARGET(kOpenCL),
VLOG(4) << "output_image_shape(w,h): " << output_image_shape["width"] << " "
<< output_image_shape["height"];
VLOG(4) << "output_tensor_w: " << output_tensor_w;
VLOG(4) << "width_: " << width_;
VLOG(4) << "global_work_size: " << global_work_size[0] << " "
<< global_work_size[1] << " " << global_work_size[2];
VLOG(4) << "kernel_func_name: " << kernel_func_name_;
Expand Down Expand Up @@ -261,25 +244,19 @@ class ConcatComputeImage : public KernelLite<TARGET(kOpenCL),
nullptr,
event_);
CL_CHECK_FATAL(status);
} else if (kernel_func_name_ == "concat2") {
} else if (kernel_func_name_ == "Concat2InputAxis1Common") {
auto* input0_image_p = GET_DATA_GPU(inputs[0]);
auto* input1_image_p = GET_DATA_GPU(inputs[1]);
int input0_axis_dims = inputs[0]->dims()[axis_];
int input0_dims_axis = inputs[0]->dims()[axis_];
cl_int status = kernel.setArg(0, *input0_image_p);
CL_CHECK_FATAL(status);
status = kernel.setArg(1, *input1_image_p);
CL_CHECK_FATAL(status);
status = kernel.setArg(2, *output_image_p);
CL_CHECK_FATAL(status);
status = kernel.setArg(3, flag_);
status = kernel.setArg(3, input0_dims_axis);
CL_CHECK_FATAL(status);
status = kernel.setArg(4, input0_axis_dims);
CL_CHECK_FATAL(status);
status = kernel.setArg(5, output_tensor_c);
CL_CHECK_FATAL(status);
status = kernel.setArg(6, output_tensor_w);
CL_CHECK_FATAL(status);
status = kernel.setArg(7, width_);
status = kernel.setArg(4, output_tensor_w);
CL_CHECK_FATAL(status);

status = EnqueueNDRangeKernel(context,
Expand Down Expand Up @@ -356,7 +333,7 @@ class ConcatComputeImage : public KernelLite<TARGET(kOpenCL),
nullptr,
event_);
CL_CHECK_FATAL(status);
} else if (kernel_func_name_ == "concat_mul_buffer") { // inputs.size() > 4
} else if (kernel_func_name_ == "concat_mul_buffer") {
// note: do image layout transform: image to buffer
size_t inputs_num = inputs.size();
std::vector<const cl::Image2D*> inputs_image_pointers(inputs_num);
Expand Down Expand Up @@ -460,7 +437,6 @@ class ConcatComputeImage : public KernelLite<TARGET(kOpenCL),
private:
int axis_ = 1;
int flag_ = 1;
int width_ = 1;
int pre_size_ = 1;
int post_size_ = 1;
int align_{true};
Expand All @@ -486,9 +462,7 @@ REGISTER_LITE_KERNEL(
PRECISION(kFP16),
DATALAYOUT(kImageDefault))})
.BindInput("AxisTensor",
{LiteType::GetTensorTy(TARGET(kOpenCL),
PRECISION(kInt32),
DATALAYOUT(kImageDefault))})
{LiteType::GetTensorTy(TARGET(kHost), PRECISION(kInt32))})
.BindOutput("Out",
{LiteType::GetTensorTy(TARGET(kOpenCL),
PRECISION(kFP16),
Expand Down
10 changes: 10 additions & 0 deletions lite/kernels/opencl/image_helper.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -75,6 +75,16 @@ static std::vector<size_t> DefaultGlobalWorkSize(const DDim& tensor_dim,
return {};
}

static DDim Broadcast2GpuShape(const DDim& src_dim) {
CHECK_LE(src_dim.size(), 4);
DDim dst_dim({1, 1, 1, 1});
size_t offset = 4 - src_dim.size();
for (auto i = 0; i < src_dim.size(); ++i) {
dst_dim[offset + i] = src_dim[i];
}
return dst_dim;
}

static const std::string GetTimeStamp() {
uint64_t usec = lite::Timer::GetCurrentUS();
return std::to_string(usec);
Expand Down
Loading