COMPMID-596: Port HOGDetector to new validation

Change-Id: I73231fc71c5166268e6c909b7930b7e034f3794e
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/118876
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Tested-by: Jenkins <bsgcomp@arm.com>

22 files changed, 871 insertions, 55 deletions

diff --git a/arm_compute/core/HOGInfo.h b/arm_compute/core/HOGInfo.h
index f55574288e..90d44e3caf 100644
--- a/arm_compute/core/HOGInfo.h
+++ b/arm_compute/core/HOGInfo.h
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2016, 2018 ARM Limited.
+ * Copyright (c) 2016-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -129,13 +129,13 @@ public:
      * @return The Size2D data object which stores the number of cells per block stride along the x and y directions
      */
     Size2D num_cells_per_block_stride() const;
-    /** Calculates the number of blocks for the given image size
+    /** Calculates the number of block positions for the given image size
      *
      * @param[in] image_size The input image size data object
      *
-     * @return The Size2D data object which stores the number of blocks along the x and y directions
+     * @return The Size2D data object which stores the number of block positions along the x and y directions
      */
-    Size2D num_blocks_per_image(const Size2D &image_size) const;
+    Size2D num_block_positions_per_image(const Size2D &image_size) const;
 
 private:
     Size2D      _cell_size;
diff --git a/src/core/CL/cl_kernels/hog.cl b/src/core/CL/cl_kernels/hog.cl
index 3d37fbcaf2..407ee2f3cf 100644
--- a/src/core/CL/cl_kernels/hog.cl
+++ b/src/core/CL/cl_kernels/hog.cl
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017, 2018 ARM Limited.
+ * Copyright (c) 2017-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -351,7 +351,7 @@ __kernel void hog_block_normalization(IMAGE_DECLARATION(src),
 }
 #endif /* NUM_CELLS_PER_BLOCK_HEIGHT and NUM_BINS_PER_BLOCK_X and NUM_BINS_PER_BLOCK and HOG_NORM_TYPE and L2_HYST_THRESHOLD */
 
-#if defined(NUM_BLOCKS_PER_DESCRIPTOR_Y) && defined(NUM_BINS_PER_DESCRIPTOR_X) && defined(THRESHOLD) && defined(MAX_NUM_DETECTION_WINDOWS) && defined(IDX_CLASS) && defined(BLOCK_STRIDE_WIDTH) && defined(BLOCK_STRIDE_HEIGHT) && defined(DETECTION_WINDOW_WIDTH) && defined(DETECTION_WINDOW_HEIGHT)
+#if defined(NUM_BLOCKS_PER_DESCRIPTOR_Y) && defined(NUM_BINS_PER_DESCRIPTOR_X) && defined(THRESHOLD) && defined(MAX_NUM_DETECTION_WINDOWS) && defined(IDX_CLASS) && defined(DETECTION_WINDOW_STRIDE_WIDTH) && defined(DETECTION_WINDOW_STRIDE_HEIGHT) && defined(DETECTION_WINDOW_WIDTH) && defined(DETECTION_WINDOW_HEIGHT)
 
 /** This OpenCL kernel computes the HOG detector using linear SVM
  *
@@ -362,8 +362,8 @@ __kernel void hog_block_normalization(IMAGE_DECLARATION(src),
  * -# -DTHRESHOLD = Threshold for the distance between features and SVM classifying plane
  * -# -DMAX_NUM_DETECTION_WINDOWS = Maximum number of possible detection windows. It is equal to the size of the DetectioWindow array
  * -# -DIDX_CLASS = Index of the class to detect
- * -# -DBLOCK_STRIDE_WIDTH = Block stride for the X direction
- * -# -DBLOCK_STRIDE_HEIGHT = Block stride for the Y direction
+ * -# -DDETECTION_WINDOW_STRIDE_WIDTH = Detection window stride for the X direction
+ * -# -DDETECTION_WINDOW_STRIDE_HEIGHT = Detection window stride for the Y direction
  * -# -DDETECTION_WINDOW_WIDTH = Width of the detection window
  * -# -DDETECTION_WINDOW_HEIGHT = Height of the detection window
  *
@@ -443,8 +443,8 @@ __kernel void hog_detector(IMAGE_DECLARATION(src),
         int id = atomic_inc(num_detection_windows);
         if(id < MAX_NUM_DETECTION_WINDOWS)
         {
-            dst[id].x         = get_global_id(0) * BLOCK_STRIDE_WIDTH;
-            dst[id].y         = get_global_id(1) * BLOCK_STRIDE_HEIGHT;
+            dst[id].x         = get_global_id(0) * DETECTION_WINDOW_STRIDE_WIDTH;
+            dst[id].y         = get_global_id(1) * DETECTION_WINDOW_STRIDE_HEIGHT;
             dst[id].width     = DETECTION_WINDOW_WIDTH;
             dst[id].height    = DETECTION_WINDOW_HEIGHT;
             dst[id].idx_class = IDX_CLASS;
@@ -453,4 +453,4 @@ __kernel void hog_detector(IMAGE_DECLARATION(src),
     }
 }
 #endif /* NUM_BLOCKS_PER_DESCRIPTOR_Y && NUM_BINS_PER_DESCRIPTOR_X && THRESHOLD && MAX_NUM_DETECTION_WINDOWS && IDX_CLASS &&
-        * BLOCK_STRIDE_WIDTH && BLOCK_STRIDE_HEIGHT && DETECTION_WINDOW_WIDTH && DETECTION_WINDOW_HEIGHT */
+        * DETECTION_WINDOW_STRIDE_WIDTH && DETECTION_WINDOW_STRIDE_HEIGHT && DETECTION_WINDOW_WIDTH && DETECTION_WINDOW_HEIGHT */
diff --git a/src/core/CL/kernels/CLHOGDescriptorKernel.cpp b/src/core/CL/kernels/CLHOGDescriptorKernel.cpp
index 87659c4ba9..a15aab1f37 100644
--- a/src/core/CL/kernels/CLHOGDescriptorKernel.cpp
+++ b/src/core/CL/kernels/CLHOGDescriptorKernel.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017 ARM Limited.
+ * Copyright (c) 2017-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -172,7 +172,7 @@ void CLHOGBlockNormalizationKernel::configure(const ICLTensor *input, ICLTensor
                               AccessWindowRectangle(input->info(), 0, 0, num_elems_read_per_iteration, num_rows_read_per_iteration),
                               output_access);
 
-    output_access.set_valid_region(win, input->info()->valid_region());
+    output_access.set_valid_region(win, ValidRegion(Coordinates(), output->info()->tensor_shape()));
 
     ICLKernel::configure(win);
 }
diff --git a/src/core/CL/kernels/CLHOGDetectorKernel.cpp b/src/core/CL/kernels/CLHOGDetectorKernel.cpp
index 0f9a98950d..caca49846f 100644
--- a/src/core/CL/kernels/CLHOGDetectorKernel.cpp
+++ b/src/core/CL/kernels/CLHOGDetectorKernel.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017 ARM Limited.
+ * Copyright (c) 2017-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -70,10 +70,10 @@ void CLHOGDetectorKernel::configure(const ICLTensor *input, const ICLHOG *hog, I
     args_str << "-DTHRESHOLD=" << threshold << " ";
     args_str << "-DMAX_NUM_DETECTION_WINDOWS=" << detection_windows->max_num_values() << " ";
     args_str << "-DIDX_CLASS=" << idx_class << " ";
-    args_str << "-DBLOCK_STRIDE_WIDTH=" << block_stride.width << " ";
-    args_str << "-DBLOCK_STRIDE_HEIGHT=" << block_stride.height << " ";
     args_str << "-DDETECTION_WINDOW_WIDTH=" << detection_window_size.width << " ";
     args_str << "-DDETECTION_WINDOW_HEIGHT=" << detection_window_size.height << " ";
+    args_str << "-DDETECTION_WINDOW_STRIDE_WIDTH=" << detection_window_stride.width << " ";
+    args_str << "-DDETECTION_WINDOW_STRIDE_HEIGHT=" << detection_window_stride.height << " ";
 
     // Construct kernel name
     std::set<std::string> build_opts = {};
@@ -102,8 +102,8 @@ void CLHOGDetectorKernel::configure(const ICLTensor *input, const ICLHOG *hog, I
 
     // Configure kernel window
     Window win;
-    win.set(Window::DimX, Window::Dimension(0, floor_to_multiple(num_blocks_x - num_blocks_per_detection_window_x, window_step_x), window_step_x));
-    win.set(Window::DimY, Window::Dimension(0, floor_to_multiple(num_blocks_y - num_blocks_per_detection_window_y, window_step_y), window_step_y));
+    win.set(Window::DimX, Window::Dimension(0, floor_to_multiple(num_blocks_x - num_blocks_per_detection_window_x, window_step_x) + window_step_x, window_step_x));
+    win.set(Window::DimY, Window::Dimension(0, floor_to_multiple(num_blocks_y - num_blocks_per_detection_window_y, window_step_y) + window_step_y, window_step_y));
 
     constexpr unsigned int num_elems_read_per_iteration = 1;
     const unsigned int     num_rows_read_per_iteration  = num_blocks_per_descriptor_y;
diff --git a/src/core/HOGInfo.cpp b/src/core/HOGInfo.cpp
index 73f4c42041..4f99455b56 100644
--- a/src/core/HOGInfo.cpp
+++ b/src/core/HOGInfo.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2016, 2018 ARM Limited.
+ * Copyright (c) 2016-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -61,7 +61,7 @@ void HOGInfo::init(const Size2D &cell_size, const Size2D &block_size, const Size
     _phase_type            = phase_type;
 
     // Compute descriptor size. +1 takes into account of the bias
-    _descriptor_size = num_cells_per_block().area() * num_blocks_per_image(_detection_window_size).area() * _num_bins + 1;
+    _descriptor_size = num_cells_per_block().area() * num_block_positions_per_image(_detection_window_size).area() * _num_bins + 1;
 }
 
 Size2D HOGInfo::num_cells_per_block() const
@@ -80,8 +80,10 @@ Size2D HOGInfo::num_cells_per_block_stride() const
                   _block_stride.height / _cell_size.height);
 }
 
-Size2D HOGInfo::num_blocks_per_image(const Size2D &image_size) const
+Size2D HOGInfo::num_block_positions_per_image(const Size2D &image_size) const
 {
+    ARM_COMPUTE_ERROR_ON(_block_stride.width == 0 || _block_stride.height == 0);
+
     return Size2D(((image_size.width - _block_size.width) / _block_stride.width) + 1,
                   ((image_size.height - _block_size.height) / _block_stride.height) + 1);
 }
diff --git a/src/core/NEON/kernels/NEHOGDescriptorKernel.cpp b/src/core/NEON/kernels/NEHOGDescriptorKernel.cpp
index 3fd81bed1c..abe224e854 100644
--- a/src/core/NEON/kernels/NEHOGDescriptorKernel.cpp
+++ b/src/core/NEON/kernels/NEHOGDescriptorKernel.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2016, 2017 ARM Limited.
+ * Copyright (c) 2016-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -764,7 +764,7 @@ void NEHOGBlockNormalizationKernel::configure(const ITensor *input, ITensor *out
                               AccessWindowRectangle(input->info(), 0, 0, num_elems_read_per_iteration, num_rows_read_per_iteration),
                               output_access);
 
-    output_access.set_valid_region(win, input->info()->valid_region());
+    output_access.set_valid_region(win, ValidRegion(Coordinates(), output->info()->tensor_shape()));
 
     INEKernel::configure(win);
 }
diff --git a/src/core/NEON/kernels/NEHOGDetectorKernel.cpp b/src/core/NEON/kernels/NEHOGDetectorKernel.cpp
index 343b0517b0..2c02ab8997 100644
--- a/src/core/NEON/kernels/NEHOGDetectorKernel.cpp
+++ b/src/core/NEON/kernels/NEHOGDetectorKernel.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2016, 2017 ARM Limited.
+ * Copyright (c) 2016-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -81,8 +81,8 @@ void NEHOGDetectorKernel::configure(const ITensor *input, const IHOG *hog, IDete
 
     // Configure kernel window
     Window win;
-    win.set(Window::DimX, Window::Dimension(0, floor_to_multiple(num_blocks_x - num_blocks_per_detection_window_x, window_step_x), window_step_x));
-    win.set(Window::DimY, Window::Dimension(0, floor_to_multiple(num_blocks_y - num_blocks_per_detection_window_y, window_step_y), window_step_y));
+    win.set(Window::DimX, Window::Dimension(0, floor_to_multiple(num_blocks_x - num_blocks_per_detection_window_x, window_step_x) + window_step_x, window_step_x));
+    win.set(Window::DimY, Window::Dimension(0, floor_to_multiple(num_blocks_y - num_blocks_per_detection_window_y, window_step_y) + window_step_y, window_step_y));
 
     constexpr unsigned int num_elems_read_per_iteration = 1;
     const unsigned int     num_rows_read_per_iteration  = _num_blocks_per_descriptor_y;
diff --git a/src/core/TensorInfo.cpp b/src/core/TensorInfo.cpp
index 2190e3415c..539d0f84b3 100644
--- a/src/core/TensorInfo.cpp
+++ b/src/core/TensorInfo.cpp
@@ -168,13 +168,13 @@ void TensorInfo::init(const HOGInfo &hog_info, unsigned int width, unsigned int
     // Number of cells for each block
     const Size2D num_cells_per_block = hog_info.num_cells_per_block();
 
-    // Tensor Size = (Number of horizontal blocks) * (Number of vertical blocks )
-    const Size2D num_blocks_per_img = hog_info.num_blocks_per_image(Size2D(width, height));
+    // Tensor Size = (Number of horizontal block positions) * (Number of vertical block positions)
+    const Size2D num_block_positions_per_img = hog_info.num_block_positions_per_image(Size2D(width, height));
 
     // Number of tensor channels = (Number of cells per block) * (Number of bins per cell)
     const size_t num_channels = num_cells_per_block.area() * hog_info.num_bins();
 
-    init(TensorShape(num_blocks_per_img.width, num_blocks_per_img.height), num_channels, DataType::F32);
+    init(TensorShape(num_block_positions_per_img.width, num_block_positions_per_img.height), num_channels, DataType::F32);
 }
 
 size_t TensorInfo::init_auto_padding(const TensorShape &tensor_shape, Format format)
@@ -212,13 +212,13 @@ size_t TensorInfo::init_auto_padding(const HOGInfo &hog_info, unsigned int width
     // Number of cells for each block
     const Size2D num_cells_per_block = hog_info.num_cells_per_block();
 
-    // Tensor Size = (Number of horizontal blocks) * (Number of vertical blocks )
-    const Size2D num_blocks_per_img = hog_info.num_blocks_per_image(Size2D(width, height));
+    // Tensor Size = (Number of horizontal block positions) * (Number of vertical block positions)
+    const Size2D num_block_positions_per_img = hog_info.num_block_positions_per_image(Size2D(width, height));
 
     // Number of tensor channels = (Number of cells per block) * (Number of bins per cell)
     const size_t num_channels = num_cells_per_block.area() * hog_info.num_bins();
 
-    return init_auto_padding(TensorShape(num_blocks_per_img.width, num_blocks_per_img.height), num_channels, DataType::F32);
+    return init_auto_padding(TensorShape(num_block_positions_per_img.width, num_block_positions_per_img.height), num_channels, DataType::F32);
 }
 
 bool TensorInfo::auto_padding()
diff --git a/tests/CL/CLHOGAccessor.h b/tests/CL/CLHOGAccessor.h
new file mode 100644
index 0000000000..7607c28142
--- /dev/null
+++ b/tests/CL/CLHOGAccessor.h
@@ -0,0 +1,74 @@
+/*
+ * Copyright (c) 2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef __ARM_COMPUTE_TEST_CLHOGACCESSOR_H__
+#define __ARM_COMPUTE_TEST_CLHOGACCESSOR_H__
+
+#include "arm_compute/runtime/CL/CLHOG.h"
+#include "tests/IHOGAccessor.h"
+
+namespace arm_compute
+{
+namespace test
+{
+/** Accessor implementation for @ref CLHOG objects. */
+class CLHOGAccessor : public IHOGAccessor
+{
+public:
+    /** Create an accessor for the given @p CLHOG. */
+    CLHOGAccessor(CLHOG &hog)
+        : _hog{ hog }
+    {
+        _hog.map();
+    }
+
+    /** Destructor that unmaps the CL memory. */
+    ~CLHOGAccessor()
+    {
+        _hog.unmap();
+    }
+
+    /** Prevent instances of this class from being copied (As this class contains references). */
+    CLHOGAccessor(const CLHOGAccessor &) = delete;
+    /** Prevent instances of this class from being copied (As this class contains references). */
+    CLHOGAccessor &operator=(const CLHOGAccessor &) = delete;
+    /** Allow instances of this class to be moved */
+    CLHOGAccessor(CLHOGAccessor &&) = default;
+    /** Allow instances of this class to be moved */
+    CLHOGAccessor &operator=(CLHOGAccessor &&) = default;
+
+    /** Pointer to the first element of the array which stores the linear SVM coefficients of HOG descriptor
+     *
+     * @return A pointer to the first element of the array which stores the linear SVM coefficients of HOG descriptor
+     */
+    float *descriptor() const override
+    {
+        return _hog.descriptor();
+    }
+
+private:
+    CLHOG &_hog;
+};
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_CLHOGACCESSOR_H__ */
diff --git a/tests/IHOGAccessor.h b/tests/IHOGAccessor.h
new file mode 100644
index 0000000000..0436cb1090
--- /dev/null
+++ b/tests/IHOGAccessor.h
@@ -0,0 +1,48 @@
+/*
+ * Copyright (c) 2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef __ARM_COMPUTE_TEST_IHOGACCESSOR_H__
+#define __ARM_COMPUTE_TEST_IHOGACCESSOR_H__
+
+namespace arm_compute
+{
+namespace test
+{
+/** Common interface to access HOG structure */
+class IHOGAccessor
+{
+public:
+    /** Virtual destructor. */
+    virtual ~IHOGAccessor() = default;
+
+    /** Pointer to the first element of the array which stores the linear SVM coefficients of HOG descriptor
+     *
+     * @note Other elements of the array can be accessed using descriptor()[idx] for idx=[0, descriptor_size() - 1]
+     *
+     * @return A pointer to the first element of the array which stores the linear SVM coefficients of HOG descriptor
+     */
+    virtual float *descriptor() const = 0;
+};
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_IHOGACCESSOR_H__ */
diff --git a/tests/NEON/HOGAccessor.h b/tests/NEON/HOGAccessor.h
new file mode 100644
index 0000000000..3250ab856a
--- /dev/null
+++ b/tests/NEON/HOGAccessor.h
@@ -0,0 +1,68 @@
+/*
+ * Copyright (c) 2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef __ARM_COMPUTE_TEST_HOGACCESSOR_H__
+#define __ARM_COMPUTE_TEST_HOGACCESSOR_H__
+
+#include "arm_compute/runtime/HOG.h"
+#include "tests/IHOGAccessor.h"
+
+namespace arm_compute
+{
+namespace test
+{
+/** Accessor implementation for @ref HOG objects. */
+class HOGAccessor : public IHOGAccessor
+{
+public:
+    /** Create an accessor for the given @p HOG.
+     */
+    HOGAccessor(HOG &hog)
+        : _hog{ hog }
+    {
+    }
+
+    /** Prevent instances of this class from being copied (As this class contains references). */
+    HOGAccessor(const HOGAccessor &) = delete;
+    /** Prevent instances of this class from being copied (As this class contains references). */
+    HOGAccessor &operator=(const HOGAccessor &) = delete;
+    /** Allow instances of this class to be moved */
+    HOGAccessor(HOGAccessor &&) = default;
+    /** Allow instances of this class to be moved */
+    HOGAccessor &operator=(HOGAccessor &&) = default;
+
+    /** Pointer to the first element of the array which stores the linear SVM coefficients of HOG descriptor
+     *
+     * @return A pointer to the first element of the array which stores the linear SVM coefficients of HOG descriptor
+     */
+    float *descriptor() const override
+    {
+        return _hog.descriptor();
+    }
+
+private:
+    HOG &_hog;
+};
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_HOGACCESSOR_H__ */
diff --git a/tests/Utils.h b/tests/Utils.h
index e9a953d727..f61b94628a 100644
--- a/tests/Utils.h
+++ b/tests/Utils.h
@@ -535,23 +535,14 @@ inline T create_multi_image(const TensorShape &shape, Format format)
 
 /** Create and initialize a HOG (Histogram of Oriented Gradients) of the given type.
  *
- * @param[in] cell_size             Cell size in pixels
- * @param[in] block_size            Block size in pixels. Must be a multiple of cell_size.
- * @param[in] detection_window_size Detection window size in pixels. Must be a multiple of block_size and block_stride.
- * @param[in] block_stride          Distance in pixels between 2 consecutive blocks along the x and y direction. Must be a multiple of cell size
- * @param[in] num_bins              Number of histogram bins for each cell
- * @param[in] normalization_type    (Optional) Normalization type to use for each block
- * @param[in] l2_hyst_threshold     (Optional) Threshold used for L2HYS_NORM normalization method
- * @param[in] phase_type            (Optional) Type of @ref PhaseType
+ * @param[in] hog_info HOGInfo object
  *
  * @return Initialized HOG of given type.
  */
 template <typename T>
-inline T create_HOG(const Size2D &cell_size, const Size2D &block_size, const Size2D &detection_window_size, const Size2D &block_stride, size_t num_bins,
-                    HOGNormType normalization_type = HOGNormType::L2HYS_NORM, float l2_hyst_threshold = 0.2f, PhaseType phase_type = PhaseType::UNSIGNED)
+inline T create_HOG(const HOGInfo &hog_info)
 {
-    T       hog;
-    HOGInfo hog_info(cell_size, block_size, block_size, block_stride, num_bins, normalization_type, l2_hyst_threshold, phase_type);
+    T hog;
     hog.init(hog_info);
 
     return hog;
@@ -603,6 +594,30 @@ inline std::vector<ROI> generate_random_rois(const TensorShape &shape, const ROI
     return rois;
 }
 
+/** Create a vector with a uniform distribution of floating point values across the specified range.
+ *
+ * @param[in] num_values The number of values to be created.
+ * @param[in] min        The minimum value in distribution (inclusive).
+ * @param[in] max        The maximum value in distribution (inclusive).
+ * @param[in] seed       The random seed to be used.
+ *
+ * @return A vector that contains the requested number of random floating point values
+ */
+template <typename T, typename = typename std::enable_if<std::is_floating_point<T>::value>::type>
+inline std::vector<T> generate_random_real(unsigned int num_values, T min, T max, std::random_device::result_type seed)
+{
+    std::vector<T>                    v(num_values);
+    std::mt19937                      gen(seed);
+    std::uniform_real_distribution<T> dist(min, max);
+
+    for(unsigned int i = 0; i < num_values; ++i)
+    {
+        v.at(i) = dist(gen);
+    }
+
+    return v;
+}
+
 template <typename T, typename ArrayAccessor_T>
 inline void fill_array(ArrayAccessor_T &&array, const std::vector<T> &v)
 {
diff --git a/tests/validation/CL/HOGDetector.cpp b/tests/validation/CL/HOGDetector.cpp
new file mode 100644
index 0000000000..6c2c18c3ea
--- /dev/null
+++ b/tests/validation/CL/HOGDetector.cpp
@@ -0,0 +1,98 @@
+/*
+ * Copyright (c) 2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/runtime/CL/CLArray.h"
+#include "arm_compute/runtime/CL/functions/CLHOGDescriptor.h"
+#include "arm_compute/runtime/CL/functions/CLHOGDetector.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/CL/CLArrayAccessor.h"
+#include "tests/CL/CLHOGAccessor.h"
+#include "tests/datasets/HOGDescriptorDataset.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/HOGDetectorFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/* Set the tolerance (percentage) used when validating the score of detection window. */
+RelativeTolerance<float> tolerance(0.01f);
+
+/* Input dataset (values must be a multiple of the HOGInfo block_size) */
+const auto DetectionWindowStrideDataset = framework::dataset::make("DetectionWindowStride", { Size2D(8, 8), Size2D(16, 16) });
+} // namespace
+
+TEST_SUITE(CL)
+TEST_SUITE(HOGDetector)
+
+// *INDENT-OFF*
+// clang-format off
+using CLHOGDetectorFixture = HOGDetectorValidationFixture<CLTensor,
+                                                          CLHOG,
+                                                          CLDetectionWindowArray,
+                                                          CLHOGDescriptor,
+                                                          CLAccessor,
+                                                          CLArrayAccessor<DetectionWindow>,
+                                                          CLHOGAccessor,
+                                                          CLHOGDetector,
+                                                          uint8_t,
+                                                          float>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, CLHOGDetectorFixture, framework::DatasetMode::PRECOMMIT,
+                       combine(combine(combine(
+                       DetectionWindowStrideDataset,
+                       datasets::SmallHOGDescriptorDataset()),
+                       framework::dataset::make("Format", Format::U8)),
+                       framework::dataset::make("BorderMode", {BorderMode::CONSTANT, BorderMode::REPLICATE})))
+
+{
+    // Validate output
+    validate_detection_windows(_target.begin(), _target.end(), _reference.begin(), _reference.end(), tolerance);
+}
+
+FIXTURE_DATA_TEST_CASE(RunLarge, CLHOGDetectorFixture, framework::DatasetMode::NIGHTLY,
+                       combine(combine(combine(
+                       DetectionWindowStrideDataset,
+                       datasets::LargeHOGDescriptorDataset()),
+                       framework::dataset::make("Format", Format::U8)),
+                       framework::dataset::make("BorderMode", {BorderMode::CONSTANT, BorderMode::REPLICATE})))
+{
+    // Validate output
+    validate_detection_windows(_target.begin(), _target.end(), _reference.begin(), _reference.end(), tolerance);
+}
+
+// clang-format on
+// *INDENT-ON*
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/NEON/HOGDetector.cpp b/tests/validation/NEON/HOGDetector.cpp
new file mode 100644
index 0000000000..c787728d2c
--- /dev/null
+++ b/tests/validation/NEON/HOGDetector.cpp
@@ -0,0 +1,98 @@
+/*
+ * Copyright (c) 2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/runtime/Array.h"
+#include "arm_compute/runtime/NEON/functions/NEHOGDescriptor.h"
+#include "arm_compute/runtime/NEON/functions/NEHOGDetector.h"
+#include "tests/NEON/Accessor.h"
+#include "tests/NEON/ArrayAccessor.h"
+#include "tests/NEON/HOGAccessor.h"
+#include "tests/datasets/HOGDescriptorDataset.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/HOGDetectorFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/* Set the tolerance (percentage) used when validating the score of detection window.
+   Note: High tolerance is required due to divergence between CL and NEON detection window scores. */
+RelativeTolerance<float> tolerance(1.0f);
+
+/* Input dataset (values must be a multiple of the HOGInfo block_size) */
+const auto DetectionWindowStrideDataset = framework::dataset::make("DetectionWindowStride", { Size2D(8, 8), Size2D(16, 16) });
+} // namespace
+
+TEST_SUITE(NEON)
+TEST_SUITE(HOGDetector)
+
+// *INDENT-OFF*
+// clang-format off
+using NEHOGDetectorFixture = HOGDetectorValidationFixture<Tensor,
+                                                          HOG,
+                                                          DetectionWindowArray,
+                                                          NEHOGDescriptor,
+                                                          Accessor,
+                                                          ArrayAccessor<DetectionWindow>,
+                                                          HOGAccessor,
+                                                          NEHOGDetector,
+                                                          uint8_t,
+                                                          float>;
+
+FIXTURE_DATA_TEST_CASE(RunSmall, NEHOGDetectorFixture, framework::DatasetMode::PRECOMMIT,
+                       combine(combine(combine(
+                       DetectionWindowStrideDataset,
+                       datasets::SmallHOGDescriptorDataset()),
+                       framework::dataset::make("Format", Format::U8)),
+                       framework::dataset::make("BorderMode", {BorderMode::CONSTANT, BorderMode::REPLICATE})))
+{
+    // Validate output
+    validate_detection_windows(_target.begin(), _target.end(), _reference.begin(), _reference.end(), tolerance);
+}
+
+FIXTURE_DATA_TEST_CASE(RunLarge, NEHOGDetectorFixture, framework::DatasetMode::NIGHTLY,
+                       combine(combine(combine(
+                       DetectionWindowStrideDataset,
+                       datasets::LargeHOGDescriptorDataset()),
+                       framework::dataset::make("Format", Format::U8)),
+                       framework::dataset::make("BorderMode", {BorderMode::CONSTANT, BorderMode::REPLICATE})))
+{
+    // Validate output
+    validate_detection_windows(_target.begin(), _target.end(), _reference.begin(), _reference.end(), tolerance);
+}
+
+// clang-format on
+// *INDENT-ON*
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/Validation.h b/tests/validation/Validation.h
index 26271c8bf3..508fb027ca 100644
--- a/tests/validation/Validation.h
+++ b/tests/validation/Validation.h
@@ -247,7 +247,11 @@ template <typename T, typename U, typename V = AbsoluteTolerance<float>>
 void validate_keypoints(T target_first, T target_last, U reference_first, U reference_last, V tolerance = AbsoluteTolerance<float>(),
                         float allowed_missing_percentage = 5.f, float allowed_mismatch_percentage = 5.f);
 
-/** Compare values with a tolerance. */
+/** Validate detection windows. */
+template <typename T, typename U, typename V = AbsoluteTolerance<float>>
+void validate_detection_windows(T target_first, T target_last, U reference_first, U reference_last, V tolerance = AbsoluteTolerance<float>(),
+                                float allowed_missing_percentage = 5.f, float allowed_mismatch_percentage = 5.f);
+
 template <typename T>
 struct compare_base
 {
@@ -733,6 +737,77 @@ void validate_keypoints(T target_first, T target_last, U reference_first, U refe
     }
 }
 
+/** Check which detection windows from [first1, last1) are missing in [first2, last2) */
+template <typename T, typename U, typename V>
+std::pair<int64_t, int64_t> compare_detection_windows(T first1, T last1, U first2, U last2, V tolerance)
+{
+    int64_t num_missing    = 0;
+    int64_t num_mismatches = 0;
+
+    while(first1 != last1)
+    {
+        const auto window = std::find_if(first2, last2, [&](DetectionWindow window)
+        {
+            return window.x == first1->x && window.y == first1->y && window.width == first1->width && window.height == first1->height && window.idx_class == first1->idx_class;
+        });
+
+        if(window == last2)
+        {
+            ++num_missing;
+            ARM_COMPUTE_TEST_INFO("Detection window not found " << *first1)
+        }
+        else
+        {
+            if(!compare<V>(window->score, first1->score, tolerance))
+            {
+                ++num_mismatches;
+                ARM_COMPUTE_TEST_INFO("Mismatching detection window")
+                ARM_COMPUTE_TEST_INFO("detection window 1= " << *first1)
+                ARM_COMPUTE_TEST_INFO("detection window 2= " << *window)
+            }
+        }
+
+        ++first1;
+    }
+
+    return std::make_pair(num_missing, num_mismatches);
+}
+
+template <typename T, typename U, typename V>
+void validate_detection_windows(T target_first, T target_last, U reference_first, U reference_last, V tolerance,
+                                float allowed_missing_percentage, float allowed_mismatch_percentage)
+{
+    const int64_t num_elements_target    = std::distance(target_first, target_last);
+    const int64_t num_elements_reference = std::distance(reference_first, reference_last);
+
+    int64_t num_missing    = 0;
+    int64_t num_mismatches = 0;
+
+    if(num_elements_reference > 0)
+    {
+        std::tie(num_missing, num_mismatches) = compare_detection_windows(reference_first, reference_last, target_first, target_last, tolerance);
+
+        const float percent_missing    = static_cast<float>(num_missing) / num_elements_reference * 100.f;
+        const float percent_mismatches = static_cast<float>(num_mismatches) / num_elements_reference * 100.f;
+
+        ARM_COMPUTE_TEST_INFO(num_missing << " detection windows (" << std::fixed << std::setprecision(2) << percent_missing << "%) are missing in target");
+        ARM_COMPUTE_EXPECT(percent_missing <= allowed_missing_percentage, framework::LogLevel::ERRORS);
+
+        ARM_COMPUTE_TEST_INFO(num_mismatches << " detection windows (" << std::fixed << std::setprecision(2) << percent_mismatches << "%) mismatched");
+        ARM_COMPUTE_EXPECT(percent_mismatches <= allowed_mismatch_percentage, framework::LogLevel::ERRORS);
+    }
+
+    if(num_elements_target > 0)
+    {
+        std::tie(num_missing, num_mismatches) = compare_detection_windows(target_first, target_last, reference_first, reference_last, tolerance);
+
+        const float percent_missing = static_cast<float>(num_missing) / num_elements_target * 100.f;
+
+        ARM_COMPUTE_TEST_INFO(num_missing << " detection windows (" << std::fixed << std::setprecision(2) << percent_missing << "%) are not part of target");
+        ARM_COMPUTE_EXPECT(percent_missing <= allowed_missing_percentage, framework::LogLevel::ERRORS);
+    }
+}
+
 } // namespace validation
 } // namespace test
 } // namespace arm_compute
diff --git a/tests/validation/fixtures/HOGDescriptorFixture.h b/tests/validation/fixtures/HOGDescriptorFixture.h
index cabee63ae3..6097059c84 100644
--- a/tests/validation/fixtures/HOGDescriptorFixture.h
+++ b/tests/validation/fixtures/HOGDescriptorFixture.h
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017, 2018 ARM Limited.
+ * Copyright (c) 2017-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -81,14 +81,7 @@ protected:
         TensorInfo tensor_info_hog_descriptor(hog_info, shape.x(), shape.y());
 
         // Create HOG
-        HOGType hog = create_HOG<HOGType>(hog_info.cell_size(),
-                                          hog_info.block_size(),
-                                          hog_info.detection_window_size(),
-                                          hog_info.block_stride(),
-                                          hog_info.num_bins(),
-                                          hog_info.normalization_type(),
-                                          hog_info.l2_hyst_threshold(),
-                                          hog_info.phase_type());
+        HOGType hog = create_HOG<HOGType>(hog_info);
 
         // Create tensors
         TensorType src = create_tensor<TensorType>(shape, data_type_from_format(format));
diff --git a/tests/validation/fixtures/HOGDetectorFixture.h b/tests/validation/fixtures/HOGDetectorFixture.h
new file mode 100644
index 0000000000..c2d0514d11
--- /dev/null
+++ b/tests/validation/fixtures/HOGDetectorFixture.h
@@ -0,0 +1,138 @@
+/*
+ * Copyright (c) 2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef ARM_COMPUTE_TEST_HOG_DETECTOR_FIXTURE
+#define ARM_COMPUTE_TEST_HOG_DETECTOR_FIXTURE
+
+#include "arm_compute/core/HOGInfo.h"
+#include "arm_compute/core/Types.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/IAccessor.h"
+#include "tests/IHOGAccessor.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Fixture.h"
+#include "tests/validation/fixtures/HOGDescriptorFixture.h"
+#include "tests/validation/reference/HOGDetector.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+template <typename TensorType,
+          typename HOGType,
+          typename DetectionWindowArrayType,
+          typename HOGDescriptorType,
+          typename AccessorType,
+          typename ArrayAccessorType,
+          typename HOGAccessorType,
+          typename HOGDetectorType,
+          typename T,
+          typename U>
+class HOGDetectorValidationFixture : public HOGDescriptorValidationFixture<TensorType, HOGType, AccessorType, HOGDescriptorType, T, U>
+{
+public:
+    template <typename...>
+    void setup(Size2D detection_window_stride, std::string image, HOGInfo hog_info, Format format, BorderMode border_mode)
+    {
+        using HDF = HOGDescriptorValidationFixture<TensorType, HOGType, AccessorType, HOGDescriptorType, T, U>;
+        HDF::setup(image, hog_info, format, border_mode);
+
+        const unsigned int max_num_detection_windows = 100000;
+
+        // Initialise descriptor (linear SVM coefficients).
+        // NOTE: Fixed values are used to keep the number of detection windows detected
+        // consistent in order to have meaningful validation tolerances.
+        // The values are "unbalanced" to reduce the number of detected objects
+        std::random_device::result_type seed       = 0;
+        std::vector<U>                  descriptor = generate_random_real(hog_info.descriptor_size(), -0.505f, 0.495f, seed);
+
+        // Compute target and reference values using feature vector from descriptor kernel
+        _target    = compute_target(HDF::_target, descriptor, max_num_detection_windows, hog_info, detection_window_stride);
+        _reference = compute_reference(HDF::_reference, descriptor, max_num_detection_windows, hog_info, detection_window_stride);
+    }
+
+protected:
+    std::vector<DetectionWindow> compute_target(const TensorType &src, const std::vector<U> &descriptor, unsigned int max_num_detection_windows,
+                                                const HOGInfo &hog_info, const Size2D &detection_window_stride)
+    {
+        // Create HOG
+        HOGType hog = create_HOG<HOGType>(hog_info);
+
+        // Create array of detection windows
+        DetectionWindowArrayType detection_windows(max_num_detection_windows);
+
+        // Copy HOG descriptor values to HOG memory
+        {
+            HOGAccessorType hog_accessor(hog);
+            std::memcpy(hog_accessor.descriptor(), descriptor.data(), descriptor.size() * sizeof(U));
+        }
+
+        // Create and configure function
+        HOGDetectorType hog_detector;
+        hog_detector.configure(&src, &hog, &detection_windows, detection_window_stride);
+
+        // Reset detection windows
+        detection_windows.clear();
+
+        // Compute function
+        hog_detector.run();
+
+        // Create array of detection windows
+        std::vector<DetectionWindow> windows;
+
+        // Copy detection windows
+        ArrayAccessorType accessor(detection_windows);
+
+        for(size_t i = 0; i < accessor.num_values(); i++)
+        {
+            DetectionWindow win;
+            win.x         = accessor.at(i).x;
+            win.y         = accessor.at(i).y;
+            win.width     = accessor.at(i).width;
+            win.height    = accessor.at(i).height;
+            win.idx_class = accessor.at(i).idx_class;
+            win.score     = accessor.at(i).score;
+
+            windows.push_back(win);
+        }
+
+        return windows;
+    }
+
+    std::vector<DetectionWindow> compute_reference(const SimpleTensor<U> &src, const std::vector<U> &descriptor, unsigned int max_num_detection_windows,
+                                                   const HOGInfo &hog_info, const Size2D &detection_window_stride)
+    {
+        // Assumes defaults value of zero for threshold and class_idx.
+        return reference::hog_detector(src, descriptor, max_num_detection_windows, hog_info, detection_window_stride);
+    }
+
+    std::vector<DetectionWindow> _target{};
+    std::vector<DetectionWindow> _reference{};
+};
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_HOG_DETECTOR_FIXTURE */
diff --git a/tests/validation/reference/HOGDescriptor.cpp b/tests/validation/reference/HOGDescriptor.cpp
index 105eb838e7..ed22695793 100644
--- a/tests/validation/reference/HOGDescriptor.cpp
+++ b/tests/validation/reference/HOGDescriptor.cpp
@@ -255,6 +255,8 @@ SimpleTensor<T> hog_descriptor(const SimpleTensor<U> &src, BorderMode border_mod
     return desc;
 }
 
+template void hog_orientation_binning(const SimpleTensor<int16_t> &mag, const SimpleTensor<uint8_t> &phase, SimpleTensor<float> &hog_space, const HOGInfo &hog_info);
+template void hog_block_normalization(SimpleTensor<float> &desc, const SimpleTensor<float> &hog_space, const HOGInfo &hog_info);
 template SimpleTensor<float> hog_descriptor(const SimpleTensor<uint8_t> &src, BorderMode border_mode, uint8_t constant_border_value, const HOGInfo &hog_info);
 } // namespace reference
 } // namespace validation
diff --git a/tests/validation/reference/HOGDescriptor.h b/tests/validation/reference/HOGDescriptor.h
index e886445ec7..6ea83fe884 100644
--- a/tests/validation/reference/HOGDescriptor.h
+++ b/tests/validation/reference/HOGDescriptor.h
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017, 2018 ARM Limited.
+ * Copyright (c) 2017-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -34,6 +34,12 @@ namespace validation
 {
 namespace reference
 {
+template <typename T, typename U, typename V>
+void hog_orientation_binning(const SimpleTensor<T> &mag, const SimpleTensor<U> &phase, SimpleTensor<V> &hog_space, const HOGInfo &hog_info);
+
+template <typename T>
+void hog_block_normalization(SimpleTensor<T> &desc, const SimpleTensor<T> &hog_space, const HOGInfo &hog_info);
+
 template <typename T, typename U>
 SimpleTensor<T> hog_descriptor(const SimpleTensor<U> &src, BorderMode border_mode, U constant_border_value, const HOGInfo &hog_info);
 } // namespace reference
diff --git a/tests/validation/reference/HOGDetector.cpp b/tests/validation/reference/HOGDetector.cpp
new file mode 100644
index 0000000000..5a5ae3700d
--- /dev/null
+++ b/tests/validation/reference/HOGDetector.cpp
@@ -0,0 +1,132 @@
+/*
+ * Copyright (c) 2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "HOGDetector.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+namespace
+{
+/** Computes the number of detection windows to iterate over in the feature vector. */
+Size2D num_detection_windows(const TensorShape &shape, const Size2D &window_step, const HOGInfo &hog_info)
+{
+    const size_t num_block_strides_width  = hog_info.detection_window_size().width / hog_info.block_stride().width;
+    const size_t num_block_strides_height = hog_info.detection_window_size().height / hog_info.block_stride().height;
+
+    return Size2D(floor_to_multiple(shape.x() - num_block_strides_width, window_step.width) + window_step.width,
+                  floor_to_multiple(shape.y() - num_block_strides_height, window_step.height) + window_step.height);
+}
+} // namespace
+
+template <typename T>
+std::vector<DetectionWindow> hog_detector(const SimpleTensor<T> &src, const std::vector<T> &descriptor, unsigned int max_num_detection_windows,
+                                          const HOGInfo &hog_info, const Size2D &detection_window_stride, float threshold, uint16_t idx_class)
+{
+    ARM_COMPUTE_ERROR_ON_MSG((detection_window_stride.width % hog_info.block_stride().width != 0),
+                             "Detection window stride width must be multiple of block stride width");
+    ARM_COMPUTE_ERROR_ON_MSG((detection_window_stride.height % hog_info.block_stride().height != 0),
+                             "Detection window stride height must be multiple of block stride height");
+
+    // Create vector for identifying each detection window
+    std::vector<DetectionWindow> windows;
+
+    // Calculate detection window step
+    const Size2D window_step(detection_window_stride.width / hog_info.block_stride().width,
+                             detection_window_stride.height / hog_info.block_stride().height);
+
+    // Calculate number of detection windows
+    const Size2D num_windows = num_detection_windows(src.shape(), window_step, hog_info);
+
+    // Calculate detection window and row offsets in feature vector
+    const size_t src_offset_x   = window_step.width * hog_info.num_bins() * hog_info.num_cells_per_block().area();
+    const size_t src_offset_y   = window_step.height * hog_info.num_bins() * hog_info.num_cells_per_block().area() * src.shape().x();
+    const size_t src_offset_row = src.num_channels() * src.shape().x();
+
+    // Calculate detection window attributes
+    const Size2D       num_block_positions_per_detection_window = hog_info.num_block_positions_per_image(hog_info.detection_window_size());
+    const unsigned int num_bins_per_descriptor_x                = num_block_positions_per_detection_window.width * src.num_channels();
+    const unsigned int num_blocks_per_descriptor_y              = num_block_positions_per_detection_window.height;
+
+    ARM_COMPUTE_ERROR_ON((num_bins_per_descriptor_x * num_blocks_per_descriptor_y + 1) != hog_info.descriptor_size());
+
+    size_t win_id = 0;
+
+    // Traverse feature vector in detection window steps
+    for(auto win_y = 0u, offset_y = 0u; win_y < num_windows.height; win_y += window_step.height, offset_y += src_offset_y)
+    {
+        for(auto win_x = 0u, offset_x = 0u; win_x < num_windows.width; win_x += window_step.width, offset_x += src_offset_x)
+        {
+            // Reset the score
+            float score = 0.0f;
+
+            // Traverse detection window
+            for(auto y = 0u, offset_row = 0u; y < num_blocks_per_descriptor_y; ++y, offset_row += src_offset_row)
+            {
+                const int bin_offset = y * num_bins_per_descriptor_x;
+
+                for(auto x = 0u; x < num_bins_per_descriptor_x; ++x)
+                {
+                    // Compute Linear SVM
+                    const float a = src[x + offset_x + offset_y + offset_row];
+                    const float b = descriptor[x + bin_offset];
+                    score += a * b;
+                }
+            }
+
+            // Add the bias. The bias is located at the position (descriptor_size() - 1)
+            score += descriptor[num_bins_per_descriptor_x * num_blocks_per_descriptor_y];
+
+            if(score > threshold)
+            {
+                DetectionWindow window;
+
+                if(win_id++ < max_num_detection_windows)
+                {
+                    window.x         = win_x * hog_info.block_stride().width;
+                    window.y         = win_y * hog_info.block_stride().height;
+                    window.width     = hog_info.detection_window_size().width;
+                    window.height    = hog_info.detection_window_size().height;
+                    window.idx_class = idx_class;
+                    window.score     = score;
+
+                    windows.push_back(window);
+                }
+            }
+        }
+    }
+
+    return windows;
+}
+
+template std::vector<DetectionWindow> hog_detector(const SimpleTensor<float> &src, const std::vector<float> &descriptor, unsigned int max_num_detection_windows,
+                                                   const HOGInfo &hog_info, const Size2D &detection_window_stride, float threshold, uint16_t idx_class);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/reference/HOGDetector.h b/tests/validation/reference/HOGDetector.h
new file mode 100644
index 0000000000..e88acb8739
--- /dev/null
+++ b/tests/validation/reference/HOGDetector.h
@@ -0,0 +1,48 @@
+/*
+ * Copyright (c) 2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef __ARM_COMPUTE_TEST_HOG_DETECTOR_H__
+#define __ARM_COMPUTE_TEST_HOG_DETECTOR_H__
+
+#include "arm_compute/core/Types.h"
+#include "arm_compute/core/Utils.h"
+#include "tests/SimpleTensor.h"
+
+#include <vector>
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+std::vector<DetectionWindow> hog_detector(const SimpleTensor<T> &src, const std::vector<T> &descriptor, unsigned int max_num_detection_windows,
+                                          const HOGInfo &hog_info, const Size2D &detection_window_stride, float threshold = 0.0f, uint16_t idx_class = 0);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_HOG_DETECTOR_H__ */
diff --git a/utils/TypePrinter.h b/utils/TypePrinter.h
index f0c1bd1329..811fb7b1e8 100644
--- a/utils/TypePrinter.h
+++ b/utils/TypePrinter.h
@@ -1454,5 +1454,24 @@ inline std::string to_string(const GPUTarget &gpu_target)
     str << gpu_target;
     return str.str();
 }
+/** Formatted output of the DetectionWindow type. */
+inline ::std::ostream &operator<<(::std::ostream &os, const DetectionWindow &detection_window)
+{
+    os << "{x=" << detection_window.x << ","
+       << "y=" << detection_window.y << ","
+       << "width=" << detection_window.width << ","
+       << "height=" << detection_window.height << ","
+       << "idx_class=" << detection_window.idx_class << ","
+       << "score=" << detection_window.score << "}";
+
+    return os;
+}
+
+inline std::string to_string(const DetectionWindow &type)
+{
+    std::stringstream str;
+    str << type;
+    return str.str();
+}
 } // namespace arm_compute
 #endif /* __ARM_COMPUTE_TEST_TYPE_PRINTER_H__ */
\ No newline at end of file