COMPMID-424 Implemented reference implementation and tests for WarpAffine

Change-Id: I4924ab1de17adc3b880a5cc22f2497abbc8e221b
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/85820
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Reviewed-by: Steven Niu <steven.niu@arm.com>

10 files changed, 741 insertions, 20 deletions

diff --git a/src/core/CL/cl_kernels/warp_affine.cl b/src/core/CL/cl_kernels/warp_affine.cl
index 0a4748f452..f41821cdca 100644
--- a/src/core/CL/cl_kernels/warp_affine.cl
+++ b/src/core/CL/cl_kernels/warp_affine.cl
@@ -84,7 +84,7 @@ __kernel void warp_affine_nearest_neighbour(
 {
     Image in  = CONVERT_TO_IMAGE_STRUCT_NO_STEP(in);
     Image out = CONVERT_TO_IMAGE_STRUCT(out);
-    vstore4(read_texels4(&in, convert_int8(clamp_to_border(apply_affine_transform(get_current_coords(), build_affine_mtx()), width, height))), 0, out.ptr);
+    vstore4(read_texels4(&in, convert_int8_rtn(clamp_to_border(apply_affine_transform(get_current_coords(), build_affine_mtx()), width, height))), 0, out.ptr);
 }
 
 /** Performs an affine transform on an image interpolating with the BILINEAR method. Input and output are single channel U8.
@@ -116,5 +116,5 @@ __kernel void warp_affine_bilinear(
 {
     Image in  = CONVERT_TO_IMAGE_STRUCT_NO_STEP(in);
     Image out = CONVERT_TO_IMAGE_STRUCT(out);
-    vstore4(bilinear_interpolate(&in, clamp_to_border(apply_affine_transform(get_current_coords(), build_affine_mtx()), width, height), width, height), 0, out.ptr);
+    vstore4(bilinear_interpolate(&in, apply_affine_transform(get_current_coords(), build_affine_mtx()), width, height), 0, out.ptr);
 }
diff --git a/src/core/CL/kernels/CLWarpAffineKernel.cpp b/src/core/CL/kernels/CLWarpAffineKernel.cpp
index e549dbc258..be095f2c34 100644
--- a/src/core/CL/kernels/CLWarpAffineKernel.cpp
+++ b/src/core/CL/kernels/CLWarpAffineKernel.cpp
@@ -88,8 +88,11 @@ void CLWarpAffineKernel::configure(const ICLTensor *input, ICLTensor *output, co
 
     Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration));
 
-    AccessWindowHorizontal input_access(input->info(), 0, num_elems_processed_per_iteration);
-    AccessWindowStatic     output_access(output->info(), 0, 0, output->info()->dimension(0), output->info()->dimension(1));
+    int       total_right  = ceil_to_multiple(input->info()->dimension(0), num_elems_processed_per_iteration);
+    const int access_right = total_right + (((total_right - input->info()->dimension(0)) == 0) ? border_size().right : 0);
+
+    AccessWindowStatic     input_access(input->info(), -border_size().left, -border_size().top, access_right, input->info()->dimension(1) + border_size().bottom);
+    AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration);
 
     update_window_and_padding(win, input_access, output_access);
 
diff --git a/src/core/NEON/kernels/NEWarpKernel.cpp b/src/core/NEON/kernels/NEWarpKernel.cpp
index 62f4e5d057..ab8ab14ae5 100644
--- a/src/core/NEON/kernels/NEWarpKernel.cpp
+++ b/src/core/NEON/kernels/NEWarpKernel.cpp
@@ -143,7 +143,11 @@ void NEWarpAffineKernel<interpolation>::warp_undefined(const Window &window)
     const float start_y0 = M10 * window.x().start();
 
     // Current row
-    int y_cur = window.y().start();
+    int y_cur  = window.y().start();
+    int z_cur  = window.z().start();
+    int d3_cur = window[3].start();
+    int d4_cur = window[4].start();
+    int d5_cur = window[5].start();
 
     // const_x0 and const_y0 are the constant parts of x0 and y0 during the row processing
     float const_x0 = M01 * y_cur + M02;
@@ -155,10 +159,14 @@ void NEWarpAffineKernel<interpolation>::warp_undefined(const Window &window)
 
     execute_window_loop(window, [&](const Coordinates & id)
     {
-        // Check if we are processing a new row. If so, update the current row (y_cur), x0 and y0
-        if(y_cur != id.y())
+        // Check if we are processing a new row. If so, update the current processed row (y_cur), x0, y0 and z0
+        if((y_cur != id.y()) || (z_cur != id.z()) || (d3_cur != id[3]) || (d4_cur != id[4]) || (d5_cur != id[5]))
         {
-            y_cur = id.y();
+            y_cur  = id.y();
+            z_cur  = id.z();
+            d3_cur = id[3];
+            d4_cur = id[4];
+            d5_cur = id[5];
 
             const_x0 = M01 * y_cur + M02;
             const_y0 = M11 * y_cur + M12;
@@ -222,7 +230,11 @@ void NEWarpAffineKernel<interpolation>::warp_constant(const Window &window)
     const float start_y0 = M10 * window.x().start();
 
     // Current row
-    int y_cur = window.y().start();
+    int y_cur  = window.y().start();
+    int z_cur  = window.z().start();
+    int d3_cur = window[3].start();
+    int d4_cur = window[4].start();
+    int d5_cur = window[5].start();
 
     // const_x0 and const_y0 are the constant parts of x0 and y0 during the row processing
     float const_x0 = M01 * y_cur + M02;
@@ -234,10 +246,14 @@ void NEWarpAffineKernel<interpolation>::warp_constant(const Window &window)
 
     execute_window_loop(window, [&](const Coordinates & id)
     {
-        // Check if we are processing a new row. If so, update the current row (y_cur), x0 and y0
-        if(y_cur != id.y())
+        // Check if we are processing a new row. If so, update the current processed row (y_cur), x0, y0 and z0
+        if((y_cur != id.y()) || (z_cur != id.z()) || (d3_cur != id[3]) || (d4_cur != id[4]) || (d5_cur != id[5]))
         {
-            y_cur = id.y();
+            y_cur  = id.y();
+            z_cur  = id.z();
+            d3_cur = id[3];
+            d4_cur = id[4];
+            d5_cur = id[5];
 
             const_x0 = M01 * y_cur + M02;
             const_y0 = M11 * y_cur + M12;
@@ -264,7 +280,34 @@ void NEWarpAffineKernel<interpolation>::warp_constant(const Window &window)
         }
         else
         {
-            *out.ptr() = _constant_border_value;
+            switch(interpolation)
+            {
+                case InterpolationPolicy::NEAREST_NEIGHBOR:
+                    *out.ptr() = _constant_border_value;
+                    break;
+                case InterpolationPolicy::BILINEAR:
+                {
+                    const auto xi   = clamp<int>(std::floor(x0), min_x - 1, max_x);
+                    const auto yi   = clamp<int>(std::floor(y0), min_y - 1, max_y);
+                    const auto xi_1 = clamp<int>(std::floor(x0 + 1), min_x - 1, max_x);
+                    const auto yi_1 = clamp<int>(std::floor(y0 + 1), min_y - 1, max_y);
+
+                    const float dx  = x0 - std::floor(x0);
+                    const float dy  = y0 - std::floor(y0);
+                    const float dx1 = 1.0f - dx;
+                    const float dy1 = 1.0f - dy;
+
+                    const float a00 = *(in.ptr() + xi + yi * stride);
+                    const float a01 = *(in.ptr() + xi_1 + yi * stride);
+                    const float a10 = *(in.ptr() + xi + yi_1 * stride);
+                    const float a11 = *(in.ptr() + xi_1 + yi_1 * stride);
+
+                    *out.ptr() = a00 * (dx1 * dy1) + a01 * (dx * dy1) + a10 * (dx1 * dy) + a11 * (dx * dy);
+                }
+                break;
+                default:
+                    ARM_COMPUTE_ERROR("Interpolation not supported");
+            }
         }
 
         x0 += M00;
@@ -292,7 +335,11 @@ void NEWarpAffineKernel<interpolation>::warp_replicate(const Window &window)
     const size_t stride = _input->info()->strides_in_bytes()[1];
 
     // Current row
-    int y_cur = window.y().start();
+    int y_cur  = window.y().start();
+    int z_cur  = window.z().start();
+    int d3_cur = window[3].start();
+    int d4_cur = window[4].start();
+    int d5_cur = window[5].start();
 
     const float M00 = _matrix[0];
     const float M10 = _matrix[1];
@@ -314,10 +361,14 @@ void NEWarpAffineKernel<interpolation>::warp_replicate(const Window &window)
 
     execute_window_loop(window, [&](const Coordinates & id)
     {
-        // Check if we are processing a new row. If so, update the current row (y_cur), x0 and y0
-        if(y_cur != id.y())
+        // Check if we are processing a new row. If so, update the current processed row (y_cur), x0, y0 and z0
+        if((y_cur != id.y()) || (z_cur != id.z()) || (d3_cur != id[3]) || (d4_cur != id[4]) || (d5_cur != id[5]))
         {
-            y_cur = id.y();
+            y_cur  = id.y();
+            z_cur  = id.z();
+            d3_cur = id[3];
+            d4_cur = id[4];
+            d5_cur = id[5];
 
             const_x0 = M01 * y_cur + M02;
             const_y0 = M11 * y_cur + M12;
@@ -345,10 +396,34 @@ void NEWarpAffineKernel<interpolation>::warp_replicate(const Window &window)
         else
         {
             // Clamp coordinates
-            const auto xi = clamp<int>(x0, min_x, max_x - 1);
-            const auto yi = clamp<int>(y0, min_y, max_y - 1);
+            const auto xi = clamp<int>(std::floor(x0), min_x, max_x - 1);
+            const auto yi = clamp<int>(std::floor(y0), min_y, max_y - 1);
+            switch(interpolation)
+            {
+                case InterpolationPolicy::NEAREST_NEIGHBOR:
+                    *out.ptr() = *(in.ptr() + xi + yi * stride);
+                    break;
+                case InterpolationPolicy::BILINEAR:
+                {
+                    const auto xi_1 = clamp<int>(std::floor(x0 + 1), min_x, max_x - 1);
+                    const auto yi_1 = clamp<int>(std::floor(y0 + 1), min_y, max_y - 1);
 
-            *out.ptr() = *(in.ptr() + xi + yi * stride);
+                    const float dx  = x0 - std::floor(x0);
+                    const float dy  = y0 - std::floor(y0);
+                    const float dx1 = 1.0f - dx;
+                    const float dy1 = 1.0f - dy;
+
+                    const float a00 = *(in.ptr() + xi + yi * stride);
+                    const float a01 = *(in.ptr() + xi_1 + yi * stride);
+                    const float a10 = *(in.ptr() + xi + yi_1 * stride);
+                    const float a11 = *(in.ptr() + xi_1 + yi_1 * stride);
+
+                    *out.ptr() = a00 * (dx1 * dy1) + a01 * (dx * dy1) + a10 * (dx1 * dy) + a11 * (dx * dy);
+                }
+                break;
+                default:
+                    ARM_COMPUTE_ERROR("Interpolation not supported");
+            }
         }
 
         x0 += M00;
diff --git a/tests/validation/CL/WarpAffine.cpp b/tests/validation/CL/WarpAffine.cpp
new file mode 100644
index 0000000000..9177fcb5a7
--- /dev/null
+++ b/tests/validation/CL/WarpAffine.cpp
@@ -0,0 +1,120 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/core/Utils.h"
+#include "arm_compute/runtime/CL/functions/CLWarpAffine.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/BorderModeDataset.h"
+#include "tests/datasets/InterpolationPolicyDataset.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/CPP/Utils.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/WarpAffineFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Tolerance */
+constexpr AbsoluteTolerance<uint8_t> tolerance(1);
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(WarpAffine)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(combine(combine(concat(datasets::SmallShapes(), datasets::LargeShapes()), framework::dataset::make("DataType", DataType::U8)),
+                                                                           framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+                                                                   datasets::BorderModes()),
+               shape, data_type, policy, border_mode)
+{
+    // Generate a random constant value if border_mode is constant
+    std::mt19937                           gen(library->seed());
+    std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+    uint8_t                                constant_border_value = distribution_u8(gen);
+
+    // Create the matrix
+    std::array<float, 6> matrix{ {} };
+    fill_warp_matrix<6>(matrix, 3, 2);
+
+    // Create tensors
+    CLTensor src = create_tensor<CLTensor>(shape, data_type);
+    CLTensor dst = create_tensor<CLTensor>(shape, data_type);
+
+    ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+    ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+    // Create and configure function
+    CLWarpAffine warp_affine;
+    warp_affine.configure(&src, &dst, matrix.data(), policy, border_mode, constant_border_value);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+
+    validate(src.info()->valid_region(), valid_region);
+    validate(dst.info()->valid_region(), valid_region);
+
+    // Validate padding
+    int               total_right  = ceil_to_multiple(shape[0], 4);
+    const int         access_right = total_right + (((total_right - shape[0]) == 0) ? 1 : 0);
+    const PaddingSize read_padding(1, access_right - shape[0], 1, 1);
+    validate(src.info()->padding(), read_padding);
+
+    PaddingCalculator calculator(shape.x(), 4);
+    validate(dst.info()->padding(), calculator.required_padding());
+}
+
+template <typename T>
+using CLWarpAffineFixture = WarpAffineValidationFixture<CLTensor, CLAccessor, CLWarpAffine, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, CLWarpAffineFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::U8)),
+                                                                                                                  framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+                                                                                                          datasets::BorderModes()))
+{
+    // Validate output
+    validate(CLAccessor(_target), _reference, _valid_mask, tolerance, 0.02f);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLWarpAffineFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", DataType::U8)),
+                                                                                                                framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+                                                                                                        datasets::BorderModes()))
+{
+    // Validate output
+    validate(CLAccessor(_target), _reference, _valid_mask, tolerance, 0.02f);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/CPP/Utils.h b/tests/validation/CPP/Utils.h
index 91d1afe1d7..e9a0f09847 100644
--- a/tests/validation/CPP/Utils.h
+++ b/tests/validation/CPP/Utils.h
@@ -100,6 +100,32 @@ void apply_2d_spatial_filter(Coordinates coord, const SimpleTensor<T> &src, Simp
 }
 
 RawTensor transpose(const RawTensor &src, int chunk_width = 1);
+
+/** Fill matrix random.
+ *
+ * @param[in,out] matrix Matrix
+ * @param[in]     cols   Columns (width) of matrix
+ * @param[in]     rows   Rows (height) of matrix
+ */
+template <std::size_t SIZE>
+inline void fill_warp_matrix(std::array<float, SIZE> &matrix, int cols, int rows)
+{
+    std::mt19937                          gen(library.get()->seed());
+    std::uniform_real_distribution<float> dist(-1, 1);
+
+    for(int v = 0, r = 0; r < rows; ++r)
+    {
+        for(int c = 0; c < cols; ++c, ++v)
+        {
+            matrix[v] = dist(gen);
+        }
+    }
+    if(SIZE == 9)
+    {
+        matrix[(cols * rows) - 1] = 1;
+    }
+}
+
 } // namespace validation
 } // namespace test
 } // namespace arm_compute
diff --git a/tests/validation/CPP/WarpAffine.cpp b/tests/validation/CPP/WarpAffine.cpp
new file mode 100644
index 0000000000..7b903b7661
--- /dev/null
+++ b/tests/validation/CPP/WarpAffine.cpp
@@ -0,0 +1,135 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "WarpAffine.h"
+
+#include "Utils.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+bool valid_bilinear_policy(float xn, float yn, int width, int height, BorderMode border_mode)
+{
+    if(border_mode != BorderMode::UNDEFINED)
+    {
+        return true;
+    }
+    if((0 <= yn + 1) && (yn + 1 < height) && (0 <= xn + 1) && (xn + 1 < width))
+    {
+        return true;
+    }
+    return false;
+}
+
+template <typename T>
+SimpleTensor<T> warp_affine(const SimpleTensor<T> &src, SimpleTensor<T> &valid_mask, const float *matrix, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
+{
+    SimpleTensor<T> dst(src.shape(), src.data_type());
+
+    // x0 = M00 * x + M01 * y + M02
+    // y0 = M10 * x + M11 * y + M12
+    const float M00 = matrix[0];
+    const float M10 = matrix[1];
+    const float M01 = matrix[0 + 1 * 2];
+    const float M11 = matrix[1 + 1 * 2];
+    const float M02 = matrix[0 + 2 * 2];
+    const float M12 = matrix[1 + 2 * 2];
+
+    const int width  = src.shape().x();
+    const int height = src.shape().y();
+
+    for(int element_idx = 0; element_idx < src.num_elements(); ++element_idx)
+    {
+        valid_mask[element_idx] = 1;
+        Coordinates id          = index2coord(src.shape(), element_idx);
+        int         idx         = id.x();
+        int         idy         = id.y();
+
+        float x0 = M00 * idx + M01 * idy + M02;
+        float y0 = M10 * idx + M11 * idy + M12;
+
+        id.set(0, static_cast<int>(std::floor(x0)));
+        id.set(1, static_cast<int>(std::floor(y0)));
+        if((0 <= y0) && (y0 < height) && (0 <= x0) && (x0 < width))
+        {
+            switch(policy)
+            {
+                case InterpolationPolicy::NEAREST_NEIGHBOR:
+                    dst[element_idx] = tensor_elem_at(src, id, border_mode, constant_border_value);
+                    break;
+                case InterpolationPolicy::BILINEAR:
+                    (valid_bilinear_policy(x0, y0, width, height, border_mode)) ? dst[element_idx] = bilinear_policy(src, id, x0, y0, border_mode, constant_border_value) :
+                                                                                                     valid_mask[element_idx] = 0;
+                    break;
+                case InterpolationPolicy::AREA:
+                default:
+                    ARM_COMPUTE_ERROR("Interpolation not supported");
+            }
+        }
+        else
+        {
+            if(border_mode == BorderMode::UNDEFINED)
+            {
+                valid_mask[element_idx] = 0;
+            }
+            else
+            {
+                switch(policy)
+                {
+                    case InterpolationPolicy::NEAREST_NEIGHBOR:
+                        if(border_mode == BorderMode::CONSTANT)
+                        {
+                            dst[element_idx] = constant_border_value;
+                        }
+                        else if(border_mode == BorderMode::REPLICATE)
+                        {
+                            id.set(0, std::max(0, std::min(static_cast<int>(x0), width - 1)));
+                            id.set(1, std::max(0, std::min(static_cast<int>(y0), height - 1)));
+                            dst[element_idx] = src[coord2index(src.shape(), id)];
+                        }
+                        break;
+                    case InterpolationPolicy::BILINEAR:
+                        dst[element_idx] = bilinear_policy(src, id, x0, y0, border_mode, constant_border_value);
+                        break;
+                    case InterpolationPolicy::AREA:
+                    default:
+                        ARM_COMPUTE_ERROR("Interpolation not supported");
+                }
+            }
+        }
+    }
+
+    return dst;
+}
+
+template SimpleTensor<uint8_t> warp_affine(const SimpleTensor<uint8_t> &src, SimpleTensor<uint8_t> &valid_mask, const float *matrix, InterpolationPolicy policy, BorderMode border_mode,
+                                           uint8_t constant_border_value);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
\ No newline at end of file
diff --git a/tests/validation/CPP/WarpAffine.h b/tests/validation/CPP/WarpAffine.h
new file mode 100644
index 0000000000..973b1b2e29
--- /dev/null
+++ b/tests/validation/CPP/WarpAffine.h
@@ -0,0 +1,43 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef __ARM_COMPUTE_TEST_WARP_AFFINE_H__
+#define __ARM_COMPUTE_TEST_WARP_AFFINE_H__
+
+#include "tests/SimpleTensor.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+SimpleTensor<T> warp_affine(const SimpleTensor<T> &src, SimpleTensor<T> &valid_mask, const float *matrix, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_WARP_AFFINE_H__ */
diff --git a/tests/validation/NEON/WarpAffine.cpp b/tests/validation/NEON/WarpAffine.cpp
new file mode 100644
index 0000000000..1b52f9f07c
--- /dev/null
+++ b/tests/validation/NEON/WarpAffine.cpp
@@ -0,0 +1,120 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEWarpAffine.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/NEON/Accessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets/BorderModeDataset.h"
+#include "tests/datasets/InterpolationPolicyDataset.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/CPP/Utils.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/WarpAffineFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Tolerance */
+constexpr AbsoluteTolerance<uint8_t> tolerance(1);
+} // namespace
+
+TEST_SUITE(NEON)
+TEST_SUITE(WarpAffine)
+
+DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(combine(combine(concat(datasets::SmallShapes(), datasets::LargeShapes()), framework::dataset::make("DataType", DataType::U8)),
+                                                                           framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+                                                                   datasets::BorderModes()),
+               shape, data_type, policy, border_mode)
+{
+    // Generate a random constant value if border_mode is constant
+    std::mt19937                           gen(library->seed());
+    std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+    uint8_t                                constant_border_value = distribution_u8(gen);
+
+    // Create the matrix
+    std::array<float, 6> matrix{ {} };
+    fill_warp_matrix<6>(matrix, 3, 2);
+
+    // Create tensors
+    Tensor src = create_tensor<Tensor>(shape, data_type);
+    Tensor dst = create_tensor<Tensor>(shape, data_type);
+
+    ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+    ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+    // Create and configure function
+    NEWarpAffine warp_affine;
+    warp_affine.configure(&src, &dst, matrix.data(), policy, border_mode, constant_border_value);
+
+    // Validate valid region
+    const ValidRegion valid_region = shape_to_valid_region(shape);
+
+    validate(dst.info()->valid_region(), valid_region);
+
+    // Validate padding
+    PaddingCalculator calculator(shape.x(), 1);
+    calculator.set_border_mode(border_mode);
+    calculator.set_border_size(1);
+
+    const PaddingSize read_padding(1);
+    const PaddingSize write_padding = calculator.required_padding();
+
+    validate(src.info()->padding(), read_padding);
+    validate(dst.info()->padding(), write_padding);
+}
+
+template <typename T>
+using NEWarpAffineFixture = WarpAffineValidationFixture<Tensor, Accessor, NEWarpAffine, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, NEWarpAffineFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::U8)),
+                                                                                                                  framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+                                                                                                          datasets::BorderModes()))
+{
+    // Validate output
+    validate(Accessor(_target), _reference, _valid_mask, tolerance, 0.02f);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, NEWarpAffineFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", DataType::U8)),
+                                                                                                                framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+                                                                                                        datasets::BorderModes()))
+{
+    // Validate output
+    validate(Accessor(_target), _reference, _valid_mask, tolerance, 0.02f);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/Validation.h b/tests/validation/Validation.h
index 5e5dab0040..f220224991 100644
--- a/tests/validation/Validation.h
+++ b/tests/validation/Validation.h
@@ -188,6 +188,19 @@ void validate(const IAccessor &tensor, const SimpleTensor<T> &reference, U toler
 template <typename T, typename U = AbsoluteTolerance<T>>
 void validate(const IAccessor &tensor, const SimpleTensor<T> &reference, const ValidRegion &valid_region, U tolerance_value = U(), float tolerance_number = 0.f);
 
+/** Validate tensors with valid mask.
+ *
+ * - Dimensionality has to be the same.
+ * - All values have to match.
+ *
+ * @note: wrap_range allows cases where reference tensor rounds up to the wrapping point, causing it to wrap around to
+ * zero while the test tensor stays at wrapping point to pass. This may permit true erroneous cases (difference between
+ * reference tensor and test tensor is multiple of wrap_range), but such errors would be detected by
+ * other test cases.
+ */
+template <typename T, typename U = AbsoluteTolerance<T>>
+void validate(const IAccessor &tensor, const SimpleTensor<T> &reference, const SimpleTensor<T> &valid_mask, U tolerance_value = U(), float tolerance_number = 0.f);
+
 /** Validate tensors against constant value.
  *
  * - All values have to match.
@@ -435,6 +448,71 @@ void validate_keypoints(T target_first, T target_last, U reference_first, U refe
 }
 
 template <typename T, typename U>
+void validate(const IAccessor &tensor, const SimpleTensor<T> &reference, const SimpleTensor<T> &valid_mask, U tolerance_value, float tolerance_number)
+{
+    int64_t num_mismatches = 0;
+    int64_t num_elements   = 0;
+
+    ARM_COMPUTE_EXPECT_EQUAL(tensor.element_size(), reference.element_size(), framework::LogLevel::ERRORS);
+    ARM_COMPUTE_EXPECT_EQUAL(tensor.data_type(), reference.data_type(), framework::LogLevel::ERRORS);
+
+    if(reference.format() != Format::UNKNOWN)
+    {
+        ARM_COMPUTE_EXPECT_EQUAL(tensor.format(), reference.format(), framework::LogLevel::ERRORS);
+    }
+
+    ARM_COMPUTE_EXPECT_EQUAL(tensor.num_channels(), reference.num_channels(), framework::LogLevel::ERRORS);
+    ARM_COMPUTE_EXPECT(compare_dimensions(tensor.shape(), reference.shape()), framework::LogLevel::ERRORS);
+
+    const int min_elements = std::min(tensor.num_elements(), reference.num_elements());
+    const int min_channels = std::min(tensor.num_channels(), reference.num_channels());
+
+    // Iterate over all elements within valid region, e.g. U8, S16, RGB888, ...
+    for(int element_idx = 0; element_idx < min_elements; ++element_idx)
+    {
+        const Coordinates id = index2coord(reference.shape(), element_idx);
+
+        if(valid_mask[element_idx] == 1)
+        {
+            // Iterate over all channels within one element
+            for(int c = 0; c < min_channels; ++c)
+            {
+                const T &target_value    = reinterpret_cast<const T *>(tensor(id))[c];
+                const T &reference_value = reinterpret_cast<const T *>(reference(id))[c];
+
+                if(!compare<U>(target_value, reference_value, tolerance_value))
+                {
+                    ARM_COMPUTE_TEST_INFO("id = " << id);
+                    ARM_COMPUTE_TEST_INFO("channel = " << c);
+                    ARM_COMPUTE_TEST_INFO("target = " << std::setprecision(5) << framework::make_printable(target_value));
+                    ARM_COMPUTE_TEST_INFO("reference = " << std::setprecision(5) << framework::make_printable(reference_value));
+                    ARM_COMPUTE_TEST_INFO("tolerance = " << std::setprecision(5) << framework::make_printable(static_cast<typename U::value_type>(tolerance_value)));
+                    framework::ARM_COMPUTE_PRINT_INFO();
+
+                    ++num_mismatches;
+                }
+
+                ++num_elements;
+            }
+        }
+        else
+        {
+            ++num_elements;
+        }
+    }
+
+    if(num_elements > 0)
+    {
+        const int64_t absolute_tolerance_number = tolerance_number * num_elements;
+        const float   percent_mismatches        = static_cast<float>(num_mismatches) / num_elements * 100.f;
+
+        ARM_COMPUTE_TEST_INFO(num_mismatches << " values (" << std::fixed << std::setprecision(2) << percent_mismatches
+                              << "%) mismatched (maximum tolerated " << std::setprecision(2) << tolerance_number << "%)");
+        ARM_COMPUTE_EXPECT(num_mismatches <= absolute_tolerance_number, framework::LogLevel::ERRORS);
+    }
+}
+
+template <typename T, typename U>
 bool validate(T target, T reference, U tolerance)
 {
     ARM_COMPUTE_TEST_INFO("reference = " << std::setprecision(5) << framework::make_printable(reference));
diff --git a/tests/validation/fixtures/WarpAffineFixture.h b/tests/validation/fixtures/WarpAffineFixture.h
new file mode 100644
index 0000000000..3b106e15d6
--- /dev/null
+++ b/tests/validation/fixtures/WarpAffineFixture.h
@@ -0,0 +1,121 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef ARM_COMPUTE_TEST_WARP_AFFINE_FIXTURE
+#define ARM_COMPUTE_TEST_WARP_AFFINE_FIXTURE
+
+#include <memory>
+
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/IAccessor.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Fixture.h"
+#include "tests/validation/CPP/Utils.h"
+#include "tests/validation/CPP/WarpAffine.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+template <typename TensorType, typename AccessorType, typename FunctionType, typename T>
+class WarpAffineValidationFixture : public framework::Fixture
+{
+public:
+    template <typename...>
+    void setup(TensorShape shape, DataType data_type, InterpolationPolicy policy, BorderMode border_mode)
+    {
+        // Generate a random constant value if border_mode is constant
+        std::mt19937                           gen(library->seed());
+        std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+        uint8_t                                constant_border_value = distribution_u8(gen);
+
+        // Create the matrix
+        std::array<float, 6> matrix{ {} };
+        fill_warp_matrix<6>(matrix, 3, 2);
+
+        _target    = compute_target(shape, data_type, matrix.data(), policy, border_mode, constant_border_value);
+        _reference = compute_reference(shape, data_type, matrix.data(), policy, border_mode, constant_border_value);
+    }
+
+protected:
+    template <typename U>
+    void fill(U &&tensor)
+    {
+        library->fill_tensor_uniform(tensor, 0);
+    }
+
+    TensorType compute_target(const TensorShape &shape, DataType data_type, const float *matrix, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
+    {
+        // Create tensors
+        TensorType src = create_tensor<TensorType>(shape, data_type);
+        TensorType dst = create_tensor<TensorType>(shape, data_type);
+
+        // Create and configure function
+        FunctionType warp_affine;
+        warp_affine.configure(&src, &dst, matrix, policy, border_mode, constant_border_value);
+
+        ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+        ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+        // Allocate tensors
+        src.allocator()->allocate();
+        dst.allocator()->allocate();
+        ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS);
+        ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+        // Fill tensors
+        fill(AccessorType(src));
+
+        // Compute function
+        warp_affine.run();
+
+        return dst;
+    }
+
+    SimpleTensor<T> compute_reference(const TensorShape &shape, DataType data_type, const float *matrix, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
+    {
+        // Create reference
+        SimpleTensor<T> src{ shape, data_type };
+
+        // Create the valid mask Tensor
+        _valid_mask = SimpleTensor<T>(shape, data_type);
+
+        // Fill reference
+        fill(src);
+
+        return reference::warp_affine<T>(src, _valid_mask, matrix, policy, border_mode, constant_border_value);
+    }
+
+    TensorType      _target{};
+    SimpleTensor<T> _reference{};
+    SimpleTensor<T> _valid_mask{};
+};
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_WARP_AFFINE_FIXTURE */
\ No newline at end of file