Common interface for all the OpenCL kernels. More...

#include <ICLKernel.h>

Collaboration diagram for ICLKernel:

Public Member Functions
	ICLKernel ()
	Constructor. More...

cl::Kernel &	kernel ()
	Returns a reference to the OpenCL kernel of this object. More...

CLKernelType	type () const
	Returns the CL kernel type. 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...

void	add_5D_tensor_argument (unsigned int &idx, const ICLTensor *tensor, const Window &window)
	Add the passed 5D tensor's parameters to the object's kernel's arguments starting from the index idx. More...

void	add_3d_tensor_nhw_argument (unsigned int &idx, const ICLTensor *tensor)
	Add the passed NHW 3D tensor's parameters to the object's kernel's arguments by passing strides, dimensions and the offset to the first valid element in bytes. More...

void	add_4d_tensor_nhwc_argument (unsigned int &idx, const ICLTensor *tensor)
	Add the passed NHWC 4D tensor's parameters to the object's kernel's arguments by passing strides, dimensions and the offset to the first valid element in bytes. More...

virtual void	run (const Window &window, cl::CommandQueue &queue)
	Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue. 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...

bool	is_window_configured () const
	Function to check if the embedded window of this kernel has been configured. More...

Static Public Member Functions
static constexpr unsigned int	num_arguments_per_3d_tensor_nhw ()
	Returns the number of arguments enqueued per NHW 3D Tensor object. More...

static constexpr unsigned int	num_arguments_per_4d_tensor_nhwc ()
	Returns the number of arguments enqueued per NHWC 4D Tensor object. More...

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...

Detailed Description

Common interface for all the OpenCL kernels.

Definition at line 67 of file ICLKernel.h.

Constructor & Destructor Documentation

◆ ICLKernel()

ICLKernel ( )

inline

Constructor.

Definition at line 126 of file ICLKernel.h.

         : _kernel(nullptr), _target(GPUTarget::MIDGARD), _config_id(arm_compute::default_config_id), _max_workgroup_size(0), _type(CLKernelType::UNKNOWN), _tuning_params_hint()
     {
     }

Member Function Documentation

◆ add_1D_array_argument()

void add_1D_array_argument	(	unsigned int &	idx,
		const ICLArray< T > *	array,
		const Strides &	strides,
		unsigned int	num_dimensions,
		const Window &	window
	)

inline

Add the passed 1D array's parameters to the object's kernel's arguments starting from the index idx.

Parameters

[in,out]	idx	Index at which to start adding the array's arguments. Will be incremented by the number of kernel arguments set.
[in]	array	Array to set as an argument of the object's kernel.
[in]	strides	Strides object containing stride of each dimension in bytes.
[in]	num_dimensions	Number of dimensions of the `array`.
[in]	window	Window the kernel will be executed on.

Definition at line 155 of file ICLKernel.h.

References IKernel::window().

     {
         add_array_argument<T, 1>(idx, array, strides, num_dimensions, window);
     }

◆ add_1D_tensor_argument()

void add_1D_tensor_argument	(	unsigned int &	idx,
		const ICLTensor *	tensor,
		const Window &	window
	)

inline

Add the passed 1D tensor's parameters to the object's kernel's arguments starting from the index idx.

Parameters

[in,out]	idx	Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]	tensor	Tensor to set as an argument of the object's kernel.
[in]	window	Window the kernel will be executed on.

Definition at line 165 of file ICLKernel.h.

References tensor, and IKernel::window().

Referenced by ICLKernel::add_1D_tensor_argument_if(), CLSelectKernel::run(), CLReverseKernel::run(), CLComputeAllAnchorsKernel::run(), CLQLSTMLayerNormalizationKernel::run(), CLGatherKernel::run(), CLFFTDigitReverseKernel::run(), CLRangeKernel::run(), CLNormalizePlanarYUVLayerKernel::run(), CLReductionOperationKernel::run(), CLBoundingBoxTransformKernel::run(), CLDepthwiseConvolutionLayerNativeKernel::run(), CLDeconvolutionReshapeOutputKernel::run(), CLBatchToSpaceLayerKernel::run(), CLBatchNormalizationLayerKernel::run(), CLFuseBatchNormalizationKernel::run(), ClTransposedConvolutionKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleByFixedPointKernel::run_op(), ClIndirectConv2dKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleByFloatKernel::run_op(), ClDirectConv2dKernel::run_op(), and ClDirectConv3dKernel::run_op().

     {
         add_tensor_argument<1>(idx, tensor, window);
     }

◆ add_1D_tensor_argument_if()

void add_1D_tensor_argument_if	(	bool	cond,
		unsigned int &	idx,
		const ICLTensor *	tensor,
		const Window &	window
	)

inline

Add the passed 1D tensor's parameters to the object's kernel's arguments starting from the index idx if the condition is true.

Parameters

[in]	cond	Condition to check
[in,out]	idx	Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]	tensor	Tensor to set as an argument of the object's kernel.
[in]	window	Window the kernel will be executed on.

Definition at line 176 of file ICLKernel.h.

References ICLKernel::add_1D_tensor_argument().

Referenced by CLSpaceToBatchLayerKernel::run(), ClGemmLowpOffsetContributionKernel::run_op(), and ClGemmLowpOffsetContributionOutputStageKernel::run_op().

     {
         if(cond)
         {
             add_1D_tensor_argument(idx, tensor, window);
         }
     }

◆ add_2D_tensor_argument()

void add_2D_tensor_argument	(	unsigned int &	idx,
		const ICLTensor *	tensor,
		const Window &	window
	)

inline

Add the passed 2D tensor's parameters to the object's kernel's arguments starting from the index idx.

Parameters

[in,out]	idx	Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]	tensor	Tensor to set as an argument of the object's kernel.
[in]	window	Window the kernel will be executed on.

Definition at line 189 of file ICLKernel.h.

References tensor, and IKernel::window().

Referenced by ICLKernel::add_2D_tensor_argument_if(), ICLSimple2DKernel::run(), CLBitwiseKernel::run(), CLQLSTMLayerNormalizationKernel::run(), CLROIPoolingLayerKernel::run(), CLMeanStdDevNormalizationKernel::run(), CLReductionOperationKernel::run(), CLPriorBoxLayerKernel::run(), CLL2NormalizeLayerKernel::run(), CLBoundingBoxTransformKernel::run(), CLArgMinMaxLayerKernel::run(), CLROIAlignLayerKernel::run(), ClTransposeKernel::run_op(), ClScaleKernel::run_op(), ClConvertFullyConnectedWeightsKernel::run_op(), ClGemmLowpMatrixMultiplyNativeKernel::run_op(), ClGemmMatrixMultiplyNativeKernel::run_op(), ClGemmLowpMatrixAReductionKernel::run_op(), ClGemmMatrixMultiplyReshapedOnlyRhsKernel::run_op(), ClIm2ColKernel::run_op(), ClGemmMatrixMultiplyReshapedKernel::run_op(), and ClGemmLowpMatrixBReductionKernel::run_op().

     {
         add_tensor_argument<2>(idx, tensor, window);
     }

◆ add_2D_tensor_argument_if()

void add_2D_tensor_argument_if	(	bool	cond,
		unsigned int &	idx,
		const ICLTensor *	tensor,
		const Window &	window
	)

inline

Add the passed 2D tensor's parameters to the object's kernel's arguments starting from the index idx if the condition is true.

Parameters

[in]	cond	Condition to check
[in,out]	idx	Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]	tensor	Tensor to set as an argument of the object's kernel.
[in]	window	Window the kernel will be executed on.

Definition at line 200 of file ICLKernel.h.

References ICLKernel::add_2D_tensor_argument().

Referenced by CLMeanStdDevNormalizationKernel::run(), CLSpaceToBatchLayerKernel::run(), ClGemmLowpOffsetContributionKernel::run_op(), ClGemmLowpOffsetContributionOutputStageKernel::run_op(), ClGemmMatrixMultiplyReshapedOnlyRhsKernel::run_op(), and ClGemmMatrixMultiplyReshapedKernel::run_op().

     {
         if(cond)
         {
             add_2D_tensor_argument(idx, tensor, window);
         }
     }

◆ add_3D_tensor_argument()

void add_3D_tensor_argument	(	unsigned int &	idx,
		const ICLTensor *	tensor,
		const Window &	window
	)

inline

Add the passed 3D tensor's parameters to the object's kernel's arguments starting from the index idx.

Parameters

[in,out]	idx	Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]	tensor	Tensor to set as an argument of the object's kernel.
[in]	window	Window the kernel will be executed on.

Definition at line 213 of file ICLKernel.h.

References tensor, and IKernel::window().

     {
         add_tensor_argument<3>(idx, tensor, window);
     }

◆ add_3d_tensor_nhw_argument()

void add_3d_tensor_nhw_argument	(	unsigned int &	idx,
		const ICLTensor *	tensor
	)

Add the passed NHW 3D tensor's parameters to the object's kernel's arguments by passing strides, dimensions and the offset to the first valid element in bytes.

Parameters

[in,out]	idx	Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]	tensor	Tensor to set as an argument of the object's kernel.

Definition at line 119 of file ICLKernel.cpp.

References ARM_COMPUTE_ERROR_ON, ICLTensor::cl_buffer(), ITensorInfo::dimension(), ITensor::info(), arm_compute::test::validation::info, ITensorInfo::offset_first_element_in_bytes(), and ITensorInfo::strides_in_bytes().

Referenced by ICLKernel::add_5D_tensor_argument(), ClGemmReshapeLhsMatrixKernel::run_op(), ClGemmMatrixMultiplyReshapedOnlyRhsMMULKernel::run_op(), and ClGemmReshapeRhsMatrixKernel::run_op().

 {
     ARM_COMPUTE_ERROR_ON(tensor == nullptr);
 
     const ITensorInfo *info = tensor->info();
     ARM_COMPUTE_ERROR_ON(info == nullptr);
     const Strides &strides = info->strides_in_bytes();
 
     // Tensor poniter
     _kernel.setArg(idx++, tensor->cl_buffer());
 
     // Add stride_y, stride_z
     _kernel.setArg<cl_uint>(idx++, strides[1]);
     _kernel.setArg<cl_uint>(idx++, strides[2]);
 
     // Tensor dimensions
     _kernel.setArg<cl_uint>(idx++, info->dimension(0));
     _kernel.setArg<cl_uint>(idx++, info->dimension(1));
     _kernel.setArg<cl_uint>(idx++, info->dimension(2));
 
     // Offset of first element
     unsigned int offset_first_element = info->offset_first_element_in_bytes();
     _kernel.setArg<cl_uint>(idx++, offset_first_element);
 }

◆ add_4D_tensor_argument()

void add_4D_tensor_argument	(	unsigned int &	idx,
		const ICLTensor *	tensor,
		const Window &	window
	)

inline

Add the passed 4D tensor's parameters to the object's kernel's arguments starting from the index idx.

Parameters

[in,out]	idx	Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]	tensor	Tensor to set as an argument of the object's kernel.
[in]	window	Window the kernel will be executed on.

Definition at line 223 of file ICLKernel.h.

References tensor, and IKernel::window().

     {
         add_tensor_argument<4>(idx, tensor, window);
     }

◆ add_4d_tensor_nhwc_argument()

void add_4d_tensor_nhwc_argument	(	unsigned int &	idx,
		const ICLTensor *	tensor
	)

Add the passed NHWC 4D tensor's parameters to the object's kernel's arguments by passing strides, dimensions and the offset to the first valid element in bytes.

Parameters

[in,out]	idx	Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]	tensor	Tensor to set as an argument of the object's kernel.

Definition at line 144 of file ICLKernel.cpp.

References ARM_COMPUTE_ERROR_ON, ICLTensor::cl_buffer(), ITensorInfo::dimension(), ITensor::info(), arm_compute::test::validation::info, ITensorInfo::offset_first_element_in_bytes(), ITensorInfo::strides_in_bytes(), tensor, and IKernel::window().

Referenced by ICLKernel::num_arguments_per_3d_tensor_nhw(), CLDepthwiseConvolutionLayerNativeKernel::run(), ClTransposedConvolutionKernel::run_op(), ClScaleKernel::run_op(), ClIndirectConv2dKernel::run_op(), ClIndirectConv2dAddressPrecalculationKernel::run_op(), and ClDirectConv2dKernel::run_op().

 {
     ARM_COMPUTE_ERROR_ON(tensor == nullptr);
 
     const ITensorInfo *info = tensor->info();
     ARM_COMPUTE_ERROR_ON(info == nullptr);
     const Strides &strides = info->strides_in_bytes();
 
     // Tensor poniter
     _kernel.setArg(idx++, tensor->cl_buffer());
 
     // Add stride_y, stride_z and stride_w
     _kernel.setArg<cl_uint>(idx++, strides[1]);
     _kernel.setArg<cl_uint>(idx++, strides[2]);
     _kernel.setArg<cl_uint>(idx++, strides[3]);
 
     // Tensor dimensions
     _kernel.setArg<cl_uint>(idx++, info->dimension(0));
     _kernel.setArg<cl_uint>(idx++, info->dimension(1));
     _kernel.setArg<cl_uint>(idx++, info->dimension(2));
     _kernel.setArg<cl_uint>(idx++, info->dimension(3));
 
     // Offset of first element
     unsigned int offset_first_element = info->offset_first_element_in_bytes();
     _kernel.setArg<cl_uint>(idx++, offset_first_element);
 }

◆ add_5D_tensor_argument()

void add_5D_tensor_argument	(	unsigned int &	idx,
		const ICLTensor *	tensor,
		const Window &	window
	)

inline

Add the passed 5D tensor's parameters to the object's kernel's arguments starting from the index idx.

Parameters

[in,out]	idx	Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]	tensor	Tensor to set as an argument of the object's kernel.
[in]	window	Window the kernel will be executed on.

Definition at line 233 of file ICLKernel.h.

References ICLKernel::add_3d_tensor_nhw_argument(), tensor, and IKernel::window().

     {
         add_tensor_argument<5>(idx, tensor, window);
     }

◆ add_argument()

void add_argument	(	unsigned int &	idx,
		T	value
	)

inline

Add the passed parameters to the object's kernel's arguments starting from the index idx.

Parameters

[in,out]	idx	Index at which to start adding the arguments. Will be incremented by the number of kernel arguments set.
[in]	value	Value to set as an argument of the object's kernel.

Definition at line 341 of file ICLKernel.h.

Referenced by CLDepthToSpaceLayerKernel::run(), CLSpaceToDepthLayerKernel::run(), CLBatchToSpaceLayerKernel::run(), CLSpaceToBatchLayerKernel::run(), and ClCropKernel::run_op().

     {
         _kernel.setArg(idx++, value);
     }

◆ add_array_argument()

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.

Parameters

[in,out]	idx	Index at which to start adding the array's arguments. Will be incremented by the number of kernel arguments set.
[in]	array	Array to set as an argument of the object's kernel.
[in]	strides	Strides object containing stride of each dimension in bytes.
[in]	num_dimensions	Number of dimensions of the `array`.
[in]	window	Window the kernel will be executed on.

Definition at line 496 of file ICLKernel.h.

References ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_MSG_VAR, ARM_COMPUTE_UNUSED, ICLArray< T >::cl_buffer(), and arm_compute::test::validation::n.

 {
     ARM_COMPUTE_ERROR_ON(array == nullptr);
 
     // Calculate offset to the start of the window
     unsigned int offset_first_element = 0;
 
     for(unsigned int n = 0; n < num_dimensions; ++n)
     {
         offset_first_element += window[n].start() * strides[n];
     }
 
     unsigned int idx_start = idx;
     _kernel.setArg(idx++, array->cl_buffer());
 
     for(unsigned int dimension = 0; dimension < dimension_size; dimension++)
     {
         _kernel.setArg<cl_uint>(idx++, strides[dimension]);
         _kernel.setArg<cl_uint>(idx++, strides[dimension] * window[dimension].step());
     }
 
     _kernel.setArg<cl_uint>(idx++, offset_first_element);
 
     ARM_COMPUTE_ERROR_ON_MSG_VAR(idx_start + num_arguments_per_array<dimension_size>() != idx,
                                  "add_%dD_array_argument() is supposed to add exactly %d arguments to the kernel", dimension_size, num_arguments_per_array<dimension_size>());
     ARM_COMPUTE_UNUSED(idx_start);
 }

◆ add_tensor_argument()

void add_tensor_argument	(	unsigned &	idx,
		const ICLTensor *	tensor,
		const Window &	window
	)

Definition at line 88 of file ICLKernel.cpp.

References ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_MSG_VAR, ARM_COMPUTE_UNUSED, ICLTensor::cl_buffer(), ITensor::info(), arm_compute::test::validation::info, Window::is_broadcasted(), arm_compute::test::validation::n, ITensorInfo::num_dimensions(), ITensorInfo::offset_first_element_in_bytes(), and ITensorInfo::strides_in_bytes().

 {
     ARM_COMPUTE_ERROR_ON(tensor == nullptr);
 
     const ITensorInfo *info    = tensor->info();
     const Strides     &strides = info->strides_in_bytes();
 
     // Calculate offset to the start of the window
     unsigned int offset_first_element = info->offset_first_element_in_bytes();
 
     for(unsigned int n = 0; n < info->num_dimensions(); ++n)
     {
         offset_first_element += (window.is_broadcasted(n) ? 0 : window[n].start()) * strides[n];
     }
 
     unsigned int idx_start = idx;
     _kernel.setArg(idx++, tensor->cl_buffer());
 
     for(unsigned int d = 0; d < dimension_size; ++d)
     {
         _kernel.setArg<cl_uint>(idx++, window.is_broadcasted(d) ? 0 : strides[d]);
         _kernel.setArg<cl_uint>(idx++, window.is_broadcasted(d) ? 0 : (strides[d] * window[d].step()));
     }
 
     _kernel.setArg<cl_uint>(idx++, offset_first_element);
 
     ARM_COMPUTE_ERROR_ON_MSG_VAR(idx_start + num_arguments_per_tensor<dimension_size>() != idx,
                                  "add_%dD_tensor_argument() is supposed to add exactly %d arguments to the kernel", dimension_size, num_arguments_per_tensor<dimension_size>());
     ARM_COMPUTE_UNUSED(idx_start);
 }

◆ config_id()

const std::string& config_id ( ) const

inline

Get the configuration ID.

Note: The configuration ID can be used by the caller to distinguish different calls of the same OpenCL kernel In particular, this method can be used by CLScheduler to keep track of the best LWS for each configuration of the same kernel. The configuration ID should be provided only for the kernels potentially affected by the LWS geometry; This method should be called after the configuration of the kernel

Returns: configuration id string

Definition at line 398 of file ICLKernel.h.

Referenced by CLTuner::tune_kernel_dynamic(), and ClSaturatedArithmeticKernel::validate().

     {
         return _config_id;
     }

◆ get_max_workgroup_size()

size_t get_max_workgroup_size ( )

Get the maximum workgroup size for the device the CLKernelLibrary uses.

Returns: The maximum workgroup size value.

Definition at line 184 of file ICLKernel.cpp.

References CLKernelLibrary::get(), and CLKernelLibrary::max_local_workgroup_size().

Referenced by arm_compute::enqueue(), and ICLKernel::get_target().

 {
     if(_max_workgroup_size == 0)
     {
         _max_workgroup_size = CLKernelLibrary::get().max_local_workgroup_size(_kernel);
     }
     return _max_workgroup_size;
 }

◆ get_target()

GPUTarget get_target ( ) const

inline

Get the targeted GPU architecture.

Returns: The targeted GPU architecture.

Definition at line 422 of file ICLKernel.h.

References arm_compute::enqueue(), CLKernelLibrary::get(), ICLKernel::get_max_workgroup_size(), ICLKernel::gws_from_window(), ICLKernel::kernel(), ICLKernel::lws_hint(), and IKernel::window().

Referenced by ClIndirectConv2dKernel::configure(), ClDirectConv2dKernel::configure(), CLDepthwiseConvolutionLayerNativeKernel::configure(), and CLTuner::tune_kernel_dynamic().

     {
         return _target;
     }

◆ gws_from_window()

cl::NDRange gws_from_window ( const Window & window )

static

Get the global work size given an execution window.

Parameters

[in] window Execution window

Returns: Global work size of the given execution window

Definition at line 193 of file ICLKernel.cpp.

References Window::Dimension::end(), Window::Dimension::start(), Window::Dimension::step(), Window::x(), Window::y(), and Window::z().

Referenced by CLTuner::add_tuning_params(), and ICLKernel::get_target().

 {
     if((window.x().end() - window.x().start()) == 0 || (window.y().end() - window.y().start()) == 0)
     {
         return cl::NullRange;
     }
 
     cl::NDRange gws((window.x().end() - window.x().start()) / window.x().step(),
                     (window.y().end() - window.y().start()) / window.y().step(),
                     (window.z().end() - window.z().start()) / window.z().step());
 
     return gws;
 }

◆ kernel()

cl::Kernel& kernel ( )

inline

Returns a reference to the OpenCL kernel of this object.

Returns: A reference to the OpenCL kernel of this object.

Definition at line 134 of file ICLKernel.h.

Referenced by arm_compute::enqueue(), and ICLKernel::get_target().

     {
         return _kernel;
     }

◆ lws_hint()

cl::NDRange lws_hint ( ) const

inline

Return the Local-Workgroup-Size hint.

Returns: Current lws hint

Definition at line 362 of file ICLKernel.h.

References CLTuningParams::get_lws().

     {
         return _tuning_params_hint.get_lws();
     }

◆ num_arguments_per_1D_array()

static constexpr unsigned int num_arguments_per_1D_array ( )

inlinestatic

Returns the number of arguments enqueued per 1D array object.

Returns: The number of arguments enqueues per 1D array object.

Definition at line 276 of file ICLKernel.h.

     {
         return num_arguments_per_array<1>();
     }

◆ num_arguments_per_1D_tensor()

static constexpr unsigned int num_arguments_per_1D_tensor ( )

inlinestatic

Returns the number of arguments enqueued per 1D tensor object.

Returns: The number of arguments enqueues per 1D tensor object.

Definition at line 284 of file ICLKernel.h.

Referenced by CLReverseKernel::configure(), CLBatchNormalizationLayerKernel::configure(), and CLSelectKernel::run().

     {
         return num_arguments_per_tensor<1>();
     }

◆ num_arguments_per_2D_tensor()

static constexpr unsigned int num_arguments_per_2D_tensor ( )

inlinestatic

Returns the number of arguments enqueued per 2D tensor object.

Returns: The number of arguments enqueues per 2D tensor object.

Definition at line 292 of file ICLKernel.h.

Referenced by CLPriorBoxLayerKernel::configure(), CLL2NormalizeLayerKernel::configure(), CLArgMinMaxLayerKernel::run(), ClGemmLowpMatrixMultiplyNativeKernel::run_op(), ClGemmMatrixMultiplyNativeKernel::run_op(), and ClIm2ColKernel::run_op().

     {
         return num_arguments_per_tensor<2>();
     }

◆ num_arguments_per_3D_tensor()

static constexpr unsigned int num_arguments_per_3D_tensor ( )

inlinestatic

Returns the number of arguments enqueued per 3D tensor object.

Returns: The number of arguments enqueues per 3D tensor object.

Definition at line 300 of file ICLKernel.h.

Referenced by ClReshapeKernel::configure(), CLFFTScaleKernel::configure(), CLFFTRadixStageKernel::configure(), CLL2NormalizeLayerKernel::configure(), ClMulKernel::configure(), CLFillBorderKernel::configure(), CLBatchNormalizationLayerKernel::configure(), CLNormalizePlanarYUVLayerKernel::run(), CLBatchNormalizationLayerKernel::run(), ClDequantizeKernel::run_op(), ClBatchConcatenateKernel::run_op(), ClDepthConcatenateKernel::run_op(), ClWinogradInputTransformKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleByFixedPointKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleByFloatKernel::run_op(), ClDirectConv2dKernel::run_op(), and ClIm2ColKernel::run_op().

     {
         return num_arguments_per_tensor<3>();
     }

◆ num_arguments_per_3d_tensor_nhw()

static constexpr unsigned int num_arguments_per_3d_tensor_nhw ( )

inlinestatic

Returns the number of arguments enqueued per NHW 3D Tensor object.

Returns: The number of arguments enqueued per NHW 3D Tensor object.

Definition at line 249 of file ICLKernel.h.

References ICLKernel::add_4d_tensor_nhwc_argument().

Referenced by ClGemmReshapeLhsMatrixKernel::configure(), and ClGemmReshapeRhsMatrixKernel::configure().

     {
         constexpr unsigned int no_args_per_3d_tensor_nhw = 7u;
         return no_args_per_3d_tensor_nhw;
     }

◆ num_arguments_per_4D_tensor()

static constexpr unsigned int num_arguments_per_4D_tensor ( )

inlinestatic

Returns the number of arguments enqueued per 4D tensor object.

Returns: The number of arguments enqueues per 4D tensor object.

Definition at line 308 of file ICLKernel.h.

Referenced by CLReverseKernel::configure(), and CLStackLayerKernel::configure().

     {
         return num_arguments_per_tensor<4>();
     }

◆ num_arguments_per_4d_tensor_nhwc()

static constexpr unsigned int num_arguments_per_4d_tensor_nhwc ( )

inlinestatic

Returns the number of arguments enqueued per NHWC 4D Tensor object.

Returns: The number of arguments enqueued per NHWC 4D Tensor object.

Definition at line 266 of file ICLKernel.h.

Referenced by ClScaleKernel::configure().

     {
         constexpr unsigned int no_args_per_4d_tensor_nhwc = 9u;
         return no_args_per_4d_tensor_nhwc;
     }

◆ run()

virtual void run	(	const Window &	window,
		cl::CommandQueue &	queue
	)

inlinevirtual

Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue.

Note: The queue is not flushed by this method, and therefore the kernel will not have been executed by the time this method returns.

Parameters

[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 in CLComputeMeanVariance, CLSpaceToBatchLayerKernel, CLFuseBatchNormalizationKernel, CLBatchNormalizationLayerKernel, CLBatchToSpaceLayerKernel, CLDeconvolutionReshapeOutputKernel, CLROIAlignLayerKernel, CLArgMinMaxLayerKernel, CLDepthwiseConvolutionLayerNativeKernel, CLStackLayerKernel, CLBoundingBoxTransformKernel, CLL2NormalizeLayerKernel, CLFFTRadixStageKernel, CLFillBorderKernel, CLPadLayerKernel, CLPriorBoxLayerKernel, CLNormalizePlanarYUVLayerKernel, CLReductionOperationKernel, CLRangeKernel, CLReorgLayerKernel, CLMeanStdDevNormalizationKernel, CLFFTDigitReverseKernel, CLTileKernel, CLComparisonKernel, CLGatherKernel, CLNormalizationLayerKernel, CLROIPoolingLayerKernel, CLQLSTMLayerNormalizationKernel, CLFFTScaleKernel, CLComputeAllAnchorsKernel, CLMaxUnpoolingLayerKernel, CLDeconvolutionLayerUpsampleKernel, CLDepthToSpaceLayerKernel, CLReverseKernel, CLSelectKernel, CLSpaceToDepthLayerKernel, CLChannelShuffleLayerKernel, CLInstanceNormalizationLayerKernel, CLBitwiseKernel, ICLSimple3DKernel, and ICLSimple2DKernel.

Definition at line 319 of file ICLKernel.h.

References ARM_COMPUTE_UNUSED.

Referenced by CLScheduler::init().

     {
         ARM_COMPUTE_UNUSED(window, queue);
     }

◆ run_op()

virtual void run_op	(	ITensorPack &	tensors,
		const Window &	window,
		cl::CommandQueue &	queue
	)

inlinevirtual

Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue.

Note: The queue is not flushed by this method, and therefore the kernel will not have been executed by the time this method returns.

Parameters

[in]	tensors	A vector containing the tensors to operato on.
[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 in ClGemmLowpMatrixBReductionKernel, ClComplexMulKernel, ClGemmMatrixMultiplyReshapedKernel, ClIm2ColKernel, ClGemmMatrixMultiplyReshapedOnlyRhsKernel, CLFillBorderKernel, ClWeightsReshapeKernel, ClGemmLowpMatrixMultiplyReshapedOnlyRhsKernel, ClGemmLowpMatrixAReductionKernel, ClMulKernel, ClDirectConv3dKernel, ClDirectConv2dKernel, ClGemmLowpMatrixMultiplyReshapedOnlyRhsMMULKernel, ClGemmLowpOffsetContributionOutputStageKernel, ClCol2ImKernel, ClGemmLowpOffsetContributionKernel, CLStridedSliceKernel, ClGemmLowpMatrixMultiplyReshapedKernel, ClGemmReshapeRhsMatrixKernel, ClGemmMatrixMultiplyReshapedOnlyRhsMMULKernel, ClGemmMatrixMultiplyNativeKernel, ClGemmLowpQuantizeDownInt32ScaleByFloatKernel, ClGemmLowpQuantizeDownInt32ScaleKernel, ClIndirectConv2dAddressPrecalculationKernel, ClWinogradOutputTransformKernel, ClIndirectConv2dKernel, ClGemmLowpQuantizeDownInt32ScaleByFixedPointKernel, ClWinogradFilterTransformKernel, ClWinogradInputTransformKernel, ClGemmLowpMatrixMultiplyNativeKernel, ClGemmReshapeLhsMatrixKernel, ClCropKernel, ClPermuteKernel, ClQuantizeKernel, ClPool2dKernel, ClPool3dKernel, ClScaleKernel, ClCopyKernel, ClElementWiseUnaryKernel, ClFillKernel, ClTransposedConvolutionKernel, ClDequantizeKernel, ClFloorKernel, ClReshapeKernel, ClTransposeKernel, and ClKernelRuntime.

Definition at line 331 of file ICLKernel.h.

References ARM_COMPUTE_UNUSED.

Referenced by CLScheduler::init().

     {
         ARM_COMPUTE_UNUSED(tensors, window, queue);
     }

◆ set_lws_hint()

void set_lws_hint ( const cl::NDRange & lws_hint )

inline

Set the Local-Workgroup-Size hint.

Note: This method should be called after the configuration of the kernel

Parameters

[in] lws_hint Local-Workgroup-Size to use

Definition at line 352 of file ICLKernel.h.

References ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, and CLTuningParams::set_lws().

Referenced by CLTuner::tune_kernel_dynamic().

     {
         ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); // lws_hint will be overwritten by configure()
         _tuning_params_hint.set_lws(lws_hint);
     }

◆ set_target() [1/2]

void set_target ( GPUTarget target )

inline

Set the targeted GPU architecture.

Parameters

[in] target The targeted GPU architecture

Definition at line 407 of file ICLKernel.h.

     {
         _target = target;
     }

◆ set_target() [2/2]

void set_target ( cl::Device & device )

Set the targeted GPU architecture according to the CL device.

Parameters

[in] device A CL device

Definition at line 179 of file ICLKernel.cpp.

References arm_compute::get_target_from_device().

 {
     _target = get_target_from_device(device);
 }

◆ set_wbsm_hint()

void set_wbsm_hint ( const cl_int & wbsm_hint )

inline

Set the workgroup batch size modifier hint.

Note: This method should be called after the configuration of the kernel

Parameters

[in] wbsm_hint workgroup batch size modifier value

Definition at line 373 of file ICLKernel.h.

References ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, and CLTuningParams::set_wbsm().

Referenced by CLTuner::tune_kernel_dynamic().

     {
         ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); // wbsm_hint will be overwritten by configure()
         _tuning_params_hint.set_wbsm(wbsm_hint);
     }

◆ type()

CLKernelType type ( ) const

inline

Returns the CL kernel type.

Returns: The CL kernel type

Definition at line 142 of file ICLKernel.h.

     {
         return _type;
     }

◆ wbsm_hint()

cl_int wbsm_hint ( ) const

inline

Return the workgroup batch size modifier hint.

Returns: Current wbsm hint

Definition at line 383 of file ICLKernel.h.

References CLTuningParams::get_wbsm().

Referenced by arm_compute::enqueue().

     {
         return _tuning_params_hint.get_wbsm();
     }

The documentation for this class was generated from the following files:

src/core/CL/ICLKernel.h
src/core/CL/ICLKernel.cpp

Public Member Functions

Static Public Member Functions

Detailed Description

Constructor & Destructor Documentation

◆ ICLKernel()

Member Function Documentation

◆ add_1D_array_argument()

◆ add_1D_tensor_argument()

◆ add_1D_tensor_argument_if()

◆ add_2D_tensor_argument()

◆ add_2D_tensor_argument_if()

◆ add_3D_tensor_argument()

◆ add_3d_tensor_nhw_argument()

◆ add_4D_tensor_argument()

◆ add_4d_tensor_nhwc_argument()

◆ add_5D_tensor_argument()

◆ add_argument()

◆ add_array_argument()

◆ add_tensor_argument()

◆ config_id()

◆ get_max_workgroup_size()

◆ get_target()

◆ gws_from_window()

◆ kernel()

◆ lws_hint()

◆ num_arguments_per_1D_array()

◆ num_arguments_per_1D_tensor()

◆ num_arguments_per_2D_tensor()

◆ num_arguments_per_3D_tensor()

◆ num_arguments_per_3d_tensor_nhw()

◆ num_arguments_per_4D_tensor()

◆ num_arguments_per_4d_tensor_nhwc()

◆ run()

◆ run_op()

◆ set_lws_hint()

◆ set_target() [1/2]

◆ set_target() [2/2]

◆ set_wbsm_hint()

◆ type()

◆ wbsm_hint()