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>
diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp
index 9df2dcb..740a98b 100644
--- a/src/core/CL/CLKernelLibrary.cpp
+++ b/src/core/CL/CLKernelLibrary.cpp
@@ -352,6 +352,7 @@
{ "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 25c129d..bd51db6 100644
--- a/src/core/CL/cl_kernels/winograd.cl
+++ b/src/core/CL/cl_kernels/winograd.cl
@@ -116,6 +116,144 @@
*(__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 655b82b..5a03332 100644
--- a/src/core/CL/kernels/CLWinogradFilterTransformKernel.cpp
+++ b/src/core/CL/kernels/CLWinogradFilterTransformKernel.cpp
@@ -44,17 +44,18 @@
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 @@
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 @@
{
}
-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 5081cba..a861e00 100644
--- a/src/runtime/CL/functions/CLWinogradConvolutionLayer.cpp
+++ b/src/runtime/CL/functions/CLWinogradConvolutionLayer.cpp
@@ -64,7 +64,7 @@
_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 @@
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();