From fe5ef38cdbc1e9a44c3786744dfc0cc915a608a6 Mon Sep 17 00:00:00 2001 From: Giorgio Arena Date: Tue, 17 Apr 2018 10:14:10 +0100 Subject: COMPMID-1037 Add support for F(4x4, 5x5) in CLWinogradInputTransformKernel Change-Id: Iac26936f46d0f7cdd9d2f8393b0092cd5a223c45 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/127675 Tested-by: Jenkins Reviewed-by: Gian Marco Iodice --- src/core/CL/CLKernelLibrary.cpp | 1 + src/core/CL/cl_kernels/winograd.cl | 169 +++++++++++++++++++++ .../CL/kernels/CLWinogradInputTransformKernel.cpp | 5 +- tests/datasets/WinogradInputTransformDataset.h | 11 ++ tests/validation/CL/Winograd.cpp | 4 +- tests/validation/reference/Winograd.cpp | 13 ++ 6 files changed, 199 insertions(+), 4 deletions(-) diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index a7b39c20e2..8405754f52 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -357,6 +357,7 @@ const std::map CLKernelLibrary::_kernel_program_map = { "winograd_filter_transform_2x2_3x3_nchw", "winograd.cl" }, { "winograd_filter_transform_4x4_3x3_nchw", "winograd.cl" }, { "winograd_filter_transform_4x4_5x5_nchw", "winograd.cl" }, + { "winograd_input_transform_4x4_5x5_stepz1_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 03ff377d52..9932119003 100644 --- a/src/core/CL/cl_kernels/winograd.cl +++ b/src/core/CL/cl_kernels/winograd.cl @@ -707,6 +707,175 @@ __kernel void winograd_input_transform_2x2_3x3_stepz2_nchw( vstore2(out32, 0, (__global float *)(dst_addr + 14 * dst_stride_z)); vstore2(out33, 0, (__global float *)(dst_addr + 15 * dst_stride_z)); } + +#define OUTPUT_ROW_4x4_5x5(out, tmp, comm_fact) \ + ({ \ + comm_fact.s0 = tmp.s2 - 4.25f * tmp.s4 + tmp.s6; \ + comm_fact.s1 = tmp.s1 - 4.25f * tmp.s3 + tmp.s5; \ + comm_fact.s2 = 2.5f * tmp.s3; \ + comm_fact.s3 = 0.5f * tmp.s1 + 2.f * tmp.s5 - comm_fact.s2; \ + comm_fact.s4 = 0.25f * tmp.s2 - 1.25f * tmp.s4 + tmp.s6; \ + comm_fact.s5 = 4.f * tmp.s2 + tmp.s6 - 5.f * tmp.s4; \ + comm_fact.s6 = 2.f * tmp.s1 + 0.5f * tmp.s5 - comm_fact.s2; \ + \ + out.s0 = tmp.s0 - tmp.s6 + 5.25f * tmp.s4 - 5.25f * tmp.s2; \ + out.s1 = comm_fact.s0 + comm_fact.s1; \ + out.s2 = comm_fact.s0 - comm_fact.s1; \ + out.s3 = comm_fact.s3 + comm_fact.s4; \ + out.s4 = comm_fact.s4 - comm_fact.s3; \ + out.s5 = comm_fact.s5 + comm_fact.s6; \ + out.s6 = comm_fact.s5 - comm_fact.s6; \ + out.s7 = tmp.s7 - tmp.s1 + 5.25f * tmp.s3 - 5.25f * tmp.s5; \ + }) + +/** This OpenCL kernel computes the input transform when the kernel size is 5x5 and the output tile is 4x4 + * + * @note The number of tiles in the x axis must be passed at compile time using -DNUM_TILES_X (i.e.-DNUM_TILES_X=5). + * @note The pad left and pad top must be passed at compile time using -DPAD_LEFT and -DPAD_TOP (i.e.-DPAD_LEFT=1 and -DPAD_TOP=0). + * + * @param[in] src_ptr Pointer to the source image. Supported data types: F32 + * @param[in] src_stride_x Stride of the source image 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 image 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_offset_first_element_in_bytes The offset of the first element in the source image + * @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 Y processed per workitem(in bytes) + * @param[in] dst_ptr Pointer to the destination tensor. Supported data types: 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] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Y 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_input_transform_4x4_5x5_stepz1_nchw( + TENSOR3D_DECLARATION(src), + TENSOR3D_DECLARATION(dst)) +{ + int x = get_global_id(0); + int y = get_global_id(1); + int z = get_global_id(2); + + // Compute input address + __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x * 4 * src_stride_x + y * 4 * src_stride_y + z * src_stride_z; + + src_addr = src_addr - ((int)PAD_LEFT * src_stride_x) - ((int)PAD_TOP * src_stride_y); + + // Load 8x8 input tile + const float8 in_row0 = vload8(0, (__global float *)(src_addr + 0 * src_stride_y)); + const float8 in_row1 = vload8(0, (__global float *)(src_addr + 1 * src_stride_y)); + const float8 in_row2 = vload8(0, (__global float *)(src_addr + 2 * src_stride_y)); + const float8 in_row3 = vload8(0, (__global float *)(src_addr + 3 * src_stride_y)); + const float8 in_row4 = vload8(0, (__global float *)(src_addr + 4 * src_stride_y)); + const float8 in_row5 = vload8(0, (__global float *)(src_addr + 5 * src_stride_y)); + const float8 in_row6 = vload8(0, (__global float *)(src_addr + 6 * src_stride_y)); + const float8 in_row7 = vload8(0, (__global float *)(src_addr + 7 * src_stride_y)); + + // Calculate common factors for intermediate tensor + float8 comm_fact0 = in_row2 + in_row6 - 4.25f * in_row4; + float8 comm_fact1 = in_row1 + in_row5 - 4.25f * in_row3; + float8 comm_fact2 = 0.25f * in_row2 - 1.25f * in_row4 + in_row6; + + // Calculate intermediate tensor and reuse common factor vectors + const float8 tmp0 = in_row0 - in_row6 + 5.25f * in_row4 - 5.25f * in_row2; + const float8 tmp1 = comm_fact0 + comm_fact1; + const float8 tmp2 = comm_fact0 - comm_fact1; + + comm_fact0 = 2.5f * in_row3; + comm_fact1 = 0.5f * in_row1 - comm_fact0 + 2.f * in_row5; + + const float8 tmp3 = comm_fact1 + comm_fact2; + const float8 tmp4 = comm_fact2 - comm_fact1; + + comm_fact1 = 2.f * in_row1 - comm_fact0 + 0.5f * in_row5; + comm_fact2 = 4.f * in_row2 - 5.f * in_row4 + in_row6; + + const float8 tmp5 = comm_fact1 + comm_fact2; + const float8 tmp6 = comm_fact2 - comm_fact1; + const float8 tmp7 = in_row7 - in_row1 + 5.25f * in_row3 - 5.25f * in_row5; + + // Calculate output rows (reuse comm_fact0 vector) + float8 out0, out1, out2, out3, out4, out5, out6, out7; + + OUTPUT_ROW_4x4_5x5(out0, tmp0, comm_fact0); + OUTPUT_ROW_4x4_5x5(out1, tmp1, comm_fact0); + OUTPUT_ROW_4x4_5x5(out2, tmp2, comm_fact0); + OUTPUT_ROW_4x4_5x5(out3, tmp3, comm_fact0); + OUTPUT_ROW_4x4_5x5(out4, tmp4, comm_fact0); + OUTPUT_ROW_4x4_5x5(out5, tmp5, comm_fact0); + OUTPUT_ROW_4x4_5x5(out6, tmp6, comm_fact0); + OUTPUT_ROW_4x4_5x5(out7, tmp7, comm_fact0); + + // Store values across the 64 channels + __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + z * dst_stride_x + (x + y * (int)NUM_TILES_X) * dst_stride_y; + + *((__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)) = out0.s6; + *((__global float *)(dst_addr + 7 * dst_stride_z)) = out0.s7; + *((__global float *)(dst_addr + 8 * dst_stride_z)) = out1.s0; + *((__global float *)(dst_addr + 9 * dst_stride_z)) = out1.s1; + *((__global float *)(dst_addr + 10 * dst_stride_z)) = out1.s2; + *((__global float *)(dst_addr + 11 * dst_stride_z)) = out1.s3; + *((__global float *)(dst_addr + 12 * dst_stride_z)) = out1.s4; + *((__global float *)(dst_addr + 13 * dst_stride_z)) = out1.s5; + *((__global float *)(dst_addr + 14 * dst_stride_z)) = out1.s6; + *((__global float *)(dst_addr + 15 * dst_stride_z)) = out1.s7; + *((__global float *)(dst_addr + 16 * dst_stride_z)) = out2.s0; + *((__global float *)(dst_addr + 17 * dst_stride_z)) = out2.s1; + *((__global float *)(dst_addr + 18 * dst_stride_z)) = out2.s2; + *((__global float *)(dst_addr + 19 * dst_stride_z)) = out2.s3; + *((__global float *)(dst_addr + 20 * dst_stride_z)) = out2.s4; + *((__global float *)(dst_addr + 21 * dst_stride_z)) = out2.s5; + *((__global float *)(dst_addr + 22 * dst_stride_z)) = out2.s6; + *((__global float *)(dst_addr + 23 * dst_stride_z)) = out2.s7; + *((__global float *)(dst_addr + 24 * dst_stride_z)) = out3.s0; + *((__global float *)(dst_addr + 25 * dst_stride_z)) = out3.s1; + *((__global float *)(dst_addr + 26 * dst_stride_z)) = out3.s2; + *((__global float *)(dst_addr + 27 * dst_stride_z)) = out3.s3; + *((__global float *)(dst_addr + 28 * dst_stride_z)) = out3.s4; + *((__global float *)(dst_addr + 29 * dst_stride_z)) = out3.s5; + *((__global float *)(dst_addr + 30 * dst_stride_z)) = out3.s6; + *((__global float *)(dst_addr + 31 * dst_stride_z)) = out3.s7; + *((__global float *)(dst_addr + 32 * dst_stride_z)) = out4.s0; + *((__global float *)(dst_addr + 33 * dst_stride_z)) = out4.s1; + *((__global float *)(dst_addr + 34 * dst_stride_z)) = out4.s2; + *((__global float *)(dst_addr + 35 * dst_stride_z)) = out4.s3; + *((__global float *)(dst_addr + 36 * dst_stride_z)) = out4.s4; + *((__global float *)(dst_addr + 37 * dst_stride_z)) = out4.s5; + *((__global float *)(dst_addr + 38 * dst_stride_z)) = out4.s6; + *((__global float *)(dst_addr + 39 * dst_stride_z)) = out4.s7; + *((__global float *)(dst_addr + 40 * dst_stride_z)) = out5.s0; + *((__global float *)(dst_addr + 41 * dst_stride_z)) = out5.s1; + *((__global float *)(dst_addr + 42 * dst_stride_z)) = out5.s2; + *((__global float *)(dst_addr + 43 * dst_stride_z)) = out5.s3; + *((__global float *)(dst_addr + 44 * dst_stride_z)) = out5.s4; + *((__global float *)(dst_addr + 45 * dst_stride_z)) = out5.s5; + *((__global float *)(dst_addr + 46 * dst_stride_z)) = out5.s6; + *((__global float *)(dst_addr + 47 * dst_stride_z)) = out5.s7; + *((__global float *)(dst_addr + 48 * dst_stride_z)) = out6.s0; + *((__global float *)(dst_addr + 49 * dst_stride_z)) = out6.s1; + *((__global float *)(dst_addr + 50 * dst_stride_z)) = out6.s2; + *((__global float *)(dst_addr + 51 * dst_stride_z)) = out6.s3; + *((__global float *)(dst_addr + 52 * dst_stride_z)) = out6.s4; + *((__global float *)(dst_addr + 53 * dst_stride_z)) = out6.s5; + *((__global float *)(dst_addr + 54 * dst_stride_z)) = out6.s6; + *((__global float *)(dst_addr + 55 * dst_stride_z)) = out6.s7; + *((__global float *)(dst_addr + 56 * dst_stride_z)) = out7.s0; + *((__global float *)(dst_addr + 57 * dst_stride_z)) = out7.s1; + *((__global float *)(dst_addr + 58 * dst_stride_z)) = out7.s2; + *((__global float *)(dst_addr + 59 * dst_stride_z)) = out7.s3; + *((__global float *)(dst_addr + 60 * dst_stride_z)) = out7.s4; + *((__global float *)(dst_addr + 61 * dst_stride_z)) = out7.s5; + *((__global float *)(dst_addr + 62 * dst_stride_z)) = out7.s6; + *((__global float *)(dst_addr + 63 * dst_stride_z)) = out7.s7; +} #endif // defined(NUM_TILES_X) && defined(PAD_LEFT) && defined(PAD_TOP) #if defined(NUM_TILES_X) diff --git a/src/core/CL/kernels/CLWinogradInputTransformKernel.cpp b/src/core/CL/kernels/CLWinogradInputTransformKernel.cpp index df7ffe83a0..a47590d20f 100644 --- a/src/core/CL/kernels/CLWinogradInputTransformKernel.cpp +++ b/src/core/CL/kernels/CLWinogradInputTransformKernel.cpp @@ -46,8 +46,9 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, c const Size2D output_tile_size = winograd_info.output_tile_size; const Size2D kernel_size = winograd_info.kernel_size; ARM_COMPUTE_RETURN_ERROR_ON_MSG(conv_info.stride().first != 1 || conv_info.stride().second != 1, "Winograd input transform only supports unit strides"); - ARM_COMPUTE_RETURN_ERROR_ON_MSG(kernel_size != Size2D(3U, 3U), "Winograd input transform only supports 3x3 kernels"); - ARM_COMPUTE_RETURN_ERROR_ON_MSG(output_tile_size != Size2D(2U, 2U), "Winograd input transform only supports 2x2 output tile size"); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(kernel_size != Size2D(3U, 3U) && kernel_size != Size2D(5U, 5U), "Winograd input transform only supports 3x3 and 5x5 kernels"); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(kernel_size == Size2D(3U, 3U) && output_tile_size != Size2D(2U, 2U), "Winograd input transform only supports 2x2 output tile for 3x3 kernels"); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(kernel_size == Size2D(5U, 5U) && output_tile_size != Size2D(4U, 4U), "Winograd input transform only supports 4x4 output tile for 5x5 kernels"); ARM_COMPUTE_UNUSED(conv_info); ARM_COMPUTE_UNUSED(output_tile_size); ARM_COMPUTE_UNUSED(kernel_size); diff --git a/tests/datasets/WinogradInputTransformDataset.h b/tests/datasets/WinogradInputTransformDataset.h index 625daa0e6e..cbe63645de 100644 --- a/tests/datasets/WinogradInputTransformDataset.h +++ b/tests/datasets/WinogradInputTransformDataset.h @@ -109,6 +109,13 @@ public: add_config(TensorShape(27U, 13U, 2U, 4U), WinogradInfo(Size2D(2U, 2U), Size2D(3U, 3U), Size2D(27U, 13U), PadStrideInfo(1, 1, 1, 1), DataLayout::NCHW)); add_config(TensorShape(9U, 9U, 3U, 5U), WinogradInfo(Size2D(2U, 2U), Size2D(3U, 3U), Size2D(9U, 9U), PadStrideInfo(1, 1, 0, 0), DataLayout::NCHW)); add_config(TensorShape(14U, 14U, 512U, 2U), WinogradInfo(Size2D(2U, 2U), Size2D(3U, 3U), Size2D(14U, 14U), PadStrideInfo(1, 1, 1, 1), DataLayout::NCHW)); + add_config(TensorShape(9U, 9U), WinogradInfo(Size2D(4U, 4U), Size2D(5U, 5U), Size2D(9U, 9U), PadStrideInfo(1, 1, 1, 1), DataLayout::NCHW)); + add_config(TensorShape(27U, 13U, 2U), WinogradInfo(Size2D(4U, 4U), Size2D(5U, 5U), Size2D(27U, 13U), PadStrideInfo(1, 1, 0, 0), DataLayout::NCHW)); + add_config(TensorShape(128U, 64U, 1U, 3U), WinogradInfo(Size2D(4U, 4U), Size2D(5U, 5U), Size2D(128U, 64U), PadStrideInfo(1, 1, 1, 1), DataLayout::NCHW)); + add_config(TensorShape(9U, 9U, 3U, 4U), WinogradInfo(Size2D(4U, 4U), Size2D(5U, 5U), Size2D(9U, 9U), PadStrideInfo(1, 1, 0, 0), DataLayout::NCHW)); + add_config(TensorShape(27U, 13U, 2U, 4U), WinogradInfo(Size2D(4U, 4U), Size2D(5U, 5U), Size2D(27U, 13U), PadStrideInfo(1, 1, 1, 1), DataLayout::NCHW)); + add_config(TensorShape(9U, 9U, 3U, 5U), WinogradInfo(Size2D(4U, 4U), Size2D(5U, 5U), Size2D(9U, 9U), PadStrideInfo(1, 1, 0, 0), DataLayout::NCHW)); + add_config(TensorShape(14U, 14U, 512U, 2U), WinogradInfo(Size2D(4U, 4U), Size2D(5U, 5U), Size2D(14U, 14U), PadStrideInfo(1, 1, 1, 1), DataLayout::NCHW)); } }; @@ -121,6 +128,10 @@ public: add_config(TensorShape(57U, 60U, 13U, 8U), WinogradInfo(Size2D(2U, 2U), Size2D(3U, 3U), Size2D(57U, 60U), PadStrideInfo(1, 1, 1, 1), DataLayout::NCHW)); add_config(TensorShape(128U, 64U, 21U, 13U), WinogradInfo(Size2D(2U, 2U), Size2D(3U, 3U), Size2D(128U, 64U), PadStrideInfo(1, 1, 0, 0), DataLayout::NCHW)); add_config(TensorShape(83U, 72U, 14U, 5U), WinogradInfo(Size2D(2U, 2U), Size2D(3U, 3U), Size2D(83U, 72U), PadStrideInfo(1, 1, 0, 0), DataLayout::NCHW)); + add_config(TensorShape(42U, 37U, 8U, 15U), WinogradInfo(Size2D(4U, 4U), Size2D(5U, 5U), Size2D(42U, 37U), PadStrideInfo(1, 1, 1, 1), DataLayout::NCHW)); + add_config(TensorShape(57U, 60U, 13U, 8U), WinogradInfo(Size2D(4U, 4U), Size2D(5U, 5U), Size2D(57U, 60U), PadStrideInfo(1, 1, 1, 1), DataLayout::NCHW)); + add_config(TensorShape(128U, 64U, 21U, 13U), WinogradInfo(Size2D(4U, 4U), Size2D(5U, 5U), Size2D(128U, 64U), PadStrideInfo(1, 1, 0, 0), DataLayout::NCHW)); + add_config(TensorShape(83U, 72U, 14U, 5U), WinogradInfo(Size2D(4U, 4U), Size2D(5U, 5U), Size2D(83U, 72U), PadStrideInfo(1, 1, 0, 0), DataLayout::NCHW)); } }; } // namespace datasets diff --git a/tests/validation/CL/Winograd.cpp b/tests/validation/CL/Winograd.cpp index 08810eaf58..4bfe4c8f42 100644 --- a/tests/validation/CL/Winograd.cpp +++ b/tests/validation/CL/Winograd.cpp @@ -126,14 +126,14 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLWinogradInputTransformFixture, framework::Dat framework::dataset::make("DataLayout", { DataLayout::NCHW })), framework::dataset::make("DataType", { DataType::F32 }))) { - validate(CLAccessor(_target), _reference); + validate(CLAccessor(_target), _reference, tolerance_f32); } FIXTURE_DATA_TEST_CASE(RunLarge, CLWinogradInputTransformFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeWinogradInputTransformDataset(), framework::dataset::make("DataLayout", { DataLayout::NCHW })), framework::dataset::make("DataType", { DataType::F32 }))) { - validate(CLAccessor(_target), _reference); + validate(CLAccessor(_target), _reference, tolerance_f32); } TEST_SUITE_END() // InputTransform diff --git a/tests/validation/reference/Winograd.cpp b/tests/validation/reference/Winograd.cpp index 757a06d847..fa6e372cde 100644 --- a/tests/validation/reference/Winograd.cpp +++ b/tests/validation/reference/Winograd.cpp @@ -62,6 +62,18 @@ void initialize_matrix_transform(SimpleTensor &src, const Size2D &output_tile 0.0f, 4.0f, 0.0f, -5.0f, 0.0f, 1.0f, }; + static const float imatrix4x4_5x5[] = + { + 1.f, 0.f, -21.f / 4.f, 0.f, 21.f / 4.f, 0.f, -1.f, 0.f, + 0.f, 1.f, 1.f, -17.f / 4.f, -17.f / 4.f, 1.f, 1.f, 0.f, + 0.f, -1.f, 1.f, 17.f / 4.f, -17.f / 4.f, -1.f, 1.f, 0.f, + 0.f, 1.f / 2.f, 1.f / 4.f, -5.f / 2.f, -5.f / 4.f, 2.f, 1.f, 0.f, + 0.f, -1.f / 2.f, 1.f / 4.f, 5.f / 2.f, -5.f / 4.f, -2.f, 1.f, 0.f, + 0.f, 2.f, 4.f, -5.f / 2.f, -5.f, 1.f / 2.f, 1.f, 0.f, + 0.f, -2.f, 4.f, 5.f / 2.f, -5.f, -1.f / 2.f, 1.f, 0.f, + 0.f, -1.f, 0.f, 21.f / 4.f, 0.f, -21.f / 4.f, 0.f, 1.f + }; + // ------------------------------------------ // Winograd filter transform matrices @@ -122,6 +134,7 @@ void initialize_matrix_transform(SimpleTensor &src, const Size2D &output_tile { { WinogradKey(std::pair(2, 2), std::pair(3, 3), WinogradTransformType::INPUT), imatrix2x2_3x3 }, { WinogradKey(std::pair(4, 4), std::pair(3, 3), WinogradTransformType::INPUT), imatrix4x4_3x3 }, + { WinogradKey(std::pair(4, 4), std::pair(5, 5), WinogradTransformType::INPUT), imatrix4x4_5x5 }, { WinogradKey(std::pair(2, 2), std::pair(3, 3), WinogradTransformType::FILTER), fmatrix2x2_3x3 }, { WinogradKey(std::pair(4, 4), std::pair(3, 3), WinogradTransformType::FILTER), fmatrix4x4_3x3 }, { WinogradKey(std::pair(4, 4), std::pair(5, 5), WinogradTransformType::FILTER), fmatrix4x4_5x5 }, -- cgit v1.2.1