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 ++++++++++++++++++++++++-- 2 files changed, 95 insertions(+), 7 deletions(-) (limited to 'src/core/CL') 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)); -- cgit v1.2.1