arm_compute v17.10
Change-Id: If1489af40eccd0219ede8946577afbf04db31b29
diff --git a/tests/validation/CL/BatchNormalizationLayer.cpp b/tests/validation/CL/BatchNormalizationLayer.cpp
index ac30c63..69f8d7b 100644
--- a/tests/validation/CL/BatchNormalizationLayer.cpp
+++ b/tests/validation/CL/BatchNormalizationLayer.cpp
@@ -43,9 +43,10 @@
{
namespace
{
-constexpr AbsoluteTolerance<float> tolerance_f(0.00001f); /**< Tolerance value for comparing reference's output against implementation's output for DataType::F32 */
-constexpr AbsoluteTolerance<float> tolerance_qs8(3.0f); /**< Tolerance value for comparing reference's output against implementation's output for DataType::QS8 */
-constexpr AbsoluteTolerance<float> tolerance_qs16(6.0f); /**< Tolerance value for comparing reference's output against implementation's output for DataType::QS16 */
+constexpr AbsoluteTolerance<float> tolerance_f32(0.00001f); /**< Tolerance value for comparing reference's output against implementation's output for DataType::F32 */
+constexpr AbsoluteTolerance<float> tolerance_f16(0.01f); /**< Tolerance value for comparing reference's output against implementation's output for DataType::F16 */
+constexpr AbsoluteTolerance<float> tolerance_qs8(3.0f); /**< Tolerance value for comparing reference's output against implementation's output for DataType::QS8 */
+constexpr AbsoluteTolerance<float> tolerance_qs16(6.0f); /**< Tolerance value for comparing reference's output against implementation's output for DataType::QS16 */
} // namespace
TEST_SUITE(CL)
@@ -54,7 +55,7 @@
template <typename T>
using CLBatchNormalizationLayerFixture = BatchNormalizationLayerValidationFixture<CLTensor, CLAccessor, CLBatchNormalizationLayer, T>;
-DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(datasets::RandomBatchNormalizationLayerDataset(), framework::dataset::make("DataType", { DataType::QS8, DataType::QS16, DataType::F32 })),
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(datasets::RandomBatchNormalizationLayerDataset(), framework::dataset::make("DataType", { DataType::QS8, DataType::QS16, DataType::F16, DataType::F32 })),
shape0, shape1, epsilon, dt)
{
// Set fixed point position data type allowed
@@ -78,14 +79,25 @@
}
TEST_SUITE(Float)
+TEST_SUITE(FP32)
FIXTURE_DATA_TEST_CASE(Random, CLBatchNormalizationLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(datasets::RandomBatchNormalizationLayerDataset(),
framework::dataset::make("DataType", DataType::F32)))
{
// Validate output
- validate(CLAccessor(_target), _reference, tolerance_f, 0);
+ validate(CLAccessor(_target), _reference, tolerance_f32, 0);
}
TEST_SUITE_END()
+TEST_SUITE(FP16)
+FIXTURE_DATA_TEST_CASE(Random, CLBatchNormalizationLayerFixture<half>, framework::DatasetMode::PRECOMMIT, combine(datasets::RandomBatchNormalizationLayerDataset(),
+ framework::dataset::make("DataType", DataType::F16)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_f16, 0);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
TEST_SUITE(Quantized)
template <typename T>
using CLBatchNormalizationLayerFixedPointFixture = BatchNormalizationLayerValidationFixedPointFixture<CLTensor, CLAccessor, CLBatchNormalizationLayer, T>;
diff --git a/tests/validation/CL/FillBorder.cpp b/tests/validation/CL/FillBorder.cpp
new file mode 100644
index 0000000..a817338
--- /dev/null
+++ b/tests/validation/CL/FillBorder.cpp
@@ -0,0 +1,103 @@
+/*
+ * 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/CL/kernels/CLFillBorderKernel.h"
+#include "arm_compute/runtime/CL/CLScheduler.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/Globals.h"
+#include "tests/datasets/BorderModeDataset.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(CL)
+TEST_SUITE(FillBorder)
+
+// *INDENT-OFF*
+// clang-format off
+const auto PaddingSizesDataset = concat(concat(
+ framework::dataset::make("PaddingSize", PaddingSize{ 0 }),
+ framework::dataset::make("PaddingSize", PaddingSize{ 1, 0, 1, 2 })),
+ framework::dataset::make("PaddingSize", PaddingSize{ 10 }));
+
+const auto BorderSizesDataset = framework::dataset::make("BorderSize", 0, 6);
+
+DATA_TEST_CASE(FillBorder, framework::DatasetMode::ALL, combine(combine(combine(combine(
+ datasets::SmallShapes(),
+ datasets::BorderModes()),
+ BorderSizesDataset),
+ PaddingSizesDataset),
+ framework::dataset::make("DataType", DataType::U8)),
+ shape, border_mode, size, padding, data_type)
+// clang-format on
+// *INDENT-ON*
+{
+ BorderSize border_size{ static_cast<unsigned int>(size) };
+
+ std::mt19937 generator(library->seed());
+ std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+ const uint8_t border_value = distribution_u8(generator);
+ const uint8_t tensor_value = distribution_u8(generator);
+
+ // Create tensors
+ CLTensor src = create_tensor<CLTensor>(shape, data_type);
+
+ 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);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/FixedPoint/FixedPointTarget.h b/tests/validation/CL/FixedPoint/FixedPointTarget.h
new file mode 100644
index 0000000..3847354
--- /dev/null
+++ b/tests/validation/CL/FixedPoint/FixedPointTarget.h
@@ -0,0 +1,131 @@
+/*
+ * 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.
+ */
+#ifndef ARM_COMPUTE_TEST_FIXED_POINT_CL_TARGET
+#define ARM_COMPUTE_TEST_FIXED_POINT_CL_TARGET
+
+#include "arm_compute/runtime/CL/CLScheduler.h"
+
+#include "tests/Globals.h"
+#include "tests/Types.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+template <typename TensorType, typename AccessorType, typename T>
+void compute_target_impl(const TensorShape &shape, DataType dt, FixedPointOp op, int fixed_point_position, TensorType &src, TensorType &dst)
+{
+ std::string fixed_point_operation_kernel;
+#ifndef EMBEDDED_KERNELS
+ std::cout << "EMBEDDED_KERNELS NOT DEFINED" << std::endl;
+
+ 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";
+
+ // 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 = (dt == DataType::QS8) ? 0x7F : 0x7FFF;
+ build_opts += " -DFIXED_POINT_OP=1";
+ break;
+ case(FixedPointOp::LOG):
+ min = (1 << (fixed_point_position - 1));
+ max = (dt == DataType::QS8) ? 0x3F : 0x3FFF;
+ build_opts += " -DFIXED_POINT_OP=2";
+ break;
+ default:
+ ARM_COMPUTE_ERROR("Fixed point operation not supported");
+ break;
+ }
+
+ std::uniform_int_distribution<> distribution(min, max);
+ library->fill(AccessorType(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);
+}
+} // namespace
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_FIXED_POINT_TARGET */
diff --git a/tests/validation/CL/FixedPoint/FixedPoint_QS8.cpp b/tests/validation/CL/FixedPoint/FixedPoint_QS8.cpp
new file mode 100644
index 0000000..0fb4641
--- /dev/null
+++ b/tests/validation/CL/FixedPoint/FixedPoint_QS8.cpp
@@ -0,0 +1,96 @@
+/*
+ * 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 "FixedPointTarget.h"
+
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.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/FixedPointFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+constexpr AbsoluteTolerance<float> tolerance_exp(1.0f); /**< Tolerance value for comparing reference's output against implementation's output (exponential)*/
+constexpr AbsoluteTolerance<float> tolerance_invsqrt(4.0f); /**< Tolerance value for comparing reference's output against implementation's output (inverse square-root) */
+constexpr AbsoluteTolerance<float> tolerance_log(5.0f); /**< Tolerance value for comparing reference's output against implementation's output (logarithm) */
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(FixedPoint)
+TEST_SUITE(QS8)
+
+template <typename T>
+using CLFixedPointFixture = FixedPointValidationFixture<CLTensor, CLAccessor, T>;
+
+TEST_SUITE(Exp)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLFixedPointFixture<int8_t>, framework::DatasetMode::ALL, combine(combine(combine(datasets::Small1DShape(), framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FixedPointOp", FixedPointOp::EXP)),
+ framework::dataset::make("FractionalBits", 1, 6)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_exp);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(Log)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLFixedPointFixture<int8_t>, framework::DatasetMode::ALL, combine(combine(combine(datasets::Small1DShape(), framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FixedPointOp", FixedPointOp::LOG)),
+ framework::dataset::make("FractionalBits", 3, 6)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_log);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(Invsqrt)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLFixedPointFixture<int8_t>, framework::DatasetMode::ALL, combine(combine(combine(datasets::Small1DShape(), framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FixedPointOp", FixedPointOp::INV_SQRT)),
+ framework::dataset::make("FractionalBits", 1, 6)))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, tolerance_invsqrt);
+}
+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/GlobalPooling.cpp b/tests/validation/CL/GlobalPoolingLayer.cpp
similarity index 97%
rename from tests/validation/CL/GlobalPooling.cpp
rename to tests/validation/CL/GlobalPoolingLayer.cpp
index c5c9d00..31e3fe0 100644
--- a/tests/validation/CL/GlobalPooling.cpp
+++ b/tests/validation/CL/GlobalPoolingLayer.cpp
@@ -47,8 +47,8 @@
const auto GlobalPoolingLayerDataset = combine(datasets::GlobalPoolingShapes(), datasets::PoolingTypes());
/** Input data set for quantized data types */
-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_f32(0.001f); /**< Tolerance value for comparing reference's output against implementation's output for FP32 types */
+constexpr AbsoluteTolerance<float> tolerance_f16(0.01f); /**< Tolerance value for comparing reference's output against implementation's output for FP16 types */
} // namespace
TEST_SUITE(CL)
diff --git a/tests/validation/CL/WarpAffine.cpp b/tests/validation/CL/WarpAffine.cpp
new file mode 100644
index 0000000..9db2cca
--- /dev/null
+++ b/tests/validation/CL/WarpAffine.cpp
@@ -0,0 +1,120 @@
+/*
+ * 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/core/Utils.h"
+#include "arm_compute/runtime/CL/functions/CLWarpAffine.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/BorderModeDataset.h"
+#include "tests/datasets/InterpolationPolicyDataset.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/CPP/Utils.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/WarpAffineFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Tolerance */
+constexpr AbsoluteTolerance<uint8_t> tolerance(1);
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(WarpAffine)
+
+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)
+{
+ // Generate a random constant value if border_mode is constant
+ std::mt19937 gen(library->seed());
+ std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+ uint8_t constant_border_value = distribution_u8(gen);
+
+ // Create the matrix
+ std::array<float, 6> matrix{ {} };
+ fill_warp_matrix<6>(matrix);
+
+ // 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
+ CLWarpAffine warp_affine;
+ warp_affine.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
+ int total_right = ceil_to_multiple(shape[0], 4);
+ const int access_right = total_right + (((total_right - shape[0]) == 0) ? 1 : 0);
+ const PaddingSize read_padding(1, access_right - shape[0], 1, 1);
+ validate(src.info()->padding(), read_padding);
+
+ PaddingCalculator calculator(shape.x(), 4);
+ validate(dst.info()->padding(), calculator.required_padding());
+}
+
+template <typename T>
+using CLWarpAffineFixture = WarpAffineValidationFixture<CLTensor, CLAccessor, CLWarpAffine, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, CLWarpAffineFixture<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()))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, _valid_mask, tolerance, 0.02f);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLWarpAffineFixture<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()))
+{
+ // Validate output
+ validate(CLAccessor(_target), _reference, _valid_mask, tolerance, 0.02f);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/WarpPerspective.cpp b/tests/validation/CL/WarpPerspective.cpp
new file mode 100644
index 0000000..2edf911
--- /dev/null
+++ b/tests/validation/CL/WarpPerspective.cpp
@@ -0,0 +1,128 @@
+/*
+ * 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/CLWarpPerspective.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/Validation.h"
+#include "tests/validation/fixtures/WarpPerspectiveFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+constexpr AbsoluteTolerance<uint8_t> tolerance_value(1);
+constexpr float tolerance_number = 0.2f;
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(WarpPerspective)
+
+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)
+{
+ uint8_t constant_border_value = 0;
+
+ // Generate a random constant value if border_mode is constant
+ if(border_mode == BorderMode::CONSTANT)
+ {
+ std::mt19937 gen(library->seed());
+ std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+ constant_border_value = distribution_u8(gen);
+ }
+
+ // Create the matrix
+ std::array<float, 9> matrix = { { 0 } };
+ fill_warp_matrix<9>(matrix);
+
+ // 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
+ 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);
+}
+
+template <typename T>
+using CLWarpPerspectiveFixture = WarpPerspectiveValidationFixture<CLTensor, CLAccessor, CLWarpPerspective, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, CLWarpPerspectiveFixture<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 the valid mask Tensor
+ RawTensor valid_mask(_reference.shape(), _reference.data_type());
+
+ validate(CLAccessor(_target), _reference, valid_mask, tolerance_value, tolerance_number);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLWarpPerspectiveFixture<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 the valid mask Tensor
+ RawTensor valid_mask{ _reference.shape(), _reference.data_type() };
+
+ validate(CLAccessor(_target), _reference, valid_mask, tolerance_value, tolerance_number);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CPP/FixedPoint.cpp b/tests/validation/CPP/FixedPoint.cpp
new file mode 100644
index 0000000..a016093
--- /dev/null
+++ b/tests/validation/CPP/FixedPoint.cpp
@@ -0,0 +1,83 @@
+/*
+ * 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 "FixedPoint.h"
+
+#include "arm_compute/core/Types.h"
+#include "tests/validation/FixedPoint.h"
+#include "tests/validation/Helpers.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+SimpleTensor<T> fixed_point_operation(const SimpleTensor<T> &src, FixedPointOp op)
+{
+ SimpleTensor<T> result(src.shape(), src.data_type());
+
+ const int p = src.fixed_point_position();
+ switch(op)
+ {
+ case FixedPointOp::EXP:
+ for(int i = 0; i < src.num_elements(); ++i)
+ {
+ result[i] = fixed_point_arithmetic::exp(fixed_point_arithmetic::fixed_point<T>(src[i], p, true)).raw();
+ }
+ break;
+ case FixedPointOp::LOG:
+ for(int i = 0; i < src.num_elements(); ++i)
+ {
+ result[i] = fixed_point_arithmetic::log(fixed_point_arithmetic::fixed_point<T>(src[i], p, true)).raw();
+ }
+ break;
+ case FixedPointOp::INV_SQRT:
+ for(int i = 0; i < src.num_elements(); ++i)
+ {
+ result[i] = fixed_point_arithmetic::inv_sqrt(fixed_point_arithmetic::fixed_point<T>(src[i], p, true)).raw();
+ }
+ break;
+ case FixedPointOp::RECIPROCAL:
+ for(int i = 0; i < src.num_elements(); ++i)
+ {
+ result[i] = fixed_point_arithmetic::div(fixed_point_arithmetic::fixed_point<T>(1, p), fixed_point_arithmetic::fixed_point<T>(src[i], p, true)).raw();
+ }
+ break;
+ default:
+ ARM_COMPUTE_ERROR("Fixed point operation not supported");
+ break;
+ }
+
+ return result;
+}
+
+template SimpleTensor<int8_t> fixed_point_operation(const SimpleTensor<int8_t> &src, FixedPointOp op);
+template SimpleTensor<int16_t> fixed_point_operation(const SimpleTensor<int16_t> &src, FixedPointOp op);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CPP/FixedPoint.h b/tests/validation/CPP/FixedPoint.h
new file mode 100644
index 0000000..f0117f9
--- /dev/null
+++ b/tests/validation/CPP/FixedPoint.h
@@ -0,0 +1,44 @@
+/*
+ * 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.
+ */
+#ifndef __ARM_COMPUTE_TEST_FIXED_POINT_OPERATION_H__
+#define __ARM_COMPUTE_TEST_FIXED_POINT_OPERATION_H__
+
+#include "tests/SimpleTensor.h"
+#include "tests/Types.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+SimpleTensor<T> fixed_point_operation(const SimpleTensor<T> &src, FixedPointOp op);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_FIXED_POINT_OPERATION_H__ */
diff --git a/tests/validation/CPP/GEMMLowp.cpp b/tests/validation/CPP/GEMMLowp.cpp
new file mode 100644
index 0000000..d172a77
--- /dev/null
+++ b/tests/validation/CPP/GEMMLowp.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 "GEMM.h"
+
+#include "arm_compute/core/Types.h"
+#include "tests/validation/FixedPoint.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+SimpleTensor<T> gemmlowp(const SimpleTensor<T> &a, const SimpleTensor<T> &b, SimpleTensor<T> &c,
+ int32_t a_offset, int32_t b_offset, int32_t c_offset, int32_t c_mult_int, int32_t out_shift)
+{
+ const int K = a.shape().x();
+ const int b_width = b.shape().x();
+ const int rows = c.shape().y(); //M
+ const int cols = c.shape().x(); //N
+ std::vector<int32_t> acc;
+ acc.resize(cols);
+ for(int i = 0; i < rows; ++i)
+ {
+ for(int j = 0; j < cols; ++j)
+ {
+ acc[j] = 0;
+ }
+ for(int k = 0; k < K; ++k)
+ {
+ const int32_t tmp_a = a_offset + static_cast<int32_t>(a[k + i * K]);
+ for(int j = 0; j < b_width; ++j)
+ {
+ const int32_t tmp_b = b_offset + static_cast<int32_t>(b[j + k * b_width]);
+ const int32_t mult_as_int = tmp_a * tmp_b;
+ acc[j] += mult_as_int;
+ }
+ }
+ for(int j = 0; j < cols; ++j)
+ {
+ const int32_t result = ((c_offset + acc[j]) * c_mult_int) >> out_shift;
+ c[j + i * cols] = static_cast<uint8_t>(std::min(255, std::max(0, result)));
+ }
+ }
+
+ return c;
+}
+
+template SimpleTensor<uint8_t> gemmlowp(const SimpleTensor<uint8_t> &a, const SimpleTensor<uint8_t> &b, SimpleTensor<uint8_t> &c,
+ int32_t a_offset, int32_t b_offset, int32_t c_offset, int32_t c_mult_int, int32_t out_shift);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CPP/GEMMLowp.h b/tests/validation/CPP/GEMMLowp.h
new file mode 100644
index 0000000..2160975
--- /dev/null
+++ b/tests/validation/CPP/GEMMLowp.h
@@ -0,0 +1,45 @@
+/*
+ * 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.
+ */
+#ifndef __ARM_COMPUTE_TEST_GEMMLOWP_H__
+#define __ARM_COMPUTE_TEST_GEMMLOWP_H__
+
+#include "tests/SimpleTensor.h"
+#include "tests/validation/Helpers.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+SimpleTensor<T> gemmlowp(const SimpleTensor<T> &a, const SimpleTensor<T> &b, SimpleTensor<T> &c,
+ int32_t a_offset, int32_t b_offset, int32_t c_offset, int32_t c_mult_int, int32_t out_shift);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_GEMMLOWP_H__ */
diff --git a/tests/validation/CPP/Utils.cpp b/tests/validation/CPP/Utils.cpp
index af3ed90..d163e84 100644
--- a/tests/validation/CPP/Utils.cpp
+++ b/tests/validation/CPP/Utils.cpp
@@ -35,8 +35,8 @@
template <typename T>
T bilinear_policy(const SimpleTensor<T> &in, Coordinates id, float xn, float yn, BorderMode border_mode, T constant_border_value)
{
- int idx = std::floor(xn);
- int idy = std::floor(yn);
+ const int idx = std::floor(xn);
+ const int idy = std::floor(yn);
const float dx = xn - idx;
const float dy = yn - idy;
@@ -94,6 +94,19 @@
return dst;
}
+
+bool valid_bilinear_policy(float xn, float yn, int width, int height, BorderMode border_mode)
+{
+ if(border_mode != BorderMode::UNDEFINED)
+ {
+ return true;
+ }
+ if((0 <= yn + 1) && (yn + 1 < height) && (0 <= xn + 1) && (xn + 1 < width))
+ {
+ return true;
+ }
+ return false;
+}
} // namespace validation
} // namespace test
} // namespace arm_compute
diff --git a/tests/validation/CPP/Utils.h b/tests/validation/CPP/Utils.h
index 91d1afe..0733411 100644
--- a/tests/validation/CPP/Utils.h
+++ b/tests/validation/CPP/Utils.h
@@ -100,6 +100,28 @@
}
RawTensor transpose(const RawTensor &src, int chunk_width = 1);
+
+/** Fill matrix random.
+ *
+ * @param[in,out] matrix Matrix
+ */
+template <std::size_t SIZE>
+inline void fill_warp_matrix(std::array<float, SIZE> &matrix)
+{
+ std::mt19937 gen(library.get()->seed());
+ std::uniform_real_distribution<float> dist(-1, 1);
+ for(auto &x : matrix)
+ {
+ x = dist(gen);
+ }
+ if(SIZE == 9)
+ {
+ // This is only used in Warp Perspective, we set M[3][3] = 1 so that Z0 is not 0 and we avoid division by 0.
+ matrix[8] = 1.f;
+ }
+}
+
+bool valid_bilinear_policy(float xn, float yn, int width, int height, BorderMode border_mode);
} // namespace validation
} // namespace test
} // namespace arm_compute
diff --git a/tests/validation/CPP/WarpAffine.cpp b/tests/validation/CPP/WarpAffine.cpp
new file mode 100644
index 0000000..7b903b7
--- /dev/null
+++ b/tests/validation/CPP/WarpAffine.cpp
@@ -0,0 +1,135 @@
+/*
+ * 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 "WarpAffine.h"
+
+#include "Utils.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+bool valid_bilinear_policy(float xn, float yn, int width, int height, BorderMode border_mode)
+{
+ if(border_mode != BorderMode::UNDEFINED)
+ {
+ return true;
+ }
+ if((0 <= yn + 1) && (yn + 1 < height) && (0 <= xn + 1) && (xn + 1 < width))
+ {
+ return true;
+ }
+ return false;
+}
+
+template <typename T>
+SimpleTensor<T> warp_affine(const SimpleTensor<T> &src, SimpleTensor<T> &valid_mask, const float *matrix, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
+{
+ SimpleTensor<T> dst(src.shape(), src.data_type());
+
+ // x0 = M00 * x + M01 * y + M02
+ // y0 = M10 * x + M11 * y + M12
+ const float M00 = matrix[0];
+ const float M10 = matrix[1];
+ const float M01 = matrix[0 + 1 * 2];
+ const float M11 = matrix[1 + 1 * 2];
+ const float M02 = matrix[0 + 2 * 2];
+ const float M12 = matrix[1 + 2 * 2];
+
+ const int width = src.shape().x();
+ const int height = src.shape().y();
+
+ for(int element_idx = 0; element_idx < src.num_elements(); ++element_idx)
+ {
+ valid_mask[element_idx] = 1;
+ Coordinates id = index2coord(src.shape(), element_idx);
+ int idx = id.x();
+ int idy = id.y();
+
+ float x0 = M00 * idx + M01 * idy + M02;
+ float y0 = M10 * idx + M11 * idy + M12;
+
+ id.set(0, static_cast<int>(std::floor(x0)));
+ id.set(1, static_cast<int>(std::floor(y0)));
+ if((0 <= y0) && (y0 < height) && (0 <= x0) && (x0 < width))
+ {
+ switch(policy)
+ {
+ case InterpolationPolicy::NEAREST_NEIGHBOR:
+ dst[element_idx] = tensor_elem_at(src, id, border_mode, constant_border_value);
+ break;
+ case InterpolationPolicy::BILINEAR:
+ (valid_bilinear_policy(x0, y0, width, height, border_mode)) ? dst[element_idx] = bilinear_policy(src, id, x0, y0, border_mode, constant_border_value) :
+ valid_mask[element_idx] = 0;
+ break;
+ case InterpolationPolicy::AREA:
+ default:
+ ARM_COMPUTE_ERROR("Interpolation not supported");
+ }
+ }
+ else
+ {
+ if(border_mode == BorderMode::UNDEFINED)
+ {
+ valid_mask[element_idx] = 0;
+ }
+ else
+ {
+ switch(policy)
+ {
+ case InterpolationPolicy::NEAREST_NEIGHBOR:
+ if(border_mode == BorderMode::CONSTANT)
+ {
+ dst[element_idx] = constant_border_value;
+ }
+ else if(border_mode == BorderMode::REPLICATE)
+ {
+ id.set(0, std::max(0, std::min(static_cast<int>(x0), width - 1)));
+ id.set(1, std::max(0, std::min(static_cast<int>(y0), height - 1)));
+ dst[element_idx] = src[coord2index(src.shape(), id)];
+ }
+ break;
+ case InterpolationPolicy::BILINEAR:
+ dst[element_idx] = bilinear_policy(src, id, x0, y0, border_mode, constant_border_value);
+ break;
+ case InterpolationPolicy::AREA:
+ default:
+ ARM_COMPUTE_ERROR("Interpolation not supported");
+ }
+ }
+ }
+ }
+
+ return dst;
+}
+
+template SimpleTensor<uint8_t> warp_affine(const SimpleTensor<uint8_t> &src, SimpleTensor<uint8_t> &valid_mask, const float *matrix, InterpolationPolicy policy, BorderMode border_mode,
+ uint8_t constant_border_value);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
\ No newline at end of file
diff --git a/tests/validation/CPP/WarpAffine.h b/tests/validation/CPP/WarpAffine.h
new file mode 100644
index 0000000..973b1b2
--- /dev/null
+++ b/tests/validation/CPP/WarpAffine.h
@@ -0,0 +1,43 @@
+/*
+ * 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.
+ */
+#ifndef __ARM_COMPUTE_TEST_WARP_AFFINE_H__
+#define __ARM_COMPUTE_TEST_WARP_AFFINE_H__
+
+#include "tests/SimpleTensor.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+SimpleTensor<T> warp_affine(const SimpleTensor<T> &src, SimpleTensor<T> &valid_mask, const float *matrix, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_WARP_AFFINE_H__ */
diff --git a/tests/validation/CPP/WarpPerspective.cpp b/tests/validation/CPP/WarpPerspective.cpp
new file mode 100644
index 0000000..7a50253
--- /dev/null
+++ b/tests/validation/CPP/WarpPerspective.cpp
@@ -0,0 +1,132 @@
+/*
+ * 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 "Utils.h"
+#include "WarpPerspective.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+SimpleTensor<T> warp_perspective(const SimpleTensor<T> &src, SimpleTensor<T> &valid_mask, const float *matrix, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
+{
+ SimpleTensor<T> dst(src.shape(), src.data_type());
+
+ // x0 = M00 * x + M01 * y + M02
+ // y0 = M10 * x + M11 * y + M12
+ // z0 = M20 * x + M21 * y + M22
+ // xn = x0 / z0
+ // yn = y0 / z0
+ const float M00 = matrix[0];
+ const float M10 = matrix[1];
+ const float M20 = matrix[2];
+ const float M01 = matrix[0 + 1 * 3];
+ const float M11 = matrix[1 + 1 * 3];
+ const float M21 = matrix[2 + 1 * 3];
+ const float M02 = matrix[0 + 2 * 3];
+ const float M12 = matrix[1 + 2 * 3];
+ const float M22 = matrix[2 + 2 * 3];
+
+ const int width = src.shape().x();
+ const int height = src.shape().y();
+
+ for(int element_idx = 0; element_idx < src.num_elements(); ++element_idx)
+ {
+ valid_mask[element_idx] = 1;
+ Coordinates id = index2coord(src.shape(), element_idx);
+ const int idx = id.x();
+ const int idy = id.y();
+ const float z0 = M20 * idx + M21 * idy + M22;
+
+ const float x0 = (M00 * idx + M01 * idy + M02);
+ const float y0 = (M10 * idx + M11 * idy + M12);
+
+ const float xn = x0 / z0;
+ const float yn = y0 / z0;
+ id.set(0, static_cast<int>(std::floor(xn)));
+ id.set(1, static_cast<int>(std::floor(yn)));
+ if((0 <= yn) && (yn < height) && (0 <= xn) && (xn < width))
+ {
+ switch(policy)
+ {
+ case InterpolationPolicy::NEAREST_NEIGHBOR:
+ dst[element_idx] = tensor_elem_at(src, id, border_mode, constant_border_value);
+ break;
+ case InterpolationPolicy::BILINEAR:
+ (valid_bilinear_policy(xn, yn, width, height, border_mode)) ? dst[element_idx] = bilinear_policy(src, id, xn, yn, border_mode, constant_border_value) : valid_mask[element_idx] = 0;
+ break;
+ case InterpolationPolicy::AREA:
+ default:
+ ARM_COMPUTE_ERROR("Interpolation not supported");
+ break;
+ }
+ }
+ else
+ {
+ if(border_mode == BorderMode::UNDEFINED)
+ {
+ valid_mask[element_idx] = 0;
+ }
+ else
+ {
+ switch(policy)
+ {
+ case InterpolationPolicy::NEAREST_NEIGHBOR:
+ if(border_mode == BorderMode::CONSTANT)
+ {
+ dst[element_idx] = constant_border_value;
+ }
+ else if(border_mode == BorderMode::REPLICATE)
+ {
+ id.set(0, std::max(0, std::min(static_cast<int>(xn), width - 1)));
+ id.set(1, std::max(0, std::min(static_cast<int>(yn), height - 1)));
+ dst[element_idx] = src[coord2index(src.shape(), id)];
+ }
+ break;
+ case InterpolationPolicy::BILINEAR:
+ dst[element_idx] = bilinear_policy(src, id, xn, yn, border_mode, constant_border_value);
+ break;
+ case InterpolationPolicy::AREA:
+ default:
+ ARM_COMPUTE_ERROR("Interpolation not supported");
+ break;
+ }
+ }
+ }
+ }
+ return dst;
+}
+
+template SimpleTensor<uint8_t> warp_perspective(const SimpleTensor<uint8_t> &src, SimpleTensor<uint8_t> &valid_mask, const float *matrix, InterpolationPolicy policy, BorderMode border_mode,
+ uint8_t constant_border_value);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CPP/WarpPerspective.h b/tests/validation/CPP/WarpPerspective.h
new file mode 100644
index 0000000..2367f4d
--- /dev/null
+++ b/tests/validation/CPP/WarpPerspective.h
@@ -0,0 +1,43 @@
+/*
+ * 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.
+ */
+#ifndef __ARM_COMPUTE_TEST_WARP_PERSPECTIVE_H__
+#define __ARM_COMPUTE_TEST_WARP_PERSPECTIVE_H__
+
+#include "tests/SimpleTensor.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+SimpleTensor<T> warp_perspective(const SimpleTensor<T> &src, SimpleTensor<T> &valid_mask, const float *matrix, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_WARP_PERSPECTIVE_H__ */
diff --git a/tests/validation/NEON/FillBorder.cpp b/tests/validation/NEON/FillBorder.cpp
new file mode 100644
index 0000000..7e0fb1a
--- /dev/null
+++ b/tests/validation/NEON/FillBorder.cpp
@@ -0,0 +1,102 @@
+/*
+ * 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/NEON/kernels/NEFillBorderKernel.h"
+#include "tests/Globals.h"
+#include "tests/NEON/Accessor.h"
+#include "tests/datasets/BorderModeDataset.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(NEON)
+TEST_SUITE(FillBorder)
+
+// *INDENT-OFF*
+// clang-format off
+const auto PaddingSizesDataset = concat(concat(
+ framework::dataset::make("PaddingSize", PaddingSize{ 0 }),
+ framework::dataset::make("PaddingSize", PaddingSize{ 1, 0, 1, 2 })),
+ framework::dataset::make("PaddingSize", PaddingSize{ 10 }));
+
+const auto BorderSizesDataset = framework::dataset::make("BorderSize", 0, 6);
+
+DATA_TEST_CASE(FillBorder, framework::DatasetMode::ALL, combine(combine(combine(combine(
+ datasets::SmallShapes(),
+ datasets::BorderModes()),
+ BorderSizesDataset),
+ PaddingSizesDataset),
+ framework::dataset::make("DataType", DataType::U8)),
+ shape, border_mode, size, padding, data_type)
+// clang-format on
+// *INDENT-ON*
+{
+ BorderSize border_size{ static_cast<unsigned int>(size) };
+
+ std::mt19937 generator(library->seed());
+ std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+ const uint8_t border_value = distribution_u8(generator);
+ const uint8_t tensor_value = distribution_u8(generator);
+
+ // Create tensors
+ Tensor src = create_tensor<Tensor>(shape, data_type);
+
+ 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(Accessor(src), distribution, 0);
+
+ // Create and configure kernel
+ NEFillBorderKernel fill_border;
+ fill_border.configure(&src, border_size, border_mode, border_value);
+
+ // Run kernel
+ fill_border.run(fill_border.window(), ThreadInfo());
+
+ // Validate border
+ border_size.limit(padding);
+ validate(Accessor(src), border_size, border_mode, &border_value);
+
+ // Validate tensor
+ validate(Accessor(src), &tensor_value);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/NEON/FixedPoint/FixedPoint.cpp b/tests/validation/NEON/FixedPoint/FixedPoint.cpp
new file mode 100644
index 0000000..3b25b05
--- /dev/null
+++ b/tests/validation/NEON/FixedPoint/FixedPoint.cpp
@@ -0,0 +1,159 @@
+/*
+ * 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 "FixedPointTarget.h"
+
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/NEON/Accessor.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/FixedPointFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+constexpr AbsoluteTolerance<float> tolerance_exp_qs8(0.0f); /**< Tolerance value for comparing reference's output against implementation's output (exponential) for DataType::QS8 */
+constexpr AbsoluteTolerance<float> tolerance_exp_qs16(1.0f); /**< Tolerance value for comparing reference's output against implementation's output (exponential) for DataType::QS16 */
+constexpr AbsoluteTolerance<float> tolerance_invsqrt_qs8(4.0f); /**< Tolerance value for comparing reference's output against implementation's output (inverse square-root) for DataType::QS8 */
+constexpr AbsoluteTolerance<float> tolerance_invsqrt_qs16(5.0f); /**< Tolerance value for comparing reference's output against implementation's output (inverse square-root) for DataType::QS16 */
+constexpr AbsoluteTolerance<float> tolerance_log_qs8(5.0f); /**< Tolerance value for comparing reference's output against implementation's output (logarithm) for DataType::QS8 */
+constexpr AbsoluteTolerance<float> tolerance_log_qs16(7.0f); /**< Tolerance value for comparing reference's output against implementation's output (logarithm) for DataType::QS16 */
+constexpr AbsoluteTolerance<float> tolerance_reciprocal_qs8(3); /**< Tolerance value for comparing reference's output against implementation's output (reciprocal) for DataType::QS8 */
+constexpr AbsoluteTolerance<float> tolerance_reciprocal_qs16(11.0f); /**< Tolerance value for comparing reference's output against implementation's output (reciprocal) for DataType::QS16 */
+} // namespace
+
+TEST_SUITE(NEON)
+TEST_SUITE(FixedPoint)
+template <typename T>
+using NEFixedPointFixture = FixedPointValidationFixture<Tensor, Accessor, T>;
+
+TEST_SUITE(QS8)
+TEST_SUITE(Exp)
+
+FIXTURE_DATA_TEST_CASE(RunSmall, NEFixedPointFixture<int8_t>, framework::DatasetMode::ALL, combine(combine(combine(datasets::Small1DShape(), framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FixedPointOp", FixedPointOp::EXP)),
+ framework::dataset::make("FractionalBits", 1, 7)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference, tolerance_exp_qs8, 0);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(Invsqrt)
+FIXTURE_DATA_TEST_CASE(RunSmall, NEFixedPointFixture<int8_t>, framework::DatasetMode::ALL, combine(combine(combine(datasets::Small1DShape(), framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FixedPointOp", FixedPointOp::INV_SQRT)),
+ framework::dataset::make("FractionalBits", 1, 6)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference, tolerance_invsqrt_qs8, 0);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(Log)
+FIXTURE_DATA_TEST_CASE(RunSmall, NEFixedPointFixture<int8_t>, framework::DatasetMode::ALL, combine(combine(combine(datasets::Small1DShape(), framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FixedPointOp", FixedPointOp::LOG)),
+ framework::dataset::make("FractionalBits", 3, 6)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference, tolerance_log_qs8, 0);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(Reciprocal)
+FIXTURE_DATA_TEST_CASE(RunSmall, NEFixedPointFixture<int8_t>, framework::DatasetMode::ALL, combine(combine(combine(datasets::Small1DShape(), framework::dataset::make("DataType",
+ DataType::QS8)),
+ framework::dataset::make("FixedPointOp", FixedPointOp::RECIPROCAL)),
+ framework::dataset::make("FractionalBits", 1, 6)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference, tolerance_reciprocal_qs8, 0);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+TEST_SUITE(QS16)
+TEST_SUITE(Exp)
+FIXTURE_DATA_TEST_CASE(RunSmall, NEFixedPointFixture<int16_t>, framework::DatasetMode::ALL, combine(combine(combine(datasets::Small1DShape(), framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FixedPointOp", FixedPointOp::EXP)),
+ framework::dataset::make("FractionalBits", 1, 15)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference, tolerance_exp_qs16, 0);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(Invsqrt)
+FIXTURE_DATA_TEST_CASE(RunSmall, NEFixedPointFixture<int16_t>, framework::DatasetMode::ALL, combine(combine(combine(framework::dataset::make("Shape", TensorShape(8192U)),
+ framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FixedPointOp", FixedPointOp::INV_SQRT)),
+ framework::dataset::make("FractionalBits", 1, 14)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference, tolerance_invsqrt_qs16, 0);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(Log)
+FIXTURE_DATA_TEST_CASE(RunSmall, NEFixedPointFixture<int16_t>, framework::DatasetMode::ALL, combine(combine(combine(datasets::Small1DShape(), framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FixedPointOp", FixedPointOp::LOG)),
+ framework::dataset::make("FractionalBits", 4, 14)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference, tolerance_log_qs16, 0);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(Reciprocal)
+FIXTURE_DATA_TEST_CASE(RunSmall, NEFixedPointFixture<int16_t>, framework::DatasetMode::ALL, combine(combine(combine(datasets::Small1DShape(), framework::dataset::make("DataType",
+ DataType::QS16)),
+ framework::dataset::make("FixedPointOp", FixedPointOp::RECIPROCAL)),
+ framework::dataset::make("FractionalBits", 1, 14)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference, tolerance_reciprocal_qs16, 0);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/NEON/FixedPoint/FixedPointTarget.h b/tests/validation/NEON/FixedPoint/FixedPointTarget.h
new file mode 100644
index 0000000..fa1d97f
--- /dev/null
+++ b/tests/validation/NEON/FixedPoint/FixedPointTarget.h
@@ -0,0 +1,223 @@
+/*
+ * 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.
+ */
+#ifndef ARM_COMPUTE_TEST_FIXED_POINT_NEON_TARGET
+#define ARM_COMPUTE_TEST_FIXED_POINT_NEON_TARGET
+
+#include "arm_compute/core/NEON/NEFixedPoint.h"
+
+#include "tests/Globals.h"
+#include "tests/Types.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+template <typename TensorType, typename AccessorType, typename T>
+void compute_target_impl(const TensorShape &shape, DataType dt, FixedPointOp op, int fixed_point_position, TensorType &src, TensorType &dst)
+{
+ Window window;
+
+ switch(dt)
+ {
+ case DataType::QS8:
+ {
+ constexpr unsigned int num_elems_processed_per_iteration = 16;
+ window = calculate_max_window(*src.info(), Steps(num_elems_processed_per_iteration));
+ AccessWindowHorizontal input_access(src.info(), 0, num_elems_processed_per_iteration);
+ AccessWindowHorizontal output_access(dst.info(), 0, num_elems_processed_per_iteration);
+ break;
+ }
+ case DataType::QS16:
+ {
+ constexpr unsigned int num_elems_processed_per_iteration = 8;
+ window = calculate_max_window(*src.info(), Steps(num_elems_processed_per_iteration));
+ AccessWindowHorizontal input_access(src.info(), 0, num_elems_processed_per_iteration);
+ AccessWindowHorizontal output_access(dst.info(), 0, num_elems_processed_per_iteration);
+ break;
+ }
+ default:
+ ARM_COMPUTE_ERROR("Not Supported");
+ break;
+ }
+
+ int min;
+ int max;
+ switch(op)
+ {
+ case FixedPointOp::EXP:
+ {
+ // Fill tensors. Keep the range between [-1.0, 1.0) so the result won't
+ // overflow.
+ min = -(1 << (fixed_point_position - 1));
+ max = (1 << (fixed_point_position - 1));
+ break;
+ }
+ case FixedPointOp::INV_SQRT:
+ {
+ if(dt == DataType::QS8)
+ {
+ // Fill tensors. Keep the range between [1, 127).
+ min = 1;
+ max = 127;
+ }
+ else
+ {
+ // Fill tensors. Keep the range between [1, 0x7FFF)
+ min = 1;
+ max = 0x7FFF;
+ }
+ break;
+ }
+ case FixedPointOp::LOG:
+ {
+ if(dt == DataType::QS8)
+ {
+ // Fill tensors. Keep the range between [(1 << (fixed_point_position - 1), 63) so the result won't
+ // overflow. E.g. for Q2.5 ln(0.001) = -6.9, which cannot be represented.
+ min = (1 << (fixed_point_position - 1));
+ max = 0x3F;
+ }
+ else
+ {
+ // Fill tensors. Keep the range between [(1 << (fixed_point_position - 1), 0x3FFF) so the result won't
+ // overflow.
+ min = (1 << (fixed_point_position - 1));
+ max = 0x3FFF;
+ }
+ break;
+ }
+ case FixedPointOp::RECIPROCAL:
+ {
+ if(dt == DataType::QS8)
+ {
+ // Fill tensors. Keep the range between [15, 100) so the result won't
+ // overflow. E.g. for Q2.5 reciprocal(0.001) = 1000, which cannot be represented.
+ min = 15;
+ max = 0x7F;
+ }
+ else
+ {
+ // Fill tensors. Keep the range between [15, 0x7FFF) so the result won't
+ // overflow.
+ min = 15;
+ max = 0x7FFF;
+ }
+ break;
+ }
+ default:
+ ARM_COMPUTE_ERROR("Not Supported");
+ break;
+ }
+
+ std::uniform_int_distribution<> distribution(min, max);
+ library->fill(AccessorType(src), distribution, 0);
+
+ Iterator input(&src, window);
+ Iterator output(&dst, window);
+
+ const auto loop_function = [&](const Coordinates & id)
+ {
+ switch(dt)
+ {
+ case DataType::QS8:
+ {
+ const qint8x16_t qs8in = vld1q_s8(reinterpret_cast<const qint8_t *>(input.ptr()));
+ switch(op)
+ {
+ case FixedPointOp::EXP:
+ {
+ // Use saturated exp
+ vst1q_s8(reinterpret_cast<qint8_t *>(output.ptr()), vqexpq_qs8(qs8in, fixed_point_position));
+ break;
+ }
+ case FixedPointOp::INV_SQRT:
+ {
+ vst1q_s8(reinterpret_cast<qint8_t *>(output.ptr()), vqinvsqrtq_qs8(qs8in, fixed_point_position));
+ break;
+ }
+ case FixedPointOp::LOG:
+ {
+ vst1q_s8(reinterpret_cast<qint8_t *>(output.ptr()), vlogq_qs8(qs8in, fixed_point_position));
+ break;
+ }
+ case FixedPointOp::RECIPROCAL:
+ {
+ vst1q_s8(reinterpret_cast<qint8_t *>(output.ptr()), vrecipq_qs8(qs8in, fixed_point_position));
+ break;
+ }
+ default:
+ ARM_COMPUTE_ERROR("Not Supported");
+ break;
+ }
+ break;
+ }
+ case DataType::QS16:
+ {
+ const qint16x8_t qs16in = vld1q_qs16(reinterpret_cast<const qint16_t *>(input.ptr()));
+ switch(op)
+ {
+ case FixedPointOp::EXP:
+ {
+ // Use saturated exp
+ vst1q_qs16(reinterpret_cast<qint16_t *>(output.ptr()), vqexpq_qs16(qs16in, fixed_point_position));
+ break;
+ }
+ case FixedPointOp::INV_SQRT:
+ {
+ vst1q_qs16(reinterpret_cast<qint16_t *>(output.ptr()), vqinvsqrtq_qs16(qs16in, fixed_point_position));
+ break;
+ }
+ case FixedPointOp::LOG:
+ {
+ vst1q_qs16(reinterpret_cast<qint16_t *>(output.ptr()), vlogq_qs16(qs16in, fixed_point_position));
+ break;
+ }
+ case FixedPointOp::RECIPROCAL:
+ {
+ vst1q_qs16(reinterpret_cast<qint16_t *>(output.ptr()), vqrecipq_qs16(qs16in, fixed_point_position));
+ break;
+ }
+ default:
+ ARM_COMPUTE_ERROR("Not Supported");
+ break;
+ }
+ break;
+ }
+ default:
+ ARM_COMPUTE_ERROR("Not Supported");
+ break;
+ }
+ };
+
+ execute_window_loop(window, loop_function, input, output);
+}
+} // namespace
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_FIXED_POINT_NEON_TARGET */
diff --git a/tests/validation/NEON/Flatten.cpp b/tests/validation/NEON/Flatten.cpp
new file mode 100644
index 0000000..f8d8301
--- /dev/null
+++ b/tests/validation/NEON/Flatten.cpp
@@ -0,0 +1,119 @@
+/*
+ * 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/Allocator.h"
+#include "arm_compute/runtime/NEON/functions/NEFlattenLayer.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "tests/NEON/Accessor.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/FlattenLayerFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(NEON)
+TEST_SUITE(FlattenLayer)
+
+template <typename T>
+using NEFlattenLayerFixture = FlattenLayerValidationFixture<Tensor, Accessor, NEFlattenLayer, T>;
+
+TEST_SUITE(Float)
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, NEFlattenLayerFixture<float>, framework::DatasetMode::ALL, combine(framework::dataset::concat(datasets::Small3DShapes(), datasets::Small4DShapes()),
+ framework::dataset::make("DataType", DataType::F32)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, NEFlattenLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(framework::dataset::concat(datasets::Large3DShapes(), datasets::Large4DShapes()),
+ framework::dataset::make("DataType", DataType::F32)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference);
+}
+TEST_SUITE_END()
+
+#ifdef ARM_COMPUTE_ENABLE_FP16
+TEST_SUITE(FP16)
+FIXTURE_DATA_TEST_CASE(RunSmall, NEFlattenLayerFixture<half>, framework::DatasetMode::ALL, combine(framework::dataset::concat(datasets::Small3DShapes(), datasets::Small4DShapes()),
+ framework::dataset::make("DataType", DataType::F16)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, NEFlattenLayerFixture<half>, framework::DatasetMode::NIGHTLY, combine(framework::dataset::concat(datasets::Large3DShapes(), datasets::Large4DShapes()),
+ framework::dataset::make("DataType", DataType::F16)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference);
+}
+TEST_SUITE_END()
+#endif // ARM_COMPUTE_ENABLE_FP16
+TEST_SUITE_END()
+
+TEST_SUITE(Quantized)
+TEST_SUITE(QS8)
+FIXTURE_DATA_TEST_CASE(RunSmall, NEFlattenLayerFixture<int8_t>, framework::DatasetMode::ALL, combine(framework::dataset::concat(datasets::Small3DShapes(), datasets::Small4DShapes()),
+ framework::dataset::make("DataType", DataType::QS8)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, NEFlattenLayerFixture<int8_t>, framework::DatasetMode::NIGHTLY, combine(framework::dataset::concat(datasets::Large3DShapes(), datasets::Large4DShapes()),
+ framework::dataset::make("DataType", DataType::QS8)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(QS16)
+FIXTURE_DATA_TEST_CASE(RunSmall, NEFlattenLayerFixture<int16_t>, framework::DatasetMode::ALL, combine(framework::dataset::concat(datasets::Small3DShapes(), datasets::Small4DShapes()),
+ framework::dataset::make("DataType", DataType::QS16)))
+{
+ // Validate output
+ validate(Accessor(_target), _reference);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, NEFlattenLayerFixture<int16_t>, framework::DatasetMode::NIGHTLY, combine(framework::dataset::concat(datasets::Large3DShapes(), datasets::Large4DShapes()),
+ framework::dataset::make("DataType", DataType::QS16)))
+{
+ // Validate output
+ validate(Accessor(_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/NEON/GEMMLowp.cpp b/tests/validation/NEON/GEMMLowp.cpp
new file mode 100644
index 0000000..3d83f80
--- /dev/null
+++ b/tests/validation/NEON/GEMMLowp.cpp
@@ -0,0 +1,68 @@
+/*
+ * 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/NEON/functions/NEGEMMLowp.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/NEON/Accessor.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/GEMMLowpFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+constexpr AbsoluteTolerance<float> tolerance_f(0.001f); /**< Tolerance value for comparing reference's output against implementation's output for floating point data types */
+
+const auto data_mnk = framework::dataset::make("M", 8, 12) * framework::dataset::make("N", 8, 12) * framework::dataset::make("K", 8, 12);
+const auto data_offsets = framework::dataset::make("a", -3, 3) * framework::dataset::make("b", -1, 2) * framework::dataset::make("c", 1, 3) * framework::dataset::make("cm", 0,
+ 3)
+ * framework::dataset::make("shift", 0, 4);
+
+} // namespace
+
+TEST_SUITE(NEON)
+TEST_SUITE(GEMMLowp)
+
+TEST_SUITE(U8)
+using NEGEMMLowpOffsetFixture = GEMMLowpOffsetValidationFixture<Tensor, Accessor, NEGEMMLowp>;
+FIXTURE_DATA_TEST_CASE(RunSmall, NEGEMMLowpOffsetFixture, framework::DatasetMode::PRECOMMIT, data_mnk *data_offsets)
+{
+ // Validate output
+ validate(Accessor(_target), _reference, tolerance_f);
+}
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CL/GlobalPooling.cpp b/tests/validation/NEON/GlobalPoolingLayer.cpp
similarity index 68%
copy from tests/validation/CL/GlobalPooling.cpp
copy to tests/validation/NEON/GlobalPoolingLayer.cpp
index c5c9d00..37950b0 100644
--- a/tests/validation/CL/GlobalPooling.cpp
+++ b/tests/validation/NEON/GlobalPoolingLayer.cpp
@@ -22,10 +22,10 @@
* 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 "arm_compute/runtime/NEON/functions/NEPoolingLayer.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/NEON/Accessor.h"
#include "tests/PaddingCalculator.h"
#include "tests/datasets/PoolingTypesDataset.h"
#include "tests/datasets/ShapeDatasets.h"
@@ -47,31 +47,21 @@
const auto GlobalPoolingLayerDataset = combine(datasets::GlobalPoolingShapes(), datasets::PoolingTypes());
/** Input data set for quantized data types */
-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_f32(0.001f); /**< Tolerance value for comparing reference's output against implementation's output for FP32 types */
} // namespace
-TEST_SUITE(CL)
+TEST_SUITE(NEON)
TEST_SUITE(GlobalPoolingLayer)
template <typename T>
-using CLGlobalPoolingLayerFixture = GlobalPoolingLayerValidationFixture<CLTensor, CLAccessor, CLPoolingLayer, T>;
+using NEGlobalPoolingLayerFixture = GlobalPoolingLayerValidationFixture<Tensor, Accessor, NEPoolingLayer, T>;
TEST_SUITE(Float)
TEST_SUITE(FP32)
-FIXTURE_DATA_TEST_CASE(RunGlobalPooling, CLGlobalPoolingLayerFixture<float>, framework::DatasetMode::ALL, combine(GlobalPoolingLayerDataset, framework::dataset::make("DataType", DataType::F32)))
+FIXTURE_DATA_TEST_CASE(RunGlobalPooling, NEGlobalPoolingLayerFixture<float>, framework::DatasetMode::ALL, combine(GlobalPoolingLayerDataset, framework::dataset::make("DataType", DataType::F32)))
{
// Validate output
- validate(CLAccessor(_target), _reference, tolerance_f32);
-}
-TEST_SUITE_END()
-
-TEST_SUITE(FP16)
-FIXTURE_DATA_TEST_CASE(RunGlobalPooling, CLGlobalPoolingLayerFixture<half>, framework::DatasetMode::ALL, combine(GlobalPoolingLayerDataset, framework::dataset::make("DataType",
- DataType::F16)))
-{
- // Validate output
- validate(CLAccessor(_target), _reference, tolerance_f16);
+ validate(Accessor(_target), _reference, tolerance_f32);
}
TEST_SUITE_END()
TEST_SUITE_END()
diff --git a/tests/validation/NEON/PoolingLayer.cpp b/tests/validation/NEON/PoolingLayer.cpp
index 27f6aab..591c4be 100644
--- a/tests/validation/NEON/PoolingLayer.cpp
+++ b/tests/validation/NEON/PoolingLayer.cpp
@@ -44,7 +44,7 @@
namespace
{
/** Input data set for float data types */
-const auto PoolingLayerDatasetFP = combine(combine(datasets::PoolingTypes(), framework::dataset::make("PoolingSize", { 2, 3, 7 })),
+const auto PoolingLayerDatasetFP = combine(combine(datasets::PoolingTypes(), framework::dataset::make("PoolingSize", { 2, 3, 7, 9 })),
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 */
diff --git a/tests/validation/NEON/WarpAffine.cpp b/tests/validation/NEON/WarpAffine.cpp
new file mode 100644
index 0000000..678549d
--- /dev/null
+++ b/tests/validation/NEON/WarpAffine.cpp
@@ -0,0 +1,121 @@
+/*
+ * 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/NEON/functions/NEWarpAffine.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/NEON/Accessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/BorderModeDataset.h"
+#include "tests/datasets/InterpolationPolicyDataset.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/CPP/Utils.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/WarpAffineFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Tolerance */
+constexpr AbsoluteTolerance<uint8_t> tolerance(1);
+} // namespace
+
+TEST_SUITE(NEON)
+TEST_SUITE(WarpAffine)
+
+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)
+{
+ // Generate a random constant value if border_mode is constant
+ std::mt19937 gen(library->seed());
+ std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+ uint8_t constant_border_value = distribution_u8(gen);
+
+ // Create the matrix
+ std::array<float, 6> matrix{ {} };
+ fill_warp_matrix<6>(matrix);
+
+ // Create tensors
+ Tensor src = create_tensor<Tensor>(shape, data_type);
+ Tensor dst = create_tensor<Tensor>(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
+ NEWarpAffine warp_affine;
+ warp_affine.configure(&src, &dst, matrix.data(), policy, border_mode, constant_border_value);
+
+ // Validate valid region
+ const ValidRegion valid_region = shape_to_valid_region(shape);
+
+ validate(dst.info()->valid_region(), valid_region);
+
+ // Validate padding
+ PaddingCalculator calculator(shape.x(), 1);
+ calculator.set_border_mode(border_mode);
+ calculator.set_border_size(1);
+
+ const PaddingSize read_padding(1);
+ const PaddingSize write_padding = calculator.required_padding();
+
+ validate(src.info()->padding(), read_padding);
+ validate(dst.info()->padding(), write_padding);
+}
+
+template <typename T>
+using NEWarpAffineFixture = WarpAffineValidationFixture<Tensor, Accessor, NEWarpAffine, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, NEWarpAffineFixture<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()))
+{
+ // Validate output
+ validate(Accessor(_target), _reference, _valid_mask, tolerance, 0.02f);
+}
+DISABLED_FIXTURE_DATA_TEST_CASE(RunLarge, NEWarpAffineFixture<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()))
+{
+ // Validate output
+ validate(Accessor(_target), _reference, _valid_mask, tolerance, 0.02f);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/NEON/WarpPerspective.cpp b/tests/validation/NEON/WarpPerspective.cpp
new file mode 100644
index 0000000..45d3a0b
--- /dev/null
+++ b/tests/validation/NEON/WarpPerspective.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/NEON/functions/NEWarpPerspective.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/NEON/Accessor.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/Validation.h"
+#include "tests/validation/fixtures/WarpPerspectiveFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+constexpr AbsoluteTolerance<uint8_t> tolerance_value(1);
+constexpr float tolerance_number = 0.2f;
+} // namespace
+
+TEST_SUITE(NEON)
+TEST_SUITE(WarpPerspective)
+
+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)
+{
+ uint8_t constant_border_value = 0;
+
+ // Generate a random constant value if border_mode is constant
+ if(border_mode == BorderMode::CONSTANT)
+ {
+ std::mt19937 gen(library->seed());
+ std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+ constant_border_value = distribution_u8(gen);
+ }
+
+ // Create the matrix
+ std::array<float, 9> matrix = { { 0 } };
+ fill_warp_matrix<9>(matrix);
+
+ // Create tensors
+ Tensor src = create_tensor<Tensor>(shape, data_type);
+ Tensor dst = create_tensor<Tensor>(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
+ NEWarpPerspective 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(), 1);
+ calculator.set_border_mode(border_mode);
+ calculator.set_border_size(1);
+
+ const PaddingSize read_padding(1);
+ const PaddingSize write_padding = calculator.required_padding();
+
+ validate(src.info()->padding(), read_padding);
+ validate(dst.info()->padding(), write_padding);
+}
+
+template <typename T>
+using NEWarpPerspectiveFixture = WarpPerspectiveValidationFixture<Tensor, Accessor, NEWarpPerspective, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, NEWarpPerspectiveFixture<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 the valid mask Tensor
+ RawTensor valid_mask(_reference.shape(), _reference.data_type());
+
+ validate(Accessor(_target), _reference, valid_mask, tolerance_value, tolerance_number);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, NEWarpPerspectiveFixture<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 the valid mask Tensor
+ RawTensor valid_mask{ _reference.shape(), _reference.data_type() };
+
+ validate(Accessor(_target), _reference, valid_mask, tolerance_value, tolerance_number);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/UNIT/FixedPoint.cpp b/tests/validation/UNIT/FixedPoint.cpp
new file mode 100644
index 0000000..251f5a8
--- /dev/null
+++ b/tests/validation/UNIT/FixedPoint.cpp
@@ -0,0 +1,210 @@
+/*
+ * 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 "tests/validation/FixedPoint.h"
+
+#include "tests/Globals.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+const auto FuncNamesDataset = framework::dataset::make("FunctionNames", { FixedPointOp::ADD,
+ FixedPointOp::SUB,
+ FixedPointOp::MUL,
+ FixedPointOp::EXP,
+ FixedPointOp::LOG,
+ FixedPointOp::INV_SQRT
+ });
+
+template <typename T>
+void load_array_from_numpy(const std::string &file, std::vector<unsigned long> &shape, std::vector<T> &data) // NOLINT
+{
+ try
+ {
+ npy::LoadArrayFromNumpy(file, shape, data);
+ }
+ catch(const std::runtime_error &e)
+ {
+ throw framework::FileNotFound("Could not load npy file: " + file + " (" + e.what() + ")");
+ }
+}
+} // namespace
+
+TEST_SUITE(UNIT)
+TEST_SUITE(FixedPoint)
+
+// *INDENT-OFF*
+// clang-format off
+DATA_TEST_CASE(FixedPointQS8Inputs, framework::DatasetMode::ALL, combine(
+ FuncNamesDataset,
+ framework::dataset::make("FractionalBits", 1, 7)),
+ func_name, frac_bits)
+// clang-format on
+// *INDENT-ON*
+{
+ std::vector<double> data;
+ std::vector<unsigned long> shape; //NOLINT
+
+ std::string func_name_lower = to_string(func_name);
+ std::transform(func_name_lower.begin(), func_name_lower.end(), func_name_lower.begin(), ::tolower);
+
+ const std::string inputs_file = library->path()
+ + "fixed_point/"
+ + func_name_lower
+ + "_Q8."
+ + support::cpp11::to_string(frac_bits)
+ + ".in.npy";
+
+ load_array_from_numpy(inputs_file, shape, data);
+
+ // Values stored as doubles so reinterpret as floats
+ const auto *float_val = reinterpret_cast<float *>(&data[0]);
+ const size_t num_elements = data.size() * sizeof(double) / sizeof(float);
+
+ for(unsigned int i = 0; i < num_elements; ++i)
+ {
+ // Convert to fixed point
+ fixed_point_arithmetic::fixed_point<int8_t> in_val(float_val[i], frac_bits);
+
+ // Check that the value didn't change
+ ARM_COMPUTE_EXPECT(static_cast<float>(in_val) == float_val[i], framework::LogLevel::ERRORS);
+ }
+}
+
+// The last input argument specifies the expected number of failures for a
+// given combination of (function name, number of fractional bits) as defined
+// by the first two arguments.
+
+// *INDENT-OFF*
+// clang-format off
+DATA_TEST_CASE(FixedPointQS8Outputs, framework::DatasetMode::ALL, zip(combine(
+ FuncNamesDataset,
+ framework::dataset::make("FractionalBits", 1, 7)),
+ framework::dataset::make("ExpectedFailures", { 0, 0, 0, 0, 0, 0,
+ 0, 0, 0, 0, 0, 0,
+ 0, 0, 0, 0, 0, 0,
+ 7, 8, 13, 2, 0, 0,
+ 0, 0, 0, 0, 0, 0,
+ 0, 0, 0, 5, 33, 96 })),
+ func_name, frac_bits, expected_failures)
+// clang-format on
+// *INDENT-ON*
+{
+ std::vector<double> in_data;
+ std::vector<unsigned long> in_shape; //NOLINT
+
+ std::vector<double> out_data;
+ std::vector<unsigned long> out_shape; //NOLINT
+
+ std::string func_name_lower = to_string(func_name);
+ std::transform(func_name_lower.begin(), func_name_lower.end(), func_name_lower.begin(), ::tolower);
+
+ const std::string base_file_name = library->path()
+ + "fixed_point/"
+ + func_name_lower
+ + "_Q8."
+ + support::cpp11::to_string(frac_bits);
+
+ const std::string inputs_file = base_file_name + ".in.npy";
+ const std::string reference_file = base_file_name + ".out.npy";
+
+ load_array_from_numpy(inputs_file, in_shape, in_data);
+ load_array_from_numpy(reference_file, out_shape, out_data);
+
+ ARM_COMPUTE_EXPECT(in_shape.front() == out_shape.front(), framework::LogLevel::ERRORS);
+
+ const float step_size = std::pow(2.f, -frac_bits);
+ int64_t num_mismatches = 0;
+
+ // Values stored as doubles so reinterpret as floats
+ const auto *float_val = reinterpret_cast<float *>(&in_data[0]);
+ const auto *ref_val = reinterpret_cast<float *>(&out_data[0]);
+
+ const size_t num_elements = in_data.size() * sizeof(double) / sizeof(float);
+
+ for(unsigned int i = 0; i < num_elements; ++i)
+ {
+ fixed_point_arithmetic::fixed_point<int8_t> in_val(float_val[i], frac_bits);
+ fixed_point_arithmetic::fixed_point<int8_t> out_val(0.f, frac_bits);
+
+ float tolerance = 0.f;
+
+ if(func_name == FixedPointOp::ADD)
+ {
+ out_val = in_val + in_val;
+ }
+ else if(func_name == FixedPointOp::SUB)
+ {
+ out_val = in_val - in_val; //NOLINT
+ }
+ else if(func_name == FixedPointOp::MUL)
+ {
+ tolerance = 1.f * step_size;
+ out_val = in_val * in_val;
+ }
+ else if(func_name == FixedPointOp::EXP)
+ {
+ tolerance = 2.f * step_size;
+ out_val = fixed_point_arithmetic::exp(in_val);
+ }
+ else if(func_name == FixedPointOp::LOG)
+ {
+ tolerance = 4.f * step_size;
+ out_val = fixed_point_arithmetic::log(in_val);
+ }
+ else if(func_name == FixedPointOp::INV_SQRT)
+ {
+ tolerance = 5.f * step_size;
+ out_val = fixed_point_arithmetic::inv_sqrt(in_val);
+ }
+
+ if(std::abs(static_cast<float>(out_val) - ref_val[i]) > tolerance)
+ {
+ ARM_COMPUTE_TEST_INFO("input = " << in_val);
+ ARM_COMPUTE_TEST_INFO("output = " << out_val);
+ ARM_COMPUTE_TEST_INFO("reference = " << ref_val[i]);
+ ARM_COMPUTE_TEST_INFO("tolerance = " << tolerance);
+
+ ARM_COMPUTE_TEST_INFO((std::abs(static_cast<float>(out_val) - ref_val[i]) <= tolerance));
+
+ ++num_mismatches;
+ }
+ }
+
+ ARM_COMPUTE_EXPECT(num_mismatches == expected_failures, framework::LogLevel::ERRORS);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/UNIT/TensorInfo.cpp b/tests/validation/UNIT/TensorInfo.cpp
new file mode 100644
index 0000000..2a6c336
--- /dev/null
+++ b/tests/validation/UNIT/TensorInfo.cpp
@@ -0,0 +1,89 @@
+/*
+ * 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/TensorInfo.h"
+#include "arm_compute/core/Types.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "utils/TypePrinter.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(UNIT)
+TEST_SUITE(TensorInfoValidation)
+
+// *INDENT-OFF*
+// clang-format off
+DATA_TEST_CASE(AutoPadding, framework::DatasetMode::ALL, zip(zip(zip(
+ framework::dataset::make("TensorShape", {
+ TensorShape{},
+ TensorShape{ 10U },
+ TensorShape{ 10U, 10U },
+ TensorShape{ 10U, 10U, 10U },
+ TensorShape{ 10U, 10U, 10U, 10U },
+ TensorShape{ 10U, 10U, 10U, 10U, 10U },
+ TensorShape{ 10U, 10U, 10U, 10U, 10U, 10U }}),
+ framework::dataset::make("PaddingSize", {
+ PaddingSize{ 0, 0, 0, 0 },
+ PaddingSize{ 0, 36, 0, 4 },
+ PaddingSize{ 4, 36, 4, 4 },
+ PaddingSize{ 4, 36, 4, 4 },
+ PaddingSize{ 4, 36, 4, 4 },
+ PaddingSize{ 4, 36, 4, 4 },
+ PaddingSize{ 4, 36, 4, 4 }})),
+ framework::dataset::make("Strides", {
+ Strides{},
+ Strides{ 1U, 50U },
+ Strides{ 1U, 50U },
+ Strides{ 1U, 50U, 900U },
+ Strides{ 1U, 50U, 900U, 9000U },
+ Strides{ 1U, 50U, 900U, 9000U, 90000U },
+ Strides{ 1U, 50U, 900U, 9000U, 90000U, 900000U }})),
+ framework::dataset::make("Offset", { 0U, 4U, 204U, 204U, 204U, 204U, 204U })),
+ shape, auto_padding, strides, offset)
+{
+ TensorInfo info{ shape, Format::U8 };
+
+ ARM_COMPUTE_EXPECT(!info.has_padding(), framework::LogLevel::ERRORS);
+
+ info.auto_padding();
+
+ validate(info.padding(), auto_padding);
+
+ ARM_COMPUTE_EXPECT(compare_dimensions(info.strides_in_bytes(), strides), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(info.offset_first_element_in_bytes() == offset, framework::LogLevel::ERRORS);
+}
+// clang-format on
+// *INDENT-ON*
+
+TEST_SUITE_END() // TensorInfoValidation
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/UNIT/TensorShape.cpp b/tests/validation/UNIT/TensorShape.cpp
new file mode 100644
index 0000000..31f95e3
--- /dev/null
+++ b/tests/validation/UNIT/TensorShape.cpp
@@ -0,0 +1,73 @@
+/*
+ * 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 "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(UNIT)
+TEST_SUITE(TensorShapeValidation)
+
+// *INDENT-OFF*
+// clang-format off
+DATA_TEST_CASE(Construction, framework::DatasetMode::ALL, zip(zip(
+ framework::dataset::make("TensorShape", {
+ TensorShape{},
+ TensorShape{ 1U },
+ TensorShape{ 2U },
+ TensorShape{ 2U, 3U },
+ TensorShape{ 2U, 3U, 5U },
+ TensorShape{ 2U, 3U, 5U, 7U },
+ TensorShape{ 2U, 3U, 5U, 7U, 11U },
+ TensorShape{ 2U, 3U, 5U, 7U, 11U, 13U }}),
+ framework::dataset::make("NumDimensions", { 0U, 1U, 1U, 2U, 3U, 4U, 5U, 6U })),
+ framework::dataset::make("TotalSize", { 0U, 1U, 2U, 6U, 30U, 210U, 2310U, 30030U })),
+ shape, num_dimensions, total_size)
+{
+ ARM_COMPUTE_EXPECT(shape.num_dimensions() == num_dimensions, framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(shape.total_size() == total_size, framework::LogLevel::ERRORS);
+}
+// clang-format on
+// *INDENT-ON*
+
+DATA_TEST_CASE(SetEmpty, framework::DatasetMode::ALL, framework::dataset::make("Dimension", { 0U, 1U, 2U, 3U, 4U, 5U }), dimension)
+{
+ TensorShape shape;
+
+ shape.set(dimension, 10);
+
+ ARM_COMPUTE_EXPECT(shape.num_dimensions() == dimension + 1, framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(shape.total_size() == 10, framework::LogLevel::ERRORS);
+}
+
+TEST_SUITE_END() // TensorShapeValidation
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/UNIT/Utils.cpp b/tests/validation/UNIT/Utils.cpp
new file mode 100644
index 0000000..71f39a0
--- /dev/null
+++ b/tests/validation/UNIT/Utils.cpp
@@ -0,0 +1,74 @@
+/*
+ * 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 "tests/Utils.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "utils/TypePrinter.h"
+
+#include <stdexcept>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+TEST_SUITE(UNIT)
+TEST_SUITE(Utils)
+
+DATA_TEST_CASE(RoundHalfUp, framework::DatasetMode::ALL, zip(framework::dataset::make("FloatIn", { 1.f, 1.2f, 1.5f, 2.5f, 2.9f, -3.f, -3.5f, -3.8f, -4.3f, -4.5f }),
+ framework::dataset::make("FloatOut", { 1.f, 1.f, 2.f, 3.f, 3.f, -3.f, -3.f, -4.f, -4.f, -4.f })),
+ value, result)
+{
+ ARM_COMPUTE_EXPECT(round_half_up(value) == result, framework::LogLevel::ERRORS);
+}
+
+DATA_TEST_CASE(RoundHalfEven, framework::DatasetMode::ALL, zip(framework::dataset::make("FloatIn", { 1.f, 1.2f, 1.5f, 2.5f, 2.9f, -3.f, -3.5f, -3.8f, -4.3f, -4.5f }),
+ framework::dataset::make("FloatOut", { 1.f, 1.f, 2.f, 2.f, 3.f, -3.f, -4.f, -4.f, -4.f, -4.f })),
+ value, result)
+{
+ ARM_COMPUTE_EXPECT(round_half_even(value) == result, framework::LogLevel::ERRORS);
+}
+
+DATA_TEST_CASE(Index2Coord, framework::DatasetMode::ALL, zip(zip(framework::dataset::make("Shape", { TensorShape{ 1U }, TensorShape{ 2U }, TensorShape{ 2U, 3U } }), framework::dataset::make("Index", { 0, 1, 2 })),
+ framework::dataset::make("Coordinates", { Coordinates{ 0 }, Coordinates{ 1 }, Coordinates{ 0, 1 } })),
+ shape, index, ref_coordinate)
+{
+ Coordinates coordinate = index2coord(shape, index);
+
+ ARM_COMPUTE_EXPECT(compare_dimensions(coordinate, ref_coordinate), framework::LogLevel::ERRORS);
+}
+
+DATA_TEST_CASE(Coord2Index, framework::DatasetMode::ALL, zip(zip(framework::dataset::make("Shape", { TensorShape{ 1U }, TensorShape{ 2U }, TensorShape{ 2U, 3U } }),
+ framework::dataset::make("Coordinates", { Coordinates{ 0 }, Coordinates{ 1 }, Coordinates{ 0, 1 } })),
+ framework::dataset::make("Index", { 0, 1, 2 })),
+ shape, coordinate, ref_index)
+{
+ int index = coord2index(shape, coordinate);
+
+ ARM_COMPUTE_EXPECT(index == ref_index, framework::LogLevel::ERRORS);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
diff --git a/tests/validation/Validation.h b/tests/validation/Validation.h
index 5e5dab0..f220224 100644
--- a/tests/validation/Validation.h
+++ b/tests/validation/Validation.h
@@ -188,6 +188,19 @@
template <typename T, typename U = AbsoluteTolerance<T>>
void validate(const IAccessor &tensor, const SimpleTensor<T> &reference, const ValidRegion &valid_region, U tolerance_value = U(), float tolerance_number = 0.f);
+/** Validate tensors with valid mask.
+ *
+ * - Dimensionality has to be the same.
+ * - All values have to match.
+ *
+ * @note: wrap_range allows cases where reference tensor rounds up to the wrapping point, causing it to wrap around to
+ * zero while the test tensor stays at wrapping point to pass. This may permit true erroneous cases (difference between
+ * reference tensor and test tensor is multiple of wrap_range), but such errors would be detected by
+ * other test cases.
+ */
+template <typename T, typename U = AbsoluteTolerance<T>>
+void validate(const IAccessor &tensor, const SimpleTensor<T> &reference, const SimpleTensor<T> &valid_mask, U tolerance_value = U(), float tolerance_number = 0.f);
+
/** Validate tensors against constant value.
*
* - All values have to match.
@@ -435,6 +448,71 @@
}
template <typename T, typename U>
+void validate(const IAccessor &tensor, const SimpleTensor<T> &reference, const SimpleTensor<T> &valid_mask, U tolerance_value, float tolerance_number)
+{
+ int64_t num_mismatches = 0;
+ int64_t num_elements = 0;
+
+ ARM_COMPUTE_EXPECT_EQUAL(tensor.element_size(), reference.element_size(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT_EQUAL(tensor.data_type(), reference.data_type(), framework::LogLevel::ERRORS);
+
+ if(reference.format() != Format::UNKNOWN)
+ {
+ ARM_COMPUTE_EXPECT_EQUAL(tensor.format(), reference.format(), framework::LogLevel::ERRORS);
+ }
+
+ ARM_COMPUTE_EXPECT_EQUAL(tensor.num_channels(), reference.num_channels(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(compare_dimensions(tensor.shape(), reference.shape()), framework::LogLevel::ERRORS);
+
+ const int min_elements = std::min(tensor.num_elements(), reference.num_elements());
+ const int min_channels = std::min(tensor.num_channels(), reference.num_channels());
+
+ // Iterate over all elements within valid region, e.g. U8, S16, RGB888, ...
+ for(int element_idx = 0; element_idx < min_elements; ++element_idx)
+ {
+ const Coordinates id = index2coord(reference.shape(), element_idx);
+
+ if(valid_mask[element_idx] == 1)
+ {
+ // Iterate over all channels within one element
+ for(int c = 0; c < min_channels; ++c)
+ {
+ const T &target_value = reinterpret_cast<const T *>(tensor(id))[c];
+ const T &reference_value = reinterpret_cast<const T *>(reference(id))[c];
+
+ if(!compare<U>(target_value, reference_value, tolerance_value))
+ {
+ ARM_COMPUTE_TEST_INFO("id = " << id);
+ ARM_COMPUTE_TEST_INFO("channel = " << c);
+ ARM_COMPUTE_TEST_INFO("target = " << std::setprecision(5) << framework::make_printable(target_value));
+ ARM_COMPUTE_TEST_INFO("reference = " << std::setprecision(5) << framework::make_printable(reference_value));
+ ARM_COMPUTE_TEST_INFO("tolerance = " << std::setprecision(5) << framework::make_printable(static_cast<typename U::value_type>(tolerance_value)));
+ framework::ARM_COMPUTE_PRINT_INFO();
+
+ ++num_mismatches;
+ }
+
+ ++num_elements;
+ }
+ }
+ else
+ {
+ ++num_elements;
+ }
+ }
+
+ if(num_elements > 0)
+ {
+ const int64_t absolute_tolerance_number = tolerance_number * num_elements;
+ const float percent_mismatches = static_cast<float>(num_mismatches) / num_elements * 100.f;
+
+ ARM_COMPUTE_TEST_INFO(num_mismatches << " values (" << std::fixed << std::setprecision(2) << percent_mismatches
+ << "%) mismatched (maximum tolerated " << std::setprecision(2) << tolerance_number << "%)");
+ ARM_COMPUTE_EXPECT(num_mismatches <= absolute_tolerance_number, framework::LogLevel::ERRORS);
+ }
+}
+
+template <typename T, typename U>
bool validate(T target, T reference, U tolerance)
{
ARM_COMPUTE_TEST_INFO("reference = " << std::setprecision(5) << framework::make_printable(reference));
diff --git a/tests/validation/fixtures/ConvolutionLayerFixture.h b/tests/validation/fixtures/ConvolutionLayerFixture.h
index 434291b..fe20699 100644
--- a/tests/validation/fixtures/ConvolutionLayerFixture.h
+++ b/tests/validation/fixtures/ConvolutionLayerFixture.h
@@ -88,7 +88,12 @@
{
// Check if its a "fully connected" convolution
const bool is_fully_connected_convolution = (output_shape.x() == 1 && output_shape.y() == 1);
- const bool is_optimised = std::is_same<FunctionType, NEConvolutionLayer>::value && NEScheduler::get().cpu_info().CPU >= CPUTarget::ARMV7 && data_type == DataType::F32;
+ bool is_optimised = false;
+#if defined(__arm__)
+ is_optimised = std::is_same<FunctionType, NEConvolutionLayer>::value && NEScheduler::get().cpu_info().CPU == CPUTarget::ARMV7 && data_type == DataType::F32;
+#elif defined(__aarch64__)
+ is_optimised = std::is_same<FunctionType, NEConvolutionLayer>::value && NEScheduler::get().cpu_info().CPU >= CPUTarget::ARMV8 && data_type == DataType::F32;
+#endif /* defined(__arm__) || defined(__aarch64__) */
reshaped_weights_shape.collapse(3);
@@ -143,7 +148,12 @@
if(!reshape_weights)
{
const bool is_fully_connected_convolution = (output_shape.x() == 1 && output_shape.y() == 1);
- const bool is_optimised = std::is_same<FunctionType, NEConvolutionLayer>::value && NEScheduler::get().cpu_info().CPU >= CPUTarget::ARMV7 && data_type == DataType::F32;
+ bool is_optimised = false;
+#if defined(__arm__)
+ is_optimised = std::is_same<FunctionType, NEConvolutionLayer>::value && NEScheduler::get().cpu_info().CPU == CPUTarget::ARMV7 && data_type == DataType::F32;
+#elif defined(__aarch64__)
+ is_optimised = std::is_same<FunctionType, NEConvolutionLayer>::value && NEScheduler::get().cpu_info().CPU >= CPUTarget::ARMV8 && data_type == DataType::F32;
+#endif /* defined(__arm__) || defined(__aarch64__) */
TensorShape tmp_weights_shape(weights_shape);
SimpleTensor<T> tmp_weights(tmp_weights_shape, data_type, 1, fixed_point_position);
diff --git a/tests/validation/fixtures/FixedPointFixture.h b/tests/validation/fixtures/FixedPointFixture.h
new file mode 100644
index 0000000..2980e2f
--- /dev/null
+++ b/tests/validation/fixtures/FixedPointFixture.h
@@ -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.
+ */
+#ifndef ARM_COMPUTE_TEST_FIXED_POINT_FIXTURE
+#define ARM_COMPUTE_TEST_FIXED_POINT_FIXTURE
+
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/IAccessor.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Fixture.h"
+#include "tests/validation/CPP/FixedPoint.h"
+#include "tests/validation/Helpers.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+template <typename TensorType, typename AccessorType, typename T>
+class FixedPointValidationFixture : public framework::Fixture
+{
+public:
+ template <typename...>
+ void setup(TensorShape shape, DataType dt, FixedPointOp op, int fractional_bits)
+ {
+ _fractional_bits = fractional_bits;
+ _target = compute_target(shape, dt, op, fractional_bits);
+ _reference = compute_reference(shape, dt, op, fractional_bits);
+ }
+
+protected:
+ template <typename U>
+ void fill(U &&tensor, int min, int max, int i)
+ {
+ std::uniform_int_distribution<> distribution(min, max);
+ library->fill(tensor, distribution, i);
+ }
+
+ TensorType compute_target(const TensorShape &shape, DataType dt, FixedPointOp op, int fixed_point_position)
+ {
+ // Create tensors
+ TensorType src = create_tensor<TensorType>(shape, dt, 1, fixed_point_position);
+ TensorType dst = create_tensor<TensorType>(shape, dt, 1, fixed_point_position);
+
+ // Allocate tensors
+ src.allocator()->allocate();
+ dst.allocator()->allocate();
+
+ ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ compute_target_impl<TensorType, AccessorType, T>(shape, dt, op, fixed_point_position, src, dst);
+
+ return dst;
+ }
+
+ SimpleTensor<T> compute_reference(const TensorShape &shape, DataType dt, FixedPointOp op, int fixed_point_position)
+ {
+ // Create reference
+ SimpleTensor<T> src{ shape, dt, 1, fixed_point_position };
+
+ // Fill reference
+ int min = 0;
+ int max = 0;
+ switch(op)
+ {
+ case FixedPointOp::EXP:
+ min = -(1 << (fixed_point_position - 1));
+ max = (1 << (fixed_point_position - 1));
+ break;
+ case FixedPointOp::INV_SQRT:
+ min = 1;
+ max = (dt == DataType::QS8) ? 0x7F : 0x7FFF;
+ break;
+ case FixedPointOp::LOG:
+ min = (1 << (fixed_point_position - 1));
+ max = (dt == DataType::QS8) ? 0x3F : 0x3FFF;
+ break;
+ case FixedPointOp::RECIPROCAL:
+ min = 15;
+ max = (dt == DataType::QS8) ? 0x7F : 0x7FFF;
+ break;
+ default:
+ ARM_COMPUTE_ERROR("Fixed point operation not supported");
+ break;
+ }
+ fill(src, min, max, 0);
+
+ return reference::fixed_point_operation<T>(src, op);
+ }
+
+ TensorType _target{};
+ SimpleTensor<T> _reference{};
+ int _fractional_bits{};
+};
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_FIXED_POINT_FIXTURE */
diff --git a/tests/validation/fixtures/GEMMLowpFixture.h b/tests/validation/fixtures/GEMMLowpFixture.h
new file mode 100644
index 0000000..c972469
--- /dev/null
+++ b/tests/validation/fixtures/GEMMLowpFixture.h
@@ -0,0 +1,126 @@
+/*
+ * 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.
+ */
+#ifndef ARM_COMPUTE_TEST_GEMMLOWP_FIXTURE
+#define ARM_COMPUTE_TEST_GEMMLOWP_FIXTURE
+
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/IAccessor.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Fixture.h"
+#include "tests/validation/CPP/GEMMLowp.h"
+#include "tests/validation/Helpers.h"
+
+#include <random>
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+template <typename TensorType, typename AccessorType, typename FunctionType>
+class GEMMLowpOffsetValidationFixture : public framework::Fixture
+{
+public:
+ template <typename...>
+ void setup(size_t m, size_t n, size_t k, int32_t a_offset, int32_t b_offset, int32_t c_offset, int32_t c_mult_int, int32_t out_shift)
+ {
+ const TensorShape shape_a(k, m);
+ const TensorShape shape_b(n, k);
+ const TensorShape shape_c(n, m);
+ _target = compute_target(shape_a, shape_b, shape_c, a_offset, b_offset, c_offset, c_mult_int, out_shift);
+ _reference = compute_reference(shape_a, shape_b, shape_c, a_offset, b_offset, c_offset, c_mult_int, out_shift);
+ }
+
+protected:
+ template <typename U>
+ void fill(U &&tensor, int i)
+ {
+ ARM_COMPUTE_ERROR_ON(tensor.data_type() != DataType::U8);
+ std::uniform_int_distribution<> distribution(0, 3);
+ library->fill(tensor, distribution, i);
+ }
+
+ TensorType compute_target(const TensorShape &shape_a, const TensorShape &shape_b, const TensorShape &shape_c,
+ int32_t a_offset, int32_t b_offset, int32_t c_offset, int32_t c_mult_int, int32_t out_shift)
+ {
+ // Create tensors
+ TensorType a = create_tensor<TensorType>(shape_a, DataType::U8, 1);
+ TensorType b = create_tensor<TensorType>(shape_b, DataType::U8, 1);
+ TensorType c = create_tensor<TensorType>(shape_c, DataType::U8, 1);
+
+ // Create and configure function
+ FunctionType gemmlowp;
+ gemmlowp.configure(&a, &b, &c, a_offset, b_offset, c_offset, c_mult_int, out_shift);
+
+ 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);
+
+ // Allocate tensors
+ a.allocator()->allocate();
+ b.allocator()->allocate();
+ c.allocator()->allocate();
+
+ 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);
+
+ // Fill tensors
+ fill(AccessorType(a), 0);
+ fill(AccessorType(b), 1);
+ fill(AccessorType(c), 2);
+
+ // Compute GEMM function
+ gemmlowp.run();
+ return c;
+ }
+
+ SimpleTensor<uint8_t> compute_reference(const TensorShape &shape_a, const TensorShape &shape_b, const TensorShape &shape_c,
+ int32_t a_offset, int32_t b_offset, int32_t c_offset, int32_t c_mult_int, int32_t out_shift)
+ {
+ // Create reference
+ SimpleTensor<uint8_t> a{ shape_a, DataType::U8, 1 };
+ SimpleTensor<uint8_t> b{ shape_b, DataType::U8, 1 };
+ SimpleTensor<uint8_t> c{ shape_c, DataType::U8, 1 };
+
+ // Fill reference
+ fill(a, 0);
+ fill(b, 1);
+ fill(c, 2);
+
+ return reference::gemmlowp<uint8_t>(a, b, c, a_offset, b_offset, c_offset, c_mult_int, out_shift);
+ }
+
+ TensorType _target{};
+ SimpleTensor<uint8_t> _reference{};
+};
+
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_GEMMLOWP_FIXTURE */
diff --git a/tests/validation/fixtures/WarpAffineFixture.h b/tests/validation/fixtures/WarpAffineFixture.h
new file mode 100644
index 0000000..fef1f6b
--- /dev/null
+++ b/tests/validation/fixtures/WarpAffineFixture.h
@@ -0,0 +1,121 @@
+/*
+ * 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.
+ */
+#ifndef ARM_COMPUTE_TEST_WARP_AFFINE_FIXTURE
+#define ARM_COMPUTE_TEST_WARP_AFFINE_FIXTURE
+
+#include <memory>
+
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/IAccessor.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Fixture.h"
+#include "tests/validation/CPP/Utils.h"
+#include "tests/validation/CPP/WarpAffine.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+template <typename TensorType, typename AccessorType, typename FunctionType, typename T>
+class WarpAffineValidationFixture : public framework::Fixture
+{
+public:
+ template <typename...>
+ void setup(TensorShape shape, DataType data_type, InterpolationPolicy policy, BorderMode border_mode)
+ {
+ // Generate a random constant value if border_mode is constant
+ std::mt19937 gen(library->seed());
+ std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+ uint8_t constant_border_value = distribution_u8(gen);
+
+ // Create the matrix
+ std::array<float, 6> matrix{ {} };
+ fill_warp_matrix<6>(matrix);
+
+ _target = compute_target(shape, data_type, matrix.data(), policy, border_mode, constant_border_value);
+ _reference = compute_reference(shape, data_type, matrix.data(), policy, border_mode, constant_border_value);
+ }
+
+protected:
+ template <typename U>
+ void fill(U &&tensor)
+ {
+ library->fill_tensor_uniform(tensor, 0);
+ }
+
+ TensorType compute_target(const TensorShape &shape, DataType data_type, const float *matrix, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
+ {
+ // Create tensors
+ TensorType src = create_tensor<TensorType>(shape, data_type);
+ TensorType dst = create_tensor<TensorType>(shape, data_type);
+
+ // Create and configure function
+ FunctionType warp_affine;
+ warp_affine.configure(&src, &dst, matrix, policy, border_mode, constant_border_value);
+
+ ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Allocate tensors
+ src.allocator()->allocate();
+ dst.allocator()->allocate();
+ ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Fill tensors
+ fill(AccessorType(src));
+
+ // Compute function
+ warp_affine.run();
+
+ return dst;
+ }
+
+ SimpleTensor<T> compute_reference(const TensorShape &shape, DataType data_type, const float *matrix, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
+ {
+ // Create reference
+ SimpleTensor<T> src{ shape, data_type };
+
+ // Create the valid mask Tensor
+ _valid_mask = SimpleTensor<T>(shape, data_type);
+
+ // Fill reference
+ fill(src);
+
+ return reference::warp_affine<T>(src, _valid_mask, matrix, policy, border_mode, constant_border_value);
+ }
+
+ TensorType _target{};
+ SimpleTensor<T> _reference{};
+ SimpleTensor<T> _valid_mask{};
+};
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_WARP_AFFINE_FIXTURE */
diff --git a/tests/validation/fixtures/WarpPerspectiveFixture.h b/tests/validation/fixtures/WarpPerspectiveFixture.h
new file mode 100644
index 0000000..c77efbd
--- /dev/null
+++ b/tests/validation/fixtures/WarpPerspectiveFixture.h
@@ -0,0 +1,132 @@
+/*
+ * 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.
+ */
+#ifndef ARM_COMPUTE_TEST_WARP_PERSPECTIVE_FIXTURE
+#define ARM_COMPUTE_TEST_WARP_PERSPECTIVE_FIXTURE
+
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/IAccessor.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Fixture.h"
+#include "tests/validation/CPP/Utils.h"
+#include "tests/validation/CPP/WarpPerspective.h"
+
+#include <random>
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+template <typename TensorType, typename AccessorType, typename FunctionType, typename T>
+class WarpPerspectiveValidationFixture : public framework::Fixture
+{
+public:
+ template <typename...>
+ void setup(TensorShape input_shape, DataType data_type, InterpolationPolicy policy, BorderMode 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(library->seed());
+ std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+ constant_border_value = distribution_u8(gen);
+ }
+
+ const TensorShape vmask_shape(input_shape);
+
+ // Create the matrix
+ std::array<float, 9> matrix = { { 0 } };
+ fill_warp_matrix<9>(matrix);
+
+ _target = compute_target(input_shape, vmask_shape, matrix.data(), policy, border_mode, constant_border_value, data_type);
+ _reference = compute_reference(input_shape, vmask_shape, matrix.data(), policy, border_mode, constant_border_value, data_type);
+ }
+
+protected:
+ template <typename U>
+ void fill(U &&tensor, int i = 0)
+ {
+ library->fill_tensor_uniform(tensor, i);
+ }
+
+ TensorType compute_target(const TensorShape &shape, const TensorShape &vmask_shape, const float *matrix, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value,
+ DataType data_type)
+ {
+ // Create tensors
+ TensorType src = create_tensor<TensorType>(shape, data_type);
+ TensorType dst = create_tensor<TensorType>(shape, data_type);
+
+ // Create and configure function
+ FunctionType warp_perspective;
+ warp_perspective.configure(&src, &dst, matrix, policy, border_mode, constant_border_value);
+
+ ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Allocate tensors
+ src.allocator()->allocate();
+ dst.allocator()->allocate();
+
+ ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+ // Fill tensors
+ fill(AccessorType(src));
+
+ // Compute function
+ warp_perspective.run();
+
+ return dst;
+ }
+
+ SimpleTensor<T> compute_reference(const TensorShape &shape, const TensorShape &vmask_shape, const float *matrix, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value,
+ DataType data_type)
+ {
+ ARM_COMPUTE_ERROR_ON(data_type != DataType::U8);
+
+ // Create reference
+ SimpleTensor<T> src{ shape, data_type };
+ SimpleTensor<T> valid_mask{ vmask_shape, data_type };
+
+ // Fill reference
+ fill(src, 0);
+ fill(valid_mask, 1);
+
+ // Compute reference
+ return reference::warp_perspective<T>(src, valid_mask, matrix, policy, border_mode, constant_border_value);
+ }
+
+ TensorType _target{};
+ SimpleTensor<T> _reference{};
+ BorderMode _border_mode{};
+};
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_WARP_PERSPECTIVE_FIXTURE */