arm_compute v17.10
Change-Id: If1489af40eccd0219ede8946577afbf04db31b29
diff --git a/tests/AssetsLibrary.cpp b/tests/AssetsLibrary.cpp
index 5841d5b..744c987 100644
--- a/tests/AssetsLibrary.cpp
+++ b/tests/AssetsLibrary.cpp
@@ -191,6 +191,11 @@
{
}
+std::string AssetsLibrary::path() const
+{
+ return _library_path;
+}
+
std::random_device::result_type AssetsLibrary::seed() const
{
return _seed;
diff --git a/tests/AssetsLibrary.h b/tests/AssetsLibrary.h
index f1504f3..7f7a26a 100644
--- a/tests/AssetsLibrary.h
+++ b/tests/AssetsLibrary.h
@@ -58,14 +58,17 @@
class AssetsLibrary final
{
public:
- /** Initialises the library with a @p path to the image directory.
+ /** Initialises the library with a @p path to the assets directory.
* Furthermore, sets the seed for the random generator to @p seed.
*
- * @param[in] path Path to load images from.
+ * @param[in] path Path to load assets from.
* @param[in] seed Seed used to initialise the random number generator.
*/
AssetsLibrary(std::string path, std::random_device::result_type seed);
+ /** Path to assets directory used to initialise library. */
+ std::string path() const;
+
/** Seed that is used to fill tensors with random values. */
std::random_device::result_type seed() const;
diff --git a/tests/SConscript b/tests/SConscript
index 6ed9e2c..a1cb9f8 100644
--- a/tests/SConscript
+++ b/tests/SConscript
@@ -51,11 +51,12 @@
if env['os'] in ['android', 'bare_metal'] or env['standalone']:
Import("arm_compute_a")
- test_env.Append(LIBS = [arm_compute_a])
+ Import("arm_compute_core_a")
+ test_env.Append(LIBS = [arm_compute_a, arm_compute_core_a])
arm_compute_lib = arm_compute_a
else:
Import("arm_compute_so")
- test_env.Append(LIBS = ["arm_compute"])
+ test_env.Append(LIBS = ["arm_compute", "arm_compute_core"])
arm_compute_lib = arm_compute_so
test_env.Append(CPPPATH = ["#3rdparty/include"])
@@ -72,9 +73,16 @@
files_benchmark = Glob('benchmark/*.cpp')
files_validation = Glob('validation/*.cpp')
+# Add unit tests
+files_validation += Glob('validation/UNIT/*.cpp')
+
# Always compile reference for validation
files_validation += Glob('validation/CPP/*.cpp')
+# Add unit tests
+files_validation += Glob('validation/UNIT/*/*.cpp')
+files_validation += Glob('validation/UNIT/*.cpp')
+
if env['opencl']:
Import('opencl')
diff --git a/tests/Types.h b/tests/Types.h
index 516f61c..8f47902 100644
--- a/tests/Types.h
+++ b/tests/Types.h
@@ -33,6 +33,9 @@
/** Fixed point operation */
enum class FixedPointOp
{
+ ADD, /**< Addition */
+ SUB, /**< Subtraction */
+ MUL, /**< Multiplication */
EXP, /**< Exponential */
LOG, /**< Logarithm */
INV_SQRT, /**< Inverse square root */
diff --git a/tests/datasets/system_tests/yolo/v2/YOLOV2ConvolutionLayerDataset.h b/tests/datasets/system_tests/yolo/v2/YOLOV2ConvolutionLayerDataset.h
index b4cb357..5e07ff8 100644
--- a/tests/datasets/system_tests/yolo/v2/YOLOV2ConvolutionLayerDataset.h
+++ b/tests/datasets/system_tests/yolo/v2/YOLOV2ConvolutionLayerDataset.h
@@ -67,7 +67,7 @@
// conv21
add_config(TensorShape(13U, 13U, 3072U), TensorShape(3U, 3U, 3072U, 1024U), TensorShape(1024U), TensorShape(13U, 13U, 1024U), PadStrideInfo(1, 1, 1, 1));
// conv22
- add_config(TensorShape(13U, 13U, 1024U), TensorShape(1U, 1U, 1024U, 425U), TensorShape(425U), TensorShape(15U, 15U, 425U), PadStrideInfo(1, 1, 1, 1));
+ add_config(TensorShape(13U, 13U, 1024U), TensorShape(1U, 1U, 1024U, 425U), TensorShape(425U), TensorShape(13U, 13U, 425U), PadStrideInfo(1, 1, 0, 0));
}
};
} // namespace datasets
diff --git a/tests/framework/datasets/CartesianProductDataset.h b/tests/framework/datasets/CartesianProductDataset.h
index f6e45dd..438a782 100644
--- a/tests/framework/datasets/CartesianProductDataset.h
+++ b/tests/framework/datasets/CartesianProductDataset.h
@@ -158,6 +158,14 @@
{
return CartesianProductDataset<T, U>(std::forward<T>(dataset1), std::forward<U>(dataset2));
}
+
+template <typename T, typename U>
+CartesianProductDataset<T, U>
+operator*(T &&dataset1, U &&dataset2)
+{
+ return CartesianProductDataset<T, U>(std::forward<T>(dataset1), std::forward<U>(dataset2));
+}
+
} // namespace dataset
} // namespace framework
} // namespace test
diff --git a/tests/framework/instruments/MaliCounter.cpp b/tests/framework/instruments/MaliCounter.cpp
index 887846e..6cc3ac5 100644
--- a/tests/framework/instruments/MaliCounter.cpp
+++ b/tests/framework/instruments/MaliCounter.cpp
@@ -23,6 +23,8 @@
*/
#include "MaliCounter.h"
+#include "arm_compute/core/Error.h"
+
namespace arm_compute
{
namespace test
@@ -46,7 +48,7 @@
if(fd < 0)
{
- throw std::runtime_error("Failed to get HW info.");
+ ARM_COMPUTE_ERROR("Failed to get HW info.");
}
{
@@ -57,7 +59,7 @@
if(mali_userspace::mali_ioctl(fd, version_check_args) != 0)
{
- throw std::runtime_error("Failed to check version.");
+ ARM_COMPUTE_ERROR("Failed to check version.");
close(fd);
}
}
@@ -70,7 +72,7 @@
if(mali_userspace::mali_ioctl(fd, flags) != 0)
{
- throw std::runtime_error("Failed settings flags ioctl.");
+ ARM_COMPUTE_ERROR("Failed settings flags ioctl.");
close(fd);
}
}
@@ -81,7 +83,7 @@
if(mali_ioctl(fd, props) != 0)
{
- throw std::runtime_error("Failed settings flags ioctl.");
+ ARM_COMPUTE_ERROR("Failed settings flags ioctl.");
close(fd);
}
@@ -141,7 +143,7 @@
if(_fd < 0)
{
- throw std::runtime_error("Failed to open /dev/mali0.");
+ ARM_COMPUTE_ERROR("Failed to open /dev/mali0.");
}
{
@@ -150,11 +152,11 @@
if(mali_userspace::mali_ioctl(_fd, check) != 0)
{
- throw std::runtime_error("Failed to get ABI version.");
+ ARM_COMPUTE_ERROR("Failed to get ABI version.");
}
else if(check.major < 10)
{
- throw std::runtime_error("Unsupported ABI version 10.");
+ ARM_COMPUTE_ERROR("Unsupported ABI version 10.");
}
}
@@ -166,7 +168,7 @@
if(mali_userspace::mali_ioctl(_fd, flags) != 0)
{
- throw std::runtime_error("Failed settings flags ioctl.");
+ ARM_COMPUTE_ERROR("Failed settings flags ioctl.");
}
}
@@ -183,7 +185,7 @@
if(mali_userspace::mali_ioctl(_fd, setup) != 0)
{
- throw std::runtime_error("Failed setting hwcnt reader ioctl.");
+ ARM_COMPUTE_ERROR("Failed setting hwcnt reader ioctl.");
}
_hwc_fd = setup.fd;
@@ -194,34 +196,34 @@
if(ioctl(_hwc_fd, mali_userspace::KBASE_HWCNT_READER_GET_API_VERSION, &api_version) != 0) // NOLINT
{
- throw std::runtime_error("Could not determine hwcnt reader API.");
+ ARM_COMPUTE_ERROR("Could not determine hwcnt reader API.");
}
else if(api_version != mali_userspace::HWCNT_READER_API)
{
- throw std::runtime_error("Invalid API version.");
+ ARM_COMPUTE_ERROR("Invalid API version.");
}
}
if(ioctl(_hwc_fd, static_cast<int>(mali_userspace::KBASE_HWCNT_READER_GET_BUFFER_SIZE), &_buffer_size) != 0) // NOLINT
{
- throw std::runtime_error("Failed to get buffer size.");
+ ARM_COMPUTE_ERROR("Failed to get buffer size.");
}
if(ioctl(_hwc_fd, static_cast<int>(mali_userspace::KBASE_HWCNT_READER_GET_HWVER), &_hw_ver) != 0) // NOLINT
{
- throw std::runtime_error("Could not determine HW version.");
+ ARM_COMPUTE_ERROR("Could not determine HW version.");
}
if(_hw_ver < 5)
{
- throw std::runtime_error("Unsupported HW version.");
+ ARM_COMPUTE_ERROR("Unsupported HW version.");
}
_sample_data = static_cast<uint8_t *>(mmap(nullptr, _buffer_count * _buffer_size, PROT_READ, MAP_PRIVATE, _hwc_fd, 0));
if(_sample_data == MAP_FAILED) // NOLINT
{
- throw std::runtime_error("Failed to map sample data.");
+ ARM_COMPUTE_ERROR("Failed to map sample data.");
}
auto product = std::find_if(std::begin(mali_userspace::products), std::end(mali_userspace::products), [&](const mali_userspace::CounterMapping & cm)
@@ -235,7 +237,7 @@
}
else
{
- throw std::runtime_error("Could not identify GPU.");
+ ARM_COMPUTE_ERROR("Could not identify GPU.");
}
_raw_counter_buffer.resize(_buffer_size / sizeof(uint32_t));
@@ -279,7 +281,7 @@
{
if(ioctl(_hwc_fd, mali_userspace::KBASE_HWCNT_READER_DUMP, 0) != 0)
{
- throw std::runtime_error("Could not sample hardware counters.");
+ ARM_COMPUTE_ERROR("Could not sample hardware counters.");
}
}
@@ -293,7 +295,7 @@
if(count < 0)
{
- throw std::runtime_error("poll() failed.");
+ ARM_COMPUTE_ERROR("poll() failed.");
}
if((poll_fd.revents & POLLIN) != 0)
@@ -302,7 +304,7 @@
if(ioctl(_hwc_fd, static_cast<int>(mali_userspace::KBASE_HWCNT_READER_GET_BUFFER), &meta) != 0) // NOLINT
{
- throw std::runtime_error("Failed READER_GET_BUFFER.");
+ ARM_COMPUTE_ERROR("Failed READER_GET_BUFFER.");
}
memcpy(_raw_counter_buffer.data(), _sample_data + _buffer_size * meta.buffer_idx, _buffer_size);
@@ -310,12 +312,12 @@
if(ioctl(_hwc_fd, mali_userspace::KBASE_HWCNT_READER_PUT_BUFFER, &meta) != 0) // NOLINT
{
- throw std::runtime_error("Failed READER_PUT_BUFFER.");
+ ARM_COMPUTE_ERROR("Failed READER_PUT_BUFFER.");
}
}
else if((poll_fd.revents & POLLHUP) != 0)
{
- throw std::runtime_error("HWC hung up.");
+ ARM_COMPUTE_ERROR("HWC hung up.");
}
}
@@ -337,7 +339,7 @@
default:
if(core < 0)
{
- std::runtime_error("Invalid core number.");
+ ARM_COMPUTE_ERROR("Invalid core number.");
}
return _raw_counter_buffer.data() + mali_userspace::MALI_NAME_BLOCK_SIZE * (3 + _core_index_remap[core]);
diff --git a/tests/framework/instruments/PMU.cpp b/tests/framework/instruments/PMU.cpp
index 0594e96..d0cacbb 100644
--- a/tests/framework/instruments/PMU.cpp
+++ b/tests/framework/instruments/PMU.cpp
@@ -74,9 +74,10 @@
ARM_COMPUTE_ERROR_ON_MSG(_fd < 0, "perf_event_open failed");
const int result = ioctl(_fd, PERF_EVENT_IOC_ENABLE, 0);
-
- ARM_COMPUTE_ERROR_ON_MSG(result == -1, "Failed to enable PMU counter: %d", errno);
- ARM_COMPUTE_UNUSED(result);
+ if(result == -1)
+ {
+ ARM_COMPUTE_ERROR("Failed to enable PMU counter: %d", errno);
+ }
}
void PMU::close()
@@ -91,8 +92,10 @@
void PMU::reset()
{
const int result = ioctl(_fd, PERF_EVENT_IOC_RESET, 0);
- ARM_COMPUTE_ERROR_ON_MSG(result == -1, "Failed to reset PMU counter: %d", errno);
- ARM_COMPUTE_UNUSED(result);
+ if(result == -1)
+ {
+ ARM_COMPUTE_ERROR("Failed to reset PMU counter: %d", errno);
+ }
}
} // namespace framework
} // namespace test
diff --git a/tests/framework/instruments/PMU.h b/tests/framework/instruments/PMU.h
index d51b2f8..e0a1870 100644
--- a/tests/framework/instruments/PMU.h
+++ b/tests/framework/instruments/PMU.h
@@ -86,8 +86,10 @@
T value{};
const ssize_t result = read(_fd, &value, sizeof(T));
- ARM_COMPUTE_ERROR_ON_MSG(result == -1, "Can't get PMU counter value: %d", errno);
- ARM_COMPUTE_UNUSED(result);
+ if(result == -1)
+ {
+ ARM_COMPUTE_ERROR("Can't get PMU counter value: %d", errno);
+ }
return value;
}
diff --git a/tests/main.cpp b/tests/main.cpp
index ee12a38..5bdd665 100644
--- a/tests/main.cpp
+++ b/tests/main.cpp
@@ -27,6 +27,7 @@
#include "tests/framework/Exceptions.h"
#include "tests/framework/Framework.h"
#include "tests/framework/Macros.h"
+#include "tests/framework/Profiler.h"
#include "tests/framework/command_line/CommandLineOptions.h"
#include "tests/framework/command_line/CommandLineParser.h"
#include "tests/framework/instruments/Instruments.h"
@@ -103,9 +104,9 @@
auto help = parser.add_option<framework::ToggleOption>("help");
help->set_help("Show this help message");
- auto dataset_mode = parser.add_option<framework::EnumOption<framework::DatasetMode>>("mode", allowed_modes, framework::DatasetMode::ALL);
+ auto dataset_mode = parser.add_option<framework::EnumOption<framework::DatasetMode>>("mode", allowed_modes, framework::DatasetMode::PRECOMMIT);
dataset_mode->set_help("For managed datasets select which group to use");
- auto instruments = parser.add_option<framework::EnumListOption<framework::InstrumentType>>("instruments", allowed_instruments, std::initializer_list<framework::InstrumentType> { framework::InstrumentType::ALL });
+ auto instruments = parser.add_option<framework::EnumListOption<framework::InstrumentType>>("instruments", allowed_instruments, std::initializer_list<framework::InstrumentType> { framework::InstrumentType::WALL_CLOCK_TIMER });
instruments->set_help("Set the profiling instruments to use");
auto iterations = parser.add_option<framework::SimpleOption<int>>("iterations", 1);
iterations->set_help("Number of iterations per test case");
@@ -131,6 +132,8 @@
color_output->set_help("Produce colored output on the console");
auto list_tests = parser.add_option<framework::ToggleOption>("list-tests", false);
list_tests->set_help("List all test names");
+ auto test_instruments = parser.add_option<framework::ToggleOption>("test-instruments", false);
+ test_instruments->set_help("Test if the instruments work on the platform");
auto error_on_missing_assets = parser.add_option<framework::ToggleOption>("error-on-missing-assets", false);
error_on_missing_assets->set_help("Mark a test as failed instead of skipping it when assets are missing");
auto assets = parser.add_positional_option<framework::SimpleOption<std::string>>("assets");
@@ -211,6 +214,18 @@
return 0;
}
+ if(test_instruments->value())
+ {
+ framework::Profiler profiler = framework.get_profiler();
+ profiler.start();
+ profiler.stop();
+ if(printer != nullptr)
+ {
+ printer->print_measurements(profiler.measurements());
+ }
+ return 0;
+ }
+
library = support::cpp14::make_unique<AssetsLibrary>(assets->value(), seed->value());
if(!parser.validate())
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 */