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