diff options
-rw-r--r-- | src/core/CL/cl_kernels/winograd_output_transform.cl | 6 | ||||
-rw-r--r-- | tests/validation/CL/Winograd.cpp | 24 | ||||
-rw-r--r-- | tests/validation/fixtures/WinogradConvolutionLayerFixture.h | 13 |
3 files changed, 24 insertions, 19 deletions
diff --git a/src/core/CL/cl_kernels/winograd_output_transform.cl b/src/core/CL/cl_kernels/winograd_output_transform.cl index bb159a7f4c..7065da1cc5 100644 --- a/src/core/CL/cl_kernels/winograd_output_transform.cl +++ b/src/core/CL/cl_kernels/winograd_output_transform.cl @@ -507,7 +507,7 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc( // Add bias Vector bias = CONVERT_TO_VECTOR_STRUCT_NO_STEP(bias); - float b = (float) * ((__global float *)(vector_offset(&bias, z_out))); + float b = (float) * ((__global float *)(vector_offset(&bias, x_out))); out00 += (float)b; out01 += (float)b; @@ -856,7 +856,7 @@ __kernel void winograd_output_transform_4x4_5x5_nhwc( // Add bias Vector bias = CONVERT_TO_VECTOR_STRUCT_NO_STEP(bias); - float b = (float) * ((__global float *)(vector_offset(&bias, z_out))); + float b = (float) * ((__global float *)(vector_offset(&bias, x_out))); out00 += (float)b; out01 += (float)b; @@ -981,7 +981,7 @@ __kernel void winograd_output_transform_4x4_5x5_nhwc( // Add bias Vector bias = CONVERT_TO_VECTOR_STRUCT_NO_STEP(bias); - float b = (float) * ((__global float *)(vector_offset(&bias, z_out))); + float b = (float) * ((__global float *)(vector_offset(&bias, x_out))); out_col0 += (float4)b; out_col1 += (float4)b; diff --git a/tests/validation/CL/Winograd.cpp b/tests/validation/CL/Winograd.cpp index 004d7c98a0..bb83f5ac3a 100644 --- a/tests/validation/CL/Winograd.cpp +++ b/tests/validation/CL/Winograd.cpp @@ -580,7 +580,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLWinogradConvolutionLayerFastMathFixture, fram combine(combine(combine(datasets::SmallWinogradConvolutionLayer3x3Dataset(), framework::dataset::make("DataType", { DataType::F32 })), framework::dataset::make("ActivationLayerInfo", { ActivationLayerInfo() })), - framework::dataset::make("DataLayout", { DataLayout::NCHW }))) + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_convolution_layer_f32); @@ -590,7 +590,7 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLWinogradConvolutionLayerFastMathFixture, fram combine(combine(combine(datasets::LargeWinogradConvolutionLayer3x3Dataset(), framework::dataset::make("DataType", { DataType::F32 })), framework::dataset::make("ActivationLayerInfo", { ActivationLayerInfo() })), - framework::dataset::make("DataLayout", { DataLayout::NCHW }))) + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_convolution_layer_f32); @@ -602,7 +602,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLWinogradConvolutionLayerFastMathFixture, fram combine(combine(combine(datasets::SmallWinogradConvolutionLayer3x1Dataset(), framework::dataset::make("DataType", { DataType::F32 })), framework::dataset::make("ActivationLayerInfo", { ActivationLayerInfo() })), - framework::dataset::make("DataLayout", { DataLayout::NCHW }))) + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_convolution_layer_f32); @@ -612,7 +612,7 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLWinogradConvolutionLayerFastMathFixture, fram combine(combine(combine(datasets::LargeWinogradConvolutionLayer3x1Dataset(), framework::dataset::make("DataType", { DataType::F32 })), framework::dataset::make("ActivationLayerInfo", { ActivationLayerInfo() })), - framework::dataset::make("DataLayout", { DataLayout::NCHW }))) + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_convolution_layer_f32); @@ -624,7 +624,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLWinogradConvolutionLayerFastMathFixture, fram combine(combine(combine(datasets::SmallWinogradConvolutionLayer1x3Dataset(), framework::dataset::make("DataType", { DataType::F32 })), framework::dataset::make("ActivationLayerInfo", { ActivationLayerInfo() })), - framework::dataset::make("DataLayout", { DataLayout::NCHW }))) + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_convolution_layer_f32); @@ -634,7 +634,7 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLWinogradConvolutionLayerFastMathFixture, fram combine(combine(combine(datasets::LargeWinogradConvolutionLayer1x3Dataset(), framework::dataset::make("DataType", { DataType::F32 })), framework::dataset::make("ActivationLayerInfo", { ActivationLayerInfo() })), - framework::dataset::make("DataLayout", { DataLayout::NCHW }))) + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_convolution_layer_f32); @@ -646,7 +646,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLWinogradConvolutionLayerFastMathFixture, fram combine(combine(combine(datasets::SmallWinogradConvolutionLayer5x5Dataset(), framework::dataset::make("DataType", { DataType::F32 })), framework::dataset::make("ActivationLayerInfo", { ActivationLayerInfo() })), - framework::dataset::make("DataLayout", { DataLayout::NCHW }))) + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC }))) { // Validate output @@ -657,7 +657,7 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLWinogradConvolutionLayerFastMathFixture, fram combine(combine(combine(datasets::LargeWinogradConvolutionLayer5x5Dataset(), framework::dataset::make("DataType", { DataType::F32 })), framework::dataset::make("ActivationLayerInfo", { ActivationLayerInfo() })), - framework::dataset::make("DataLayout", { DataLayout::NCHW }))) + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC }))) { // Validate output @@ -670,7 +670,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLWinogradConvolutionLayerFastMathFixture, fram combine(combine(combine(datasets::SmallWinogradConvolutionLayer5x1Dataset(), framework::dataset::make("DataType", { DataType::F32 })), framework::dataset::make("ActivationLayerInfo", { ActivationLayerInfo() })), - framework::dataset::make("DataLayout", { DataLayout::NCHW }))) + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC }))) { // Validate output @@ -681,7 +681,7 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLWinogradConvolutionLayerFastMathFixture, fram combine(combine(combine(datasets::LargeWinogradConvolutionLayer5x1Dataset(), framework::dataset::make("DataType", { DataType::F32 })), framework::dataset::make("ActivationLayerInfo", { ActivationLayerInfo() })), - framework::dataset::make("DataLayout", { DataLayout::NCHW }))) + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC }))) { // Validate output @@ -694,7 +694,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLWinogradConvolutionLayerFastMathFixture, fram combine(combine(combine(datasets::SmallWinogradConvolutionLayer1x5Dataset(), framework::dataset::make("DataType", { DataType::F32 })), framework::dataset::make("ActivationLayerInfo", { ActivationLayerInfo() })), - framework::dataset::make("DataLayout", { DataLayout::NCHW }))) + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC }))) { // Validate output @@ -705,7 +705,7 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLWinogradConvolutionLayerFastMathFixture, fram combine(combine(combine(datasets::LargeWinogradConvolutionLayer1x5Dataset(), framework::dataset::make("DataType", { DataType::F32 })), framework::dataset::make("ActivationLayerInfo", { ActivationLayerInfo() })), - framework::dataset::make("DataLayout", { DataLayout::NCHW }))) + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC }))) { // Validate output diff --git a/tests/validation/fixtures/WinogradConvolutionLayerFixture.h b/tests/validation/fixtures/WinogradConvolutionLayerFixture.h index f1660e6e90..aba3efffa8 100644 --- a/tests/validation/fixtures/WinogradConvolutionLayerFixture.h +++ b/tests/validation/fixtures/WinogradConvolutionLayerFixture.h @@ -501,25 +501,30 @@ protected: TensorShape output_shape = compute_winograd_output_transform_shape(TensorInfo(input_shape, 1, data_type), winograd_info); // Create tensors - TensorType src = create_tensor<TensorType>(input_shape, data_type); - TensorType dst = create_tensor<TensorType>(output_shape, data_type, 1, QuantizationInfo(), winograd_info.output_data_layout); + TensorType src = create_tensor<TensorType>(input_shape, data_type); + TensorType bias = create_tensor<TensorType>(output_shape[get_data_layout_dimension_index(winograd_info.output_data_layout, DataLayoutDimension::CHANNEL)], data_type); + TensorType dst = create_tensor<TensorType>(output_shape, data_type, 1, QuantizationInfo(), winograd_info.output_data_layout); // Create and configure function FunctionType output_transform; - output_transform.configure(&src, nullptr, &dst, winograd_info); + output_transform.configure(&src, &bias, &dst, winograd_info); ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(bias.info()->is_resizable(), framework::LogLevel::ERRORS); ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); // Allocate tensors src.allocator()->allocate(); + bias.allocator()->allocate(); dst.allocator()->allocate(); ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!bias.info()->is_resizable(), framework::LogLevel::ERRORS); ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS); // Fill tensors fill(AccessorType(src), 0, -1.f, 1.f); + fill(AccessorType(bias), 1, -1.f, 1.f); output_transform.run(); @@ -537,7 +542,7 @@ protected: // Fill reference fill(src, 0, -1.f, 1.f); - fill(bias, 1, 0.0f, 0.0f); // Fill with zeros as we validate just the output transform without bias contribution + fill(bias, 1, -1.f, 1.f); return reference::winograd_output_transform<T>(src, bias, output_shape, winograd_info); } |