COMPMID-415: Rename and move tests

The boost validation is now "standalone" in validation_old and builds as
arm_compute_validation_old. The new validation builds now as
arm_compute_validation.

Change-Id: Ib93ba848a25680ac60afb92b461d574a0757150d
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/86187
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>

31 files changed, 6191 insertions, 0 deletions

diff --git a/tests/validation_old/NEON/AbsoluteDifference.cpp b/tests/validation_old/NEON/AbsoluteDifference.cpp
new file mode 100644
index 0000000000..aa866fff49
--- /dev/null
+++ b/tests/validation_old/NEON/AbsoluteDifference.cpp
@@ -0,0 +1,200 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEAbsoluteDifference.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute Neon absolute difference function.
+ *
+ * @param[in] shape  Shape of the input and output tensors.
+ * @param[in] dt_in0 Data type of first input tensor.
+ * @param[in] dt_in1 Data type of second input tensor.
+ * @param[in] dt_out Data type of the output tensor.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_absolute_difference(const TensorShape &shape, DataType dt_in0, DataType dt_in1, DataType dt_out)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, dt_in0);
+    Tensor src2 = create_tensor<Tensor>(shape, dt_in1);
+    Tensor dst  = create_tensor<Tensor>(shape, dt_out);
+
+    // Create and configure function
+    NEAbsoluteDifference abs_d;
+    abs_d.configure(&src1, &src2, &dst);
+
+    // Allocate tensors
+    src1.allocator()->allocate();
+    src2.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src1.info()->is_resizable());
+    BOOST_TEST(!src2.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src1), 0);
+    library->fill_tensor_uniform(Accessor(src2), 1);
+
+    // Compute function
+    abs_d.run();
+
+    return dst;
+}
+
+void validate_configuration(const Tensor &src1, const Tensor &src2, Tensor &dst, TensorShape shape)
+{
+    BOOST_TEST(src1.info()->is_resizable());
+    BOOST_TEST(src2.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create and configure function
+    NEAbsoluteDifference abs_d;
+    abs_d.configure(&src1, &src2, &dst);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+    validate(src1.info()->valid_region(), valid_region);
+    validate(src2.info()->valid_region(), valid_region);
+    validate(dst.info()->valid_region(), valid_region);
+
+    // Validate padding
+    const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+    validate(src1.info()->padding(), padding);
+    validate(src2.info()->padding(), padding);
+    validate(dst.info()->padding(), padding);
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(AbsoluteDifference)
+
+BOOST_AUTO_TEST_SUITE(U8)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()),
+                     shape)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor src2 = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::U8);
+
+    validate_configuration(src1, src2, dst, shape);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(),
+                     shape)
+{
+    // Compute function
+    Tensor dst = compute_absolute_difference(shape, DataType::U8, DataType::U8, DataType::U8);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_absolute_difference(shape, DataType::U8, DataType::U8, DataType::U8);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes(),
+                     shape)
+{
+    // Compute function
+    Tensor dst = compute_absolute_difference(shape, DataType::U8, DataType::U8, DataType::U8);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_absolute_difference(shape, DataType::U8, DataType::U8, DataType::U8);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(S16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }),
+                     shape, dt)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, dt);
+    Tensor src2 = create_tensor<Tensor>(shape, DataType::S16);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::S16);
+
+    validate_configuration(src1, src2, dst, shape);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }),
+                     shape, dt)
+{
+    // Compute function
+    Tensor dst = compute_absolute_difference(shape, dt, DataType::S16, DataType::S16);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_absolute_difference(shape, dt, DataType::S16, DataType::S16);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }),
+                     shape, dt)
+{
+    // Compute function
+    Tensor dst = compute_absolute_difference(shape, dt, DataType::S16, DataType::S16);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_absolute_difference(shape, dt, DataType::S16, DataType::S16);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Accumulate.cpp b/tests/validation_old/NEON/Accumulate.cpp
new file mode 100644
index 0000000000..eb680a383d
--- /dev/null
+++ b/tests/validation_old/NEON/Accumulate.cpp
@@ -0,0 +1,145 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEAccumulate.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute Neon accumulate function.
+ *
+ * @param[in] shape Shape of the input and output tensors.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_accumulate(const TensorShape &shape)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::S16);
+
+    // Create and configure function
+    NEAccumulate acc;
+    acc.configure(&src, &dst);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src), 0);
+    library->fill_tensor_uniform(Accessor(dst), 1);
+
+    // Compute function
+    acc.run();
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(Accumulate)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()),
+                     shape)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::S16);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create and configure function
+    NEAccumulate acc;
+    acc.configure(&src, &dst);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+    validate(src.info()->valid_region(), valid_region);
+    validate(dst.info()->valid_region(), valid_region);
+
+    // Validate padding
+    const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+    validate(src.info()->padding(), padding);
+    validate(dst.info()->padding(), padding);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(),
+                     shape)
+{
+    // Compute function
+    Tensor dst = compute_accumulate(shape);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_accumulate(shape);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes(),
+                     shape)
+{
+    // Compute function
+    Tensor dst = compute_accumulate(shape);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_accumulate(shape);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/AccumulateSquared.cpp b/tests/validation_old/NEON/AccumulateSquared.cpp
new file mode 100644
index 0000000000..29b5edf41b
--- /dev/null
+++ b/tests/validation_old/NEON/AccumulateSquared.cpp
@@ -0,0 +1,146 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEAccumulate.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute Neon accumulate squared function.
+ *
+ * @param[in] shape Shape of the input and output tensors.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_accumulate_squared(const TensorShape &shape, uint32_t shift)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::S16);
+
+    // Create and configure function
+    NEAccumulateSquared acc;
+    acc.configure(&src, shift, &dst);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    // dst tensor filled with non-negative values
+    library->fill_tensor_uniform(Accessor(src), 0);
+    library->fill_tensor_uniform(Accessor(dst), 1, static_cast<int16_t>(0), std::numeric_limits<int16_t>::max());
+
+    // Compute function
+    acc.run();
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(AccumulateSquared)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::xrange(0U, 16U),
+                     shape, shift)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::S16);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create and configure function
+    NEAccumulateSquared acc;
+    acc.configure(&src, shift, &dst);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+    validate(src.info()->valid_region(), valid_region);
+    validate(dst.info()->valid_region(), valid_region);
+
+    // Validate padding
+    const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+    validate(src.info()->padding(), padding);
+    validate(dst.info()->padding(), padding);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::xrange(0U, 16U),
+                     shape, shift)
+{
+    // Compute function
+    Tensor dst = compute_accumulate_squared(shape, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_accumulate_squared(shape, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ 0U, 1U, 15U }),
+                     shape, shift)
+{
+    // Compute function
+    Tensor dst = compute_accumulate_squared(shape, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_accumulate_squared(shape, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/AccumulateWeighted.cpp b/tests/validation_old/NEON/AccumulateWeighted.cpp
new file mode 100644
index 0000000000..c59c1edbc8
--- /dev/null
+++ b/tests/validation_old/NEON/AccumulateWeighted.cpp
@@ -0,0 +1,145 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEAccumulate.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute Neon accumulate weighted function.
+ *
+ * @param[in] shape Shape of the input and output tensors.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_accumulate_weighted(const TensorShape &shape, float alpha)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U8);
+
+    // Create and configure function
+    NEAccumulateWeighted acc;
+    acc.configure(&src, alpha, &dst);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src), 0);
+    library->fill_tensor_uniform(Accessor(dst), 1);
+
+    // Compute function
+    acc.run();
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(AccumulateWeighted)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ 0.f, 0.5f, 1.f }),
+                     shape, alpha)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U8);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create and configure function
+    NEAccumulateWeighted acc;
+    acc.configure(&src, alpha, &dst);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+    validate(src.info()->valid_region(), valid_region);
+    validate(dst.info()->valid_region(), valid_region);
+
+    // Validate padding
+    const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+    validate(src.info()->padding(), padding);
+    validate(dst.info()->padding(), padding);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ 0.f, 0.5f, 1.f }),
+                     shape, alpha)
+{
+    // Compute function
+    Tensor dst = compute_accumulate_weighted(shape, alpha);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_accumulate_weighted(shape, alpha);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ 0.f, 0.5f, 1.f }),
+                     shape, alpha)
+{
+    // Compute function
+    Tensor dst = compute_accumulate_weighted(shape, alpha);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_accumulate_weighted(shape, alpha);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/ArithmeticAddition.cpp b/tests/validation_old/NEON/ArithmeticAddition.cpp
new file mode 100644
index 0000000000..490f124422
--- /dev/null
+++ b/tests/validation_old/NEON/ArithmeticAddition.cpp
@@ -0,0 +1,304 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEArithmeticAddition.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute Neon arithmetic addition function.
+ *
+ * @param[in] shape                Shape of the input and output tensors.
+ * @param[in] dt_in0               Data type of first input tensor.
+ * @param[in] dt_in1               Data type of second input tensor.
+ * @param[in] dt_out               Data type of the output tensor.
+ * @param[in] policy               Overflow policy of the operation.
+ * @param[in] fixed_point_position (Optional) Fixed point position that expresses the number of bits for the fractional part of the number when the tensor's data type is QS8 or QS16 (default = 0).
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_arithmetic_addition(const TensorShape &shape, DataType dt_in0, DataType dt_in1, DataType dt_out, ConvertPolicy policy, int fixed_point_position = 0)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, dt_in0, 1, fixed_point_position);
+    Tensor src2 = create_tensor<Tensor>(shape, dt_in1, 1, fixed_point_position);
+    Tensor dst  = create_tensor<Tensor>(shape, dt_out, 1, fixed_point_position);
+
+    // Create and configure function
+    NEArithmeticAddition add;
+    add.configure(&src1, &src2, &dst, policy);
+
+    // Allocate tensors
+    src1.allocator()->allocate();
+    src2.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src1.info()->is_resizable());
+    BOOST_TEST(!src2.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src1), 0);
+    library->fill_tensor_uniform(Accessor(src2), 1);
+
+    // Compute function
+    add.run();
+
+    return dst;
+}
+
+void validate_configuration(const Tensor &src1, const Tensor &src2, Tensor &dst, TensorShape shape, ConvertPolicy policy)
+{
+    BOOST_TEST(src1.info()->is_resizable());
+    BOOST_TEST(src2.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create and configure function
+    NEArithmeticAddition add;
+    add.configure(&src1, &src2, &dst, policy);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+    validate(src1.info()->valid_region(), valid_region);
+    validate(src2.info()->valid_region(), valid_region);
+    validate(dst.info()->valid_region(), valid_region);
+
+    // Validate padding
+    const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+    validate(src1.info()->padding(), padding);
+    validate(src2.info()->padding(), padding);
+    validate(dst.info()->padding(), padding);
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(ArithmeticAddition)
+
+BOOST_AUTO_TEST_SUITE(U8)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
+                     shape, policy)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor src2 = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::U8);
+
+    validate_configuration(src1, src2, dst, shape, policy);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
+                     shape, policy)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_addition(shape, DataType::U8, DataType::U8, DataType::U8, policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::U8, DataType::U8, DataType::U8, policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(S16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
+                     shape, dt, policy)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, dt);
+    Tensor src2 = create_tensor<Tensor>(shape, DataType::S16);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::S16);
+
+    validate_configuration(src1, src2, dst, shape, policy);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
+                     shape, dt, policy)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_addition(shape, dt, DataType::S16, DataType::S16, policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, dt, DataType::S16, DataType::S16, policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
+                     shape, dt, policy)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_addition(shape, dt, DataType::S16, DataType::S16, policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, dt, DataType::S16, DataType::S16, policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(Quantized)
+BOOST_AUTO_TEST_SUITE(QS8)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 7),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_addition(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 7),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_addition(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(QS16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 15),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_addition(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 15),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_addition(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+
+#ifdef ARM_COMPUTE_ENABLE_FP16
+BOOST_AUTO_TEST_SUITE(F16)
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), shape)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_addition(shape, DataType::F16, DataType::F16, DataType::F16, ConvertPolicy::WRAP);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::F16, DataType::F16, DataType::F16, ConvertPolicy::WRAP);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* ARM_COMPUTE_ENABLE_FP16 */
+
+BOOST_AUTO_TEST_SUITE(F32)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
+                     shape, policy)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, DataType::F32);
+    Tensor src2 = create_tensor<Tensor>(shape, DataType::F32);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::F32);
+
+    validate_configuration(src1, src2, dst, shape, policy);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), shape)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_addition(shape, DataType::F32, DataType::F32, DataType::F32, ConvertPolicy::WRAP);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::F32, DataType::F32, DataType::F32, ConvertPolicy::WRAP);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
+                     shape, policy)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_addition(shape, DataType::F32, DataType::F32, DataType::F32, policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::F32, DataType::F32, DataType::F32, policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/ArithmeticSubtraction.cpp b/tests/validation_old/NEON/ArithmeticSubtraction.cpp
new file mode 100644
index 0000000000..86aa124f00
--- /dev/null
+++ b/tests/validation_old/NEON/ArithmeticSubtraction.cpp
@@ -0,0 +1,306 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEArithmeticSubtraction.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute Neon arithmetic subtraction function.
+ *
+ * @param[in] shape                Shape of the input and output tensors.
+ * @param[in] dt_in0               Data type of first input tensor.
+ * @param[in] dt_in1               Data type of second input tensor.
+ * @param[in] dt_out               Data type of the output tensor.
+ * @param[in] policy               Overflow policy of the operation.
+ * @param[in] fixed_point_position (Optional) Fixed point position that expresses the number of bits for the fractional part of the number when the tensor's data type is QS8 or QS16 (default = 0).
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_arithmetic_subtraction(const TensorShape &shape, DataType dt_in0, DataType dt_in1, DataType dt_out, ConvertPolicy policy, int fixed_point_position = 0)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, dt_in0, 1, fixed_point_position);
+    Tensor src2 = create_tensor<Tensor>(shape, dt_in1, 1, fixed_point_position);
+    Tensor dst  = create_tensor<Tensor>(shape, dt_out, 1, fixed_point_position);
+
+    // Create and configure function
+    NEArithmeticSubtraction sub;
+    sub.configure(&src1, &src2, &dst, policy);
+
+    // Allocate tensors
+    src1.allocator()->allocate();
+    src2.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src1.info()->is_resizable());
+    BOOST_TEST(!src2.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src1), 0);
+    library->fill_tensor_uniform(Accessor(src2), 1);
+
+    // Compute function
+    sub.run();
+
+    return dst;
+}
+
+void validate_configuration(const Tensor &src1, const Tensor &src2, Tensor &dst, TensorShape shape, ConvertPolicy policy)
+{
+    BOOST_TEST(src1.info()->is_resizable());
+    BOOST_TEST(src2.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create and configure function
+    NEArithmeticSubtraction sub;
+    sub.configure(&src1, &src2, &dst, policy);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+    validate(src1.info()->valid_region(), valid_region);
+    validate(src2.info()->valid_region(), valid_region);
+    validate(dst.info()->valid_region(), valid_region);
+
+    // Validate padding
+    const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+    validate(src1.info()->padding(), padding);
+    validate(src2.info()->padding(), padding);
+    validate(dst.info()->padding(), padding);
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(ArithmeticSubtraction)
+
+BOOST_AUTO_TEST_SUITE(U8)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
+                     shape, policy)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor src2 = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::U8);
+
+    validate_configuration(src1, src2, dst, shape, policy);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
+                     shape, policy)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_subtraction(shape, DataType::U8, DataType::U8, DataType::U8, policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::U8, DataType::U8, DataType::U8, policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(S16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
+                     shape, dt, policy)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, dt);
+    Tensor src2 = create_tensor<Tensor>(shape, DataType::S16);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::S16);
+
+    validate_configuration(src1, src2, dst, shape, policy);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
+                     shape, dt, policy)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_subtraction(shape, dt, DataType::S16, DataType::S16, policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, dt, DataType::S16, DataType::S16, policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
+                     shape, dt, policy)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_subtraction(shape, dt, DataType::S16, DataType::S16, policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, dt, DataType::S16, DataType::S16, policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(Quantized)
+BOOST_AUTO_TEST_SUITE(QS8)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 7),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_subtraction(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 7),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_subtraction(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(QS16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 15),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_subtraction(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 15),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_subtraction(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+
+#ifdef ARM_COMPUTE_ENABLE_FP16
+BOOST_AUTO_TEST_SUITE(Float16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), shape)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_subtraction(shape, DataType::F16, DataType::F16, DataType::F16, ConvertPolicy::WRAP);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::F16, DataType::F16, DataType::F16, ConvertPolicy::WRAP);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* ARM_COMPUTE_ENABLE_FP16 */
+
+BOOST_AUTO_TEST_SUITE(Float)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
+                     shape, policy)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, DataType::F32);
+    Tensor src2 = create_tensor<Tensor>(shape, DataType::F32);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::F32);
+
+    validate_configuration(src1, src2, dst, shape, policy);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), shape)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_subtraction(shape, DataType::F32, DataType::F32, DataType::F32, ConvertPolicy::WRAP);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::F32, DataType::F32, DataType::F32, ConvertPolicy::WRAP);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }),
+                     shape, policy)
+{
+    // Compute function
+    Tensor dst = compute_arithmetic_subtraction(shape, DataType::F32, DataType::F32, DataType::F32, policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::F32, DataType::F32, DataType::F32, policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/BatchNormalizationLayer.cpp b/tests/validation_old/NEON/BatchNormalizationLayer.cpp
new file mode 100644
index 0000000000..d98f99a63c
--- /dev/null
+++ b/tests/validation_old/NEON/BatchNormalizationLayer.cpp
@@ -0,0 +1,258 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "TypePrinter.h"
+#include "tests/Globals.h"
+#include "tests/NEON/Helper.h"
+#include "tests/Utils.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Helpers.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+#include "tests/validation_old/dataset/BatchNormalizationLayerDataset.h"
+
+#include "arm_compute/runtime/NEON/functions/NEBatchNormalizationLayer.h"
+
+#include <random>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+const float tolerance_qs8  = 6;      /**< Tolerance value for comparing reference's output against quantized implementation's output */
+const float tolerance_qs16 = 6;      /**< Tolerance value for comparing reference's output against quantized implementation's output */
+const float tolerance_f32  = 1e-05f; /**< Tolerance value for comparing reference's output against floating point implementation's output */
+#ifdef ARM_COMPUTE_ENABLE_FP16
+const float tolerance_f16 = 0.01f; /**< Tolerance value for comparing reference's output against half precision floating point implementation's output */
+#endif                             /* ARM_COMPUTE_ENABLE_FP16 */
+
+/** Compute Neon batch normalization function.
+ *
+ * @param[in] shape     Shape of the input and output tensors.
+ * @param[in] dt        Data type of input and output tensors.
+ * @param[in] norm_info Normalization Layer information.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_reference_batch_normalization_layer(const TensorShape &shape0, const TensorShape &shape1, DataType dt, float epsilon, int fixed_point_position = 0)
+{
+    // Create tensors
+    Tensor src   = create_tensor<Tensor>(shape0, dt, 1, fixed_point_position);
+    Tensor dst   = create_tensor<Tensor>(shape0, dt, 1, fixed_point_position);
+    Tensor mean  = create_tensor<Tensor>(shape1, dt, 1, fixed_point_position);
+    Tensor var   = create_tensor<Tensor>(shape1, dt, 1, fixed_point_position);
+    Tensor beta  = create_tensor<Tensor>(shape1, dt, 1, fixed_point_position);
+    Tensor gamma = create_tensor<Tensor>(shape1, dt, 1, fixed_point_position);
+
+    // Create and configure function
+    NEBatchNormalizationLayer norm;
+    norm.configure(&src, &dst, &mean, &var, &beta, &gamma, epsilon);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+    mean.allocator()->allocate();
+    var.allocator()->allocate();
+    beta.allocator()->allocate();
+    gamma.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+    BOOST_TEST(!mean.info()->is_resizable());
+    BOOST_TEST(!var.info()->is_resizable());
+    BOOST_TEST(!beta.info()->is_resizable());
+    BOOST_TEST(!gamma.info()->is_resizable());
+
+    // Fill tensors
+    switch(dt)
+    {
+        case DataType::QS8:
+        {
+            const std::pair<int8_t, int8_t> bounds = get_batchnormalization_layer_test_bounds<int8_t>(fixed_point_position);
+            std::uniform_int_distribution<> distribution(bounds.first, bounds.second);
+            std::uniform_int_distribution<> distribution_var(0, bounds.second);
+            test::fill_tensors(distribution, { 0, 1, 3, 4 }, &src, &mean, &beta, &gamma);
+            test::fill_tensors(distribution_var, { 0 }, &var);
+            break;
+        }
+        case DataType::QS16:
+        {
+            const std::pair<int16_t, int16_t> bounds = get_batchnormalization_layer_test_bounds<int16_t>(fixed_point_position);
+            std::uniform_int_distribution<> distribution(bounds.first, bounds.second);
+            std::uniform_int_distribution<> distribution_var(0, bounds.second);
+            test::fill_tensors(distribution, { 0, 1, 3, 4 }, &src, &mean, &beta, &gamma);
+            test::fill_tensors(distribution_var, { 0 }, &var);
+            break;
+        }
+#ifdef ARM_COMPUTE_ENABLE_FP16
+        case DataType::F16:
+        {
+            const std::pair<half_float::half, half_float::half> bounds = get_batchnormalization_layer_test_bounds<half_float::half>();
+            std::uniform_real_distribution<> distribution(bounds.first, bounds.second);
+            std::uniform_real_distribution<> distribution_var(0, bounds.second);
+            test::fill_tensors(distribution, { 0, 1, 3, 4 }, &src, &mean, &beta, &gamma);
+            test::fill_tensors(distribution_var, { 0 }, &var);
+            break;
+        }
+#endif /* ARM_COMPUTE_ENABLE_FP16 */
+        case DataType::F32:
+        {
+            const std::pair<float, float> bounds = get_batchnormalization_layer_test_bounds<float>();
+            std::uniform_real_distribution<> distribution(bounds.first, bounds.second);
+            std::uniform_real_distribution<> distribution_var(0, bounds.second);
+            test::fill_tensors(distribution, { 0, 1, 3, 4 }, &src, &mean, &beta, &gamma);
+            test::fill_tensors(distribution_var, { 0 }, &var);
+            break;
+        }
+        default:
+        {
+            ARM_COMPUTE_ERROR("Not supported");
+            break;
+        }
+    }
+
+    // Compute function
+    norm.run();
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(BatchNormalizationLayer)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, RandomBatchNormalizationLayerDataset() * boost::unit_test::data::make({ DataType::QS8, DataType::QS16, DataType::F32 }), obj, dt)
+{
+    // Set fixed point position data type allowed
+    int fixed_point_position = (arm_compute::is_data_type_fixed_point(dt)) ? 3 : 0;
+
+    // Create tensors
+    Tensor src   = create_tensor<Tensor>(obj.shape0, dt, 1, fixed_point_position);
+    Tensor dst   = create_tensor<Tensor>(obj.shape0, dt, 1, fixed_point_position);
+    Tensor mean  = create_tensor<Tensor>(obj.shape1, dt, 1, fixed_point_position);
+    Tensor var   = create_tensor<Tensor>(obj.shape1, dt, 1, fixed_point_position);
+    Tensor beta  = create_tensor<Tensor>(obj.shape1, dt, 1, fixed_point_position);
+    Tensor gamma = create_tensor<Tensor>(obj.shape1, dt, 1, fixed_point_position);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+    BOOST_TEST(mean.info()->is_resizable());
+    BOOST_TEST(var.info()->is_resizable());
+    BOOST_TEST(beta.info()->is_resizable());
+    BOOST_TEST(gamma.info()->is_resizable());
+
+    // Create and configure function
+    NEBatchNormalizationLayer norm;
+    norm.configure(&src, &dst, &mean, &var, &beta, &gamma, obj.epsilon);
+
+    // Validate valid region
+    const ValidRegion valid_region     = shape_to_valid_region(obj.shape0);
+    const ValidRegion valid_region_vec = shape_to_valid_region(obj.shape1);
+    validate(src.info()->valid_region(), valid_region);
+    validate(dst.info()->valid_region(), valid_region);
+    validate(mean.info()->valid_region(), valid_region_vec);
+    validate(var.info()->valid_region(), valid_region_vec);
+    validate(beta.info()->valid_region(), valid_region_vec);
+    validate(gamma.info()->valid_region(), valid_region_vec);
+}
+
+BOOST_AUTO_TEST_SUITE(Float)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(Random,
+                     RandomBatchNormalizationLayerDataset() * boost::unit_test::data::make(DataType::F32),
+                     obj, dt)
+{
+    // Compute function
+    Tensor dst = compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, tolerance_f32, 0);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+#ifdef ARM_COMPUTE_ENABLE_FP16
+BOOST_AUTO_TEST_SUITE(Float16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(Random,
+                     RandomBatchNormalizationLayerDataset() * boost::unit_test::data::make(DataType::F16),
+                     obj, dt)
+{
+    // Compute function
+    Tensor dst = compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, tolerance_f16, 0);
+}
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* ARM_COMPUTE_ENABLE_FP16 */
+
+BOOST_AUTO_TEST_SUITE(Quantized)
+BOOST_AUTO_TEST_SUITE(QS8)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(Random,
+                     RandomBatchNormalizationLayerDataset() * boost::unit_test::data::make(DataType::QS8) * boost::unit_test::data::xrange(1, 6),
+                     obj, dt, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, tolerance_qs8);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(QS16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(Random,
+                     RandomBatchNormalizationLayerDataset() * boost::unit_test::data::make(DataType::QS16) * boost::unit_test::data::xrange(1, 14),
+                     obj, dt, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, tolerance_qs16);
+}
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Box3x3.cpp b/tests/validation_old/NEON/Box3x3.cpp
new file mode 100644
index 0000000000..708b7de204
--- /dev/null
+++ b/tests/validation_old/NEON/Box3x3.cpp
@@ -0,0 +1,167 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+#include "tests/validation_old/ValidationUserConfiguration.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEBox3x3.h"
+#include "arm_compute/runtime/SubTensor.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+constexpr unsigned int filter_size = 3;              /** Size of the kernel/filter in number of elements. */
+constexpr BorderSize   border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */
+
+/** Compute Neon box3x3 filter.
+ *
+ * @param[in] shape                 Shape of the input and output tensors.
+ * @param[in] border_mode           BorderMode used by the input tensor.
+ * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_box3x3(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U8);
+
+    // Create and configure function
+    NEBox3x3 box3x3;
+    box3x3.configure(&src, &dst, border_mode, constant_border_value);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src), 0);
+
+    // Compute function
+    box3x3.run();
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(Box3x3)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U8);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create and configure function
+    NEBox3x3 box3x3;
+    box3x3.configure(&src, &dst, border_mode);
+
+    // Validate valid region
+    const ValidRegion src_valid_region = shape_to_valid_region(shape);
+    const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
+    validate(src.info()->valid_region(), src_valid_region);
+    validate(dst.info()->valid_region(), dst_valid_region);
+
+    // Validate padding
+    PaddingCalculator calculator(shape.x(), 8);
+    calculator.set_border_size(1);
+    calculator.set_border_mode(border_mode);
+
+    const PaddingSize dst_padding = calculator.required_padding();
+
+    calculator.set_accessed_elements(16);
+    calculator.set_access_offset(-1);
+
+    const PaddingSize src_padding = calculator.required_padding();
+
+    validate(src.info()->padding(), src_padding);
+    validate(dst.info()->padding(), dst_padding);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode)
+{
+    std::mt19937                           gen(user_config.seed.get());
+    std::uniform_int_distribution<uint8_t> distribution(0, 255);
+    const uint8_t                          border_value = distribution(gen);
+
+    // Compute function
+    Tensor dst = compute_box3x3(shape, border_mode, border_value);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_box3x3(shape, border_mode, border_value);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size));
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode)
+{
+    std::mt19937                           gen(user_config.seed.get());
+    std::uniform_int_distribution<uint8_t> distribution(0, 255);
+    const uint8_t                          border_value = distribution(gen);
+
+    // Compute function
+    Tensor dst = compute_box3x3(shape, border_mode, border_value);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_box3x3(shape, border_mode, border_value);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size));
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/DepthConvert.cpp b/tests/validation_old/NEON/DepthConvert.cpp
new file mode 100644
index 0000000000..48a2b6d3f0
--- /dev/null
+++ b/tests/validation_old/NEON/DepthConvert.cpp
@@ -0,0 +1,637 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEDepthConvert.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute Neon depth convert function.
+ *
+ * @param[in] shape                    Shape of the input and output tensors.
+ * @param[in] dt_in                    Data type of input tensor.
+ * @param[in] dt_out                   Data type of the output tensor.
+ * @param[in] policy                   Conversion policy.
+ * @param[in] shift                    Value for down/up conversions. Must be 0 <= shift < 8.
+ * @param[in] fixed_point_position_in  (Optional) Fixed point position for the input tensor.
+ * @param[in] fixed_point_position_out (Optional) Fixed point position for the output tensor.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_depth_convert(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy,
+                             uint32_t shift, uint32_t fixed_point_position_in = 0, uint32_t fixed_point_position_out = 0)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, dt_in, 1, fixed_point_position_in);
+    Tensor dst = create_tensor<Tensor>(shape, dt_out, 1, fixed_point_position_out);
+
+    // Create and configure function
+    NEDepthConvert depth_convert;
+    depth_convert.configure(&src, &dst, policy, shift);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src), 0);
+
+    // Compute function
+    depth_convert.run();
+
+    return dst;
+}
+/** Configure and validate region/padding function.
+ *
+ * @param[in] shape                    Shape of the input and output tensors.
+ * @param[in] dt_in                    Data type of input tensor.
+ * @param[in] dt_out                   Data type of the output tensor.
+ * @param[in] policy                   Conversion policy.
+ * @param[in] shift                    Value for down/up conversions. Must be 0 <= shift < 8.
+ * @param[in] fixed_point_position_in  (Optional) Fixed point position for the input tensor.
+ * @param[in] fixed_point_position_out (Optional) Fixed point position for the output tensor.
+ *
+ */
+
+void compute_configure_validate(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy,
+                                uint32_t shift, uint32_t fixed_point_position_in = 0, uint32_t fixed_point_position_out = 0)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, dt_in, 1, fixed_point_position_in);
+    Tensor dst = create_tensor<Tensor>(shape, dt_out, 1, fixed_point_position_out);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create and configure function
+    NEDepthConvert depth_convert;
+    depth_convert.configure(&src, &dst, policy, shift);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+    validate(src.info()->valid_region(), valid_region);
+    validate(dst.info()->valid_region(), valid_region);
+
+    // Validate padding
+    const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+    validate(src.info()->padding(), padding);
+    validate(dst.info()->padding(), padding);
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(DepthConvert)
+
+BOOST_AUTO_TEST_SUITE(QS8_to_QS8)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * (boost::unit_test::data::make({ 1, 3, 5, 6 }) ^ boost::unit_test::data::make({ 6, 5, 1, 3 })),
+                     shape, policy, fixed_point_position_in, fixed_point_position_out)
+{
+    // Compute configure and validate region/padding
+    compute_configure_validate(shape, DataType::QS8, DataType::QS8, policy, 0, fixed_point_position_in, fixed_point_position_out);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * (boost::unit_test::data::make({ 1, 3, 5, 6 }) ^ boost::unit_test::data::make({ 6, 5, 1, 3 })),
+                     shape, policy, fixed_point_position_in, fixed_point_position_out)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::QS8, DataType::QS8, policy, 0, fixed_point_position_in, fixed_point_position_out);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS8, DataType::QS8, policy, 0, fixed_point_position_in, fixed_point_position_out);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(QS8_to_F32)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * boost::unit_test::data::xrange(1, 7, 1),
+                     shape, policy, fixed_point_position)
+{
+    // Compute configure and validate region/padding
+    compute_configure_validate(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * boost::unit_test::data::xrange(1, 7, 1),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * boost::unit_test::data::xrange(1, 7, 1),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(F32_to_QS8)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * boost::unit_test::data::xrange(1, 7, 1),
+                     shape, policy, fixed_point_position)
+{
+    // Compute configure and validate region/padding
+    compute_configure_validate(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * boost::unit_test::data::xrange(1, 7, 1),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * boost::unit_test::data::xrange(1, 7, 1),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(QS16_to_QS16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * (boost::unit_test::data::make({ 3, 6, 7, 13, 14 }) ^ boost::unit_test::data::make({ 5, 10, 14, 4, 7 })),
+                     shape, policy, fixed_point_position_in, fixed_point_position_out)
+{
+    // Compute configure and validate region/padding
+    compute_configure_validate(shape, DataType::QS16, DataType::QS16, policy, 0, fixed_point_position_in, fixed_point_position_out);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * (boost::unit_test::data::make({ 3, 6, 7, 13, 14 }) ^ boost::unit_test::data::make({ 5, 10, 14, 4, 7 })),
+                     shape, policy, fixed_point_position_in, fixed_point_position_out)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::QS16, DataType::QS16, policy, 0, fixed_point_position_in, fixed_point_position_out);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS16, DataType::QS16, policy, 0, fixed_point_position_in, fixed_point_position_out);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(QS16_to_F32)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * boost::unit_test::data::xrange(1, 15, 1),
+                     shape, policy, fixed_point_position)
+{
+    // Compute configure and validate region/padding
+    compute_configure_validate(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * boost::unit_test::data::xrange(1, 15, 1),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * boost::unit_test::data::xrange(1, 15, 1),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(F32_to_QS16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * boost::unit_test::data::xrange(1, 7, 1),
+                     shape, policy, fixed_point_position)
+{
+    // Compute configure and validate region/padding
+    compute_configure_validate(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * boost::unit_test::data::xrange(1, 15, 1),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE })
+                     * boost::unit_test::data::xrange(1, 15, 1),
+                     shape, policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(U8_to_U16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute configure and validate region/padding
+    compute_configure_validate(shape, DataType::U8, DataType::U16, policy, shift, 0);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::U16, policy, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::U16, policy, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::U16, policy, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::U16, policy, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(U8_to_S16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute configure and validate region/padding
+    compute_configure_validate(shape, DataType::U8, DataType::S16, policy, shift);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S16, policy, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S16, policy, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S16, policy, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S16, policy, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(U8_to_S32)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute configure and validate region/padding
+    compute_configure_validate(shape, DataType::U8, DataType::S32, policy, shift);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S32, policy, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S32, policy, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S32, policy, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S32, policy, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(U16_to_U8)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute configure and validate region/padding
+    compute_configure_validate(shape, DataType::U16, DataType::U8, policy, shift);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U8, policy, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U8, policy, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U8, policy, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U8, policy, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(U16_to_U32)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute configure and validate region/padding
+    compute_configure_validate(shape, DataType::U16, DataType::U32, policy, shift);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U32, policy, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U32, policy, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U32, policy, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U32, policy, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(S16_to_U8)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute configure and validate region/padding
+    compute_configure_validate(shape, DataType::S16, DataType::U8, policy, shift);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::U8, policy, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::U8, policy, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::U8, policy, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::U8, policy, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(S16_to_S32)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute configure and validate region/padding
+    compute_configure_validate(shape, DataType::S16, DataType::S32, policy, shift);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::S32, policy, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::S32, policy, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP })
+                     * boost::unit_test::data::xrange(0, 7, 1),
+                     shape, policy, shift)
+{
+    // Compute function
+    Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::S32, policy, shift);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::S32, policy, shift);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/FillBorder.cpp b/tests/validation_old/NEON/FillBorder.cpp
new file mode 100644
index 0000000000..ad703d97fb
--- /dev/null
+++ b/tests/validation_old/NEON/FillBorder.cpp
@@ -0,0 +1,88 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/NEON/kernels/NEFillBorderKernel.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(FillBorder, BorderModes() * boost::unit_test::data::make({ PaddingSize{ 0 }, PaddingSize{ 1, 0, 1, 2 }, PaddingSize{ 10 } }), border_mode, padding)
+{
+    constexpr uint8_t border_value = 42U;
+    constexpr uint8_t tensor_value = 89U;
+    BorderSize        border_size{ 5 };
+
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(TensorShape{ 10U, 10U, 2U }, DataType::U8);
+
+    src.info()->extend_padding(padding);
+
+    // Allocate tensor
+    src.allocator()->allocate();
+
+    // Check padding is as required
+    validate(src.info()->padding(), padding);
+
+    // Fill tensor with constant value
+    std::uniform_int_distribution<uint8_t> distribution{ tensor_value, tensor_value };
+    library->fill(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());
+
+    // Validate border
+    border_size.limit(padding);
+    validate(Accessor(src), border_size, border_mode, &border_value);
+
+    // Validate tensor
+    validate(Accessor(src), &tensor_value);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Fixedpoint/Exp_QS16.cpp b/tests/validation_old/NEON/Fixedpoint/Exp_QS16.cpp
new file mode 100644
index 0000000000..66115879aa
--- /dev/null
+++ b/tests/validation_old/NEON/Fixedpoint/Exp_QS16.cpp
@@ -0,0 +1,122 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/ReferenceCPP.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/NEON/NEFixedPoint.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+const float tolerance = 1.0f; /**< Tolerance value for comparing reference's output against implementation's output */
+
+/** Compute Neon exponential function for signed 16 bit fixed point.
+ *
+ * @param[in] shape Shape of the input and output tensors.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_exp_qs16(const TensorShape &shape, int fixed_point_position)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::QS16, 1, fixed_point_position);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::QS16, 1, fixed_point_position);
+
+    constexpr unsigned int num_elems_processed_per_iteration = 8;
+    Window                 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);
+
+    update_window_and_padding(window, input_access, output_access);
+    output_access.set_valid_region(window, src.info()->valid_region());
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors. Keep the range between [-1.0, 1.0) so the result won't
+    // overflow.
+    std::uniform_int_distribution<> distribution(-(1 << (fixed_point_position - 1)), (1 << (fixed_point_position - 1)));
+    library->fill(Accessor(src), distribution, 0);
+
+    Iterator input(&src, window);
+    Iterator output(&dst, window);
+
+    execute_window_loop(window, [&](const Coordinates & id)
+    {
+        qint16x8_t in = vld1q_qs16(reinterpret_cast<const qint16_t *>(input.ptr()));
+        // Use saturated exp
+        vst1q_qs16(reinterpret_cast<qint16_t *>(output.ptr()), vqexpq_qs16(in, fixed_point_position));
+    },
+    input, output);
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(FixedPoint)
+BOOST_AUTO_TEST_SUITE(QS16)
+BOOST_AUTO_TEST_SUITE(Exp)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(1, 15), shape, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_exp_qs16(shape, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS16, DataType::QS16, FixedPointOp::EXP, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, tolerance, 0);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Fixedpoint/Exp_QS8.cpp b/tests/validation_old/NEON/Fixedpoint/Exp_QS8.cpp
new file mode 100644
index 0000000000..9e8096fa75
--- /dev/null
+++ b/tests/validation_old/NEON/Fixedpoint/Exp_QS8.cpp
@@ -0,0 +1,122 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/ReferenceCPP.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/NEON/NEFixedPoint.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+const float tolerance = 0.0f; /**< Tolerance value for comparing reference's output against implementation's output */
+
+/** Compute Neon exponential function for signed 8bit fixed point.
+ *
+ * @param[in] shape Shape of the input and output tensors.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_exp_qs8(const TensorShape &shape, int fixed_point_position)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::QS8, 1, fixed_point_position);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::QS8, 1, fixed_point_position);
+
+    constexpr unsigned int num_elems_processed_per_iteration = 16;
+    Window                 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);
+
+    update_window_and_padding(window, input_access, output_access);
+    output_access.set_valid_region(window, src.info()->valid_region());
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors. Keep the range between [-1.0, 1.0) so the result won't
+    // overflow. E.g. e^7 = 1096, which cannot be represented in QS8
+    std::uniform_int_distribution<> distribution(-(1 << (fixed_point_position - 1)), (1 << (fixed_point_position - 1)));
+    library->fill(Accessor(src), distribution, 0);
+
+    Iterator input(&src, window);
+    Iterator output(&dst, window);
+
+    execute_window_loop(window, [&](const Coordinates & id)
+    {
+        qint8x16_t in = vld1q_s8(reinterpret_cast<const qint8_t *>(input.ptr()));
+        // Use saturated exp
+        vst1q_s8(reinterpret_cast<qint8_t *>(output.ptr()), vqexpq_qs8(in, fixed_point_position));
+    },
+    input, output);
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(FixedPoint)
+BOOST_AUTO_TEST_SUITE(QS8)
+BOOST_AUTO_TEST_SUITE(Exp)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(1, 7), shape, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_exp_qs8(shape, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS8, DataType::QS8, FixedPointOp::EXP, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, tolerance, 0);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Fixedpoint/Invsqrt_QS16.cpp b/tests/validation_old/NEON/Fixedpoint/Invsqrt_QS16.cpp
new file mode 100644
index 0000000000..f56707a93d
--- /dev/null
+++ b/tests/validation_old/NEON/Fixedpoint/Invsqrt_QS16.cpp
@@ -0,0 +1,122 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/ReferenceCPP.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/NEON/NEFixedPoint.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+const float tolerance = 5.0f; /**< Tolerance value for comparing reference's output against implementation's output */
+
+/** Compute Neon inverse square root function for signed 16 bit fixed point.
+ *
+ * @param[in] shape Shape of the input and output tensors.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_invsqrt_qs16(const TensorShape &shape, int fixed_point_position)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::QS16, 1, fixed_point_position);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::QS16, 1, fixed_point_position);
+
+    constexpr unsigned int num_elems_processed_per_iteration = 8;
+    Window                 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);
+
+    update_window_and_padding(window, input_access, output_access);
+    output_access.set_valid_region(window, src.info()->valid_region());
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors. Keep the range between [1, 0x7FFF)
+    std::uniform_int_distribution<> distribution(1, 0x7FFF);
+    library->fill(Accessor(src), distribution, 0);
+
+    Iterator input(&src, window);
+    Iterator output(&dst, window);
+
+    execute_window_loop(window, [&](const Coordinates & id)
+    {
+        qint16x8_t in = vld1q_qs16(reinterpret_cast<const qint16_t *>(input.ptr()));
+        vst1q_qs16(reinterpret_cast<qint16_t *>(output.ptr()), vqinvsqrtq_qs16(in, fixed_point_position));
+    },
+    input, output);
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(FixedPoint)
+BOOST_AUTO_TEST_SUITE(QS16)
+BOOST_AUTO_TEST_SUITE(Invsqrt)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunSmall, boost::unit_test::data::xrange(1, 14), fixed_point_position)
+{
+    TensorShape shape(8192U);
+
+    // Compute function
+    Tensor dst = compute_invsqrt_qs16(shape, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS16, DataType::QS16, FixedPointOp::INV_SQRT, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, tolerance, 0);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Fixedpoint/Invsqrt_QS8.cpp b/tests/validation_old/NEON/Fixedpoint/Invsqrt_QS8.cpp
new file mode 100644
index 0000000000..fb33fd4632
--- /dev/null
+++ b/tests/validation_old/NEON/Fixedpoint/Invsqrt_QS8.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 "NEON/Accessor.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/ReferenceCPP.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/NEON/NEFixedPoint.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+const float tolerance = 4.0f; /**< Tolerance value for comparing reference's output against implementation's output */
+
+/** Compute Neon inverse square root function for signed 8bit fixed point.
+ *
+ * @param[in] shape Shape of the input and output tensors.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_invsqrt_qs8(const TensorShape &shape, int fixed_point_position)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::QS8, 1, fixed_point_position);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::QS8, 1, fixed_point_position);
+
+    constexpr unsigned int num_elems_processed_per_iteration = 16;
+    Window                 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);
+
+    update_window_and_padding(window, input_access, output_access);
+    output_access.set_valid_region(window, src.info()->valid_region());
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors. Keep the range between [1, 127).
+    std::uniform_int_distribution<> distribution(1, 127);
+    library->fill(Accessor(src), distribution, 0);
+
+    Iterator input(&src, window);
+    Iterator output(&dst, window);
+
+    execute_window_loop(window, [&](const Coordinates & id)
+    {
+        qint8x16_t in = vld1q_s8(reinterpret_cast<const qint8_t *>(input.ptr()));
+        vst1q_s8(reinterpret_cast<qint8_t *>(output.ptr()), vqinvsqrtq_qs8(in, fixed_point_position));
+    },
+    input, output);
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(FixedPoint)
+BOOST_AUTO_TEST_SUITE(QS8)
+BOOST_AUTO_TEST_SUITE(Invsqrt)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Small1DShape, SmallShapes() * boost::unit_test::data::xrange(1, 6), shape, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_invsqrt_qs8(shape, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS8, DataType::QS8, FixedPointOp::INV_SQRT, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, tolerance, 0);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Fixedpoint/Log_QS16.cpp b/tests/validation_old/NEON/Fixedpoint/Log_QS16.cpp
new file mode 100644
index 0000000000..6485b2031c
--- /dev/null
+++ b/tests/validation_old/NEON/Fixedpoint/Log_QS16.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 "NEON/Accessor.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/ReferenceCPP.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/NEON/NEFixedPoint.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+const float tolerance = 7.0f; /**< Tolerance value for comparing reference's output against implementation's output */
+
+/** Compute Neon logarithm function for signed 16 bit fixed point.
+ *
+ * @param[in] shape Shape of the input and output tensors.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_log_qs16(const TensorShape &shape, int fixed_point_position)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::QS16, 1, fixed_point_position);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::QS16, 1, fixed_point_position);
+
+    constexpr unsigned int num_elems_processed_per_iteration = 8;
+    Window                 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);
+
+    update_window_and_padding(window, input_access, output_access);
+    output_access.set_valid_region(window, src.info()->valid_region());
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors. Keep the range between [(1 << (fixed_point_position - 1), 0x3FFF) so the result won't
+    // overflow.
+    std::uniform_int_distribution<> distribution((1 << (fixed_point_position - 1)), 0x3FFF);
+    library->fill(Accessor(src), distribution, 0);
+
+    Iterator input(&src, window);
+    Iterator output(&dst, window);
+
+    execute_window_loop(window, [&](const Coordinates & id)
+    {
+        qint16x8_t in = vld1q_qs16(reinterpret_cast<const qint16_t *>(input.ptr()));
+        vst1q_qs16(reinterpret_cast<qint16_t *>(output.ptr()), vlogq_qs16(in, fixed_point_position));
+    },
+    input, output);
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(FixedPoint)
+BOOST_AUTO_TEST_SUITE(QS16)
+BOOST_AUTO_TEST_SUITE(Log)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(4, 14), shape, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_log_qs16(shape, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS16, DataType::QS16, FixedPointOp::LOG, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, tolerance, 0);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Fixedpoint/Log_QS8.cpp b/tests/validation_old/NEON/Fixedpoint/Log_QS8.cpp
new file mode 100644
index 0000000000..21012c52b0
--- /dev/null
+++ b/tests/validation_old/NEON/Fixedpoint/Log_QS8.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 "NEON/Accessor.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/ReferenceCPP.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/NEON/NEFixedPoint.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+const float tolerance = 5; /**< Tolerance value for comparing reference's output against implementation's output */
+
+/** Compute Neon logarithm function for signed 8bit fixed point.
+ *
+ * @param[in] shape Shape of the input and output tensors.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_log_qs8(const TensorShape &shape, int fixed_point_position)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::QS8, 1, fixed_point_position);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::QS8, 1, fixed_point_position);
+
+    constexpr unsigned int num_elems_processed_per_iteration = 16;
+    Window                 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);
+
+    update_window_and_padding(window, input_access, output_access);
+    output_access.set_valid_region(window, src.info()->valid_region());
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // 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.
+    std::uniform_int_distribution<> distribution((1 << (fixed_point_position - 1)), 0x3F);
+    library->fill(Accessor(src), distribution, 0);
+
+    Iterator input(&src, window);
+    Iterator output(&dst, window);
+
+    execute_window_loop(window, [&](const Coordinates & id)
+    {
+        qint8x16_t in = vld1q_s8(reinterpret_cast<const qint8_t *>(input.ptr()));
+        vst1q_s8(reinterpret_cast<qint8_t *>(output.ptr()), vlogq_qs8(in, fixed_point_position));
+    },
+    input, output);
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(FixedPoint)
+BOOST_AUTO_TEST_SUITE(QS8)
+BOOST_AUTO_TEST_SUITE(Log)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(3, 6), shape, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_log_qs8(shape, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS8, DataType::QS8, FixedPointOp::LOG, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, tolerance, 0);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Fixedpoint/Reciprocal_QS16.cpp b/tests/validation_old/NEON/Fixedpoint/Reciprocal_QS16.cpp
new file mode 100644
index 0000000000..5630a3391a
--- /dev/null
+++ b/tests/validation_old/NEON/Fixedpoint/Reciprocal_QS16.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 "NEON/Accessor.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/ReferenceCPP.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/NEON/NEFixedPoint.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+const float tolerance = 11.0f; /**< Tolerance value for comparing reference's output against implementation's output. */
+
+/** Compute Neon reciprocal function for signed 16 bit fixed point.
+ *
+ * @param[in] shape Shape of the input and output tensors.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_reciprocal_qs16(const TensorShape &shape, int fixed_point_position)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::QS16, 1, fixed_point_position);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::QS16, 1, fixed_point_position);
+
+    constexpr unsigned int num_elems_processed_per_iteration = 8;
+    Window                 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);
+
+    update_window_and_padding(window, input_access, output_access);
+    output_access.set_valid_region(window, src.info()->valid_region());
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors. Keep the range between [15, 0x7FFF) so the result won't
+    // overflow.
+    std::uniform_int_distribution<> distribution(15, 0x7FFF);
+    library->fill(Accessor(src), distribution, 0);
+
+    Iterator input(&src, window);
+    Iterator output(&dst, window);
+
+    execute_window_loop(window, [&](const Coordinates & id)
+    {
+        qint16x8_t in = vld1q_qs16(reinterpret_cast<const qint16_t *>(input.ptr()));
+        vst1q_qs16(reinterpret_cast<qint16_t *>(output.ptr()), vqrecipq_qs16(in, fixed_point_position));
+    },
+    input, output);
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(FixedPoint)
+BOOST_AUTO_TEST_SUITE(QS16)
+BOOST_AUTO_TEST_SUITE(Reciprocal)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(1, 14), shape, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_reciprocal_qs16(shape, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS16, DataType::QS16, FixedPointOp::RECIPROCAL, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, tolerance, 0);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Fixedpoint/Reciprocal_QS8.cpp b/tests/validation_old/NEON/Fixedpoint/Reciprocal_QS8.cpp
new file mode 100644
index 0000000000..23f98acc40
--- /dev/null
+++ b/tests/validation_old/NEON/Fixedpoint/Reciprocal_QS8.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 "NEON/Accessor.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/ReferenceCPP.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/NEON/NEFixedPoint.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+const float tolerance = 3; /**< Tolerance value for comparing reference's output against implementation's output */
+
+/** Compute Neon reciprocal function for signed 8bit fixed point.
+ *
+ * @param[in] shape Shape of the input and output tensors.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_reciprocal_qs8(const TensorShape &shape, int fixed_point_position)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::QS8, 1, fixed_point_position);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::QS8, 1, fixed_point_position);
+
+    constexpr unsigned int num_elems_processed_per_iteration = 16;
+    Window                 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);
+
+    update_window_and_padding(window, input_access, output_access);
+    output_access.set_valid_region(window, src.info()->valid_region());
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // 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.
+    std::uniform_int_distribution<> distribution(15, 0x7F);
+    library->fill(Accessor(src), distribution, 0);
+
+    Iterator input(&src, window);
+    Iterator output(&dst, window);
+
+    execute_window_loop(window, [&](const Coordinates & id)
+    {
+        qint8x16_t in = vld1q_s8(reinterpret_cast<const qint8_t *>(input.ptr()));
+        vst1q_s8(reinterpret_cast<qint8_t *>(output.ptr()), vrecipq_qs8(in, fixed_point_position));
+    },
+    input, output);
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(FixedPoint)
+BOOST_AUTO_TEST_SUITE(QS8)
+BOOST_AUTO_TEST_SUITE(Reciprocal)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(1, 6), shape, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_reciprocal_qs8(shape, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS8, DataType::QS8, FixedPointOp::RECIPROCAL, fixed_point_position);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, tolerance, 0);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Gaussian3x3.cpp b/tests/validation_old/NEON/Gaussian3x3.cpp
new file mode 100644
index 0000000000..becd9196ea
--- /dev/null
+++ b/tests/validation_old/NEON/Gaussian3x3.cpp
@@ -0,0 +1,167 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+#include "tests/validation_old/ValidationUserConfiguration.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEGaussian3x3.h"
+#include "arm_compute/runtime/SubTensor.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+constexpr unsigned int filter_size = 3;              /** Size of the kernel/filter in number of elements. */
+constexpr BorderSize   border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */
+
+/** Compute Neon gaussian3x3 filter.
+ *
+ * @param[in] shape                 Shape of the input and output tensors.
+ * @param[in] border_mode           BorderMode used by the input tensor.
+ * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_gaussian3x3(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U8);
+
+    // Create and configure function
+    NEGaussian3x3 gaussian3x3;
+    gaussian3x3.configure(&src, &dst, border_mode, constant_border_value);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src), 0);
+
+    // Compute function
+    gaussian3x3.run();
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(Gaussian3x3)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U8);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create and configure function
+    NEGaussian3x3 gaussian3x3;
+    gaussian3x3.configure(&src, &dst, border_mode);
+
+    // Validate valid region
+    const ValidRegion src_valid_region = shape_to_valid_region(shape);
+    const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
+    validate(src.info()->valid_region(), src_valid_region);
+    validate(dst.info()->valid_region(), dst_valid_region);
+
+    // Validate padding
+    PaddingCalculator calculator(shape.x(), 8);
+    calculator.set_border_size(1);
+    calculator.set_border_mode(border_mode);
+
+    const PaddingSize dst_padding = calculator.required_padding();
+
+    calculator.set_accessed_elements(16);
+    calculator.set_access_offset(-1);
+
+    const PaddingSize src_padding = calculator.required_padding();
+
+    validate(src.info()->padding(), src_padding);
+    validate(dst.info()->padding(), dst_padding);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode)
+{
+    std::mt19937                           gen(user_config.seed.get());
+    std::uniform_int_distribution<uint8_t> distribution(0, 255);
+    const uint8_t                          border_value = distribution(gen);
+
+    // Compute function
+    Tensor dst = compute_gaussian3x3(shape, border_mode, border_value);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_gaussian3x3(shape, border_mode, border_value);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size));
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode)
+{
+    std::mt19937                           gen(user_config.seed.get());
+    std::uniform_int_distribution<uint8_t> distribution(0, 255);
+    const uint8_t                          border_value = distribution(gen);
+
+    // Compute function
+    Tensor dst = compute_gaussian3x3(shape, border_mode, border_value);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_gaussian3x3(shape, border_mode, border_value);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size));
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Gaussian5x5.cpp b/tests/validation_old/NEON/Gaussian5x5.cpp
new file mode 100644
index 0000000000..240285afb6
--- /dev/null
+++ b/tests/validation_old/NEON/Gaussian5x5.cpp
@@ -0,0 +1,167 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+#include "tests/validation_old/ValidationUserConfiguration.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEGaussian5x5.h"
+#include "arm_compute/runtime/SubTensor.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+constexpr unsigned int filter_size = 5;              /** Size of the kernel/filter in number of elements. */
+constexpr BorderSize   border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */
+
+/** Compute Neon gaussian5x5 filter.
+ *
+ * @param[in] shape                 Shape of the input and output tensors.
+ * @param[in] border_mode           BorderMode used by the input tensor.
+ * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_gaussian5x5(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U8);
+
+    // Create and configure function
+    NEGaussian5x5 gaussian5x5;
+    gaussian5x5.configure(&src, &dst, border_mode, constant_border_value);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src), 0);
+
+    // Compute function
+    gaussian5x5.run();
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(Gaussian5x5)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U8);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create and configure function
+    NEGaussian5x5 gaussian5x5;
+    gaussian5x5.configure(&src, &dst, border_mode);
+
+    // Validate valid region
+    const ValidRegion src_valid_region = shape_to_valid_region(shape);
+    const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
+    validate(src.info()->valid_region(), src_valid_region);
+    validate(dst.info()->valid_region(), dst_valid_region);
+
+    // Validate padding
+    PaddingCalculator calculator(shape.x(), 16);
+    calculator.set_border_size(2);
+    calculator.set_border_mode(border_mode);
+
+    const PaddingSize dst_padding = calculator.required_padding();
+
+    calculator.set_processed_elements(8);
+    calculator.set_access_offset(-2);
+
+    const PaddingSize src_padding = calculator.required_padding();
+
+    validate(src.info()->padding(), src_padding);
+    validate(dst.info()->padding(), dst_padding);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode)
+{
+    std::mt19937                           gen(user_config.seed.get());
+    std::uniform_int_distribution<uint8_t> distribution(0, 255);
+    const uint8_t                          border_value = distribution(gen);
+
+    // Compute function
+    Tensor dst = compute_gaussian5x5(shape, border_mode, border_value);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_gaussian5x5(shape, border_mode, border_value);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size));
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode)
+{
+    std::mt19937                           gen(user_config.seed.get());
+    std::uniform_int_distribution<uint8_t> distribution(0, 255);
+    const uint8_t                          border_value = distribution(gen);
+
+    // Compute function
+    Tensor dst = compute_gaussian5x5(shape, border_mode, border_value);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_gaussian5x5(shape, border_mode, border_value);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size));
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/HarrisCorners.cpp b/tests/validation_old/NEON/HarrisCorners.cpp
new file mode 100644
index 0000000000..809e61c053
--- /dev/null
+++ b/tests/validation_old/NEON/HarrisCorners.cpp
@@ -0,0 +1,229 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "NEON/Helper.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+#include "tests/validation_old/ValidationUserConfiguration.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEHarrisCorners.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "PaddingCalculator.h"
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute Neon Harris corners function.
+ *
+ * @param[in] shape                 Shape of input tensor
+ * @param[in] threshold             Minimum threshold with which to eliminate Harris Corner scores (computed using the normalized Sobel kernel).
+ * @param[in] min_dist              Radial Euclidean distance for the euclidean distance stage
+ * @param[in] sensitivity           Sensitivity threshold k from the Harris-Stephens equation
+ * @param[in] gradient_size         The gradient window size to use on the input. The implementation supports 3, 5, and 7
+ * @param[in] block_size            The block window size used to compute the Harris Corner score. The implementation supports 3, 5, and 7.
+ * @param[in] border_mode           Border mode to use
+ * @param[in] constant_border_value Constant value to use for borders if border_mode is set to CONSTANT.
+ * @param[in] use_fp16              If true the FP16 kernels will be used. If false F32 kernels are used.
+ *
+ * @return Computed corners' keypoints.
+ */
+KeyPointArray compute_harris_corners(const TensorShape &shape, float threshold, float min_dist, float sensitivity,
+                                     int32_t gradient_size, int32_t block_size, BorderMode border_mode, uint8_t constant_border_value, bool use_fp16)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    src.info()->set_format(Format::U8);
+
+    // Create array of keypoints
+    KeyPointArray corners(shape.total_size());
+
+    // Create harris corners configure function
+    NEHarrisCorners harris_corners;
+    harris_corners.configure(&src, threshold, min_dist, sensitivity, gradient_size, block_size, &corners, border_mode, constant_border_value, use_fp16);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src), 0);
+
+    // Compute function
+    harris_corners.run();
+
+    return corners;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(HarrisCorners)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (Small2DShapes() + Large2DShapes()) * BorderModes()
+                     * boost::unit_test::data::make({ 3, 5, 7 }) * boost::unit_test::data::make({ 3, 5, 7 }),
+                     shape, border_mode, gradient, block)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    src.info()->set_format(Format::U8);
+
+    KeyPointArray corners;
+
+    uint8_t constant_border_value = 0;
+
+    std::mt19937                          gen(user_config.seed.get());
+    std::uniform_real_distribution<float> real_dist(0.01, std::numeric_limits<float>::min());
+
+    const float threshold              = real_dist(gen);
+    const float sensitivity            = real_dist(gen);
+    const float max_euclidean_distance = 30.f;
+
+    real_dist            = std::uniform_real_distribution<float>(0.f, max_euclidean_distance);
+    const float min_dist = real_dist(gen);
+
+    // 50% chance to use fp16
+    bool use_fp16 = real_dist(gen) < max_euclidean_distance / 2 ? true : false;
+
+    // Generate a random constant value if border_mode is constant
+    if(border_mode == BorderMode::CONSTANT)
+    {
+        std::uniform_int_distribution<uint8_t> int_dist(0, 255);
+        constant_border_value = int_dist(gen);
+    }
+
+    BOOST_TEST(src.info()->is_resizable());
+
+    // Create harris corners configure function
+    NEHarrisCorners harris_corners;
+    harris_corners.configure(&src, threshold, min_dist, sensitivity, gradient, block, &corners, border_mode, constant_border_value, use_fp16);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+
+    validate(src.info()->valid_region(), valid_region);
+
+    // Validate padding
+    PaddingCalculator calculator(shape.x(), 8);
+
+    calculator.set_border_mode(border_mode);
+    calculator.set_border_size(gradient / 2);
+    calculator.set_access_offset(-gradient / 2);
+    calculator.set_accessed_elements(16);
+
+    const PaddingSize padding = calculator.required_padding();
+
+    validate(src.info()->padding(), padding);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes() * BorderModes() * boost::unit_test::data::make({ 3, 5, 7 }) * boost::unit_test::data::make({ 3, 5, 7 }), shape, border_mode, gradient, block)
+{
+    uint8_t constant_border_value = 0;
+
+    std::mt19937                          gen(user_config.seed.get());
+    std::uniform_real_distribution<float> real_dist(0.01, std::numeric_limits<float>::min());
+
+    const float threshold              = real_dist(gen);
+    const float sensitivity            = real_dist(gen);
+    const float max_euclidean_distance = 30.f;
+
+    real_dist            = std::uniform_real_distribution<float>(0.f, max_euclidean_distance);
+    const float min_dist = real_dist(gen);
+
+    // 50% chance to use fp16
+    bool use_fp16 = real_dist(gen) < max_euclidean_distance / 2 ? true : false;
+
+    // Generate a random constant value if border_mode is constant
+    if(border_mode == BorderMode::CONSTANT)
+    {
+        std::uniform_int_distribution<uint8_t> int_dist(0, 255);
+        constant_border_value = int_dist(gen);
+    }
+
+    // Compute function
+    KeyPointArray dst = compute_harris_corners(shape, threshold, min_dist, sensitivity, gradient, block, border_mode, constant_border_value, use_fp16);
+
+    // Compute reference
+    KeyPointArray ref_dst = Reference::compute_reference_harris_corners(shape, threshold, min_dist, sensitivity, gradient, block, border_mode, constant_border_value);
+
+    // Validate output
+    validate(dst, ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, Large2DShapes() * BorderModes() * boost::unit_test::data::make({ 3, 5, 7 }) * boost::unit_test::data::make({ 3, 5, 7 }), shape, border_mode, gradient, block)
+{
+    uint8_t constant_border_value = 0;
+
+    std::mt19937                          gen(user_config.seed.get());
+    std::uniform_real_distribution<float> real_dist(0.01, std::numeric_limits<float>::min());
+
+    const float threshold              = real_dist(gen);
+    const float sensitivity            = real_dist(gen);
+    const float max_euclidean_distance = 30.f;
+
+    real_dist      = std::uniform_real_distribution<float>(0.f, max_euclidean_distance);
+    float min_dist = real_dist(gen);
+
+    // 50% chance to use fp16
+    bool use_fp16 = real_dist(gen) < max_euclidean_distance / 2 ? true : false;
+
+    // Generate a random constant value if border_mode is constant
+    if(border_mode == BorderMode::CONSTANT)
+    {
+        std::uniform_int_distribution<uint8_t> int_dist(0, 255);
+        constant_border_value = int_dist(gen);
+    }
+
+    // Compute function
+    KeyPointArray dst = compute_harris_corners(shape, threshold, min_dist, sensitivity, gradient, block, border_mode, constant_border_value, use_fp16);
+
+    // Compute reference
+    KeyPointArray ref_dst = Reference::compute_reference_harris_corners(shape, threshold, min_dist, sensitivity, gradient, block, border_mode, constant_border_value);
+
+    // Validate output
+    validate(dst, ref_dst);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/IntegralImage.cpp b/tests/validation_old/NEON/IntegralImage.cpp
new file mode 100644
index 0000000000..69654b2585
--- /dev/null
+++ b/tests/validation_old/NEON/IntegralImage.cpp
@@ -0,0 +1,144 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEIntegralImage.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute Neon integral image function.
+ *
+ * @param[in] shape Shape of the input and output tensors.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_integral_image(const TensorShape &shape)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U32);
+
+    // Create integral image configure function
+    NEIntegralImage integral_image;
+    integral_image.configure(&src, &dst);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src), 0);
+
+    // Compute function
+    integral_image.run();
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(IntegralImage)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, SmallShapes() + LargeShapes(), shape)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U32);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create integral image configure function
+    NEIntegralImage integral_image;
+    integral_image.configure(&src, &dst);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+    validate(src.info()->valid_region(), valid_region);
+    validate(dst.info()->valid_region(), valid_region);
+
+    // Validate padding
+    const PaddingSize src_padding = PaddingCalculator(shape.x(), 16).required_padding();
+    const PaddingSize dst_padding(1, src_padding.right, 0, 1);
+
+    validate(src.info()->padding(), src_padding);
+    validate(dst.info()->padding(), dst_padding);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), shape)
+{
+    // Compute function
+    Tensor dst = compute_integral_image(shape);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_integral_image(shape);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes(), shape)
+{
+    // Compute function
+    Tensor dst = compute_integral_image(shape);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_integral_image(shape);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/MinMaxLocation.cpp b/tests/validation_old/NEON/MinMaxLocation.cpp
new file mode 100644
index 0000000000..c41745a636
--- /dev/null
+++ b/tests/validation_old/NEON/MinMaxLocation.cpp
@@ -0,0 +1,224 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "NEON/Helper.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEMinMaxLocation.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute Neon MinMaxLocation function.
+ *
+ * @param[in]  shape     Shape of the input and output tensors.
+ * @param[in]  dt_in     Data type of first input tensor.
+ * @param[out] min       Minimum value of tensor
+ * @param[out] max       Maximum value of tensor
+ * @param[out] min_loc   Array with locations of minimum values
+ * @param[out] max_loc   Array with locations of maximum values
+ * @param[out] min_count Number of minimum values found
+ * @param[out] max_count Number of maximum values found
+ *
+ * @return Computed output tensor.
+ */
+
+void compute_min_max_location(const TensorShape &shape, DataType dt_in, void *min, void *max,
+                              Coordinates2DArray &min_loc, Coordinates2DArray &max_loc, uint32_t &min_count, uint32_t &max_count)
+{
+    // Create tensor
+    Tensor src = create_tensor<Tensor>(shape, dt_in);
+
+    // Create and configure min_max_location configure function
+    NEMinMaxLocation min_max_loc;
+    min_max_loc.configure(&src, min, max, &min_loc, &max_loc, &min_count, &max_count);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src), 0);
+
+    // Compute function
+    min_max_loc.run();
+}
+
+void validate_configuration(const Tensor &src, TensorShape shape)
+{
+    BOOST_TEST(src.info()->is_resizable());
+
+    // Create output storage
+    int32_t            min;
+    int32_t            max;
+    Coordinates2DArray min_loc;
+    Coordinates2DArray max_loc;
+    uint32_t           min_count;
+    uint32_t           max_count;
+
+    // Create and configure function
+    NEMinMaxLocation min_max_loc;
+    min_max_loc.configure(&src, &min, &max, &min_loc, &max_loc, &min_count, &max_count);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+    validate(src.info()->valid_region(), valid_region);
+
+    // Validate padding
+    const PaddingSize padding = PaddingCalculator(shape.x(), 1).required_padding();
+    validate(src.info()->padding(), padding);
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(MinMaxLocation)
+
+BOOST_AUTO_TEST_SUITE(Integer)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (Small2DShapes() + Large2DShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }),
+                     shape, dt)
+{
+    // Create tensor
+    Tensor src = create_tensor<Tensor>(shape, dt);
+    src.info()->set_format(dt == DataType::U8 ? Format::U8 : Format::S16);
+
+    validate_configuration(src, shape);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }),
+                     shape, dt)
+{
+    // Create output storage
+    int32_t            min;
+    int32_t            max;
+    Coordinates2DArray min_loc(shape.total_size());
+    Coordinates2DArray max_loc(shape.total_size());
+    uint32_t           min_count;
+    uint32_t           max_count;
+
+    int32_t            ref_min;
+    int32_t            ref_max;
+    Coordinates2DArray ref_min_loc(shape.total_size());
+    Coordinates2DArray ref_max_loc(shape.total_size());
+    uint32_t           ref_min_count;
+    uint32_t           ref_max_count;
+
+    // Compute function
+    compute_min_max_location(shape, dt, &min, &max, min_loc, max_loc, min_count, max_count);
+
+    // Compute reference
+    Reference::compute_reference_min_max_location(shape, dt, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count);
+
+    // Validate output
+    validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, Large2DShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }),
+                     shape, dt)
+{
+    // Create output storage
+    int32_t            min;
+    int32_t            max;
+    Coordinates2DArray min_loc(shape.total_size());
+    Coordinates2DArray max_loc(shape.total_size());
+    uint32_t           min_count;
+    uint32_t           max_count;
+
+    int32_t            ref_min;
+    int32_t            ref_max;
+    Coordinates2DArray ref_min_loc(shape.total_size());
+    Coordinates2DArray ref_max_loc(shape.total_size());
+    uint32_t           ref_min_count;
+    uint32_t           ref_max_count;
+
+    // Compute function
+    compute_min_max_location(shape, dt, &min, &max, min_loc, max_loc, min_count, max_count);
+
+    // Compute reference
+    Reference::compute_reference_min_max_location(shape, dt, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count);
+
+    // Validate output
+    validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(Float)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes() * DataType::F32,
+                     shape, dt)
+{
+    // Create output storage
+    float              min;
+    float              max;
+    Coordinates2DArray min_loc(shape.total_size());
+    Coordinates2DArray max_loc(shape.total_size());
+    uint32_t           min_count;
+    uint32_t           max_count;
+
+    float              ref_min;
+    float              ref_max;
+    Coordinates2DArray ref_min_loc(shape.total_size());
+    Coordinates2DArray ref_max_loc(shape.total_size());
+    uint32_t           ref_min_count;
+    uint32_t           ref_max_count;
+
+    // Compute function
+    compute_min_max_location(shape, dt, &min, &max, min_loc, max_loc, min_count, max_count);
+
+    // Compute reference
+    Reference::compute_reference_min_max_location(shape, dt, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count);
+
+    // Validate output
+    validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/NonLinearFilter.cpp b/tests/validation_old/NEON/NonLinearFilter.cpp
new file mode 100644
index 0000000000..acc90a436a
--- /dev/null
+++ b/tests/validation_old/NEON/NonLinearFilter.cpp
@@ -0,0 +1,203 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Helpers.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+#include "tests/validation_old/ValidationUserConfiguration.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NENonLinearFilter.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute NonLinearFilter function.
+ *
+     * @param[in] input                 Shape of the input and output tensors.
+     * @param[in] function              Non linear function to perform
+     * @param[in] mask_size             Mask size. Supported sizes: 3, 5
+     * @param[in] pattern               Mask pattern
+     * @param[in] mask                  The given mask. Will be used only if pattern is specified to PATTERN_OTHER
+     * @param[in] border_mode           Strategy to use for borders.
+     * @param[in] constant_border_value (Optional) Constant value to use for borders if border_mode is set to CONSTANT.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_non_linear_filter(const TensorShape &shape, NonLinearFilterFunction function, unsigned int mask_size,
+                                 MatrixPattern pattern, const uint8_t *mask, BorderMode border_mode,
+                                 uint8_t constant_border_value)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U8);
+
+    // Create and configure function
+    NENonLinearFilter filter;
+    filter.configure(&src, &dst, function, mask_size, pattern, mask, border_mode, constant_border_value);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src), 0);
+
+    // Compute function
+    filter.run();
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(NonLinearFilter)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes())
+                     * NonLinearFilterFunctions() * boost::unit_test::data::make({ 3U, 5U })
+                     * MatrixPatterns() * BorderModes(),
+                     shape, function, mask_size, pattern, border_mode)
+{
+    std::mt19937                           generator(user_config.seed.get());
+    std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+    const uint8_t                          constant_border_value = distribution_u8(generator);
+
+    // Create the mask
+    uint8_t mask[mask_size * mask_size];
+    fill_mask_from_pattern(mask, mask_size, mask_size, pattern);
+    const auto half_mask_size = static_cast<int>(mask_size / 2);
+
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U8);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create and configure function
+    NENonLinearFilter filter;
+    filter.configure(&src, &dst, function, mask_size, pattern, mask, border_mode, constant_border_value);
+
+    // Validate valid region
+    const ValidRegion src_valid_region = shape_to_valid_region(shape);
+    const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, BorderSize(half_mask_size));
+
+    validate(src.info()->valid_region(), src_valid_region);
+    validate(dst.info()->valid_region(), dst_valid_region);
+
+    // Validate padding
+    PaddingCalculator calculator(shape.x(), ((MatrixPattern::OTHER == pattern) ? 1 : 8));
+    calculator.set_border_mode(border_mode);
+    calculator.set_border_size(half_mask_size);
+
+    const PaddingSize write_padding = calculator.required_padding(PaddingCalculator::Option::EXCLUDE_BORDER);
+
+    calculator.set_accessed_elements(16);
+    calculator.set_access_offset(-half_mask_size);
+
+    const PaddingSize read_padding = calculator.required_padding(PaddingCalculator::Option::INCLUDE_BORDER);
+
+    validate(src.info()->padding(), read_padding);
+    validate(dst.info()->padding(), write_padding);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes()
+                     * NonLinearFilterFunctions() * boost::unit_test::data::make({ 3U, 5U })
+                     * MatrixPatterns() * BorderModes(),
+                     shape, function, mask_size, pattern, border_mode)
+{
+    std::mt19937                           generator(user_config.seed.get());
+    std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+    const uint8_t                          constant_border_value = distribution_u8(generator);
+
+    // Create the mask
+    uint8_t mask[mask_size * mask_size];
+    fill_mask_from_pattern(mask, mask_size, mask_size, pattern);
+
+    // Compute function
+    Tensor dst = compute_non_linear_filter(shape, function, mask_size, pattern, mask, border_mode, constant_border_value);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_non_linear_filter(shape, function, mask_size, pattern, mask, border_mode, constant_border_value);
+
+    // Calculate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, BorderSize(static_cast<int>(mask_size / 2)));
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, valid_region);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes()
+                     * NonLinearFilterFunctions() * boost::unit_test::data::make({ 3U, 5U })
+                     * MatrixPatterns() * BorderModes(),
+                     shape, function, mask_size, pattern, border_mode)
+{
+    std::mt19937                           generator(user_config.seed.get());
+    std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+    const uint8_t                          constant_border_value = distribution_u8(generator);
+
+    // Create the mask
+    uint8_t mask[mask_size * mask_size];
+    fill_mask_from_pattern(mask, mask_size, mask_size, pattern);
+
+    // Compute function
+    Tensor dst = compute_non_linear_filter(shape, function, mask_size, pattern, mask, border_mode, constant_border_value);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_non_linear_filter(shape, function, mask_size, pattern, mask, border_mode, constant_border_value);
+
+    // Calculate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, BorderSize(static_cast<int>(mask_size / 2)));
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, valid_region);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/PixelWiseMultiplication.cpp b/tests/validation_old/NEON/PixelWiseMultiplication.cpp
new file mode 100644
index 0000000000..60eb82ef84
--- /dev/null
+++ b/tests/validation_old/NEON/PixelWiseMultiplication.cpp
@@ -0,0 +1,583 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEPixelWiseMultiplication.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute Neon pixel-wise multiplication function.
+ *
+ * @param[in] shape                Shape of the input and output tensors.
+ * @param[in] dt_in0               Data type of first input tensor.
+ * @param[in] dt_in1               Data type of second input tensor.
+ * @param[in] dt_out               Data type of the output tensor.
+ * @param[in] scale                Non-negative scale.
+ * @param[in] convert_policy       Overflow policy of the operation.
+ * @param[in] rounding_policy      Rounding policy of the operation.
+ * @param[in] fixed_point_position (Optional) Fixed point position that expresses the number of bits for the fractional part of the number.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_pixel_wise_multiplication(const TensorShape &shape, DataType dt_in0, DataType dt_in1, DataType dt_out, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy,
+                                         int fixed_point_position = 0)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, dt_in0, 1, fixed_point_position);
+    Tensor src2 = create_tensor<Tensor>(shape, dt_in1, 1, fixed_point_position);
+    Tensor dst  = create_tensor<Tensor>(shape, dt_out, 1, fixed_point_position);
+
+    // Create and configure function
+    NEPixelWiseMultiplication multiply;
+    multiply.configure(&src1, &src2, &dst, scale, convert_policy, rounding_policy);
+
+    // Allocate tensors
+    src1.allocator()->allocate();
+    src2.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src1.info()->is_resizable());
+    BOOST_TEST(!src2.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src1), 0);
+    library->fill_tensor_uniform(Accessor(src2), 1);
+
+    // Compute function
+    multiply.run();
+
+    return dst;
+}
+
+void validate_configuration(const Tensor &src1, const Tensor &src2, Tensor &dst, TensorShape shape, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy)
+{
+    BOOST_TEST(src1.info()->is_resizable());
+    BOOST_TEST(src2.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create and configure function
+    NEPixelWiseMultiplication multiply;
+    multiply.configure(&src1, &src2, &dst, scale, convert_policy, rounding_policy);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+    validate(src1.info()->valid_region(), valid_region);
+    validate(src2.info()->valid_region(), valid_region);
+    validate(dst.info()->valid_region(), valid_region);
+
+    // Validate padding
+    const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+    validate(src1.info()->padding(), padding);
+    validate(src2.info()->padding(), padding);
+    validate(dst.info()->padding(), padding);
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(PixelWiseMultiplication)
+
+BOOST_AUTO_TEST_SUITE(U8)
+BOOST_AUTO_TEST_SUITE(Scale255)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * (1.f / 255.f) * ConvertPolicies()
+                     * RoundingPolicy::TO_NEAREST_UP,
+                     shape, scale, convert_policy, rounding_policy)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor src2 = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::U8);
+
+    validate_configuration(src1, src2, dst, shape, scale, convert_policy, rounding_policy);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * (1.f / 255.f) * ConvertPolicies() * RoundingPolicy::TO_NEAREST_UP,
+                     shape, scale, convert_policy, rounding_policy)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, DataType::U8, DataType::U8, DataType::U8, scale, convert_policy,
+                                                   rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::U8, DataType::U8,
+                                                                               DataType::U8, scale, convert_policy, rounding_policy);
+
+    // Validate output
+    // Allow tolerance value of 1.f to counteract imprecision due to 32-bit float conversion
+    validate(Accessor(dst), ref_dst, 1.f, 0.f, std::numeric_limits<uint8_t>::max());
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * (1.f / 255.f) * ConvertPolicies() * RoundingPolicy::TO_NEAREST_UP,
+                     shape, scale, convert_policy, rounding_policy)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, DataType::U8, DataType::U8, DataType::U8, scale, convert_policy,
+                                                   rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::U8, DataType::U8,
+                                                                               DataType::U8, scale, convert_policy, rounding_policy);
+
+    // Validate output
+    // Allow tolerance value of 1.f to counteract imprecision due to 32-bit float conversion
+    validate(Accessor(dst), ref_dst, 1.f, 0.f, std::numeric_limits<uint8_t>::max());
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(ScaleOther)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ 1.f, 1.f / 32768.f })
+                     * ConvertPolicies()
+                     * RoundingPolicy::TO_ZERO,
+                     shape, scale, convert_policy, rounding_policy)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor src2 = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::U8);
+
+    validate_configuration(src1, src2, dst, shape, scale, convert_policy, rounding_policy);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) * ConvertPolicies()
+                     * RoundingPolicy::TO_ZERO,
+                     shape, scale, convert_policy, rounding_policy)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, DataType::U8, DataType::U8, DataType::U8, scale, convert_policy,
+                                                   rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::U8, DataType::U8,
+                                                                               DataType::U8, scale, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) * ConvertPolicies()
+                     * RoundingPolicy::TO_ZERO,
+                     shape, scale, convert_policy, rounding_policy)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, DataType::U8, DataType::U8, DataType::U8, scale, convert_policy,
+                                                   rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::U8, DataType::U8,
+                                                                               DataType::U8, scale, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(S16)
+BOOST_AUTO_TEST_SUITE(Scale255)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * (1.f / 255.f) * ConvertPolicies()
+                     * RoundingPolicy::TO_NEAREST_UP,
+                     shape, dt, scale, convert_policy, rounding_policy)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, dt);
+    Tensor src2 = create_tensor<Tensor>(shape, DataType::S16);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::S16);
+
+    validate_configuration(src1, src2, dst, shape, scale, convert_policy, rounding_policy);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * (1.f / 255.f) * ConvertPolicies()
+                     * RoundingPolicy::TO_NEAREST_UP,
+                     shape, dt, scale, convert_policy, rounding_policy)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16, scale, convert_policy, rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16, scale, convert_policy, rounding_policy);
+
+    // Validate output
+    // Allow tolerance value of 2.f to counteract imprecision due to 32-bit float conversion
+    validate(Accessor(dst), ref_dst, 2.f, 0.f, std::numeric_limits<int16_t>::max());
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * (1.f / 255.f) * ConvertPolicies()
+                     * RoundingPolicy::TO_NEAREST_UP,
+                     shape, dt, scale, convert_policy, rounding_policy)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16, scale, convert_policy, rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16,
+                                                                               scale, convert_policy, rounding_policy);
+
+    // Validate output
+    // Allow tolerance value of 2.f to counteract imprecision due to 32-bit float conversion
+    validate(Accessor(dst), ref_dst, 2.f, 0.f, std::numeric_limits<int16_t>::max());
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(ScaleOther)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ 1.f, 1.f / 32768.f })
+                     * ConvertPolicies()
+                     * RoundingPolicy::TO_ZERO,
+                     shape, dt, scale, convert_policy, rounding_policy)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, dt);
+    Tensor src2 = create_tensor<Tensor>(shape, DataType::S16);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::S16);
+
+    validate_configuration(src1, src2, dst, shape, scale, convert_policy, rounding_policy);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) * ConvertPolicies()
+                     * RoundingPolicy::TO_ZERO,
+                     shape, dt, scale, convert_policy, rounding_policy)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16, scale, convert_policy, rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16, scale, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) * ConvertPolicies()
+                     * RoundingPolicy::TO_ZERO,
+                     shape, dt, scale, convert_policy, rounding_policy)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16, scale, convert_policy, rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16,
+                                                                               scale, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+
+#ifdef ARM_COMPUTE_ENABLE_FP16
+BOOST_AUTO_TEST_SUITE(F16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * (1.f / 255.f) * ConvertPolicies() * RoundingPolicy::TO_NEAREST_UP,
+                     shape, scale, convert_policy, rounding_policy)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, DataType::F16, DataType::F16, DataType::F16, scale, convert_policy, rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::F16, DataType::F16, DataType::F16, scale, convert_policy, rounding_policy);
+
+    // Validate output
+    // Allow tolerance value of 1.f to counteract imprecision due to 32-bit float conversion
+    validate(Accessor(dst), ref_dst, 1.f, 0.f, std::numeric_limits<int16_t>::max());
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* ARM_COMPUTE_ENABLE_FP16 */
+
+BOOST_AUTO_TEST_SUITE(F32)
+BOOST_AUTO_TEST_SUITE(Scale255)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * (1.f / 255.f) * ConvertPolicies()
+                     * RoundingPolicy::TO_NEAREST_UP,
+                     shape, scale, convert_policy, rounding_policy)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, DataType::F32);
+    Tensor src2 = create_tensor<Tensor>(shape, DataType::F32);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::F32);
+
+    validate_configuration(src1, src2, dst, shape, scale, convert_policy, rounding_policy);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * (1.f / 255.f) * ConvertPolicies()
+                     * RoundingPolicy::TO_NEAREST_UP,
+                     shape, scale, convert_policy, rounding_policy)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32, scale, convert_policy, rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32, scale, convert_policy, rounding_policy);
+
+    // Validate output
+    // Allow tolerance value of 1.f to counteract imprecision due to 32-bit float conversion
+    validate(Accessor(dst), ref_dst, 1.f, 0.f, std::numeric_limits<int16_t>::max());
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * (1.f / 255.f) * ConvertPolicies()
+                     * RoundingPolicy::TO_NEAREST_UP,
+                     shape, scale, convert_policy, rounding_policy)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32, scale, convert_policy, rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32,
+                                                                               scale, convert_policy, rounding_policy);
+
+    // Validate output
+    // Allow tolerance value of 1.f to counteract imprecision due to 32-bit float conversion
+    validate(Accessor(dst), ref_dst, 1.f, 0.f, std::numeric_limits<int16_t>::max());
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(ScaleOther)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ 1.f, 1.f / 32768.f })
+                     * ConvertPolicies()
+                     * RoundingPolicy::TO_ZERO,
+                     shape, scale, convert_policy, rounding_policy)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, DataType::F32);
+    Tensor src2 = create_tensor<Tensor>(shape, DataType::F32);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::F32);
+
+    validate_configuration(src1, src2, dst, shape, scale, convert_policy, rounding_policy);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) * ConvertPolicies()
+                     * RoundingPolicy::TO_ZERO,
+                     shape, scale, convert_policy, rounding_policy)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32, scale, convert_policy, rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32, scale, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) * ConvertPolicies()
+                     * RoundingPolicy::TO_ZERO,
+                     shape, scale, convert_policy, rounding_policy)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32, scale, convert_policy, rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32,
+                                                                               scale, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(Quantized)
+BOOST_AUTO_TEST_SUITE(QS8)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * DataType::QS8 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange<int>(1, 7),
+                     shape, dt, convert_policy, rounding_policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, convert_policy, rounding_policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, fixed_point_position, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmallScale255, SmallShapes() * DataType::QS8 * (1.f / 255.f) * ConvertPolicies() * RoundingPolicy::TO_NEAREST_UP * boost::unit_test::data::xrange(1, 7),
+                     shape, dt, scale, convert_policy, rounding_policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, scale, fixed_point_position, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmallScaleOther, SmallShapes() * DataType::QS8 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange(1, 7),
+                     shape, dt, convert_policy, rounding_policy, fixed_point_position)
+{
+    const float scale = 1.f / static_cast<float>(1 << fixed_point_position);
+
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, scale, fixed_point_position, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, 1.f);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * DataType::QS8 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange<int>(1, 7),
+                     shape, dt, convert_policy, rounding_policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, convert_policy, rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLargeScale255, LargeShapes() * DataType::QS8 * (1.f / 255.f) * ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange(1, 7),
+                     shape, dt, scale, convert_policy, rounding_policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, scale, fixed_point_position, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLargeScaleOther, LargeShapes() * DataType::QS8 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange(1, 7),
+                     shape, dt, convert_policy, rounding_policy, fixed_point_position)
+{
+    const float scale = 1.f / static_cast<float>(1 << fixed_point_position);
+
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, scale, fixed_point_position, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, 1.f);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE(QS16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * DataType::QS16 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange<int>(1, 15),
+                     shape, dt, convert_policy, rounding_policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, convert_policy, rounding_policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, fixed_point_position, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmallScale255, SmallShapes() * DataType::QS16 * (1.f / 255.f) * ConvertPolicies() * RoundingPolicy::TO_NEAREST_UP * boost::unit_test::data::xrange(1, 15),
+                     shape, dt, scale, convert_policy, rounding_policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, scale, fixed_point_position, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmallScaleOther, SmallShapes() * DataType::QS16 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange(1, 15),
+                     shape, dt, convert_policy, rounding_policy, fixed_point_position)
+{
+    const float scale = 1.f / static_cast<float>(1 << fixed_point_position);
+
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy, fixed_point_position);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, scale, fixed_point_position, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, 1.f);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * DataType::QS16 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange<int>(1, 15),
+                     shape, dt, convert_policy, rounding_policy, fixed_point_position)
+{
+    // Compute function
+    Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, convert_policy, rounding_policy);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, convert_policy, rounding_policy);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/ROIPoolingLayer.cpp b/tests/validation_old/NEON/ROIPoolingLayer.cpp
new file mode 100644
index 0000000000..2046beb196
--- /dev/null
+++ b/tests/validation_old/NEON/ROIPoolingLayer.cpp
@@ -0,0 +1,110 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "NEON/ArrayAccessor.h"
+#include "TypePrinter.h"
+#include "arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h"
+#include "tests/Globals.h"
+#include "tests/Utils.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+#include "tests/validation_old/ValidationUserConfiguration.h"
+
+#include <random>
+#include <vector>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+Tensor compute_roi_pooling_layer(const TensorShape &shape, DataType dt, const std::vector<ROI> &rois, ROIPoolingLayerInfo pool_info)
+{
+    TensorShape shape_dst;
+    shape_dst.set(0, pool_info.pooled_width());
+    shape_dst.set(1, pool_info.pooled_height());
+    shape_dst.set(2, shape.z());
+    shape_dst.set(3, rois.size());
+
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, dt);
+    Tensor dst = create_tensor<Tensor>(shape_dst, dt);
+
+    // Create ROI array
+    Array<ROI> rois_array(rois.size());
+    fill_array(ArrayAccessor<ROI>(rois_array), rois);
+
+    // Create and configure function
+    NEROIPoolingLayer roi_pool;
+    roi_pool.configure(&src, &rois_array, &dst, pool_info);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    std::uniform_real_distribution<> distribution(-1, 1);
+    library->fill(Accessor(src), distribution, 0);
+
+    // Compute function
+    roi_pool.run();
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(ROIPoolingLayer)
+
+BOOST_AUTO_TEST_SUITE(Float)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, CNNFloatDataTypes() * boost::unit_test::data::make({ 10, 20, 40 }) * boost::unit_test::data::make({ 7, 9 }) * boost::unit_test::data::make({ 1.f / 8.f, 1.f / 16.f }),
+                     dt, num_rois, roi_pool_size, roi_scale)
+{
+    TensorShape         shape(50U, 47U, 2U, 3U);
+    ROIPoolingLayerInfo pool_info(roi_pool_size, roi_pool_size, roi_scale);
+
+    // Construct ROI vector
+    std::vector<ROI> rois = generate_random_rois(shape, pool_info, num_rois, user_config.seed);
+
+    // Compute function
+    Tensor dst = compute_roi_pooling_layer(shape, dt, rois, pool_info);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_roi_pooling_layer(shape, dt, rois, pool_info);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Sobel3x3.cpp b/tests/validation_old/NEON/Sobel3x3.cpp
new file mode 100644
index 0000000000..cb249e1a58
--- /dev/null
+++ b/tests/validation_old/NEON/Sobel3x3.cpp
@@ -0,0 +1,203 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+#include "tests/validation_old/ValidationUserConfiguration.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NESobel3x3.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "PaddingCalculator.h"
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+constexpr unsigned int filter_size = 3;              /** Size of the kernel/filter in number of elements. */
+constexpr BorderSize   border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */
+
+/** Compute Neon Sobel 3x3 function.
+ *
+ * @param[in] shape                 Shape of the input and output tensors.
+ * @param[in] border_mode           BorderMode used by the input tensor
+ * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT
+ *
+ * @return Computed output tensor.
+ */
+std::pair<Tensor, Tensor> compute_sobel_3x3(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value)
+{
+    // Create tensors
+    Tensor src   = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst_x = create_tensor<Tensor>(shape, DataType::S16);
+    Tensor dst_y = create_tensor<Tensor>(shape, DataType::S16);
+
+    src.info()->set_format(Format::U8);
+    dst_x.info()->set_format(Format::S16);
+    dst_y.info()->set_format(Format::S16);
+
+    // Create sobel image configure function
+    NESobel3x3 sobel_3x3;
+    sobel_3x3.configure(&src, &dst_x, &dst_y, border_mode, constant_border_value);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst_x.allocator()->allocate();
+    dst_y.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst_x.info()->is_resizable());
+    BOOST_TEST(!dst_y.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src), 0);
+
+    // Compute function
+    sobel_3x3.run();
+
+    return std::make_pair(std::move(dst_x), std::move(dst_y));
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(Sobel3x3)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode)
+{
+    // Create tensors
+    Tensor src   = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst_x = create_tensor<Tensor>(shape, DataType::S16);
+    Tensor dst_y = create_tensor<Tensor>(shape, DataType::S16);
+
+    src.info()->set_format(Format::U8);
+    dst_x.info()->set_format(Format::S16);
+    dst_y.info()->set_format(Format::S16);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst_x.info()->is_resizable());
+    BOOST_TEST(dst_y.info()->is_resizable());
+
+    // Create sobel 3x3 configure function
+    NESobel3x3 sobel_3x3;
+    sobel_3x3.configure(&src, &dst_x, &dst_y, border_mode);
+
+    // Validate valid region
+    const ValidRegion src_valid_region = shape_to_valid_region(shape);
+    const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
+
+    validate(src.info()->valid_region(), src_valid_region);
+    validate(dst_x.info()->valid_region(), dst_valid_region);
+    validate(dst_y.info()->valid_region(), dst_valid_region);
+
+    // Validate padding
+    PaddingCalculator calculator(shape.x(), 8);
+
+    calculator.set_border_mode(border_mode);
+    calculator.set_border_size(1);
+
+    const PaddingSize dst_padding = calculator.required_padding();
+
+    calculator.set_accessed_elements(16);
+    calculator.set_access_offset(-1);
+
+    const PaddingSize src_padding = calculator.required_padding();
+
+    validate(src.info()->padding(), src_padding);
+    validate(dst_x.info()->padding(), dst_padding);
+    validate(dst_y.info()->padding(), dst_padding);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode)
+{
+    uint8_t constant_border_value = 0;
+
+    // Generate a random constant value if border_mode is constant
+    if(border_mode == BorderMode::CONSTANT)
+    {
+        std::mt19937                           gen(user_config.seed.get());
+        std::uniform_int_distribution<uint8_t> distribution(0, 255);
+        constant_border_value = distribution(gen);
+    }
+
+    // Compute function
+    std::pair<Tensor, Tensor> dst = compute_sobel_3x3(shape, border_mode, constant_border_value);
+
+    // Compute reference
+    std::pair<RawTensor, RawTensor> ref_dst = Reference::compute_reference_sobel_3x3(shape, border_mode, constant_border_value);
+
+    // Calculate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
+
+    // Validate output
+    validate(Accessor(dst.first), ref_dst.first, valid_region);
+    validate(Accessor(dst.second), ref_dst.second, valid_region);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode)
+{
+    uint8_t constant_border_value = 0;
+
+    // Generate a random constant value if border_mode is constant
+    if(border_mode == BorderMode::CONSTANT)
+    {
+        std::mt19937                           gen(user_config.seed.get());
+        std::uniform_int_distribution<uint8_t> distribution(0, 255);
+        constant_border_value = distribution(gen);
+    }
+
+    // Compute function
+    std::pair<Tensor, Tensor> dst = compute_sobel_3x3(shape, border_mode, constant_border_value);
+
+    // Compute reference
+    std::pair<RawTensor, RawTensor> ref_dst = Reference::compute_reference_sobel_3x3(shape, border_mode, constant_border_value);
+
+    // Calculate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
+
+    // Validate output
+    validate(Accessor(dst.first), ref_dst.first, valid_region);
+    validate(Accessor(dst.second), ref_dst.second, valid_region);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Sobel5x5.cpp b/tests/validation_old/NEON/Sobel5x5.cpp
new file mode 100644
index 0000000000..2f26e62e83
--- /dev/null
+++ b/tests/validation_old/NEON/Sobel5x5.cpp
@@ -0,0 +1,204 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+#include "tests/validation_old/ValidationUserConfiguration.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NESobel5x5.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+constexpr unsigned int filter_size = 5;              /** Size of the kernel/filter in number of elements. */
+constexpr BorderSize   border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */
+
+/** Compute Neon Sobel 5x5 function.
+ *
+ * @param[in] shape                 Shape of the input and output tensors.
+ * @param[in] border_mode           BorderMode used by the input tensor
+ * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT
+ *
+ * @return Computed output tensor.
+ */
+std::pair<Tensor, Tensor> compute_sobel_5x5(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value)
+{
+    // Create tensors
+    Tensor src   = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst_x = create_tensor<Tensor>(shape, DataType::S16);
+    Tensor dst_y = create_tensor<Tensor>(shape, DataType::S16);
+
+    src.info()->set_format(Format::U8);
+    dst_x.info()->set_format(Format::S16);
+    dst_y.info()->set_format(Format::S16);
+
+    // Create sobel image configure function
+    NESobel5x5 sobel_5x5;
+    sobel_5x5.configure(&src, &dst_x, &dst_y, border_mode, constant_border_value);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst_x.allocator()->allocate();
+    dst_y.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst_x.info()->is_resizable());
+    BOOST_TEST(!dst_y.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src), 0);
+
+    // Compute function
+    sobel_5x5.run();
+
+    return std::make_pair(std::move(dst_x), std::move(dst_y));
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(Sobel5x5)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode)
+{
+    // Create tensors
+    Tensor src   = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst_x = create_tensor<Tensor>(shape, DataType::S16);
+    Tensor dst_y = create_tensor<Tensor>(shape, DataType::S16);
+
+    src.info()->set_format(Format::U8);
+    dst_x.info()->set_format(Format::S16);
+    dst_y.info()->set_format(Format::S16);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst_x.info()->is_resizable());
+    BOOST_TEST(dst_y.info()->is_resizable());
+
+    // Create sobel 5x5 configure function
+    NESobel5x5 sobel_5x5;
+    sobel_5x5.configure(&src, &dst_x, &dst_y, border_mode);
+
+    // Validate valid region
+    const ValidRegion src_valid_region = shape_to_valid_region(shape);
+    const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
+
+    validate(src.info()->valid_region(), src_valid_region);
+    validate(dst_x.info()->valid_region(), dst_valid_region);
+    validate(dst_y.info()->valid_region(), dst_valid_region);
+
+    // Validate padding
+    PaddingCalculator calculator(shape.x(), 16);
+
+    calculator.set_border_mode(border_mode);
+    calculator.set_border_size(2);
+
+    const PaddingSize dst_padding = calculator.required_padding();
+
+    calculator.set_processed_elements(8);
+    calculator.set_access_offset(-2);
+
+    const PaddingSize src_padding = calculator.required_padding();
+
+    validate(src.info()->padding(), src_padding);
+    validate(dst_x.info()->padding(), dst_padding);
+    validate(dst_y.info()->padding(), dst_padding);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode)
+{
+    uint8_t constant_border_value = 0;
+
+    // Generate a random constant value if border_mode is constant
+    if(border_mode == BorderMode::CONSTANT)
+    {
+        std::mt19937                           gen(user_config.seed.get());
+        std::uniform_int_distribution<uint8_t> distribution(0, 255);
+        constant_border_value = distribution(gen);
+    }
+
+    // Compute function
+    std::pair<Tensor, Tensor> dst = compute_sobel_5x5(shape, border_mode, constant_border_value);
+
+    // Compute reference
+    std::pair<RawTensor, RawTensor> ref_dst = Reference::compute_reference_sobel_5x5(shape, border_mode, constant_border_value);
+
+    // Calculate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
+
+    // Validate output
+    validate(Accessor(dst.first), ref_dst.first, valid_region);
+    validate(Accessor(dst.second), ref_dst.second, valid_region);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode)
+{
+    uint8_t constant_border_value = 0;
+
+    // Generate a random constant value if border_mode is constant
+    if(border_mode == BorderMode::CONSTANT)
+    {
+        std::mt19937                           gen(user_config.seed.get());
+        std::uniform_int_distribution<uint8_t> distribution(0, 255);
+        constant_border_value = distribution(gen);
+    }
+
+    // Compute function
+    std::pair<Tensor, Tensor> dst = compute_sobel_5x5(shape, border_mode, constant_border_value);
+
+    // Compute reference
+    std::pair<RawTensor, RawTensor> ref_dst = Reference::compute_reference_sobel_5x5(shape, border_mode, constant_border_value);
+
+    // Calculate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size);
+
+    // Validate output
+    validate(Accessor(dst.first), ref_dst.first, valid_region);
+    validate(Accessor(dst.second), ref_dst.second, valid_region);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/TableLookup.cpp b/tests/validation_old/NEON/TableLookup.cpp
new file mode 100644
index 0000000000..f134e5d417
--- /dev/null
+++ b/tests/validation_old/NEON/TableLookup.cpp
@@ -0,0 +1,229 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "NEON/Helper.h"
+#include "NEON/LutAccessor.h"
+#include "PaddingCalculator.h"
+#include "RawLutAccessor.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Helpers.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NETableLookup.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute Table Lookup function.
+ *
+ * @param[in] shape     Shape of the input tensors
+ * @param[in] data_type Datatype of the input/output tensors
+ * @param[in] lut       The input LUT.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_table_lookup(const TensorShape &shape, DataType data_type, Lut &lut)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, data_type);
+    Tensor dst = create_tensor<Tensor>(shape, data_type);
+
+    // Create and configure function
+    NETableLookup table_lookup;
+    table_lookup.configure(&src, &lut, &dst);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src), 0);
+
+    // Compute function
+    table_lookup.run();
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(TableLookup)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }),
+                     shape, data_type)
+{
+    //Create Lut
+    const int num_elem = (data_type == DataType::U8) ? std::numeric_limits<uint8_t>::max() + 1 : std::numeric_limits<int16_t>::max() - std::numeric_limits<int16_t>::lowest() + 1;
+    Lut       lut(num_elem, data_type);
+
+    if(data_type == DataType::U8)
+    {
+        fill_lookuptable(LutAccessor<uint8_t>(lut));
+    }
+    else
+    {
+        fill_lookuptable(LutAccessor<int16_t>(lut));
+    }
+
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, data_type);
+    Tensor dst = create_tensor<Tensor>(shape, data_type);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create and configure function
+    NETableLookup table_lookup;
+    table_lookup.configure(&src, &lut, &dst);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+    validate(src.info()->valid_region(), valid_region);
+    validate(dst.info()->valid_region(), valid_region);
+
+    // Validate padding
+    const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+    validate(src.info()->padding(), padding);
+    validate(dst.info()->padding(), padding);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall,
+                     SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }),
+                     shape, data_type)
+{
+    //Create Lut
+    const int num_elem = (data_type == DataType::U8) ? std::numeric_limits<uint8_t>::max() + 1 : std::numeric_limits<int16_t>::max() - std::numeric_limits<int16_t>::lowest() + 1;
+    Lut       lut(num_elem, data_type);
+
+    if(data_type == DataType::U8)
+    {
+        //Create rawLut
+        std::map<uint8_t, uint8_t> rawlut;
+
+        //Fill the Lut
+        fill_lookuptable(LutAccessor<uint8_t>(lut));
+        fill_lookuptable(RawLutAccessor<uint8_t>(rawlut));
+
+        // Compute function
+        Tensor dst = compute_table_lookup(shape, data_type, lut);
+
+        // Compute reference
+        RawTensor ref_dst = Reference::compute_reference_table_lookup(shape, data_type, rawlut);
+
+        // Validate output
+        validate(Accessor(dst), ref_dst);
+    }
+    else
+    {
+        //Create rawLut
+        std::map<int16_t, int16_t> rawlut;
+
+        //Fill the Lut
+        fill_lookuptable(LutAccessor<int16_t>(lut));
+        fill_lookuptable(RawLutAccessor<int16_t>(rawlut));
+
+        // Compute function
+        Tensor dst = compute_table_lookup(shape, data_type, lut);
+
+        // Compute reference
+        RawTensor ref_dst = Reference::compute_reference_table_lookup(shape, data_type, rawlut);
+
+        // Validate output
+        validate(Accessor(dst), ref_dst);
+    }
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge,
+                     LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }),
+                     shape, data_type)
+{
+    //Create Lut
+    const int num_elem = (data_type == DataType::U8) ? std::numeric_limits<uint8_t>::max() + 1 : std::numeric_limits<int16_t>::max() - std::numeric_limits<int16_t>::lowest() + 1;
+    Lut       lut(num_elem, data_type);
+
+    if(data_type == DataType::U8)
+    {
+        //Create rawLut
+        std::map<uint8_t, uint8_t> rawlut;
+
+        //Fill the Lut
+        fill_lookuptable(LutAccessor<uint8_t>(lut));
+        fill_lookuptable(RawLutAccessor<uint8_t>(rawlut));
+
+        // Compute function
+        Tensor dst = compute_table_lookup(shape, data_type, lut);
+
+        // Compute reference
+        RawTensor ref_dst = Reference::compute_reference_table_lookup(shape, data_type, rawlut);
+
+        // Validate output
+        validate(Accessor(dst), ref_dst);
+    }
+    else
+    {
+        //Create rawLut
+        std::map<int16_t, int16_t> rawlut;
+
+        //Fill the Lut
+        fill_lookuptable(LutAccessor<int16_t>(lut));
+        fill_lookuptable(RawLutAccessor<int16_t>(rawlut));
+
+        // Compute function
+        Tensor dst = compute_table_lookup(shape, data_type, lut);
+
+        // Compute reference
+        RawTensor ref_dst = Reference::compute_reference_table_lookup(shape, data_type, rawlut);
+
+        // Validate output
+        validate(Accessor(dst), ref_dst);
+    }
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/Threshold.cpp b/tests/validation_old/NEON/Threshold.cpp
new file mode 100644
index 0000000000..d56ec5eb42
--- /dev/null
+++ b/tests/validation_old/NEON/Threshold.cpp
@@ -0,0 +1,153 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+#include "tests/validation_old/dataset/ThresholdDataset.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEThreshold.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+
+#include "tests/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute Threshold function.
+ *
+ * @param[in] shape       Shape of the input and output tensors.
+ * @param[in] threshold   Threshold. When the threshold type is RANGE, this is used as the lower threshold.
+ * @param[in] false_value value to set when the condition is not respected.
+ * @param[in] true_value  value to set when the condition is respected.
+ * @param[in] type        Thresholding type. Either RANGE or BINARY.
+ * @param[in] upper       Upper threshold. Only used when the thresholding type is RANGE.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_threshold(const TensorShape &shape, uint8_t threshold, uint8_t false_value, uint8_t true_value, ThresholdType type, uint8_t upper)
+{
+    // Create tensors
+    Tensor src1 = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst  = create_tensor<Tensor>(shape, DataType::U8);
+
+    // Create and configure function
+    NEThreshold thrsh;
+    thrsh.configure(&src1, &dst, threshold, false_value, true_value, type, upper);
+
+    // Allocate tensors
+    src1.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src1.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src1), 0);
+
+    // Compute function
+    thrsh.run();
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(Threshold)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration,
+                     (SmallShapes() + LargeShapes()) * ThresholdDataset(),
+                     shape, thrshConf)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U8);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // Create and configure function
+    NEThreshold thrsh;
+    thrsh.configure(&src, &dst, thrshConf.threshold, thrshConf.false_value, thrshConf.true_value, thrshConf.type, thrshConf.upper);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+    validate(src.info()->valid_region(), valid_region);
+    validate(dst.info()->valid_region(), valid_region);
+
+    // Validate padding
+    const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding();
+    validate(src.info()->padding(), padding);
+    validate(dst.info()->padding(), padding);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall,
+                     SmallShapes() * ThresholdDataset(),
+                     shape, thrshConf)
+{
+    // Compute function
+    Tensor dst = compute_threshold(shape, thrshConf.threshold, thrshConf.false_value, thrshConf.true_value, thrshConf.type, thrshConf.upper);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_threshold(shape, thrshConf.threshold, thrshConf.false_value, thrshConf.true_value, thrshConf.type, thrshConf.upper);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge,
+                     LargeShapes() * ThresholdDataset(),
+                     shape, thrshConf)
+{
+    // Compute function
+    Tensor dst = compute_threshold(shape, thrshConf.threshold, thrshConf.false_value, thrshConf.true_value, thrshConf.type, thrshConf.upper);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_threshold(shape, thrshConf.threshold, thrshConf.false_value, thrshConf.true_value, thrshConf.type, thrshConf.upper);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation_old/NEON/WarpPerspective.cpp b/tests/validation_old/NEON/WarpPerspective.cpp
new file mode 100644
index 0000000000..5a15591261
--- /dev/null
+++ b/tests/validation_old/NEON/WarpPerspective.cpp
@@ -0,0 +1,209 @@
+/*
+ * 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 "NEON/Accessor.h"
+#include "PaddingCalculator.h"
+#include "TypePrinter.h"
+#include "Utils.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/validation_old/Datasets.h"
+#include "tests/validation_old/Helpers.h"
+#include "tests/validation_old/Reference.h"
+#include "tests/validation_old/Validation.h"
+#include "tests/validation_old/ValidationUserConfiguration.h"
+
+#include "arm_compute/core/Helpers.h"
+#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/validation_old/boost_wrapper.h"
+
+#include <random>
+#include <string>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+/** Compute Warp Perspective function.
+ *
+     * @param[in] input                 Shape of the input and output tensors.
+     * @param[in] matrix                The perspective matrix. Must be 3x3 of type float.
+     * @param[in] policy                The interpolation type.
+     * @param[in] border_mode           Strategy to use for borders.
+     * @param[in] constant_border_value Constant value to use for borders if border_mode is set to CONSTANT.
+ *
+ * @return Computed output tensor.
+ */
+Tensor compute_warp_perspective(const TensorShape &shape, const float *matrix, InterpolationPolicy policy,
+                                BorderMode border_mode, uint8_t constant_border_value)
+{
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U8);
+
+    // Create and configure function
+    NEWarpPerspective warp_perspective;
+    warp_perspective.configure(&src, &dst, matrix, policy, border_mode, constant_border_value);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    library->fill_tensor_uniform(Accessor(src), 0);
+
+    // Compute function
+    warp_perspective.run();
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(NEON)
+BOOST_AUTO_TEST_SUITE(WarpPerspective)
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes())
+                     * boost::unit_test::data::make({ InterpolationPolicy::BILINEAR, InterpolationPolicy::NEAREST_NEIGHBOR }) * BorderModes(),
+                     shape, policy, border_mode)
+{
+    uint8_t constant_border_value = 0;
+
+    // Generate a random constant value if border_mode is constant
+    if(border_mode == BorderMode::CONSTANT)
+    {
+        std::mt19937                           gen(user_config.seed.get());
+        std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+        constant_border_value = distribution_u8(gen);
+    }
+
+    // Create the matrix
+    std::array<float, 9> matrix;
+    fill_warp_matrix<9>(matrix, 3, 3);
+
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, DataType::U8);
+    Tensor dst = create_tensor<Tensor>(shape, DataType::U8);
+
+    BOOST_TEST(src.info()->is_resizable());
+    BOOST_TEST(dst.info()->is_resizable());
+
+    // 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);
+}
+
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes()
+                     * boost::unit_test::data::make({ InterpolationPolicy::BILINEAR, InterpolationPolicy::NEAREST_NEIGHBOR })
+                     * BorderModes(),
+                     shape, policy, border_mode)
+{
+    uint8_t constant_border_value = 0;
+
+    // Generate a random constant value if border_mode is constant
+    if(border_mode == BorderMode::CONSTANT)
+    {
+        std::mt19937                           gen(user_config.seed.get());
+        std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+        constant_border_value = distribution_u8(gen);
+    }
+
+    // Create the valid mask Tensor
+    RawTensor valid_mask(shape, DataType::U8);
+
+    // Create the matrix
+    std::array<float, 9> matrix;
+    fill_warp_matrix<9>(matrix, 3, 3);
+
+    // Compute function
+    Tensor dst = compute_warp_perspective(shape, matrix.data(), policy, border_mode, constant_border_value);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_warp_perspective(shape, valid_mask, matrix.data(), policy, border_mode, constant_border_value);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, valid_mask, 1, 0.2f);
+}
+BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly"))
+BOOST_DATA_TEST_CASE(RunLarge, LargeShapes()
+                     * boost::unit_test::data::make({ InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR }) * BorderModes(),
+                     shape, policy, border_mode)
+{
+    uint8_t constant_border_value = 0;
+
+    // Generate a random constant value if border_mode is constant
+    if(border_mode == BorderMode::CONSTANT)
+    {
+        std::mt19937                           gen(user_config.seed.get());
+        std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+        constant_border_value = distribution_u8(gen);
+    }
+
+    // Create the valid mask Tensor
+    RawTensor valid_mask(shape, DataType::U8);
+
+    // Create the matrix
+    std::array<float, 9> matrix;
+    fill_warp_matrix<9>(matrix, 3, 3);
+
+    // Compute function
+    Tensor dst = compute_warp_perspective(shape, matrix.data(), policy, border_mode, constant_border_value);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_warp_perspective(shape, valid_mask, matrix.data(), policy, border_mode, constant_border_value);
+
+    // Validate output
+    validate(Accessor(dst), ref_dst, valid_mask, 1, 0.2f);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */