From 487d6089904b496e9604c945ca55c6b1a3bdbaf6 Mon Sep 17 00:00:00 2001 From: StarryRain <36948762+StarryRain@users.noreply.github.com> Date: Wed, 11 Sep 2019 22:05:07 +0800 Subject: [PATCH] Optimal OPENCL conv_trans_3x3s2 Performance (#2015) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * add opencl depthwise_conv_trans_op * test=develop * Optimal OPENCL conv_3x3S2 Performance * test=develop * Optimal OPENCL conv_trans_3x3s2 Performance , test=develop --- .../kernel/cl/cl-kernel-func/conv_func.cpp | 118 +++++++++++ .../kernel/cl/cl-kernel-func/conv_func.h | 6 + .../cl/cl_kernel/conv_transpose_kernel.cl | 195 ++++++++++++++++++ .../kernel/cl/conv_transpose_kernel.cpp | 5 +- 4 files changed, 322 insertions(+), 2 deletions(-) diff --git a/mobile/src/operators/kernel/cl/cl-kernel-func/conv_func.cpp b/mobile/src/operators/kernel/cl/cl-kernel-func/conv_func.cpp index 08cae42762..13c4d58885 100644 --- a/mobile/src/operators/kernel/cl/cl-kernel-func/conv_func.cpp +++ b/mobile/src/operators/kernel/cl/cl-kernel-func/conv_func.cpp @@ -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 ¶m, + 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(1), + static_cast(60), + static_cast(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 diff --git a/mobile/src/operators/kernel/cl/cl-kernel-func/conv_func.h b/mobile/src/operators/kernel/cl/cl-kernel-func/conv_func.h index a0dcd99d9e..6254455eac 100644 --- a/mobile/src/operators/kernel/cl/cl-kernel-func/conv_func.h +++ b/mobile/src/operators/kernel/cl/cl-kernel-func/conv_func.h @@ -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 ¶m, + 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 diff --git a/mobile/src/operators/kernel/cl/cl_kernel/conv_transpose_kernel.cl b/mobile/src/operators/kernel/cl/cl_kernel/conv_transpose_kernel.cl index e13f5debba..a67ad9a017 100644 --- a/mobile/src/operators/kernel/cl/cl_kernel/conv_transpose_kernel.cl +++ b/mobile/src/operators/kernel/cl/cl_kernel/conv_transpose_kernel.cl @@ -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]); + } +} diff --git a/mobile/src/operators/kernel/cl/conv_transpose_kernel.cpp b/mobile/src/operators/kernel/cl/conv_transpose_kernel.cpp index f5be81eefd..8d66b50a99 100644 --- a/mobile/src/operators/kernel/cl/conv_transpose_kernel.cpp +++ b/mobile/src/operators/kernel/cl/conv_transpose_kernel.cpp @@ -40,7 +40,8 @@ bool ConvTransposeKernel::Init( param->ExecMode() = ConvTransposeParam::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::Compute( DWConvTransposeAddBnRelu(&this->cl_helper_, param); break; case ConvTransposeParam::EXEC_CONVTRANS3x3s2_FLOAT: - ConvTransposeAddBnRelu(&this->cl_helper_, param); + ConvTranspose3x3s2AddBnRelu(&this->cl_helper_, param); break; default: PADDLE_MOBILE_THROW_EXCEPTION( -- GitLab