21.02
|
OpenCL kernel to perform HOG detector kernel using linear SVM. More...
#include <CLHOGDetectorKernel.h>
Public Member Functions | |
CLHOGDetectorKernel () | |
Default constructor. More... | |
CLHOGDetectorKernel (const CLHOGDetectorKernel &)=delete | |
Prevent instances of this class from being copied (As this class contains pointers) More... | |
CLHOGDetectorKernel & | operator= (const CLHOGDetectorKernel &)=delete |
Prevent instances of this class from being copied (As this class contains pointers) More... | |
CLHOGDetectorKernel (CLHOGDetectorKernel &&)=default | |
Allow instances of this class to be moved. More... | |
CLHOGDetectorKernel & | operator= (CLHOGDetectorKernel &&)=default |
Allow instances of this class to be moved. More... | |
~CLHOGDetectorKernel ()=default | |
Default destructor. More... | |
void | configure (const ICLTensor *input, const ICLHOG *hog, ICLDetectionWindowArray *detection_windows, cl::Buffer *num_detection_windows, const Size2D &detection_window_stride, float threshold=0.0f, uint16_t idx_class=0) |
Initialise the kernel's input, HOG data-object, detection window, the stride of the detection window, the threshold and index of the object to detect. More... | |
void | configure (const CLCompileContext &compile_context, const ICLTensor *input, const ICLHOG *hog, ICLDetectionWindowArray *detection_windows, cl::Buffer *num_detection_windows, const Size2D &detection_window_stride, float threshold=0.0f, uint16_t idx_class=0) |
Initialise the kernel's input, HOG data-object, detection window, the stride of the detection window, the threshold and index of the object to detect. More... | |
void | run (const Window &window, cl::CommandQueue &queue) |
Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue. More... | |
Public Member Functions inherited from ICLKernel | |
ICLKernel () | |
Constructor. More... | |
cl::Kernel & | kernel () |
Returns a reference to the OpenCL kernel of this object. More... | |
template<typename T > | |
void | add_1D_array_argument (unsigned int &idx, const ICLArray< T > *array, const Strides &strides, unsigned int num_dimensions, const Window &window) |
Add the passed 1D array's parameters to the object's kernel's arguments starting from the index idx. More... | |
void | add_1D_tensor_argument (unsigned int &idx, const ICLTensor *tensor, const Window &window) |
Add the passed 1D tensor's parameters to the object's kernel's arguments starting from the index idx. More... | |
void | add_1D_tensor_argument_if (bool cond, unsigned int &idx, const ICLTensor *tensor, const Window &window) |
Add the passed 1D tensor's parameters to the object's kernel's arguments starting from the index idx if the condition is true. More... | |
void | add_2D_tensor_argument (unsigned int &idx, const ICLTensor *tensor, const Window &window) |
Add the passed 2D tensor's parameters to the object's kernel's arguments starting from the index idx. More... | |
void | add_2D_tensor_argument_if (bool cond, unsigned int &idx, const ICLTensor *tensor, const Window &window) |
Add the passed 2D tensor's parameters to the object's kernel's arguments starting from the index idx if the condition is true. More... | |
void | add_3D_tensor_argument (unsigned int &idx, const ICLTensor *tensor, const Window &window) |
Add the passed 3D tensor's parameters to the object's kernel's arguments starting from the index idx. More... | |
void | add_4D_tensor_argument (unsigned int &idx, const ICLTensor *tensor, const Window &window) |
Add the passed 4D tensor's parameters to the object's kernel's arguments starting from the index idx. More... | |
virtual void | run_op (ITensorPack &tensors, const Window &window, cl::CommandQueue &queue) |
Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue. More... | |
template<typename T > | |
void | add_argument (unsigned int &idx, T value) |
Add the passed parameters to the object's kernel's arguments starting from the index idx. More... | |
void | set_lws_hint (const cl::NDRange &lws_hint) |
Set the Local-Workgroup-Size hint. More... | |
cl::NDRange | lws_hint () const |
Return the Local-Workgroup-Size hint. More... | |
void | set_wbsm_hint (const cl_int &wbsm_hint) |
Set the workgroup batch size modifier hint. More... | |
cl_int | wbsm_hint () const |
Return the workgroup batch size modifier hint. More... | |
const std::string & | config_id () const |
Get the configuration ID. More... | |
void | set_target (GPUTarget target) |
Set the targeted GPU architecture. More... | |
void | set_target (cl::Device &device) |
Set the targeted GPU architecture according to the CL device. More... | |
GPUTarget | get_target () const |
Get the targeted GPU architecture. More... | |
size_t | get_max_workgroup_size () |
Get the maximum workgroup size for the device the CLKernelLibrary uses. More... | |
template<unsigned int dimension_size> | |
void | add_tensor_argument (unsigned &idx, const ICLTensor *tensor, const Window &window) |
template<typename T , unsigned int dimension_size> | |
void | add_array_argument (unsigned &idx, const ICLArray< T > *array, const Strides &strides, unsigned int num_dimensions, const Window &window) |
Add the passed array's parameters to the object's kernel's arguments starting from the index idx. More... | |
Public Member Functions inherited from IKernel | |
IKernel () | |
Constructor. More... | |
virtual | ~IKernel ()=default |
Destructor. More... | |
virtual bool | is_parallelisable () const |
Indicates whether or not the kernel is parallelisable. More... | |
virtual BorderSize | border_size () const |
The size of the border for that kernel. More... | |
const Window & | window () const |
The maximum window the kernel can be executed on. More... | |
Additional Inherited Members | |
Static Public Member Functions inherited from ICLKernel | |
static constexpr unsigned int | num_arguments_per_1D_array () |
Returns the number of arguments enqueued per 1D array object. More... | |
static constexpr unsigned int | num_arguments_per_1D_tensor () |
Returns the number of arguments enqueued per 1D tensor object. More... | |
static constexpr unsigned int | num_arguments_per_2D_tensor () |
Returns the number of arguments enqueued per 2D tensor object. More... | |
static constexpr unsigned int | num_arguments_per_3D_tensor () |
Returns the number of arguments enqueued per 3D tensor object. More... | |
static constexpr unsigned int | num_arguments_per_4D_tensor () |
Returns the number of arguments enqueued per 4D tensor object. More... | |
static cl::NDRange | gws_from_window (const Window &window) |
Get the global work size given an execution window. More... | |
OpenCL kernel to perform HOG detector kernel using linear SVM.
Definition at line 42 of file CLHOGDetectorKernel.h.
Default constructor.
Definition at line 38 of file CLHOGDetectorKernel.cpp.
|
delete |
Prevent instances of this class from being copied (As this class contains pointers)
|
default |
Allow instances of this class to be moved.
|
default |
Default destructor.
void configure | ( | const ICLTensor * | input, |
const ICLHOG * | hog, | ||
ICLDetectionWindowArray * | detection_windows, | ||
cl::Buffer * | num_detection_windows, | ||
const Size2D & | detection_window_stride, | ||
float | threshold = 0.0f , |
||
uint16_t | idx_class = 0 |
||
) |
Initialise the kernel's input, HOG data-object, detection window, the stride of the detection window, the threshold and index of the object to detect.
[in] | input | Input tensor which stores the HOG descriptor obtained with CLHOGOrientationBinningKernel. Data type supported: F32. Number of channels supported: equal to the number of histogram bins per block |
[in] | hog | HOG data object used by CLHOGOrientationBinningKernel and CLHOGBlockNormalizationKernel |
[out] | detection_windows | Array of DetectionWindow. This array stores all the detected objects |
[in] | num_detection_windows | Number of detected objects |
[in] | detection_window_stride | Distance in pixels between 2 consecutive detection windows in x and y directions. It must be multiple of the hog->info()->block_stride() |
[in] | threshold | (Optional) Threshold for the distance between features and SVM classifying plane |
[in] | idx_class | (Optional) Index of the class used for evaluating which class the detection window belongs to |
Definition at line 43 of file CLHOGDetectorKernel.cpp.
References CLKernelLibrary::get().
void configure | ( | const CLCompileContext & | compile_context, |
const ICLTensor * | input, | ||
const ICLHOG * | hog, | ||
ICLDetectionWindowArray * | detection_windows, | ||
cl::Buffer * | num_detection_windows, | ||
const Size2D & | detection_window_stride, | ||
float | threshold = 0.0f , |
||
uint16_t | idx_class = 0 |
||
) |
Initialise the kernel's input, HOG data-object, detection window, the stride of the detection window, the threshold and index of the object to detect.
[in] | compile_context | The compile context to be used. |
[in] | input | Input tensor which stores the HOG descriptor obtained with CLHOGOrientationBinningKernel. Data type supported: F32. Number of channels supported: equal to the number of histogram bins per block |
[in] | hog | HOG data object used by CLHOGOrientationBinningKernel and CLHOGBlockNormalizationKernel |
[out] | detection_windows | Array of DetectionWindow. This array stores all the detected objects |
[in] | num_detection_windows | Number of detected objects |
[in] | detection_window_stride | Distance in pixels between 2 consecutive detection windows in x and y directions. It must be multiple of the hog->info()->block_stride() |
[in] | threshold | (Optional) Threshold for the distance between features and SVM classifying plane |
[in] | idx_class | (Optional) Index of the class used for evaluating which class the detection window belongs to |
Definition at line 49 of file CLHOGDetectorKernel.cpp.
References ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_DATA_TYPE_NOT_IN, HOGInfo::block_size(), HOGInfo::block_stride(), ICLHOG::cl_buffer(), ICLArray< T >::cl_buffer(), arm_compute::create_kernel(), ITensorInfo::data_type(), HOGInfo::descriptor_size(), HOGInfo::detection_window_size(), ITensorInfo::dimension(), Window::DimX, Window::DimY, arm_compute::F32, arm_compute::floor_to_multiple(), Size2D::height, IHOG::info(), ITensor::info(), arm_compute::test::validation::input, kernel_name, arm_compute::lower_string(), IArray< T >::max_num_values(), ICLKernel::num_arguments_per_2D_tensor(), ITensorInfo::num_channels(), Window::set(), ValidRegion::shape, arm_compute::string_from_data_type(), arm_compute::support::cpp11::to_string(), arm_compute::update_window_and_padding(), arm_compute::test::validation::valid_region, ITensorInfo::valid_region(), and Size2D::width.
|
delete |
Prevent instances of this class from being copied (As this class contains pointers)
|
default |
Allow instances of this class to be moved.
|
virtual |
Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue.
[in] | window | Region on which to execute the kernel. (Must be a valid region of the window returned by window()). |
[in,out] | queue | Command queue on which to enqueue the kernel. |
Reimplemented from ICLKernel.
Definition at line 132 of file CLHOGDetectorKernel.cpp.
References ICLKernel::add_2D_tensor_argument(), ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW, ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, arm_compute::enqueue(), Window::first_slice_window_2D(), ICLKernel::lws_hint(), arm_compute::test::validation::reference::slice(), Window::slide_window_slice_2D(), and IKernel::window().