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>

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 <typename T>
     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<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);
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<CLTensor>(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<CLTensor>(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<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<RawTensor, RawTensor> 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<Coordinates2D> &min_loc, IArray<Coordinates2D> &max_loc,
+void Reference::compute_reference_min_max_location(const TensorShape &shape, DataType dt_in, void *min, void *max, IArray<Coordinates2D> &min_loc, IArray<Coordinates2D> &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<Coordinates2D> &min_loc, IArray<Coordinates2D> &max_loc,
+    static void compute_reference_min_max_location(const TensorShape &shape, DataType dt_in, void *min, void *max, IArray<Coordinates2D> &min_loc, IArray<Coordinates2D> &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<Coordinates2D> &min_loc, IArray<Coordinates2D> &max_loc, uint32_t &min_count, uint32_t &max_count)
+void ReferenceCPP::min_max_location(const RawTensor &src, void *min, void *max, IArray<Coordinates2D> &min_loc, IArray<Coordinates2D> &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<Coordinates2D> &min_loc, IArray<Coordinates2D> &max_loc, uint32_t &min_count, uint32_t &max_count);
+    static void min_max_location(const RawTensor &src, void *min, void *max, IArray<Coordinates2D> &min_loc, IArray<Coordinates2D> &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<T1> &in, Tensor<T2> &out_x, Tensor<T2> &out_y, BorderMode
     }
 }
 
-// Min max location
-template <typename T1>
-void min_max_location(const Tensor<T1> &in, int32_t &min, int32_t &max, IArray<Coordinates2D> &min_loc, IArray<Coordinates2D> &max_loc, uint32_t &min_count, uint32_t &max_count)
+template <typename T>
+void compute_min_max(const Tensor<T> &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<std::is_same<T, float>::value, float, int32_t>::type;
 
-    const size_t width = in.shape().x();
+    // Set min and max to first pixel
+    type tmp_min = static_cast<type>(in[0]);
+    type tmp_max = static_cast<type>(in[0]);
 
     // Look for min and max values
     for(int i = 1; i < in.num_elements(); ++i)
     {
-        if(static_cast<int32_t>(in[i]) < min)
+        if(static_cast<type>(in[i]) < tmp_min)
         {
-            min = in[i];
+            tmp_min = static_cast<type>(in[i]);
         }
-        if(static_cast<int32_t>(in[i]) > max)
+        if(static_cast<type>(in[i]) > tmp_max)
         {
-            max = in[i];
+            tmp_max = static_cast<type>(in[i]);
         }
     }
 
+    *static_cast<type *>(min) = tmp_min;
+    *static_cast<type *>(max) = tmp_max;
+}
+
+// Min max location
+template <typename T1>
+void min_max_location(const Tensor<T1> &in, void *min, void *max, IArray<Coordinates2D> &min_loc, IArray<Coordinates2D> &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<std::is_same<T1, float>::value, float, int32_t>::type;
+
+    type min_value = *static_cast<type *>(min);
+    type max_value = *static_cast<type *>(max);
+
+    min_count = 0;
+    max_count = 0;
     for(int i = 0; i < in.num_elements(); ++i)
     {
-        if(static_cast<int32_t>(in[i]) == min)
+        if(static_cast<type>(in[i]) == min_value)
         {
             Coordinates2D min_coord;
             min_coord.x = static_cast<int32_t>(i % width);
@@ -342,7 +358,7 @@ void min_max_location(const Tensor<T1> &in, int32_t &min, int32_t &max, IArray<C
 
             min_count++;
         }
-        if(static_cast<int32_t>(in[i]) == max)
+        if(static_cast<type>(in[i]) == max_value)
         {
             Coordinates2D max_coord;
             max_coord.x = static_cast<int32_t>(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<Coordinates2D> &min_loc, IArray<Coordinates2D> &max_loc, uint32_t &min_count, uint32_t &max_count)
+    explicit min_max_location_visitor(void *min, void *max, IArray<Coordinates2D> &min_loc, IArray<Coordinates2D> &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<Coordinates2D> &_min_loc;
     IArray<Coordinates2D> &_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<Coordinates2D> &min_loc, IArray<Coordinates2D> &ref_min_loc, IArray<Coordinates2D> &max_loc, IArray<Coordinates2D> &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 <typename T>
+void validate_min_max_loc(T min, T ref_min, T max, T ref_max,
                           IArray<Coordinates2D> &min_loc, IArray<Coordinates2D> &ref_min_loc, IArray<Coordinates2D> &max_loc, IArray<Coordinates2D> &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