From 684cb0f29d23fbe418e5e5347234abf9eccef363 Mon Sep 17 00:00:00 2001 From: John Richardson Date: Tue, 9 Jan 2018 11:17:00 +0000 Subject: COMPMID-596: Port HOGDetector to new validation Change-Id: I73231fc71c5166268e6c909b7930b7e034f3794e Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/118876 Reviewed-by: Anthony Barbier Tested-by: Jenkins --- arm_compute/core/HOGInfo.h | 8 +- src/core/CL/cl_kernels/hog.cl | 14 +-- src/core/CL/kernels/CLHOGDescriptorKernel.cpp | 4 +- src/core/CL/kernels/CLHOGDetectorKernel.cpp | 10 +- src/core/HOGInfo.cpp | 8 +- src/core/NEON/kernels/NEHOGDescriptorKernel.cpp | 4 +- src/core/NEON/kernels/NEHOGDetectorKernel.cpp | 6 +- src/core/TensorInfo.cpp | 12 +- tests/CL/CLHOGAccessor.h | 74 ++++++++++++ tests/IHOGAccessor.h | 48 ++++++++ tests/NEON/HOGAccessor.h | 68 +++++++++++ tests/Utils.h | 39 +++++-- tests/validation/CL/HOGDetector.cpp | 98 ++++++++++++++++ tests/validation/NEON/HOGDetector.cpp | 98 ++++++++++++++++ tests/validation/Validation.h | 77 ++++++++++++- tests/validation/fixtures/HOGDescriptorFixture.h | 11 +- tests/validation/fixtures/HOGDetectorFixture.h | 138 +++++++++++++++++++++++ tests/validation/reference/HOGDescriptor.cpp | 2 + tests/validation/reference/HOGDescriptor.h | 8 +- tests/validation/reference/HOGDetector.cpp | 132 ++++++++++++++++++++++ tests/validation/reference/HOGDetector.h | 48 ++++++++ utils/TypePrinter.h | 19 ++++ 22 files changed, 871 insertions(+), 55 deletions(-) create mode 100644 tests/CL/CLHOGAccessor.h create mode 100644 tests/IHOGAccessor.h create mode 100644 tests/NEON/HOGAccessor.h create mode 100644 tests/validation/CL/HOGDetector.cpp create mode 100644 tests/validation/NEON/HOGDetector.cpp create mode 100644 tests/validation/fixtures/HOGDetectorFixture.h create mode 100644 tests/validation/reference/HOGDetector.cpp create mode 100644 tests/validation/reference/HOGDetector.h 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 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 -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 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 ::value>::type> +inline std::vector generate_random_real(unsigned int num_values, T min, T max, std::random_device::result_type seed) +{ + std::vector v(num_values); + std::mt19937 gen(seed); + std::uniform_real_distribution dist(min, max); + + for(unsigned int i = 0; i < num_values; ++i) + { + v.at(i) = dist(gen); + } + + return v; +} + template inline void fill_array(ArrayAccessor_T &&array, const std::vector &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 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, + 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 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, + 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 > void validate_keypoints(T target_first, T target_last, U reference_first, U reference_last, V tolerance = AbsoluteTolerance(), float allowed_missing_percentage = 5.f, float allowed_mismatch_percentage = 5.f); -/** Compare values with a tolerance. */ +/** Validate detection windows. */ +template > +void validate_detection_windows(T target_first, T target_last, U reference_first, U reference_last, V tolerance = AbsoluteTolerance(), + float allowed_missing_percentage = 5.f, float allowed_mismatch_percentage = 5.f); + template 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 +std::pair 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(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 +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(num_missing) / num_elements_reference * 100.f; + const float percent_mismatches = static_cast(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(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(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(hog_info); // Create tensors TensorType src = create_tensor(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 +class HOGDetectorValidationFixture : public HOGDescriptorValidationFixture +{ +public: + template + void setup(Size2D detection_window_stride, std::string image, HOGInfo hog_info, Format format, BorderMode border_mode) + { + using HDF = HOGDescriptorValidationFixture; + 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 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 compute_target(const TensorType &src, const std::vector &descriptor, unsigned int max_num_detection_windows, + const HOGInfo &hog_info, const Size2D &detection_window_stride) + { + // Create HOG + HOGType hog = create_HOG(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 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 compute_reference(const SimpleTensor &src, const std::vector &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 _target{}; + std::vector _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 hog_descriptor(const SimpleTensor &src, BorderMode border_mod return desc; } +template void hog_orientation_binning(const SimpleTensor &mag, const SimpleTensor &phase, SimpleTensor &hog_space, const HOGInfo &hog_info); +template void hog_block_normalization(SimpleTensor &desc, const SimpleTensor &hog_space, const HOGInfo &hog_info); template SimpleTensor hog_descriptor(const SimpleTensor &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 +void hog_orientation_binning(const SimpleTensor &mag, const SimpleTensor &phase, SimpleTensor &hog_space, const HOGInfo &hog_info); + +template +void hog_block_normalization(SimpleTensor &desc, const SimpleTensor &hog_space, const HOGInfo &hog_info); + template SimpleTensor hog_descriptor(const SimpleTensor &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 +std::vector hog_detector(const SimpleTensor &src, const std::vector &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 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 hog_detector(const SimpleTensor &src, const std::vector &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 + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +std::vector hog_detector(const SimpleTensor &src, const std::vector &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 -- cgit v1.2.1