aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--src/core/CL/cl_kernels/winograd_output_transform.cl6
-rw-r--r--tests/validation/CL/Winograd.cpp24
-rw-r--r--tests/validation/fixtures/WinogradConvolutionLayerFixture.h13
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);
}
generated by cgit v1.2.1 (git 2.23.0) at 2024-06-05 04:23:03 +0000