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 --- 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 +- 5 files changed, 213 insertions(+), 38 deletions(-) (limited to 'src') 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); -- cgit v1.2.1