aboutsummaryrefslogtreecommitdiff
path: root/tests/validation/CL
diff options
context:
space:
mode:
authorMoritz Pflanzer <moritz.pflanzer@arm.com>2017-09-01 20:41:12 +0100
committerAnthony Barbier <anthony.barbier@arm.com>2018-11-02 16:35:24 +0000
commita09de0c8b2ed0f1481502d3b023375609362d9e3 (patch)
treee34b56d9ca69b025d7d9b943cc4df59cd458f6cb /tests/validation/CL
parent5280071b336d53aff94ca3a6c70ebbe6bf03f4c3 (diff)
downloadComputeLibrary-a09de0c8b2ed0f1481502d3b023375609362d9e3.tar.gz
COMPMID-415: Rename and move tests
The boost validation is now "standalone" in validation_old and builds as arm_compute_validation_old. The new validation builds now as arm_compute_validation. Change-Id: Ib93ba848a25680ac60afb92b461d574a0757150d Reviewed-on: http://mpd-gerrit.cambridge.arm.com/86187 Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Diffstat (limited to 'tests/validation/CL')
-rw-r--r--tests/validation/CL/ActivationLayer.cpp247
-rw-r--r--tests/validation/CL/ArithmeticAddition.cpp302
-rw-r--r--tests/validation/CL/ArithmeticSubtraction.cpp288
-rw-r--r--tests/validation/CL/BatchNormalizationLayer.cpp227
-rw-r--r--tests/validation/CL/BitwiseAnd.cpp94
-rw-r--r--tests/validation/CL/BitwiseNot.cpp90
-rw-r--r--tests/validation/CL/BitwiseOr.cpp94
-rw-r--r--tests/validation/CL/BitwiseXor.cpp94
-rw-r--r--tests/validation/CL/Box3x3.cpp166
-rw-r--r--tests/validation/CL/CLFixture.cpp32
-rw-r--r--tests/validation/CL/CMakeLists.txt68
-rw-r--r--tests/validation/CL/ConvolutionLayer.cpp186
-rw-r--r--tests/validation/CL/DepthConcatenateLayer.cpp123
-rw-r--r--tests/validation/CL/DepthConvert.cpp495
-rw-r--r--tests/validation/CL/DepthwiseConvolution.cpp69
-rw-r--r--tests/validation/CL/DepthwiseSeparableConvolutionLayer.cpp64
-rw-r--r--tests/validation/CL/DirectConvolutionLayer.cpp129
-rw-r--r--tests/validation/CL/FillBorder.cpp89
-rw-r--r--tests/validation/CL/FixedPoint/FixedPoint_QS8.cpp224
-rw-r--r--tests/validation/CL/Floor.cpp (renamed from tests/validation/CL/CLFixture.h)43
-rw-r--r--tests/validation/CL/FullyConnectedLayer.cpp195
-rw-r--r--tests/validation/CL/GEMM.cpp167
-rw-r--r--tests/validation/CL/Gaussian3x3.cpp166
-rw-r--r--tests/validation/CL/Gaussian5x5.cpp166
-rw-r--r--tests/validation/CL/HarrisCorners.cpp224
-rw-r--r--tests/validation/CL/IntegralImage.cpp142
-rw-r--r--tests/validation/CL/L2Normalize.cpp78
-rw-r--r--tests/validation/CL/MeanStdDev.cpp97
-rw-r--r--tests/validation/CL/MinMaxLocation.cpp397
-rw-r--r--tests/validation/CL/NonLinearFilter.cpp203
-rw-r--r--tests/validation/CL/NormalizationLayer.cpp141
-rw-r--r--tests/validation/CL/PixelWiseMultiplication.cpp164
-rw-r--r--tests/validation/CL/PoolingLayer.cpp142
-rw-r--r--tests/validation/CL/ROIPoolingLayer.cpp112
-rw-r--r--tests/validation/CL/ReductionOperation.cpp78
-rw-r--r--tests/validation/CL/Scale.cpp129
-rw-r--r--tests/validation/CL/Sobel3x3.cpp205
-rw-r--r--tests/validation/CL/Sobel5x5.cpp204
-rw-r--r--tests/validation/CL/SoftmaxLayer.cpp168
-rw-r--r--tests/validation/CL/TableLookup.cpp229
-rw-r--r--tests/validation/CL/Threshold.cpp153
-rw-r--r--tests/validation/CL/WarpPerspective.cpp208
42 files changed, 2417 insertions, 4475 deletions
diff --git a/tests/validation/CL/ActivationLayer.cpp b/tests/validation/CL/ActivationLayer.cpp
new file mode 100644
index 0000000000..097fb638f9
--- /dev/null
+++ b/tests/validation/CL/ActivationLayer.cpp
@@ -0,0 +1,247 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLActivationLayer.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/ActivationFunctionsDataset.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/ActivationLayerFixture.h"
+#include "tests/validation/half.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Define tolerance of the activation layer.
+ *
+ * @param[in] activation The activation function used.
+ * @param[in] data_type Data type.
+ *
+ * @return Tolerance depending on the activation function.
+ */
+AbsoluteTolerance<float> tolerance(ActivationLayerInfo::ActivationFunction activation, DataType data_type)
+{
+ constexpr float epsilon = 1e-6f;
+
+ switch(activation)
+ {
+ case ActivationLayerInfo::ActivationFunction::LINEAR:
+ return AbsoluteTolerance<float>(data_type == DataType::F16 ? 0.2f : epsilon);
+ case ActivationLayerInfo::ActivationFunction::SQUARE:
+ return AbsoluteTolerance<float>(data_type == DataType::F16 ? 0.1f : epsilon);
+ case ActivationLayerInfo::ActivationFunction::LOGISTIC:
+ if(is_data_type_fixed_point(data_type))
+ {
+ return AbsoluteTolerance<float>(5.f);
+ }
+ else
+ {
+ return AbsoluteTolerance<float>(data_type == DataType::F16 ? 0.001f : epsilon);
+ }
+ case ActivationLayerInfo::ActivationFunction::LEAKY_RELU:
+ return AbsoluteTolerance<float>(data_type == DataType::F16 ? 0.00001f : epsilon);
+ case ActivationLayerInfo::ActivationFunction::SOFT_RELU:
+ case ActivationLayerInfo::ActivationFunction::SQRT:
+ if(is_data_type_fixed_point(data_type))
+ {
+ return AbsoluteTolerance<float>(5.f);
+ }
+ else
+ {
+ return AbsoluteTolerance<float>(data_type == DataType::F16 ? 0.01f : 0.00001f);
+ }
+ case ActivationLayerInfo::ActivationFunction::TANH:
+ if(is_data_type_fixed_point(data_type))
+ {
+ return AbsoluteTolerance<float>(5.f);
+ }
+ else
+ {
+ return AbsoluteTolerance<float>(data_type == DataType::F16 ? 0.001f : 0.00001f);
+ }
+ default:
+ return AbsoluteTolerance<float>(epsilon);
+ }
+}
+
+/** CNN data types */
+const auto CNNDataTypes = framework::dataset::make("DataType",
+{
+ DataType::F16,
+ DataType::F32,
+ DataType::QS8,
+ DataType::QS16,
+});
+
+/** Input data sets. */
+const auto ActivationDataset = combine(combine(framework::dataset::make("InPlace", { false, true }), datasets::ActivationFunctions()), framework::dataset::make("AlphaBeta", { 0.5f, 1.f }));
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(ActivationLayer)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(combine(concat(datasets::SmallShapes(), datasets::LargeShapes()), CNNDataTypes), framework::dataset::make("InPlace", { false, true })),
+ shape, data_type, in_place)
+{
+ // Set fixed point position data type allowed
+ const int fixed_point_position = is_data_type_fixed_point(data_type) ? 3 : 0;
+
+ // Create tensors
+ CLTensor src = create_tensor<CLTensor>(shape, data_type, 1, fixed_point_position);
+ CLTensor dst = create_tensor<CLTensor>(shape, data_type, 1, fixed_point_position);
+
+ ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Create and configure function
+ CLActivationLayer act_layer;
+
+ if(in_place)
+ {
+ act_layer.configure(&src, nullptr, ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::ABS));
+ }
+ else
+ {
+ act_layer.configure(&src, &dst, ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::ABS));
+ }
+
+ // Validate valid region
+ const ValidRegion valid_region = shape_to_valid_region(shape);
+ validate(src.info()->valid_region(), valid_region);
+
+ if(!in_place)
+ {
+ validate(dst.info()->valid_region(), valid_region);
+ }
+
+ // Validate padding
+ const int step = 16 / arm_compute::data_size_from_type(data_type);
+ const PaddingSize padding = PaddingCalculator(shape.x(), step).required_padding();
+ validate(src.info()->padding(), padding);
+
+ if(!in_place)
+ {
+ validate(dst.info()->padding(), padding);
+ }
+}
+
+template <typename T>
+using CLActivationLayerFixture = ActivationValidationFixture<CLTensor, CLAccessor, CLActivationLayer, T>;
+
+TEST_SUITE(Float)
+TEST_SUITE(FP16)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLActivationLayerFixture<half_float::half>, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), ActivationDataset),
+ framework::dataset::make("DataType",
+ DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance(_function, _data_type));
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLActivationLayerFixture<half_float::half>, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), ActivationDataset),
+ framework::dataset::make("DataType",
+ DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance(_function, _data_type));
+}
+TEST_SUITE_END()
+
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLActivationLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), ActivationDataset), framework::dataset::make("DataType",
+ DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance(_function, _data_type));
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLActivationLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), ActivationDataset), framework::dataset::make("DataType",
+ DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance(_function, _data_type));
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+template <typename T>
+using CLActivationLayerFixedPointFixture = ActivationValidationFixedPointFixture<CLTensor, CLAccessor, CLActivationLayer, T>;
+
+TEST_SUITE(Quantized)
+TEST_SUITE(QS8)
+// We test for fixed point precision [3,5] because [1,2] and [6,7] ranges cause
+// overflowing issues in most of the transcendentals functions.
+FIXTURE_DATA_TEST_CASE(RunSmall, CLActivationLayerFixedPointFixture<int8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), ActivationDataset),
+ framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FractionalBits", 3, 6)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance(_function, _data_type));
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLActivationLayerFixedPointFixture<int8_t>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), ActivationDataset),
+ framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FractionalBits", 3, 6)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance(_function, _data_type));
+}
+TEST_SUITE_END()
+
+TEST_SUITE(QS16)
+// Testing for fixed point position [1,14) as reciprocal limits the maximum fixed point position to 14
+FIXTURE_DATA_TEST_CASE(RunSmall, CLActivationLayerFixedPointFixture<int16_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), ActivationDataset),
+ framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FractionalBits", 1, 14)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance(_function, _data_type));
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLActivationLayerFixedPointFixture<int16_t>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), ActivationDataset),
+ framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FractionalBits", 1, 14)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance(_function, _data_type));
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/ArithmeticAddition.cpp b/tests/validation/CL/ArithmeticAddition.cpp
deleted file mode 100644
index fc1bf5905d..0000000000
--- a/tests/validation/CL/ArithmeticAddition.cpp
+++ /dev/null
@@ -1,302 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "PaddingCalculator.h"
-#include "TypePrinter.h"
-#include "tests/Globals.h"
-#include "tests/Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/functions/CLArithmeticAddition.h"
-#include "arm_compute/runtime/Tensor.h"
-#include "arm_compute/runtime/TensorAllocator.h"
-
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-/** Compute Neon arithmetic addition function.
- *
- * @param[in] shape Shape of the input and output tensors.
- * @param[in] dt_in0 Data type of first input tensor.
- * @param[in] dt_in1 Data type of second input tensor.
- * @param[in] dt_out Data type of the output tensor.
- * @param[in] policy Overflow policy of the operation.
- * @param[in] fixed_point_position (Optional) Fixed point position that expresses the number of bits for the fractional part of the number when the tensor's data type is QS8 or QS16 (default = 0).
- *
- * @return Computed output tensor.
- */
-CLTensor compute_arithmetic_addition(const TensorShape &shape, DataType dt_in0, DataType dt_in1, DataType dt_out, ConvertPolicy policy, int fixed_point_position = 0)
-{
- // Create tensors
- CLTensor src1 = create_tensor<CLTensor>(shape, dt_in0, 1, fixed_point_position);
- CLTensor src2 = create_tensor<CLTensor>(shape, dt_in1, 1, fixed_point_position);
- CLTensor dst = create_tensor<CLTensor>(shape, dt_out, 1, fixed_point_position);
-
- // Create and configure function
- CLArithmeticAddition add;
- add.configure(&src1, &src2, &dst, policy);
-
- // Allocate tensors
- src1.allocator()->allocate();
- src2.allocator()->allocate();
- dst.allocator()->allocate();
-
- BOOST_TEST(!src1.info()->is_resizable());
- BOOST_TEST(!src2.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src1), 0);
- library->fill_tensor_uniform(CLAccessor(src2), 1);
-
- // Compute function
- add.run();
-
- return dst;
-}
-
-void validate_configuration(const CLTensor &src1, const CLTensor &src2, CLTensor &dst, TensorShape shape, ConvertPolicy policy)
-{
- BOOST_TEST(src1.info()->is_resizable());
- BOOST_TEST(src2.info()->is_resizable());
- BOOST_TEST(dst.info()->is_resizable());
-
- // Create and configure function
- CLArithmeticAddition add;
- add.configure(&src1, &src2, &dst, policy);
-
- // Validate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape);
- validate(src1.info()->valid_region(), valid_region);
- validate(src2.info()->valid_region(), valid_region);
- validate(dst.info()->valid_region(), valid_region);
-
- // Validate padding
- const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
- validate(src1.info()->padding(), padding);
- validate(src2.info()->padding(), padding);
- validate(dst.info()->padding(), padding);
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(ArithmeticAddition)
-
-BOOST_AUTO_TEST_SUITE(U8)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
- shape, policy)
-{
- // Create tensors
- CLTensor src1 = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor src2 = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U8);
-
- validate_configuration(src1, src2, dst, shape, policy);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
- shape, policy)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_addition(shape, DataType::U8, DataType::U8, DataType::U8, policy);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::U8, DataType::U8, DataType::U8, policy);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(S16)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
- shape, dt, policy)
-{
- // Create tensors
- CLTensor src1 = create_tensor<CLTensor>(shape, dt);
- CLTensor src2 = create_tensor<CLTensor>(shape, DataType::S16);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::S16);
-
- validate_configuration(src1, src2, dst, shape, policy);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
- shape, dt, policy)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_addition(shape, dt, DataType::S16, DataType::S16, policy);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, dt, DataType::S16, DataType::S16, policy);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
- shape, dt, policy)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_addition(shape, dt, DataType::S16, DataType::S16, policy);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, dt, DataType::S16, DataType::S16, policy);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(Quantized)
-BOOST_AUTO_TEST_SUITE(QS8)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 7),
- shape, policy, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_addition(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 7),
- shape, policy, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_addition(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(QS16)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 15),
- shape, policy, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_addition(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 15),
- shape, policy, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_addition(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(F16)
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), shape)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_addition(shape, DataType::F16, DataType::F16, DataType::F16, ConvertPolicy::WRAP);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::F16, DataType::F16, DataType::F16, ConvertPolicy::WRAP);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(F32)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
- shape, policy)
-{
- // Create tensors
- CLTensor src1 = create_tensor<CLTensor>(shape, DataType::F32);
- CLTensor src2 = create_tensor<CLTensor>(shape, DataType::F32);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::F32);
-
- validate_configuration(src1, src2, dst, shape, policy);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), shape)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_addition(shape, DataType::F32, DataType::F32, DataType::F32, ConvertPolicy::WRAP);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::F32, DataType::F32, DataType::F32, ConvertPolicy::WRAP);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
- shape, policy)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_addition(shape, DataType::F32, DataType::F32, DataType::F32, policy);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::F32, DataType::F32, DataType::F32, policy);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/ArithmeticSubtraction.cpp b/tests/validation/CL/ArithmeticSubtraction.cpp
deleted file mode 100644
index 086281cdda..0000000000
--- a/tests/validation/CL/ArithmeticSubtraction.cpp
+++ /dev/null
@@ -1,288 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "PaddingCalculator.h"
-#include "TypePrinter.h"
-#include "tests/Globals.h"
-#include "tests/Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/functions/CLArithmeticSubtraction.h"
-#include "arm_compute/runtime/Tensor.h"
-#include "arm_compute/runtime/TensorAllocator.h"
-
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-/** Compute Neon arithmetic subtraction function.
- *
- * @param[in] shape Shape of the input and output tensors.
- * @param[in] dt_in0 Data type of first input tensor.
- * @param[in] dt_in1 Data type of second input tensor.
- * @param[in] dt_out Data type of the output tensor.
- * @param[in] policy Overflow policy of the operation.
- * @param[in] fixed_point_position (Optional) Fixed point position that expresses the number of bits for the fractional part of the number when the tensor's data type is QS8 or QS16 (default = 0).
- *
- * @return Computed output tensor.
- */
-CLTensor compute_arithmetic_subtraction(const TensorShape &shape, DataType dt_in0, DataType dt_in1, DataType dt_out, ConvertPolicy policy, int fixed_point_position = 0)
-{
- // Create tensors
- CLTensor src1 = create_tensor<CLTensor>(shape, dt_in0, 1, fixed_point_position);
- CLTensor src2 = create_tensor<CLTensor>(shape, dt_in1, 1, fixed_point_position);
- CLTensor dst = create_tensor<CLTensor>(shape, dt_out, 1, fixed_point_position);
-
- // Create and configure function
- CLArithmeticSubtraction sub;
- sub.configure(&src1, &src2, &dst, policy);
-
- // Allocate tensors
- src1.allocator()->allocate();
- src2.allocator()->allocate();
- dst.allocator()->allocate();
-
- BOOST_TEST(!src1.info()->is_resizable());
- BOOST_TEST(!src2.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src1), 0);
- library->fill_tensor_uniform(CLAccessor(src2), 1);
-
- // Compute function
- sub.run();
-
- return dst;
-}
-
-void validate_configuration(const CLTensor &src1, const CLTensor &src2, CLTensor &dst, TensorShape shape, ConvertPolicy policy)
-{
- BOOST_TEST(src1.info()->is_resizable());
- BOOST_TEST(src2.info()->is_resizable());
- BOOST_TEST(dst.info()->is_resizable());
-
- // Create and configure function
- CLArithmeticSubtraction sub;
- sub.configure(&src1, &src2, &dst, policy);
-
- // Validate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape);
- validate(src1.info()->valid_region(), valid_region);
- validate(src2.info()->valid_region(), valid_region);
- validate(dst.info()->valid_region(), valid_region);
-
- // Validate padding
- const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
- validate(src1.info()->padding(), padding);
- validate(src2.info()->padding(), padding);
- validate(dst.info()->padding(), padding);
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(ArithmeticSubtraction)
-
-BOOST_AUTO_TEST_SUITE(U8)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
- shape, policy)
-{
- // Create tensors
- CLTensor src1 = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor src2 = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U8);
-
- validate_configuration(src1, src2, dst, shape, policy);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
- shape, policy)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_subtraction(shape, DataType::U8, DataType::U8, DataType::U8, policy);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::U8, DataType::U8, DataType::U8, policy);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(S16)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
- shape, dt, policy)
-{
- // Create tensors
- CLTensor src1 = create_tensor<CLTensor>(shape, dt);
- CLTensor src2 = create_tensor<CLTensor>(shape, DataType::S16);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::S16);
-
- validate_configuration(src1, src2, dst, shape, policy);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
- shape, dt, policy)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_subtraction(shape, dt, DataType::S16, DataType::S16, policy);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, dt, DataType::S16, DataType::S16, policy);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
- shape, dt, policy)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_subtraction(shape, dt, DataType::S16, DataType::S16, policy);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, dt, DataType::S16, DataType::S16, policy);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(Quantized)
-BOOST_AUTO_TEST_SUITE(QS8)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 7),
- shape, policy, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_subtraction(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 7),
- shape, policy, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_subtraction(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(QS16)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 15),
- shape, policy, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_subtraction(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 15),
- shape, policy, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_subtraction(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(Float)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
- shape, policy)
-{
- // Create tensors
- CLTensor src1 = create_tensor<CLTensor>(shape, DataType::F32);
- CLTensor src2 = create_tensor<CLTensor>(shape, DataType::F32);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::F32);
-
- validate_configuration(src1, src2, dst, shape, policy);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), shape)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_subtraction(shape, DataType::F32, DataType::F32, DataType::F32, ConvertPolicy::WRAP);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::F32, DataType::F32, DataType::F32, ConvertPolicy::WRAP);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
- shape, policy)
-{
- // Compute function
- CLTensor dst = compute_arithmetic_subtraction(shape, DataType::F32, DataType::F32, DataType::F32, policy);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::F32, DataType::F32, DataType::F32, policy);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/BatchNormalizationLayer.cpp b/tests/validation/CL/BatchNormalizationLayer.cpp
deleted file mode 100644
index abcc619cb8..0000000000
--- a/tests/validation/CL/BatchNormalizationLayer.cpp
+++ /dev/null
@@ -1,227 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "dataset/BatchNormalizationLayerDataset.h"
-#include "tests/validation/Helpers.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/functions/CLBatchNormalizationLayer.h"
-#include "arm_compute/runtime/Tensor.h"
-#include "arm_compute/runtime/TensorAllocator.h"
-
-#include <random>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-const float tolerance_f = 1e-05; /**< Tolerance value for comparing reference's output against floating point implementation's output */
-const float tolerance_qs8 = 3; /**< Tolerance value for comparing reference's output against quantized implementation's output */
-const float tolerance_qs16 = 6; /**< Tolerance value for comparing reference's output against quantized implementation's output */
-
-/** Compute Neon batch normalization function.
- *
- * @param[in] shape Shape of the input and output tensors.
- * @param[in] dt Data type of input and output tensors.
- * @param[in] norm_info Normalization Layer information.
- *
- * @return Computed output tensor.
- */
-CLTensor compute_reference_batch_normalization_layer(const TensorShape &shape0, const TensorShape &shape1, DataType dt, float epsilon, int fixed_point_position = 0)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape0, dt, 1, fixed_point_position);
- CLTensor dst = create_tensor<CLTensor>(shape0, dt, 1, fixed_point_position);
- CLTensor mean = create_tensor<CLTensor>(shape1, dt, 1, fixed_point_position);
- CLTensor var = create_tensor<CLTensor>(shape1, dt, 1, fixed_point_position);
- CLTensor beta = create_tensor<CLTensor>(shape1, dt, 1, fixed_point_position);
- CLTensor gamma = create_tensor<CLTensor>(shape1, dt, 1, fixed_point_position);
-
- // Create and configure function
- CLBatchNormalizationLayer norm;
- norm.configure(&src, &dst, &mean, &var, &beta, &gamma, epsilon);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst.allocator()->allocate();
- mean.allocator()->allocate();
- var.allocator()->allocate();
- beta.allocator()->allocate();
- gamma.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
- BOOST_TEST(!mean.info()->is_resizable());
- BOOST_TEST(!var.info()->is_resizable());
- BOOST_TEST(!beta.info()->is_resizable());
- BOOST_TEST(!gamma.info()->is_resizable());
-
- // Fill tensors
- if(dt == DataType::F32)
- {
- float min_bound = 0.f;
- float max_bound = 0.f;
- std::tie(min_bound, max_bound) = get_batchnormalization_layer_test_bounds<float>();
- std::uniform_real_distribution<> distribution(min_bound, max_bound);
- std::uniform_real_distribution<> distribution_var(0, max_bound);
- library->fill(CLAccessor(src), distribution, 0);
- library->fill(CLAccessor(mean), distribution, 1);
- library->fill(CLAccessor(var), distribution_var, 0);
- library->fill(CLAccessor(beta), distribution, 3);
- library->fill(CLAccessor(gamma), distribution, 4);
- }
- else
- {
- int min_bound = 0;
- int max_bound = 0;
- if(dt == DataType::QS8)
- {
- std::tie(min_bound, max_bound) = get_batchnormalization_layer_test_bounds<int8_t>(fixed_point_position);
- }
- else
- {
- std::tie(min_bound, max_bound) = get_batchnormalization_layer_test_bounds<int16_t>(fixed_point_position);
- }
- std::uniform_int_distribution<> distribution(min_bound, max_bound);
- std::uniform_int_distribution<> distribution_var(0, max_bound);
- library->fill(CLAccessor(src), distribution, 0);
- library->fill(CLAccessor(mean), distribution, 1);
- library->fill(CLAccessor(var), distribution_var, 0);
- library->fill(CLAccessor(beta), distribution, 3);
- library->fill(CLAccessor(gamma), distribution, 4);
- }
-
- // Compute function
- norm.run();
-
- return dst;
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(BatchNormalizationLayer)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, RandomBatchNormalizationLayerDataset() * boost::unit_test::data::make({ DataType::QS8, DataType::QS16, DataType::F32 }), obj, dt)
-{
- // Set fixed point position data type allowed
- int fixed_point_position = (arm_compute::is_data_type_fixed_point(dt)) ? 3 : 0;
-
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(obj.shape0, dt, 1, fixed_point_position);
- CLTensor dst = create_tensor<CLTensor>(obj.shape0, dt, 1, fixed_point_position);
- CLTensor mean = create_tensor<CLTensor>(obj.shape1, dt, 1, fixed_point_position);
- CLTensor var = create_tensor<CLTensor>(obj.shape1, dt, 1, fixed_point_position);
- CLTensor beta = create_tensor<CLTensor>(obj.shape1, dt, 1, fixed_point_position);
- CLTensor gamma = create_tensor<CLTensor>(obj.shape1, dt, 1, fixed_point_position);
-
- BOOST_TEST(src.info()->is_resizable());
- BOOST_TEST(dst.info()->is_resizable());
- BOOST_TEST(mean.info()->is_resizable());
- BOOST_TEST(var.info()->is_resizable());
- BOOST_TEST(beta.info()->is_resizable());
- BOOST_TEST(gamma.info()->is_resizable());
-
- // Create and configure function
- CLBatchNormalizationLayer norm;
- norm.configure(&src, &dst, &mean, &var, &beta, &gamma, obj.epsilon);
-
- // Validate valid region
- const ValidRegion valid_region = shape_to_valid_region(obj.shape0);
- const ValidRegion valid_region_vec = shape_to_valid_region(obj.shape1);
- validate(src.info()->valid_region(), valid_region);
- validate(dst.info()->valid_region(), valid_region);
- validate(mean.info()->valid_region(), valid_region_vec);
- validate(var.info()->valid_region(), valid_region_vec);
- validate(beta.info()->valid_region(), valid_region_vec);
- validate(gamma.info()->valid_region(), valid_region_vec);
-}
-
-BOOST_AUTO_TEST_SUITE(Float)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(Random,
- RandomBatchNormalizationLayerDataset() * boost::unit_test::data::make(DataType::F32),
- obj, dt)
-{
- // Compute function
- CLTensor dst = compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, tolerance_f, 0);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(Quantized)
-
-BOOST_AUTO_TEST_SUITE(QS8)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(Random,
- RandomBatchNormalizationLayerDataset() * boost::unit_test::data::make(DataType::QS8) * boost::unit_test::data::xrange(1, 6),
- obj, dt, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, tolerance_qs8, 0);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(QS16)
-BOOST_DATA_TEST_CASE(Random,
- RandomBatchNormalizationLayerDataset() * boost::unit_test::data::make(DataType::QS16) * boost::unit_test::data::xrange(1, 14),
- obj, dt, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, tolerance_qs16, 0);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/BitwiseAnd.cpp b/tests/validation/CL/BitwiseAnd.cpp
new file mode 100644
index 0000000000..3e458a4c9b
--- /dev/null
+++ b/tests/validation/CL/BitwiseAnd.cpp
@@ -0,0 +1,94 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/functions/CLBitwiseAnd.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/BitwiseAndFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(CL)
+TEST_SUITE(BitwiseAnd)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(concat(datasets::SmallShapes(), datasets::LargeShapes()), framework::dataset::make("DataType", DataType::U8)), shape, data_type)
+{
+ // Create tensors
+ CLTensor src1 = create_tensor<CLTensor>(shape, data_type);
+ CLTensor src2 = create_tensor<CLTensor>(shape, data_type);
+ CLTensor dst = create_tensor<CLTensor>(shape, data_type);
+
+ ARM_COMPUTE_EXPECT(src1.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(src2.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Create and configure function
+ CLBitwiseAnd bitwise_and;
+ bitwise_and.configure(&src1, &src2, &dst);
+
+ // Validate valid region
+ const ValidRegion valid_region = shape_to_valid_region(shape);
+ validate(src1.info()->valid_region(), valid_region);
+ validate(src2.info()->valid_region(), valid_region);
+ validate(dst.info()->valid_region(), valid_region);
+
+ // Validate padding
+ const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+ validate(src1.info()->padding(), padding);
+ validate(src2.info()->padding(), padding);
+ validate(dst.info()->padding(), padding);
+}
+
+template <typename T>
+using CLBitwiseAndFixture = BitwiseAndValidationFixture<CLTensor, CLAccessor, CLBitwiseAnd, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, CLBitwiseAndFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), framework::dataset::make("DataType",
+ DataType::U8)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLBitwiseAndFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapes(), framework::dataset::make("DataType",
+ DataType::U8)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/BitwiseNot.cpp b/tests/validation/CL/BitwiseNot.cpp
new file mode 100644
index 0000000000..376bde9c78
--- /dev/null
+++ b/tests/validation/CL/BitwiseNot.cpp
@@ -0,0 +1,90 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/functions/CLBitwiseNot.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/BitwiseNotFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(CL)
+TEST_SUITE(BitwiseNot)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(concat(datasets::SmallShapes(), datasets::LargeShapes()), framework::dataset::make("DataType", DataType::U8)), shape, data_type)
+{
+ // Create tensors
+ CLTensor src = create_tensor<CLTensor>(shape, data_type);
+ CLTensor dst = create_tensor<CLTensor>(shape, data_type);
+
+ ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Create and configure function
+ CLBitwiseNot bitwise_not;
+ bitwise_not.configure(&src, &dst);
+
+ // Validate valid region
+ const ValidRegion valid_region = shape_to_valid_region(shape);
+ validate(src.info()->valid_region(), valid_region);
+ validate(dst.info()->valid_region(), valid_region);
+
+ // Validate padding
+ const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+ validate(src.info()->padding(), padding);
+ validate(dst.info()->padding(), padding);
+}
+
+template <typename T>
+using CLBitwiseNotFixture = BitwiseNotValidationFixture<CLTensor, CLAccessor, CLBitwiseNot, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, CLBitwiseNotFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), framework::dataset::make("DataType",
+ DataType::U8)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLBitwiseNotFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapes(), framework::dataset::make("DataType",
+ DataType::U8)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/BitwiseOr.cpp b/tests/validation/CL/BitwiseOr.cpp
new file mode 100644
index 0000000000..ecff0be6c0
--- /dev/null
+++ b/tests/validation/CL/BitwiseOr.cpp
@@ -0,0 +1,94 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/functions/CLBitwiseOr.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/BitwiseOrFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(CL)
+TEST_SUITE(BitwiseOr)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(concat(datasets::SmallShapes(), datasets::LargeShapes()), framework::dataset::make("DataType", DataType::U8)), shape, data_type)
+{
+ // Create tensors
+ CLTensor src1 = create_tensor<CLTensor>(shape, data_type);
+ CLTensor src2 = create_tensor<CLTensor>(shape, data_type);
+ CLTensor dst = create_tensor<CLTensor>(shape, data_type);
+
+ ARM_COMPUTE_EXPECT(src1.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(src2.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Create and configure function
+ CLBitwiseOr bitwise_or;
+ bitwise_or.configure(&src1, &src2, &dst);
+
+ // Validate valid region
+ const ValidRegion valid_region = shape_to_valid_region(shape);
+ validate(src1.info()->valid_region(), valid_region);
+ validate(src2.info()->valid_region(), valid_region);
+ validate(dst.info()->valid_region(), valid_region);
+
+ // Validate padding
+ const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+ validate(src1.info()->padding(), padding);
+ validate(src2.info()->padding(), padding);
+ validate(dst.info()->padding(), padding);
+}
+
+template <typename T>
+using CLBitwiseOrFixture = BitwiseOrValidationFixture<CLTensor, CLAccessor, CLBitwiseOr, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, CLBitwiseOrFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), framework::dataset::make("DataType",
+ DataType::U8)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLBitwiseOrFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapes(), framework::dataset::make("DataType",
+ DataType::U8)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/BitwiseXor.cpp b/tests/validation/CL/BitwiseXor.cpp
new file mode 100644
index 0000000000..3104894c29
--- /dev/null
+++ b/tests/validation/CL/BitwiseXor.cpp
@@ -0,0 +1,94 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/functions/CLBitwiseXor.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/BitwiseXorFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(CL)
+TEST_SUITE(BitwiseXor)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(concat(datasets::SmallShapes(), datasets::LargeShapes()), framework::dataset::make("DataType", DataType::U8)), shape, data_type)
+{
+ // Create tensors
+ CLTensor src1 = create_tensor<CLTensor>(shape, data_type);
+ CLTensor src2 = create_tensor<CLTensor>(shape, data_type);
+ CLTensor dst = create_tensor<CLTensor>(shape, data_type);
+
+ ARM_COMPUTE_EXPECT(src1.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(src2.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Create and configure function
+ CLBitwiseXor bitwise_xor;
+ bitwise_xor.configure(&src1, &src2, &dst);
+
+ // Validate valid region
+ const ValidRegion valid_region = shape_to_valid_region(shape);
+ validate(src1.info()->valid_region(), valid_region);
+ validate(src2.info()->valid_region(), valid_region);
+ validate(dst.info()->valid_region(), valid_region);
+
+ // Validate padding
+ const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+ validate(src1.info()->padding(), padding);
+ validate(src2.info()->padding(), padding);
+ validate(dst.info()->padding(), padding);
+}
+
+template <typename T>
+using CLBitwiseXorFixture = BitwiseXorValidationFixture<CLTensor, CLAccessor, CLBitwiseXor, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, CLBitwiseXorFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), framework::dataset::make("DataType",
+ DataType::U8)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLBitwiseXorFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapes(), framework::dataset::make("DataType",
+ DataType::U8)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/Box3x3.cpp b/tests/validation/CL/Box3x3.cpp
deleted file mode 100644
index f2df5e6511..0000000000
--- a/tests/validation/CL/Box3x3.cpp
+++ /dev/null
@@ -1,166 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "PaddingCalculator.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-#include "validation/ValidationUserConfiguration.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/CLTensor.h"
-#include "arm_compute/runtime/CL/CLTensorAllocator.h"
-#include "arm_compute/runtime/CL/functions/CLBox3x3.h"
-
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-constexpr unsigned int filter_size = 3; /** Size of the kernel/filter in number of elements. */
-constexpr BorderSize border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */
-
-/** Compute CL box3x3 filter.
- *
- * @param[in] shape Shape of the input and output tensors.
- * @param[in] border_mode BorderMode used by the input tensor.
- * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT.
- *
- * @return Computed output tensor.
- */
-CLTensor compute_box3x3(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U8);
-
- // Create and configure function
- CLBox3x3 box3x3;
- box3x3.configure(&src, &dst, border_mode, constant_border_value);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src), 0);
-
- // Compute function
- box3x3.run();
-
- return dst;
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(Box3x3)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U8);
-
- BOOST_TEST(src.info()->is_resizable());
- BOOST_TEST(dst.info()->is_resizable());
-
- // Create and configure function
- CLBox3x3 box3x3;
- box3x3.configure(&src, &dst, border_mode);
-
- // Validate valid region
- const ValidRegion src_valid_region = shape_to_valid_region(shape);
- const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
- validate(src.info()->valid_region(), src_valid_region);
- validate(dst.info()->valid_region(), dst_valid_region);
-
- // Validate padding
- PaddingCalculator calculator(shape.x(), 8);
- calculator.set_border_size(1);
- calculator.set_border_mode(border_mode);
-
- const PaddingSize dst_padding = calculator.required_padding();
-
- calculator.set_accessed_elements(16);
- calculator.set_access_offset(-1);
-
- const PaddingSize src_padding = calculator.required_padding();
-
- validate(src.info()->padding(), src_padding);
- validate(dst.info()->padding(), dst_padding);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode)
-{
- std::mt19937 gen(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution(0, 255);
- const uint8_t border_value = distribution(gen);
-
- // Compute function
- CLTensor dst = compute_box3x3(shape, border_mode, border_value);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_box3x3(shape, border_mode, border_value);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size));
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode)
-{
- std::mt19937 gen(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution(0, 255);
- const uint8_t border_value = distribution(gen);
-
- // Compute function
- CLTensor dst = compute_box3x3(shape, border_mode, border_value);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_box3x3(shape, border_mode, border_value);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size));
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/CLFixture.cpp b/tests/validation/CL/CLFixture.cpp
deleted file mode 100644
index 38e52c31f6..0000000000
--- a/tests/validation/CL/CLFixture.cpp
+++ /dev/null
@@ -1,32 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "validation/CL/CLFixture.h"
-
-#include "boost_wrapper.h"
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-BOOST_GLOBAL_FIXTURE(CLFixture);
diff --git a/tests/validation/CL/CMakeLists.txt b/tests/validation/CL/CMakeLists.txt
deleted file mode 100644
index f4477f621f..0000000000
--- a/tests/validation/CL/CMakeLists.txt
+++ /dev/null
@@ -1,68 +0,0 @@
-# Copyright (c) 2017 ARM Limited.
-#
-# SPDX-License-Identifier: MIT
-#
-# Permission is hereby granted, free of charge, to any person obtaining a copy
-# of this software and associated documentation files (the "Software"), to
-# deal in the Software without restriction, including without limitation the
-# rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
-# sell copies of the Software, and to permit persons to whom the Software is
-# furnished to do so, subject to the following conditions:
-#
-# The above copyright notice and this permission notice shall be included in all
-# copies or substantial portions of the Software.
-#
-# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
-# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
-# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
-# SOFTWARE.
-cmake_minimum_required (VERSION 3.1)
-
-include_directories(${CMAKE_SOURCE_DIR}/../include)
-
-set(arm_compute_test_validation_OPENCL_SOURCE_FILES
- ${CMAKE_SOURCE_DIR}/CL/CLAccessor.h
- ${CMAKE_CURRENT_SOURCE_DIR}/CLFixture.h
- ${CMAKE_CURRENT_SOURCE_DIR}/CLFixture.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/ActivationLayer.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/BitwiseAnd.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/Box3x3.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/ConvolutionLayer.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/DepthConvert.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/FillBorder.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/FixedPoint/FixedPoint_QS8.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/FullyConnectedLayer.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/Gaussian3x3.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/Gaussian5x5.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/GEMM.cpp
-` ${CMAKE_CURRENT_SOURCE_DIR}/IntegralImage.cpp
-` ${CMAKE_CURRENT_SOURCE_DIR}/NonLinearFilter.cpp
-` ${CMAKE_CURRENT_SOURCE_DIR}/PoolingLayer.cpp
-` ${CMAKE_CURRENT_SOURCE_DIR}/ROIPoolingLayer.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/Sobel3x3.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/Sobel5x5.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/SoftmaxLayer.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/Threshold.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/DirectConvolutionLayer.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/MeanStdDev.cpp
- ${CMAKE_CURRENT_SOURCE_DIR}/HarrisCorners.cpp
-)
-
-add_library(arm_compute_test_validation_OPENCL OBJECT
- ${arm_compute_test_validation_OPENCL_SOURCE_FILES}
-)
-
-set(arm_compute_test_validation_TARGET_OBJECTS
- ${arm_compute_test_validation_TARGET_OBJECTS}
- $<TARGET_OBJECTS:arm_compute_test_validation_OPENCL>
- PARENT_SCOPE
-)
-
-set(arm_compute_test_validation_TARGET_LIBRARIES
- ${arm_compute_test_validation_TARGET_LIBRARIES}
- OpenCL
- PARENT_SCOPE
-)
diff --git a/tests/validation/CL/ConvolutionLayer.cpp b/tests/validation/CL/ConvolutionLayer.cpp
new file mode 100644
index 0000000000..d7123842e9
--- /dev/null
+++ b/tests/validation/CL/ConvolutionLayer.cpp
@@ -0,0 +1,186 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONCLCTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLConvolutionLayer.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/LargeConvolutionLayerDataset.h"
+#include "tests/datasets/SmallConvolutionLayerDataset.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/ConvolutionLayerFixture.h"
+#include "tests/validation/half.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+constexpr AbsoluteTolerance<float> tolerance_f32(0.001f); /**< Tolerance value for comparing reference's output against implementation's output for DataType::F32 */
+constexpr AbsoluteTolerance<float> tolerance_f16(0.1f); /**< Tolerance value for comparing reference's output against implementation's output for DataType::F16 */
+constexpr AbsoluteTolerance<float> tolerance_q(1.0f); /**< Tolerance value for comparing reference's output against implementation's output for fixed point data types */
+
+/** CNN data types */
+const auto CNNDataTypes = framework::dataset::make("DataType",
+{
+ DataType::F16,
+ DataType::F32,
+ DataType::QS8,
+ DataType::QS16,
+});
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(ConvolutionLayer)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(framework::dataset::concat(datasets::SmallConvolutionLayerDataset(), datasets::LargeConvolutionLayerDataset()), CNNDataTypes),
+ input_shape, weights_shape, bias_shape, output_shape, info, data_type)
+{
+ // Set fixed point position data type allowed
+ int fixed_point_position = is_data_type_fixed_point(data_type) ? 3 : 0;
+
+ // Create tensors
+ CLTensor src = create_tensor<CLTensor>(input_shape, data_type, 1, fixed_point_position);
+ CLTensor weights = create_tensor<CLTensor>(weights_shape, data_type, 1, fixed_point_position);
+ CLTensor bias = create_tensor<CLTensor>(bias_shape, data_type, 1, fixed_point_position);
+ CLTensor dst = create_tensor<CLTensor>(output_shape, data_type, 1, fixed_point_position);
+
+ ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(weights.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);
+
+ // Create and configure function
+ CLConvolutionLayer conv;
+ conv.configure(&src, &weights, &bias, &dst, info);
+
+ // Validate valid region
+ const ValidRegion src_valid_region = shape_to_valid_region(input_shape);
+ const ValidRegion weights_valid_region = shape_to_valid_region(weights_shape);
+ const ValidRegion bias_valid_region = shape_to_valid_region(bias_shape);
+ const ValidRegion dst_valid_region = shape_to_valid_region(output_shape);
+
+ validate(src.info()->valid_region(), src_valid_region);
+ validate(weights.info()->valid_region(), weights_valid_region);
+ validate(bias.info()->valid_region(), bias_valid_region);
+ validate(dst.info()->valid_region(), dst_valid_region);
+
+ // Validate padding
+ //TODO(COMPMID-415) Need to validate padding?
+}
+
+template <typename T>
+using CLConvolutionLayerFixture = ConvolutionValidationFixture<CLTensor, CLAccessor, CLConvolutionLayer, T>;
+
+TEST_SUITE(Float)
+TEST_SUITE(FP16)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLConvolutionLayerFixture<half_float::half>, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallConvolutionLayerDataset(),
+ framework::dataset::make("DataType",
+ DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f16);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLConvolutionLayerFixture<half_float::half>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeConvolutionLayerDataset(),
+ framework::dataset::make("DataType",
+ DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f16);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLConvolutionLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallConvolutionLayerDataset(), framework::dataset::make("DataType",
+ DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLConvolutionLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeConvolutionLayerDataset(), framework::dataset::make("DataType",
+ DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+template <typename T>
+using CLConvolutionLayerFixedPointFixture = ConvolutionValidationFixedPointFixture<CLTensor, CLAccessor, CLConvolutionLayer, T>;
+
+TEST_SUITE(Quantized)
+TEST_SUITE(QS8)
+// We test for fixed point precision [4,6]
+FIXTURE_DATA_TEST_CASE(RunSmall, CLConvolutionLayerFixedPointFixture<int8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallConvolutionLayerDataset(),
+ framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FractionalBits", 4, 7)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_q);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLConvolutionLayerFixedPointFixture<int8_t>, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeConvolutionLayerDataset(),
+ framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FractionalBits", 4, 7)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_q);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(QS16)
+// Testing for fixed point position [1,14)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLConvolutionLayerFixedPointFixture<int16_t>, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallConvolutionLayerDataset(),
+ framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FractionalBits", 1, 14)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_q);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLConvolutionLayerFixedPointFixture<int16_t>, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeConvolutionLayerDataset(),
+ framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FractionalBits", 1, 14)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_q);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/DepthConcatenateLayer.cpp b/tests/validation/CL/DepthConcatenateLayer.cpp
new file mode 100644
index 0000000000..a8ef1c37c7
--- /dev/null
+++ b/tests/validation/CL/DepthConcatenateLayer.cpp
@@ -0,0 +1,123 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLDepthConcatenate.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/DepthConcatenateLayerFixture.h"
+#include "tests/validation/half.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(CL)
+TEST_SUITE(DepthConcatenateLayer)
+
+//TODO(COMPMID-415): Add configuration test?
+
+template <typename T>
+using CLDepthConcatenateLayerFixture = DepthConcatenateValidationFixture<CLTensor, CLAccessor, CLDepthConcatenate, T>;
+
+TEST_SUITE(Float)
+TEST_SUITE(FP16)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLDepthConcatenateLayerFixture<half_float::half>, framework::DatasetMode::PRECOMMIT, combine(datasets::Small2DShapes(), framework::dataset::make("DataType",
+ DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLDepthConcatenateLayerFixture<half_float::half>, framework::DatasetMode::NIGHTLY, combine(datasets::Large2DShapes(), framework::dataset::make("DataType",
+ DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLDepthConcatenateLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(datasets::Small2DShapes(), framework::dataset::make("DataType",
+ DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLDepthConcatenateLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(datasets::Large2DShapes(), framework::dataset::make("DataType",
+ DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+TEST_SUITE(Quantized)
+TEST_SUITE(QS8)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLDepthConcatenateLayerFixture<int8_t>, framework::DatasetMode::PRECOMMIT, combine(datasets::Small2DShapes(),
+ framework::dataset::make("DataType",
+ DataType::QS8)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLDepthConcatenateLayerFixture<int8_t>, framework::DatasetMode::NIGHTLY, combine(datasets::Large2DShapes(),
+ framework::dataset::make("DataType",
+ DataType::QS8)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(QS16)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLDepthConcatenateLayerFixture<int16_t>, framework::DatasetMode::PRECOMMIT, combine(datasets::Small2DShapes(),
+ framework::dataset::make("DataType",
+ DataType::QS16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLDepthConcatenateLayerFixture<int16_t>, framework::DatasetMode::NIGHTLY, combine(datasets::Large2DShapes(),
+ framework::dataset::make("DataType",
+ DataType::QS16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/DepthConvert.cpp b/tests/validation/CL/DepthConvert.cpp
deleted file mode 100644
index 57a21e89c7..0000000000
--- a/tests/validation/CL/DepthConvert.cpp
+++ /dev/null
@@ -1,495 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "PaddingCalculator.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/functions/CLDepthConvert.h"
-#include "arm_compute/runtime/Tensor.h"
-#include "arm_compute/runtime/TensorAllocator.h"
-
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-/** Compute CL depth convert function.
- *
- * @param[in] shape Shape of the input and output tensors.
- * @param[in] dt_in Data type of input tensor.
- * @param[in] dt_out Data type of the output tensor.
- * @param[in] policy Conversion policy.
- * @param[in] shift Value for down/up conversions. Must be 0 <= shift < 8.
- *
- * @return Computed output CLtensor.
- */
-CLTensor compute_depth_convert(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy, uint32_t shift, uint32_t fixed_point_position = 0)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, dt_in, 1, fixed_point_position);
- CLTensor dst = create_tensor<CLTensor>(shape, dt_out, 1, fixed_point_position);
-
- // Create and configure function
- CLDepthConvert depth_convert;
- depth_convert.configure(&src, &dst, policy, shift);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src), 0);
-
- // Compute function
- depth_convert.run();
-
- return dst;
-}
-/** Configure and validate region/padding function.
- *
- * @param[in] shape Shape of the input and output tensors.
- * @param[in] dt_in Data type of input tensor.
- * @param[in] dt_out Data type of the output tensor.
- * @param[in] policy Conversion policy.
- * @param[in] shift Value for down/up conversions. Must be 0 <= shift < 8.
- * @param[in] fixed_point_position Fixed point position.
- *
- */
-void compute_configure_validate(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy, uint32_t shift, uint32_t fixed_point_position = 0)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, dt_in, 1, fixed_point_position);
- CLTensor dst = create_tensor<CLTensor>(shape, dt_out, 1, fixed_point_position);
-
- BOOST_TEST(src.info()->is_resizable());
- BOOST_TEST(dst.info()->is_resizable());
-
- // Create and configure function
- CLDepthConvert depth_convert;
- depth_convert.configure(&src, &dst, policy, shift);
-
- // Validate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape);
- validate(src.info()->valid_region(), valid_region);
- validate(dst.info()->valid_region(), valid_region);
-
- // Validate padding
- const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
- validate(src.info()->padding(), padding);
- validate(dst.info()->padding(), padding);
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(DepthConvert)
-
-BOOST_AUTO_TEST_SUITE(U8_to_U16)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute configure and validate region/padding
- compute_configure_validate(shape, DataType::U8, DataType::U16, policy, shift);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::U8, DataType::U16, policy, shift);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::U16, policy, shift, 0);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::U8, DataType::U16, policy, shift);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::U16, policy, shift, 0);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(U8_to_S16)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute configure and validate region/padding
- compute_configure_validate(shape, DataType::U8, DataType::S16, policy, shift);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::U8, DataType::S16, policy, shift);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S16, policy, shift, 0);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::U8, DataType::S16, policy, shift);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S16, policy, shift, 0);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(U8_to_S32)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute configure and validate region/padding
- compute_configure_validate(shape, DataType::U8, DataType::S32, policy, shift);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::U8, DataType::S32, policy, shift);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S32, policy, shift, 0);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::U8, DataType::S32, policy, shift);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S32, policy, shift, 0);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(U16_to_U8)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute configure and validate region/padding
- compute_configure_validate(shape, DataType::U16, DataType::U8, policy, shift);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::U16, DataType::U8, policy, shift);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U8, policy, shift, 0);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::U16, DataType::U8, policy, shift);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U8, policy, shift, 0);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(U16_to_U32)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute configure and validate region/padding
- compute_configure_validate(shape, DataType::U16, DataType::U32, policy, shift);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::U16, DataType::U32, policy, shift);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U32, policy, shift, 0);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::U16, DataType::U32, policy, shift);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U32, policy, shift, 0);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(S16_to_U8)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute configure and validate region/padding
- compute_configure_validate(shape, DataType::S16, DataType::U8, policy, shift);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::S16, DataType::U8, policy, shift);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::U8, policy, shift, 0);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::S16, DataType::U8, policy, shift);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::U8, policy, shift, 0);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(S16_to_S32)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute configure and validate region/padding
- compute_configure_validate(shape, DataType::S16, DataType::S32, policy, shift);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::S16, DataType::S32, policy, shift);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::S32, policy, shift, 0);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
- * boost::unit_test::data::xrange(0, 7, 1),
- shape, policy, shift)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::S16, DataType::S32, policy, shift);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::S32, policy, shift, 0);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(Quantized_to_F32)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::QS8, DataType::QS16 })
- * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) * boost::unit_test::data::xrange(1, 7, 1),
- shape, dt, policy, fixed_point_position)
-{
- // Compute configure and validate region/padding
- compute_configure_validate(shape, dt, DataType::F32, policy, 0, fixed_point_position);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
- * boost::unit_test::data::make({ DataType::QS8, DataType::QS16 }) * boost::unit_test::data::xrange(1, 7, 1),
- shape, policy, dt, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, dt, DataType::F32, policy, 0, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, dt, DataType::F32, policy, 0, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ DataType::QS8, DataType::QS16 })
- * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) * boost::unit_test::data::xrange(1, 7, 1),
- shape, dt, policy, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, dt, DataType::F32, policy, 0, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, dt, DataType::F32, policy, 0, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(F32_to_Quantized)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::QS8, DataType::QS16 })
- * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) * boost::unit_test::data::xrange(1, 7, 1),
- shape, dt, policy, fixed_point_position)
-{
- // Compute configure and validate region/padding
- compute_configure_validate(shape, DataType::F32, dt, policy, 0, fixed_point_position);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ DataType::QS8, DataType::QS16 })
- * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) * boost::unit_test::data::xrange(1, 7, 1),
- shape, dt, policy, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::F32, dt, policy, 0, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, dt, policy, 0, fixed_point_position, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ DataType::QS8, DataType::QS16 })
- * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) * boost::unit_test::data::xrange(1, 7, 1),
- shape, dt, policy, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_depth_convert(shape, DataType::F32, dt, policy, 0, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, dt, policy, 0, fixed_point_position, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/DepthwiseConvolution.cpp b/tests/validation/CL/DepthwiseConvolution.cpp
new file mode 100644
index 0000000000..ccee9607d3
--- /dev/null
+++ b/tests/validation/CL/DepthwiseConvolution.cpp
@@ -0,0 +1,69 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONCLCTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLDepthwiseConvolution.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/LargeDepthwiseConvolutionDataset.h"
+#include "tests/datasets/SmallDepthwiseConvolutionDataset.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/DepthwiseConvolutionFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+constexpr AbsoluteTolerance<float> tolerance_f32(0.001f); /**< Tolerance value for comparing reference's output against implementation's output for DataType::F32 */
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(DepthwiseConvolution)
+
+template <typename T>
+using CLDepthwiseConvolutionFixture = DepthwiseConvolutionValidationFixture<CLTensor, CLAccessor, CLDepthwiseConvolution, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, CLDepthwiseConvolutionFixture<float>, framework::DatasetMode::PRECOMMIT, datasets::SmallDepthwiseConvolutionDataset())
+{
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+
+FIXTURE_DATA_TEST_CASE(RunLarge, CLDepthwiseConvolutionFixture<float>, framework::DatasetMode::NIGHTLY, datasets::LargeDepthwiseConvolutionDataset())
+{
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/DepthwiseSeparableConvolutionLayer.cpp b/tests/validation/CL/DepthwiseSeparableConvolutionLayer.cpp
new file mode 100644
index 0000000000..4fac9b2d0d
--- /dev/null
+++ b/tests/validation/CL/DepthwiseSeparableConvolutionLayer.cpp
@@ -0,0 +1,64 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONCLCTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLDepthwiseSeparableConvolutionLayer.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/MobileNetDepthwiseSeparableConvolutionLayerDataset.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/DepthwiseSeparableConvolutionLayerFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+constexpr AbsoluteTolerance<float> tolerance_f32(0.001f); /**< Tolerance value for comparing reference's output against implementation's output for DataType::F32 */
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(DepthwiseSeparableConvolutionLayer)
+
+// Configuration test to do
+
+template <typename T>
+using CLDepthwiseSeparableConvolutionLayerFixture = DepthwiseSeparableConvolutionValidationFixture<CLTensor, CLAccessor, CLDepthwiseSeparableConvolutionLayer, T>;
+FIXTURE_DATA_TEST_CASE(RunSmall, CLDepthwiseSeparableConvolutionLayerFixture<float>, framework::DatasetMode::PRECOMMIT, datasets::MobileNetDepthwiseSeparableConvolutionLayerDataset())
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/DirectConvolutionLayer.cpp b/tests/validation/CL/DirectConvolutionLayer.cpp
new file mode 100644
index 0000000000..d2d2cd1419
--- /dev/null
+++ b/tests/validation/CL/DirectConvolutionLayer.cpp
@@ -0,0 +1,129 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLDirectConvolutionLayer.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/DirectConvolutionLayerFixture.h"
+#include "tests/validation/half.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+constexpr AbsoluteTolerance<float> tolerance_fp16(0.1f); /**< Tolerance for floating point tests */
+constexpr AbsoluteTolerance<float> tolerance_fp32(0.001f); /**< Tolerance for floating point tests */
+
+constexpr AbsoluteTolerance<int8_t> tolerance_qs8(0); /**< Tolerance for fixed point tests */
+constexpr AbsoluteTolerance<int16_t> tolerance_qs16(0); /**< Tolerance for fixed point tests */
+
+/** Direct convolution data set. */
+const auto data_quantized = combine(datasets::SmallDirectConvolutionShapes(),
+ combine(framework::dataset::make("StrideX", 1, 3),
+ combine(framework::dataset::make("StrideY", 1, 3),
+ combine(concat(combine(framework::dataset::make("PadX", 0),
+ combine(framework::dataset::make("PadY", 0),
+ framework::dataset::make("KernelSize", 1))),
+ combine(framework::dataset::make("PadX", 0, 2),
+ combine(framework::dataset::make("PadY", 0, 2),
+ framework::dataset::make("KernelSize", { 3 })))),
+ framework::dataset::make("NumKernels", { 1, 4, 8, 16 })))));
+
+const auto data = combine(datasets::SmallDirectConvolutionShapes(),
+ combine(framework::dataset::make("StrideX", 1, 3),
+ combine(framework::dataset::make("StrideY", 1, 3),
+ combine(concat(combine(framework::dataset::make("PadX", 0),
+ combine(framework::dataset::make("PadY", 0),
+ framework::dataset::make("KernelSize", 1))),
+ combine(framework::dataset::make("PadX", 0, 2),
+ combine(framework::dataset::make("PadY", 0, 2),
+ framework::dataset::make("KernelSize", { 3, 5 })))),
+ framework::dataset::make("NumKernels", { 1, 4, 8, 16 })))));
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(DirectConvolutionLayer)
+
+//TODO(COMPMID-415): Configuration tests?
+
+template <typename T>
+using CLDirectConvolutionLayerFixture = DirectConvolutionValidationFixture<CLTensor, CLAccessor, CLDirectConvolutionLayer, T>;
+
+TEST_SUITE(Float)
+TEST_SUITE(FP16)
+FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixture<half_float::half>, framework::DatasetMode::ALL, combine(data, framework::dataset::make("DataType", DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_fp16);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixture<float>, framework::DatasetMode::ALL, combine(data, framework::dataset::make("DataType", DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_fp32);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+template <typename T>
+using CLDirectConvolutionLayerFixedPointFixture = DirectConvolutionValidationFixedPointFixture<CLTensor, CLAccessor, CLDirectConvolutionLayer, T>;
+
+TEST_SUITE(Quantized)
+TEST_SUITE(QS8)
+FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixedPointFixture<int8_t>, framework::DatasetMode::ALL, combine(combine(data_quantized, framework::dataset::make("DataType", DataType::QS8)),
+ framework::dataset::make("FractionalBits", 2, 7)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_qs8);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(QS16)
+FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixedPointFixture<int16_t>, framework::DatasetMode::ALL, combine(combine(data_quantized, framework::dataset::make("DataType", DataType::QS16)),
+ framework::dataset::make("FractionalBits", 2, 15)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_qs16);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/FillBorder.cpp b/tests/validation/CL/FillBorder.cpp
deleted file mode 100644
index dc522c5c3b..0000000000
--- a/tests/validation/CL/FillBorder.cpp
+++ /dev/null
@@ -1,89 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Validation.h"
-
-#include "arm_compute/core/CL/kernels/CLFillBorderKernel.h"
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/CLScheduler.h"
-#include "arm_compute/runtime/CL/CLTensor.h"
-#include "arm_compute/runtime/CL/CLTensorAllocator.h"
-
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(FillBorder, BorderModes() * boost::unit_test::data::make({ PaddingSize{ 0 }, PaddingSize{ 1, 0, 1, 2 }, PaddingSize{ 10 } }), border_mode, padding)
-{
- constexpr uint8_t border_value = 42U;
- constexpr uint8_t tensor_value = 89U;
- BorderSize border_size{ 5 };
-
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(TensorShape{ 10U, 10U, 2U }, DataType::U8);
-
- src.info()->extend_padding(padding);
-
- // Allocate tensor
- src.allocator()->allocate();
-
- // Check padding is as required
- validate(src.info()->padding(), padding);
-
- // Fill tensor with constant value
- std::uniform_int_distribution<uint8_t> distribution{ tensor_value, tensor_value };
- library->fill(CLAccessor(src), distribution, 0);
-
- // Create and configure kernel
- CLFillBorderKernel fill_border;
- fill_border.configure(&src, border_size, border_mode, border_value);
-
- // Run kernel
- fill_border.run(fill_border.window(), CLScheduler::get().queue());
-
- // Validate border
- border_size.limit(padding);
- validate(CLAccessor(src), border_size, border_mode, &border_value);
-
- // Validate tensor
- validate(CLAccessor(src), &tensor_value);
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/FixedPoint/FixedPoint_QS8.cpp b/tests/validation/CL/FixedPoint/FixedPoint_QS8.cpp
deleted file mode 100644
index d4ccc2ba2b..0000000000
--- a/tests/validation/CL/FixedPoint/FixedPoint_QS8.cpp
+++ /dev/null
@@ -1,224 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-
-#include "arm_compute/core/CL/CLKernelLibrary.h"
-#include "arm_compute/core/CL/ICLKernel.h"
-#include "arm_compute/core/CL/OpenCL.h"
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/core/Window.h"
-#include "arm_compute/runtime/CL/CLScheduler.h"
-#include "arm_compute/runtime/CL/CLSubTensor.h"
-#include "arm_compute/runtime/CL/CLTensor.h"
-#include "arm_compute/runtime/CL/CLTensorAllocator.h"
-
-#include "arm_compute/core/CL/ICLTensor.h"
-
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-const float tolerance_exp = 1.0f; /**< Tolerance value for comparing reference's output against implementation's output (exponential)*/
-const float tolerance_invsqrt = 4.0f; /**< Tolerance value for comparing reference's output against implementation's output (inverse square-root) */
-const float tolerance_log = 5.0f; /**< Tolerance value for comparing reference's output against implementation's output (logarithm) */
-
-/** Compute Neon fixed point operation for signed 8bit fixed point.
- *
- * @param[in] shape Shape of the input and output tensors.
- *
- * @return Computed output tensor.
- */
-CLTensor compute_fixed_point_op(const TensorShape &shape, int fixed_point_position, FixedPointOp op)
-{
- std::string fixed_point_operation_kernel;
-#ifndef EMBEDDED_KERNELS
- fixed_point_operation_kernel += "#include \"fixed_point.h\"\n";
-#endif /* EMBEDDED_KERNELS */
- fixed_point_operation_kernel +=
- "__kernel void fixed_point_operation_qs8( \n"
- " __global char* src, \n"
- " __global char* dst) \n"
- "{ \n"
- " char16 in = vload16(0, src + get_global_id(0) * 16); \n"
- " if(FIXED_POINT_OP == 0) \n"
- " { \n"
- " vstore16(EXP_OP_EXPAND(in, DATA_TYPE, 16, FIXED_POINT_POS), 0, dst + get_global_id(0) * 16); \n"
- " } \n"
- " else if(FIXED_POINT_OP == 1) \n"
- " { \n"
- " vstore16(INVSQRT_OP_EXPAND(in, DATA_TYPE, 16, FIXED_POINT_POS), 0, dst + get_global_id(0) * 16); \n"
- " } \n"
- " else \n"
- " { \n"
- " vstore16(LOG_OP_EXPAND(in, DATA_TYPE, 16, FIXED_POINT_POS), 0, dst + get_global_id(0) * 16); \n"
- " } \n"
- "} \n"
- "\n";
-
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::QS8, 1, fixed_point_position);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::QS8, 1, fixed_point_position);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
-
- // Set build options
- std::string build_opts = "-DFIXED_POINT_POS=" + support::cpp11::to_string(fixed_point_position);
- build_opts += " -DDATA_TYPE=qs8";
-
- // Fill tensors.
- int min = 0;
- int max = 0;
- switch(op)
- {
- case FixedPointOp::EXP:
- min = -(1 << (fixed_point_position - 1));
- max = (1 << (fixed_point_position - 1));
- build_opts += " -DFIXED_POINT_OP=0";
- break;
- case FixedPointOp::INV_SQRT:
- min = 1;
- max = 0x7F;
- build_opts += " -DFIXED_POINT_OP=1";
- break;
- case FixedPointOp::LOG:
- min = (1 << (fixed_point_position - 1));
- max = 0x3F;
- build_opts += " -DFIXED_POINT_OP=2";
- break;
- default:
- ARM_COMPUTE_ERROR("Operation not supported");
- }
-
- std::uniform_int_distribution<> distribution(min, max);
- library->fill(CLAccessor(src), distribution, 0);
-
- std::vector<std::string> sources;
-
-#ifndef EMBEDDED_KERNELS
- build_opts += " -I" + CLKernelLibrary::get().get_kernel_path();
-#else /* EMBEDDED_KERNELS */
- sources.push_back(CLKernelLibrary::get().get_program_source("fixed_point.h"));
-#endif /* EMBEDDED_KERNELS */
-
- sources.push_back(fixed_point_operation_kernel);
-
- // Create program
- ::cl::Program program(sources);
-
- // Build program
- program.build(build_opts.c_str());
-
- ::cl::Kernel kernel(program, "fixed_point_operation_qs8", nullptr);
-
- unsigned int idx = 0;
- kernel.setArg(idx++, src.cl_buffer());
- kernel.setArg(idx++, dst.cl_buffer());
-
- ::cl::NDRange gws(shape[0] / 16, 1, 1);
- CLScheduler::get().queue().enqueueNDRangeKernel(kernel, 0, gws);
-
- return dst;
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(FixedPoint)
-BOOST_AUTO_TEST_SUITE(QS8)
-
-BOOST_AUTO_TEST_SUITE(Exp)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(1, 6), shape, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_fixed_point_op(shape, fixed_point_position, FixedPointOp::EXP);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS8, DataType::QS8, FixedPointOp::EXP, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, tolerance_exp);
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(Log)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(3, 6), shape, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_fixed_point_op(shape, fixed_point_position, FixedPointOp::LOG);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS8, DataType::QS8, FixedPointOp::LOG, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, tolerance_log);
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(Invsqrt)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(1, 6), shape, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_fixed_point_op(shape, fixed_point_position, FixedPointOp::INV_SQRT);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS8, DataType::QS8, FixedPointOp::INV_SQRT, fixed_point_position);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, tolerance_invsqrt);
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/CLFixture.h b/tests/validation/CL/Floor.cpp
index 77538be8f4..81495e8337 100644
--- a/tests/validation/CL/CLFixture.h
+++ b/tests/validation/CL/Floor.cpp
@@ -21,10 +21,18 @@
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*/
-#ifndef __ARM_COMPUTE_TEST_VALIDATION_CL_CLFIXTURE_H__
-#define __ARM_COMPUTE_TEST_VALIDATION_CL_CLFIXTURE_H__
-
-#include "arm_compute/runtime/CL/CLScheduler.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLFloor.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/FloorFixture.h"
namespace arm_compute
{
@@ -32,14 +40,27 @@ namespace test
{
namespace validation
{
-struct CLFixture
+TEST_SUITE(CL)
+TEST_SUITE(Floor)
+
+template <typename T>
+using CLFloorFixture = FloorValidationFixture<CLTensor, CLAccessor, CLFloor, T>;
+
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLFloorFixture<float>, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::F32)))
{
- CLFixture()
- {
- CLScheduler::get().default_init();
- }
-};
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLFloorFixture<float>, framework::DatasetMode::NIGHTLY, combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference);
+}
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
} // namespace validation
} // namespace test
} // namespace arm_compute
-#endif /* __ARM_COMPUTE_TEST_VALIDATION_CL_CLFIXTURE_H__ */
diff --git a/tests/validation/CL/FullyConnectedLayer.cpp b/tests/validation/CL/FullyConnectedLayer.cpp
new file mode 100644
index 0000000000..7a8734b5a3
--- /dev/null
+++ b/tests/validation/CL/FullyConnectedLayer.cpp
@@ -0,0 +1,195 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLFullyConnectedLayer.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/FullyConnectedLayerDataset.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/FullyConnectedLayerFixture.h"
+#include "tests/validation/half.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Tolerance for float operations */
+constexpr AbsoluteTolerance<float> tolerance_f32(0.001f);
+constexpr AbsoluteTolerance<float> tolerance_f16(0.4f);
+/** Tolerance for fixed point operations */
+constexpr AbsoluteTolerance<float> tolerance_fixed_point(1.f);
+
+/** CNN data types */
+const auto CNNDataTypes = framework::dataset::make("DataType",
+{
+ DataType::F16,
+ DataType::F32,
+ DataType::QS8,
+ DataType::QS16,
+});
+
+const auto FullyConnectedParameters = combine(framework::dataset::make("TransposeWeights", { false, true }), framework::dataset::make("ReshapeWeights", { false, true }));
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(FullyConnectedLayer)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(combine(framework::dataset::concat(datasets::SmallFullyConnectedLayerDataset(), datasets::LargeFullyConnectedLayerDataset()),
+ FullyConnectedParameters),
+ CNNDataTypes),
+ src_shape, weights_shape, bias_shape, dst_shape, transpose_weights, reshape_weights, data_type)
+{
+ // Set fixed point position data type allowed
+ int fixed_point_position = is_data_type_fixed_point(data_type) ? 3 : 0;
+
+ TensorShape ws(weights_shape);
+
+ // Transpose weights if not done in the function
+ if(!reshape_weights || !transpose_weights)
+ {
+ const size_t shape_x = ws.x();
+ ws.set(0, ws.y());
+ ws.set(1, shape_x);
+ }
+
+ // Create tensors
+ CLTensor src = create_tensor<CLTensor>(src_shape, data_type, 1, fixed_point_position);
+ CLTensor weights = create_tensor<CLTensor>(ws, data_type, 1, fixed_point_position);
+ CLTensor bias = create_tensor<CLTensor>(bias_shape, data_type, 1, fixed_point_position);
+ CLTensor dst = create_tensor<CLTensor>(dst_shape, data_type, 1, fixed_point_position);
+
+ ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(weights.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);
+
+ // Create and configure function.
+ CLFullyConnectedLayer fc;
+ fc.configure(&src, &weights, &bias, &dst, transpose_weights, !reshape_weights);
+
+ // Validate valid region
+ const ValidRegion dst_valid_region = shape_to_valid_region(dst_shape);
+ validate(dst.info()->valid_region(), dst_valid_region);
+}
+
+template <typename T>
+using CLFullyConnectedLayerFixture = FullyConnectedLayerValidationFixture<CLTensor, CLAccessor, CLFullyConnectedLayer, T, false>;
+
+TEST_SUITE(Float)
+TEST_SUITE(FP16)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLFullyConnectedLayerFixture<half_float::half>, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallFullyConnectedLayerDataset(),
+ FullyConnectedParameters),
+ framework::dataset::make("DataType", DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f16);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLFullyConnectedLayerFixture<half_float::half>, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeFullyConnectedLayerDataset(),
+ FullyConnectedParameters),
+ framework::dataset::make("DataType", DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f16);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLFullyConnectedLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallFullyConnectedLayerDataset(), FullyConnectedParameters),
+ framework::dataset::make("DataType", DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLFullyConnectedLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeFullyConnectedLayerDataset(), FullyConnectedParameters),
+ framework::dataset::make("DataType", DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+template <typename T>
+using CLFullyConnectedLayerFixedPointFixture = FullyConnectedLayerValidationFixedPointFixture<CLTensor, CLAccessor, CLFullyConnectedLayer, T, false>;
+
+TEST_SUITE(Quantized)
+TEST_SUITE(QS8)
+// Testing for fixed point position [1,6) as reciprocal limits the maximum fixed point position to 5
+FIXTURE_DATA_TEST_CASE(RunSmall, CLFullyConnectedLayerFixedPointFixture<int8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallFullyConnectedLayerDataset(),
+ FullyConnectedParameters),
+ framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FractionalBits", 1, 6)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_fixed_point);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLFullyConnectedLayerFixedPointFixture<int8_t>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeFullyConnectedLayerDataset(),
+ FullyConnectedParameters),
+ framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FractionalBits", 1, 6)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_fixed_point);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(QS16)
+// Testing for fixed point position [1,14) as reciprocal limits the maximum fixed point position to 14
+FIXTURE_DATA_TEST_CASE(RunSmall, CLFullyConnectedLayerFixedPointFixture<int16_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallFullyConnectedLayerDataset(),
+ FullyConnectedParameters),
+ framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FractionalBits", 1, 14)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_fixed_point);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLFullyConnectedLayerFixedPointFixture<int16_t>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeFullyConnectedLayerDataset(),
+ FullyConnectedParameters),
+ framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FractionalBits", 1, 14)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_fixed_point);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/GEMM.cpp b/tests/validation/CL/GEMM.cpp
new file mode 100644
index 0000000000..6b2b2b41b1
--- /dev/null
+++ b/tests/validation/CL/GEMM.cpp
@@ -0,0 +1,167 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLGEMM.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/LargeGEMMDataset.h"
+#include "tests/datasets/SmallGEMMDataset.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/GEMMFixture.h"
+#include "tests/validation/half.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+constexpr AbsoluteTolerance<float> tolerance_f32(0.001f); /**< Tolerance value for comparing reference's output against implementation's output for floating point data types */
+constexpr AbsoluteTolerance<float> tolerance_f16(0.1f); /**< Tolerance value for comparing reference's output against implementation's output for floating point data types */
+constexpr AbsoluteTolerance<float> tolerance_q(1.0f); /**< Tolerance value for comparing reference's output against implementation's output for fixed point data types */
+
+/** CNN data types */
+const auto CNNDataTypes = framework::dataset::make("DataType",
+{
+ DataType::F16,
+ DataType::F32,
+ DataType::QS8,
+ DataType::QS16,
+});
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(GEMM)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(framework::dataset::concat(datasets::SmallGEMMDataset(), datasets::LargeGEMMDataset()), CNNDataTypes),
+ shape_a, shape_b, shape_c, output_shape, alpha, beta, data_type)
+{
+ // Set fixed point position data type allowed
+ const int fixed_point_position = is_data_type_fixed_point(data_type) ? 3 : 0;
+
+ // Create tensors
+ CLTensor a = create_tensor<CLTensor>(shape_a, data_type, 1, fixed_point_position);
+ CLTensor b = create_tensor<CLTensor>(shape_b, data_type, 1, fixed_point_position);
+ CLTensor c = create_tensor<CLTensor>(shape_c, data_type, 1, fixed_point_position);
+ CLTensor dst = create_tensor<CLTensor>(output_shape, data_type, 1, fixed_point_position);
+
+ ARM_COMPUTE_EXPECT(a.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(b.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(c.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Create and configure function
+ CLGEMM gemm;
+ gemm.configure(&a, &b, &c, &dst, alpha, beta);
+
+ //TODO(COMPMID-415): Validate valid region
+}
+
+template <typename T>
+using CLGEMMFixture = GEMMValidationFixture<CLTensor, CLAccessor, CLGEMM, T>;
+
+TEST_SUITE(Float)
+TEST_SUITE(FP16)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLGEMMFixture<half_float::half>, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallGEMMDataset(), framework::dataset::make("DataType", DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f16);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLGEMMFixture<half_float::half>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeGEMMDataset(), framework::dataset::make("DataType",
+ DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f16);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLGEMMFixture<float>, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallGEMMDataset(), framework::dataset::make("DataType", DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLGEMMFixture<float>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeGEMMDataset(), framework::dataset::make("DataType", DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+template <typename T>
+using CLGEMMFixedPointFixture = GEMMValidationFixedPointFixture<CLTensor, CLAccessor, CLGEMM, T>;
+
+TEST_SUITE(Quantized)
+TEST_SUITE(QS8)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLGEMMFixedPointFixture<int8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallGEMMDataset(),
+ framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FractionalBits", 1, 7)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_q);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLGEMMFixedPointFixture<int8_t>, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeGEMMDataset(),
+ framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FractionalBits", 1, 7)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_q);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(QS16)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLGEMMFixedPointFixture<int16_t>, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallGEMMDataset(),
+ framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FractionalBits", 1, 14)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_q);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLGEMMFixedPointFixture<int16_t>, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeGEMMDataset(),
+ framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FractionalBits", 1, 14)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_q);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/Gaussian3x3.cpp b/tests/validation/CL/Gaussian3x3.cpp
deleted file mode 100644
index 2ef077c005..0000000000
--- a/tests/validation/CL/Gaussian3x3.cpp
+++ /dev/null
@@ -1,166 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "PaddingCalculator.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-#include "validation/ValidationUserConfiguration.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/CLTensor.h"
-#include "arm_compute/runtime/CL/CLTensorAllocator.h"
-#include "arm_compute/runtime/CL/functions/CLGaussian3x3.h"
-
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-constexpr unsigned int filter_size = 3; /** Size of the kernel/filter in number of elements. */
-constexpr BorderSize border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */
-
-/** Compute CL gaussian3x3 filter.
- *
- * @param[in] shape Shape of the input and output tensors.
- * @param[in] border_mode BorderMode used by the input tensor.
- * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT.
- *
- * @return Computed output tensor.
- */
-CLTensor compute_gaussian3x3(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U8);
-
- // Create and configure function
- CLGaussian3x3 gaussian3x3;
- gaussian3x3.configure(&src, &dst, border_mode, constant_border_value);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src), 0);
-
- // Compute function
- gaussian3x3.run();
-
- return dst;
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(Gaussian3x3)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U8);
-
- BOOST_TEST(src.info()->is_resizable());
- BOOST_TEST(dst.info()->is_resizable());
-
- // Create and configure function
- CLGaussian3x3 gaussian3x3;
- gaussian3x3.configure(&src, &dst, border_mode);
-
- // Validate valid region
- const ValidRegion src_valid_region = shape_to_valid_region(shape);
- const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
- validate(src.info()->valid_region(), src_valid_region);
- validate(dst.info()->valid_region(), dst_valid_region);
-
- // Validate padding
- PaddingCalculator calculator(shape.x(), 8);
- calculator.set_border_size(1);
- calculator.set_border_mode(border_mode);
-
- const PaddingSize dst_padding = calculator.required_padding();
-
- calculator.set_accessed_elements(16);
- calculator.set_access_offset(-1);
-
- const PaddingSize src_padding = calculator.required_padding();
-
- validate(src.info()->padding(), src_padding);
- validate(dst.info()->padding(), dst_padding);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode)
-{
- std::mt19937 gen(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution(0, 255);
- const uint8_t border_value = distribution(gen);
-
- // Compute function
- CLTensor dst = compute_gaussian3x3(shape, border_mode, border_value);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_gaussian3x3(shape, border_mode, border_value);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size));
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode)
-{
- std::mt19937 gen(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution(0, 255);
- const uint8_t border_value = distribution(gen);
-
- // Compute function
- CLTensor dst = compute_gaussian3x3(shape, border_mode, border_value);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_gaussian3x3(shape, border_mode, border_value);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size));
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/Gaussian5x5.cpp b/tests/validation/CL/Gaussian5x5.cpp
deleted file mode 100644
index fb21ed06af..0000000000
--- a/tests/validation/CL/Gaussian5x5.cpp
+++ /dev/null
@@ -1,166 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "PaddingCalculator.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-#include "validation/ValidationUserConfiguration.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/CLTensor.h"
-#include "arm_compute/runtime/CL/CLTensorAllocator.h"
-#include "arm_compute/runtime/CL/functions/CLGaussian5x5.h"
-
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-constexpr unsigned int filter_size = 5; /** Size of the kernel/filter in number of elements. */
-constexpr BorderSize border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */
-
-/** Compute CL gaussian5x5 filter.
- *
- * @param[in] shape Shape of the input and output tensors.
- * @param[in] border_mode BorderMode used by the input tensor.
- * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT.
- *
- * @return Computed output tensor.
- */
-CLTensor compute_gaussian5x5(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U8);
-
- // Create and configure function
- CLGaussian5x5 gaussian5x5;
- gaussian5x5.configure(&src, &dst, border_mode, constant_border_value);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src), 0);
-
- // Compute function
- gaussian5x5.run();
-
- return dst;
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(Gaussian5x5)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U8);
-
- BOOST_TEST(src.info()->is_resizable());
- BOOST_TEST(dst.info()->is_resizable());
-
- // Create and configure function
- CLGaussian5x5 gaussian5x5;
- gaussian5x5.configure(&src, &dst, border_mode);
-
- // Validate valid region
- const ValidRegion src_valid_region = shape_to_valid_region(shape);
- const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
- validate(src.info()->valid_region(), src_valid_region);
- validate(dst.info()->valid_region(), dst_valid_region);
-
- // Validate padding
- PaddingCalculator calculator(shape.x(), 8);
- calculator.set_border_size(2);
- calculator.set_border_mode(border_mode);
-
- const PaddingSize dst_padding = calculator.required_padding();
-
- calculator.set_accessed_elements(16);
- calculator.set_access_offset(-2);
-
- const PaddingSize src_padding = calculator.required_padding();
-
- validate(src.info()->padding(), src_padding);
- validate(dst.info()->padding(), dst_padding);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode)
-{
- std::mt19937 gen(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution(0, 255);
- const uint8_t border_value = distribution(gen);
-
- // Compute function
- CLTensor dst = compute_gaussian5x5(shape, border_mode, border_value);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_gaussian5x5(shape, border_mode, border_value);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size));
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode)
-{
- std::mt19937 gen(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution(0, 255);
- const uint8_t border_value = distribution(gen);
-
- // Compute function
- CLTensor dst = compute_gaussian5x5(shape, border_mode, border_value);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_gaussian5x5(shape, border_mode, border_value);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size));
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/HarrisCorners.cpp b/tests/validation/CL/HarrisCorners.cpp
deleted file mode 100644
index 6370c4d972..0000000000
--- a/tests/validation/CL/HarrisCorners.cpp
+++ /dev/null
@@ -1,224 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-#include "validation/ValidationUserConfiguration.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/CLArray.h"
-#include "arm_compute/runtime/CL/functions/CLHarrisCorners.h"
-#include "arm_compute/runtime/Tensor.h"
-#include "arm_compute/runtime/TensorAllocator.h"
-
-#include "PaddingCalculator.h"
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-/** Compute CL Harris corners function.
- *
- * @param[in] shape Shape of input tensor
- * @param[in] threshold Minimum threshold with which to eliminate Harris Corner scores (computed using the normalized Sobel kernel).
- * @param[in] min_dist Radial Euclidean distance for the euclidean distance stage
- * @param[in] sensitivity Sensitivity threshold k from the Harris-Stephens equation
- * @param[in] gradient_size The gradient window size to use on the input. The implementation supports 3, 5, and 7
- * @param[in] block_size The block window size used to compute the Harris Corner score. The implementation supports 3, 5, and 7.
- * @param[in] border_mode Border mode to use
- * @param[in] constant_border_value Constant value to use for borders if border_mode is set to CONSTANT.
- *
- * @return Computed corners' keypoints.
- */
-void compute_harris_corners(const TensorShape &shape, CLKeyPointArray &corners, float threshold, float min_dist, float sensitivity,
- int32_t gradient_size, int32_t block_size, BorderMode border_mode, uint8_t constant_border_value)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- src.info()->set_format(Format::U8);
-
- // Create harris corners configure function
- CLHarrisCorners harris_corners;
- harris_corners.configure(&src, threshold, min_dist, sensitivity, gradient_size, block_size, &corners, border_mode, constant_border_value);
-
- // Allocate tensors
- src.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src), 0);
-
- // Compute function
- harris_corners.run();
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(HarrisCorners)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (Small2DShapes() + Large2DShapes()) * BorderModes()
- * boost::unit_test::data::make({ 3, 5, 7 }) * boost::unit_test::data::make({ 3, 5, 7 }),
- shape, border_mode, gradient, block)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- src.info()->set_format(Format::U8);
-
- CLKeyPointArray corners(shape.total_size());
-
- uint8_t constant_border_value = 0;
-
- std::mt19937 gen(user_config.seed.get());
- std::uniform_real_distribution<float> real_dist(0.01, std::numeric_limits<float>::min());
-
- const float threshold = real_dist(gen);
- const float sensitivity = real_dist(gen);
- const float max_euclidean_distance = 30.f;
-
- real_dist = std::uniform_real_distribution<float>(0.f, max_euclidean_distance);
- float min_dist = real_dist(gen);
-
- // Generate a random constant value if border_mode is constant
- if(border_mode == BorderMode::CONSTANT)
- {
- std::uniform_int_distribution<uint8_t> int_dist(0, 255);
- constant_border_value = int_dist(gen);
- }
-
- BOOST_TEST(src.info()->is_resizable());
-
- // Create harris corners configure function
- CLHarrisCorners harris_corners;
- harris_corners.configure(&src, threshold, min_dist, sensitivity, gradient, block, &corners, border_mode, constant_border_value);
-
- // Validate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape);
-
- validate(src.info()->valid_region(), valid_region);
-
- // Validate padding
- PaddingCalculator calculator(shape.x(), 8);
-
- calculator.set_border_mode(border_mode);
- calculator.set_border_size(gradient / 2);
- calculator.set_access_offset(-gradient / 2);
- calculator.set_accessed_elements(16);
-
- const PaddingSize padding = calculator.required_padding();
-
- validate(src.info()->padding(), padding);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes() * BorderModes() * boost::unit_test::data::make({ 3, 5, 7 }) * boost::unit_test::data::make({ 3, 5, 7 }), shape, border_mode, gradient, block)
-{
- uint8_t constant_border_value = 0;
-
- std::mt19937 gen(user_config.seed.get());
- std::uniform_real_distribution<float> real_dist(0.01, std::numeric_limits<float>::min());
-
- const float threshold = real_dist(gen);
- const float sensitivity = real_dist(gen);
- const float max_euclidean_distance = 30.f;
-
- real_dist = std::uniform_real_distribution<float>(0.f, max_euclidean_distance);
- const float min_dist = real_dist(gen);
-
- // Generate a random constant value if border_mode is constant
- if(border_mode == BorderMode::CONSTANT)
- {
- std::uniform_int_distribution<uint8_t> int_dist(0, 255);
- constant_border_value = int_dist(gen);
- }
-
- // Create array of keypoints
- CLKeyPointArray dst(shape.total_size());
-
- // Compute function
- compute_harris_corners(shape, dst, threshold, min_dist, sensitivity, gradient, block, border_mode, constant_border_value);
-
- // Compute reference
- KeyPointArray ref_dst = Reference::compute_reference_harris_corners(shape, threshold, min_dist, sensitivity, gradient, block, border_mode, constant_border_value);
-
- // Validate output
- dst.map();
- validate(dst, ref_dst, 1);
- dst.unmap();
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, Large2DShapes() * BorderModes() * boost::unit_test::data::make({ 3, 5, 7 }) * boost::unit_test::data::make({ 3, 5, 7 }), shape, border_mode, gradient, block)
-{
- uint8_t constant_border_value = 0;
-
- std::mt19937 gen(user_config.seed.get());
- std::uniform_real_distribution<float> real_dist(0.01, std::numeric_limits<float>::min());
-
- const float threshold = real_dist(gen);
- const float sensitivity = real_dist(gen);
- const float max_euclidean_distance = 30.f;
-
- real_dist = std::uniform_real_distribution<float>(0.f, max_euclidean_distance);
- const float min_dist = real_dist(gen);
-
- // Generate a random constant value if border_mode is constant
- if(border_mode == BorderMode::CONSTANT)
- {
- std::uniform_int_distribution<uint8_t> int_dist(0, 255);
- constant_border_value = int_dist(gen);
- }
-
- // Create array of keypoints
- CLKeyPointArray dst(shape.total_size());
-
- // Compute function
- compute_harris_corners(shape, dst, threshold, min_dist, sensitivity, gradient, block, border_mode, constant_border_value);
-
- // Compute reference
- KeyPointArray ref_dst = Reference::compute_reference_harris_corners(shape, threshold, min_dist, sensitivity, gradient, block, border_mode, constant_border_value);
-
- // Validate output
- dst.map();
- validate(dst, ref_dst);
- dst.unmap();
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/IntegralImage.cpp b/tests/validation/CL/IntegralImage.cpp
deleted file mode 100644
index dc325a1f8a..0000000000
--- a/tests/validation/CL/IntegralImage.cpp
+++ /dev/null
@@ -1,142 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "PaddingCalculator.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/functions/CLIntegralImage.h"
-#include "arm_compute/runtime/Tensor.h"
-#include "arm_compute/runtime/TensorAllocator.h"
-
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-/** Compute CL integral image function.
- *
- * @param[in] shape Shape of the input and output tensors.
- *
- * @return Computed output tensor.
- */
-CLTensor compute_integral_image(const TensorShape &shape)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U32);
-
- // Create integral image configure function
- CLIntegralImage integral_image;
- integral_image.configure(&src, &dst);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src), 0);
-
- // Compute function
- integral_image.run();
-
- return dst;
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(IntegralImage)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, SmallShapes() + LargeShapes(), shape)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U32);
-
- BOOST_TEST(src.info()->is_resizable());
- BOOST_TEST(dst.info()->is_resizable());
-
- // Create integral image configure function
- CLIntegralImage integral_image;
- integral_image.configure(&src, &dst);
-
- // Validate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape);
- validate(src.info()->valid_region(), valid_region);
- validate(dst.info()->valid_region(), valid_region);
-
- // Validate padding
- const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
- validate(src.info()->padding(), padding);
- validate(dst.info()->padding(), padding);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), shape)
-{
- // Compute function
- CLTensor dst = compute_integral_image(shape);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_integral_image(shape);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes(), shape)
-{
- // Compute function
- CLTensor dst = compute_integral_image(shape);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_integral_image(shape);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/L2Normalize.cpp b/tests/validation/CL/L2Normalize.cpp
new file mode 100644
index 0000000000..bd9bf17322
--- /dev/null
+++ b/tests/validation/CL/L2Normalize.cpp
@@ -0,0 +1,78 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLL2Normalize.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/L2NormalizeFixture.h"
+#include "tests/validation/half.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Tolerance for float operations */
+constexpr AbsoluteTolerance<float> tolerance_f32(0.00001f);
+
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(L2Normalize)
+
+template <typename T>
+using CLL2NormalizeFixture = L2NormalizeValidationFixture<CLTensor, CLAccessor, CLL2Normalize, T>;
+
+TEST_SUITE(Float)
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLL2NormalizeFixture<float>, framework::DatasetMode::PRECOMMIT,
+ combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Axis", { 0 })), framework::dataset::make("Epsilon", { 1e-12 })))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLL2NormalizeFixture<float>, framework::DatasetMode::NIGHTLY,
+ combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Axis", { 0 })), framework::dataset::make("Epsilon", { 1e-12 })))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/MeanStdDev.cpp b/tests/validation/CL/MeanStdDev.cpp
new file mode 100644
index 0000000000..ff8a087c6b
--- /dev/null
+++ b/tests/validation/CL/MeanStdDev.cpp
@@ -0,0 +1,97 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/runtime/CL/functions/CLMeanStdDev.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Macros.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/MeanStdDevFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+constexpr RelativeTolerance tolerance_rel_high_error(0.05f);
+constexpr RelativeTolerance tolerance_rel_low_error(0.0005f);
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(MeanStdDev)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(concat(datasets::Small2DShapes(), datasets::Large2DShapes()), framework::dataset::make("DataType", DataType::U8)), shape, data_type)
+{
+ // Create tensors
+ CLTensor src = create_tensor<CLTensor>(shape, data_type);
+
+ // Create output variables
+ float mean = 0.f;
+ float std_dev = 0.f;
+
+ ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Create configure function
+ CLMeanStdDev mean_std_dev_image;
+ mean_std_dev_image.configure(&src, &mean, &std_dev);
+
+ // Validate valid region
+ const ValidRegion valid_region = shape_to_valid_region(shape);
+ validate(src.info()->valid_region(), valid_region);
+
+ // Validate padding
+ const PaddingSize padding = PaddingCalculator(shape.x(), 8).required_padding();
+ validate(src.info()->padding(), padding);
+}
+
+template <typename T>
+using CLMeanStdDevFixture = MeanStdDevValidationFixture<CLTensor, CLAccessor, CLMeanStdDev, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, CLMeanStdDevFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(datasets::Small2DShapes(), framework::dataset::make("DataType",
+ DataType::U8)))
+{
+ // Validate mean output
+ validate(_target.first, _reference.first);
+
+ // Validate std_dev output
+ validate(_target.second, _reference.second, tolerance_rel_high_error);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLMeanStdDevFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(datasets::Large2DShapes(), framework::dataset::make("DataType",
+ DataType::U8)))
+{
+ // Validate mean output
+ validate(_target.first, _reference.first, tolerance_rel_low_error);
+
+ // Validate std_dev output
+ validate(_target.second, _reference.second, tolerance_rel_high_error);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/MinMaxLocation.cpp b/tests/validation/CL/MinMaxLocation.cpp
deleted file mode 100644
index 5f8be433cd..0000000000
--- a/tests/validation/CL/MinMaxLocation.cpp
+++ /dev/null
@@ -1,397 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/functions/CLMinMaxLocation.h"
-#include "arm_compute/runtime/Tensor.h"
-#include "arm_compute/runtime/TensorAllocator.h"
-
-#include "PaddingCalculator.h"
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-/** Compute CL MinMaxLocation function.
- *
- * @param[in] shape Shape of the input and output tensors.
- * @param[in] dt_in Data type of first input tensor.
- * @param[out] min Minimum value of tensor
- * @param[out] max Maximum value of tensor
- * @param[out] min_loc Array with locations of minimum values
- * @param[out] max_loc Array with locations of maximum values
- * @param[out] min_count Number of minimum values found
- * @param[out] max_count Number of maximum values found
- *
- * @return Computed output tensor.
- */
-void compute_min_max_location(const TensorShape &shape, DataType dt_in, void *min, void *max,
- CLCoordinates2DArray &min_loc, CLCoordinates2DArray &max_loc, uint32_t &min_count, uint32_t &max_count)
-{
- // Create tensor
- CLTensor src = create_tensor<CLTensor>(shape, dt_in);
-
- // Create and configure min_max_location configure function
- CLMinMaxLocation min_max_loc;
- min_max_loc.configure(&src, min, max, &min_loc, &max_loc, &min_count, &max_count);
-
- // Allocate tensors
- src.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src), 0);
-
- // Compute function
- min_max_loc.run();
-}
-
-void validate_configuration(const CLTensor &src, TensorShape shape)
-{
- BOOST_TEST(src.info()->is_resizable());
-
- // Create output storage
- int32_t min;
- int32_t max;
- CLCoordinates2DArray min_loc(shape.total_size());
- CLCoordinates2DArray max_loc(shape.total_size());
- uint32_t min_count;
- uint32_t max_count;
-
- // Create and configure function
- CLMinMaxLocation min_max_loc;
- min_max_loc.configure(&src, &min, &max, &min_loc, &max_loc, &min_count, &max_count);
-
- // Validate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape);
- validate(src.info()->valid_region(), valid_region);
-
- // Validate padding
- const PaddingSize padding = PaddingCalculator(shape.x(), src.info()->dimension(0)).required_padding();
- validate(src.info()->padding(), padding);
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(MinMaxLocation)
-BOOST_AUTO_TEST_SUITE(U8)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (Small2DShapes() + Large2DShapes()),
- shape)
-{
- // Create tensor
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- src.info()->set_format(Format::U8);
-
- validate_configuration(src, shape);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes(),
- shape)
-{
- // Create output storage
- int32_t min;
- int32_t max;
- CLCoordinates2DArray min_loc(shape.total_size());
- CLCoordinates2DArray max_loc(shape.total_size());
- uint32_t min_count;
- uint32_t max_count;
-
- int32_t ref_min;
- int32_t ref_max;
- CLCoordinates2DArray ref_min_loc(shape.total_size());
- CLCoordinates2DArray ref_max_loc(shape.total_size());
- uint32_t ref_min_count;
- uint32_t ref_max_count;
-
- // Compute function
- compute_min_max_location(shape, DataType::U8, &min, &max, min_loc, max_loc, min_count, max_count);
-
- // Compute reference
- ref_min_loc.map();
- ref_max_loc.map();
-
- Reference::compute_reference_min_max_location(shape, DataType::U8, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count);
-
- min_loc.map();
- max_loc.map();
-
- // Validate output
- validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count);
-
- ref_min_loc.unmap();
- ref_max_loc.unmap();
- min_loc.unmap();
- max_loc.unmap();
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, Large2DShapes(),
- shape)
-{
- // Create output storage
- int32_t min;
- int32_t max;
- CLCoordinates2DArray min_loc(shape.total_size());
- CLCoordinates2DArray max_loc(shape.total_size());
- uint32_t min_count;
- uint32_t max_count;
-
- int32_t ref_min;
- int32_t ref_max;
- CLCoordinates2DArray ref_min_loc(shape.total_size());
- CLCoordinates2DArray ref_max_loc(shape.total_size());
- uint32_t ref_min_count;
- uint32_t ref_max_count;
-
- // Compute function
- compute_min_max_location(shape, DataType::U8, &min, &max, min_loc, max_loc, min_count, max_count);
-
- // Compute reference
- ref_min_loc.map();
- ref_max_loc.map();
-
- Reference::compute_reference_min_max_location(shape, DataType::U8, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count);
-
- min_loc.map();
- max_loc.map();
-
- // Validate output
- validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count);
-
- ref_min_loc.unmap();
- ref_max_loc.unmap();
- min_loc.unmap();
- max_loc.unmap();
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(S16)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (Small2DShapes() + Large2DShapes()),
- shape)
-{
- // Create tensor
- CLTensor src = create_tensor<CLTensor>(shape, DataType::S16);
- src.info()->set_format(Format::S16);
-
- validate_configuration(src, shape);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes(),
- shape)
-{
- // Create output storage
- int32_t min;
- int32_t max;
- CLCoordinates2DArray min_loc(shape.total_size());
- CLCoordinates2DArray max_loc(shape.total_size());
- uint32_t min_count;
- uint32_t max_count;
-
- int32_t ref_min;
- int32_t ref_max;
- CLCoordinates2DArray ref_min_loc(shape.total_size());
- CLCoordinates2DArray ref_max_loc(shape.total_size());
- uint32_t ref_min_count;
- uint32_t ref_max_count;
-
- // Compute function
- compute_min_max_location(shape, DataType::S16, &min, &max, min_loc, max_loc, min_count, max_count);
-
- // Compute reference
- ref_min_loc.map();
- ref_max_loc.map();
-
- Reference::compute_reference_min_max_location(shape, DataType::S16, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count);
-
- min_loc.map();
- max_loc.map();
-
- // Validate output
- validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count);
-
- ref_min_loc.unmap();
- ref_max_loc.unmap();
- min_loc.unmap();
- max_loc.unmap();
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, Large2DShapes(),
- shape)
-{
- // Create output storage
- int32_t min;
- int32_t max;
- CLCoordinates2DArray min_loc(shape.total_size());
- CLCoordinates2DArray max_loc(shape.total_size());
- uint32_t min_count;
- uint32_t max_count;
-
- int32_t ref_min;
- int32_t ref_max;
- CLCoordinates2DArray ref_min_loc(shape.total_size());
- CLCoordinates2DArray ref_max_loc(shape.total_size());
- uint32_t ref_min_count;
- uint32_t ref_max_count;
-
- // Compute function
- compute_min_max_location(shape, DataType::S16, &min, &max, min_loc, max_loc, min_count, max_count);
-
- // Compute reference
- ref_min_loc.map();
- ref_max_loc.map();
-
- Reference::compute_reference_min_max_location(shape, DataType::S16, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count);
-
- min_loc.map();
- max_loc.map();
-
- // Validate output
- validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count);
-
- ref_min_loc.unmap();
- ref_max_loc.unmap();
- min_loc.unmap();
- max_loc.unmap();
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(Float)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (Small2DShapes() + Large2DShapes()),
- shape)
-{
- // Create tensor
- CLTensor src = create_tensor<CLTensor>(shape, DataType::F32);
-
- validate_configuration(src, shape);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes(),
- shape)
-{
- // Create output storage
- float min;
- float max;
- CLCoordinates2DArray min_loc(shape.total_size());
- CLCoordinates2DArray max_loc(shape.total_size());
- uint32_t min_count;
- uint32_t max_count;
-
- float ref_min;
- float ref_max;
- CLCoordinates2DArray ref_min_loc(shape.total_size());
- CLCoordinates2DArray ref_max_loc(shape.total_size());
- uint32_t ref_min_count;
- uint32_t ref_max_count;
-
- // Compute function
- compute_min_max_location(shape, DataType::F32, &min, &max, min_loc, max_loc, min_count, max_count);
-
- // Compute reference
- ref_min_loc.map();
- ref_max_loc.map();
-
- Reference::compute_reference_min_max_location(shape, DataType::F32, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count);
-
- min_loc.map();
- max_loc.map();
-
- // Validate output
- validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count);
-
- ref_min_loc.unmap();
- ref_max_loc.unmap();
- min_loc.unmap();
- max_loc.unmap();
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, Large2DShapes(),
- shape)
-{
- // Create output storage
- float min;
- float max;
- CLCoordinates2DArray min_loc(shape.total_size());
- CLCoordinates2DArray max_loc(shape.total_size());
- uint32_t min_count;
- uint32_t max_count;
-
- float ref_min;
- float ref_max;
- CLCoordinates2DArray ref_min_loc(shape.total_size());
- CLCoordinates2DArray ref_max_loc(shape.total_size());
- uint32_t ref_min_count;
- uint32_t ref_max_count;
-
- // Compute function
- compute_min_max_location(shape, DataType::F32, &min, &max, min_loc, max_loc, min_count, max_count);
-
- // Compute reference
- ref_min_loc.map();
- ref_max_loc.map();
-
- Reference::compute_reference_min_max_location(shape, DataType::F32, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count);
-
- min_loc.map();
- max_loc.map();
-
- // Validate output
- validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count);
-
- ref_min_loc.unmap();
- ref_max_loc.unmap();
- min_loc.unmap();
- max_loc.unmap();
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/NonLinearFilter.cpp b/tests/validation/CL/NonLinearFilter.cpp
deleted file mode 100644
index f453f27e8a..0000000000
--- a/tests/validation/CL/NonLinearFilter.cpp
+++ /dev/null
@@ -1,203 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "PaddingCalculator.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Helpers.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-#include "validation/ValidationUserConfiguration.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/functions/CLNonLinearFilter.h"
-#include "arm_compute/runtime/Tensor.h"
-#include "arm_compute/runtime/TensorAllocator.h"
-
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-/** Compute NonLinearFilter function.
- *
- * @param[in] input Shape of the input and output tensors.
- * @param[in] function Non linear function to perform
- * @param[in] mask_size Mask size. Supported sizes: 3, 5
- * @param[in] pattern Mask pattern
- * @param[in] mask The given mask. Will be used only if pattern is specified to PATTERN_OTHER
- * @param[in] border_mode Strategy to use for borders.
- * @param[in] constant_border_value (Optional) Constant value to use for borders if border_mode is set to CONSTANT.
- *
- * @return Computed output CL tensor.
- */
-CLTensor compute_non_linear_filter(const TensorShape &shape, NonLinearFilterFunction function, unsigned int mask_size,
- MatrixPattern pattern, const uint8_t *mask, BorderMode border_mode,
- uint8_t constant_border_value)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U8);
-
- // Create and configure function
- CLNonLinearFilter filter;
- filter.configure(&src, &dst, function, mask_size, pattern, mask, border_mode, constant_border_value);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src), 0);
-
- // Compute function
- filter.run();
-
- return dst;
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(NonLinearFilter)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes())
- * NonLinearFilterFunctions() * boost::unit_test::data::make({ 3U, 5U })
- * boost::unit_test::data::make({ MatrixPattern::BOX, MatrixPattern::CROSS, MatrixPattern::DISK }) * BorderModes(),
- shape, function, mask_size, pattern, border_mode)
-{
- std::mt19937 generator(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
- const uint8_t constant_border_value = distribution_u8(generator);
-
- // Create the mask
- uint8_t mask[mask_size * mask_size];
- fill_mask_from_pattern(mask, mask_size, mask_size, pattern);
- const auto half_mask_size = static_cast<int>(mask_size / 2);
-
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U8);
-
- BOOST_TEST(src.info()->is_resizable());
- BOOST_TEST(dst.info()->is_resizable());
-
- // Create and configure function
- CLNonLinearFilter filter;
- filter.configure(&src, &dst, function, mask_size, pattern, mask, border_mode, constant_border_value);
-
- // Validate valid region
- const ValidRegion src_valid_region = shape_to_valid_region(shape);
- const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, BorderSize(half_mask_size));
-
- validate(src.info()->valid_region(), src_valid_region);
- validate(dst.info()->valid_region(), dst_valid_region);
-
- // Validate padding
- PaddingCalculator calculator(shape.x(), ((MatrixPattern::OTHER == pattern) ? 1 : 8));
- calculator.set_border_mode(border_mode);
- calculator.set_border_size(half_mask_size);
-
- const PaddingSize write_padding = calculator.required_padding(PaddingCalculator::Option::EXCLUDE_BORDER);
-
- calculator.set_accessed_elements(16);
- calculator.set_access_offset(-half_mask_size);
-
- const PaddingSize read_padding = calculator.required_padding(PaddingCalculator::Option::INCLUDE_BORDER);
-
- validate(src.info()->padding(), read_padding);
- validate(dst.info()->padding(), write_padding);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes()
- * NonLinearFilterFunctions() * boost::unit_test::data::make({ 3U, 5U })
- * boost::unit_test::data::make({ MatrixPattern::BOX, MatrixPattern::CROSS, MatrixPattern::DISK }) * BorderModes(),
- shape, function, mask_size, pattern, border_mode)
-{
- std::mt19937 generator(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
- const uint8_t constant_border_value = distribution_u8(generator);
-
- // Create the mask
- uint8_t mask[mask_size * mask_size];
- fill_mask_from_pattern(mask, mask_size, mask_size, pattern);
-
- // Compute function
- CLTensor dst = compute_non_linear_filter(shape, function, mask_size, pattern, mask, border_mode, constant_border_value);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_non_linear_filter(shape, function, mask_size, pattern, mask, border_mode, constant_border_value);
-
- // Calculate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, BorderSize(static_cast<int>(mask_size / 2)));
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, valid_region);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes()
- * NonLinearFilterFunctions() * boost::unit_test::data::make({ 3U, 5U })
- * boost::unit_test::data::make({ MatrixPattern::BOX, MatrixPattern::CROSS, MatrixPattern::DISK }) * BorderModes(),
- shape, function, mask_size, pattern, border_mode)
-{
- std::mt19937 generator(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
- const uint8_t constant_border_value = distribution_u8(generator);
-
- // Create the mask
- uint8_t mask[mask_size * mask_size];
- fill_mask_from_pattern(mask, mask_size, mask_size, pattern);
-
- // Compute function
- CLTensor dst = compute_non_linear_filter(shape, function, mask_size, pattern, mask, border_mode, constant_border_value);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_non_linear_filter(shape, function, mask_size, pattern, mask, border_mode, constant_border_value);
-
- // Calculate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, BorderSize(static_cast<int>(mask_size / 2)));
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, valid_region);
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/NormalizationLayer.cpp b/tests/validation/CL/NormalizationLayer.cpp
new file mode 100644
index 0000000000..4d14649a91
--- /dev/null
+++ b/tests/validation/CL/NormalizationLayer.cpp
@@ -0,0 +1,141 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLNormalizationLayer.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/NormalizationTypesDataset.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/NormalizationLayerFixture.h"
+#include "tests/validation/half.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Tolerance for float operations */
+constexpr AbsoluteTolerance<float> tolerance_f16(0.01f);
+constexpr AbsoluteTolerance<float> tolerance_f32(0.00001f);
+/** Tolerance for fixed point operations */
+constexpr AbsoluteTolerance<int8_t> tolerance_qs8(2);
+constexpr AbsoluteTolerance<int16_t> tolerance_qs16(2);
+
+/** Input data set. */
+const auto NormalizationDataset = combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("NormType", { NormType::IN_MAP_1D, NormType::CROSS_MAP })),
+ framework::dataset::make("NormalizationSize", 3, 9, 2)),
+ framework::dataset::make("Beta", { 0.5f, 1.f, 2.f }));
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(NormalizationLayer)
+
+//TODO(COMPMID-415): Missing configuration?
+
+template <typename T>
+using CLNormalizationLayerFixture = NormalizationValidationFixture<CLTensor, CLAccessor, CLNormalizationLayer, T>;
+
+TEST_SUITE(Float)
+TEST_SUITE(FP16)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLNormalizationLayerFixture<half_float::half>, framework::DatasetMode::PRECOMMIT, combine(NormalizationDataset, framework::dataset::make("DataType", DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f16);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLNormalizationLayerFixture<half_float::half>, framework::DatasetMode::NIGHTLY, combine(NormalizationDataset, framework::dataset::make("DataType", DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f16);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLNormalizationLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(NormalizationDataset, framework::dataset::make("DataType", DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLNormalizationLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(NormalizationDataset, framework::dataset::make("DataType", DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+template <typename T>
+using CLNormalizationLayerFixedPointFixture = NormalizationValidationFixedPointFixture<CLTensor, CLAccessor, CLNormalizationLayer, T>;
+
+TEST_SUITE(Quantized)
+TEST_SUITE(QS8)
+// Testing for fixed point position [1,6) as reciprocal limits the maximum fixed point position to 5
+FIXTURE_DATA_TEST_CASE(RunSmall, CLNormalizationLayerFixedPointFixture<int8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(NormalizationDataset, framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FractionalBits", 1, 6)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_qs8);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLNormalizationLayerFixedPointFixture<int8_t>, framework::DatasetMode::NIGHTLY, combine(combine(NormalizationDataset, framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FractionalBits", 1, 6)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_qs8);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(QS16)
+// Testing for fixed point position [1,14) as reciprocal limits the maximum fixed point position to 5
+FIXTURE_DATA_TEST_CASE(RunSmall, CLNormalizationLayerFixedPointFixture<int16_t>, framework::DatasetMode::PRECOMMIT, combine(combine(NormalizationDataset, framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FractionalBits", 1, 14)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_qs16);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLNormalizationLayerFixedPointFixture<int16_t>, framework::DatasetMode::NIGHTLY, combine(combine(NormalizationDataset, framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FractionalBits", 1, 14)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_qs16);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/PixelWiseMultiplication.cpp b/tests/validation/CL/PixelWiseMultiplication.cpp
deleted file mode 100644
index 375c77dedf..0000000000
--- a/tests/validation/CL/PixelWiseMultiplication.cpp
+++ /dev/null
@@ -1,164 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "CL/CLAccessor.h"
-#include "TypePrinter.h"
-#include "tests/Globals.h"
-#include "tests/Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-
-#include "arm_compute/runtime/CL/functions/CLPixelWiseMultiplication.h"
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-const float tolerance_f32 = 1.f; /**< Tolerance value for comparing reference's output against implementation's output for float input */
-const float tolerance_f16 = 1.f; /**< Tolerance value for comparing reference's output against implementation's output for float input */
-
-/** Compute CL pixel-wise multiplication function.
- *
- * @param[in] shape Shape of the input and output tensors.
- * @param[in] dt_in0 Data type of first input tensor.
- * @param[in] dt_in1 Data type of second input tensor.
- * @param[in] dt_out Data type of the output tensor.
- * @param[in] scale Non-negative scale.
- * @param[in] convert_policy Overflow policy of the operation.
- * @param[in] rounding_policy Rounding policy of the operation.
- * @param[in] fixed_point_position (Optional) Fixed point position that expresses the number of bits for the fractional part of the number.
- *
- * @return Computed output tensor.
- */
-CLTensor compute_pixel_wise_multiplication(const TensorShape &shape, DataType dt_in0, DataType dt_in1, DataType dt_out, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy,
- int fixed_point_position = 0)
-{
- // Create tensors
- CLTensor src1 = create_tensor<CLTensor>(shape, dt_in0, 1, fixed_point_position);
- CLTensor src2 = create_tensor<CLTensor>(shape, dt_in1, 1, fixed_point_position);
- CLTensor dst = create_tensor<CLTensor>(shape, dt_out, 1, fixed_point_position);
-
- // Create and configure function
- CLPixelWiseMultiplication multiply;
- multiply.configure(&src1, &src2, &dst, scale, convert_policy, rounding_policy);
-
- // Allocate tensors
- src1.allocator()->allocate();
- src2.allocator()->allocate();
- dst.allocator()->allocate();
-
- BOOST_TEST(!src1.info()->is_resizable());
- BOOST_TEST(!src2.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src1), 0);
- library->fill_tensor_uniform(CLAccessor(src2), 1);
-
- // Compute function
- multiply.run();
-
- return dst;
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(PixelWiseMultiplication)
-
-BOOST_AUTO_TEST_SUITE(Float16)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * DataType::F16 *ConvertPolicies() * RoundingPolicy::TO_NEAREST_UP,
- shape, dt, convert_policy, rounding_policy)
-{
- constexpr float scale = 1.f / 255.f;
-
- // Compute function
- CLTensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, tolerance_f16);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(Float)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * DataType::F32 *ConvertPolicies() * RoundingPolicy::TO_NEAREST_UP,
- shape, dt, convert_policy, rounding_policy)
-{
- constexpr float scale = 1.f / 255.f;
-
- // Compute function
- CLTensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, tolerance_f32);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(Quantized)
-BOOST_AUTO_TEST_SUITE(QS8)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * DataType::QS8 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange<int>(1, 7),
- shape, dt, convert_policy, rounding_policy, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, convert_policy, rounding_policy, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, fixed_point_position, convert_policy, rounding_policy);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE(QS16)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * DataType::QS16 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange<int>(1, 15),
- shape, dt, convert_policy, rounding_policy, fixed_point_position)
-{
- // Compute function
- CLTensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, convert_policy, rounding_policy, fixed_point_position);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, fixed_point_position, convert_policy, rounding_policy);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif // DOXYGEN_SKIP_THIS
diff --git a/tests/validation/CL/PoolingLayer.cpp b/tests/validation/CL/PoolingLayer.cpp
new file mode 100644
index 0000000000..24380cb1f0
--- /dev/null
+++ b/tests/validation/CL/PoolingLayer.cpp
@@ -0,0 +1,142 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLPoolingLayer.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/PoolingTypesDataset.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/PoolingLayerFixture.h"
+#include "tests/validation/half.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Input data set for float data types */
+const auto PoolingLayerDatasetFP = combine(combine(datasets::PoolingTypes(), framework::dataset::make("PoolingSize", { 2, 3, 7 })),
+ framework::dataset::make("PadStride", { PadStrideInfo(1, 1, 0, 0), PadStrideInfo(2, 1, 0, 0), PadStrideInfo(1, 2, 1, 1), PadStrideInfo(2, 2, 1, 0) }));
+
+/** Input data set for quantized data types */
+const auto PoolingLayerDatasetQS = combine(combine(datasets::PoolingTypes(), framework::dataset::make("PoolingSize", { 2, 3 })),
+ framework::dataset::make("PadStride", { PadStrideInfo(1, 1, 0, 0), PadStrideInfo(2, 1, 0, 0), PadStrideInfo(1, 2, 1, 1), PadStrideInfo(2, 2, 1, 0) }));
+
+constexpr AbsoluteTolerance<float> tolerance_f32(0.001f); /**< Tolerance value for comparing reference's output against implementation's output for float types */
+constexpr AbsoluteTolerance<float> tolerance_f16(0.01f); /**< Tolerance value for comparing reference's output against implementation's output for float types */
+constexpr AbsoluteTolerance<float> tolerance_qs8(3); /**< Tolerance value for comparing reference's output against implementation's output for quantized input */
+constexpr AbsoluteTolerance<float> tolerance_qs16(6); /**< Tolerance value for comparing reference's output against implementation's output for quantized input */
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(PoolingLayer)
+
+template <typename T>
+using CLPoolingLayerFixture = PoolingLayerValidationFixture<CLTensor, CLAccessor, CLPoolingLayer, T>;
+
+TEST_SUITE(Float)
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLPoolingLayerFixture<float>, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), combine(PoolingLayerDatasetFP, framework::dataset::make("DataType",
+ DataType::F32))))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLPoolingLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapes(), combine(PoolingLayerDatasetFP, framework::dataset::make("DataType",
+ DataType::F32))))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(FP16)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLPoolingLayerFixture<half_float::half>, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), combine(PoolingLayerDatasetFP,
+ framework::dataset::make("DataType", DataType::F16))))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f16);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLPoolingLayerFixture<half_float::half>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapes(), combine(PoolingLayerDatasetFP,
+ framework::dataset::make("DataType", DataType::F16))))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f16);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+template <typename T>
+using CLPoolingLayerFixedPointFixture = PoolingLayerValidationFixedPointFixture<CLTensor, CLAccessor, CLPoolingLayer, T>;
+
+TEST_SUITE(Quantized)
+TEST_SUITE(QS8)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLPoolingLayerFixedPointFixture<int8_t>, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), combine(PoolingLayerDatasetQS,
+ framework::dataset::make("DataType", DataType::QS8))),
+ framework::dataset::make("FractionalBits", 1, 4)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_qs8);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLPoolingLayerFixedPointFixture<int8_t>, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), combine(PoolingLayerDatasetQS,
+ framework::dataset::make("DataType", DataType::QS8))),
+ framework::dataset::make("FractionalBits", 1, 4)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_qs8);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(QS16)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLPoolingLayerFixedPointFixture<int16_t>, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), combine(PoolingLayerDatasetQS,
+ framework::dataset::make("DataType", DataType::QS16))),
+ framework::dataset::make("FractionalBits", 1, 12)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_qs16);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLPoolingLayerFixedPointFixture<int16_t>, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), combine(PoolingLayerDatasetQS,
+ framework::dataset::make("DataType", DataType::QS16))),
+ framework::dataset::make("FractionalBits", 1, 12)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_qs16);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/ROIPoolingLayer.cpp b/tests/validation/CL/ROIPoolingLayer.cpp
deleted file mode 100644
index 19d7903128..0000000000
--- a/tests/validation/CL/ROIPoolingLayer.cpp
+++ /dev/null
@@ -1,112 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "CL/CLAccessor.h"
-#include "CL/CLArrayAccessor.h"
-#include "TypePrinter.h"
-#include "arm_compute/runtime/CL/CLArray.h"
-#include "arm_compute/runtime/CL/functions/CLROIPoolingLayer.h"
-#include "tests/Globals.h"
-#include "tests/Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-#include "validation/ValidationUserConfiguration.h"
-
-#include <random>
-#include <vector>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-CLTensor compute_roi_pooling_layer(const TensorShape &shape, DataType dt, const std::vector<ROI> &rois, ROIPoolingLayerInfo pool_info)
-{
- TensorShape shape_dst;
- shape_dst.set(0, pool_info.pooled_width());
- shape_dst.set(1, pool_info.pooled_height());
- shape_dst.set(2, shape.z());
- shape_dst.set(3, rois.size());
-
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, dt);
- CLTensor dst = create_tensor<CLTensor>(shape_dst, dt);
-
- // Create ROI array
- CLArray<ROI> rois_array(rois.size());
- fill_array(CLArrayAccessor<ROI>(rois_array), rois);
-
- // Create and configure function
- CLROIPoolingLayer roi_pool;
- roi_pool.configure(&src, &rois_array, &dst, pool_info);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
-
- // Fill tensors
- std::uniform_real_distribution<> distribution(-1, 1);
- library->fill(CLAccessor(src), distribution, 0);
-
- // Compute function
- roi_pool.run();
-
- return dst;
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(ROIPoolingLayer)
-
-BOOST_AUTO_TEST_SUITE(Float)
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, boost::unit_test::data::make({ DataType::F16, DataType::F32 }) * boost::unit_test::data::make({ 10, 20, 40 }) * boost::unit_test::data::make({ 7, 9 }) *
- boost::unit_test::data::make({ 1.f / 8.f, 1.f / 16.f }),
- dt, num_rois, roi_pool_size, roi_scale)
-{
- TensorShape shape(50U, 47U, 2U, 3U);
- ROIPoolingLayerInfo pool_info(roi_pool_size, roi_pool_size, roi_scale);
-
- // Construct ROI vector
- std::vector<ROI> rois = generate_random_rois(shape, pool_info, num_rois, user_config.seed);
-
- // Compute function
- CLTensor dst = compute_roi_pooling_layer(shape, dt, rois, pool_info);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_roi_pooling_layer(shape, dt, rois, pool_info);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-BOOST_AUTO_TEST_SUITE_END()
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/ReductionOperation.cpp b/tests/validation/CL/ReductionOperation.cpp
new file mode 100644
index 0000000000..5896add68f
--- /dev/null
+++ b/tests/validation/CL/ReductionOperation.cpp
@@ -0,0 +1,78 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLReductionOperation.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/ReductionOperationDataset.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/ReductionOperationFixture.h"
+#include "tests/validation/half.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Tolerance for float operations */
+constexpr RelativeTolerance tolerance_f32(0.00001f);
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(ReductionOperation)
+
+template <typename T>
+using CLReductionOperationFixture = ReductionOperationValidationFixture<CLTensor, CLAccessor, CLReductionOperation, T>;
+
+TEST_SUITE(Float)
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLReductionOperationFixture<float>, framework::DatasetMode::PRECOMMIT,
+ combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Axis", { 0 })), datasets::ReductionOperations()))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLReductionOperationFixture<float>, framework::DatasetMode::NIGHTLY,
+ combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Axis", { 0 })), datasets::ReductionOperations()))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/Scale.cpp b/tests/validation/CL/Scale.cpp
new file mode 100644
index 0000000000..d5866fad97
--- /dev/null
+++ b/tests/validation/CL/Scale.cpp
@@ -0,0 +1,129 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/functions/CLScale.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/BorderModeDataset.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Helpers.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/ScaleFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Tolerance */
+constexpr AbsoluteTolerance<uint8_t> tolerance(1);
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(Scale)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(combine(combine(concat(datasets::SmallShapes(), datasets::LargeShapes()), framework::dataset::make("DataType", DataType::U8)),
+ framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+ datasets::BorderModes()),
+ shape, data_type, policy, border_mode)
+{
+ std::mt19937 generator(library->seed());
+ std::uniform_real_distribution<float> distribution_float(0.25, 2);
+ const float scale_x = distribution_float(generator);
+ const float scale_y = distribution_float(generator);
+ std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+ uint8_t constant_border_value = distribution_u8(generator);
+
+ // Create tensors
+ CLTensor src = create_tensor<CLTensor>(shape, data_type);
+ TensorShape shape_scaled(shape);
+ shape_scaled.set(0, shape[0] * scale_x);
+ shape_scaled.set(1, shape[1] * scale_y);
+ CLTensor dst = create_tensor<CLTensor>(shape_scaled, data_type);
+
+ ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Create and configure function
+ CLScale clscale;
+ clscale.configure(&src, &dst, policy, border_mode, constant_border_value);
+
+ // Validate valid region
+ const ValidRegion dst_valid_region = calculate_valid_region_scale(*(src.info()), shape_scaled, policy, BorderSize(1), (border_mode == BorderMode::UNDEFINED));
+
+ validate(dst.info()->valid_region(), dst_valid_region);
+
+ // Validate padding
+ PaddingCalculator calculator(shape_scaled.x(), 4);
+ calculator.set_border_mode(border_mode);
+
+ const PaddingSize read_padding(1);
+ const PaddingSize write_padding = calculator.required_padding(PaddingCalculator::Option::EXCLUDE_BORDER);
+ validate(src.info()->padding(), read_padding);
+ validate(dst.info()->padding(), write_padding);
+}
+
+template <typename T>
+using CLScaleFixture = ScaleValidationFixture<CLTensor, CLAccessor, CLScale, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, CLScaleFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType",
+ DataType::U8)),
+ framework::dataset::make("InterpolationPolicy",
+{ InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+datasets::BorderModes()))
+{
+ //Create valid region
+ TensorInfo src_info(_shape, 1, _data_type);
+ const ValidRegion valid_region = calculate_valid_region_scale(src_info, _reference.shape(), _policy, BorderSize(1), (_border_mode == BorderMode::UNDEFINED));
+
+ // Validate output
+ validate(CLAccessor(_target), _reference, valid_region, tolerance);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLScaleFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType",
+ DataType::U8)),
+ framework::dataset::make("InterpolationPolicy",
+{ InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+datasets::BorderModes()))
+{
+ //Create valid region
+ TensorInfo src_info(_shape, 1, _data_type);
+ const ValidRegion valid_region = calculate_valid_region_scale(src_info, _reference.shape(), _policy, BorderSize(1), (_border_mode == BorderMode::UNDEFINED));
+
+ // Validate output
+ validate(CLAccessor(_target), _reference, valid_region, tolerance);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/Sobel3x3.cpp b/tests/validation/CL/Sobel3x3.cpp
deleted file mode 100644
index 9e32a3da66..0000000000
--- a/tests/validation/CL/Sobel3x3.cpp
+++ /dev/null
@@ -1,205 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "PaddingCalculator.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-#include "validation/ValidationUserConfiguration.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/CLSubTensor.h"
-#include "arm_compute/runtime/CL/CLTensor.h"
-#include "arm_compute/runtime/CL/CLTensorAllocator.h"
-#include "arm_compute/runtime/CL/functions/CLSobel3x3.h"
-
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-constexpr unsigned int filter_size = 3; /** Size of the kernel/filter in number of elements. */
-constexpr BorderSize border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */
-
-/** Compute CL Sobel 3x3 function.
- *
- * @param[in] shape Shape of the input and output tensors.
- * @param[in] border_mode BorderMode used by the input tensor
- * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT
- *
- * @return Computed output tensor.
- */
-std::pair<CLTensor, CLTensor> compute_sobel_3x3(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst_x = create_tensor<CLTensor>(shape, DataType::S16);
- CLTensor dst_y = create_tensor<CLTensor>(shape, DataType::S16);
-
- src.info()->set_format(Format::U8);
- dst_x.info()->set_format(Format::S16);
- dst_y.info()->set_format(Format::S16);
-
- // Create sobel image configure function
- CLSobel3x3 sobel_3x3;
- sobel_3x3.configure(&src, &dst_x, &dst_y, border_mode, constant_border_value);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst_x.allocator()->allocate();
- dst_y.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
- BOOST_TEST(!dst_x.info()->is_resizable());
- BOOST_TEST(!dst_y.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src), 0);
-
- // Compute function
- sobel_3x3.run();
-
- return std::make_pair(std::move(dst_x), std::move(dst_y));
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(Sobel3x3)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst_x = create_tensor<CLTensor>(shape, DataType::S16);
- CLTensor dst_y = create_tensor<CLTensor>(shape, DataType::S16);
-
- src.info()->set_format(Format::U8);
- dst_x.info()->set_format(Format::S16);
- dst_y.info()->set_format(Format::S16);
-
- BOOST_TEST(src.info()->is_resizable());
- BOOST_TEST(dst_x.info()->is_resizable());
- BOOST_TEST(dst_y.info()->is_resizable());
-
- // Create sobel 3x3 configure function
- CLSobel3x3 sobel_3x3;
- sobel_3x3.configure(&src, &dst_x, &dst_y, border_mode);
-
- // Validate valid region
- const ValidRegion src_valid_region = shape_to_valid_region(shape);
- const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
-
- validate(src.info()->valid_region(), src_valid_region);
- validate(dst_x.info()->valid_region(), dst_valid_region);
- validate(dst_y.info()->valid_region(), dst_valid_region);
-
- // Validate padding
- PaddingCalculator calculator(shape.x(), 8);
-
- calculator.set_border_mode(border_mode);
- calculator.set_border_size(1);
-
- const PaddingSize dst_padding = calculator.required_padding();
-
- calculator.set_accessed_elements(16);
- calculator.set_access_offset(-1);
-
- const PaddingSize src_padding = calculator.required_padding();
-
- validate(src.info()->padding(), src_padding);
- validate(dst_x.info()->padding(), dst_padding);
- validate(dst_y.info()->padding(), dst_padding);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode)
-{
- uint8_t constant_border_value = 0;
-
- // Generate a random constant value if border_mode is constant
- if(border_mode == BorderMode::CONSTANT)
- {
- std::mt19937 gen(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution(0, 255);
- constant_border_value = distribution(gen);
- }
-
- // Compute function
- std::pair<CLTensor, CLTensor> dst = compute_sobel_3x3(shape, border_mode, constant_border_value);
-
- // Compute reference
- std::pair<RawTensor, RawTensor> ref_dst = Reference::compute_reference_sobel_3x3(shape, border_mode, constant_border_value);
-
- // Calculate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
-
- // Validate output
- validate(CLAccessor(dst.first), ref_dst.first, valid_region);
- validate(CLAccessor(dst.second), ref_dst.second, valid_region);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode)
-{
- uint8_t constant_border_value = 0;
-
- // Generate a random constant value if border_mode is constant
- if(border_mode == BorderMode::CONSTANT)
- {
- std::mt19937 gen(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution(0, 255);
- constant_border_value = distribution(gen);
- }
-
- // Compute function
- std::pair<CLTensor, CLTensor> dst = compute_sobel_3x3(shape, border_mode, constant_border_value);
-
- // Compute reference
- std::pair<RawTensor, RawTensor> ref_dst = Reference::compute_reference_sobel_3x3(shape, border_mode, constant_border_value);
-
- // Calculate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
-
- // Validate output
- validate(CLAccessor(dst.first), ref_dst.first, valid_region);
- validate(CLAccessor(dst.second), ref_dst.second, valid_region);
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/Sobel5x5.cpp b/tests/validation/CL/Sobel5x5.cpp
deleted file mode 100644
index a7c971aa3a..0000000000
--- a/tests/validation/CL/Sobel5x5.cpp
+++ /dev/null
@@ -1,204 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "PaddingCalculator.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-#include "validation/ValidationUserConfiguration.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/CLSubTensor.h"
-#include "arm_compute/runtime/CL/CLTensor.h"
-#include "arm_compute/runtime/CL/CLTensorAllocator.h"
-#include "arm_compute/runtime/CL/functions/CLSobel5x5.h"
-
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-constexpr unsigned int filter_size = 5; /** Size of the kernel/filter in number of elements. */
-constexpr BorderSize border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */
-
-/** Compute CL Sobel 5x5 function.
- *
- * @param[in] shape Shape of the input and output tensors.
- * @param[in] border_mode BorderMode used by the input tensor
- * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT
- *
- * @return Computed output tensor.
- */
-std::pair<CLTensor, CLTensor> compute_sobel_5x5(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst_x = create_tensor<CLTensor>(shape, DataType::S16);
- CLTensor dst_y = create_tensor<CLTensor>(shape, DataType::S16);
-
- src.info()->set_format(Format::U8);
- dst_x.info()->set_format(Format::S16);
- dst_y.info()->set_format(Format::S16);
-
- // Create sobel image configure function
- CLSobel5x5 sobel_5x5;
- sobel_5x5.configure(&src, &dst_x, &dst_y, border_mode, constant_border_value);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst_x.allocator()->allocate();
- dst_y.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
- BOOST_TEST(!dst_x.info()->is_resizable());
- BOOST_TEST(!dst_y.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src), 0);
-
- // Compute function
- sobel_5x5.run();
-
- return std::make_pair(std::move(dst_x), std::move(dst_y));
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(Sobel5x5)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst_x = create_tensor<CLTensor>(shape, DataType::S16);
- CLTensor dst_y = create_tensor<CLTensor>(shape, DataType::S16);
-
- src.info()->set_format(Format::U8);
- dst_x.info()->set_format(Format::S16);
- dst_y.info()->set_format(Format::S16);
-
- BOOST_TEST(src.info()->is_resizable());
- BOOST_TEST(dst_x.info()->is_resizable());
- BOOST_TEST(dst_y.info()->is_resizable());
-
- // Create sobel 5x5 configure function
- CLSobel5x5 sobel_5x5;
- sobel_5x5.configure(&src, &dst_x, &dst_y, border_mode);
-
- // Validate valid region
- const ValidRegion src_valid_region = shape_to_valid_region(shape);
- const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
-
- validate(src.info()->valid_region(), src_valid_region);
- validate(dst_x.info()->valid_region(), dst_valid_region);
- validate(dst_y.info()->valid_region(), dst_valid_region);
-
- // Validate padding
- PaddingCalculator calculator(shape.x(), 8);
- calculator.set_border_mode(border_mode);
- calculator.set_border_size(2);
-
- const PaddingSize dst_padding = calculator.required_padding();
-
- calculator.set_accessed_elements(16);
- calculator.set_access_offset(-2);
-
- const PaddingSize src_padding = calculator.required_padding();
-
- validate(src.info()->padding(), src_padding);
- validate(dst_x.info()->padding(), dst_padding);
- validate(dst_y.info()->padding(), dst_padding);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode)
-{
- uint8_t constant_border_value = 0;
-
- // Generate a random constant value if border_mode is constant
- if(border_mode == BorderMode::CONSTANT)
- {
- std::mt19937 gen(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution(0, 255);
- constant_border_value = distribution(gen);
- }
-
- // Compute function
- std::pair<CLTensor, CLTensor> dst = compute_sobel_5x5(shape, border_mode, constant_border_value);
-
- // Compute reference
- std::pair<RawTensor, RawTensor> ref_dst = Reference::compute_reference_sobel_5x5(shape, border_mode, constant_border_value);
-
- // Calculate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
-
- // Validate output
- validate(CLAccessor(dst.first), ref_dst.first, valid_region);
- validate(CLAccessor(dst.second), ref_dst.second, valid_region);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode)
-{
- uint8_t constant_border_value = 0;
-
- // Generate a random constant value if border_mode is constant
- if(border_mode == BorderMode::CONSTANT)
- {
- std::mt19937 gen(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution(0, 255);
- constant_border_value = distribution(gen);
- }
-
- // Compute function
- std::pair<CLTensor, CLTensor> dst = compute_sobel_5x5(shape, border_mode, constant_border_value);
-
- // Compute reference
- std::pair<RawTensor, RawTensor> ref_dst = Reference::compute_reference_sobel_5x5(shape, border_mode, constant_border_value);
-
- // Calculate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
-
- // Validate output
- validate(CLAccessor(dst.first), ref_dst.first, valid_region);
- validate(CLAccessor(dst.second), ref_dst.second, valid_region);
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/SoftmaxLayer.cpp b/tests/validation/CL/SoftmaxLayer.cpp
new file mode 100644
index 0000000000..c4a9970b78
--- /dev/null
+++ b/tests/validation/CL/SoftmaxLayer.cpp
@@ -0,0 +1,168 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLSoftmaxLayer.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/SoftmaxLayerFixture.h"
+#include "tests/validation/half.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Tolerance for float operations */
+constexpr AbsoluteTolerance<float> tolerance_f16(0.002f);
+constexpr AbsoluteTolerance<float> tolerance_f32(0.000001f);
+/** Tolerance for fixed point operations */
+constexpr AbsoluteTolerance<int8_t> tolerance_fixed_point(2);
+
+/** CNN data types */
+const auto CNNDataTypes = framework::dataset::make("DataType",
+{
+ DataType::F16,
+ DataType::F32,
+ DataType::QS8,
+ DataType::QS16,
+});
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(SoftmaxLayer)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(concat(datasets::SmallShapes(), datasets::LargeShapes()), CNNDataTypes), shape, data_type)
+{
+ // Set fixed point position data type allowed
+ const int fixed_point_position = is_data_type_fixed_point(data_type) ? 3 : 0;
+
+ // Create tensors
+ CLTensor src = create_tensor<CLTensor>(shape, data_type, 1, fixed_point_position);
+ CLTensor dst = create_tensor<CLTensor>(shape, data_type, 1, fixed_point_position);
+
+ ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Create and configure function
+ CLSoftmaxLayer smx_layer;
+ smx_layer.configure(&src, &dst);
+
+ // Validate valid region
+ const ValidRegion valid_region = shape_to_valid_region(shape);
+ validate(src.info()->valid_region(), valid_region);
+ validate(dst.info()->valid_region(), valid_region);
+
+ // Validate padding
+ const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+ validate(src.info()->padding(), padding);
+ validate(dst.info()->padding(), padding);
+}
+
+template <typename T>
+using CLSoftmaxLayerFixture = SoftmaxValidationFixture<CLTensor, CLAccessor, CLSoftmaxLayer, T>;
+
+TEST_SUITE(Float)
+TEST_SUITE(FP16)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLSoftmaxLayerFixture<half_float::half>, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f16);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLSoftmaxLayerFixture<half_float::half>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapes(), framework::dataset::make("DataType", DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f16);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLSoftmaxLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLSoftmaxLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapes(), framework::dataset::make("DataType", DataType::F32)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f32);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+template <typename T>
+using CLSoftmaxLayerFixedPointFixture = SoftmaxValidationFixedPointFixture<CLTensor, CLAccessor, CLSoftmaxLayer, T>;
+
+TEST_SUITE(Quantized)
+TEST_SUITE(QS8)
+// Testing for fixed point position [1,6) as reciprocal limits the maximum fixed point position to 5
+FIXTURE_DATA_TEST_CASE(RunSmall, CLSoftmaxLayerFixedPointFixture<int8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FractionalBits", 1, 6)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_fixed_point);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLSoftmaxLayerFixedPointFixture<int8_t>, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FractionalBits", 1, 6)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_fixed_point);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(QS16)
+// Testing for fixed point position [1,14) as reciprocal limits the maximum fixed point position to 14
+FIXTURE_DATA_TEST_CASE(RunSmall, CLSoftmaxLayerFixedPointFixture<int16_t>, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FractionalBits", 1, 14)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_fixed_point);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLSoftmaxLayerFixedPointFixture<int16_t>, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FractionalBits", 1, 14)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_fixed_point);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/TableLookup.cpp b/tests/validation/CL/TableLookup.cpp
deleted file mode 100644
index 40e847f1cd..0000000000
--- a/tests/validation/CL/TableLookup.cpp
+++ /dev/null
@@ -1,229 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "CL/CLLutAccessor.h"
-#include "Globals.h"
-#include "PaddingCalculator.h"
-#include "RawLutAccessor.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Helpers.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/CLTensor.h"
-#include "arm_compute/runtime/CL/CLTensorAllocator.h"
-#include "arm_compute/runtime/CL/functions/CLTableLookup.h"
-
-#include "boost_wrapper.h"
-
-#include <map>
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-/** Compute Table Lookup function.
- *
- * @param[in] shape Shape of the input tensors
- * @param[in] data_type Type of the input/output tensor
- * @param[in] lut The input LUT.
- *
- * @return Computed output cl tensor.
- */
-CLTensor compute_table_lookup(const TensorShape &shape, DataType data_type, CLLut &lut)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, data_type);
- CLTensor dst = create_tensor<CLTensor>(shape, data_type);
-
- // Create and configure function
- CLTableLookup table_lookup;
- table_lookup.configure(&src, &lut, &dst);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src), 0);
-
- // Compute function
- table_lookup.run();
-
- return dst;
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(TableLookup)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }),
- shape, data_type)
-{
- //Create Lut
- const int num_elem = (data_type == DataType::U8) ? std::numeric_limits<uint8_t>::max() + 1 : std::numeric_limits<int16_t>::max() - std::numeric_limits<int16_t>::lowest() + 1;
- CLLut cllut(num_elem, data_type);
-
- if(data_type == DataType::U8)
- {
- fill_lookuptable(CLLutAccessor<uint8_t>(cllut));
- }
- else
- {
- fill_lookuptable(CLLutAccessor<int16_t>(cllut));
- }
-
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, data_type);
- CLTensor dst = create_tensor<CLTensor>(shape, data_type);
-
- BOOST_TEST(src.info()->is_resizable());
- BOOST_TEST(dst.info()->is_resizable());
-
- // Create and configure function
- CLTableLookup table_lookup;
- table_lookup.configure(&src, &cllut, &dst);
-
- // Validate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape);
- validate(src.info()->valid_region(), valid_region);
- validate(dst.info()->valid_region(), valid_region);
-
- // Validate padding
- const PaddingSize padding = PaddingCalculator(shape.x(), 8).required_padding();
- validate(src.info()->padding(), padding);
- validate(dst.info()->padding(), padding);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall,
- SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }),
- shape, data_type)
-{
- //Create Lut
- const int num_elem = (data_type == DataType::U8) ? std::numeric_limits<uint8_t>::max() + 1 : std::numeric_limits<int16_t>::max() - std::numeric_limits<int16_t>::lowest() + 1;
- CLLut cllut(num_elem, data_type);
-
- if(data_type == DataType::U8)
- {
- //Create rawLut
- std::map<uint8_t, uint8_t> rawlut;
-
- //Fill the Lut
- fill_lookuptable(CLLutAccessor<uint8_t>(cllut));
- fill_lookuptable(RawLutAccessor<uint8_t>(rawlut));
-
- // Compute function
- CLTensor dst = compute_table_lookup(shape, data_type, cllut);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_table_lookup(shape, data_type, rawlut);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
- }
- else
- {
- //Create rawLut
- std::map<int16_t, int16_t> rawlut;
-
- //Fill the Lut
- fill_lookuptable(CLLutAccessor<int16_t>(cllut));
- fill_lookuptable(RawLutAccessor<int16_t>(rawlut));
-
- // Compute function
- CLTensor dst = compute_table_lookup(shape, data_type, cllut);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_table_lookup(shape, data_type, rawlut);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
- }
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge,
- LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }),
- shape, data_type)
-{
- //Create Lut
- const int num_elem = (data_type == DataType::U8) ? std::numeric_limits<uint8_t>::max() + 1 : std::numeric_limits<int16_t>::max() - std::numeric_limits<int16_t>::lowest() + 1;
- CLLut cllut(num_elem, data_type);
-
- if(data_type == DataType::U8)
- {
- //Create rawLut
- std::map<uint8_t, uint8_t> rawlut;
-
- //Fill the Lut
- fill_lookuptable(CLLutAccessor<uint8_t>(cllut));
- fill_lookuptable(RawLutAccessor<uint8_t>(rawlut));
-
- // Compute function
- CLTensor dst = compute_table_lookup(shape, data_type, cllut);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_table_lookup(shape, data_type, rawlut);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
- }
- else
- {
- //Create rawLut
- std::map<int16_t, int16_t> rawlut;
-
- //Fill the Lut
- fill_lookuptable(CLLutAccessor<int16_t>(cllut));
- fill_lookuptable(RawLutAccessor<int16_t>(rawlut));
-
- // Compute function
- CLTensor dst = compute_table_lookup(shape, data_type, cllut);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_table_lookup(shape, data_type, rawlut);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
- }
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/Threshold.cpp b/tests/validation/CL/Threshold.cpp
deleted file mode 100644
index a2e7b7b4ba..0000000000
--- a/tests/validation/CL/Threshold.cpp
+++ /dev/null
@@ -1,153 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "PaddingCalculator.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "dataset/ThresholdDataset.h"
-#include "validation/Datasets.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/CLTensor.h"
-#include "arm_compute/runtime/CL/CLTensorAllocator.h"
-#include "arm_compute/runtime/CL/functions/CLThreshold.h"
-
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-/** Compute Threshold function.
- *
- * @param[in] shape Shape of the input and output tensors.
- * @param[in] threshold Threshold. When the threshold type is RANGE, this is used as the lower threshold.
- * @param[in] false_value value to set when the condition is not respected.
- * @param[in] true_value value to set when the condition is respected.
- * @param[in] type Thresholding type. Either RANGE or BINARY.
- * @param[in] upper Upper threshold. Only used when the thresholding type is RANGE.
- *
- * @return Computed output tensor.
- */
-CLTensor compute_threshold(const TensorShape &shape, uint8_t threshold, uint8_t false_value, uint8_t true_value, ThresholdType type, uint8_t upper)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U8);
-
- // Create and configure function
- CLThreshold thrsh;
- thrsh.configure(&src, &dst, threshold, false_value, true_value, type, upper);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src), 0);
-
- // Compute function
- thrsh.run();
-
- return dst;
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(Threshold)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration,
- (SmallShapes() + LargeShapes()) * ThresholdDataset(),
- shape, threshold_conf)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U8);
-
- BOOST_TEST(src.info()->is_resizable());
- BOOST_TEST(dst.info()->is_resizable());
-
- // Create and configure function
- CLThreshold cl_threshold;
- cl_threshold.configure(&src, &dst, threshold_conf.threshold, threshold_conf.false_value, threshold_conf.true_value, threshold_conf.type, threshold_conf.upper);
-
- // Validate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape);
- validate(src.info()->valid_region(), valid_region);
- validate(dst.info()->valid_region(), valid_region);
-
- // Validate padding
- const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
- validate(src.info()->padding(), padding);
- validate(dst.info()->padding(), padding);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall,
- SmallShapes() * ThresholdDataset(),
- shape, threshold_conf)
-{
- // Compute function
- CLTensor dst = compute_threshold(shape, threshold_conf.threshold, threshold_conf.false_value, threshold_conf.true_value, threshold_conf.type, threshold_conf.upper);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_threshold(shape, threshold_conf.threshold, threshold_conf.false_value, threshold_conf.true_value, threshold_conf.type, threshold_conf.upper);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge,
- LargeShapes() * ThresholdDataset(),
- shape, threshold_conf)
-{
- // Compute function
- CLTensor dst = compute_threshold(shape, threshold_conf.threshold, threshold_conf.false_value, threshold_conf.true_value, threshold_conf.type, threshold_conf.upper);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_threshold(shape, threshold_conf.threshold, threshold_conf.false_value, threshold_conf.true_value, threshold_conf.type, threshold_conf.upper);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst);
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/CL/WarpPerspective.cpp b/tests/validation/CL/WarpPerspective.cpp
deleted file mode 100644
index 260b22be03..0000000000
--- a/tests/validation/CL/WarpPerspective.cpp
+++ /dev/null
@@ -1,208 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "AssetsLibrary.h"
-#include "CL/CLAccessor.h"
-#include "Globals.h"
-#include "PaddingCalculator.h"
-#include "TypePrinter.h"
-#include "Utils.h"
-#include "validation/Datasets.h"
-#include "validation/Helpers.h"
-#include "validation/Reference.h"
-#include "validation/Validation.h"
-#include "validation/ValidationUserConfiguration.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/functions/CLWarpPerspective.h"
-#include "arm_compute/runtime/Tensor.h"
-#include "arm_compute/runtime/TensorAllocator.h"
-
-#include "boost_wrapper.h"
-
-#include <random>
-#include <string>
-
-using namespace arm_compute;
-using namespace arm_compute::test;
-using namespace arm_compute::test::validation;
-
-namespace
-{
-/** Compute Warp Perspective function.
- *
- * @param[in] input Shape of the input and output tensors.
- * @param[in] matrix The perspective matrix. Must be 3x3 of type float.
- * @param[in] policy The interpolation type.
- * @param[in] border_mode Strategy to use for borders.
- * @param[in] constant_border_value Constant value to use for borders if border_mode is set to CONSTANT.
- *
- * @return Computed output tensor.
- */
-CLTensor compute_warp_perspective(const TensorShape &shape, const float *matrix, InterpolationPolicy policy,
- BorderMode border_mode, uint8_t constant_border_value)
-{
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U8);
-
- // Create and configure function
- CLWarpPerspective warp_perspective;
- warp_perspective.configure(&src, &dst, matrix, policy, border_mode, constant_border_value);
-
- // Allocate tensors
- src.allocator()->allocate();
- dst.allocator()->allocate();
-
- BOOST_TEST(!src.info()->is_resizable());
- BOOST_TEST(!dst.info()->is_resizable());
-
- // Fill tensors
- library->fill_tensor_uniform(CLAccessor(src), 0);
-
- // Compute function
- warp_perspective.run();
-
- return dst;
-}
-} // namespace
-
-#ifndef DOXYGEN_SKIP_THIS
-BOOST_AUTO_TEST_SUITE(CL)
-BOOST_AUTO_TEST_SUITE(WarpPerspective)
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes())
- * boost::unit_test::data::make({ InterpolationPolicy::BILINEAR, InterpolationPolicy::NEAREST_NEIGHBOR }) * BorderModes(),
- shape, policy, border_mode)
-{
- uint8_t constant_border_value = 0;
-
- // Generate a random constant value if border_mode is constant
- if(border_mode == BorderMode::CONSTANT)
- {
- std::mt19937 gen(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
- constant_border_value = distribution_u8(gen);
- }
-
- std::array<float, 9> matrix;
- fill_warp_matrix<9>(matrix, 3, 3);
-
- // Create tensors
- CLTensor src = create_tensor<CLTensor>(shape, DataType::U8);
- CLTensor dst = create_tensor<CLTensor>(shape, DataType::U8);
-
- BOOST_TEST(src.info()->is_resizable());
- BOOST_TEST(dst.info()->is_resizable());
-
- // Create and configure function
- CLWarpPerspective warp_perspective;
- warp_perspective.configure(&src, &dst, matrix.data(), policy, border_mode, constant_border_value);
-
- // Validate valid region
- const ValidRegion valid_region = shape_to_valid_region(shape);
-
- validate(src.info()->valid_region(), valid_region);
- validate(dst.info()->valid_region(), valid_region);
-
- // Validate padding
- PaddingCalculator calculator(shape.x(), 4);
- calculator.set_border_mode(border_mode);
-
- const PaddingSize read_padding(1);
- const PaddingSize write_padding = calculator.required_padding(PaddingCalculator::Option::EXCLUDE_BORDER);
-
- validate(src.info()->padding(), read_padding);
- validate(dst.info()->padding(), write_padding);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
-BOOST_DATA_TEST_CASE(RunSmall, SmallShapes()
- * boost::unit_test::data::make({ InterpolationPolicy::BILINEAR, InterpolationPolicy::NEAREST_NEIGHBOR })
- * BorderModes(),
- shape, policy, border_mode)
-{
- uint8_t constant_border_value = 0;
-
- // Generate a random constant value if border_mode is constant
- if(border_mode == BorderMode::CONSTANT)
- {
- std::mt19937 gen(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
- constant_border_value = distribution_u8(gen);
- }
-
- // Create the valid mask Tensor
- RawTensor valid_mask(shape, DataType::U8);
-
- // Create the matrix
- std::array<float, 9> matrix;
- fill_warp_matrix<9>(matrix, 3, 3);
-
- // Compute function
- CLTensor dst = compute_warp_perspective(shape, matrix.data(), policy, border_mode, constant_border_value);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_warp_perspective(shape, valid_mask, matrix.data(), policy, border_mode, constant_border_value);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, valid_mask, 1, 0.2f);
-}
-
-BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
-BOOST_DATA_TEST_CASE(RunLarge, LargeShapes()
- * boost::unit_test::data::make({ InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR }) * BorderModes(),
- shape, policy, border_mode)
-{
- uint8_t constant_border_value = 0;
-
- // Generate a random constant value if border_mode is constant
- if(border_mode == BorderMode::CONSTANT)
- {
- std::mt19937 gen(user_config.seed.get());
- std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
- constant_border_value = distribution_u8(gen);
- }
-
- // Create the valid mask Tensor
- RawTensor valid_mask(shape, DataType::U8);
-
- // Create the matrix
- std::array<float, 9> matrix;
- fill_warp_matrix<9>(matrix, 3, 3);
-
- // Compute function
- CLTensor dst = compute_warp_perspective(shape, matrix.data(), policy, border_mode, constant_border_value);
-
- // Compute reference
- RawTensor ref_dst = Reference::compute_reference_warp_perspective(shape, valid_mask, matrix.data(), policy, border_mode, constant_border_value);
-
- // Validate output
- validate(CLAccessor(dst), ref_dst, valid_mask, 1, 0.2f);
-}
-
-BOOST_AUTO_TEST_SUITE_END()
-BOOST_AUTO_TEST_SUITE_END()
-#endif /* DOXYGEN_SKIP_THIS */