From ef4b4ae784f7533ed6d9e7b51827a894c32ed48e Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Tue, 4 Jul 2017 17:19:43 +0100 Subject: COMPMID-438: Add support for floating point Min-Max Location layer. Change-Id: I84ae564a40fc7320a6f94a84d53906ba51404f51 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/79797 Reviewed-by: Anthony Barbier Tested-by: Kaizen --- .../core/CL/kernels/CLMinMaxLocationKernel.h | 8 +- .../core/NEON/kernels/NEMinMaxLocationKernel.h | 35 ++--- .../runtime/CL/functions/CLMinMaxLocation.h | 12 +- .../runtime/NEON/functions/NEMinMaxLocation.h | 8 +- src/core/CL/cl_kernels/minmaxloc.cl | 35 ++++- src/core/CL/kernels/CLMinMaxLocationKernel.cpp | 67 +++++++++- src/core/NEON/kernels/NEMinMaxLocationKernel.cpp | 141 +++++++++++++++++---- src/runtime/CL/functions/CLMinMaxLocation.cpp | 6 +- src/runtime/NEON/functions/NEMinMaxLocation.cpp | 2 +- tests/validation/CL/MinMaxLocation.cpp | 133 ++++++++++++++++--- tests/validation/NEON/MinMaxLocation.cpp | 75 ++++++++--- tests/validation/Reference.cpp | 2 +- tests/validation/Reference.h | 2 +- tests/validation/ReferenceCPP.cpp | 2 +- tests/validation/ReferenceCPP.h | 2 +- tests/validation/TensorOperations.h | 46 ++++--- tests/validation/TensorVisitors.h | 6 +- tests/validation/Validation.cpp | 37 ------ tests/validation/Validation.h | 36 +++++- 19 files changed, 487 insertions(+), 168 deletions(-) diff --git a/arm_compute/core/CL/kernels/CLMinMaxLocationKernel.h b/arm_compute/core/CL/kernels/CLMinMaxLocationKernel.h index 6a31f3cf18..03ae3c1b1f 100644 --- a/arm_compute/core/CL/kernels/CLMinMaxLocationKernel.h +++ b/arm_compute/core/CL/kernels/CLMinMaxLocationKernel.h @@ -51,8 +51,8 @@ public: CLMinMaxKernel &operator=(CLMinMaxKernel &&) = default; /** Initialise the kernel's input and output. * - * @param[in] input Input Image. Data types supported: U8 or S16. - * @param[out] min_max Buffer of 2 elements to store the min value at position 0 and the max value at position 1. Data type supported: S32. + * @param[in] input Input Image. Data types supported: U8/S16/F32. + * @param[out] min_max Buffer of 2 elements to store the min value at position 0 and the max value at position 1. Data type supported: S32 if input type is U8/S16, F32 if input type is F32. */ void configure(const ICLImage *input, cl::Buffer *min_max); @@ -84,8 +84,8 @@ public: * * @note When locations of min and max occurrences are requested, the reported number of locations is limited to the given array size. * - * @param[in] input Input image. Data types supported: U8 or S16. - * @param[in] min_max Buffer of 2 elements which contains the min value at position 0 and the max value at position 1. Data type supported: S32 + * @param[in] input Input image. Data types supported: U8/S16/F32. + * @param[out] min_max Buffer of 2 elements to store the min value at position 0 and the max value at position 1. Data type supported: S32 if input type is U8/S16, F32 if input type is F32. * @param[out] min_max_count Buffer of 2 elements to store the min value occurrences at position 0 and the max value occurrences at position 1. Data type supported: S32 * @param[out] min_loc (Optional) Array of Coordinates2D used to store minimum value locations. * @param[out] max_loc (Optional) Array of Coordinates2D used to store maximum value locations. diff --git a/arm_compute/core/NEON/kernels/NEMinMaxLocationKernel.h b/arm_compute/core/NEON/kernels/NEMinMaxLocationKernel.h index b5375f613b..7b2748deee 100644 --- a/arm_compute/core/NEON/kernels/NEMinMaxLocationKernel.h +++ b/arm_compute/core/NEON/kernels/NEMinMaxLocationKernel.h @@ -54,11 +54,11 @@ public: /** Initialise the kernel's input and outputs. * - * @param[in] input Input Image. Data types supported: U8/S16. - * @param[out] min Minimum value of image. - * @param[out] max Maximum value of image. + * @param[in] input Input Image. Data types supported: U8/S16/F32. + * @param[out] min Minimum value of image. Data types supported: S32 if input type is U8/S16, F32 if input type is F32. + * @param[out] max Maximum value of image. Data types supported: S32 if input type is U8/S16, F32 if input type is F32. */ - void configure(const IImage *input, int32_t *min, int32_t *max); + void configure(const IImage *input, void *min, void *max); /** Resets global minimum and maximum. */ void reset(); @@ -76,6 +76,11 @@ private: * @param win The window to run the algorithm on. */ void minmax_S16(Window win); + /** Performs the min/max algorithm on F32 images on a given window. + * + * @param win The window to run the algorithm on. + */ + void minmax_F32(Window win); /** Common signature for all the specialised MinMax functions * * @param[in] window Region on which to execute the kernel. @@ -87,12 +92,10 @@ private: template void update_min_max(T min, T max); - const IImage *_input; /**< Input image. */ - int32_t *_min; /**< Minimum value. */ - int32_t *_max; /**< Maximum value. */ - int32_t _min_init; /**< Value to initialise global minimum value. */ - int32_t _max_init; /**< Value to initialise global maximum value. */ - std::mutex _mtx; /**< Mutex used for result reduction. */ + const IImage *_input; /**< Input image. */ + void *_min; /**< Minimum value. */ + void *_max; /**< Maximum value. */ + std::mutex _mtx; /**< Mutex used for result reduction. */ }; /** Interface for the kernel to find min max locations of an image. */ @@ -114,15 +117,15 @@ public: /** Initialise the kernel's input and outputs. * - * @param[in] input Input Image. Data types supported: U8 or S16. - * @param[out] min Minimum value of image. - * @param[out] max Maximum value of image. + * @param[in] input Input Image. Data types supported: U8/S16/F32. + * @param[out] min Minimum value of image. Data types supported: S32 if input type is U8/S16, F32 if input type is F32. + * @param[out] max Maximum value of image. Data types supported: S32 if input type is U8/S16, F32 if input type is F32. * @param[out] min_loc Array of minimum value locations. * @param[out] max_loc Array of maximum value locations. * @param[out] min_count Number of minimum value encounters. * @param[out] max_count Number of maximum value encounters. */ - void configure(const IImage *input, int32_t *min, int32_t *max, + void configure(const IImage *input, void *min, void *max, ICoordinates2DArray *min_loc = nullptr, ICoordinates2DArray *max_loc = nullptr, uint32_t *min_count = nullptr, uint32_t *max_count = nullptr); @@ -149,8 +152,8 @@ private: struct create_func_table; const IImage *_input; /**< Input image. */ - int32_t *_min; /**< Minimum value. */ - int32_t *_max; /**< Maximum value. */ + void *_min; /**< Minimum value. */ + void *_max; /**< Maximum value. */ uint32_t *_min_count; /**< Count of minimum value encounters. */ uint32_t *_max_count; /**< Count of maximum value encounters. */ ICoordinates2DArray *_min_loc; /**< Locations of minimum values. */ diff --git a/arm_compute/runtime/CL/functions/CLMinMaxLocation.h b/arm_compute/runtime/CL/functions/CLMinMaxLocation.h index 84fd67515b..52a2857d0b 100644 --- a/arm_compute/runtime/CL/functions/CLMinMaxLocation.h +++ b/arm_compute/runtime/CL/functions/CLMinMaxLocation.h @@ -55,15 +55,15 @@ public: * * @note When locations of min and max occurrences are requested, the reported number of locations is limited to the given array size. * - * @param[in] input Input image. Data types supported: U8 or S16. - * @param[out] min Minimum value of image. - * @param[out] max Maximum value of image. + * @param[in] input Input image. Data types supported: U8/S16/F32. + * @param[out] min Minimum value of image. Data types supported: S32 if input type is U8/S16, F32 if input type is F32. + * @param[out] max Maximum value of image. Data types supported: S32 if input type is U8/S16, F32 if input type is F32. * @param[out] min_loc (Optional) Array of Coordinates2D used to store minimum value locations. * @param[out] max_loc (Optional) Array of Coordinates2D used to store maximum value locations. * @param[out] min_count (Optional) Number of minimum value encounters. * @param[out] max_count (Optional) Number of maximum value encounters. */ - void configure(const ICLImage *input, int32_t *min, int32_t *max, + void configure(const ICLImage *input, void *min, void *max, CLCoordinates2DArray *min_loc = nullptr, CLCoordinates2DArray *max_loc = nullptr, uint32_t *min_count = nullptr, uint32_t *max_count = nullptr); @@ -75,8 +75,8 @@ private: CLMinMaxLocationKernel _min_max_loc_kernel; /**< Kernel that counts min/max occurrences and identifies their positions */ cl::Buffer _min_max_vals; /**< Buffer to collect min, max values */ cl::Buffer _min_max_count_vals; /**< Buffer to collect min, max values */ - int32_t *_min; /**< Minimum value. */ - int32_t *_max; /**< Maximum value. */ + void *_min; /**< Minimum value. */ + void *_max; /**< Maximum value. */ uint32_t *_min_count; /**< Minimum value occurrences. */ uint32_t *_max_count; /**< Maximum value occurrences. */ CLCoordinates2DArray *_min_loc; /**< Minimum value occurrences coordinates. */ diff --git a/arm_compute/runtime/NEON/functions/NEMinMaxLocation.h b/arm_compute/runtime/NEON/functions/NEMinMaxLocation.h index 82e75ee48b..e658d22a66 100644 --- a/arm_compute/runtime/NEON/functions/NEMinMaxLocation.h +++ b/arm_compute/runtime/NEON/functions/NEMinMaxLocation.h @@ -48,15 +48,15 @@ public: NEMinMaxLocation(); /** Initialise the kernel's inputs and outputs. * - * @param[in] input Input image. Data types supported: U8/S16. - * @param[out] min Minimum value of image. - * @param[out] max Maximum value of image. + * @param[in] input Input image. Data types supported: U8/S16/F32. + * @param[out] min Minimum value of image. Data types supported: S32 if input type is U8/S16, F32 if input type is F32. + * @param[out] max Maximum value of image. Data types supported: S32 if input type is U8/S16, F32 if input type is F32. * @param[out] min_loc (Optional) Array of minimum value locations. * @param[out] max_loc (Optional) Array of maximum value locations. * @param[out] min_count (Optional) Number of minimum value encounters. * @param[out] max_count (Optional) Number of maximum value encounters. */ - void configure(const IImage *input, int32_t *min, int32_t *max, + void configure(const IImage *input, void *min, void *max, ICoordinates2DArray *min_loc = nullptr, ICoordinates2DArray *max_loc = nullptr, uint32_t *min_count = nullptr, uint32_t *max_count = nullptr); diff --git a/src/core/CL/cl_kernels/minmaxloc.cl b/src/core/CL/cl_kernels/minmaxloc.cl index e628e9bd5b..05fc78dcb5 100644 --- a/src/core/CL/cl_kernels/minmaxloc.cl +++ b/src/core/CL/cl_kernels/minmaxloc.cl @@ -32,6 +32,17 @@ #define DATA_TYPE_MAX 0xFF #endif /* DATA_TYPE_MAX */ +inline int FloatFlip(float val) +{ + union + { + int int_val; + float flt_val; + } u_val; + u_val.flt_val = val; + return (u_val.int_val >= 0) ? u_val.int_val : u_val.int_val ^ 0x7FFFFFFF; +} + __constant VEC_DATA_TYPE(DATA_TYPE, 16) type_min = (VEC_DATA_TYPE(DATA_TYPE, 16))(DATA_TYPE_MIN); __constant VEC_DATA_TYPE(DATA_TYPE, 16) type_max = (VEC_DATA_TYPE(DATA_TYPE, 16))(DATA_TYPE_MAX); __constant uint16 idx16 = (uint16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15); @@ -78,8 +89,12 @@ __kernel void minmax( // Handle non multiple of 16 VEC_DATA_TYPE(DATA_TYPE, 16) data = vload16(0, (__global DATA_TYPE *)offset(&src, width4 << 4, 0)); +#ifdef IS_DATA_TYPE_FLOAT + int16 widx = convert_int16(((uint16)(width4 << 4) + idx16) < width); +#else /* IS_DATA_TYPE_FLOAT */ VEC_DATA_TYPE(DATA_TYPE, 16) - widx = CONVERT(((uint16)(width4 << 4) + idx16) < width, VEC_DATA_TYPE(DATA_TYPE, 16)); + widx = CONVERT(((uint16)(width4 << 4) + idx16) < width, VEC_DATA_TYPE(DATA_TYPE, 16)); +#endif /* IS_DATA_TYPE_FLOAT */ local_max = max(local_max, select(type_min, data, widx)); local_min = min(local_min, select(type_max, data, widx)); #endif /* NON_MULTIPLE_OF_16 */ @@ -98,8 +113,13 @@ __kernel void minmax( local_max.s0 = max(local_max.s0, local_max.s1); // Update global min/max +#ifdef IS_DATA_TYPE_FLOAT + atomic_min(&min_max[0], FloatFlip(local_min.s0)); + atomic_max(&min_max[1], FloatFlip(local_max.s0)); +#else /* IS_DATA_TYPE_FLOAT */ atomic_min(&min_max[0], local_min.s0); atomic_max(&min_max[1], local_max.s0); +#endif /* IS_DATA_TYPE_FLOAT */ } /** This function counts the min and max occurrences in an image and tags their position. @@ -136,9 +156,18 @@ __kernel void minmaxloc( { Image src = CONVERT_TO_IMAGE_STRUCT(src); +#ifdef IS_DATA_TYPE_FLOAT + __global float *min_max_ptr = (__global float *)min_max; + float min_value = min_max_ptr[0]; + float max_value = min_max_ptr[1]; +#else /* IS_DATA_TYPE_FLOAT */ + int min_value = min_max[0]; + int max_value = min_max[1]; +#endif /* IS_DATA_TYPE_FLOAT */ + DATA_TYPE value = *((__global DATA_TYPE *)src.ptr); #ifdef COUNT_MIN_MAX - if(value == min_max[0]) + if(value == min_value) { uint idx = atomic_inc(&min_max_count[0]); #ifdef LOCATE_MIN @@ -149,7 +178,7 @@ __kernel void minmaxloc( } #endif /* LOCATE_MIN */ } - if(value == min_max[1]) + if(value == max_value) { uint idx = atomic_inc(&min_max_count[1]); #ifdef LOCATE_MAX diff --git a/src/core/CL/kernels/CLMinMaxLocationKernel.cpp b/src/core/CL/kernels/CLMinMaxLocationKernel.cpp index 8a493209ca..be0034382c 100644 --- a/src/core/CL/kernels/CLMinMaxLocationKernel.cpp +++ b/src/core/CL/kernels/CLMinMaxLocationKernel.cpp @@ -34,6 +34,26 @@ using namespace arm_compute; +inline int32_t FloatFlip(float val) +{ + static_assert(sizeof(float) == sizeof(int32_t), "Float must be same size as int32_t"); + int32_t int_val = 0; + + memcpy(&int_val, &val, sizeof(float)); + int_val = (int_val >= 0) ? int_val : int_val ^ 0x7FFFFFFF; + return int_val; +} + +inline float IFloatFlip(int32_t val) +{ + static_assert(sizeof(float) == sizeof(int32_t), "Float must be same size as int32_t"); + float flt_val = 0.f; + + val = (val >= 0) ? val : val ^ 0x7FFFFFFF; + memcpy(&flt_val, &val, sizeof(float)); + return flt_val; +} + CLMinMaxKernel::CLMinMaxKernel() : _input(nullptr), _min_max(), _data_type_max_min() { @@ -41,7 +61,7 @@ CLMinMaxKernel::CLMinMaxKernel() void CLMinMaxKernel::configure(const ICLImage *input, cl::Buffer *min_max) { - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8, DataType::S16); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8, DataType::S16, DataType::F32); ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D(input); ARM_COMPUTE_ERROR_ON(min_max == nullptr); @@ -59,6 +79,10 @@ void CLMinMaxKernel::configure(const ICLImage *input, cl::Buffer *min_max) _data_type_max_min[0] = SHRT_MAX; _data_type_max_min[1] = SHRT_MIN; break; + case DataType::F32: + _data_type_max_min[0] = FloatFlip(std::numeric_limits::max()); + _data_type_max_min[1] = FloatFlip(std::numeric_limits::lowest()); + break; default: ARM_COMPUTE_ERROR("You called with the wrong image data types"); } @@ -66,9 +90,18 @@ void CLMinMaxKernel::configure(const ICLImage *input, cl::Buffer *min_max) // Set kernel build options std::set build_opts; build_opts.emplace("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())); - build_opts.emplace("-DDATA_TYPE_MAX=" + support::cpp11::to_string(_data_type_max_min[0])); - build_opts.emplace("-DDATA_TYPE_MIN=" + support::cpp11::to_string(_data_type_max_min[1])); build_opts.emplace((0 != (num_elems_processed_per_iteration % max_cl_vector_width)) ? "-DNON_MULTIPLE_OF_16" : ""); + if(input->info()->data_type() == DataType::F32) + { + build_opts.emplace("-DDATA_TYPE_MAX=" + support::cpp11::to_string(std::numeric_limits::max())); + build_opts.emplace("-DDATA_TYPE_MIN=" + support::cpp11::to_string(std::numeric_limits::lowest())); + build_opts.emplace("-DIS_DATA_TYPE_FLOAT"); + } + else + { + build_opts.emplace("-DDATA_TYPE_MAX=" + support::cpp11::to_string(_data_type_max_min[0])); + build_opts.emplace("-DDATA_TYPE_MIN=" + support::cpp11::to_string(_data_type_max_min[1])); + } // Create kernel _kernel = static_cast(CLKernelLibrary::get().create_kernel("minmax", build_opts)); @@ -100,6 +133,28 @@ void CLMinMaxKernel::run(const Window &window, cl::CommandQueue &queue) enqueue(queue, *this, slice); } while(window.slide_window_slice_2D(slice)); + + cl_int min = 0; + cl_int max = 0; + queue.enqueueReadBuffer(*_min_max, CL_TRUE /* blocking */, 0 * sizeof(cl_int), sizeof(cl_int), static_cast(&min)); + queue.enqueueReadBuffer(*_min_max, CL_TRUE /* blocking */, 1 * sizeof(cl_int), sizeof(cl_int), static_cast(&max)); + + if(_input->info()->data_type() == DataType::F32) + { + std::array min_max = + { + { + IFloatFlip(min), + IFloatFlip(max) + } + }; + queue.enqueueWriteBuffer(*_min_max, CL_TRUE /* blocking */, 0, min_max.size() * sizeof(float), min_max.data()); + } + else + { + std::array min_max = { { min, max } }; + queue.enqueueWriteBuffer(*_min_max, CL_TRUE /* blocking */, 0, min_max.size() * sizeof(int32_t), min_max.data()); + } } CLMinMaxLocationKernel::CLMinMaxLocationKernel() @@ -109,7 +164,7 @@ CLMinMaxLocationKernel::CLMinMaxLocationKernel() void CLMinMaxLocationKernel::configure(const ICLImage *input, cl::Buffer *min_max, cl::Buffer *min_max_count, ICLCoordinates2DArray *min_loc, ICLCoordinates2DArray *max_loc) { - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8, DataType::S16); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8, DataType::S16, DataType::F32); ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D(input); ARM_COMPUTE_ERROR_ON(min_max == nullptr); ARM_COMPUTE_ERROR_ON(min_max_count == nullptr && min_loc == nullptr && max_loc == nullptr); @@ -123,6 +178,10 @@ void CLMinMaxLocationKernel::configure(const ICLImage *input, cl::Buffer *min_ma build_opts.emplace((min_max_count != nullptr) ? "-DCOUNT_MIN_MAX" : ""); build_opts.emplace((min_loc != nullptr) ? "-DLOCATE_MIN" : ""); build_opts.emplace((max_loc != nullptr) ? "-DLOCATE_MAX" : ""); + if(input->info()->data_type() == DataType::F32) + { + build_opts.emplace("-DIS_DATA_TYPE_FLOAT"); + } // Create kernel _kernel = static_cast(CLKernelLibrary::get().create_kernel("minmaxloc", build_opts)); diff --git a/src/core/NEON/kernels/NEMinMaxLocationKernel.cpp b/src/core/NEON/kernels/NEMinMaxLocationKernel.cpp index 1e41ddcf80..a6da7f415d 100644 --- a/src/core/NEON/kernels/NEMinMaxLocationKernel.cpp +++ b/src/core/NEON/kernels/NEMinMaxLocationKernel.cpp @@ -41,14 +41,14 @@ namespace arm_compute { NEMinMaxKernel::NEMinMaxKernel() - : _func(), _input(nullptr), _min(), _max(), _min_init(), _max_init(), _mtx() + : _func(), _input(nullptr), _min(), _max(), _mtx() { } -void NEMinMaxKernel::configure(const IImage *input, int32_t *min, int32_t *max) +void NEMinMaxKernel::configure(const IImage *input, void *min, void *max) { ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D(input); - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8, DataType::S16); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8, DataType::S16, DataType::F32); ARM_COMPUTE_ERROR_ON(nullptr == min); ARM_COMPUTE_ERROR_ON(nullptr == max); @@ -56,20 +56,19 @@ void NEMinMaxKernel::configure(const IImage *input, int32_t *min, int32_t *max) _min = min; _max = max; - switch(input->info()->format()) + switch(_input->info()->data_type()) { - case Format::U8: - _min_init = UCHAR_MAX; - _max_init = 0; - _func = &NEMinMaxKernel::minmax_U8; + case DataType::U8: + _func = &NEMinMaxKernel::minmax_U8; break; - case Format::S16: - _min_init = SHRT_MAX; - _max_init = SHRT_MIN; - _func = &NEMinMaxKernel::minmax_S16; + case DataType::S16: + _func = &NEMinMaxKernel::minmax_S16; + break; + case DataType::F32: + _func = &NEMinMaxKernel::minmax_F32; break; default: - ARM_COMPUTE_ERROR("You called with the wrong img formats"); + ARM_COMPUTE_ERROR("Unsupported data type"); break; } @@ -93,8 +92,24 @@ void NEMinMaxKernel::run(const Window &window) void NEMinMaxKernel::reset() { ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); - *_min = _min_init; - *_max = _max_init; + switch(_input->info()->data_type()) + { + case DataType::U8: + *static_cast(_min) = UCHAR_MAX; + *static_cast(_max) = 0; + break; + case DataType::S16: + *static_cast(_min) = SHRT_MAX; + *static_cast(_max) = SHRT_MIN; + break; + case DataType::F32: + *static_cast(_min) = std::numeric_limits::max(); + *static_cast(_max) = std::numeric_limits::lowest(); + break; + default: + ARM_COMPUTE_ERROR("Unsupported data type"); + break; + } } template @@ -102,14 +117,19 @@ void NEMinMaxKernel::update_min_max(const T min, const T max) { std::lock_guard lock(_mtx); - if(min < *_min) + using type = typename std::conditional::value, float, int32_t>::type; + + auto min_ptr = static_cast(_min); + auto max_ptr = static_cast(_max); + + if(min < *min_ptr) { - *_min = min; + *min_ptr = min; } - if(max > *_max) + if(max > *max_ptr) { - *_max = max; + *max_ptr = max; } } @@ -229,6 +249,65 @@ void NEMinMaxKernel::minmax_S16(Window win) update_min_max(min_i, max_i); } +void NEMinMaxKernel::minmax_F32(Window win) +{ + float32x2_t carry_min = vdup_n_f32(std::numeric_limits::max()); + float32x2_t carry_max = vdup_n_f32(std::numeric_limits::lowest()); + + float carry_min_scalar = std::numeric_limits::max(); + float carry_max_scalar = std::numeric_limits::lowest(); + + const int x_start = win.x().start(); + const int x_end = win.x().end(); + + // Handle X dimension manually to split into two loops + // First one will use vector operations, second one processes the left over pixels + win.set(Window::DimX, Window::Dimension(0, 1, 1)); + + Iterator input(_input, win); + + execute_window_loop(win, [&](const Coordinates & id) + { + int x = x_start; + const auto in_ptr = reinterpret_cast(input.ptr()); + + // Vector loop + for(; x <= x_end - 8; x += 8) + { + const float32x4x2_t pixels = vld2q_f32(in_ptr + x); + const float32x4_t tmp_min1 = vminq_f32(pixels.val[0], pixels.val[1]); + const float32x4_t tmp_max1 = vmaxq_f32(pixels.val[0], pixels.val[1]); + const float32x2_t tmp_min2 = vmin_f32(vget_high_f32(tmp_min1), vget_low_f32(tmp_min1)); + const float32x2_t tmp_max2 = vmax_f32(vget_high_f32(tmp_max1), vget_low_f32(tmp_max1)); + carry_min = vmin_f32(tmp_min2, carry_min); + carry_max = vmax_f32(tmp_max2, carry_max); + } + + // Process leftover pixels + for(; x < x_end; ++x) + { + const float pixel = in_ptr[x]; + carry_min_scalar = std::min(pixel, carry_min_scalar); + carry_max_scalar = std::max(pixel, carry_max_scalar); + } + + }, + input); + + // Reduce result + carry_min = vpmin_f32(carry_min, carry_min); + carry_max = vpmax_f32(carry_max, carry_max); + carry_min = vpmin_f32(carry_min, carry_min); + carry_max = vpmax_f32(carry_max, carry_max); + + // Extract max/min values + const float min_i = std::min(vget_lane_f32(carry_min, 0), carry_min_scalar); + const float max_i = std::max(vget_lane_f32(carry_max, 0), carry_max_scalar); + + // Perform reduction of local min/max values + update_min_max(min_i, max_i); +} + NEMinMaxLocationKernel::NEMinMaxLocationKernel() : _func(nullptr), _input(nullptr), _min(nullptr), _max(nullptr), _min_count(nullptr), _max_count(nullptr), _min_loc(nullptr), _max_loc(nullptr) { @@ -271,12 +350,12 @@ const NEMinMaxLocationKernel::MinMaxLocFunction NEMinMaxLocationKernel::create_f &NEMinMaxLocationKernel::minmax_loc... }; -void NEMinMaxLocationKernel::configure(const IImage *input, int32_t *min, int32_t *max, +void NEMinMaxLocationKernel::configure(const IImage *input, void *min, void *max, ICoordinates2DArray *min_loc, ICoordinates2DArray *max_loc, uint32_t *min_count, uint32_t *max_count) { ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D(input); - ARM_COMPUTE_ERROR_ON_FORMAT_NOT_IN(input, Format::U8, Format::S16); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8, DataType::S16, DataType::F32); ARM_COMPUTE_ERROR_ON(nullptr == min); ARM_COMPUTE_ERROR_ON(nullptr == max); @@ -295,16 +374,19 @@ void NEMinMaxLocationKernel::configure(const IImage *input, int32_t *min, int32_ unsigned int table_idx = (count_min << 3) | (count_max << 2) | (loc_min << 1) | loc_max; - switch(input->info()->format()) + switch(input->info()->data_type()) { - case Format::U8: + case DataType::U8: _func = create_func_table::type>::func_table[table_idx]; break; - case Format::S16: + case DataType::S16: _func = create_func_table::type>::func_table[table_idx]; break; + case DataType::F32: + _func = create_func_table::type>::func_table[table_idx]; + break; default: - ARM_COMPUTE_ERROR("You called with the wrong img formats"); + ARM_COMPUTE_ERROR("Unsupported data type"); break; } @@ -349,6 +431,11 @@ void NEMinMaxLocationKernel::minmax_loc(const Window &win) _max_loc->clear(); } + using type = typename std::conditional::value, float, int32_t>::type; + + auto min_ptr = static_cast(_min); + auto max_ptr = static_cast(_max); + execute_window_loop(win, [&](const Coordinates & id) { auto in_ptr = reinterpret_cast(input.ptr()); @@ -360,7 +447,7 @@ void NEMinMaxLocationKernel::minmax_loc(const Window &win) if(count_min || loc_min) { - if(*_min == pixel) + if(*min_ptr == pixel) { if(count_min) { @@ -376,7 +463,7 @@ void NEMinMaxLocationKernel::minmax_loc(const Window &win) if(count_max || loc_max) { - if(*_max == pixel) + if(*max_ptr == pixel) { if(count_max) { diff --git a/src/runtime/CL/functions/CLMinMaxLocation.cpp b/src/runtime/CL/functions/CLMinMaxLocation.cpp index ad783d8a53..bc70ceb621 100644 --- a/src/runtime/CL/functions/CLMinMaxLocation.cpp +++ b/src/runtime/CL/functions/CLMinMaxLocation.cpp @@ -41,7 +41,7 @@ CLMinMaxLocation::CLMinMaxLocation() { } -void CLMinMaxLocation::configure(const ICLImage *input, int32_t *min, int32_t *max, CLCoordinates2DArray *min_loc, CLCoordinates2DArray *max_loc, uint32_t *min_count, uint32_t *max_count) +void CLMinMaxLocation::configure(const ICLImage *input, void *min, void *max, CLCoordinates2DArray *min_loc, CLCoordinates2DArray *max_loc, uint32_t *min_count, uint32_t *max_count) { ARM_COMPUTE_ERROR_ON(nullptr == min); ARM_COMPUTE_ERROR_ON(nullptr == max); @@ -67,8 +67,8 @@ void CLMinMaxLocation::run() CLScheduler::get().enqueue(_min_max_loc_kernel, false); // Update min and max - q.enqueueReadBuffer(_min_max_vals, CL_FALSE, 0 * sizeof(int32_t), sizeof(int32_t), _min); - q.enqueueReadBuffer(_min_max_vals, CL_FALSE, 1 * sizeof(int32_t), sizeof(int32_t), _max); + q.enqueueReadBuffer(_min_max_vals, CL_FALSE, 0 * sizeof(int32_t), sizeof(int32_t), static_cast(_min)); + q.enqueueReadBuffer(_min_max_vals, CL_FALSE, 1 * sizeof(int32_t), sizeof(int32_t), static_cast(_max)); // Update min and max count if(_min_count != nullptr) diff --git a/src/runtime/NEON/functions/NEMinMaxLocation.cpp b/src/runtime/NEON/functions/NEMinMaxLocation.cpp index cab9200cf8..54e89abe24 100644 --- a/src/runtime/NEON/functions/NEMinMaxLocation.cpp +++ b/src/runtime/NEON/functions/NEMinMaxLocation.cpp @@ -32,7 +32,7 @@ NEMinMaxLocation::NEMinMaxLocation() { } -void NEMinMaxLocation::configure(const IImage *input, int32_t *min, int32_t *max, ICoordinates2DArray *min_loc, ICoordinates2DArray *max_loc, uint32_t *min_count, uint32_t *max_count) +void NEMinMaxLocation::configure(const IImage *input, void *min, void *max, ICoordinates2DArray *min_loc, ICoordinates2DArray *max_loc, uint32_t *min_count, uint32_t *max_count) { _min_max.configure(input, min, max); _min_max_loc.configure(input, min, max, min_loc, max_loc, min_count, max_count); diff --git a/tests/validation/CL/MinMaxLocation.cpp b/tests/validation/CL/MinMaxLocation.cpp index 0646ad9bbf..5f8be433cd 100644 --- a/tests/validation/CL/MinMaxLocation.cpp +++ b/tests/validation/CL/MinMaxLocation.cpp @@ -50,22 +50,27 @@ using namespace arm_compute::test::validation; namespace { /** Compute CL MinMaxLocation function. -* -* @param[in] shape Shape of the input and output tensors. -* @param[in] dt_in Data type of first input tensor. -* -* @return Computed output tensor. -*/ -void compute_min_max_location(const TensorShape &shape, DataType dt_in, int32_t &min, int32_t &max, + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt_in Data type of first input tensor. + * @param[out] min Minimum value of tensor + * @param[out] max Maximum value of tensor + * @param[out] min_loc Array with locations of minimum values + * @param[out] max_loc Array with locations of maximum values + * @param[out] min_count Number of minimum values found + * @param[out] max_count Number of maximum values found + * + * @return Computed output tensor. + */ +void compute_min_max_location(const TensorShape &shape, DataType dt_in, void *min, void *max, CLCoordinates2DArray &min_loc, CLCoordinates2DArray &max_loc, uint32_t &min_count, uint32_t &max_count) { // Create tensor CLTensor src = create_tensor(shape, dt_in); - src.info()->set_format((dt_in == DataType::U8) ? Format::U8 : Format::S16); // Create and configure min_max_location configure function CLMinMaxLocation min_max_loc; - min_max_loc.configure(&src, &min, &max, &min_loc, &max_loc, &min_count, &max_count); + min_max_loc.configure(&src, min, max, &min_loc, &max_loc, &min_count, &max_count); // Allocate tensors src.allocator()->allocate(); @@ -141,13 +146,13 @@ BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes(), uint32_t ref_max_count; // Compute function - compute_min_max_location(shape, DataType::U8, min, max, min_loc, max_loc, min_count, max_count); + compute_min_max_location(shape, DataType::U8, &min, &max, min_loc, max_loc, min_count, max_count); // Compute reference ref_min_loc.map(); ref_max_loc.map(); - Reference::compute_reference_min_max_location(shape, DataType::U8, ref_min, ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); + Reference::compute_reference_min_max_location(shape, DataType::U8, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); min_loc.map(); max_loc.map(); @@ -181,13 +186,13 @@ BOOST_DATA_TEST_CASE(RunLarge, Large2DShapes(), uint32_t ref_max_count; // Compute function - compute_min_max_location(shape, DataType::U8, min, max, min_loc, max_loc, min_count, max_count); + compute_min_max_location(shape, DataType::U8, &min, &max, min_loc, max_loc, min_count, max_count); // Compute reference ref_min_loc.map(); ref_max_loc.map(); - Reference::compute_reference_min_max_location(shape, DataType::U8, ref_min, ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); + Reference::compute_reference_min_max_location(shape, DataType::U8, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); min_loc.map(); max_loc.map(); @@ -234,13 +239,13 @@ BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes(), uint32_t ref_max_count; // Compute function - compute_min_max_location(shape, DataType::S16, min, max, min_loc, max_loc, min_count, max_count); + compute_min_max_location(shape, DataType::S16, &min, &max, min_loc, max_loc, min_count, max_count); // Compute reference ref_min_loc.map(); ref_max_loc.map(); - Reference::compute_reference_min_max_location(shape, DataType::S16, ref_min, ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); + Reference::compute_reference_min_max_location(shape, DataType::S16, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); min_loc.map(); max_loc.map(); @@ -274,13 +279,105 @@ BOOST_DATA_TEST_CASE(RunLarge, Large2DShapes(), uint32_t ref_max_count; // Compute function - compute_min_max_location(shape, DataType::S16, min, max, min_loc, max_loc, min_count, max_count); + compute_min_max_location(shape, DataType::S16, &min, &max, min_loc, max_loc, min_count, max_count); + + // Compute reference + ref_min_loc.map(); + ref_max_loc.map(); + + Reference::compute_reference_min_max_location(shape, DataType::S16, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); + + min_loc.map(); + max_loc.map(); + + // Validate output + validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count); + + ref_min_loc.unmap(); + ref_max_loc.unmap(); + min_loc.unmap(); + max_loc.unmap(); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (Small2DShapes() + Large2DShapes()), + shape) +{ + // Create tensor + CLTensor src = create_tensor(shape, DataType::F32); + + validate_configuration(src, shape); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes(), + shape) +{ + // Create output storage + float min; + float max; + CLCoordinates2DArray min_loc(shape.total_size()); + CLCoordinates2DArray max_loc(shape.total_size()); + uint32_t min_count; + uint32_t max_count; + + float ref_min; + float ref_max; + CLCoordinates2DArray ref_min_loc(shape.total_size()); + CLCoordinates2DArray ref_max_loc(shape.total_size()); + uint32_t ref_min_count; + uint32_t ref_max_count; + + // Compute function + compute_min_max_location(shape, DataType::F32, &min, &max, min_loc, max_loc, min_count, max_count); + + // Compute reference + ref_min_loc.map(); + ref_max_loc.map(); + + Reference::compute_reference_min_max_location(shape, DataType::F32, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); + + min_loc.map(); + max_loc.map(); + + // Validate output + validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count); + + ref_min_loc.unmap(); + ref_max_loc.unmap(); + min_loc.unmap(); + max_loc.unmap(); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, Large2DShapes(), + shape) +{ + // Create output storage + float min; + float max; + CLCoordinates2DArray min_loc(shape.total_size()); + CLCoordinates2DArray max_loc(shape.total_size()); + uint32_t min_count; + uint32_t max_count; + + float ref_min; + float ref_max; + CLCoordinates2DArray ref_min_loc(shape.total_size()); + CLCoordinates2DArray ref_max_loc(shape.total_size()); + uint32_t ref_min_count; + uint32_t ref_max_count; + + // Compute function + compute_min_max_location(shape, DataType::F32, &min, &max, min_loc, max_loc, min_count, max_count); // Compute reference ref_min_loc.map(); ref_max_loc.map(); - Reference::compute_reference_min_max_location(shape, DataType::S16, ref_min, ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); + Reference::compute_reference_min_max_location(shape, DataType::F32, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); min_loc.map(); max_loc.map(); @@ -297,4 +394,4 @@ BOOST_AUTO_TEST_SUITE_END() BOOST_AUTO_TEST_SUITE_END() BOOST_AUTO_TEST_SUITE_END() -#endif /* DOXYGEN_SKIP_THIS */ \ No newline at end of file +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation/NEON/MinMaxLocation.cpp b/tests/validation/NEON/MinMaxLocation.cpp index b1fa8cd7a9..a467172550 100644 --- a/tests/validation/NEON/MinMaxLocation.cpp +++ b/tests/validation/NEON/MinMaxLocation.cpp @@ -51,29 +51,28 @@ using namespace arm_compute::test::validation; namespace { /** Compute Neon MinMaxLocation function. - * - * @param[in] shape Shape of the input and output tensors. - * @param[in] dt_in Data type of first input tensor. - * @param[out] min Minimum value of tensor - * @param[out] max Maximum value of tensor - * @param[out] min_loc Array with locations of minimum values - * @param[out] max_loc Array with locations of maximum values - * @param[out] min_count Number of minimum values found - * @param[out] max_count Number of maximum values found - * - * @return Computed output tensor. - */ - -void compute_min_max_location(const TensorShape &shape, DataType dt_in, int32_t &min, int32_t &max, + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt_in Data type of first input tensor. + * @param[out] min Minimum value of tensor + * @param[out] max Maximum value of tensor + * @param[out] min_loc Array with locations of minimum values + * @param[out] max_loc Array with locations of maximum values + * @param[out] min_count Number of minimum values found + * @param[out] max_count Number of maximum values found + * + * @return Computed output tensor. + */ + +void compute_min_max_location(const TensorShape &shape, DataType dt_in, void *min, void *max, Coordinates2DArray &min_loc, Coordinates2DArray &max_loc, uint32_t &min_count, uint32_t &max_count) { // Create tensor Tensor src = create_tensor(shape, dt_in); - src.info()->set_format((dt_in == DataType::U8) ? Format::U8 : Format::S16); // Create and configure min_max_location configure function NEMinMaxLocation min_max_loc; - min_max_loc.configure(&src, &min, &max, &min_loc, &max_loc, &min_count, &max_count); + min_max_loc.configure(&src, min, max, &min_loc, &max_loc, &min_count, &max_count); // Allocate tensors src.allocator()->allocate(); @@ -117,6 +116,7 @@ void validate_configuration(const Tensor &src, TensorShape shape) BOOST_AUTO_TEST_SUITE(NEON) BOOST_AUTO_TEST_SUITE(MinMaxLocation) +BOOST_AUTO_TEST_SUITE(Integer) BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) BOOST_DATA_TEST_CASE(Configuration, (Small2DShapes() + Large2DShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }), shape, dt) @@ -148,10 +148,10 @@ BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes() * boost::unit_test::data::make({ uint32_t ref_max_count; // Compute function - compute_min_max_location(shape, dt, min, max, min_loc, max_loc, min_count, max_count); + compute_min_max_location(shape, dt, &min, &max, min_loc, max_loc, min_count, max_count); // Compute reference - Reference::compute_reference_min_max_location(shape, dt, ref_min, ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); + Reference::compute_reference_min_max_location(shape, dt, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); // Validate output validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count); @@ -177,15 +177,48 @@ BOOST_DATA_TEST_CASE(RunLarge, Large2DShapes() * boost::unit_test::data::make({ uint32_t ref_max_count; // Compute function - compute_min_max_location(shape, dt, min, max, min_loc, max_loc, min_count, max_count); + compute_min_max_location(shape, dt, &min, &max, min_loc, max_loc, min_count, max_count); + + // Compute reference + Reference::compute_reference_min_max_location(shape, dt, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); + + // Validate output + validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes() * DataType::F32, + shape, dt) +{ + // Create output storage + float min; + float max; + Coordinates2DArray min_loc(shape.total_size()); + Coordinates2DArray max_loc(shape.total_size()); + uint32_t min_count; + uint32_t max_count; + + float ref_min; + float ref_max; + Coordinates2DArray ref_min_loc(shape.total_size()); + Coordinates2DArray ref_max_loc(shape.total_size()); + uint32_t ref_min_count; + uint32_t ref_max_count; + + // Compute function + compute_min_max_location(shape, dt, &min, &max, min_loc, max_loc, min_count, max_count); // Compute reference - Reference::compute_reference_min_max_location(shape, dt, ref_min, ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); + Reference::compute_reference_min_max_location(shape, dt, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); // Validate output validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count); } +BOOST_AUTO_TEST_SUITE_END() + BOOST_AUTO_TEST_SUITE_END() BOOST_AUTO_TEST_SUITE_END() -#endif /* DOXYGEN_SKIP_THIS */ \ No newline at end of file +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation/Reference.cpp b/tests/validation/Reference.cpp index e9ddea78cb..f9052f1dba 100644 --- a/tests/validation/Reference.cpp +++ b/tests/validation/Reference.cpp @@ -71,7 +71,7 @@ std::pair Reference::compute_reference_sobel_5x5(const Ten return std::make_pair(ref_dst_x, ref_dst_y); } -void Reference::compute_reference_min_max_location(const TensorShape &shape, DataType dt_in, int32_t &min, int32_t &max, IArray &min_loc, IArray &max_loc, +void Reference::compute_reference_min_max_location(const TensorShape &shape, DataType dt_in, void *min, void *max, IArray &min_loc, IArray &max_loc, uint32_t &min_count, uint32_t &max_count) { // Create reference diff --git a/tests/validation/Reference.h b/tests/validation/Reference.h index 778e7b0b2b..eeaa55c739 100644 --- a/tests/validation/Reference.h +++ b/tests/validation/Reference.h @@ -72,7 +72,7 @@ public: * * @return Computed minimum, maximum values and their locations. */ - static void compute_reference_min_max_location(const TensorShape &shape, DataType dt_in, int32_t &min, int32_t &max, IArray &min_loc, IArray &max_loc, + static void compute_reference_min_max_location(const TensorShape &shape, DataType dt_in, void *min, void *max, IArray &min_loc, IArray &max_loc, uint32_t &min_count, uint32_t &max_count); /** Compute reference mean and standard deviation. diff --git a/tests/validation/ReferenceCPP.cpp b/tests/validation/ReferenceCPP.cpp index 13f4b90a82..81ec60d5b9 100644 --- a/tests/validation/ReferenceCPP.cpp +++ b/tests/validation/ReferenceCPP.cpp @@ -71,7 +71,7 @@ void ReferenceCPP::sobel_5x5(RawTensor &src, RawTensor &dst_x, RawTensor &dst_y, } // Minimum maximum location -void ReferenceCPP::min_max_location(const RawTensor &src, int32_t &min, int32_t &max, IArray &min_loc, IArray &max_loc, uint32_t &min_count, uint32_t &max_count) +void ReferenceCPP::min_max_location(const RawTensor &src, void *min, void *max, IArray &min_loc, IArray &max_loc, uint32_t &min_count, uint32_t &max_count) { const TensorVariant s = TensorFactory::get_tensor(src); boost::apply_visitor(tensor_visitors::min_max_location_visitor(min, max, min_loc, max_loc, min_count, max_count), s); diff --git a/tests/validation/ReferenceCPP.h b/tests/validation/ReferenceCPP.h index 3f5e4aeaf5..97e573cfa2 100644 --- a/tests/validation/ReferenceCPP.h +++ b/tests/validation/ReferenceCPP.h @@ -74,7 +74,7 @@ public: * @param[out] min_count Number of minimum values found * @param[out] max_count Number of maximum values found */ - static void min_max_location(const RawTensor &src, int32_t &min, int32_t &max, IArray &min_loc, IArray &max_loc, uint32_t &min_count, uint32_t &max_count); + static void min_max_location(const RawTensor &src, void *min, void *max, IArray &min_loc, IArray &max_loc, uint32_t &min_count, uint32_t &max_count); /** Function to compute the mean and standard deviation of a tensor. * * @param[in] src Input tensor. diff --git a/tests/validation/TensorOperations.h b/tests/validation/TensorOperations.h index a5039a4641..b472e3d5cf 100644 --- a/tests/validation/TensorOperations.h +++ b/tests/validation/TensorOperations.h @@ -305,34 +305,50 @@ void sobel_5x5(Tensor &in, Tensor &out_x, Tensor &out_y, BorderMode } } -// Min max location -template -void min_max_location(const Tensor &in, int32_t &min, int32_t &max, IArray &min_loc, IArray &max_loc, uint32_t &min_count, uint32_t &max_count) +template +void compute_min_max(const Tensor &in, void *min, void *max) { - // Set min and max to first pixel - min = in[0]; - max = in[0]; - min_count = 0; - max_count = 0; + using type = typename std::conditional::value, float, int32_t>::type; - const size_t width = in.shape().x(); + // Set min and max to first pixel + type tmp_min = static_cast(in[0]); + type tmp_max = static_cast(in[0]); // Look for min and max values for(int i = 1; i < in.num_elements(); ++i) { - if(static_cast(in[i]) < min) + if(static_cast(in[i]) < tmp_min) { - min = in[i]; + tmp_min = static_cast(in[i]); } - if(static_cast(in[i]) > max) + if(static_cast(in[i]) > tmp_max) { - max = in[i]; + tmp_max = static_cast(in[i]); } } + *static_cast(min) = tmp_min; + *static_cast(max) = tmp_max; +} + +// Min max location +template +void min_max_location(const Tensor &in, void *min, void *max, IArray &min_loc, IArray &max_loc, uint32_t &min_count, uint32_t &max_count) +{ + const size_t width = in.shape().x(); + + compute_min_max(in, min, max); + + using type = typename std::conditional::value, float, int32_t>::type; + + type min_value = *static_cast(min); + type max_value = *static_cast(max); + + min_count = 0; + max_count = 0; for(int i = 0; i < in.num_elements(); ++i) { - if(static_cast(in[i]) == min) + if(static_cast(in[i]) == min_value) { Coordinates2D min_coord; min_coord.x = static_cast(i % width); @@ -342,7 +358,7 @@ void min_max_location(const Tensor &in, int32_t &min, int32_t &max, IArray(in[i]) == max) + if(static_cast(in[i]) == max_value) { Coordinates2D max_coord; max_coord.x = static_cast(i % width); diff --git a/tests/validation/TensorVisitors.h b/tests/validation/TensorVisitors.h index fa9c3ecbb8..44ae6f13e8 100644 --- a/tests/validation/TensorVisitors.h +++ b/tests/validation/TensorVisitors.h @@ -49,7 +49,7 @@ namespace tensor_visitors struct min_max_location_visitor : public boost::static_visitor<> { public: - explicit min_max_location_visitor(int32_t &min, int32_t &max, IArray &min_loc, IArray &max_loc, uint32_t &min_count, uint32_t &max_count) + explicit min_max_location_visitor(void *min, void *max, IArray &min_loc, IArray &max_loc, uint32_t &min_count, uint32_t &max_count) : _min(min), _max(max), _min_loc(min_loc), _max_loc(max_loc), _min_count(min_count), _max_count(max_count) { } @@ -60,8 +60,8 @@ public: } private: - int32_t &_min; - int32_t &_max; + void *_min; + void *_max; IArray &_min_loc; IArray &_max_loc; uint32_t &_min_count; diff --git a/tests/validation/Validation.cpp b/tests/validation/Validation.cpp index a13eeb0b85..eac4105b21 100644 --- a/tests/validation/Validation.cpp +++ b/tests/validation/Validation.cpp @@ -403,43 +403,6 @@ void validate(float target, float ref, float tolerance_abs_error, float toleranc BOOST_TEST_INFO("target = " << std::setprecision(5) << target); BOOST_TEST(equal); } - -void validate_min_max_loc(int32_t min, int32_t ref_min, int32_t max, int32_t ref_max, - IArray &min_loc, IArray &ref_min_loc, IArray &max_loc, IArray &ref_max_loc, - uint32_t min_count, uint32_t ref_min_count, uint32_t max_count, uint32_t ref_max_count) -{ - BOOST_TEST(min == ref_min); - BOOST_TEST(max == ref_max); - - BOOST_TEST(min_count == min_loc.num_values()); - BOOST_TEST(max_count == max_loc.num_values()); - BOOST_TEST(ref_min_count == ref_min_loc.num_values()); - BOOST_TEST(ref_max_count == ref_max_loc.num_values()); - - BOOST_TEST(min_count == ref_min_count); - BOOST_TEST(max_count == ref_max_count); - - for(uint32_t i = 0; i < min_count; i++) - { - Coordinates2D *same_coords = std::find_if(ref_min_loc.buffer(), ref_min_loc.buffer() + min_count, [&min_loc, i](Coordinates2D coord) - { - return coord.x == min_loc.at(i).x && coord.y == min_loc.at(i).y; - }); - - BOOST_TEST(same_coords != ref_min_loc.buffer() + min_count); - } - - for(uint32_t i = 0; i < max_count; i++) - { - Coordinates2D *same_coords = std::find_if(ref_max_loc.buffer(), ref_max_loc.buffer() + max_count, [&max_loc, i](Coordinates2D coord) - { - return coord.x == max_loc.at(i).x && coord.y == max_loc.at(i).y; - }); - - BOOST_TEST(same_coords != ref_max_loc.buffer() + max_count); - } -} - } // namespace validation } // namespace test } // namespace arm_compute diff --git a/tests/validation/Validation.h b/tests/validation/Validation.h index 217ec63658..66bb2be2ca 100644 --- a/tests/validation/Validation.h +++ b/tests/validation/Validation.h @@ -133,10 +133,42 @@ void validate(float target, float ref, float tolerance_abs_error = std::numeric_ * * - All values should match */ -void validate_min_max_loc(int32_t min, int32_t ref_min, int32_t max, int32_t ref_max, +template +void validate_min_max_loc(T min, T ref_min, T max, T ref_max, IArray &min_loc, IArray &ref_min_loc, IArray &max_loc, IArray &ref_max_loc, - uint32_t min_count, uint32_t ref_min_count, uint32_t max_count, uint32_t ref_max_count); + uint32_t min_count, uint32_t ref_min_count, uint32_t max_count, uint32_t ref_max_count) +{ + BOOST_TEST(min == ref_min); + BOOST_TEST(max == ref_max); + + BOOST_TEST(min_count == min_loc.num_values()); + BOOST_TEST(max_count == max_loc.num_values()); + BOOST_TEST(ref_min_count == ref_min_loc.num_values()); + BOOST_TEST(ref_max_count == ref_max_loc.num_values()); + + BOOST_TEST(min_count == ref_min_count); + BOOST_TEST(max_count == ref_max_count); + + for(uint32_t i = 0; i < min_count; i++) + { + Coordinates2D *same_coords = std::find_if(ref_min_loc.buffer(), ref_min_loc.buffer() + min_count, [&min_loc, i](Coordinates2D coord) + { + return coord.x == min_loc.at(i).x && coord.y == min_loc.at(i).y; + }); + BOOST_TEST(same_coords != ref_min_loc.buffer() + min_count); + } + + for(uint32_t i = 0; i < max_count; i++) + { + Coordinates2D *same_coords = std::find_if(ref_max_loc.buffer(), ref_max_loc.buffer() + max_count, [&max_loc, i](Coordinates2D coord) + { + return coord.x == max_loc.at(i).x && coord.y == max_loc.at(i).y; + }); + + BOOST_TEST(same_coords != ref_max_loc.buffer() + max_count); + } +} } // namespace validation } // namespace test } // namespace arm_compute -- cgit v1.2.1