aboutsummaryrefslogtreecommitdiff
path: root/src/core/CL/cl_kernels/remap.cl
diff options
context:
space:
mode:
authorFreddie Liardet <frederick.liardet@arm.com>2021-06-10 16:45:58 +0100
committerfrederick.liardet <frederick.liardet@arm.com>2021-06-22 12:39:24 +0000
commitef5aac6c1e119e8db16a33332b5551829f409786 (patch)
tree78b7d34c9da20e8c6e4393981ada8e1253f239b2 /src/core/CL/cl_kernels/remap.cl
parent8266ae50a7da14ce27592f89181287be81969fd0 (diff)
downloadComputeLibrary-ef5aac6c1e119e8db16a33332b5551829f409786.tar.gz
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>
Diffstat (limited to 'src/core/CL/cl_kernels/remap.cl')
-rw-r--r--src/core/CL/cl_kernels/remap.cl16
1 files changed, 8 insertions, 8 deletions
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),