COMPMID-424 Implemented reference implementation, new output valid region and validation tests (NEON and CL) for Scale

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

18 files changed, 873 insertions, 73 deletions

diff --git a/arm_compute/core/Helpers.h b/arm_compute/core/Helpers.h
index dfcca96eec..f3702e7c93 100644
--- a/arm_compute/core/Helpers.h
+++ b/arm_compute/core/Helpers.h
@@ -461,6 +461,17 @@ bool set_data_type_if_unknown(ITensorInfo &info, DataType data_type);
  * @return True if the fixed point position has been changed.
  */
 bool set_fixed_point_position_if_zero(ITensorInfo &info, int fixed_point_position);
+/** Helper function to calculate the Valid Region for Scale.
+ *
+ * @param[in] src_info         Input tensor info used to check.
+ * @param[in] dst_shape        Shape of the output.
+ * @param[in] policy           Interpolation policy.
+ * @param[in] border_size      Size of the border.
+ * @param[in] border_undefined True if the border is undefined.
+ *
+ * @return The corrispondent valid region
+ */
+ValidRegion calculate_valid_region_scale(const ITensorInfo &src_info, const TensorShape &dst_shape, InterpolationPolicy policy, BorderSize border_size, bool border_undefined);
 } // namespace arm_compute
 
 #include "arm_compute/core/Helpers.inl"
diff --git a/arm_compute/core/Helpers.inl b/arm_compute/core/Helpers.inl
index 78e0c70e1b..90a4618fcc 100644
--- a/arm_compute/core/Helpers.inl
+++ b/arm_compute/core/Helpers.inl
@@ -303,4 +303,30 @@ inline bool set_fixed_point_position_if_zero(ITensorInfo &info, int fixed_point_
 
     return false;
 }
+
+inline ValidRegion calculate_valid_region_scale(const ITensorInfo &src_info, const TensorShape &dst_shape, InterpolationPolicy policy, BorderSize border_size, bool border_undefined)
+{
+    const auto  wr = static_cast<float>(dst_shape[0]) / static_cast<float>(src_info.tensor_shape()[0]);
+    const auto  hr = static_cast<float>(dst_shape[1]) / static_cast<float>(src_info.tensor_shape()[1]);
+    Coordinates anchor;
+    anchor.set_num_dimensions(src_info.tensor_shape().num_dimensions());
+    TensorShape new_dst_shape(dst_shape);
+    anchor.set(0, (policy == InterpolationPolicy::BILINEAR && border_undefined) ? ((static_cast<int>(src_info.valid_region().anchor[0]) + border_size.left + 0.5f) * wr - 0.5f) :
+               ((static_cast<int>(src_info.valid_region().anchor[0]) + 0.5f) * wr - 0.5f));
+    anchor.set(1, (policy == InterpolationPolicy::BILINEAR && border_undefined) ? ((static_cast<int>(src_info.valid_region().anchor[1]) + border_size.top + 0.5f) * hr - 0.5f) :
+               ((static_cast<int>(src_info.valid_region().anchor[1]) + 0.5f) * hr - 0.5f));
+    float shape_out_x = (policy == InterpolationPolicy::BILINEAR
+                         && border_undefined) ?
+                        ((static_cast<int>(src_info.valid_region().anchor[0]) + static_cast<int>(src_info.valid_region().shape[0]) - 1) - 1 + 0.5f) * wr - 0.5f :
+                        ((static_cast<int>(src_info.valid_region().anchor[0]) + static_cast<int>(src_info.valid_region().shape[0])) + 0.5f) * wr - 0.5f;
+    float shape_out_y = (policy == InterpolationPolicy::BILINEAR
+                         && border_undefined) ?
+                        ((static_cast<int>(src_info.valid_region().anchor[1]) + static_cast<int>(src_info.valid_region().shape[1]) - 1) - 1 + 0.5f) * hr - 0.5f :
+                        ((static_cast<int>(src_info.valid_region().anchor[1]) + static_cast<int>(src_info.valid_region().shape[1])) + 0.5f) * hr - 0.5f;
+
+    new_dst_shape.set(0, shape_out_x - anchor[0]);
+    new_dst_shape.set(1, shape_out_y - anchor[1]);
+
+    return ValidRegion(std::move(anchor), std::move(new_dst_shape));
+}
 } // namespace arm_compute
diff --git a/src/core/CL/cl_kernels/scale.cl b/src/core/CL/cl_kernels/scale.cl
index 9ef33b83ce..b3398bd11c 100644
--- a/src/core/CL/cl_kernels/scale.cl
+++ b/src/core/CL/cl_kernels/scale.cl
@@ -118,6 +118,6 @@ __kernel void scale_bilinear(
     Image        in  = CONVERT_TO_IMAGE_STRUCT_NO_STEP(in);
     Image        out = CONVERT_TO_IMAGE_STRUCT(out);
     const float2 r   = (float2)(input_width / output_width, input_height / output_height);
-    const float8 tc  = clamp_to_border(transform_bilinear(get_current_coords(), r), input_width, input_height);
+    const float8 tc  = transform_bilinear(get_current_coords(), r);
     vstore4(bilinear_interpolate(&in, tc, input_width, input_height), 0, (__global DATA_TYPE *)out.ptr);
 }
diff --git a/src/core/CL/kernels/CLScaleKernel.cpp b/src/core/CL/kernels/CLScaleKernel.cpp
index d74e837ace..23ce89aba2 100644
--- a/src/core/CL/kernels/CLScaleKernel.cpp
+++ b/src/core/CL/kernels/CLScaleKernel.cpp
@@ -76,17 +76,23 @@ void CLScaleKernel::configure(const ICLTensor *input, ICLTensor *output, Interpo
 
     // Configure kernel window
     constexpr unsigned int num_elems_processed_per_iteration = 4;
-    const int              border_offset                     = (border_undefined) ? 0 : border_size().left;
 
     Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration));
 
-    AccessWindowStatic input_access(input->info(), -border_offset, -border_offset,
-                                    input->info()->dimension(0) + border_offset, input->info()->dimension(1) + border_offset);
+    const ValidRegion &input_valid_region = input->info()->valid_region();
+
+    // Reads can occur within the valid region of the input
+    AccessWindowStatic input_access(input->info(),
+                                    input_valid_region.anchor[0] - border_size().left, input_valid_region.anchor[1] - border_size().top,
+                                    input_valid_region.anchor[0] + input_valid_region.shape[0] + border_size().right,
+                                    input_valid_region.anchor[1] + input_valid_region.shape[1] + border_size().bottom);
+
     AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration);
 
     update_window_and_padding(win, input_access, output_access);
 
-    output_access.set_valid_region(win, ValidRegion(Coordinates(), output->info()->tensor_shape()));
+    output_access.set_valid_region(win, calculate_valid_region_scale(*(input->info()), output->info()->tensor_shape(), policy, border_size(),
+                                                                     border_undefined));
 
     ICLKernel::configure(win);
 
diff --git a/src/core/NEON/kernels/NEScaleKernel.cpp b/src/core/NEON/kernels/NEScaleKernel.cpp
index fd2978de1c..ae164eb979 100644
--- a/src/core/NEON/kernels/NEScaleKernel.cpp
+++ b/src/core/NEON/kernels/NEScaleKernel.cpp
@@ -79,6 +79,16 @@ void NEScaleKernel::configure(const ITensor *input, const ITensor *dx, const ITe
     _dx      = dx;
     _dy      = dy;
 
+    /* Compute the ratio between source width/height and destination width/height */
+    const auto wr = static_cast<float>(input->info()->dimension(0)) / static_cast<float>(output->info()->dimension(0));
+    const auto hr = static_cast<float>(input->info()->dimension(1)) / static_cast<float>(output->info()->dimension(1));
+
+    /* Area interpolation behaves as Nearest Neighbour in case of up-sampling */
+    if(policy == InterpolationPolicy::AREA && wr <= 1.f && hr <= 1.f)
+    {
+        policy = InterpolationPolicy::NEAREST_NEIGHBOR;
+    }
+
     switch(policy)
     {
         case InterpolationPolicy::NEAREST_NEIGHBOR:
@@ -104,13 +114,18 @@ void NEScaleKernel::configure(const ITensor *input, const ITensor *dx, const ITe
     }
 
     constexpr unsigned int num_elems_processed_per_iteration = 16;
-    const int              border_offset                     = (border_undefined) ? 0 : border_size().left;
 
     // Configure kernel window
     Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration));
 
-    AccessWindowStatic     input_access(input->info(), -border_offset, -border_offset, input->info()->dimension(0) + border_offset, input->info()->dimension(1) + border_offset);
-    AccessWindowHorizontal offsets_access(offsets->info(), 0, num_elems_processed_per_iteration);
+    const ValidRegion &input_valid_region = input->info()->valid_region();
+
+    // Reads can occur within the valid region of the input
+    AccessWindowStatic input_access(input->info(),
+                                    input_valid_region.anchor[0] - border_size().left, input_valid_region.anchor[1] - border_size().top,
+                                    input_valid_region.anchor[0] + input_valid_region.shape[0] + border_size().right,
+                                    input_valid_region.anchor[1] + input_valid_region.shape[1] + border_size().bottom);
+    AccessWindowHorizontal offsets_access(offsets == nullptr ? nullptr : offsets->info(), 0, num_elems_processed_per_iteration);
     AccessWindowHorizontal dx_access(dx == nullptr ? nullptr : dx->info(), 0, num_elems_processed_per_iteration);
     AccessWindowHorizontal dy_access(dy == nullptr ? nullptr : dy->info(), 0, num_elems_processed_per_iteration);
     AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration);
@@ -122,8 +137,7 @@ void NEScaleKernel::configure(const ITensor *input, const ITensor *dx, const ITe
                               dy_access,
                               output_access);
 
-    output_access.set_valid_region(win, ValidRegion(Coordinates(), output->info()->tensor_shape()));
-
+    output_access.set_valid_region(win, calculate_valid_region_scale(*(input->info()), output->info()->tensor_shape(), policy, border_size(), border_undefined));
     INEKernel::configure(win);
 }
 
@@ -164,8 +178,8 @@ void NEScaleKernel::scale_nearest(const Window &window)
                 const auto           offsets_ptr = reinterpret_cast<const int32_t *>(offsets.ptr());
                 const uint8_t *const in_ptr      = in.ptr();
 
-                const size_t in_yi      = (id.y() + 0.5f) * hr;
-                const size_t offset_row = in_yi * input_stride;
+                const int in_yi      = std::floor((id.y() + 0.5f) * hr);
+                const int offset_row = in_yi * input_stride;
 
                 tmp = vsetq_lane_u8(in_ptr[offsets_ptr[0] + offset_row], tmp, 0);
                 tmp = vsetq_lane_u8(in_ptr[offsets_ptr[1] + offset_row], tmp, 1);
@@ -203,8 +217,8 @@ void NEScaleKernel::scale_nearest(const Window &window)
             {
                 const auto offsets_ptr = reinterpret_cast<const int32_t *>(offsets.ptr());
 
-                const size_t in_yi      = (id.y() + 0.5f) * hr;
-                const size_t offset_row = in_yi * input_stride;
+                const int in_yi      = (id.y() + 0.5f) * hr;
+                const int offset_row = in_yi * input_stride;
 
                 tmp.val[0] = vsetq_lane_s16(*reinterpret_cast<const int16_t *>(in.ptr() + offsets_ptr[0] + offset_row), tmp.val[0], 0);
                 tmp.val[0] = vsetq_lane_s16(*reinterpret_cast<const int16_t *>(in.ptr() + offsets_ptr[2] + offset_row), tmp.val[0], 1);
@@ -273,8 +287,8 @@ void NEScaleKernel::scale_bilinear(const Window &window)
         const auto dy_ptr      = reinterpret_cast<const float *>(dy.ptr());
         const auto in_ptr      = reinterpret_cast<const uint8_t *>(in.ptr());
 
-        const size_t in_yi      = std::floor((id.y() + 0.5f) * hr - 0.5f);
-        const size_t offset_row = in_yi * in_stride;
+        const int in_yi      = std::floor((id.y() + 0.5f) * hr - 0.5f);
+        const int offset_row = in_yi * in_stride;
 
         uint8x8_t tmp0 = vdup_n_u8(0);
         tmp0           = vset_lane_u8(delta_bilinear_c1u8(&in_ptr[offsets_ptr[0] + offset_row], in_stride, dx_ptr[0], dy_ptr[0]), tmp0, 0);
diff --git a/tests/TypePrinter.h b/tests/TypePrinter.h
index 49e717a48a..394ee9d87c 100644
--- a/tests/TypePrinter.h
+++ b/tests/TypePrinter.h
@@ -121,6 +121,13 @@ inline ::std::ostream &operator<<(::std::ostream &os, const BorderMode &mode)
     return os;
 }
 
+inline std::string to_string(const BorderMode &mode)
+{
+    std::stringstream str;
+    str << mode;
+    return str.str();
+}
+
 /** Formatted output of the NonLinearFilterFunction type. */
 inline ::std::ostream &operator<<(::std::ostream &os, const NonLinearFilterFunction &function)
 {
@@ -187,6 +194,13 @@ inline ::std::ostream &operator<<(::std::ostream &os, const InterpolationPolicy
     return os;
 }
 
+inline std::string to_string(const InterpolationPolicy &policy)
+{
+    std::stringstream str;
+    str << policy;
+    return str.str();
+}
+
 /** Formatted output of the ConversionPolicy type. */
 inline ::std::ostream &operator<<(::std::ostream &os, const ConvertPolicy &policy)
 {
diff --git a/tests/datasets_new/BorderModeDataset.h b/tests/datasets_new/BorderModeDataset.h
new file mode 100644
index 0000000000..b45e7835c1
--- /dev/null
+++ b/tests/datasets_new/BorderModeDataset.h
@@ -0,0 +1,51 @@
+/*
+ * 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_BORDER_MODE_DATASET_H__
+#define __ARM_COMPUTE_TEST_BORDER_MODE_DATASET_H__
+
+#include "arm_compute/core/Types.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace datasets
+{
+class BorderModes final : public framework::dataset::ContainerDataset<std::vector<BorderMode>>
+{
+public:
+    BorderModes()
+        : ContainerDataset("BorderMode",
+    {
+        BorderMode::UNDEFINED,
+                   BorderMode::CONSTANT,
+                   BorderMode::REPLICATE
+    })
+    {
+    }
+};
+} // namespace datasets
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_BORDER_MODE_DATASET_H__ */
diff --git a/tests/datasets_new/InterpolationPolicyDataset.h b/tests/datasets_new/InterpolationPolicyDataset.h
new file mode 100644
index 0000000000..154f887962
--- /dev/null
+++ b/tests/datasets_new/InterpolationPolicyDataset.h
@@ -0,0 +1,51 @@
+/*
+ * 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_INTEPOLATIONPOLICY_DATASET_H__
+#define __ARM_COMPUTE_TEST_INTEPOLATIONPOLICY_DATASET_H__
+
+#include "arm_compute/core/Types.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace datasets
+{
+class InterpolationPolicies final : public framework::dataset::ContainerDataset<std::vector<InterpolationPolicy>>
+{
+public:
+    InterpolationPolicies()
+        : ContainerDataset("InterpolationPolicy",
+    {
+        InterpolationPolicy::NEAREST_NEIGHBOR,
+                            InterpolationPolicy::BILINEAR,
+                            InterpolationPolicy::AREA
+    })
+    {
+    }
+};
+} // namespace datasets
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_INTEPOLATIONPOLICY_DATASET_H__ */
diff --git a/tests/validation/TensorOperations.h b/tests/validation/TensorOperations.h
index e68a344112..f157671c18 100644
--- a/tests/validation/TensorOperations.h
+++ b/tests/validation/TensorOperations.h
@@ -24,6 +24,8 @@
 #ifndef __ARM_COMPUTE_TEST_TENSOR_OPERATIONS_H__
 #define __ARM_COMPUTE_TEST_TENSOR_OPERATIONS_H__
 
+#include "arm_compute/core/FixedPoint.h"
+#include "arm_compute/core/Helpers.h"
 #include "arm_compute/core/Types.h"
 #include "support/ToolchainSupport.h"
 #include "tests/Types.h"
@@ -116,6 +118,46 @@ void apply_2d_spatial_filter(Coordinates coord, const Tensor<T1> &in, Tensor<T3>
 }
 } // namespace
 
+template <typename T>
+T bilinear_policy(const Tensor<T> &in, Coordinates id, float xn, float yn, BorderMode border_mode, uint8_t constant_border_value)
+{
+    int idx = std::floor(xn);
+    int idy = std::floor(yn);
+
+    const float dx   = xn - idx;
+    const float dy   = yn - idy;
+    const float dx_1 = 1.0f - dx;
+    const float dy_1 = 1.0f - dy;
+
+    id.set(0, idx);
+    id.set(1, idy);
+    const T tl = tensor_elem_at(in, id, border_mode, constant_border_value);
+    id.set(0, idx + 1);
+    id.set(1, idy);
+    const T tr = tensor_elem_at(in, id, border_mode, constant_border_value);
+    id.set(0, idx);
+    id.set(1, idy + 1);
+    const T bl = tensor_elem_at(in, id, border_mode, constant_border_value);
+    id.set(0, idx + 1);
+    id.set(1, idy + 1);
+    const T br = tensor_elem_at(in, id, border_mode, constant_border_value);
+
+    return tl * (dx_1 * dy_1) + tr * (dx * dy_1) + bl * (dx_1 * dy) + br * (dx * dy);
+}
+
+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;
+}
+
 // Sobel 3x3
 template <typename T1, typename T2>
 void sobel_3x3(Tensor<T1> &in, Tensor<T2> &out_x, Tensor<T2> &out_y, BorderMode border_mode, uint8_t constant_border_value)
@@ -881,46 +923,6 @@ void threshold(const Tensor<T> &in, Tensor<T> &out, uint8_t threshold, uint8_t f
     }
 }
 
-template <typename T>
-T bilinear_policy(const Tensor<T> &in, Coordinates id, float xn, float yn, BorderMode border_mode, uint8_t constant_border_value)
-{
-    int idx = std::floor(xn);
-    int idy = std::floor(yn);
-
-    const float dx   = xn - idx;
-    const float dy   = yn - idy;
-    const float dx_1 = 1.0f - dx;
-    const float dy_1 = 1.0f - dy;
-
-    id.set(0, idx);
-    id.set(1, idy);
-    const T tl = tensor_elem_at(in, id, border_mode, constant_border_value);
-    id.set(0, idx + 1);
-    id.set(1, idy);
-    const T tr = tensor_elem_at(in, id, border_mode, constant_border_value);
-    id.set(0, idx);
-    id.set(1, idy + 1);
-    const T bl = tensor_elem_at(in, id, border_mode, constant_border_value);
-    id.set(0, idx + 1);
-    id.set(1, idy + 1);
-    const T br = tensor_elem_at(in, id, border_mode, constant_border_value);
-
-    return tl * (dx_1 * dy_1) + tr * (dx * dy_1) + bl * (dx_1 * dy) + br * (dx * dy);
-}
-
-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;
-}
-
 // Warp Perspective
 template <typename T>
 void warp_perspective(const Tensor<T> &in, Tensor<T> &out, Tensor<T> &valid_mask, const float *matrix, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
diff --git a/tests/validation_new/CL/Scale.cpp b/tests/validation_new/CL/Scale.cpp
new file mode 100644
index 0000000000..e28555ac21
--- /dev/null
+++ b/tests/validation_new/CL/Scale.cpp
@@ -0,0 +1,129 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/functions/CLScale.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "framework/Asserts.h"
+#include "framework/Macros.h"
+#include "framework/datasets/Datasets.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets_new/BorderModeDataset.h"
+#include "tests/datasets_new/ShapeDatasets.h"
+#include "tests/validation_new/Helpers.h"
+#include "tests/validation_new/Validation.h"
+#include "tests/validation_new/fixtures/ScaleFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Tolerance */
+constexpr AbsoluteTolerance<uint8_t> tolerance(1);
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(Scale)
+
+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)
+{
+    std::mt19937                           generator(library->seed());
+    std::uniform_real_distribution<float>  distribution_float(0.25, 2);
+    const float                            scale_x = distribution_float(generator);
+    const float                            scale_y = distribution_float(generator);
+    std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+    uint8_t                                constant_border_value = distribution_u8(generator);
+
+    // Create tensors
+    CLTensor    src = create_tensor<CLTensor>(shape, data_type);
+    TensorShape shape_scaled(shape);
+    shape_scaled.set(0, shape[0] * scale_x);
+    shape_scaled.set(1, shape[1] * scale_y);
+    CLTensor dst = create_tensor<CLTensor>(shape_scaled, 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
+    CLScale clscale;
+    clscale.configure(&src, &dst, policy, border_mode, constant_border_value);
+
+    // Validate valid region
+    const ValidRegion dst_valid_region = calculate_valid_region_scale(*(src.info()), shape_scaled, policy, BorderSize(1), (border_mode == BorderMode::UNDEFINED));
+
+    validate(dst.info()->valid_region(), dst_valid_region);
+
+    // Validate padding
+    PaddingCalculator calculator(shape_scaled.x(), 4);
+    calculator.set_border_mode(border_mode);
+
+    const PaddingSize read_padding(1);
+    const PaddingSize write_padding = calculator.required_padding(PaddingCalculator::Option::EXCLUDE_BORDER);
+    validate(src.info()->padding(), read_padding);
+    validate(dst.info()->padding(), write_padding);
+}
+
+template <typename T>
+using CLScaleFixture = ScaleValidationFixture<CLTensor, CLAccessor, CLScale, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, CLScaleFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType",
+                                                                                                                     DataType::U8)),
+                                                                                                             framework::dataset::make("InterpolationPolicy",
+{ InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+datasets::BorderModes()))
+{
+    //Create valid region
+    TensorInfo        src_info(_shape, 1, _data_type);
+    const ValidRegion valid_region = calculate_valid_region_scale(src_info, _reference.shape(), _policy, BorderSize(1), (_border_mode == BorderMode::UNDEFINED));
+
+    // Validate output
+    validate(CLAccessor(_target), _reference, valid_region, tolerance);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, CLScaleFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType",
+                                                                                                                   DataType::U8)),
+                                                                                                           framework::dataset::make("InterpolationPolicy",
+{ InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+datasets::BorderModes()))
+{
+    //Create valid region
+    TensorInfo        src_info(_shape, 1, _data_type);
+    const ValidRegion valid_region = calculate_valid_region_scale(src_info, _reference.shape(), _policy, BorderSize(1), (_border_mode == BorderMode::UNDEFINED));
+
+    // Validate output
+    validate(CLAccessor(_target), _reference, valid_region, tolerance);
+}
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation_new/CPP/DepthwiseConvolution.cpp b/tests/validation_new/CPP/DepthwiseConvolution.cpp
index be18ffb911..8c5cec596e 100644
--- a/tests/validation_new/CPP/DepthwiseConvolution.cpp
+++ b/tests/validation_new/CPP/DepthwiseConvolution.cpp
@@ -24,7 +24,7 @@
 #include "DepthwiseConvolution.h"
 
 #include "ConvolutionLayer.h"
-#include "TensorElementAt.h"
+#include "Utils.h"
 
 #include "tests/validation_new/Helpers.h"
 #include "tests/validation_new/half.h"
diff --git a/tests/validation_new/CPP/DepthwiseSeparableConvolutionLayer.cpp b/tests/validation_new/CPP/DepthwiseSeparableConvolutionLayer.cpp
index 88584275cf..eba0a19189 100644
--- a/tests/validation_new/CPP/DepthwiseSeparableConvolutionLayer.cpp
+++ b/tests/validation_new/CPP/DepthwiseSeparableConvolutionLayer.cpp
@@ -26,7 +26,7 @@
 #include "DepthwiseSeparableConvolutionLayer.h"
 
 #include "ConvolutionLayer.h"
-#include "TensorElementAt.h"
+#include "Utils.h"
 
 #include "tests/validation_new/Helpers.h"
 #include "tests/validation_new/half.h"
diff --git a/tests/validation_new/CPP/Scale.cpp b/tests/validation_new/CPP/Scale.cpp
new file mode 100644
index 0000000000..a1119f33b9
--- /dev/null
+++ b/tests/validation_new/CPP/Scale.cpp
@@ -0,0 +1,166 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+
+#include "arm_compute/core/Helpers.h"
+
+#include "Scale.h"
+#include "Utils.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+SimpleTensor<T> scale(const SimpleTensor<T> &in, float scale_x, float scale_y, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
+{
+    TensorShape shape_scaled(in.shape());
+    shape_scaled.set(0, in.shape()[0] * scale_x);
+    shape_scaled.set(1, in.shape()[1] * scale_y);
+    SimpleTensor<T> out(shape_scaled, in.data_type());
+
+    // Compute the ratio between source width/height and destination width/height
+    const auto wr = static_cast<float>(in.shape()[0]) / static_cast<float>(out.shape()[0]);
+    const auto hr = static_cast<float>(in.shape()[1]) / static_cast<float>(out.shape()[1]);
+
+    const auto width  = static_cast<int>(in.shape().x());
+    const auto height = static_cast<int>(in.shape().y());
+
+    // Area interpolation behaves as Nearest Neighbour in case of up-sampling
+    if(policy == InterpolationPolicy::AREA && wr <= 1.f && hr <= 1.f)
+    {
+        policy = InterpolationPolicy::NEAREST_NEIGHBOR;
+    }
+
+    for(int element_idx = 0, count = 0; element_idx < out.num_elements(); ++element_idx, ++count)
+    {
+        Coordinates id    = index2coord(out.shape(), element_idx);
+        int         idx   = id.x();
+        int         idy   = id.y();
+        float       x_src = (idx + 0.5f) * wr - 0.5f;
+        float       y_src = (idy + 0.5f) * hr - 0.5f;
+
+        switch(policy)
+        {
+            case InterpolationPolicy::NEAREST_NEIGHBOR:
+            {
+                //Calculate the source coords without -0.5f is equivalent to round the x_scr/y_src coords
+                x_src = (idx + 0.5f) * wr;
+                y_src = (idy + 0.5f) * hr;
+                id.set(0, x_src);
+                id.set(1, y_src);
+
+                // If coordinates in range of tensor's width or height
+                if(x_src >= -1 || y_src >= -1 || x_src <= width || y_src <= height)
+                {
+                    out[element_idx] = tensor_elem_at(in, id, border_mode, constant_border_value);
+                }
+                else
+                {
+                    if(border_mode == BorderMode::CONSTANT)
+                    {
+                        out[element_idx] = constant_border_value;
+                    }
+                    else if(border_mode == BorderMode::REPLICATE)
+                    {
+                        id.set(0, clamp(static_cast<int>(x_src), 0, width - 1));
+                        id.set(1, clamp(static_cast<int>(y_src), 0, height - 1));
+                        out[element_idx] = in[coord2index(in.shape(), id)];
+                    }
+                }
+                break;
+            }
+            case InterpolationPolicy::BILINEAR:
+            {
+                id.set(0, std::floor(x_src));
+                id.set(1, std::floor(y_src));
+                if(x_src >= -1 || y_src >= -1 || x_src <= width || y_src <= height)
+                {
+                    out[element_idx] = bilinear_policy(in, id, x_src, y_src, border_mode, constant_border_value);
+                }
+                else
+                {
+                    if(border_mode == BorderMode::CONSTANT)
+                    {
+                        out[element_idx] = constant_border_value;
+                    }
+                    else if(border_mode == BorderMode::REPLICATE)
+                    {
+                        id.set(0, clamp(static_cast<int>(x_src), 0, width - 1));
+                        id.set(1, clamp(static_cast<int>(y_src), 0, height - 1));
+                        out[element_idx] = in[coord2index(in.shape(), id)];
+                    }
+                }
+                break;
+            }
+            case InterpolationPolicy::AREA:
+            {
+                int       x_from = std::floor(idx * wr - 0.5f - x_src);
+                int       y_from = std::floor(idy * hr - 0.5f - y_src);
+                int       x_to   = std::ceil((idx + 1) * wr - 0.5f - x_src);
+                int       y_to   = std::ceil((idy + 1) * hr - 0.5f - y_src);
+                const int xi     = std::floor(x_src);
+                const int yi     = std::floor(y_src);
+
+                // Clamp position to borders
+                x_src = std::max(-1.f, std::min(x_src, static_cast<float>(width)));
+                y_src = std::max(-1.f, std::min(y_src, static_cast<float>(height)));
+
+                // Clamp bounding box offsets to borders
+                x_from = ((x_src + x_from) < -1) ? -1 : x_from;
+                y_from = ((y_src + y_from) < -1) ? -1 : y_from;
+                x_to   = ((x_src + x_to) > width) ? (width - x_src) : x_to;
+                y_to   = ((y_src + y_to) > height) ? (height - y_src) : y_to;
+                ARM_COMPUTE_ERROR_ON((x_to - x_from + 1) == 0 || (y_to - y_from + 1) == 0);
+
+                float sum = 0;
+                for(int j = yi + y_from, je = yi + y_to; j <= je; ++j)
+                {
+                    for(int i = xi + x_from, ie = xi + x_to; i <= ie; ++i)
+                    {
+                        id.set(0, static_cast<int>(i));
+                        id.set(1, static_cast<int>(j));
+                        sum += tensor_elem_at(in, id, border_mode, constant_border_value);
+                    }
+                }
+                out[element_idx] = sum / ((x_to - x_from + 1) * (y_to - y_from + 1));
+
+                break;
+            }
+            default:
+                ARM_COMPUTE_ERROR("Unsupported interpolation mode");
+        }
+    }
+
+    return out;
+}
+
+template SimpleTensor<uint8_t> scale(const SimpleTensor<uint8_t> &src, float scale_x, float scale_y, 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_new/CPP/TensorElementAt.h b/tests/validation_new/CPP/Scale.h
index d0d68ab6af..b882915946 100644
--- a/tests/validation_new/CPP/TensorElementAt.h
+++ b/tests/validation_new/CPP/Scale.h
@@ -21,11 +21,10 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  * SOFTWARE.
  */
-#ifndef __ARM_COMPUTE_TEST_TENSOR_ELEMENT_AT_H__
-#define __ARM_COMPUTE_TEST_TENSOR_ELEMENT_AT_H__
+#ifndef __ARM_COMPUTE_TEST_SCALE_H__
+#define __ARM_COMPUTE_TEST_SCALE_H__
 
 #include "tests/SimpleTensor.h"
-#include "tests/validation_new/Helpers.h"
 
 namespace arm_compute
 {
@@ -36,9 +35,9 @@ namespace validation
 namespace reference
 {
 template <typename T>
-T tensor_elem_at(const SimpleTensor<T> &in, Coordinates coord, BorderMode border_mode, T constant_border_value);
+SimpleTensor<T> scale(const SimpleTensor<T> &in, float scale_x, float scale_y, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value = 0);
 } // namespace reference
 } // namespace validation
 } // namespace test
 } // namespace arm_compute
-#endif /* __ARM_COMPUTE_TEST_TENSOR_ELEMENT_AT_H__ */
+#endif /* __ARM_COMPUTE_TEST_SCALE_H__ */
diff --git a/tests/validation_new/CPP/TensorElementAt.cpp b/tests/validation_new/CPP/Utils.cpp
index ca4cca9650..c89807b69a 100644
--- a/tests/validation_new/CPP/TensorElementAt.cpp
+++ b/tests/validation_new/CPP/Utils.cpp
@@ -21,7 +21,7 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  * SOFTWARE.
  */
-#include "TensorElementAt.h"
+#include "Utils.h"
 
 #include "tests/validation_new/Helpers.h"
 #include "tests/validation_new/half.h"
@@ -32,8 +32,6 @@ namespace test
 {
 namespace validation
 {
-namespace reference
-{
 // Return a tensor element at a specified coordinate with different border modes
 template <typename T>
 T tensor_elem_at(const SimpleTensor<T> &in, Coordinates coord, BorderMode border_mode, T constant_border_value)
@@ -50,21 +48,46 @@ T tensor_elem_at(const SimpleTensor<T> &in, Coordinates coord, BorderMode border
         {
             coord.set(0, std::max(0, std::min(x, width - 1)));
             coord.set(1, std::max(0, std::min(y, height - 1)));
-
-            return in[coord2index(in.shape(), coord)];
         }
         else
         {
             return constant_border_value;
         }
     }
-    else
-    {
-        return in[coord2index(in.shape(), coord)];
-    }
+    return in[coord2index(in.shape(), coord)];
 }
 template float tensor_elem_at(const SimpleTensor<float> &in, Coordinates coord, BorderMode border_mode, float constant_border_value);
-} // namespace reference
+template uint8_t tensor_elem_at(const SimpleTensor<uint8_t> &in, Coordinates coord, BorderMode border_mode, uint8_t constant_border_value);
+
+// Return the bilinear value at a specified coordinate with different border modes
+template <typename T>
+T bilinear_policy(const SimpleTensor<T> &in, Coordinates id, float xn, float yn, BorderMode border_mode, uint8_t constant_border_value)
+{
+    int idx = std::floor(xn);
+    int idy = std::floor(yn);
+
+    const float dx   = xn - idx;
+    const float dy   = yn - idy;
+    const float dx_1 = 1.0f - dx;
+    const float dy_1 = 1.0f - dy;
+
+    id.set(0, idx);
+    id.set(1, idy);
+    const T tl = tensor_elem_at(in, id, border_mode, constant_border_value);
+    id.set(0, idx + 1);
+    id.set(1, idy);
+    const T tr = tensor_elem_at(in, id, border_mode, constant_border_value);
+    id.set(0, idx);
+    id.set(1, idy + 1);
+    const T bl = tensor_elem_at(in, id, border_mode, constant_border_value);
+    id.set(0, idx + 1);
+    id.set(1, idy + 1);
+    const T br = tensor_elem_at(in, id, border_mode, constant_border_value);
+
+    return tl * (dx_1 * dy_1) + tr * (dx * dy_1) + bl * (dx_1 * dy) + br * (dx * dy);
+}
+template uint8_t bilinear_policy(const SimpleTensor<uint8_t> &in, Coordinates id, float xn, float yn, BorderMode border_mode, uint8_t constant_border_value);
+
 } // namespace validation
 } // namespace test
 } // namespace arm_compute
diff --git a/tests/validation_new/CPP/Utils.h b/tests/validation_new/CPP/Utils.h
new file mode 100644
index 0000000000..4e3deb4d86
--- /dev/null
+++ b/tests/validation_new/CPP/Utils.h
@@ -0,0 +1,54 @@
+/*
+ * 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_VALIDATION_UTILS_H__
+#define __ARM_COMPUTE_TEST_VALIDATION_UTILS_H__
+
+#include "arm_compute/core/Types.h"
+#include "tests/Globals.h"
+#include "tests/ILutAccessor.h"
+#include "tests/Types.h"
+#include "tests/validation/ValidationUserConfiguration.h"
+#include "tests/validation/half.h"
+
+#include <array>
+#include <random>
+#include <type_traits>
+#include <utility>
+#include <vector>
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+template <typename T>
+T tensor_elem_at(const SimpleTensor<T> &in, Coordinates coord, BorderMode border_mode, T constant_border_value);
+
+template <typename T>
+T bilinear_policy(const SimpleTensor<T> &in, Coordinates id, float xn, float yn, BorderMode border_mode, uint8_t constant_border_value);
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_VALIDATION_UTILS_H__ */
diff --git a/tests/validation_new/NEON/Scale.cpp b/tests/validation_new/NEON/Scale.cpp
new file mode 100644
index 0000000000..1036f98c23
--- /dev/null
+++ b/tests/validation_new/NEON/Scale.cpp
@@ -0,0 +1,127 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEScale.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "framework/Asserts.h"
+#include "framework/Macros.h"
+#include "framework/datasets/Datasets.h"
+#include "tests/NEON/Accessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets_new/BorderModeDataset.h"
+#include "tests/datasets_new/InterpolationPolicyDataset.h"
+#include "tests/datasets_new/ShapeDatasets.h"
+#include "tests/validation_new/Helpers.h"
+#include "tests/validation_new/Validation.h"
+#include "tests/validation_new/fixtures/ScaleFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(NEON)
+TEST_SUITE(Scale)
+
+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)
+{
+    std::mt19937                          generator(library->seed());
+    std::uniform_real_distribution<float> distribution_float(0.25, 2);
+    const float                           scale_x               = distribution_float(generator);
+    const float                           scale_y               = distribution_float(generator);
+    uint8_t                               constant_border_value = 0;
+    if(border_mode == BorderMode::CONSTANT)
+    {
+        std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+        constant_border_value = distribution_u8(generator);
+    }
+
+    // Create tensors
+    Tensor      src = create_tensor<Tensor>(shape, data_type);
+    TensorShape shape_scaled(shape);
+    shape_scaled.set(0, shape[0] * scale_x);
+    shape_scaled.set(1, shape[1] * scale_y);
+    Tensor dst = create_tensor<Tensor>(shape_scaled, 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
+    NEScale nescale;
+    nescale.configure(&src, &dst, policy, border_mode, constant_border_value);
+
+    // Validate valid region
+    const ValidRegion dst_valid_region = calculate_valid_region_scale(*(src.info()), shape_scaled, policy, BorderSize(1), (border_mode == BorderMode::UNDEFINED));
+
+    validate(dst.info()->valid_region(), dst_valid_region);
+
+    // Validate padding
+    PaddingCalculator calculator(shape_scaled.x(), 16);
+    calculator.set_border_mode(border_mode);
+
+    const PaddingSize read_padding(1);
+    const PaddingSize write_padding = calculator.required_padding(PaddingCalculator::Option::EXCLUDE_BORDER);
+    validate(src.info()->padding(), read_padding);
+    validate(dst.info()->padding(), write_padding);
+}
+
+template <typename T>
+using NEScaleFixture = ScaleValidationFixture<Tensor, Accessor, NEScale, T>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, NEScaleFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType",
+                                                                                                                     DataType::U8)),
+                                                                                                             framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+                                                                                                     datasets::BorderModes()))
+{
+    //Create valid region
+    TensorInfo  src_info(_shape, 1, _data_type);
+    ValidRegion valid_region = calculate_valid_region_scale(src_info, _reference.shape(), _policy, BorderSize(1), (_border_mode == BorderMode::UNDEFINED));
+
+    // Validate output
+    validate(Accessor(_target), _reference, valid_region);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, NEScaleFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType",
+                                                                                                                   DataType::U8)),
+                                                                                                           framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+                                                                                                   datasets::BorderModes()))
+{
+    //Create valid region
+    TensorInfo  src_info(_shape, 1, _data_type);
+    ValidRegion valid_region = calculate_valid_region_scale(src_info, _reference.shape(), _policy, BorderSize(1), (_border_mode == BorderMode::UNDEFINED));
+
+    // Validate output
+    validate(Accessor(_target), _reference, valid_region);
+}
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation_new/fixtures/ScaleFixture.h b/tests/validation_new/fixtures/ScaleFixture.h
new file mode 100644
index 0000000000..74dc0d6816
--- /dev/null
+++ b/tests/validation_new/fixtures/ScaleFixture.h
@@ -0,0 +1,127 @@
+/*
+ * 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_SCALE_FIXTURE
+#define ARM_COMPUTE_TEST_SCALE_FIXTURE
+
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+#include "framework/Asserts.h"
+#include "framework/Fixture.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/IAccessor.h"
+#include "tests/validation_new/CPP/Scale.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+template <typename TensorType, typename AccessorType, typename FunctionType, typename T>
+class ScaleValidationFixture : public framework::Fixture
+{
+public:
+    template <typename...>
+    void setup(TensorShape shape, DataType data_type, InterpolationPolicy policy, BorderMode border_mode)
+    {
+        _shape       = shape;
+        _policy      = policy;
+        _border_mode = border_mode;
+        _data_type   = data_type;
+
+        std::mt19937                           generator(library->seed());
+        std::uniform_real_distribution<float>  distribution_float(0.25, 4);
+        const float                            scale_x = distribution_float(generator);
+        const float                            scale_y = distribution_float(generator);
+        std::uniform_int_distribution<uint8_t> distribution_u8(0, 255);
+        uint8_t                                constant_border_value = distribution_u8(generator);
+
+        _target    = compute_target(shape, scale_x, scale_y, policy, border_mode, constant_border_value);
+        _reference = compute_reference(shape, scale_x, scale_y, 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, const float scale_x, const float scale_y,
+                              InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
+    {
+        // Create tensors
+        TensorType  src = create_tensor<TensorType>(shape, _data_type);
+        TensorShape shape_scaled(shape);
+        shape_scaled.set(0, shape[0] * scale_x);
+        shape_scaled.set(1, shape[1] * scale_y);
+        TensorType dst = create_tensor<TensorType>(shape_scaled, _data_type);
+
+        // Create and configure function
+        FunctionType scale;
+
+        scale.configure(&src, &dst, 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
+        scale.run();
+
+        return dst;
+    }
+
+    SimpleTensor<T> compute_reference(const TensorShape &shape, const float scale_x, const float scale_y,
+                                      InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
+    {
+        // Create reference
+        SimpleTensor<T> src{ shape, _data_type };
+
+        // Fill reference
+        fill(src);
+
+        return reference::scale<T>(src, scale_x, scale_y, policy, border_mode, constant_border_value);
+    }
+
+    TensorType          _target{};
+    SimpleTensor<T>     _reference{};
+    TensorShape         _shape{};
+    InterpolationPolicy _policy{};
+    BorderMode          _border_mode{};
+    DataType            _data_type{};
+};
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_SCALE_FIXTURE */
\ No newline at end of file