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>
diff --git a/src/core/CL/cl_kernels/hog.cl b/src/core/CL/cl_kernels/hog.cl
index 3d37fbc..407ee2f 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 @@
}
#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 @@
* -# -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 @@
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 @@
}
}
#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 87659c4..a15aab1 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 @@
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 0f9a989..caca498 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 @@
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 @@
// 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 73f4c42..4f99455 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 @@
_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 @@
_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 3fd81be..abe224e 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 @@
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 343b051..2c02ab8 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 @@
// 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 2190e34..539d0f8 100644
--- a/src/core/TensorInfo.cpp
+++ b/src/core/TensorInfo.cpp
@@ -168,13 +168,13 @@
// 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 @@
// 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()