diff options
| -rw-r--r-- | arm_compute/runtime/CL/functions/CLRemap.h | 47 | ||||
| -rw-r--r-- | arm_compute/runtime/NEON/functions/NERemap.h | 16 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/remap.cl | 16 | ||||
| -rw-r--r-- | src/core/CL/kernels/CLRemapKernel.cpp | 32 | ||||
| -rw-r--r-- | src/core/CL/kernels/CLRemapKernel.h | 18 | ||||
| -rw-r--r-- | src/runtime/CL/functions/CLRemap.cpp | 32 | ||||
| -rw-r--r-- | src/runtime/NEON/functions/NERemap.cpp | 6 | ||||
| -rw-r--r-- | tests/validation/CL/Remap.cpp | 69 | ||||
| -rw-r--r-- | tests/validation/fixtures/RemapFixture.h | 40 | ||||
| -rw-r--r-- | 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 <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 |