21.02
|
Interface for the im2col reshape kernel. More...
#include <CLIm2ColKernel.h>
Public Member Functions | |
CLIm2ColKernel () | |
Default constructor. More... | |
CLIm2ColKernel (const CLIm2ColKernel &)=delete | |
Prevent instances of this class from being copied (As this class contains pointers) More... | |
CLIm2ColKernel & | operator= (const CLIm2ColKernel &)=delete |
Prevent instances of this class from being copied (As this class contains pointers) More... | |
CLIm2ColKernel (CLIm2ColKernel &&)=default | |
Allow instances of this class to be moved. More... | |
CLIm2ColKernel & | operator= (CLIm2ColKernel &&)=default |
Allow instances of this class to be moved. More... | |
void | configure (const ICLTensor *input, ICLTensor *output, const Size2D &kernel_dims, const PadStrideInfo &conv_info, bool has_bias, const Size2D &dilation=Size2D(1U, 1U), unsigned int num_groups=1) |
Set the input and output of the kernel. More... | |
void | configure (const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output, const Size2D &kernel_dims, const PadStrideInfo &conv_info, bool has_bias, const Size2D &dilation=Size2D(1U, 1U), unsigned int num_groups=1) |
Set the input and output of the kernel. More... | |
void | run (const Window &window, cl::CommandQueue &queue) override |
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... | |
Static Public Member Functions | |
static Status | validate (const ITensorInfo *input, const ITensorInfo *output, const Size2D &kernel_dims, const PadStrideInfo &conv_info, bool has_bias, const Size2D &dilation=Size2D(1U, 1U), unsigned int num_groups=1) |
Static function to check if given info will lead to a valid configuration of CLIm2ColKernel. More... | |
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... | |
Data Fields | |
const ICLTensor * | _input |
ICLTensor * | _output |
DataLayout | _data_layout |
std::pair< unsigned int, unsigned int > | _convolved_dims |
unsigned int | _num_elems_processed_per_iteration |
Size2D | _kernel_dims |
PadStrideInfo | _conv_info |
unsigned int | _num_groups |
Interface for the im2col reshape kernel.
Rearranges image blocks into columns. It is used to strip out each convolution block to a single column. It is used to transform a convolution to a plain matrix multiplication.
For example taking into account the image below and assuming 3x3 image blocks with stride of 1 we have:
\[ \left( \begin{array}{cccc} a00 & a01 & a02 & a03 \\ a10 & a11 & a12 & a13 \\ a20 & a21 & a22 & a23 \\ a30 & a31 & a32 & a33 \\ \end{array} \right) = \left( \begin{array}{ccccccccc} a00 & a01 & a02 & a10 & a11 & a12 & a20 & a21 & a22 \\ a01 & a02 & a03 & a11 & a12 & a13 & a21 & a22 & a23 \\ a10 & a11 & a12 & a20 & a21 & a22 & a30 & a31 & a32 \\ a11 & a12 & a13 & a21 & a22 & a23 & a31 & a32 & a33 \\ \end{array} \right) \]
Definition at line 56 of file CLIm2ColKernel.h.
CLIm2ColKernel | ( | ) |
Default constructor.
Definition at line 302 of file CLIm2ColKernel.cpp.
|
delete |
Prevent instances of this class from being copied (As this class contains pointers)
|
default |
Allow instances of this class to be moved.
void configure | ( | const ICLTensor * | input, |
ICLTensor * | output, | ||
const Size2D & | kernel_dims, | ||
const PadStrideInfo & | conv_info, | ||
bool | has_bias, | ||
const Size2D & | dilation = Size2D(1U, 1U) , |
||
unsigned int | num_groups = 1 |
||
) |
Set the input and output of the kernel.
[in] | input | The input tensor to convert. 3 lower dimensions represent a single input [width, height, IFM], while every optional dimension from 4 and above represent a batch of inputs. Data types supported: QASYMM8/QASYMM8_SIGNED/F16/F32 |
[out] | output | The output tensor. First 2 lower dimensions represent a transform of each 3D input, while every dimension above represents a batch. Data types supported: Same as input |
[in] | kernel_dims | The kernel dimensions (width and height). |
[in] | conv_info | Contains padding and stride information described in PadStrideInfo. |
[in] | has_bias | In case biases are provided expands the matrix with 1. This is valid only for non-quantized inputs. |
[in] | dilation | (Optional) Dilation, in elements, across x and y. Defaults to (1, 1). |
[in] | num_groups | (Optional) Number of groups when performing a grouped convolution. Number of groups other than 1 is only supported for NCHW data layout. Number of groups should be multiple to the number of channels. |
Definition at line 307 of file CLIm2ColKernel.cpp.
References CLKernelLibrary::get().
void configure | ( | const CLCompileContext & | compile_context, |
const ICLTensor * | input, | ||
ICLTensor * | output, | ||
const Size2D & | kernel_dims, | ||
const PadStrideInfo & | conv_info, | ||
bool | has_bias, | ||
const Size2D & | dilation = Size2D(1U, 1U) , |
||
unsigned int | num_groups = 1 |
||
) |
Set the input and output of the kernel.
[in] | compile_context | The compile context to be used. |
[in] | input | The input tensor to convert. 3 lower dimensions represent a single input [width, height, IFM], while every optional dimension from 4 and above represent a batch of inputs. Data types supported: QASYMM8/QASYMM8_SIGNED/F16/F32 |
[out] | output | The output tensor. First 2 lower dimensions represent a transform of each 3D input, while every dimension above represents a batch. Data types supported: Same as input |
[in] | kernel_dims | The kernel dimensions (width and height). |
[in] | conv_info | Contains padding and stride information described in PadStrideInfo. |
[in] | has_bias | In case biases are provided expands the matrix with 1. |
[in] | dilation | (Optional) Dilation, in elements, across x and y. Defaults to (1, 1). |
[in] | num_groups | (Optional) Number of groups when performing a grouped convolution. num_groups != 1 is only supported for NCHW data layout |
Definition at line 313 of file CLIm2ColKernel.cpp.
References CLIm2ColKernel::_conv_info, CLIm2ColKernel::_convolved_dims, CLIm2ColKernel::_data_layout, CLIm2ColKernel::_input, CLIm2ColKernel::_kernel_dims, CLIm2ColKernel::_num_elems_processed_per_iteration, CLIm2ColKernel::_num_groups, CLIm2ColKernel::_output, ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_NULLPTR, ARM_COMPUTE_ERROR_THROW_ON, arm_compute::test::validation::conv_info, arm_compute::create_kernel(), ITensorInfo::data_layout(), ITensorInfo::data_type(), ITensorInfo::dimension(), arm_compute::get_data_layout_dimension_index(), arm_compute::get_padding_info(), arm_compute::test::validation::has_bias, arm_compute::has_padding_changed(), Size2D::height, arm_compute::HEIGHT, ITensor::info(), arm_compute::test::validation::input, input_height, input_width, arm_compute::lower_string(), arm_compute::NHWC, arm_compute::test::validation::num_groups, arm_compute::scaled_dimensions(), arm_compute::string_from_data_layout(), arm_compute::string_from_data_type(), arm_compute::support::cpp11::to_string(), arm_compute::validate_arguments(), Size2D::width, and arm_compute::WIDTH.
|
delete |
Prevent instances of this class from being copied (As this class contains pointers)
|
default |
Allow instances of this class to be moved.
|
overridevirtual |
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 377 of file CLIm2ColKernel.cpp.
References CLIm2ColKernel::_convolved_dims, CLIm2ColKernel::_data_layout, CLIm2ColKernel::_input, CLIm2ColKernel::_num_elems_processed_per_iteration, CLIm2ColKernel::_num_groups, CLIm2ColKernel::_output, ICLKernel::add_2D_tensor_argument(), ICLKernel::add_3D_tensor_argument(), ARM_COMPUTE_ERROR_ON_MISMATCHING_WINDOWS, ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, arm_compute::ceil_to_multiple(), Window::collapse_if_possible(), Window::DimX, Window::DimY, Window::DimZ, arm_compute::enqueue(), Window::first_slice_window_2D(), Window::first_slice_window_3D(), ITensor::info(), ICLKernel::lws_hint(), arm_compute::NHWC, ICLKernel::num_arguments_per_2D_tensor(), ICLKernel::num_arguments_per_3D_tensor(), Window::set(), Window::set_dimension_step(), arm_compute::test::validation::reference::slice(), Window::slide_window_slice_2D(), Window::slide_window_slice_3D(), ITensorInfo::strides_in_bytes(), ITensorInfo::tensor_shape(), Window::use_tensor_dimensions(), and IKernel::window().
|
static |
Static function to check if given info will lead to a valid configuration of CLIm2ColKernel.
[in] | input | The input tensor to convert. 3 lower dimensions represent a single input [width, height, IFM], while every optional dimension from 4 and above represent a batch of inputs. Data types supported: QASYMM8/QASYMM8_SIGNED/F16/F32 |
[in] | output | The output tensor. First 2 lower dimensions represent a transform of each 3D input, while every dimension above represents a batch. Data types supported: Same as input |
[in] | kernel_dims | The kernel dimensions (width and height). |
[in] | conv_info | Contains padding and stride information described in PadStrideInfo. |
[in] | has_bias | In case biases are provided expands the matrix with 1. This is valid only for non-quantized inputs. |
[in] | dilation | (Optional) Dilation, in elements, across x and y. Defaults to (1, 1). |
[in] | num_groups | (Optional) Number of groups when performing a grouped convolution. Number of groups other than 1 is only supported for NCHW data layout. Number of groups should be multiple to the number of channels. |
Definition at line 366 of file CLIm2ColKernel.cpp.
References ARM_COMPUTE_RETURN_ON_ERROR, ICloneable< T >::clone(), arm_compute::test::validation::conv_info, arm_compute::test::validation::has_bias, arm_compute::test::validation::num_groups, and arm_compute::validate_arguments().
Referenced by CLGEMMConvolutionLayer::validate().
PadStrideInfo _conv_info |
Definition at line 132 of file CLIm2ColKernel.h.
Referenced by CLIm2ColKernel::configure().
std::pair<unsigned int, unsigned int> _convolved_dims |
Definition at line 129 of file CLIm2ColKernel.h.
Referenced by CLIm2ColKernel::configure(), and CLIm2ColKernel::run().
DataLayout _data_layout |
Definition at line 128 of file CLIm2ColKernel.h.
Referenced by CLIm2ColKernel::configure(), and CLIm2ColKernel::run().
const ICLTensor* _input |
Definition at line 126 of file CLIm2ColKernel.h.
Referenced by CLIm2ColKernel::configure(), and CLIm2ColKernel::run().
Size2D _kernel_dims |
Definition at line 131 of file CLIm2ColKernel.h.
Referenced by CLIm2ColKernel::configure().
unsigned int _num_elems_processed_per_iteration |
Definition at line 130 of file CLIm2ColKernel.h.
Referenced by CLIm2ColKernel::configure(), and CLIm2ColKernel::run().
unsigned int _num_groups |
Definition at line 133 of file CLIm2ColKernel.h.
Referenced by CLIm2ColKernel::configure(), and CLIm2ColKernel::run().
ICLTensor* _output |
Definition at line 127 of file CLIm2ColKernel.h.
Referenced by CLIm2ColKernel::configure(), and CLIm2ColKernel::run().