Add FP16 support to CLRemap

Add FP16 support to CLRemap when data layout is NHWC.
Add relevant tests for FP16 and validation.
Update NERemap function level to be consistent with CLRemap.
Add depreciation notice for uint_8 only function level methods.

Resolves: COMPMID-4335

Signed-off-by: Freddie Liardet <frederick.liardet@arm.com>
Change-Id: If05f06801aef7a169b73ff1ebe760a42f11ca05c
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5816
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>

10 files changed, 227 insertions, 54 deletions

diff --git a/arm_compute/runtime/CL/functions/CLRemap.h b/arm_compute/runtime/CL/functions/CLRemap.h
index f69b045c9b..8322f7a8a9 100644
--- a/arm_compute/runtime/CL/functions/CLRemap.h
+++ b/arm_compute/runtime/CL/functions/CLRemap.h
@@ -24,6 +24,7 @@
 #ifndef ARM_COMPUTE_CLREMAP_H
 #define ARM_COMPUTE_CLREMAP_H
 
+#include "arm_compute/core/PixelValue.h"
 #include "arm_compute/core/Types.h"
 #include "arm_compute/runtime/CL/ICLSimpleFunction.h"
 
@@ -33,6 +34,7 @@ namespace arm_compute
 {
 class CLCompileContext;
 class ICLTensor;
+class ITensorInfo;
 
 /** Basic function to execute remap. This function calls the following OpenCL kernels:
  *
@@ -51,8 +53,10 @@ public:
      * Valid data type configurations:
      * |src0   |src1   |src2   |dst    |
      * |:------|:------|:------|:------|
-     * |U8     |F32    |F32    |U 8    |
+     * |U8     |F32    |F32    |U8     |
+     * |F16    |F32    |F32    |F16    |
      *
+     * @param[in]     compile_context       The compile context to be used.
      * @param[in,out] input                 Source tensor. Data types supported: U8. (Written to only for @p border_mode != UNDEFINED)
      * @param[in]     map_x                 Map for X coords. Data types supported: F32.
      * @param[in]     map_y                 Map for Y coords. Data types supported: F32.
@@ -61,23 +65,58 @@ public:
      * @param[in]     border_mode           Border mode to use on the input tensor. Only CONSTANT and UNDEFINED are supported.
      * @param[in]     constant_border_value (Optional) Constant value to use for borders if border_mode is set to CONSTANT.
      *
+     * @deprecated This function is deprecated and is intended to be removed in 22.02 release
+     *
+     */
+    ARM_COMPUTE_DEPRECATED_REL(22.02)
+    void configure(const CLCompileContext &compile_context, ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output,
+                   InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value = 0);
+
+    /** Initialise the function's sources, destination, interpolation policy and border mode.
+     *
+     *  Similar to @ref CLRemap::configure(const CLCompileContext &compile_context, ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output,
+                                           InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
+     *
+     * @deprecated This function is deprecated and is intended to be removed in 22.02 release
+     *
      */
+    ARM_COMPUTE_DEPRECATED_REL(22.02)
     void configure(ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output,
                    InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value = 0);
+
     /** Initialise the function's sources, destination, interpolation policy and border mode.
      *
      * @param[in]     compile_context       The compile context to be used.
-     * @param[in,out] input                 Source tensor. Data types supported: U8. (Written to only for @p border_mode != UNDEFINED)
+     * @param[in,out] input                 Source tensor. Data types supported: U8,(or F16 when layout is NHWC). (Written to only for @p border_mode != UNDEFINED)
      * @param[in]     map_x                 Map for X coords. Data types supported: F32.
      * @param[in]     map_y                 Map for Y coords. Data types supported: F32.
-     * @param[out]    output                Output tensor. Data types supported: U8.
+     * @param[out]    output                Output tensor. Data types supported: Same as @p input.
      * @param[in]     policy                Interpolation policy to use. Only NEAREST and BILINEAR are supported.
      * @param[in]     border_mode           Border mode to use on the input tensor. Only CONSTANT and UNDEFINED are supported.
      * @param[in]     constant_border_value (Optional) Constant value to use for borders if border_mode is set to CONSTANT.
      *
      */
     void configure(const CLCompileContext &compile_context, ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output,
-                   InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value = 0);
+                   InterpolationPolicy policy, BorderMode border_mode, PixelValue constant_border_value = PixelValue{});
+
+    /** Initialise the function's sources, destination, interpolation policy and border mode.
+     *
+     * Similar to @ref CLRemap::configure(const CLCompileContext &compile_context, ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output,
+                                          InterpolationPolicy policy, BorderMode border_mode, PixelValue constant_border_value)
+     *
+     */
+    void configure(ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output,
+                   InterpolationPolicy policy, BorderMode border_mode, PixelValue constant_border_value = PixelValue{});
+
+    /** Checks if the kernel's input, output and border mode will lead to a valid configuration of @ref CLRemap
+     *
+     * Similar to @ref CLRemap::configure(const CLCompileContext &compile_context, ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output,
+                                          InterpolationPolicy policy, BorderMode border_mode, PixelValue constant_border_value)
+     *
+     * @return a Status
+     */
+    static Status validate(const ITensorInfo *input, const ITensorInfo *map_x, const ITensorInfo *map_y, const ITensorInfo *output,
+                           const InterpolationPolicy policy, const BorderMode border_mode, PixelValue constant_border_value = PixelValue{});
 };
 }
 #endif /*ARM_COMPUTE_CLREMAP_H */
diff --git a/arm_compute/runtime/NEON/functions/NERemap.h b/arm_compute/runtime/NEON/functions/NERemap.h
index 271ac9739b..140c20bd2f 100644
--- a/arm_compute/runtime/NEON/functions/NERemap.h
+++ b/arm_compute/runtime/NEON/functions/NERemap.h
@@ -59,9 +59,25 @@ public:
      * @param[in]      border_mode           Border mode to use on the input tensor.
      * @param[in]      constant_border_value (Optional) Constant value to use for borders if border_mode is set to CONSTANT. Defaults to 0.
      *
+     * @deprecated This function is deprecated and is intended to be removed in 22.02 release
+     *
      */
+    ARM_COMPUTE_DEPRECATED_REL(22.02)
     void configure(ITensor *input, const ITensor *map_x, const ITensor *map_y, ITensor *output,
                    InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value = 0);
+    /** Initialise the function's sources, destination, interpolation policy and border mode.
+     *
+     * @param[in, out] input                 Source tensor. Data type supported: U8. (Written to only for @p border_mode != UNDEFINED)
+     * @param[in]      map_x                 Map for X coordinates. Data type supported: F32.
+     * @param[in]      map_y                 Map for Y coordinates. Data type supported: F32.
+     * @param[out]     output                Output tensor. Data type supported: U8.
+     * @param[in]      policy                Interpolation policy to use. Only NEAREST and BILINEAR are supported.
+     * @param[in]      border_mode           Border mode to use on the input tensor.
+     * @param[in]      constant_border_value Constant value to use for borders if border_mode is set to CONSTANT.
+     *
+     */
+    void configure(ITensor *input, const ITensor *map_x, const ITensor *map_y, ITensor *output,
+                   InterpolationPolicy policy, BorderMode border_mode, PixelValue constant_border_value);
 };
 }
 #endif /*ARM_COMPUTE_NEREMAP_H */
diff --git a/src/core/CL/cl_kernels/remap.cl b/src/core/CL/cl_kernels/remap.cl
index 8ea4e84e96..cb67c2df1e 100644
--- a/src/core/CL/cl_kernels/remap.cl
+++ b/src/core/CL/cl_kernels/remap.cl
@@ -24,6 +24,7 @@
 #include "helpers.h"
 #include "warp_helpers.h"
 
+#ifndef DEPTH_OUT
 /** Performs a remapping of an input image to an output given two remapping image using nearest neighbor as interpolation.
  *
  * This kernel performs remapping with this method of pixel coordinate translation:
@@ -129,20 +130,20 @@ __kernel void remap_bilinear_nchw(
 
     vstore4(bilinear_interpolate(&in, clamp_to_border(map_coords, width, height), width, height), 0, out.ptr);
 }
-
+#else // DEPTH_OUT
 /** Performs a remapping of an input image to an output given two remapping image using nearest neighbor as interpolation.
  *  Also applies constant border value, "border_val", if "CONSTANT_BORDER" is set.
  *
  * This kernel performs remapping with this method of pixel coordinate translation:
  *     out(x,y) = in(mapx(x,y), mapy(x,y));
  *
- * @param[in]  in_ptr                             Pointer to the source image. Supported data types: U8.
+ * @param[in]  in_ptr                             Pointer to the source image. Supported data types: U8,F16.
  * @param[in]  in_stride_x                        Stride of the source image in X dimension (in bytes)
  * @param[in]  in_step_x                          in_stride_x * number of elements along X processed per work item (in bytes)
  * @param[in]  in_stride_y                        Stride of the source image in Y dimension (in bytes)
  * @param[in]  in_step_y                          in_stride_y * number of elements along Y processed per work item (in bytes)
  * @param[in]  in_offset_first_element_in_bytes   Offset of the first element in the source image
- * @param[out] out_ptr                            Pointer to the destination image. Supported data types: U8.
+ * @param[out] out_ptr                            Pointer to the destination image. Supported data types: U8,F16.
  * @param[in]  out_stride_x                       Stride of the destination image in X dimension (in bytes)
  * @param[in]  out_step_x                         out_stride_x * number of elements along X processed per work item (in bytes)
  * @param[in]  out_stride_y                       Stride of the destination image in Y dimension (in bytes)
@@ -162,10 +163,8 @@ __kernel void remap_bilinear_nchw(
  * @param[in]  mapy_offset_first_element_in_bytes Offset of the first element in the remapping image
  * @param[in]  width                              Width of the input image
  * @param[in]  height                             Height of the input image
+ * @param[in]  border_val                         Value to use for border around input tensor when in CONSTANT border is selected
  */
-
-#if defined(DEPTH_OUT)
-
 __kernel void remap_nearest_neighbour_nhwc(
     TENSOR4D_DECLARATION(in),
     TENSOR4D_DECLARATION(out),
@@ -206,13 +205,13 @@ __kernel void remap_nearest_neighbour_nhwc(
  * This kernel performs remapping with this method of pixel coordinate translation:
  *     out(x,y) = in(mapx(x,y), mapy(x,y));
  *
- * @param[in]  in_ptr                             Pointer to the source image. Supported data types: U8.
+ * @param[in]  in_ptr                             Pointer to the source image. Supported data types: U8,F16.
  * @param[in]  in_stride_x                        Stride of the source image in X dimension (in bytes)
  * @param[in]  in_step_x                          in_stride_x * number of elements along X processed per work item (in bytes)
  * @param[in]  in_stride_y                        Stride of the source image in Y dimension (in bytes)
  * @param[in]  in_step_y                          in_stride_y * number of elements along Y processed per work item (in bytes)
  * @param[in]  in_offset_first_element_in_bytes   Offset of the first element in the source image
- * @param[out] out_ptr                            Pointer to the destination image. Supported data types: U8.
+ * @param[out] out_ptr                            Pointer to the destination image. Supported data types: U8,F16.
  * @param[in]  out_stride_x                       Stride of the destination image in X dimension (in bytes)
  * @param[in]  out_step_x                         out_stride_x * number of elements along X processed per work item (in bytes)
  * @param[in]  out_stride_y                       Stride of the destination image in Y dimension (in bytes)
@@ -232,6 +231,7 @@ __kernel void remap_nearest_neighbour_nhwc(
  * @param[in]  mapy_offset_first_element_in_bytes Offset of the first element in the remapping image
  * @param[in]  width                              Width of the input image
  * @param[in]  height                             Height of the input image
+ * @param[in]  border_val                         Value to use for border around input tensor when in CONSTANT border is selected
  */
 __kernel void remap_bilinear_nhwc(
     TENSOR4D_DECLARATION(in),
diff --git a/src/core/CL/kernels/CLRemapKernel.cpp b/src/core/CL/kernels/CLRemapKernel.cpp
index 6edd744db7..7e3157c99d 100644
--- a/src/core/CL/kernels/CLRemapKernel.cpp
+++ b/src/core/CL/kernels/CLRemapKernel.cpp
@@ -32,8 +32,6 @@
 #include "src/core/AccessWindowStatic.h"
 #include "src/core/helpers/WindowHelpers.h"
 
-#include <algorithm>
-
 namespace arm_compute
 {
 CLRemapKernel::CLRemapKernel()
@@ -54,11 +52,18 @@ void CLRemapKernel::set_constant_border(unsigned int idx, const PixelValue &cons
     ICLKernel::add_argument<T>(idx, static_cast<T>(value));
 }
 
-Status CLRemapKernel::validate(const ITensorInfo *input, const ITensorInfo *map_x, const ITensorInfo *map_y, ITensorInfo *output, RemapInfo info)
+Status CLRemapKernel::validate(const ITensorInfo *input, const ITensorInfo *map_x, const ITensorInfo *map_y, const ITensorInfo *output, RemapInfo info)
 {
     ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, map_x, map_y, output);
-    ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8);
-    ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::U8);
+    if(input->data_layout() == DataLayout::NCHW)
+    {
+        ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8);
+    }
+    else
+    {
+        ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8, DataType::F16);
+    }
+    ARM_COMPUTE_RETURN_ERROR_ON_MSG(input->data_type() != output->data_type(), "Input/output have different data types");
     ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(map_x, 1, DataType::F32);
     ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(map_y, 1, DataType::F32);
     ARM_COMPUTE_RETURN_ERROR_ON_MSG(info.policy == InterpolationPolicy::AREA, "Area interpolation is not supported!");
@@ -68,7 +73,8 @@ Status CLRemapKernel::validate(const ITensorInfo *input, const ITensorInfo *map_
 
 void CLRemapKernel::configure(const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output, RemapInfo info)
 {
-    CLRemapKernel::validate(input->info(), map_x->info(), map_y->info(), output->info(), info);
+    ARM_COMPUTE_ERROR_ON_NULLPTR(input, map_x, map_y, output);
+    ARM_COMPUTE_ERROR_THROW_ON(CLRemapKernel::validate(input->info(), map_x->info(), map_y->info(), output->info(), info));
 
     _input       = input;
     _output      = output;
@@ -118,7 +124,19 @@ void CLRemapKernel::configure(const CLCompileContext &compile_context, const ICL
     _kernel.setArg<cl_float>(idx++, input_height);
     if(is_nhwc && is_constant_border)
     {
-        set_constant_border<uint8_t>(idx, info.constant_border_value);
+        switch(input->info()->data_type())
+        {
+            case DataType::U8:
+                set_constant_border<uint8_t>(idx, info.constant_border_value);
+                break;
+            case DataType::F16:
+                static_assert(sizeof(cl_half) == sizeof(half), "Half must be same size as cl_half");
+                static_assert(sizeof(cl_half) == 2, "Half must be 16 bit");
+                set_constant_border<half>(idx, info.constant_border_value);
+                break;
+            default:
+                ARM_COMPUTE_ERROR("Data Type not handled");
+        }
     }
 }
 
diff --git a/src/core/CL/kernels/CLRemapKernel.h b/src/core/CL/kernels/CLRemapKernel.h
index 1e3a4ad13f..93b0b4e660 100644
--- a/src/core/CL/kernels/CLRemapKernel.h
+++ b/src/core/CL/kernels/CLRemapKernel.h
@@ -49,28 +49,22 @@ public:
     /** Initialize the kernel's input, output and border mode.
      *
      * @param[in]  compile_context The compile context to be used.
-     * @param[in]  input           Source tensor. Data types supported: U8.
+     * @param[in]  input           Source tensor. Data types supported: U8 (or F16 when layout is NHWC).
      * @param[in]  map_x           Map for X coordinates. Data types supported: F32.
      * @param[in]  map_y           Map for Y coordinates. Data types supported: F32.
-     * @param[out] output          Destination tensor. Data types supported: U8. All but the lowest two dimensions must be the same size as in the input tensor, i.e. remapping is only performed within the XY-plane.
+     * @param[out] output          Destination tensor. Data types supported: Same as @p input. All but the lowest two dimensions must be the same size as in the input tensor, i.e. remapping is only performed within the XY-plane.
      * @param[in]  info            RemapInfo struct:
      *                                   - policy                   Interpolation policy to use. Only NEAREST and BILINEAR are supported.
      *                                   - border_mode              Border mode to use on the input tensor. Only CONSTANT and UNDEFINED are supported.
      *                                   - constant_border_value    Constant value to use for borders if border_mode is set to CONSTANT.
      */
     void configure(const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output, RemapInfo info);
-    /** Validate the kernel's input, output and border mode.
+    /** Checks if the kernel's input, output and border mode will lead to a valid configuration of @ref CLRemapKernel
+     *
+     * Similar to @ref CLRemapKernel::configure(const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output, RemapInfo info)
      *
-     * @param[in]  input  Source tensor. Data types supported: U8.
-     * @param[in]  map_x  Map for X coordinates. Data types supported: F32.
-     * @param[in]  map_y  Map for Y coordinates. Data types supported: F32.
-     * @param[out] output Destination tensor. Data types supported: U8. All but the lowest two dimensions must be the same size as in the input tensor, i.e. remapping is only performed within the XY-plane.
-     * @param[in]  info   RemapInfo struct:
-     *                         - policy                   Interpolation policy to use. Only NEAREST and BILINEAR are supported.
-     *                         - border_mode              Border mode to use on the input tensor. Only CONSTANT and UNDEFINED are supported.
-     *                         - constant_border_value    Constant value to use for borders if border_mode is set to CONSTANT.
      */
-    static Status validate(const ITensorInfo *input, const ITensorInfo *map_x, const ITensorInfo *map_y, ITensorInfo *output, RemapInfo info);
+    static Status validate(const ITensorInfo *input, const ITensorInfo *map_x, const ITensorInfo *map_y, const ITensorInfo *output, RemapInfo info);
     /** Function to set the constant value on fill border kernel depending on type.
      *
      * @param[in] idx                   Index of the kernel argument to set.
diff --git a/src/runtime/CL/functions/CLRemap.cpp b/src/runtime/CL/functions/CLRemap.cpp
index 0a1f864543..de9f0a0148 100644
--- a/src/runtime/CL/functions/CLRemap.cpp
+++ b/src/runtime/CL/functions/CLRemap.cpp
@@ -21,31 +21,45 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  * SOFTWARE.
  */
-#include "arm_compute/runtime/CL/functions/CLRemap.h"
 
+#include "arm_compute/runtime/CL/functions/CLRemap.h"
 #include "arm_compute/core/CL/ICLTensor.h"
 #include "arm_compute/core/Error.h"
-#include "arm_compute/core/PixelValue.h"
 #include "arm_compute/core/TensorInfo.h"
 #include "arm_compute/core/Validate.h"
 #include "src/core/CL/kernels/CLFillBorderKernel.h"
 #include "src/core/CL/kernels/CLRemapKernel.h"
 
-#include <utility>
-
-using namespace arm_compute;
+namespace arm_compute
+{
+void CLRemap::configure(const CLCompileContext &compile_context, ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output, InterpolationPolicy policy,
+                        BorderMode border_mode, uint8_t constant_border_value)
+{
+    configure(compile_context, input, map_x, map_y, output, policy, border_mode, PixelValue{ constant_border_value });
+}
 
 void CLRemap::configure(ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
 {
+    configure(CLKernelLibrary::get().get_compile_context(), input, map_x, map_y, output, policy, border_mode, PixelValue{ constant_border_value });
+}
+
+void CLRemap::configure(ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output, InterpolationPolicy policy, BorderMode border_mode, PixelValue constant_border_value)
+{
     configure(CLKernelLibrary::get().get_compile_context(), input, map_x, map_y, output, policy, border_mode, constant_border_value);
 }
 
 void CLRemap::configure(const CLCompileContext &compile_context, ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output, InterpolationPolicy policy,
-                        BorderMode border_mode,
-                        uint8_t    constant_border_value)
+                        BorderMode border_mode, PixelValue constant_border_value)
 {
     auto k = std::make_unique<CLRemapKernel>();
-    k->configure(compile_context, input, map_x, map_y, output, RemapInfo{ policy, border_mode, PixelValue(constant_border_value) });
+    k->configure(compile_context, input, map_x, map_y, output, RemapInfo{ policy, border_mode, constant_border_value });
     _kernel = std::move(k);
-    _border_handler->configure(compile_context, input, _kernel->border_size(), border_mode, PixelValue(constant_border_value));
+    _border_handler->configure(compile_context, input, _kernel->border_size(), border_mode, constant_border_value);
+}
+
+Status CLRemap::validate(const ITensorInfo *input, const ITensorInfo *map_x, const ITensorInfo *map_y, const ITensorInfo *output, const InterpolationPolicy policy, const BorderMode border_mode,
+                         PixelValue constant_border_value)
+{
+    return CLRemapKernel::validate(input, map_x, map_y, output, RemapInfo{ policy, border_mode, constant_border_value });
 }
+} // namespace arm_compute
diff --git a/src/runtime/NEON/functions/NERemap.cpp b/src/runtime/NEON/functions/NERemap.cpp
index d9fd987480..dc2c4b4f91 100644
--- a/src/runtime/NEON/functions/NERemap.cpp
+++ b/src/runtime/NEON/functions/NERemap.cpp
@@ -36,6 +36,10 @@ namespace arm_compute
 {
 void NERemap::configure(ITensor *input, const ITensor *map_x, const ITensor *map_y, ITensor *output, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
 {
+    configure(input, map_x, map_y, output, policy, border_mode, PixelValue{ constant_border_value });
+}
+void NERemap::configure(ITensor *input, const ITensor *map_x, const ITensor *map_y, ITensor *output, InterpolationPolicy policy, BorderMode border_mode, PixelValue constant_border_value)
+{
     ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8);
     ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::U8);
     ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(map_x, 1, DataType::F32);
@@ -43,7 +47,7 @@ void NERemap::configure(ITensor *input, const ITensor *map_x, const ITensor *map
     ARM_COMPUTE_ERROR_ON_MSG(policy == InterpolationPolicy::AREA, "Area interpolation is not supported");
 
     auto k = std::make_unique<NERemapKernel>();
-    k->configure(input, map_x, map_y, output, policy, border_mode, constant_border_value);
+    k->configure(input, map_x, map_y, output, policy, border_mode, constant_border_value.get<uint8_t>());
     _kernel = std::move(k);
 }
 } // namespace arm_compute
diff --git a/tests/validation/CL/Remap.cpp b/tests/validation/CL/Remap.cpp
index bbb3cecea9..7849d77394 100644
--- a/tests/validation/CL/Remap.cpp
+++ b/tests/validation/CL/Remap.cpp
@@ -48,15 +48,64 @@ constexpr AbsoluteTolerance<uint8_t> tolerance_value(1);
 
 TEST_SUITE(CL)
 TEST_SUITE(Remap)
+
+// *INDENT-OFF*
+// clang-format off
+
+DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip(zip(zip(
+               framework::dataset::make("input", { TensorInfo(TensorShape(10U, 10U), 1, DataType::U8, DataLayout::NCHW),
+                                                   TensorInfo(TensorShape(10U, 10U), 1, DataType::U8, DataLayout::NHWC),
+                                                   TensorInfo(TensorShape(10U, 10U), 1, DataType::F16, DataLayout::NCHW),
+                                                   TensorInfo(TensorShape(10U, 10U), 1, DataType::F16, DataLayout::NHWC)
+                                                      }),
+               framework::dataset::make("map_x", { TensorInfo(TensorShape(10U, 10U), 1, DataType::F32, DataLayout::NCHW),
+                                                   TensorInfo(TensorShape(10U, 10U), 1, DataType::F32, DataLayout::NHWC),
+                                                   TensorInfo(TensorShape(10U, 10U), 1, DataType::F32, DataLayout::NCHW),
+                                                   TensorInfo(TensorShape(10U, 10U), 1, DataType::F32, DataLayout::NHWC)
+
+                                                      })),
+               framework::dataset::make("map_y", { TensorInfo(TensorShape(10U, 10U), 1, DataType::F32, DataLayout::NCHW),
+                                                   TensorInfo(TensorShape(10U, 10U), 1, DataType::F32, DataLayout::NHWC),
+                                                   TensorInfo(TensorShape(10U, 10U), 1, DataType::F32, DataLayout::NCHW),
+                                                   TensorInfo(TensorShape(10U, 10U), 1, DataType::F32, DataLayout::NHWC)
+                                                      })),
+               framework::dataset::make("output", { TensorInfo(TensorShape(10U, 10U), 1, DataType::U8, DataLayout::NCHW),
+                                                    TensorInfo(TensorShape(10U, 10U), 1, DataType::U8, DataLayout::NHWC),
+                                                    TensorInfo(TensorShape(10U, 10U), 1, DataType::F16, DataLayout::NCHW),
+                                                    TensorInfo(TensorShape(10U, 10U), 1, DataType::F16, DataLayout::NHWC)
+                                                      })),
+               framework::dataset::make("policy",{ InterpolationPolicy::NEAREST_NEIGHBOR,
+                                                   InterpolationPolicy::NEAREST_NEIGHBOR,
+                                                   InterpolationPolicy::NEAREST_NEIGHBOR,
+                                                   InterpolationPolicy::NEAREST_NEIGHBOR
+                                                      })),
+               framework::dataset::make("border_mode",{ BorderMode::CONSTANT,
+                                                        BorderMode::CONSTANT,
+                                                        BorderMode::CONSTANT,
+                                                        BorderMode::CONSTANT
+                                                      })),
+               framework::dataset::make("Expected", { true, // NCHW, U8
+                                                      true, // NHWC, U8
+                                                      false, // NCHW, F16
+                                                      true // NHWC, F16
+                                                    })),
+               input, map_x, map_y, output, policy, border_mode, expected)
+{
+    ARM_COMPUTE_EXPECT(bool(CLRemap::validate(&input, &map_x, &map_y, &output, policy, border_mode, PixelValue{})) == expected, framework::LogLevel::ERRORS);
+}
+// clang-format on
+// *INDENT-ON*
 template <typename T>
 using CLRemapFixture = RemapValidationFixture<CLTensor, CLAccessor, CLRemap, T>;
 template <typename T>
 using CLRemapLayoutFixture = RemapValidationMixedLayoutFixture<CLTensor, CLAccessor, CLRemap, T>;
 
-FIXTURE_DATA_TEST_CASE(RunSmall, CLRemapLayoutFixture<uint8_t>, framework::DatasetMode::ALL, combine(combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
-                                                                                                                     framework::dataset::make("DataType", DataType::U8)),
-                                                                                                             framework::dataset::make("BorderModes", { BorderMode::UNDEFINED, BorderMode::CONSTANT })),
-                                                                                                     framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC })))
+TEST_SUITE(U8)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLRemapLayoutFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(),
+                                                                                                                   framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+                                                                                                                   framework::dataset::make("DataType", DataType::U8)),
+                                                                                                                   framework::dataset::make("BorderModes", { BorderMode::UNDEFINED, BorderMode::CONSTANT })),
+                                                                                                           framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC })))
 {
     // Validate output
     validate(CLAccessor(_target), _reference, _valid_mask, tolerance_value);
@@ -69,7 +118,19 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLRemapFixture<uint8_t>, framework::DatasetMode
     // Validate output
     validate(CLAccessor(_target), _reference, _valid_mask, tolerance_value);
 }
+TEST_SUITE_END() // U8
 
+TEST_SUITE(F16)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLRemapLayoutFixture<half>, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(),
+                                                                                                                        framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })),
+                                                                                                                        framework::dataset::make("DataType", DataType::F16)),
+                                                                                                                framework::dataset::make("BorderModes", { BorderMode::UNDEFINED, BorderMode::CONSTANT })),
+                                                                                                        framework::dataset::make("DataLayout", DataLayout::NHWC)))
+{
+    // Validate output
+    validate(CLAccessor(_target), _reference, _valid_mask, tolerance_value);
+}
+TEST_SUITE_END() // F16
 TEST_SUITE_END()
 TEST_SUITE_END()
 } // namespace validation
diff --git a/tests/validation/fixtures/RemapFixture.h b/tests/validation/fixtures/RemapFixture.h
index 2cb8e67f62..03cb6aef42 100644
--- a/tests/validation/fixtures/RemapFixture.h
+++ b/tests/validation/fixtures/RemapFixture.h
@@ -50,7 +50,7 @@ public:
     {
         std::mt19937                           gen(library->seed());
         std::uniform_int_distribution<uint8_t> distribution(0, 255);
-        const T                                constant_border_value = static_cast<T>(distribution(gen));
+        PixelValue                             constant_border_value{ static_cast<T>(distribution(gen)) };
 
         _data_layout = data_layout;
         _target      = compute_target(shape, policy, data_type, border_mode, constant_border_value);
@@ -59,13 +59,35 @@ public:
 
 protected:
     template <typename U>
-    void fill(U &&tensor, int i, float min, float max)
+    void fill(U &&tensor, int i, int min, int max)
     {
-        std::uniform_int_distribution<> distribution((int)min, (int)max);
-        library->fill(tensor, distribution, i);
+        switch(tensor.data_type())
+        {
+            case DataType::F32:
+            {
+                // map_x,y as integer values
+                std::uniform_int_distribution<int> distribution(min, max);
+                library->fill(tensor, distribution, i);
+                break;
+            }
+            case DataType::F16:
+            {
+                arm_compute::utils::uniform_real_distribution_16bit<half> distribution(static_cast<float>(min), static_cast<float>(max));
+                library->fill(tensor, distribution, i);
+                break;
+            }
+            case DataType::U8:
+            {
+                std::uniform_int_distribution<uint8_t> distribution(min, max);
+                library->fill(tensor, distribution, i);
+                break;
+            }
+            default:
+                ARM_COMPUTE_ERROR("DataType for Remap not supported");
+        }
     }
 
-    TensorType compute_target(TensorShape shape, InterpolationPolicy policy, DataType data_type, BorderMode border_mode, T constant_border_value)
+    TensorType compute_target(TensorShape shape, InterpolationPolicy policy, DataType data_type, BorderMode border_mode, PixelValue constant_border_value)
     {
         if(_data_layout == DataLayout::NHWC)
         {
@@ -111,14 +133,16 @@ protected:
         return dst;
     }
 
-    SimpleTensor<T> compute_reference(const TensorShape shape, InterpolationPolicy policy, DataType data_type, BorderMode border_mode, T constant_border_value)
+    SimpleTensor<T> compute_reference(const TensorShape shape, InterpolationPolicy policy, DataType data_type, BorderMode border_mode, PixelValue constant_border_value)
     {
-        ARM_COMPUTE_ERROR_ON(data_type != DataType::U8);
+        ARM_COMPUTE_ERROR_ON(data_type != DataType::U8 && data_type != DataType::F16);
 
         // Create reference
         SimpleTensor<T>     src{ shape, data_type };
         SimpleTensor<float> map_x{ shape, DataType::F32 };
         SimpleTensor<float> map_y{ shape, DataType::F32 };
+        T                   border_value{};
+        constant_border_value.get(border_value);
 
         // Create the valid mask Tensor
         _valid_mask = SimpleTensor<T> { shape, data_type };
@@ -131,7 +155,7 @@ protected:
         fill(map_y, 2, -5, max_val);
 
         // Compute reference
-        return reference::remap<T>(src, map_x, map_y, _valid_mask, policy, border_mode, constant_border_value);
+        return reference::remap<T>(src, map_x, map_y, _valid_mask, policy, border_mode, border_value);
     }
 
     TensorType      _target{};
diff --git a/tests/validation/reference/Remap.cpp b/tests/validation/reference/Remap.cpp
index 33c5a7de68..dfbe1aa12b 100644
--- a/tests/validation/reference/Remap.cpp
+++ b/tests/validation/reference/Remap.cpp
@@ -99,13 +99,16 @@ SimpleTensor<T> remap(const SimpleTensor<T> &in, SimpleTensor<float> &map_x, Sim
             }
         }
     }
-
     return out;
 }
 
 template SimpleTensor<uint8_t> remap(const SimpleTensor<uint8_t> &src, SimpleTensor<float> &map_x, SimpleTensor<float> &map_y, SimpleTensor<uint8_t> &valid_mask, InterpolationPolicy policy,
                                      BorderMode border_mode,
                                      uint8_t    constant_border_value);
+
+template SimpleTensor<half> remap(const SimpleTensor<half> &src, SimpleTensor<float> &map_x, SimpleTensor<float> &map_y, SimpleTensor<half> &valid_mask, InterpolationPolicy policy,
+                                  BorderMode border_mode,
+                                  half       constant_border_value);
 } // namespace reference
 } // namespace validation
 } // namespace test