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 <anthony.barbier@arm.com>
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>

5 files changed, 213 insertions, 38 deletions

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<float>::max());
+            _data_type_max_min[1] = FloatFlip(std::numeric_limits<float>::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<std::string> 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<float>::max()));
+        build_opts.emplace("-DDATA_TYPE_MIN=" + support::cpp11::to_string(std::numeric_limits<float>::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<cl::Kernel>(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<int *>(&min));
+    queue.enqueueReadBuffer(*_min_max, CL_TRUE /* blocking */, 1 * sizeof(cl_int), sizeof(cl_int), static_cast<int *>(&max));
+
+    if(_input->info()->data_type() == DataType::F32)
+    {
+        std::array<float, 2> 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<int32_t, 2> 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<cl::Kernel>(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<int32_t *>(_min) = UCHAR_MAX;
+            *static_cast<int32_t *>(_max) = 0;
+            break;
+        case DataType::S16:
+            *static_cast<int32_t *>(_min) = SHRT_MAX;
+            *static_cast<int32_t *>(_max) = SHRT_MIN;
+            break;
+        case DataType::F32:
+            *static_cast<float *>(_min) = std::numeric_limits<float>::max();
+            *static_cast<float *>(_max) = std::numeric_limits<float>::lowest();
+            break;
+        default:
+            ARM_COMPUTE_ERROR("Unsupported data type");
+            break;
+    }
 }
 
 template <typename T>
@@ -102,14 +117,19 @@ void NEMinMaxKernel::update_min_max(const T min, const T max)
 {
     std::lock_guard<std::mutex> lock(_mtx);
 
-    if(min < *_min)
+    using type = typename std::conditional<std::is_same<T, float>::value, float, int32_t>::type;
+
+    auto min_ptr = static_cast<type *>(_min);
+    auto max_ptr = static_cast<type *>(_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<float>::max());
+    float32x2_t carry_max = vdup_n_f32(std::numeric_limits<float>::lowest());
+
+    float carry_min_scalar = std::numeric_limits<float>::max();
+    float carry_max_scalar = std::numeric_limits<float>::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<const float *const>(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<T, bool(N & 8), bool(N & 4), bool(N & 2), bool(N & 1)>...
 };
 
-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<uint8_t, gen_index_seq<16>::type>::func_table[table_idx];
             break;
-        case Format::S16:
+        case DataType::S16:
             _func = create_func_table<int16_t, gen_index_seq<16>::type>::func_table[table_idx];
             break;
+        case DataType::F32:
+            _func = create_func_table<float, gen_index_seq<16>::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<std::is_same<T, float>::value, float, int32_t>::type;
+
+        auto min_ptr = static_cast<type *>(_min);
+        auto max_ptr = static_cast<type *>(_max);
+
         execute_window_loop(win, [&](const Coordinates & id)
         {
             auto    in_ptr = reinterpret_cast<const T *>(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<int32_t *>(_min));
+    q.enqueueReadBuffer(_min_max_vals, CL_FALSE, 1 * sizeof(int32_t), sizeof(int32_t), static_cast<int32_t *>(_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);