COMPMID-578: Implement FAST corners for CL/NEON

Change-Id: Ifa74e2bf05546de9a49aa185e22fba50438d8ad6
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/113946
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Pablo Tello <pablo.tello@arm.com>

12 files changed, 847 insertions, 12 deletions

diff --git a/arm_compute/runtime/CL/functions/CLFastCorners.h b/arm_compute/runtime/CL/functions/CLFastCorners.h
index 9afec71bc3..404b561e48 100644
--- a/arm_compute/runtime/CL/functions/CLFastCorners.h
+++ b/arm_compute/runtime/CL/functions/CLFastCorners.h
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2016, 2017 ARM Limited.
+ * Copyright (c) 2016-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -69,7 +69,7 @@ public:
      * @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.
      */
-    void configure(const ICLImage *input, float threshold, bool nonmax_suppression, CLKeyPointArray *corners, unsigned int *num_corners,
+    void configure(const ICLImage *input, float threshold, bool nonmax_suppression, ICLKeyPointArray *corners, unsigned int *num_corners,
                    BorderMode border_mode, uint8_t constant_border_value = 0);
     // Inherited methods overridden:
     void run() override;
@@ -85,7 +85,7 @@ private:
     bool                      _non_max;
     unsigned int             *_num_corners;
     cl::Buffer                _num_buffer;
-    CLKeyPointArray          *_corners;
+    ICLKeyPointArray         *_corners;
     uint8_t                   _constant_border_value;
 };
 }
diff --git a/src/core/CL/cl_kernels/fast_corners.cl b/src/core/CL/cl_kernels/fast_corners.cl
index 3e1929c637..76b35b931b 100644
--- a/src/core/CL/cl_kernels/fast_corners.cl
+++ b/src/core/CL/cl_kernels/fast_corners.cl
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2016, 2017 ARM Limited.
+ * Copyright (c) 2016-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -254,6 +254,9 @@ __kernel void copy_to_keypoint(
             out[id].x               = get_global_id(0) + offset;
             out[id].y               = get_global_id(1) + offset;
             out[id].tracking_status = 1;
+            out[id].scale           = 0.f;
+            out[id].orientation     = 0.f;
+            out[id].error           = 0.f;
         }
     }
 }
diff --git a/src/core/CL/kernels/CLFastCornersKernel.cpp b/src/core/CL/kernels/CLFastCornersKernel.cpp
index 1d4d776730..616e41b5fc 100644
--- a/src/core/CL/kernels/CLFastCornersKernel.cpp
+++ b/src/core/CL/kernels/CLFastCornersKernel.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2016, 2017 ARM Limited.
+ * Copyright (c) 2016-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -138,7 +138,7 @@ void CLCopyToArrayKernel::configure(const ICLImage *input, bool update_number, I
 
     // Set static kernel arguments
     unsigned int idx = num_arguments_per_2D_tensor(); // Skip the input and output parameters
-    _kernel.setArg<unsigned int>(idx++, corners->max_num_values());
+    _kernel.setArg<unsigned int>(idx++, _corners->max_num_values());
     _kernel.setArg<cl_uint>(idx++, offset);
     _kernel.setArg(idx++, *_num_buffer);
     _kernel.setArg(idx++, _corners->cl_buffer());
diff --git a/src/core/NEON/kernels/NEFillArrayKernel.cpp b/src/core/NEON/kernels/NEFillArrayKernel.cpp
index 5a2e1a0aa4..c6c6c452b7 100644
--- a/src/core/NEON/kernels/NEFillArrayKernel.cpp
+++ b/src/core/NEON/kernels/NEFillArrayKernel.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2016, 2017 ARM Limited.
+ * Copyright (c) 2016-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -81,6 +81,9 @@ void NEFillArrayKernel::run(const Window &window, const ThreadInfo &info)
             p.y               = id.y();
             p.strength        = value;
             p.tracking_status = 1;
+            p.scale           = 0.f;
+            p.orientation     = 0.f;
+            p.error           = 0.f;
 
             if(!_output->push_back(p))
             {
diff --git a/src/runtime/CL/functions/CLFastCorners.cpp b/src/runtime/CL/functions/CLFastCorners.cpp
index 7a0dd09fbe..d6cda91cea 100644
--- a/src/runtime/CL/functions/CLFastCorners.cpp
+++ b/src/runtime/CL/functions/CLFastCorners.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2016, 2017 ARM Limited.
+ * Copyright (c) 2016-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -52,7 +52,7 @@ CLFastCorners::CLFastCorners(std::shared_ptr<IMemoryManager> memory_manager)
 {
 }
 
-void CLFastCorners::configure(const ICLImage *input, float threshold, bool nonmax_suppression, CLKeyPointArray *const corners,
+void CLFastCorners::configure(const ICLImage *input, float threshold, bool nonmax_suppression, ICLKeyPointArray *corners,
                               unsigned int *num_corners, BorderMode border_mode, uint8_t constant_border_value)
 {
     ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D(input);
@@ -76,7 +76,7 @@ void CLFastCorners::configure(const ICLImage *input, float threshold, bool nonma
 
     if(!_non_max)
     {
-        _copy_array_kernel.configure(&_output, update_number, corners, &_num_buffer);
+        _copy_array_kernel.configure(&_output, update_number, _corners, &_num_buffer);
     }
     else
     {
@@ -84,7 +84,7 @@ void CLFastCorners::configure(const ICLImage *input, float threshold, bool nonma
         _memory_group.manage(&_suppr);
 
         _suppr_func.configure(&_output, &_suppr, border_mode);
-        _copy_array_kernel.configure(&_suppr, update_number, corners, &_num_buffer);
+        _copy_array_kernel.configure(&_suppr, update_number, _corners, &_num_buffer);
 
         _suppr.allocator()->allocate();
     }
diff --git a/tests/validation/CL/FastCorners.cpp b/tests/validation/CL/FastCorners.cpp
new file mode 100644
index 0000000000..b5086ef75d
--- /dev/null
+++ b/tests/validation/CL/FastCorners.cpp
@@ -0,0 +1,119 @@
+/*
+ * Copyright (c) 2017-2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/functions/CLFastCorners.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/CL/CLArrayAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/FastValidation.h"
+#include "tests/validation/fixtures/FastCornersFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/* Radius of the Bresenham circle around the candidate point */
+const unsigned int bresenham_radius = 3;
+/* Allowed percentage of keypoints missing for target */
+const float allowed_missing = 10.f;
+/* Allowed percentage of keypoints mismatching between target and reference */
+const float allowed_mismatching = 10.f;
+/* Tolerance used to compare corner strengths */
+const AbsoluteTolerance<float> tolerance(0.5f);
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(FastCorners)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(combine(combine(concat(datasets::Small2DShapes(), datasets::Large2DShapes()),
+                                                                                   framework::dataset::make("Format", Format::U8)),
+                                                                           framework::dataset::make("SuppressNonMax", { false, true })),
+                                                                   framework::dataset::make("BorderMode", BorderMode::UNDEFINED)),
+               shape, format, suppress_nonmax, border_mode)
+{
+    std::mt19937                           gen(library->seed());
+    std::uniform_int_distribution<uint8_t> int_dist(0, 255);
+    std::uniform_real_distribution<float>  real_dist(0, 255);
+
+    const uint8_t constant_border_value = int_dist(gen);
+    const float   threshold             = real_dist(gen);
+
+    // Create tensors
+    CLTensor src = create_tensor<CLTensor>(shape, data_type_from_format(format));
+    src.info()->set_format(format);
+
+    ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+    CLKeyPointArray corners;
+    unsigned int    num_corners;
+
+    // Create and configure function
+    CLFastCorners fast_corners;
+    fast_corners.configure(&src, threshold, suppress_nonmax, &corners, &num_corners, border_mode, constant_border_value);
+
+    // Validate padding
+    PaddingCalculator calculator(shape.x(), 1); // elems_processed
+
+    calculator.set_border_size(bresenham_radius);
+    calculator.set_access_offset(-bresenham_radius);
+    calculator.set_accessed_elements(7); // elems_read
+
+    validate(src.info()->padding(), calculator.required_padding());
+}
+
+template <typename T>
+using CLFastCornersFixture = FastCornersValidationFixture<CLTensor, CLAccessor, CLKeyPointArray, CLFastCorners, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, CLFastCornersFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::Small2DShapes(), framework::dataset::make("Format", Format::U8)),
+                                                                                                                   framework::dataset::make("SuppressNonMax", { false, true })),
+                                                                                                           framework::dataset::make("BorderMode", BorderMode::UNDEFINED)))
+{
+    // Validate output
+    CLArrayAccessor<KeyPoint> array(_target);
+    fast_validate_keypoints(array.buffer(), array.buffer() + array.num_values(), _reference.begin(), _reference.end(), tolerance, allowed_missing, allowed_mismatching);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLFastCornersFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::Large2DShapes(), framework::dataset::make("Format", Format::U8)),
+                                                                                                                 framework::dataset::make("SuppressNonMax", { false, true })),
+                                                                                                         framework::dataset::make("BorderMode", BorderMode::UNDEFINED)))
+{
+    // Validate output
+    CLArrayAccessor<KeyPoint> array(_target);
+    fast_validate_keypoints(array.buffer(), array.buffer() + array.num_values(), _reference.begin(), _reference.end(), tolerance, allowed_missing, allowed_mismatching);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/FastValidation.h b/tests/validation/FastValidation.h
new file mode 100644
index 0000000000..10757cfdc2
--- /dev/null
+++ b/tests/validation/FastValidation.h
@@ -0,0 +1,161 @@
+/*
+ * Copyright (c) 2017-2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef __ARM_COMPUTE_TEST_FAST_VALIDATION_H__
+#define __ARM_COMPUTE_TEST_FAST_VALIDATION_H__
+
+#include "Validation.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+/** Check which keypoints from [first1, last1) are missing in [first2, last2) */
+template <typename T, typename U, typename V>
+std::pair<int64_t, int64_t> fast_compare_keypoints(T first1, T last1, U first2, U last2, V tolerance, bool check_mismatches = true)
+{
+    /* Keypoint (x,y) should have similar strength (within tolerance) and other properties in both reference and target */
+    const auto compare_props_eq = [&](const KeyPoint & lhs, const KeyPoint & rhs)
+    {
+        return compare<V>(lhs.strength, rhs.strength, tolerance)
+               && lhs.tracking_status == rhs.tracking_status
+               && lhs.scale == rhs.scale
+               && lhs.orientation == rhs.orientation
+               && lhs.error == rhs.error;
+    };
+
+    /* Used to sort KeyPoints by coordinates (x, y) */
+    const auto compare_coords_lt = [](const KeyPoint & lhs, const KeyPoint & rhs)
+    {
+        return std::tie(lhs.x, lhs.y) < std::tie(rhs.x, rhs.y);
+    };
+
+    std::sort(first1, last1, compare_coords_lt);
+    std::sort(first2, last2, compare_coords_lt);
+
+    if(check_mismatches)
+    {
+        std::cout << "ref count = " << std::distance(first1, last1) << " \ttarget count = " << std::distance(first2, last2) << std::endl;
+    }
+
+    int64_t num_missing    = 0;
+    int64_t num_mismatches = 0;
+    bool    rest_missing   = false;
+
+    while(first1 != last1)
+    {
+        if(first2 == last2)
+        {
+            // num_missing += std::distance(first1, last1);
+            rest_missing = true;
+            ARM_COMPUTE_TEST_INFO("All key points from (" << first1->x << "," << first1->y << ") onwards not found");
+            break;
+        }
+
+        if(compare_coords_lt(*first1, *first2))
+        {
+            ++num_missing;
+            ARM_COMPUTE_TEST_INFO("Key point not found");
+            ARM_COMPUTE_TEST_INFO("keypoint1 = " << *first1++);
+        }
+        else
+        {
+            if(!compare_coords_lt(*first2, *first1)) // Equal coordinates
+            {
+                if(check_mismatches && !compare_props_eq(*first1, *first2)) // Check other properties
+                {
+                    ++num_mismatches;
+                    ARM_COMPUTE_TEST_INFO("Mismatching keypoint");
+                    ARM_COMPUTE_TEST_INFO("keypoint1 [ref] = " << *first1);
+                    ARM_COMPUTE_TEST_INFO("keypoint2 [tgt] = " << *first2);
+                }
+                ++first1;
+            }
+            ++first2;
+        }
+    }
+
+    if(rest_missing)
+    {
+        while(first1 != last1)
+        {
+            ++num_missing;
+            ARM_COMPUTE_TEST_INFO("Key point not found");
+            ARM_COMPUTE_TEST_INFO("keypoint1 = " << *first1++);
+        }
+    }
+
+    return std::make_pair(num_missing, num_mismatches);
+}
+
+template <typename T, typename U, typename V>
+void fast_validate_keypoints(T target_first, T target_last, U reference_first, U reference_last, V tolerance,
+                             float allowed_missing_percentage, float allowed_mismatch_percentage)
+{
+    const int64_t num_elements_target    = std::distance(target_first, target_last);
+    const int64_t num_elements_reference = std::distance(reference_first, reference_last);
+
+    int64_t num_missing    = 0;
+    int64_t num_mismatches = 0;
+
+    if(num_elements_reference > 0)
+    {
+        std::tie(num_missing, num_mismatches) = fast_compare_keypoints(reference_first, reference_last, target_first, target_last, tolerance);
+
+        const float percent_missing    = static_cast<float>(num_missing) / num_elements_reference * 100.f;
+        const float percent_mismatches = static_cast<float>(num_mismatches) / num_elements_reference * 100.f;
+
+        ARM_COMPUTE_TEST_INFO(num_missing << " keypoints (" << std::fixed << std::setprecision(2) << percent_missing << "%) in ref are missing from target");
+        ARM_COMPUTE_EXPECT(percent_missing <= allowed_missing_percentage, framework::LogLevel::ERRORS);
+
+        ARM_COMPUTE_TEST_INFO(num_mismatches << " keypoints (" << std::fixed << std::setprecision(2) << percent_mismatches << "%) mismatched");
+        ARM_COMPUTE_EXPECT(percent_mismatches <= allowed_mismatch_percentage, framework::LogLevel::ERRORS);
+
+        std::cout << "Mismatched keypoints: " << num_mismatches << "/" << num_elements_reference << " = " << std::fixed << std::setprecision(2) << percent_mismatches
+                  << "% \tMax allowed: " << allowed_mismatch_percentage << "%" << std::endl;
+        std::cout << "Missing (not in tgt): " << num_missing << "/" << num_elements_reference << " = " << std::fixed << std::setprecision(2) << percent_missing
+                  << "% \tMax allowed: " << allowed_missing_percentage << "%" << std::endl;
+    }
+
+    if(num_elements_target > 0)
+    {
+        // Note: no need to check for mismatches a second time (last argument is 'false')
+        std::tie(num_missing, num_mismatches) = fast_compare_keypoints(target_first, target_last, reference_first, reference_last, tolerance, false);
+
+        const float percent_missing = static_cast<float>(num_missing) / num_elements_target * 100.f;
+
+        ARM_COMPUTE_TEST_INFO(num_missing << " keypoints (" << std::fixed << std::setprecision(2) << percent_missing << "%) in target are missing from ref");
+        ARM_COMPUTE_EXPECT(percent_missing <= allowed_missing_percentage, framework::LogLevel::ERRORS);
+
+        std::cout << "Missing (not in ref): " << num_missing << "/" << num_elements_target << " = " << std::fixed << std::setprecision(2) << percent_missing
+                  << "% \tMax allowed: " << allowed_missing_percentage << "%\n"
+                  << std::endl;
+    }
+}
+
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_FAST_VALIDATION_H__ */
diff --git a/tests/validation/NEON/FastCorners.cpp b/tests/validation/NEON/FastCorners.cpp
new file mode 100644
index 0000000000..4d42aa5349
--- /dev/null
+++ b/tests/validation/NEON/FastCorners.cpp
@@ -0,0 +1,118 @@
+/*
+ * Copyright (c) 2017-2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEFastCorners.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/NEON/Accessor.h"
+#include "tests/NEON/ArrayAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/FastValidation.h"
+#include "tests/validation/fixtures/FastCornersFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/* Radius of the Bresenham circle around the candidate point */
+const unsigned int bresenham_radius = 3;
+/* Allowed percentage of keypoints missing for target */
+const float allowed_missing = 10.f;
+/* Allowed percentage of keypoints mismatching between target and reference */
+const float allowed_mismatching = 10.f;
+/* Tolerance used to compare corner strengths */
+const AbsoluteTolerance<float> tolerance(0.5f);
+} // namespace
+
+TEST_SUITE(NEON)
+TEST_SUITE(FastCorners)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(combine(combine(concat(datasets::Small2DShapes(), datasets::Large2DShapes()),
+                                                                                   framework::dataset::make("Format", Format::U8)),
+                                                                           framework::dataset::make("SuppressNonMax", { false, true })),
+                                                                   framework::dataset::make("BorderMode", BorderMode::UNDEFINED)),
+               shape, format, suppress_nonmax, border_mode)
+{
+    std::mt19937                           gen(library->seed());
+    std::uniform_int_distribution<uint8_t> int_dist(0, 255);
+    std::uniform_real_distribution<float>  real_dist(0, 255);
+
+    const uint8_t constant_border_value = int_dist(gen);
+    const float   threshold             = real_dist(gen);
+
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, data_type_from_format(format));
+    src.info()->set_format(format);
+
+    ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+    KeyPointArray corners;
+
+    // Create and configure function
+    NEFastCorners fast_corners;
+    fast_corners.configure(&src, threshold, suppress_nonmax, &corners, border_mode, constant_border_value);
+
+    // Validate padding
+    PaddingCalculator calculator(shape.x(), 1); // elems_processed
+
+    calculator.set_border_size(bresenham_radius);
+    calculator.set_access_offset(-bresenham_radius);
+    calculator.set_accessed_elements(8); // elems_read
+
+    validate(src.info()->padding(), calculator.required_padding());
+}
+
+template <typename T>
+using NEFastCornersFixture = FastCornersValidationFixture<Tensor, Accessor, KeyPointArray, NEFastCorners, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, NEFastCornersFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::Small2DShapes(), framework::dataset::make("Format", Format::U8)),
+                                                                                                                   framework::dataset::make("SuppressNonMax", { false, true })),
+                                                                                                           framework::dataset::make("BorderMode", BorderMode::UNDEFINED)))
+{
+    // Validate output
+    ArrayAccessor<KeyPoint> array(_target);
+    fast_validate_keypoints(array.buffer(), array.buffer() + array.num_values(), _reference.begin(), _reference.end(), tolerance, allowed_missing, allowed_mismatching);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, NEFastCornersFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::Large2DShapes(), framework::dataset::make("Format", Format::U8)),
+                                                                                                                 framework::dataset::make("SuppressNonMax", { false, true })),
+                                                                                                         framework::dataset::make("BorderMode", BorderMode::UNDEFINED)))
+{
+    // Validate output
+    ArrayAccessor<KeyPoint> array(_target);
+    fast_validate_keypoints(array.buffer(), array.buffer() + array.num_values(), _reference.begin(), _reference.end(), tolerance, allowed_missing, allowed_mismatching);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/fixtures/FastCornersFixture.h b/tests/validation/fixtures/FastCornersFixture.h
new file mode 100644
index 0000000000..0b827f74b5
--- /dev/null
+++ b/tests/validation/fixtures/FastCornersFixture.h
@@ -0,0 +1,134 @@
+/*
+ * Copyright (c) 2017-2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef ARM_COMPUTE_TEST_FAST_CORNERS_FIXTURE
+#define ARM_COMPUTE_TEST_FAST_CORNERS_FIXTURE
+
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/IAccessor.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Fixture.h"
+#include "tests/validation/reference/FastCorners.h"
+
+#include <random>
+
+namespace arm_compute
+{
+class CLFastCorners;
+class NEFastCorners;
+
+namespace test
+{
+namespace validation
+{
+template <typename TensorType, typename AccessorType, typename ArrayType, typename FunctionType, typename T>
+class FastCornersValidationFixture : public framework::Fixture
+{
+public:
+    template <typename...>
+    void setup(TensorShape shape, Format format, bool suppress_nonmax, BorderMode border_mode)
+    {
+        std::mt19937                           gen(library->seed());
+        std::uniform_int_distribution<uint8_t> int_dist(0, 255);
+        std::uniform_real_distribution<float>  real_dist(0, 255);
+
+        const uint8_t constant_border_value = int_dist(gen);
+        const float   threshold             = real_dist(gen);
+
+        _target    = compute_target(shape, format, threshold, suppress_nonmax, border_mode, constant_border_value);
+        _reference = compute_reference(shape, format, threshold, suppress_nonmax, border_mode, constant_border_value);
+    }
+
+protected:
+    template <typename U>
+    void fill(U &&tensor)
+    {
+        library->fill_tensor_uniform(tensor, 0);
+    }
+
+    template <typename F, typename std::enable_if<std::is_same<F, CLFastCorners>::value, int>::type = 0>
+    void configure_target(F &func, TensorType &src, ArrayType &corners, unsigned int *num_corners, float threshold, bool suppress_nonmax, BorderMode border_mode, uint8_t constant_border_value)
+    {
+        func.configure(&src, threshold, suppress_nonmax, &corners, num_corners, border_mode, constant_border_value);
+    }
+
+    template <typename F, typename std::enable_if<std::is_same<F, NEFastCorners>::value, int>::type = 0>
+    void configure_target(F &func, TensorType &src, ArrayType &corners, unsigned int *num_corners, float threshold, bool suppress_nonmax, BorderMode border_mode, uint8_t constant_border_value)
+    {
+        ARM_COMPUTE_UNUSED(num_corners);
+        // ARM_COMPUTE_ERROR_ON(num_corners);
+        func.configure(&src, threshold, suppress_nonmax, &corners, border_mode, constant_border_value);
+    }
+
+    ArrayType compute_target(const TensorShape &shape, Format format, float threshold, bool suppress_nonmax, BorderMode border_mode, uint8_t constant_border_value)
+    {
+        // Create tensors
+        TensorType src = create_tensor<TensorType>(shape, data_type_from_format(format));
+        src.info()->set_format(format);
+
+        // Create array of keypoints
+        ArrayType    corners(shape.total_size());
+        unsigned int num_corners = shape.total_size();
+
+        // Create and configure function
+        FunctionType fast_corners;
+        configure_target<FunctionType>(fast_corners, src, corners, &num_corners, threshold, suppress_nonmax, border_mode, constant_border_value);
+
+        ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+        // Allocate tensors
+        src.allocator()->allocate();
+
+        ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+        // Fill tensors
+        fill(AccessorType(src));
+
+        // Compute function
+        fast_corners.run();
+
+        return corners;
+    }
+
+    std::vector<KeyPoint> compute_reference(const TensorShape &shape, Format format, float threshold, bool suppress_nonmax, BorderMode border_mode, uint8_t constant_border_value)
+    {
+        // Create reference
+        SimpleTensor<T> src{ shape, format };
+
+        // Fill reference
+        fill(src);
+
+        // Compute reference
+        return reference::fast_corners<T>(src, threshold, suppress_nonmax, border_mode, constant_border_value);
+    }
+
+    ArrayType             _target{};
+    std::vector<KeyPoint> _reference{};
+};
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_FAST_CORNERS_FIXTURE */
diff --git a/tests/validation/reference/FastCorners.cpp b/tests/validation/reference/FastCorners.cpp
new file mode 100644
index 0000000000..ae70e365c3
--- /dev/null
+++ b/tests/validation/reference/FastCorners.cpp
@@ -0,0 +1,252 @@
+/*
+ * Copyright (c) 2017-2018 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 "FastCorners.h"
+
+#include "Utils.h"
+#include "tests/validation/Helpers.h"
+#include "tests/validation/reference/NonMaximaSuppression.h"
+
+#include "tests/framework/Asserts.h"
+#include <iomanip>
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+namespace
+{
+constexpr unsigned int bresenham_radius = 3;
+constexpr unsigned int bresenham_count  = 16;
+
+/*
+    Offsets of the 16 pixels in the Bresenham circle of radius 3 centered on P
+        . . . . . . . . .
+        . . . F 0 1 . . .
+        . . E . . . 2 . .
+        . D . . . . . 3 .
+        . C . . P . . 4 .
+        . B . . . . . 5 .
+        . . A . . . 6 . .
+        . . . 9 8 7 . . .
+        . . . . . . . . .
+*/
+const std::array<std::array<int, 2>, 16> circle_offsets =
+{
+    {
+        { { 0, -3 } },  // 0 - pixel #1
+        { { 1, -3 } },  // 1 - pixel #2
+        { { 2, -2 } },  // 2 - pixel #3
+        { { 3, -1 } },  // 3 - pixel #4
+        { { 3, 0 } },   // 4 - pixel #5
+        { { 3, 1 } },   // 5 - pixel #6
+        { { 2, 2 } },   // 6 - pixel #7
+        { { 1, 3 } },   // 7 - pixel #8
+        { { 0, 3 } },   // 8 - pixel #9
+        { { -1, 3 } },  // 9 - pixel #10
+        { { -2, 2 } },  // A - pixel #11
+        { { -3, 1 } },  // B - pixel #12
+        { { -3, 0 } },  // C - pixel #13
+        { { -3, -1 } }, // D - pixel #14
+        { { -2, -2 } }, // E - pixel #15
+        { { -1, -3 } }  // F - pixel #16
+    }
+};
+
+/*
+    FAST-9 bit masks for consecutive points surrounding a corner candidate
+
+    // Speed-up rejection of non-corners by checking pixels 1, 9, then 5, 13...
+    const std::array<unsigned int, 16> fast9_order = { { 0, 8, 4, 12, 1, 2, 3, 5, 6, 7, 9, 10, 11, 13, 14, 15 } };
+*/
+const std::array<uint16_t, 16> fast9_masks =
+{
+    {
+        0x01FF, // 0000 0001 1111 1111
+        0x03FE, // 0000 0011 1111 1110
+        0x07FC, // 0000 0111 1111 1100
+        0x0FF8, // 0000 1111 1111 1000
+        0x1FF0, // 0001 1111 1111 0000
+        0x3FE0, // 0011 1111 1110 0000
+        0x7FC0, // 0111 1111 1100 0000
+        0xFF80, // 1111 1111 1000 0000
+        0xFF01, // 1111 1111 0000 0001
+        0xFE03, // 1111 1110 0000 0011
+        0xFC07, // 1111 1100 0000 0111
+        0xF80F, // 1111 1000 0000 1111
+        0xF01F, // 1111 0000 0001 1111
+        0xE03F, // 1110 0000 0011 1111
+        0xC07F, // 1100 0000 0111 1111
+        0x80FF  // 1000 0000 1111 1111
+    }
+};
+
+inline bool in_range(const uint8_t low, const uint8_t high, const uint8_t val)
+{
+    return low <= val && val <= high;
+}
+
+template <typename T, typename F>
+bool is_a_corner(const Coordinates &candidate, const SimpleTensor<T> &src, uint8_t threshold, BorderMode border_mode, T constant_border_value, F intensity_at)
+{
+    const auto intensity_p   = tensor_elem_at(src, candidate, border_mode, constant_border_value);
+    const auto thresh_bright = intensity_p + threshold;
+    const auto thresh_dark   = intensity_p - threshold;
+
+    // Quicker rejection of non-corner points by checking pixels 1, 9 then 5, 13 around the candidate
+    const auto p1  = intensity_at(candidate, 0);
+    const auto p9  = intensity_at(candidate, 8);
+    const auto p5  = intensity_at(candidate, 4);
+    const auto p13 = intensity_at(candidate, 12);
+
+    if((in_range(thresh_dark, thresh_bright, p1) && in_range(thresh_dark, thresh_bright, p9))
+       || (in_range(thresh_dark, thresh_bright, p5) && in_range(thresh_dark, thresh_bright, p13)))
+    {
+        return false;
+    }
+
+    uint16_t mask_bright = 0;
+    uint16_t mask_dark   = 0;
+
+    // Set bits of the brighter/darker pixels mask accordingly
+    for(unsigned int n = 0; n < bresenham_count; ++n)
+    {
+        T intensity_n = intensity_at(candidate, n);
+        mask_bright |= (intensity_n > thresh_bright) << n;
+        mask_dark |= (intensity_n < thresh_dark) << n;
+    }
+
+    // Mark as corner candidate if brighter/darker pixel sequence satisfies any one of the FAST-9 masks
+    const auto found = std::find_if(fast9_masks.begin(), fast9_masks.end(), [&](decltype(fast9_masks[0]) mask)
+    {
+        return (mask_bright & mask) == mask || (mask_dark & mask) == mask;
+    });
+
+    return found != fast9_masks.end();
+}
+} // namespace
+
+template <typename T>
+std::vector<KeyPoint> fast_corners(const SimpleTensor<T> &src, float input_thresh, bool suppress_nonmax, BorderMode border_mode, T constant_border_value)
+{
+    // Get intensity of pixel at given index on the Bresenham circle around a candidate point
+    const auto intensity_at = [&](const Coordinates & point, const unsigned int idx)
+    {
+        const auto  offs = circle_offsets[idx];
+        Coordinates px{ point.x() + offs[0], point.y() + offs[1] };
+        return tensor_elem_at(src, px, border_mode, constant_border_value);
+    };
+
+    const auto            threshold = static_cast<uint8_t>(input_thresh);
+    std::vector<KeyPoint> corners;
+
+    // 1. Detect potential corners (the segment test)
+    std::vector<Coordinates> corner_candidates;
+    SimpleTensor<float>      scores(src.shape(), DataType::F32);
+    ValidRegion              valid_region = shape_to_valid_region(src.shape(), BorderMode::UNDEFINED == border_mode, BorderSize(bresenham_radius));
+
+    for(int i = 0; i < src.num_elements(); ++i)
+    {
+        Coordinates candidate = index2coord(src.shape(), i);
+        scores[i]             = 0.f;
+        if(!is_in_valid_region(valid_region, candidate))
+        {
+            continue;
+        }
+
+        if(is_a_corner(candidate, src, threshold, border_mode, constant_border_value, intensity_at))
+        {
+            corner_candidates.emplace_back(candidate);
+            scores[i] = 1.f;
+        }
+    }
+
+    // 2. Calculate corner scores if non-maxima suppression
+    // The corner response Cp function is defined as the largest threshold t for which the pixel p remains a corner
+    if(suppress_nonmax)
+    {
+        for(const auto &candidate : corner_candidates)
+        {
+            const auto index = coord2index(scores.shape(), candidate);
+
+#ifdef CALC_CORNER_RESPONSE_BY_ITERATION
+            auto response = threshold;
+            while(is_a_corner(candidate, src, response, border_mode, constant_border_value, intensity_at))
+            {
+                response += 1;
+            }
+            scores[index] = response - 1;
+#else  // CALC_CORNER_RESPONSE_BY_ITERATION
+            uint8_t thresh_max = UINT8_MAX;
+            uint8_t thresh_min = threshold;
+            uint8_t response   = (thresh_min + thresh_max) / 2;
+
+            while(thresh_max - thresh_min > 1)
+            {
+                response = (thresh_min + thresh_max) / 2;
+                if(is_a_corner(candidate, src, response, border_mode, constant_border_value, intensity_at))
+                {
+                    thresh_min = response; // raise threshold
+                }
+                else
+                {
+                    thresh_max = response; // lower threshold
+                }
+            }
+            scores[index] = thresh_min;
+#endif // CALC_CORNER_RESPONSE_BY_ITERATION
+        }
+
+        scores       = non_maxima_suppression(scores, border_mode, static_cast<float>(constant_border_value));
+        valid_region = shape_to_valid_region(scores.shape(), BorderMode::UNDEFINED == border_mode, BorderSize(bresenham_radius + 1));
+    }
+
+    for(const auto &candidate : corner_candidates)
+    {
+        const auto index = coord2index(scores.shape(), candidate);
+        if(scores[index] > 0.f && is_in_valid_region(valid_region, candidate))
+        {
+            KeyPoint corner;
+            corner.x               = candidate.x();
+            corner.y               = candidate.y();
+            corner.strength        = scores[index];
+            corner.tracking_status = 1;
+            corner.scale           = 0.f;
+            corner.orientation     = 0.f;
+            corner.error           = 0.f;
+            corners.emplace_back(corner);
+        }
+    }
+
+    return corners;
+}
+
+template std::vector<KeyPoint> fast_corners(const SimpleTensor<uint8_t> &src, float threshold, bool suppress_nonmax, BorderMode border_mode, uint8_t constant_border_value);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/reference/FastCorners.h b/tests/validation/reference/FastCorners.h
new file mode 100644
index 0000000000..3d070d97b4
--- /dev/null
+++ b/tests/validation/reference/FastCorners.h
@@ -0,0 +1,44 @@
+/*
+ * Copyright (c) 2017-2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef __ARM_COMPUTE_TEST_FAST_CORNERS_H__
+#define __ARM_COMPUTE_TEST_FAST_CORNERS_H__
+
+#include "arm_compute/core/Types.h"
+#include "tests/SimpleTensor.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+std::vector<KeyPoint> fast_corners(const SimpleTensor<T> &src, float input_thresh, bool suppress_nonmax, BorderMode border_mode, T constant_border_value = 0);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_FAST_CORNERS_H__ */
diff --git a/tests/validation/reference/NonMaximaSuppression.cpp b/tests/validation/reference/NonMaximaSuppression.cpp
index eab5cecfc8..34c6c07abc 100644
--- a/tests/validation/reference/NonMaximaSuppression.cpp
+++ b/tests/validation/reference/NonMaximaSuppression.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017 ARM Limited.
+ * Copyright (c) 2017-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -69,6 +69,7 @@ SimpleTensor<T> non_maxima_suppression(const SimpleTensor<T> &src, BorderMode bo
 }
 
 template SimpleTensor<float> non_maxima_suppression(const SimpleTensor<float> &src, BorderMode border_mode, float constant_border_value);
+template SimpleTensor<uint8_t> non_maxima_suppression(const SimpleTensor<uint8_t> &src, BorderMode border_mode, uint8_t constant_border_value);
 } // namespace reference
 } // namespace validation
 } // namespace test