21.02
|
OpenCL kernel to perform reshaping on the weights used by convolution and locally connected layer. More...
#include <CLWeightsReshapeKernel.h>
Public Member Functions | |
CLWeightsReshapeKernel () | |
Constructor. More... | |
CLWeightsReshapeKernel (const CLWeightsReshapeKernel &)=delete | |
Prevent instances of this class from being copied (As this class contains pointers) More... | |
CLWeightsReshapeKernel & | operator= (const CLWeightsReshapeKernel &)=delete |
Prevent instances of this class from being copied (As this class contains pointers) More... | |
CLWeightsReshapeKernel (CLWeightsReshapeKernel &&)=default | |
Allow instances of this class to be moved. More... | |
CLWeightsReshapeKernel & | operator= (CLWeightsReshapeKernel &&)=default |
Allow instances of this class to be moved. More... | |
~CLWeightsReshapeKernel ()=default | |
Default destructor. More... | |
void | configure (const ICLTensor *input, const ICLTensor *biases, ICLTensor *output, unsigned int num_groups=1) |
Set the input and output of the kernel. More... | |
void | configure (const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *biases, ICLTensor *output, 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 *biases, const ITensorInfo *output, unsigned int num_groups=1) |
Static function to check if given info will lead to a valid configuration of CLWeightsReshapeKernel. 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... | |
OpenCL kernel to perform reshaping on the weights used by convolution and locally connected layer.
Rearranges each 3-dimensional kernel to a single row leading to a matrix with linearized kernels. In combination with the CLIm2ColKernel can transform a convolution to a matrix multiplication.
For example assuming a 3D weight kernel of 3x3 dimensions and depth of 2 we have:
\[ \left( \begin{array}{ccc} a000 & a001 & a002 \\ a010 & a011 & a012 \\ a020 & a021 & a022 \\ \end{array} \right) \left( \begin{array}{ccc} a100 & a101 & a102 \\ a110 & a111 & a112 \\ a120 & a121 & a122 \\ \end{array} \right) \rightarrow \left( \begin{array}{ccccccccc} a000 & a001 & a002 & a010 & a011 & a012 & a020 & a021 & a022 & a100 & a101 & a102 & a110 & a111 & a112 & a120 & a121 & a122 \\ \end{array} \right) \]
Definition at line 54 of file CLWeightsReshapeKernel.h.
Constructor.
Definition at line 69 of file CLWeightsReshapeKernel.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 ICLTensor * | biases, | ||
ICLTensor * | output, | ||
unsigned int | num_groups = 1 |
||
) |
Set the input and output of the kernel.
[in] | input | The input tensor to convert. Weights are 4D tensor with dimensions [kernel_x, kernel_y, IFM, OFM] if shared, and 5D tensor with dimensions [kernel_x, kernel_y, IFM, OFM, num_patches] if unshared. Data types supported: All |
[in] | biases | The shared biases tensor to append. Bias is 1D tensor with dimensions [OFM] if shared and 2D tensor with dimensions [OFM, num_patches] if unshared. Data types supported: F16/F32, for quantized types this must be nullptr. |
[out] | output | The output tensor. Should be a 2D Tensor if there are no groups and the weights are not shared; a 3D Tensor otherwise. Data types supported: Same as input |
[in] | num_groups | (Optional) Number of groups when performing a grouped convolution. num_groups != 1 is only supported for NCHW data layout Number of groups greater than one are only supported for NCHW data layout, and the number of weights must be a multiple of it. |
Definition at line 74 of file CLWeightsReshapeKernel.cpp.
References CLKernelLibrary::get().
void configure | ( | const CLCompileContext & | compile_context, |
const ICLTensor * | input, | ||
const ICLTensor * | biases, | ||
ICLTensor * | output, | ||
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. Weights are 4D tensor with dimensions [kernel_x, kernel_y, IFM, OFM] if shared, and 5D tensor with dimensions [kernel_x, kernel_y, IFM, OFM, num_patches] if unshared. Data types supported: All |
[in] | biases | The shared biases tensor to append. Bias is 1D tensor with dimensions [OFM] if shared and 2D tensor with dimensions [OFM, num_patches] if unshared. Data types supported: F16/F32, for quantized types this must be nullptr. |
[out] | output | The output tensor. Should be a 2D Tensor if there are no groups and the weights are not shared; a 3D Tensor otherwise. Data types supported: Same as input |
[in] | num_groups | (Optional) Number of groups when performing a grouped convolution. num_groups != 1 is only supported for NCHW data layout Number of groups greater than one are only supported for NCHW data layout, and the number of weights must be a multiple of it. |
Definition at line 79 of file CLWeightsReshapeKernel.cpp.
References CLBuildOptions::add_option(), CLBuildOptions::add_option_if(), ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_NULLPTR, ARM_COMPUTE_ERROR_THROW_ON, arm_compute::auto_init_if_empty(), arm_compute::calculate_max_window(), ICloneable< T >::clone(), arm_compute::misc::shape_calculator::compute_weights_reshaped_shape(), arm_compute::create_kernel(), arm_compute::data_size_from_type(), arm_compute::test::validation::data_type, ITensorInfo::data_type(), arm_compute::get_cl_unsigned_type_from_element_size(), arm_compute::get_padding_info(), arm_compute::has_padding_changed(), ITensor::info(), arm_compute::test::validation::input, arm_compute::test::validation::num_groups, CLBuildOptions::options(), ITensorInfo::set_valid_region(), ITensorInfo::tensor_shape(), arm_compute::support::cpp11::to_string(), and arm_compute::validate_arguments().
|
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 123 of file CLWeightsReshapeKernel.cpp.
References ICLKernel::add_1D_tensor_argument(), ICLKernel::add_2D_tensor_argument(), ICLKernel::add_3D_tensor_argument(), ARM_COMPUTE_ERROR_ON_MISMATCHING_WINDOWS, ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, ARM_COMPUTE_UNUSED, ITensorInfo::dimension(), arm_compute::enqueue(), Window::first_slice_window_1D(), Window::first_slice_window_2D(), Window::first_slice_window_3D(), ITensor::info(), ICLKernel::lws_hint(), ICLKernel::num_arguments_per_1D_tensor(), ICLKernel::num_arguments_per_2D_tensor(), ICLKernel::num_arguments_per_3D_tensor(), Window::slide_window_slice_1D(), Window::slide_window_slice_2D(), Window::slide_window_slice_4D(), ITensorInfo::strides_in_bytes(), ITensorInfo::tensor_shape(), Window::use_tensor_dimensions(), IKernel::window(), and Dimensions< T >::z().
|
static |
Static function to check if given info will lead to a valid configuration of CLWeightsReshapeKernel.
[in] | input | The input tensor to convert. Weights are 4D tensor with dimensions [kernel_x, kernel_y, IFM, OFM] if shared, and 5D tensor with dimensions [kernel_x, kernel_y, IFM, OFM, num_patches] if unshared. Data types supported: All |
[in] | biases | The shared biases tensor to append. Bias is 1D tensor with dimensions [OFM] if shared and 2D tensor with dimensions [OFM, num_patches] if unshared. Data types supported: F16/F32, for quantized types this must be nullptr. |
[in] | output | The output tensor. Should be a 2D Tensor if there are no groups and the weights are not shared; a 3D Tensor otherwise. Data types supported: Same as input |
[in] | num_groups | (Optional) Number of groups when performing a grouped convolution. num_groups != 1 is only supported for NCHW data layout Number of groups greater than one are only supported for NCHW data layout, and the number of weights must be a multiple of it. |
Definition at line 117 of file CLWeightsReshapeKernel.cpp.
References ARM_COMPUTE_RETURN_ON_ERROR, and arm_compute::validate_arguments().
Referenced by CLConvolutionLayerReshapeWeights::validate().