Skip to content

P
Paddle-Lite

PaddlePaddle / Paddle-Lite

通知 331
Star 4
Fork 1
P
Paddle-Lite
提交 487d6089 编写于 作者: S StarryRain 提交者： Yanzhan Yang
浏览文件
操作

Optimal OPENCL conv_trans_3x3s2 Performance (#2015) 

* add opencl  depthwise_conv_trans_op

* test=develop

* Optimal OPENCL conv_3x3S2 Performance

* test=develop

* Optimal OPENCL conv_trans_3x3s2 Performance , test=develop
上级 8ca10db8
隐藏空白更改
内联 并排
Showing with 322 addition and 2 deletion
mobile/src/operators/kernel/cl/cl-kernel-func/conv_func.cpp
浏览文件 @ 487d6089
......@@ -638,5 +638,123 @@ void ConvTransposeAddBnRelu(framework::CLHelper *cl_helper,
clEnqueueNDRangeKernel(cl_helper->CLCommandQueue(), kernel, 3, NULL,
work_size, NULL, 0, NULL, NULL);
}
void ConvTranspose3x3s2AddBnRelu(framework::CLHelper *cl_helper,
const ConvTransposeParam<GPU_CL> &param,
bool ifRelu, const framework::CLImage *biase,
const framework::CLImage *new_scale,
const framework::CLImage *new_bias) {
auto kernel = cl_helper->KernelAt(0);
auto default_work_size = cl_helper->DefaultWorkSize(*param.Output());
int c_block = default_work_size[0];
int w = default_work_size[1];
int nh = default_work_size[2];
int w_blk_size = 5;
int w_blk = (w + w_blk_size - 1 + 5) / w_blk_size / 2 * 2;
default_work_size[1] = w_blk;
int h_blk_size = 1;
int h_blk = (nh + h_blk_size - 1) / h_blk_size;
default_work_size[2] = h_blk;
auto input = param.Input()->GetCLImage();
auto filter = param.Filter()->GetCLImage();
auto output = param.Output()->GetCLImage();
int stride = param.Strides()[0];
int pad = param.Paddings()[0];
int dilation = param.Dilations()[0];
int input_channel = param.Input()->dims()[1];
int input_height = param.Input()->dims()[2];
int input_width = param.Input()->dims()[3];
int output_height = param.Output()->dims()[2];
int output_width = param.Output()->dims()[3];
int filter_height = param.Filter()->dims()[2];
int filter_width = param.Filter()->dims()[3];
cl_int status;
int index = 0;
status = clSetKernelArg(kernel, index++, sizeof(int), &c_block);
CL_CHECK_ERRORS(status);
status = clSetKernelArg(kernel, index++, sizeof(int), &w_blk);
CL_CHECK_ERRORS(status);
status = clSetKernelArg(kernel, index++, sizeof(int), &h_blk);
CL_CHECK_ERRORS(status);
status = clSetKernelArg(kernel, index++, sizeof(cl_mem), &input);
CL_CHECK_ERRORS(status);
status = clSetKernelArg(kernel, index++, sizeof(cl_mem), &filter);
CL_CHECK_ERRORS(status);
if (biase) {
auto bias_mem = biase->GetCLImage();
status = clSetKernelArg(kernel, index++, sizeof(cl_mem), &bias_mem);
CL_CHECK_ERRORS(status);
}
if (new_scale && new_bias) {
auto new_scale_mem = new_scale->GetCLImage();
status = clSetKernelArg(kernel, index++, sizeof(cl_mem), &new_scale_mem);
CL_CHECK_ERRORS(status);
auto new_bias_mem = new_bias->GetCLImage();
status = clSetKernelArg(kernel, index++, sizeof(cl_mem), &new_bias_mem);
CL_CHECK_ERRORS(status);
}
status = clSetKernelArg(kernel, index++, sizeof(cl_mem), &output);
CL_CHECK_ERRORS(status);
status = clSetKernelArg(kernel, index++, sizeof(int), &stride);
CL_CHECK_ERRORS(status);
status = clSetKernelArg(kernel, index++, sizeof(int), &pad);
CL_CHECK_ERRORS(status);
status = clSetKernelArg(kernel, index++, sizeof(int), &dilation);
CL_CHECK_ERRORS(status);
status = clSetKernelArg(kernel, index++, sizeof(int), &input_channel);
CL_CHECK_ERRORS(status);
status = clSetKernelArg(kernel, index++, sizeof(int), &input_width);
CL_CHECK_ERRORS(status);
status = clSetKernelArg(kernel, index++, sizeof(int), &input_height);
CL_CHECK_ERRORS(status);
status = clSetKernelArg(kernel, index++, sizeof(int), &output_width);
CL_CHECK_ERRORS(status);
status = clSetKernelArg(kernel, index++, sizeof(int), &output_height);
CL_CHECK_ERRORS(status);
status = clSetKernelArg(kernel, index++, sizeof(int), &filter_width);
CL_CHECK_ERRORS(status);
status = clSetKernelArg(kernel, index++, sizeof(int), &filter_height);
CL_CHECK_ERRORS(status);
if (default_work_size.data()[1] % 60 == 0 && use_lws) {
const size_t local_work_size[3] = {static_cast<const uint32_t>(1),
static_cast<const uint32_t>(60),
static_cast<const uint32_t>(1)};
status = clEnqueueNDRangeKernel(
cl_helper->CLCommandQueue(), kernel, default_work_size.size(), NULL,
default_work_size.data(), local_work_size, 0, NULL, NULL);
} else {
status = clEnqueueNDRangeKernel(
cl_helper->CLCommandQueue(), kernel, default_work_size.size(), NULL,
default_work_size.data(), NULL, 0, NULL, NULL);
}
CL_CHECK_ERRORS(status);
}
} // namespace operators
} // namespace paddle_mobile
mobile/src/operators/kernel/cl/cl-kernel-func/conv_func.h
浏览文件 @ 487d6089
......@@ -64,6 +64,12 @@ void ConvTransposeAddBnRelu(framework::CLHelper *cl_helper,
const framework::CLImage *biase = nullptr,
const framework::CLImage *new_scale = nullptr,
const framework::CLImage *new_bias = nullptr);
void ConvTranspose3x3s2AddBnRelu(framework::CLHelper *cl_helper,
const ConvTransposeParam<GPU_CL> &param,
bool ifRelu = false,
const framework::CLImage *biase = nullptr,
const framework::CLImage *new_scale = nullptr,
const framework::CLImage *new_bias = nullptr);
} // namespace operators
} // namespace paddle_mobile
......
mobile/src/operators/kernel/cl/cl_kernel/conv_transpose_kernel.cl
浏览文件 @ 487d6089
......@@ -242,7 +242,202 @@ __read_only image2d_t new_scale,
}
/* batch == 1 pad(output) == 1 out_w % 2 == 0 */
__kernel void conv_transpose3x3s2(__private const int item_ch,
__private const int item_w,
__private const int item_h,
__read_only image2d_t input_image,
__read_only image2d_t filter_image,
#if defined(BIASE_CH) || defined(BIASE_ELE)
__read_only image2d_t bias,
#endif
#ifdef BATCH_NORM
__read_only image2d_t new_scale,
__read_only image2d_t new_biase,
#endif
__write_only image2d_t output_image,
__private const int stride,
__private const int pad,
__private const int dilation,
__private const int in_ch,
__private const int in_w,
__private const int in_h,
__private const int out_w,
__private const int out_h,
__private const int filter_w,
__private const int filter_h) {
const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE |
CLK_ADDRESS_CLAMP |
CLK_FILTER_NEAREST;
// item_id
const int item_ch_id = get_global_id(0);
const int item_w_id = get_global_id(1);
const int item_h_id = get_global_id(2);
// out_id
int out_w_id_per_ch_blk = item_w_id / 2 * 10 + item_w_id % 2;
int out_h_id = item_h_id;
int out_w_id0 = item_ch_id * out_w + out_w_id_per_ch_blk;
int out_w_id1 = out_w_id0 + 2;
int out_w_id2 = out_w_id1 + 2;
int out_w_id3 = out_w_id2 + 2;
int out_w_id4 = out_w_id3 + 2;
// in_id
int in_w_id_per_ch_blk = (out_w_id_per_ch_blk) / 2;
in_w_id_per_ch_blk = in_w_id_per_ch_blk > 0 ? in_w_id_per_ch_blk : 0;
int in_h_id_per_batch = (out_h_id) / 2;
in_h_id_per_batch = in_h_id_per_batch > 0 ? in_h_id_per_batch : 0;
// filter_id
int align_w_i = out_w_id_per_ch_blk - 1;
int align_w = align_w_i % 2 > 0 ?
align_w_i % 2 - 2 : align_w_i % 2;
int filter_w_id_per_ch_blk = out_w_id_per_ch_blk + 1 < 3 ? out_w_id_per_ch_blk + 1 : 2 + align_w;
int align_h_i = out_h_id - 1;
int align_h = align_h_i % 2 > 0 ?
align_h_i % 2 - 2 : align_h_i % 2;
int filter_h_id_per_out_ch = out_h_id + 1 < 3 ? out_h_id + 1 : 2 + align_h;
#ifdef BIASE_CH
half4 output[5];
output[0] = read_imageh(bias, sampler, (int2)(item_ch_id, 0));
output[1] = output[0];
output[2] = output[0];
output[3] = output[0];
output[4] = output[0];
#elif defined(BIASE_ELE)
half4 output[5];
output[0] = read_imageh(bias, sampler, (int2)(out_w_id0, item_h_id));
if (out_w_id_per_ch_blk + 2 < out_w) {
output[1] = read_imageh(bias, sampler, (int2)(out_w_id1, item_h_id));
}
if (out_w_id_per_ch_blk + 4 < out_w) {
output[2] = read_imageh(bias, sampler, (int2)(out_w_id2, item_h_id));
}
if (out_w_id_per_ch_blk + 6 < out_w) {
output[3] = read_imageh(bias, sampler, (int2)(out_w_id3, item_h_id));
}
if (out_w_id_per_ch_blk + 8 < out_w) {
output[4] = read_imageh(bias, sampler, (int2)(out_w_id4, item_h_id));
}
#else
half4 output[5] = {0.0f};
#endif
half4 filter[4] = {0.0f};
half4 filter_trans[4] = {0.0f};
half4 input[5] = {0.0f};
for (int ch = 0; ch < (in_ch + 3) / 4; ch++) {
int filter_w_id = ch * 3;
int h_idx = 0;
for (int h = filter_h_id_per_out_ch; h >= 0; h -= 2) {
int in_h_id = select(in_h_id_per_batch + h_idx, -1,
in_h_id_per_batch + h_idx < 0 || in_h_id_per_batch + h_idx >= in_h);
int filter_h_id = item_ch_id * 12 + h;
int w_idx = 0;
for (int w = filter_w_id_per_ch_blk; w >= 0; w -= 2) {
int in_w_id0 = select(ch * in_w + in_w_id_per_ch_blk + w_idx, -1,
in_w_id_per_ch_blk + w_idx < 0 || in_w_id_per_ch_blk + w_idx >= in_w);
int in_w_id1 = select(ch * in_w + in_w_id_per_ch_blk + 1 + w_idx, -1,
in_w_id_per_ch_blk + 1 + w_idx < 0 || in_w_id_per_ch_blk + 1 + w_idx >= in_w);
int in_w_id2 = select(ch * in_w + in_w_id_per_ch_blk + 2 + w_idx, -1,
in_w_id_per_ch_blk + 2 + w_idx < 0 || in_w_id_per_ch_blk + 2 + w_idx >= in_w);
int in_w_id3 = select(ch * in_w + in_w_id_per_ch_blk + 3 + w_idx, -1,
in_w_id_per_ch_blk + 3 + w_idx < 0 || in_w_id_per_ch_blk + 3 + w_idx >= in_w);
int in_w_id4 = select(ch * in_w + in_w_id_per_ch_blk + 4 + w_idx, -1,
in_w_id_per_ch_blk + 4 + w_idx < 0 || in_w_id_per_ch_blk + 4 + w_idx >= in_w);
input[0] = read_imageh(input_image, sampler, (int2)(in_w_id0, in_h_id));
input[1] = read_imageh(input_image, sampler, (int2)(in_w_id1, in_h_id));
input[2] = read_imageh(input_image, sampler, (int2)(in_w_id2, in_h_id));
input[3] = read_imageh(input_image, sampler, (int2)(in_w_id3, in_h_id));
input[4] = read_imageh(input_image, sampler, (int2)(in_w_id4, in_h_id));
filter[0] = read_imageh(filter_image, sampler, (int2)(filter_w_id + w, filter_h_id)); // in_ch:0-3,out_ch:0
filter[1] = read_imageh(filter_image, sampler, (int2)(filter_w_id + w, filter_h_id + 3)); // in_ch:0-3,out_ch:1
filter[2] = read_imageh(filter_image, sampler, (int2)(filter_w_id + w, filter_h_id + 6)); // in_ch:0-3,out_ch:2
filter[3] = read_imageh(filter_image, sampler, (int2)(filter_w_id + w, filter_h_id + 9)); // in_ch:0-3,out_ch:3
filter_trans[0] = (half4)(filter[0].x, filter[1].x, filter[2].x, filter[3].x); // in_ch:0,out_ch:0-3
filter_trans[1] = (half4)(filter[0].y, filter[1].y, filter[2].y, filter[3].y); // in_ch:1,out_ch:0-3
filter_trans[2] = (half4)(filter[0].z, filter[1].z, filter[2].z, filter[3].z); // in_ch:2,out_ch:0-3
filter_trans[3] = (half4)(filter[0].w, filter[1].w, filter[2].w, filter[3].w); // in_ch:3,out_ch:0-3
output[0] = mad(input[0].x, filter_trans[0], output[0]);
output[0] = mad(input[0].y, filter_trans[1], output[0]);
output[0] = mad(input[0].z, filter_trans[2], output[0]);
output[0] = mad(input[0].w, filter_trans[3], output[0]);
output[1] = mad(input[1].x, filter_trans[0], output[1]);
output[1] = mad(input[1].y, filter_trans[1], output[1]);
output[1] = mad(input[1].z, filter_trans[2], output[1]);
output[1] = mad(input[1].w, filter_trans[3], output[1]);
output[2] = mad(input[2].x, filter_trans[0], output[2]);
output[2] = mad(input[2].y, filter_trans[1], output[2]);
output[2] = mad(input[2].z, filter_trans[2], output[2]);
output[2] = mad(input[2].w, filter_trans[3], output[2]);
output[3] = mad(input[3].x, filter_trans[0], output[3]);
output[3] = mad(input[3].y, filter_trans[1], output[3]);
output[3] = mad(input[3].z, filter_trans[2], output[3]);
output[3] = mad(input[3].w, filter_trans[3], output[3]);
output[4] = mad(input[4].x, filter_trans[0], output[4]);
output[4] = mad(input[4].y, filter_trans[1], output[4]);
output[4] = mad(input[4].z, filter_trans[2], output[4]);
output[4] = mad(input[4].w, filter_trans[3], output[4]);
w_idx++;
}
h_idx++;
}
}
#ifdef BATCH_NORM
half4 scale = read_imageh(new_scale, sampler, (int2)(item_ch_id, 0));
half4 biase = read_imageh(new_biase, sampler, (int2)(item_ch_id, 0));
output[0] = mad(scale, output[0], biase);
if (out_w_id_per_ch_blk + 2 < out_w) {
output[1] = mad(scale, output[1], biase);
}
if (out_w_id_per_ch_blk + 4 < out_w) {
output[2] = mad(scale, output[2], biase);
}
if (out_w_id_per_ch_blk + 6 < out_w) {
output[3] = mad(scale, output[3], biase);
}
if (out_w_id_per_ch_blk + 8 < out_w) {
output[4] = mad(scale, output[4], biase);
}
#endif
#ifdef RELU
output[0] = activation(output[0]);
output[1] = activation(output[1]);
output[2] = activation(output[2]);
output[3] = activation(output[3]);
output[4] = activation(output[4]);
#endif
write_imageh(output_image, (int2)(out_w_id0, item_h_id), output[0]);
if (out_w_id_per_ch_blk + 2 < out_w) {
write_imageh(output_image, (int2)(out_w_id1, item_h_id), output[1]);
}
if (out_w_id_per_ch_blk + 4 < out_w) {
write_imageh(output_image, (int2)(out_w_id2, item_h_id), output[2]);
}
if (out_w_id_per_ch_blk + 6 < out_w) {
write_imageh(output_image, (int2)(out_w_id3, item_h_id), output[3]);
}
if (out_w_id_per_ch_blk + 8 < out_w) {
write_imageh(output_image, (int2)(out_w_id4, item_h_id), output[4]);
}
}
mobile/src/operators/kernel/cl/conv_transpose_kernel.cpp
浏览文件 @ 487d6089
......@@ -40,7 +40,8 @@ bool ConvTransposeKernel<GPU_CL, float>::Init(
param->ExecMode() = ConvTransposeParam<GPU_CL>::EXEC_CONVTRANS3x3s2_FLOAT;
param->Filter()->InitConv2dTransposeFilterCLImage(
cl_helper_.CLContext(), cl_helper_.CLCommandQueue());
this->cl_helper_.AddKernel("conv_transpose", "conv_transpose_kernel.cl");
this->cl_helper_.AddKernel("conv_transpose3x3s2",
"conv_transpose_kernel.cl");
} else {
PADDLE_MOBILE_THROW_EXCEPTION(" not support ");
}
......@@ -55,7 +56,7 @@ void ConvTransposeKernel<GPU_CL, float>::Compute(
DWConvTransposeAddBnRelu(&this->cl_helper_, param);
break;
case ConvTransposeParam<GPU_CL>::EXEC_CONVTRANS3x3s2_FLOAT:
ConvTransposeAddBnRelu(&this->cl_helper_, param);
ConvTranspose3x3s2AddBnRelu(&this->cl_helper_, param);
break;
default:
PADDLE_MOBILE_THROW_EXCEPTION(
......
Markdown is supported
0% 或 .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑！
想要评论请 注册 或