From ef5aac6c1e119e8db16a33332b5551829f409786 Mon Sep 17 00:00:00 2001 From: Freddie Liardet Date: Thu, 10 Jun 2021 16:45:58 +0100 Subject: 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 Change-Id: If05f06801aef7a169b73ff1ebe760a42f11ca05c Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5816 Reviewed-by: Michele Di Giorgio Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins --- arm_compute/runtime/CL/functions/CLRemap.h | 47 +++++++++++++++++-- arm_compute/runtime/NEON/functions/NERemap.h | 16 +++++++ src/core/CL/cl_kernels/remap.cl | 16 +++---- src/core/CL/kernels/CLRemapKernel.cpp | 32 ++++++++++--- src/core/CL/kernels/CLRemapKernel.h | 18 +++----- src/runtime/CL/functions/CLRemap.cpp | 32 +++++++++---- src/runtime/NEON/functions/NERemap.cpp | 6 ++- tests/validation/CL/Remap.cpp | 69 ++++++++++++++++++++++++++-- tests/validation/fixtures/RemapFixture.h | 40 ++++++++++++---- tests/validation/reference/Remap.cpp | 5 +- 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 - namespace arm_compute { CLRemapKernel::CLRemapKernel() @@ -54,11 +52,18 @@ void CLRemapKernel::set_constant_border(unsigned int idx, const PixelValue &cons ICLKernel::add_argument(idx, static_cast(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(idx++, input_height); if(is_nhwc && is_constant_border) { - set_constant_border(idx, info.constant_border_value); + switch(input->info()->data_type()) + { + case DataType::U8: + set_constant_border(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(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 - -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(); - 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 @@ -35,6 +35,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); @@ -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(); - 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()); _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 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 using CLRemapFixture = RemapValidationFixture; template using CLRemapLayoutFixture = RemapValidationMixedLayoutFixture; -FIXTURE_DATA_TEST_CASE(RunSmall, CLRemapLayoutFixture, 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, 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, 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, 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 distribution(0, 255); - const T constant_border_value = static_cast(distribution(gen)); + PixelValue constant_border_value{ static_cast(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 - 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 distribution(min, max); + library->fill(tensor, distribution, i); + break; + } + case DataType::F16: + { + arm_compute::utils::uniform_real_distribution_16bit distribution(static_cast(min), static_cast(max)); + library->fill(tensor, distribution, i); + break; + } + case DataType::U8: + { + std::uniform_int_distribution 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 compute_reference(const TensorShape shape, InterpolationPolicy policy, DataType data_type, BorderMode border_mode, T constant_border_value) + SimpleTensor 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 src{ shape, data_type }; SimpleTensor map_x{ shape, DataType::F32 }; SimpleTensor map_y{ shape, DataType::F32 }; + T border_value{}; + constant_border_value.get(border_value); // Create the valid mask Tensor _valid_mask = SimpleTensor { shape, data_type }; @@ -131,7 +155,7 @@ protected: fill(map_y, 2, -5, max_val); // Compute reference - return reference::remap(src, map_x, map_y, _valid_mask, policy, border_mode, constant_border_value); + return reference::remap(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 remap(const SimpleTensor &in, SimpleTensor &map_x, Sim } } } - return out; } template SimpleTensor remap(const SimpleTensor &src, SimpleTensor &map_x, SimpleTensor &map_y, SimpleTensor &valid_mask, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value); + +template SimpleTensor remap(const SimpleTensor &src, SimpleTensor &map_x, SimpleTensor &map_y, SimpleTensor &valid_mask, InterpolationPolicy policy, + BorderMode border_mode, + half constant_border_value); } // namespace reference } // namespace validation } // namespace test -- cgit v1.2.1