diff options
author | Giorgio Arena <giorgio.arena@arm.com> | 2018-03-15 17:58:20 +0000 |
---|---|---|
committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:49:16 +0000 |
commit | 2d9de0a3fa6ad858e70040124f362799a962bb6a (patch) | |
tree | 0a055c5100438a929b3b04945821665d2fef8751 /src | |
parent | ed99f411d52949720a4d64d91664cd71e46b79d5 (diff) | |
download | ComputeLibrary-2d9de0a3fa6ad858e70040124f362799a962bb6a.tar.gz |
COMPMID-1009 Support 4x4 output tile for Winograd Filter Transform on OpenCL.
Change-Id: I68c6453e0f192de659582404f109a89616b9fbb9
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/124811
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Diffstat (limited to 'src')
-rw-r--r-- | src/core/CL/CLKernelLibrary.cpp | 1 | ||||
-rw-r--r-- | src/core/CL/cl_kernels/winograd.cl | 138 | ||||
-rw-r--r-- | src/core/CL/kernels/CLWinogradFilterTransformKernel.cpp | 25 | ||||
-rw-r--r-- | src/runtime/CL/functions/CLWinogradConvolutionLayer.cpp | 6 |
4 files changed, 156 insertions, 14 deletions
diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index 9df2dcbacd..740a98bbac 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -352,6 +352,7 @@ const std::map<std::string, std::string> CLKernelLibrary::_kernel_program_map = { "warp_perspective_nearest_neighbour", "warp_perspective.cl" }, { "warp_perspective_bilinear", "warp_perspective.cl" }, { "winograd_filter_transform_2x2_3x3_nchw", "winograd.cl" }, + { "winograd_filter_transform_4x4_3x3_nchw", "winograd.cl" }, { "winograd_input_transform_2x2_3x3_stepz1_nchw", "winograd.cl" }, { "winograd_input_transform_2x2_3x3_stepz2_nchw", "winograd.cl" }, { "winograd_output_transform_2x2_3x3_nchw", "winograd.cl" }, diff --git a/src/core/CL/cl_kernels/winograd.cl b/src/core/CL/cl_kernels/winograd.cl index 25c129d0aa..bd51db6b03 100644 --- a/src/core/CL/cl_kernels/winograd.cl +++ b/src/core/CL/cl_kernels/winograd.cl @@ -116,6 +116,144 @@ __kernel void winograd_filter_transform_2x2_3x3_nchw( *(__global float *)(dst_addr + 14 * dst_stride_z) = out3.s2; *(__global float *)(dst_addr + 15 * dst_stride_z) = out3.s3; } + +/** This OpenCL kernel performs Winograd filter transform 3x3 when the data format is NCHW and the output tile is 4x4 + * + * @note The number of channels must be passed at compile time using -DNUM_CHANNELS: e.g. -DNUM_CHANNELS=64 + * + * @param[in] src_ptr Pointer to the source tensor. Supported data types: F32 + * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] src_step_w src_stride_w * number of elements along W processed per workitem(in bytes) + * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr + * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) + * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + */ +__kernel void winograd_filter_transform_4x4_3x3_nchw( + TENSOR4D_DECLARATION(src), + TENSOR3D_DECLARATION(dst)) +{ + Tensor4D src = CONVERT_TO_TENSOR4D_STRUCT(src, NUM_CHANNELS); + + const __global uchar *src_addr = tensor4D_offset(&src, 0, 0, 0, 0); + + // Load the values from the input tensor + float3 w0 = vload3(0, (__global float *)(src_addr + 0 * src_stride_y)); + float3 w1 = vload3(0, (__global float *)(src_addr + 1 * src_stride_y)); + float3 w2 = vload3(0, (__global float *)(src_addr + 2 * src_stride_y)); + + // Transform the 3x3 tile in a 6x6 tile + float8 out0 = 0.0f; + float8 out1 = 0.0f; + float8 out2 = 0.0f; + float8 out3 = 0.0f; + float8 out4 = 0.0f; + float8 out5 = 0.0f; + + // Row 0 + out0.s0 = (w0.s0) / 16.f; + out0.s1 = (-w0.s0 - w0.s1 - w0.s2) / 24.f; + out0.s2 = (-w0.s0 + w0.s1 - w0.s2) / 24.f; + out0.s3 = (w0.s0 + 2 * w0.s1 + 4 * w0.s2) / 96.f; + out0.s4 = (w0.s0 - 2 * w0.s1 + 4 * w0.s2) / 96.f; + out0.s5 = (w0.s2) / 4.f; + + // Row 1 + out1.s0 = (-w0.s0 - w1.s0 - w2.s0) / 24.f; + out1.s1 = (w0.s0 + w1.s0 + w2.s0 + w0.s1 + w1.s1 + w2.s1 + w0.s2 + w1.s2 + w2.s2) / 36.f; + out1.s2 = (w0.s0 + w1.s0 + w2.s0 - w0.s1 - w1.s1 - w2.s1 + w0.s2 + w1.s2 + w2.s2) / 36.f; + out1.s3 = (-w0.s0 - w1.s0 - w2.s0 + 2 * (-w0.s1 - w1.s1 - w2.s1) + 4 * (-w0.s2 - w1.s2 - w2.s2)) / 144.f; + out1.s4 = (-w0.s0 - w1.s0 - w2.s0 + 2 * (w0.s1 + w1.s1 + w2.s1) + 4 * (-w0.s2 - w1.s2 - w2.s2)) / 144.f; + out1.s5 = (-w0.s2 - w1.s2 - w2.s2) / 6.f; + + // Row 2 + out2.s0 = (-w0.s0 + w1.s0 - w2.s0) / 24.f; + out2.s1 = (w0.s0 - w1.s0 + w2.s0 + w0.s1 - w1.s1 + w2.s1 + w0.s2 - w1.s2 + w2.s2) / 36.f; + out2.s2 = (w0.s0 - w1.s0 + w2.s0 - w0.s1 + w1.s1 - w2.s1 + w0.s2 - w1.s2 + w2.s2) / 36.f; + out2.s3 = (-w0.s0 + w1.s0 - w2.s0 + 2 * (-w0.s1 + w1.s1 - w2.s1) + 4 * (-w0.s2 + w1.s2 - w2.s2)) / 144.f; + out2.s4 = (-w0.s0 + w1.s0 - w2.s0 + 2 * (w0.s1 - w1.s1 + w2.s1) + 4 * (-w0.s2 + w1.s2 - w2.s2)) / 144.f; + out2.s5 = (-w0.s2 + w1.s2 - w2.s2) / 6.f; + + // Row 3 + out3.s0 = (w0.s0 + 2 * w1.s0 + 4 * w2.s0) / 96.f; + out3.s1 = (-w0.s0 - 2 * w1.s0 - 4 * w2.s0 - w0.s1 - 2 * w1.s1 - 4 * w2.s1 - w0.s2 - 2 * w1.s2 - 4 * w2.s2) / 144.f; + out3.s2 = (-w0.s0 - 2 * w1.s0 - 4 * w2.s0 + w0.s1 + 2 * w1.s1 + 4 * w2.s1 - w0.s2 - 2 * w1.s2 - 4 * w2.s2) / 144.f; + out3.s3 = ((w0.s0 + 2 * w1.s0 + 4 * w2.s0) + 2 * (w0.s1 + 2 * w1.s1 + 4 * w2.s1) + 4 * (w0.s2 + 2 * w1.s2 + 4 * w2.s2)) / 576.f; + out3.s4 = ((w0.s0 + 2 * w1.s0 + 4 * w2.s0) + 2 * (-w0.s1 - 2 * w1.s1 - 4 * w2.s1) + 4 * (w0.s2 + 2 * w1.s2 + 4 * w2.s2)) / 576.f; + out3.s5 = (w0.s2 + 2 * w1.s2 + 4 * w2.s2) / 24.f; + + // Row 4 + out4.s0 = (w0.s0 - 2 * w1.s0 + 4 * w2.s0) / 96.f; + out4.s1 = (-w0.s0 + 2 * w1.s0 - 4 * w2.s0 - w0.s1 + 2 * w1.s1 - 4 * w2.s1 - w0.s2 + 2 * w1.s2 - 4 * w2.s2) / 144.f; + out4.s2 = (-w0.s0 + 2 * w1.s0 - 4 * w2.s0 + w0.s1 - 2 * w1.s1 + 4 * w2.s1 - w0.s2 + 2 * w1.s2 - 4 * w2.s2) / 144.f; + out4.s3 = ((w0.s0 - 2 * w1.s0 + 4 * w2.s0) + 2 * (w0.s1 - 2 * w1.s1 + 4 * w2.s1) + 4 * (w0.s2 - 2 * w1.s2 + 4 * w2.s2)) / 576.f; + out4.s4 = ((w0.s0 - 2 * w1.s0 + 4 * w2.s0) + 2 * (-w0.s1 + 2 * w1.s1 - 4 * w2.s1) + 4 * (w0.s2 - 2 * w1.s2 + 4 * w2.s2)) / 576.f; + out4.s5 = (w0.s2 - 2 * w1.s2 + 4 * w2.s2) / 24.f; + + // Row 5 + out5.s0 = (w2.s0) / 4.f; + out5.s1 = (-w2.s0 - w2.s1 - w2.s2) / 6.f; + out5.s2 = (-w2.s0 + w2.s1 - w2.s2) / 6.f; + out5.s3 = (w2.s0 + 2 * w2.s1 + 4 * w2.s2) / 24.f; + out5.s4 = (w2.s0 - 2 * w2.s1 + 4 * w2.s2) / 24.f; + out5.s5 = (w2.s2); + + int z = get_global_id(2); + int x0 = z / NUM_CHANNELS; // idx filter + int y0 = z % NUM_CHANNELS; // idx channel + + // Get output address + __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x0 * dst_stride_x + y0 * dst_stride_y; + + // Store the 36 values across the 36 channels + *(__global float *)(dst_addr + 0 * dst_stride_z) = out0.s0; + *(__global float *)(dst_addr + 1 * dst_stride_z) = out0.s1; + *(__global float *)(dst_addr + 2 * dst_stride_z) = out0.s2; + *(__global float *)(dst_addr + 3 * dst_stride_z) = out0.s3; + *(__global float *)(dst_addr + 4 * dst_stride_z) = out0.s4; + *(__global float *)(dst_addr + 5 * dst_stride_z) = out0.s5; + *(__global float *)(dst_addr + 6 * dst_stride_z) = out1.s0; + *(__global float *)(dst_addr + 7 * dst_stride_z) = out1.s1; + *(__global float *)(dst_addr + 8 * dst_stride_z) = out1.s2; + *(__global float *)(dst_addr + 9 * dst_stride_z) = out1.s3; + *(__global float *)(dst_addr + 10 * dst_stride_z) = out1.s4; + *(__global float *)(dst_addr + 11 * dst_stride_z) = out1.s5; + *(__global float *)(dst_addr + 12 * dst_stride_z) = out2.s0; + *(__global float *)(dst_addr + 13 * dst_stride_z) = out2.s1; + *(__global float *)(dst_addr + 14 * dst_stride_z) = out2.s2; + *(__global float *)(dst_addr + 15 * dst_stride_z) = out2.s3; + *(__global float *)(dst_addr + 16 * dst_stride_z) = out2.s4; + *(__global float *)(dst_addr + 17 * dst_stride_z) = out2.s5; + *(__global float *)(dst_addr + 18 * dst_stride_z) = out3.s0; + *(__global float *)(dst_addr + 19 * dst_stride_z) = out3.s1; + *(__global float *)(dst_addr + 20 * dst_stride_z) = out3.s2; + *(__global float *)(dst_addr + 21 * dst_stride_z) = out3.s3; + *(__global float *)(dst_addr + 22 * dst_stride_z) = out3.s4; + *(__global float *)(dst_addr + 23 * dst_stride_z) = out3.s5; + *(__global float *)(dst_addr + 24 * dst_stride_z) = out4.s0; + *(__global float *)(dst_addr + 25 * dst_stride_z) = out4.s1; + *(__global float *)(dst_addr + 26 * dst_stride_z) = out4.s2; + *(__global float *)(dst_addr + 27 * dst_stride_z) = out4.s3; + *(__global float *)(dst_addr + 28 * dst_stride_z) = out4.s4; + *(__global float *)(dst_addr + 29 * dst_stride_z) = out4.s5; + *(__global float *)(dst_addr + 30 * dst_stride_z) = out5.s0; + *(__global float *)(dst_addr + 31 * dst_stride_z) = out5.s1; + *(__global float *)(dst_addr + 32 * dst_stride_z) = out5.s2; + *(__global float *)(dst_addr + 33 * dst_stride_z) = out5.s3; + *(__global float *)(dst_addr + 34 * dst_stride_z) = out5.s4; + *(__global float *)(dst_addr + 35 * dst_stride_z) = out5.s5; +} #endif // defined(NUM_CHANNELS) #if defined(NUM_TILES_X) && defined(PAD_LEFT) && defined(PAD_TOP) diff --git a/src/core/CL/kernels/CLWinogradFilterTransformKernel.cpp b/src/core/CL/kernels/CLWinogradFilterTransformKernel.cpp index 655b82bf66..5a03332e99 100644 --- a/src/core/CL/kernels/CLWinogradFilterTransformKernel.cpp +++ b/src/core/CL/kernels/CLWinogradFilterTransformKernel.cpp @@ -44,17 +44,18 @@ using namespace arm_compute::misc::shape_calculator; namespace { -Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output) +Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, const Size2D &output_tile) { ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32); ARM_COMPUTE_RETURN_ERROR_ON(input->dimension(0) != 3); ARM_COMPUTE_RETURN_ERROR_ON(input->dimension(0) != input->dimension(1)); ARM_COMPUTE_RETURN_ERROR_ON(input->num_dimensions() > 4); + ARM_COMPUTE_RETURN_ERROR_ON(output_tile != Size2D(2U, 2U) && output_tile != Size2D(4U, 4U)); // Checks performed when output is configured if(output->total_size() != 0) { - const TensorInfo tensor_info_output = input->clone()->set_tensor_shape(compute_winograd_filter_transform_shape(*input)); + const TensorInfo tensor_info_output = input->clone()->set_tensor_shape(compute_winograd_filter_transform_shape(*input, output_tile)); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(output, &tensor_info_output); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); @@ -63,8 +64,9 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output) return Status{}; } -std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITensorInfo *output) +std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITensorInfo *output, const Size2D &output_tile) { + ARM_COMPUTE_UNUSED(output_tile); ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); constexpr unsigned int num_elems_processed_per_iteration_x = 3; @@ -90,35 +92,36 @@ CLWinogradFilterTransformKernel::CLWinogradFilterTransformKernel() { } -void CLWinogradFilterTransformKernel::configure(const ICLTensor *input, ICLTensor *output) +void CLWinogradFilterTransformKernel::configure(const ICLTensor *input, ICLTensor *output, const Size2D &output_tile) { ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); // Output tensor auto inizialitation if not yet initialized - auto_init_if_empty(*output->info(), input->info()->clone()->set_tensor_shape(compute_winograd_filter_transform_shape(*input->info()))); + auto_init_if_empty(*output->info(), input->info()->clone()->set_tensor_shape(compute_winograd_filter_transform_shape(*input->info(), output_tile))); - ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), output->info())); + ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), output->info(), output_tile)); // Set build options CLBuildOptions build_opts; build_opts.add_option("-DNUM_CHANNELS=" + support::cpp11::to_string(input->info()->dimension(2))); // Create kernel - _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel("winograd_filter_transform_2x2_3x3_nchw", build_opts.options())); + std::string kernel_name = std::string("winograd_filter_transform_") + output_tile.to_string() + std::string("_3x3_nchw"); + _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options())); _input = input; _output = output; // Configure kernel window - auto win_config = validate_and_configure_window(input->info(), output->info()); + auto win_config = validate_and_configure_window(input->info(), output->info(), output_tile); ARM_COMPUTE_ERROR_THROW_ON(win_config.first); ICLKernel::configure(win_config.second); } -Status CLWinogradFilterTransformKernel::validate(const ITensorInfo *input, const ITensorInfo *output) +Status CLWinogradFilterTransformKernel::validate(const ITensorInfo *input, const ITensorInfo *output, const Size2D &output_tile) { - ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output)); - ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(input->clone().get(), output->clone().get()).first); + ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output, output_tile)); + ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(input->clone().get(), output->clone().get(), output_tile).first); return Status{}; } diff --git a/src/runtime/CL/functions/CLWinogradConvolutionLayer.cpp b/src/runtime/CL/functions/CLWinogradConvolutionLayer.cpp index 5081cbac4e..a861e0072e 100644 --- a/src/runtime/CL/functions/CLWinogradConvolutionLayer.cpp +++ b/src/runtime/CL/functions/CLWinogradConvolutionLayer.cpp @@ -64,7 +64,7 @@ void CLWinogradConvolutionLayer::configure(ICLTensor *input, const ICLTensor *we _input_transform.configure(input, &_input0, conv_info, Size2D(kernel_w, kernel_h)); // Configure filter transform - _filter_transform.configure(weights, &_input1); + _filter_transform.configure(weights, &_input1, Size2D(2U, 2U)); // Configure batched matrix multiply _batched_mm.configure(&_input0, &_input1, nullptr, &_batched_mm_output, 1.0f, 0.0f, GEMMInfo(false, false, true /* Reshape weights only for the first run*/)); @@ -103,9 +103,9 @@ Status CLWinogradConvolutionLayer::validate(const ITensorInfo *input, const ITen ARM_COMPUTE_RETURN_ON_ERROR(CLWinogradInputTransform::validate(input, &input0, conv_info, Size2D(kernel_w, kernel_h))); // Validate filter transform - const TensorShape input1_shape = misc::shape_calculator::compute_winograd_filter_transform_shape(*weights); + const TensorShape input1_shape = misc::shape_calculator::compute_winograd_filter_transform_shape(*weights, Size2D(2U, 2U)); const TensorInfo input1 = weights->clone()->set_tensor_shape(input1_shape); - ARM_COMPUTE_RETURN_ON_ERROR(CLWinogradFilterTransformKernel::validate(weights, &input1)); + ARM_COMPUTE_RETURN_ON_ERROR(CLWinogradFilterTransformKernel::validate(weights, &input1, Size2D(2U, 2U))); // Configure batched matrix multiply TensorShape batched_mm_output_shape = input0.tensor_shape(); |