21.02
|
Interface for the direct convolution kernel. More...
#include <CLDirectConvolutionLayerKernel.h>
Public Member Functions | |
CLDirectConvolutionLayerKernel () | |
Default constructor. More... | |
CLDirectConvolutionLayerKernel (const CLDirectConvolutionLayerKernel &)=delete | |
Prevent instances of this class from being copied (As this class contains pointers) More... | |
CLDirectConvolutionLayerKernel & | operator= (const CLDirectConvolutionLayerKernel &)=delete |
Prevent instances of this class from being copied (As this class contains pointers) More... | |
CLDirectConvolutionLayerKernel (CLDirectConvolutionLayerKernel &&)=default | |
Allow instances of this class to be moved. More... | |
CLDirectConvolutionLayerKernel & | operator= (CLDirectConvolutionLayerKernel &&)=default |
Allow instances of this class to be moved. More... | |
~CLDirectConvolutionLayerKernel ()=default | |
Default destructor. More... | |
void | configure (const ICLTensor *input, const ICLTensor *weights, const ICLTensor *biases, ICLTensor *output, const PadStrideInfo &conv_info) |
Set the input, weights, biases and output tensors. More... | |
void | configure (const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *weights, const ICLTensor *biases, ICLTensor *output, const PadStrideInfo &conv_info) |
Set the input, weights, biases and output tensors. 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... | |
BorderSize | border_size () const override |
The size of the border for that kernel. 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... | |
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 *weights, const ITensorInfo *biases, const ITensorInfo *output, const PadStrideInfo &conv_info, const GPUTarget target) |
Static function to check if given info will lead to a valid configuration of CLDirectConvolutionLayerKernel. 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 |
const ICLTensor * | _biases |
const ICLTensor * | _weights |
ICLTensor * | _output |
DataLayout | _data_layout |
BorderSize | _border_size |
int | _conv_stride_x |
int | _conv_stride_y |
PadStrideInfo | _conv_info |
Interface for the direct convolution kernel.
Definition at line 36 of file CLDirectConvolutionLayerKernel.h.
Default constructor.
Definition at line 319 of file CLDirectConvolutionLayerKernel.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.
|
overridevirtual |
The size of the border for that kernel.
Reimplemented from IKernel.
Definition at line 324 of file CLDirectConvolutionLayerKernel.cpp.
References CLDirectConvolutionLayerKernel::_border_size.
Referenced by CLDirectConvolutionLayerKernel::configure().
void configure | ( | const ICLTensor * | input, |
const ICLTensor * | weights, | ||
const ICLTensor * | biases, | ||
ICLTensor * | output, | ||
const PadStrideInfo & | conv_info | ||
) |
Set the input, weights, biases and output tensors.
[in] | input | The input tensor to convolve. 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_SIGNED/QASYMM8/F16/F32. |
[in] | weights | Weights tensor. Weights are 4D tensor with dimensions [kernel_x, kernel_y, IFM, OFM]. The 3rd dimension must be the same as the input's volume 3rd dimension. Data type supported:Same as input . |
[in] | biases | Biases tensor. Biases are 1D tensor with dimension [OFM]. Data type supported: Should match input data type, except for input of QASYMM8 and QASYMM8_SIGNED type where biases should be of S32 type |
[out] | output | Output tensor. The 3rd dimensions must be equal to the 4th dimension of the kernels tensor. Data types supported: Same as input . |
[in] | conv_info | Contains padding and stride information described in PadStrideInfo. |
Definition at line 329 of file CLDirectConvolutionLayerKernel.cpp.
References CLKernelLibrary::get().
Referenced by arm_compute::test::validation::TEST_CASE().
void configure | ( | const CLCompileContext & | compile_context, |
const ICLTensor * | input, | ||
const ICLTensor * | weights, | ||
const ICLTensor * | biases, | ||
ICLTensor * | output, | ||
const PadStrideInfo & | conv_info | ||
) |
Set the input, weights, biases and output tensors.
[in] | compile_context | The compile context to be used. |
[in] | input | The input tensor to convolve. 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_SIGNED/QASYMM8/F16/F32. |
[in] | weights | Weights tensor. Weights are 4D tensor with dimensions [kernel_x, kernel_y, IFM, OFM]. The 3rd dimension must be the same as the input's volume 3rd dimension. Data type supported:Same as input . |
[in] | biases | Biases tensor. Biases are 1D tensor with dimension [OFM]. Data type supported: Should match input data type, except for input of QASYMM8 and QASYMM8_SIGNED type where biases should be of S32 type |
[out] | output | Output tensor. The 3rd dimensions must be equal to the 4th dimension of the kernels tensor. Data types supported: Same as input . |
[in] | conv_info | Contains padding and stride information described in PadStrideInfo. |
Definition at line 334 of file CLDirectConvolutionLayerKernel.cpp.
References CLDirectConvolutionLayerKernel::_biases, CLDirectConvolutionLayerKernel::_border_size, CLDirectConvolutionLayerKernel::_conv_info, CLDirectConvolutionLayerKernel::_conv_stride_x, CLDirectConvolutionLayerKernel::_conv_stride_y, CLDirectConvolutionLayerKernel::_data_layout, CLDirectConvolutionLayerKernel::_input, CLDirectConvolutionLayerKernel::_output, CLDirectConvolutionLayerKernel::_weights, CLBuildOptions::add_option(), CLBuildOptions::add_option_if(), arm_compute::adjust_vec_size(), ARM_COMPUTE_ERROR_ON_NULLPTR, ARM_COMPUTE_ERROR_THROW_ON, CLDirectConvolutionLayerKernel::border_size(), build_options, arm_compute::quantization::calculate_quantized_multiplier(), arm_compute::CHANNEL, arm_compute::test::validation::conv_info, arm_compute::create_kernel(), ITensorInfo::data_layout(), ITensorInfo::data_type(), ITensorInfo::dimension(), PixelValue::get(), arm_compute::get_cl_type_from_data_type(), arm_compute::get_data_layout_dimension_index(), arm_compute::get_data_size_from_data_type(), ICLKernel::get_target(), arm_compute::HEIGHT, ITensor::info(), arm_compute::test::validation::input, arm_compute::is_data_type_quantized(), kernel_name, arm_compute::lower_string(), arm_compute::NHWC, UniformQuantizationInfo::offset, CLBuildOptions::options(), PadStrideInfo::pad_left(), PadStrideInfo::pad_top(), ITensorInfo::padding(), ITensorInfo::quantization_info(), arm_compute::S32, UniformQuantizationInfo::scale, PadStrideInfo::stride(), arm_compute::string_from_data_layout(), arm_compute::string_from_data_type(), arm_compute::support::cpp11::to_string(), QuantizationInfo::uniform(), arm_compute::validate_arguments(), 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 525 of file CLDirectConvolutionLayerKernel.cpp.
References CLDirectConvolutionLayerKernel::_biases, CLDirectConvolutionLayerKernel::_conv_info, CLDirectConvolutionLayerKernel::_conv_stride_x, CLDirectConvolutionLayerKernel::_conv_stride_y, CLDirectConvolutionLayerKernel::_data_layout, CLDirectConvolutionLayerKernel::_input, CLDirectConvolutionLayerKernel::_output, CLDirectConvolutionLayerKernel::_weights, ICLKernel::add_1D_tensor_argument(), ICLKernel::add_3D_tensor_argument(), Window::adjust(), ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW, ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, ITensorInfo::dimension(), Window::DimX, Window::DimY, Window::DimZ, arm_compute::enqueue(), Window::first_slice_window_3D(), arm_compute::get_data_layout_dimension_index(), arm_compute::HEIGHT, ITensor::info(), ICLKernel::lws_hint(), arm_compute::NHWC, ICLKernel::num_arguments_per_3D_tensor(), PadStrideInfo::pad_left(), PadStrideInfo::pad_top(), Window::set(), Window::set_dimension_step(), arm_compute::test::validation::reference::slice(), Window::slide_window_slice_3D(), arm_compute::cpu::step, ITensorInfo::strides_in_bytes(), ITensorInfo::tensor_shape(), Window::use_tensor_dimensions(), arm_compute::WIDTH, and IKernel::window().
|
static |
Static function to check if given info will lead to a valid configuration of CLDirectConvolutionLayerKernel.
[in] | input | The input tensor to convolve. 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_SIGNED/QASYMM8/F16/F32. |
[in] | weights | Weights tensor. Weights are 4D tensor with dimensions [kernel_x, kernel_y, IFM, OFM]. The 3rd dimension must be the same as the input's volume 3rd dimension. Data type supported:Same as input . |
[in] | biases | Biases tensor. Biases are 1D tensor with dimension [OFM]. Data type supported: Should match input data type, except for input of QASYMM8 and QASYMM8_SIGNED type where biases should be of S32 type. |
[in] | output | Output tensor. The 3rd dimensions must be equal to the 4th dimension of the kernels tensor. Data types supported: Same as input . |
[in] | conv_info | Contains padding and stride information described in PadStrideInfo. |
[in] | target | Target GPU architecture. |
Definition at line 516 of file CLDirectConvolutionLayerKernel.cpp.
References ARM_COMPUTE_RETURN_ON_ERROR, ICloneable< T >::clone(), arm_compute::test::validation::conv_info, and arm_compute::validate_arguments().
Referenced by CLDirectConvolutionLayer::validate().
const ICLTensor* _biases |
Definition at line 116 of file CLDirectConvolutionLayerKernel.h.
Referenced by CLDirectConvolutionLayerKernel::configure(), and CLDirectConvolutionLayerKernel::run().
BorderSize _border_size |
Definition at line 120 of file CLDirectConvolutionLayerKernel.h.
Referenced by CLDirectConvolutionLayerKernel::border_size(), and CLDirectConvolutionLayerKernel::configure().
PadStrideInfo _conv_info |
Definition at line 123 of file CLDirectConvolutionLayerKernel.h.
Referenced by CLDirectConvolutionLayerKernel::configure(), and CLDirectConvolutionLayerKernel::run().
int _conv_stride_x |
Definition at line 121 of file CLDirectConvolutionLayerKernel.h.
Referenced by CLDirectConvolutionLayerKernel::configure(), and CLDirectConvolutionLayerKernel::run().
int _conv_stride_y |
Definition at line 122 of file CLDirectConvolutionLayerKernel.h.
Referenced by CLDirectConvolutionLayerKernel::configure(), and CLDirectConvolutionLayerKernel::run().
DataLayout _data_layout |
Definition at line 119 of file CLDirectConvolutionLayerKernel.h.
Referenced by CLDirectConvolutionLayerKernel::configure(), and CLDirectConvolutionLayerKernel::run().
const ICLTensor* _input |
Definition at line 115 of file CLDirectConvolutionLayerKernel.h.
Referenced by CLDirectConvolutionLayerKernel::configure(), and CLDirectConvolutionLayerKernel::run().
ICLTensor* _output |
Definition at line 118 of file CLDirectConvolutionLayerKernel.h.
Referenced by CLDirectConvolutionLayerKernel::configure(), and CLDirectConvolutionLayerKernel::run().
const ICLTensor* _weights |
Definition at line 117 of file CLDirectConvolutionLayerKernel.h.
Referenced by CLDirectConvolutionLayerKernel::configure(), and CLDirectConvolutionLayerKernel::run().