From 7317e393de848cb20dd326a45a9efb18a3ee7701 Mon Sep 17 00:00:00 2001 From: Michalis Spyrou Date: Fri, 17 Jan 2020 11:27:49 +0000 Subject: COMPMID-2941 Add QAsymm8 and S32 support for ArgMinMax to CL Select operation does not work with any combination of datatypes. Added a new OpenCL build option, DATA_TYPE_SELECT, in oder to specify the select result. Change-Id: I5c73c2a6b1d121288b7f6edba3e0cb1104636844 Signed-off-by: Michalis Spyrou Reviewed-on: https://review.mlplatform.org/c/2601 Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins Reviewed-by: Michele Di Giorgio --- arm_compute/core/CL/CLHelpers.h | 10 +++- .../runtime/CL/functions/CLArgMinMaxLayer.h | 6 +- src/core/CL/CLHelpers.cpp | 20 ++++++- src/core/CL/cl_kernels/arg_min_max.cl | 65 +++++++++++++--------- src/core/CL/cl_kernels/helpers.h | 15 +---- src/core/CL/cl_kernels/reduction_operation.cl | 13 +++++ src/core/CL/kernels/CLArgMinMaxLayerKernel.cpp | 13 ++--- tests/validation/CL/ArgMinMax.cpp | 51 ++++++++++------- tests/validation/NEON/ArgMinMax.cpp | 22 +------- 9 files changed, 120 insertions(+), 95 deletions(-) diff --git a/arm_compute/core/CL/CLHelpers.h b/arm_compute/core/CL/CLHelpers.h index 4c9cfff2bf..4b9ca4d253 100644 --- a/arm_compute/core/CL/CLHelpers.h +++ b/arm_compute/core/CL/CLHelpers.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2019 ARM Limited. + * Copyright (c) 2016-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -66,6 +66,14 @@ std::string get_cl_promoted_type_from_data_type(const DataType &dt); */ std::string get_cl_unsigned_type_from_element_size(size_t element_size); +/** Translates the element size to an signed integer data type + * + * @param[in] element_size Size in bytes of an element. + * + * @return The string specifying the OpenCL type to be used. + */ +std::string get_cl_signed_type_from_element_size(size_t element_size); + /** Translates a tensor data type to the appropriate OpenCL select type. * * @param[in] dt @ref DataType to be translated to OpenCL select type. diff --git a/arm_compute/runtime/CL/functions/CLArgMinMaxLayer.h b/arm_compute/runtime/CL/functions/CLArgMinMaxLayer.h index 21cded0417..a26fcfda56 100644 --- a/arm_compute/runtime/CL/functions/CLArgMinMaxLayer.h +++ b/arm_compute/runtime/CL/functions/CLArgMinMaxLayer.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -55,7 +55,7 @@ public: CLArgMinMaxLayer(std::shared_ptr memory_manager = nullptr); /** Set the input and output tensors. * - * @param[in] input Input source tensor. Data types supported: F16/F32. + * @param[in] input Input source tensor. Data types supported: QASYMM8/F16/F32. * @param[in] axis Axis to find max/min index. * @param[out] output Output source tensor. Data types supported: U32/S32. * @param[in] op Reduction operation to perform. Operations supported: ARG_IDX_MAX, ARG_IDX_MIN @@ -63,7 +63,7 @@ public: void configure(const ICLTensor *input, int axis, ICLTensor *output, const ReductionOperation &op); /** Static function to check if given info will lead to a valid configuration of @ref CLArgMinMaxLayer * - * @param[in] input Input source tensor info. Data types supported: F16/F32. + * @param[in] input Input source tensor info. Data types supported: QASYMM8/F16/F32. * @param[in] axis Axis to find max/min index. * @param[in] output Output source tensor info. Data types supported: U32/S32. * @param[in] op Reduction operation to perform. Operations supported: ARG_IDX_MAX, ARG_IDX_MIN diff --git a/src/core/CL/CLHelpers.cpp b/src/core/CL/CLHelpers.cpp index 9754bebd18..84de380cc9 100644 --- a/src/core/CL/CLHelpers.cpp +++ b/src/core/CL/CLHelpers.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2019 ARM Limited. + * Copyright (c) 2016-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -118,6 +118,24 @@ std::string get_cl_unsigned_type_from_element_size(size_t element_size) } } +std::string get_cl_signed_type_from_element_size(size_t element_size) +{ + switch(element_size) + { + case 1: + return "char"; + case 2: + return "short"; + case 4: + return "int"; + case 8: + return "long"; + default: + ARM_COMPUTE_ERROR("Data type not supported"); + return ""; + } +} + std::string get_cl_select_type_from_data_type(const DataType &dt) { switch(dt) diff --git a/src/core/CL/cl_kernels/arg_min_max.cl b/src/core/CL/cl_kernels/arg_min_max.cl index 3f75377636..06dcc8ddde 100644 --- a/src/core/CL/cl_kernels/arg_min_max.cl +++ b/src/core/CL/cl_kernels/arg_min_max.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019 ARM Limited. + * Copyright (c) 2019-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -23,6 +23,19 @@ */ #include "helpers.h" +#if defined(FLOAT_DATA_TYPE) +#define ISGREATER(x, y) isgreater(x, y) +#define ISLESS(x, y) isless(x, y) +#else // !FLOAT_DATA_TYPE +#if defined(WIDTH) +#define ISGREATER(x, y) (x > y) ? 1 : 0 +#define ISLESS(x, y) (x < y) ? 1 : 0 +#else // !defined(WIDTH) +#define ISGREATER(x, y) select((VEC_DATA_TYPE(DATA_TYPE_SELECT, 16))0, (VEC_DATA_TYPE(DATA_TYPE_SELECT, 16)) - 1, x > y) +#define ISLESS(x, y) select((VEC_DATA_TYPE(DATA_TYPE_SELECT, 16))0, (VEC_DATA_TYPE(DATA_TYPE_SELECT, 16)) - 1, x < y) +#endif // defined(WIDTH) +#endif // defined(FLOAT_DATA_TYPE) + #if defined(ARG_MAX) #define CONDITION_TO_USE(x, y) ISGREATER(x, y) #elif defined(ARG_MIN) @@ -31,7 +44,7 @@ #error "Unsupported reduction operation!" #endif // defined(ARG_MAX) -#if defined(DATA_TYPE_OUTPUT) +#if defined(DATA_TYPE_OUTPUT) && defined(DATA_TYPE_SELECT) #if defined(WIDTH) #if defined(ARG_MIN) #if defined(PREV_OUTPUT) @@ -81,20 +94,20 @@ inline DATA_TYPE_OUTPUT arg_idx_min(__global const DATA_TYPE *input, const int x VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16) res = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 }; - VEC_DATA_TYPE(COND_DATA_TYPE, 8) + VEC_DATA_TYPE(DATA_TYPE_SELECT, 8) idx_sel = (in.s01234567 <= in.s89abcdef); in.s01234567 = select(in.s89abcdef, in.s01234567, idx_sel); res.s01234567 = select(res.s89abcdef, res.s01234567, CONVERT(idx_sel, int8)); - idx_sel.s0123 = (in.s0123 < in.s4567) || (in.s0123 == in.s4567 && CONVERT((res.s0123 < res.s4567), VEC_DATA_TYPE(COND_DATA_TYPE, 4))); + idx_sel.s0123 = (in.s0123 < in.s4567) || (in.s0123 == in.s4567 && CONVERT((res.s0123 < res.s4567), VEC_DATA_TYPE(DATA_TYPE_SELECT, 4))); in.s0123 = select(in.s4567, in.s0123, idx_sel.s0123); res.s0123 = select(res.s4567, res.s0123, CONVERT(idx_sel.s0123, int4)); - idx_sel.s01 = (in.s01 < in.s23) || (in.s01 == in.s23 && CONVERT((res.s01 < res.s23), VEC_DATA_TYPE(COND_DATA_TYPE, 2))); + idx_sel.s01 = (in.s01 < in.s23) || (in.s01 == in.s23 && CONVERT((res.s01 < res.s23), VEC_DATA_TYPE(DATA_TYPE_SELECT, 2))); in.s01 = select(in.s23, in.s01, idx_sel.s01); res.s01 = select(res.s23, res.s01, CONVERT(idx_sel.s01, int2)); - idx_sel.s0 = (in.s0 < in.s1) || (in.s0 == in.s1 && CONVERT((res.s0 < res.s1), COND_DATA_TYPE)); + idx_sel.s0 = (in.s0 < in.s1) || (in.s0 == in.s1 && CONVERT((res.s0 < res.s1), DATA_TYPE_SELECT)); res.s0 = select(res.s1, res.s0, CONVERT(idx_sel.s0, int)); return res.s0 + x_elem; @@ -150,20 +163,20 @@ inline DATA_TYPE_OUTPUT arg_idx_max(__global const DATA_TYPE *input, const int x VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16) res = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 }; - VEC_DATA_TYPE(COND_DATA_TYPE, 8) + VEC_DATA_TYPE(DATA_TYPE_SELECT, 8) idx_sel = (in.s01234567 >= in.s89abcdef); in.s01234567 = select(in.s89abcdef, in.s01234567, idx_sel); res.s01234567 = select(res.s89abcdef, res.s01234567, CONVERT(idx_sel, int8)); - idx_sel.s0123 = (in.s0123 > in.s4567) || (in.s0123 == in.s4567 && CONVERT((res.s0123 < res.s4567), VEC_DATA_TYPE(COND_DATA_TYPE, 4))); + idx_sel.s0123 = (in.s0123 > in.s4567) || (in.s0123 == in.s4567 && CONVERT((res.s0123 < res.s4567), VEC_DATA_TYPE(DATA_TYPE_SELECT, 4))); in.s0123 = select(in.s4567, in.s0123, idx_sel.s0123); res.s0123 = select(res.s4567, res.s0123, CONVERT(idx_sel.s0123, int4)); - idx_sel.s01 = (in.s01 > in.s23) || (in.s01 == in.s23 && CONVERT((res.s01 < res.s23), VEC_DATA_TYPE(COND_DATA_TYPE, 2))); + idx_sel.s01 = (in.s01 > in.s23) || (in.s01 == in.s23 && CONVERT((res.s01 < res.s23), VEC_DATA_TYPE(DATA_TYPE_SELECT, 2))); in.s01 = select(in.s23, in.s01, idx_sel.s01); res.s01 = select(res.s23, res.s01, CONVERT(idx_sel.s01, int2)); - idx_sel.s0 = (in.s0 > in.s1) || (in.s0 == in.s1 && CONVERT((res.s0 < res.s1), COND_DATA_TYPE)); + idx_sel.s0 = (in.s0 > in.s1) || (in.s0 == in.s1 && CONVERT((res.s0 < res.s1), DATA_TYPE_SELECT)); res.s0 = select(res.s1, res.s0, CONVERT(idx_sel.s0, int)); return res.s0 + x_elem; @@ -276,7 +289,7 @@ __kernel void arg_min_max_x( * * @note The input data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float * @note The data type of the output must be passed at compile time using -DDATA_TYPE_OUTPUT: e.g. -DDATA_TYPE_OUTPUT=uint - * @note The data type of the intermediate results must be passed at compile time using -DDATA_TYPE_PROMOTED: e.g. -DDATA_TYPE_PROMOTED=uint + * @note The data type of the select results must be passed at compile time using -DDATA_TYPE_SELECT: e.g. -DDATA_TYPE_SELECT=int * @note The height size must be passed at compile time using -DHEIGHT e.g. -DHEIGHT=128 * * @param[in] src_ptr Pointer to the source tensor. Supported data types: S32/F16/F32 @@ -299,15 +312,15 @@ __kernel void arg_min_max_y( Image src = CONVERT_TO_IMAGE_STRUCT(src); Image output = CONVERT_TO_IMAGE_STRUCT(output); - VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16) - res = CONVERT(vload16(0, (__global DATA_TYPE *)offset(&src, 0, 0)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)); + VEC_DATA_TYPE(DATA_TYPE, 16) + res = CONVERT(vload16(0, (__global DATA_TYPE *)offset(&src, 0, 0)), VEC_DATA_TYPE(DATA_TYPE, 16)); VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16) indx = 0; for(unsigned int y = 1; y < HEIGHT; ++y) { - VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16) - in = CONVERT(vload16(0, (__global DATA_TYPE *)offset(&src, 0, y)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)); + VEC_DATA_TYPE(DATA_TYPE, 16) + in = CONVERT(vload16(0, (__global DATA_TYPE *)offset(&src, 0, y)), VEC_DATA_TYPE(DATA_TYPE, 16)); VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16) cond_conv = CONVERT(CONDITION_TO_USE(in, res), VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16)); @@ -324,7 +337,7 @@ __kernel void arg_min_max_y( /** This kernel performs reduction on z-axis. * * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float - * @note The data type of the intermediate results must be passed at compile time using -DDATA_TYPE_PROMOTED: e.g. -DDATA_TYPE_PROMOTED=uint + * @note The data type of the select results must be passed at compile time using -DDATA_TYPE_SELECT: e.g. -DDATA_TYPE_SELECT=int * @note The depth size must be passed at compile time using -DDEPTH e.g. -DDEPTH=128 * * @param[in] input_ptr Pointer to the source tensor. Supported data types: S32/F16/F32 @@ -351,15 +364,15 @@ __kernel void arg_min_max_z( Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input); Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); - VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16) - res = CONVERT(vload16(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)); + VEC_DATA_TYPE(DATA_TYPE, 16) + res = CONVERT(vload16(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0)), VEC_DATA_TYPE(DATA_TYPE, 16)); VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16) indx = 0; for(DATA_TYPE_OUTPUT z = 1; z < DEPTH; ++z) { - VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16) - in = CONVERT(vload16(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, z)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)); + VEC_DATA_TYPE(DATA_TYPE, 16) + in = CONVERT(vload16(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, z)), VEC_DATA_TYPE(DATA_TYPE, 16)); VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16) cond_conv = CONVERT(CONDITION_TO_USE(in, res), VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16)); @@ -376,7 +389,7 @@ __kernel void arg_min_max_z( /** This kernel performs reduction on w-axis. * * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float - * @note The data type of the intermediate results must be passed at compile time using -DDATA_TYPE_PROMOTED: e.g. -DDATA_TYPE_PROMOTED=uint + * @note The data type of the select results must be passed at compile time using -DDATA_TYPE_SELECT: e.g. -DDATA_TYPE_SELECT=int * @note The batch size must be passed at compile time using -DBATCH e.g. -DBATCH=128 * @note The depth size must be passed at compile time using -DBATCH e.g. -DDEPTH=128 * @@ -408,15 +421,15 @@ __kernel void arg_min_max_w( Tensor4D input = CONVERT_TO_TENSOR4D_STRUCT(input, DEPTH); Tensor4D output = CONVERT_TO_TENSOR4D_STRUCT(output, DEPTH); - VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16) - res = CONVERT(vload16(0, (__global DATA_TYPE *)tensor4D_offset(&input, 0, 0, 0, 0)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)); + VEC_DATA_TYPE(DATA_TYPE, 16) + res = CONVERT(vload16(0, (__global DATA_TYPE *)tensor4D_offset(&input, 0, 0, 0, 0)), VEC_DATA_TYPE(DATA_TYPE, 16)); VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16) indx = 0; for(DATA_TYPE_OUTPUT w = 1; w < BATCH; ++w) { - VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16) - in = CONVERT(vload16(0, (__global DATA_TYPE *)tensor4D_offset(&input, 0, 0, 0, w)), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 16)); + VEC_DATA_TYPE(DATA_TYPE, 16) + in = CONVERT(vload16(0, (__global DATA_TYPE *)tensor4D_offset(&input, 0, 0, 0, w)), VEC_DATA_TYPE(DATA_TYPE, 16)); VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16) cond_conv = CONVERT(CONDITION_TO_USE(in, res), VEC_DATA_TYPE(DATA_TYPE_OUTPUT, 16)); @@ -428,4 +441,4 @@ __kernel void arg_min_max_w( vstore16(indx, 0, (__global DATA_TYPE_OUTPUT *)output.ptr); } #endif /* defined(BATCH) && defined(DEPTH) */ -#endif // defined(DATA_TYPE_OUTPUT) \ No newline at end of file +#endif /* defined(DATA_TYPE_OUTPUT) && defined(DATA_TYPE_SELECT) */ \ No newline at end of file diff --git a/src/core/CL/cl_kernels/helpers.h b/src/core/CL/cl_kernels/helpers.h index ec5701dc69..c4cbf77e96 100644 --- a/src/core/CL/cl_kernels/helpers.h +++ b/src/core/CL/cl_kernels/helpers.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2019 ARM Limited. + * Copyright (c) 2016-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -266,19 +266,6 @@ #define CONVERT_SAT_ROUND_STR(x, type, round) (convert_##type##_sat_##round((x))) #define CONVERT_SAT_ROUND(x, type, round) CONVERT_SAT_ROUND_STR(x, type, round) -#if FLOAT_DATA_TYPE -#define ISGREATER(x, y) isgreater(x, y) -#define ISLESS(x, y) isless(x, y) -#else // !FLOAT_DATA_TYPE -#if defined(WIDTH) -#define ISGREATER(x, y) (x > y) ? 1 : 0 -#define ISLESS(x, y) (x < y) ? 1 : 0 -#else // !defined(WIDTH) -#define ISGREATER(x, y) select((int16)0, (int16)-1, x > y) -#define ISLESS(x, y) select((int16)0, (int16)-1, x < y) -#endif // defined(WIDTH) -#endif // FLOAT_DATA_TYPE - #define VECTOR_DECLARATION(name) \ __global uchar *name##_ptr, \ uint name##_stride_x, \ diff --git a/src/core/CL/cl_kernels/reduction_operation.cl b/src/core/CL/cl_kernels/reduction_operation.cl index 451b962b01..79d0d0683e 100644 --- a/src/core/CL/cl_kernels/reduction_operation.cl +++ b/src/core/CL/cl_kernels/reduction_operation.cl @@ -23,6 +23,19 @@ */ #include "helpers.h" +#if defined(FLOAT_DATA_TYPE) +#define ISGREATER(x, y) isgreater(x, y) +#define ISLESS(x, y) isless(x, y) +#else // !FLOAT_DATA_TYPE +#if defined(WIDTH) +#define ISGREATER(x, y) (x > y) ? 1 : 0 +#define ISLESS(x, y) (x < y) ? 1 : 0 +#else // !defined(WIDTH) +#define ISGREATER(x, y) select(int16)0, int16) - 1, x > y) +#define ISLESS(x, y) select(int16)0, int16) - 1, x < y) +#endif // defined(WIDTH) +#endif // defined(FLOAT_DATA_TYPE) + /** Calculate square sum of a vector * * @param[in] input Pointer to the first pixel. diff --git a/src/core/CL/kernels/CLArgMinMaxLayerKernel.cpp b/src/core/CL/kernels/CLArgMinMaxLayerKernel.cpp index c8e87ba5ce..bdb7ab76d5 100644 --- a/src/core/CL/kernels/CLArgMinMaxLayerKernel.cpp +++ b/src/core/CL/kernels/CLArgMinMaxLayerKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019 ARM Limited. + * Copyright (c) 2019-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -46,7 +46,7 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *prev_outp { ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, output); ARM_COMPUTE_RETURN_ERROR_ON_F16_UNSUPPORTED(input); - ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::S32, DataType::F16, DataType::F32); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QASYMM8, DataType::S32, DataType::F16, DataType::F32); ARM_COMPUTE_RETURN_ERROR_ON_MSG(op != ReductionOperation::ARG_IDX_MAX && op != ReductionOperation::ARG_IDX_MIN, "Only ARG_IDX_MAX and ARG_IDX_MIN are supported"); ARM_COMPUTE_RETURN_ERROR_ON_MSG(axis >= TensorShape::num_max_dimensions, "Reduction axis greater than max number of dimensions"); ARM_COMPUTE_RETURN_ERROR_ON_MSG(axis > 3, "Unsupported reduction axis"); @@ -128,17 +128,14 @@ void CLArgMinMaxLayerKernel::configure(const ICLTensor *input, const ICLTensor * _op = op; // Set build options - CLBuildOptions build_opts; - const std::string data_type_promoted = get_cl_type_from_data_type(input->info()->data_type()); + CLBuildOptions build_opts; build_opts.add_option_if(_prev_output != nullptr, "-DPREV_OUTPUT"); build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())); - build_opts.add_option("-DDATA_TYPE_PROMOTED=" + data_type_promoted); build_opts.add_option_if(is_data_type_float(input->info()->data_type()), "-DFLOAT_DATA_TYPE"); - build_opts.add_option_if(op == ReductionOperation::ARG_IDX_MAX, "-DARG_MAX"); - build_opts.add_option_if(op == ReductionOperation::ARG_IDX_MIN, "-DARG_MIN"); - build_opts.add_option("-DCOND_DATA_TYPE=" + get_cl_select_type_from_data_type(input->info()->data_type())); + build_opts.add_option_if_else(op == ReductionOperation::ARG_IDX_MAX, "-DARG_MAX", "-DARG_MIN"); build_opts.add_option("-DDATA_TYPE_OUTPUT=" + get_cl_type_from_data_type(output->info()->data_type())); + build_opts.add_option("-DDATA_TYPE_SELECT=" + get_cl_signed_type_from_element_size(input->info()->element_size())); // Create kernel cl::NDRange lws_hint = CLKernelLibrary::get().default_ndrange(); diff --git a/tests/validation/CL/ArgMinMax.cpp b/tests/validation/CL/ArgMinMax.cpp index 275641cb35..4bee942d70 100644 --- a/tests/validation/CL/ArgMinMax.cpp +++ b/tests/validation/CL/ArgMinMax.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -83,26 +83,6 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip( // clang-format on // *INDENT-ON* -DATA_TEST_CASE(Configuration, - framework::DatasetMode::ALL, - combine(datasets::SmallShapes(), framework::dataset::make("DataType", { DataType::F16, DataType::F32 })), - shape, data_type) -{ - // Create tensors - CLTensor ref_src = create_tensor(shape, data_type); - CLTensor dst; - - constexpr int axis = 1; - - // Create and Configure function - CLArgMinMaxLayer arg_min_max_layer; - arg_min_max_layer.configure(&ref_src, axis, &dst, ReductionOperation::ARG_IDX_MAX); - - const auto output_shape = arm_compute::misc::shape_calculator::compute_reduced_shape(shape, axis, false); - const ValidRegion valid_region = shape_to_valid_region(output_shape); - validate(dst.info()->valid_region(), valid_region); -} - template using CLArgMinMaxValidationFixture = ArgMinMaxValidationFixture; @@ -172,6 +152,35 @@ FIXTURE_DATA_TEST_CASE(RunLarge, } TEST_SUITE_END() // FP32 TEST_SUITE_END() // Float + +template +using CLArgMinMaxQuantizedValidationFixture = ArgMinMaxValidationQuantizedFixture; + +TEST_SUITE(Quantized) +TEST_SUITE(QASYMM8) +FIXTURE_DATA_TEST_CASE(RunSmall, + CLArgMinMaxQuantizedValidationFixture, + framework::DatasetMode::PRECOMMIT, + combine(combine(combine(combine(ArgMinMaxSmallDataset, framework::dataset::make("DataType", DataType::QASYMM8)), framework::dataset::make("Axis", { 0, 1, 2, 3 })), + framework::dataset::make("Operation", { ReductionOperation::ARG_IDX_MIN, ReductionOperation::ARG_IDX_MAX })), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(5.f / 255.f, 20) }))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} + +FIXTURE_DATA_TEST_CASE(RunLarge, + CLArgMinMaxQuantizedValidationFixture, + framework::DatasetMode::NIGHTLY, + combine(combine(combine(combine(ArgMinMaxLargeDataset, framework::dataset::make("DataType", DataType::QASYMM8)), framework::dataset::make("Axis", { 0, 1, 2, 3 })), + framework::dataset::make("Operation", { ReductionOperation::ARG_IDX_MIN, ReductionOperation::ARG_IDX_MAX })), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(5.f / 255.f, 20) }))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() // QASYMM8 +TEST_SUITE_END() // Quantized TEST_SUITE_END() // ArgMinMax TEST_SUITE_END() // CL } // namespace validation diff --git a/tests/validation/NEON/ArgMinMax.cpp b/tests/validation/NEON/ArgMinMax.cpp index 174cb19b42..ec90ab0db0 100644 --- a/tests/validation/NEON/ArgMinMax.cpp +++ b/tests/validation/NEON/ArgMinMax.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -70,26 +70,6 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip( // clang-format on // *INDENT-ON* -DATA_TEST_CASE(Configuration, - framework::DatasetMode::ALL, - combine(datasets::SmallShapes(), framework::dataset::make("DataType", { DataType::F32 })), - shape, data_type) -{ - // Create tensors - Tensor ref_src = create_tensor(shape, data_type); - Tensor dst; - const int axis = 1; - - // Create and Configure function - NEArgMinMaxLayer arg_min_max_layer; - arg_min_max_layer.configure(&ref_src, axis, &dst, ReductionOperation::ARG_IDX_MAX); - - // Validate valid region - const auto expected_output_shape = arm_compute::misc::shape_calculator::compute_reduced_shape(shape, axis, false); - const ValidRegion valid_region = shape_to_valid_region(expected_output_shape); - validate(dst.info()->valid_region(), valid_region); -} - template using NEArgMinMaxValidationFixture = ArgMinMaxValidationFixture; -- cgit v1.2.1