Compute Library
 19.08
ICLKernel Class Referenceabstract

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

#include <ICLKernel.h>

Collaboration diagram for ICLKernel:
[legend]

Public Member Functions

 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 (const Window &window, cl::CommandQueue &queue)=0
 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...
 
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<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...
 
template<unsigned int dimension_size>
void add_tensor_argument (unsigned &idx, const ICLTensor *tensor, const Window &window)
 
- 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 Windowwindow () const
 The maximum window the kernel can be executed on. More...
 

Static Public Member Functions

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 43 of file ICLKernel.h.

Constructor & Destructor Documentation

◆ ICLKernel()

ICLKernel ( )
inline

Constructor.

Definition at line 79 of file ICLKernel.h.

80  : _kernel(nullptr), _target(GPUTarget::MIDGARD), _config_id(arm_compute::default_config_id), _max_workgroup_size(0), _lws_hint()
81  {
82  }

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]idxIndex at which to start adding the array's arguments. Will be incremented by the number of kernel arguments set.
[in]arrayArray to set as an argument of the object's kernel.
[in]stridesStrides object containing stride of each dimension in bytes.
[in]num_dimensionsNumber of dimensions of the array.
[in]windowWindow the kernel will be executed on.

Definition at line 100 of file ICLKernel.h.

101  {
102  add_array_argument<T, 1>(idx, array, strides, num_dimensions, window);
103  }
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28

References IKernel::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]idxIndex at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]tensorTensor to set as an argument of the object's kernel.
[in]windowWindow the kernel will be executed on.

Definition at line 110 of file ICLKernel.h.

111  {
112  add_tensor_argument<1>(idx, tensor, window);
113  }
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28

References IKernel::window().

Referenced by ICLKernel::add_1D_tensor_argument_if(), CLElementWiseUnaryLayerKernel::run(), CLGEMMMatrixAccumulateBiasesKernel::run(), CLGEMMMatrixVectorMultiplyKernel::run(), CLReverseKernel::run(), CLComputeAllAnchorsKernel::run(), CLDepthwiseConvolutionLayerReshapeWeightsGenericKernel::run(), CLGatherKernel::run(), CLDepthwiseVectorToTensorKernel::run(), CLMinMaxLayerKernel::run(), CLDepthwiseConvolutionLayer3x3NCHWKernel::run(), CLFFTDigitReverseKernel::run(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), CLRangeKernel::run(), CLNormalizePlanarYUVLayerKernel::run(), CLSelectKernel::run(), CLBoundingBoxTransformKernel::run(), CLReductionOperationKernel::run(), CLDirectConvolutionLayerOutputStageKernel::run(), CLGEMMLowpQuantizeDownInt32ToInt16ScaleByFixedPointKernel::run(), CLDeconvolutionReshapeOutputKernel::run(), CLBatchToSpaceLayerKernel::run(), CLGEMMLowpQuantizeDownInt32ToUint8ScaleByFloatKernel::run(), CLGEMMLowpQuantizeDownInt32ToUint8ScaleByFixedPointKernel::run(), CLGEMMLowpQuantizeDownInt32ToUint8ScaleKernel::run(), CLBatchNormalizationLayerKernel::run(), CLDirectConvolutionLayerKernel::run(), CLFuseBatchNormalizationKernel::run(), CLWinogradOutputTransformKernel::run(), and CLWeightsReshapeKernel::run().

◆ 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]condCondition to check
[in,out]idxIndex at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]tensorTensor to set as an argument of the object's kernel.
[in]windowWindow the kernel will be executed on.

Definition at line 121 of file ICLKernel.h.

122  {
123  if(cond)
124  {
125  add_1D_tensor_argument(idx, tensor, window);
126  }
127  }
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28
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.
Definition: ICLKernel.h:110

References ICLKernel::add_1D_tensor_argument(), and IKernel::window().

Referenced by CLGEMMLowpOffsetContributionOutputStageKernel::run(), CLGEMMLowpOffsetContributionKernel::run(), and CLSpaceToBatchLayerKernel::run().

◆ 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]idxIndex at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]tensorTensor to set as an argument of the object's kernel.
[in]windowWindow the kernel will be executed on.

Definition at line 134 of file ICLKernel.h.

135  {
136  add_tensor_argument<2>(idx, tensor, window);
137  }
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28

References IKernel::window().

Referenced by ICLKernel::add_2D_tensor_argument_if(), ICLSimple2DKernel::run(), CLGaussianPyramidHorKernel::run(), CLGradientKernel::run(), CLHistogramKernel::run(), CLBitwiseAndKernel::run(), CLBitwiseOrKernel::run(), CLBitwiseXorKernel::run(), CLMinMaxKernel::run(), CLRemapKernel::run(), CLDerivativeKernel::run(), CLSobel3x3Kernel::run(), CLHOGOrientationBinningKernel::run(), CLSobel5x5HorKernel::run(), CLSobel7x7HorKernel::run(), CLAbsoluteDifferenceKernel::run(), CLGEMMMatrixAccumulateBiasesKernel::run(), CLGEMMMatrixVectorMultiplyKernel::run(), CLMagnitudePhaseKernel::run(), CLDepthwiseConvolutionLayerReshapeWeightsKernel::run(), CLIntegralImageVertKernel::run(), CLROIPoolingLayerKernel::run(), CLLocallyConnectedMatrixMultiplyKernel::run(), CLChannelExtractKernel::run(), CLDepthwiseConvolutionLayerReshapeWeightsGenericKernel::run(), CLFastCornersKernel::run(), CLMeanStdDevNormalizationKernel::run(), CLHarrisScoreKernel::run(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), CLHOGDetectorKernel::run(), CLChannelCombineKernel::run(), CLPriorBoxLayerKernel::run(), CLConvertFullyConnectedWeightsKernel::run(), CLScaleKernel::run(), CLScharr3x3Kernel::run(), CLMeanStdDevKernel::run(), CLReductionOperationKernel::run(), CLBoundingBoxTransformKernel::run(), CLL2NormalizeLayerKernel::run(), CLGEMMLowpMatrixMultiplyKernel::run(), CLGEMMLowpMatrixMultiplyNativeKernel::run(), CLColorConvertKernel::run(), CLROIAlignLayerKernel::run(), CLGEMMLowpMatrixAReductionKernel::run(), CLGEMMLowpMatrixMultiplyReshapedOnlyRHSKernel::run(), CLHistogramBorderKernel::run(), CLGaussianPyramidVertKernel::run(), CLGEMMMatrixMultiplyNativeKernel::run(), CLGEMMMatrixMultiplyKernel::run(), CLGEMMLowpMatrixMultiplyReshapedKernel::run(), CLGEMMMatrixMultiplyReshapedOnlyRHSKernel::run(), CLEdgeNonMaxSuppressionKernel::run(), CLMinMaxLocationKernel::run(), CLHOGBlockNormalizationKernel::run(), CLWeightsReshapeKernel::run(), CLIm2ColKernel::run(), CLGEMMMatrixMultiplyReshapedKernel::run(), CLCopyToArrayKernel::run(), CLGEMMLowpMatrixBReductionKernel::run(), CLEdgeTraceKernel::run(), CLLKTrackerStage0Kernel::run(), CLConvolutionRectangleKernel::run(), and CLLKTrackerStage1Kernel::run().

◆ 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]condCondition to check
[in,out]idxIndex at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]tensorTensor to set as an argument of the object's kernel.
[in]windowWindow the kernel will be executed on.

Definition at line 145 of file ICLKernel.h.

146  {
147  if(cond)
148  {
149  add_2D_tensor_argument(idx, tensor, window);
150  }
151  }
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28
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.
Definition: ICLKernel.h:134

References ICLKernel::add_2D_tensor_argument(), and IKernel::window().

Referenced by CLDerivativeKernel::run(), CLSobel3x3Kernel::run(), CLSobel5x5HorKernel::run(), CLSobel7x7HorKernel::run(), CLMagnitudePhaseKernel::run(), CLMeanStdDevNormalizationKernel::run(), CLChannelCombineKernel::run(), CLScharr3x3Kernel::run(), CLGEMMLowpOffsetContributionOutputStageKernel::run(), CLGEMMLowpOffsetContributionKernel::run(), CLSpaceToBatchLayerKernel::run(), CLGEMMMatrixMultiplyReshapedOnlyRHSKernel::run(), CLGEMMMatrixMultiplyReshapedKernel::run(), CLSobel7x7VertKernel::run(), and CLSobel5x5VertKernel::run().

◆ 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]idxIndex at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]tensorTensor to set as an argument of the object's kernel.
[in]windowWindow the kernel will be executed on.

Definition at line 158 of file ICLKernel.h.

159  {
160  add_tensor_argument<3>(idx, tensor, window);
161  }
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28

References IKernel::window().

Referenced by ICLSimple3DKernel::run(), CLElementwiseOperationKernel::run(), CLDequantizationLayerKernel::run(), CLGEMMMatrixVectorMultiplyKernel::run(), CLFloorKernel::run(), CLDepthwiseConvolutionLayerReshapeWeightsKernel::run(), CLFlattenLayerKernel::run(), CLReshapeLayerKernel::run(), CLCopyKernel::run(), CLQuantizationLayerKernel::run(), CLDepthToSpaceLayerKernel::run(), CLROIPoolingLayerKernel::run(), CLSpaceToDepthLayerKernel::run(), CLDeconvolutionLayerUpsampleKernel::run(), CLLocallyConnectedMatrixMultiplyKernel::run(), CLUpsampleLayerKernel::run(), CLMemsetKernel::run(), CLNormalizationLayerKernel::run(), CLFFTScaleKernel::run(), CLFillBorderKernel::run(), CLActivationLayerKernel::run(), CLDepthwiseConvolutionLayerReshapeWeightsGenericKernel::run(), CLGEMMMatrixAdditionKernel::run(), CLPoolingLayerKernel::run(), CLComparisonKernel::run(), CLReorgLayerKernel::run(), CLDepthwiseVectorToTensorKernel::run(), CLMinMaxLayerKernel::run(), CLFFTDigitReverseKernel::run(), CLDepthwiseConvolutionLayer3x3NCHWKernel::run(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), CLNormalizePlanarYUVLayerKernel::run(), CLDepthConcatenateLayerKernel::run(), CLBatchConcatenateLayerKernel::run(), CLSelectKernel::run(), CLReductionOperationKernel::run(), CLPixelWiseMultiplicationKernel::run(), CLYOLOLayerKernel::run(), CLCropKernel::run(), CLL2NormalizeLayerKernel::run(), CLGEMMReshapeRHSMatrixKernel::run(), CLFFTRadixStageKernel::run(), CLDirectConvolutionLayerOutputStageKernel::run(), CLGEMMReshapeLHSMatrixKernel::run(), CLDepthwiseIm2ColKernel::run(), CLROIAlignLayerKernel::run(), CLGEMMLowpMatrixAReductionKernel::run(), CLDeconvolutionReshapeOutputKernel::run(), CLGEMMLowpQuantizeDownInt32ToInt16ScaleByFixedPointKernel::run(), CLBatchToSpaceLayerKernel::run(), CLGEMMLowpQuantizeDownInt32ToUint8ScaleByFloatKernel::run(), CLWinogradInputTransformKernel::run(), CLCol2ImKernel::run(), CLGEMMLowpQuantizeDownInt32ToUint8ScaleKernel::run(), CLGEMMLowpQuantizeDownInt32ToUint8ScaleByFixedPointKernel::run(), CLGEMMLowpOffsetContributionOutputStageKernel::run(), CLWinogradFilterTransformKernel::run(), CLDirectConvolutionLayerKernel::run(), CLBatchNormalizationLayerKernel::run(), CLGEMMLowpOffsetContributionKernel::run(), CLFuseBatchNormalizationKernel::run(), CLSpaceToBatchLayerKernel::run(), CLWeightsReshapeKernel::run(), CLIm2ColKernel::run(), CLGEMMLowpMatrixBReductionKernel::run(), CLComplexPixelWiseMultiplicationKernel::run(), CLLogits1DMaxShiftExpSumKernel::run(), and CLLogits1DNormKernel::run().

◆ 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]idxIndex at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
[in]tensorTensor to set as an argument of the object's kernel.
[in]windowWindow the kernel will be executed on.

Definition at line 168 of file ICLKernel.h.

169  {
170  add_tensor_argument<4>(idx, tensor, window);
171  }
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28

References IKernel::window().

Referenced by CLChannelShuffleLayerKernel::run(), CLFlattenLayerKernel::run(), CLDepthToSpaceLayerKernel::run(), CLPermuteKernel::run(), CLReverseKernel::run(), CLSpaceToDepthLayerKernel::run(), CLTileKernel::run(), CLWidthConcatenate2TensorsKernel::run(), CLPoolingLayerKernel::run(), CLGatherKernel::run(), CLHeightConcatenateLayerKernel::run(), CLWidthConcatenateLayerKernel::run(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), CLScaleKernel::run(), CLWidthConcatenate4TensorsKernel::run(), CLReductionOperationKernel::run(), CLStackLayerKernel::run(), CLBatchToSpaceLayerKernel::run(), CLCol2ImKernel::run(), CLWinogradFilterTransformKernel::run(), CLStridedSliceKernel::run(), CLSpaceToBatchLayerKernel::run(), and CLWinogradOutputTransformKernel::run().

◆ add_argument()

void add_argument ( unsigned int &  idx,
value 
)
inline

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

Parameters
[in,out]idxIndex at which to start adding the arguments. Will be incremented by the number of kernel arguments set.
[in]valueValue to set as an argument of the object's kernel.

Definition at line 226 of file ICLKernel.h.

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

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

◆ 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]idxIndex at which to start adding the array's arguments. Will be incremented by the number of kernel arguments set.
[in]arrayArray to set as an argument of the object's kernel.
[in]stridesStrides object containing stride of each dimension in bytes.
[in]num_dimensionsNumber of dimensions of the array.
[in]windowWindow the kernel will be executed on.

Definition at line 359 of file ICLKernel.h.

360 {
361  ARM_COMPUTE_ERROR_ON(array == nullptr);
362 
363  // Calculate offset to the start of the window
364  unsigned int offset_first_element = 0;
365 
366  for(unsigned int n = 0; n < num_dimensions; ++n)
367  {
368  offset_first_element += window[n].start() * strides[n];
369  }
370 
371  unsigned int idx_start = idx;
372  _kernel.setArg(idx++, array->cl_buffer());
373 
374  for(unsigned int dimension = 0; dimension < dimension_size; dimension++)
375  {
376  _kernel.setArg<cl_uint>(idx++, strides[dimension]);
377  _kernel.setArg<cl_uint>(idx++, strides[dimension] * window[dimension].step());
378  }
379 
380  _kernel.setArg<cl_uint>(idx++, offset_first_element);
381 
382  ARM_COMPUTE_ERROR_ON_MSG(idx_start + num_arguments_per_array<dimension_size>() != idx,
383  "add_%dD_array_argument() is supposed to add exactly %d arguments to the kernel", dimension_size, num_arguments_per_array<dimension_size>());
384  ARM_COMPUTE_UNUSED(idx_start);
385 }
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28
#define ARM_COMPUTE_ERROR_ON(cond)
If the condition is true then an error message is printed and an exception thrown.
Definition: Error.h:337
#define ARM_COMPUTE_UNUSED(...)
To avoid unused variables warnings.
Definition: Error.h:160
#define ARM_COMPUTE_ERROR_ON_MSG(cond,...)
Definition: Error.h:328

References ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_MSG, ARM_COMPUTE_UNUSED, ICLArray< T >::cl_buffer(), and IKernel::window().

◆ add_tensor_argument()

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

Definition at line 89 of file ICLKernel.cpp.

90 {
91  ARM_COMPUTE_ERROR_ON(tensor == nullptr);
92 
93  const ITensorInfo *info = tensor->info();
94  const Strides &strides = info->strides_in_bytes();
95 
96  // Calculate offset to the start of the window
97  unsigned int offset_first_element = info->offset_first_element_in_bytes();
98 
99  for(unsigned int n = 0; n < info->num_dimensions(); ++n)
100  {
101  offset_first_element += window[n].start() * strides[n];
102  }
103 
104  unsigned int idx_start = idx;
105  _kernel.setArg(idx++, tensor->cl_buffer());
106 
107  for(unsigned int d = 0; d < dimension_size; ++d)
108  {
109  _kernel.setArg<cl_uint>(idx++, strides[d]);
110  _kernel.setArg<cl_uint>(idx++, strides[d] * window[d].step());
111  }
112 
113  _kernel.setArg<cl_uint>(idx++, offset_first_element);
114 
115  ARM_COMPUTE_ERROR_ON_MSG(idx_start + num_arguments_per_tensor<dimension_size>() != idx,
116  "add_%dD_tensor_argument() is supposed to add exactly %d arguments to the kernel", dimension_size, num_arguments_per_tensor<dimension_size>());
117  ARM_COMPUTE_UNUSED(idx_start);
118 }
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28
#define ARM_COMPUTE_ERROR_ON(cond)
If the condition is true then an error message is printed and an exception thrown.
Definition: Error.h:337
Store the tensor's metadata.
Definition: ITensorInfo.h:40
#define ARM_COMPUTE_UNUSED(...)
To avoid unused variables warnings.
Definition: Error.h:160
virtual ITensorInfo * info() const =0
Interface to be implemented by the child class to return the tensor's metadata.
Strides of an item in bytes.
Definition: Strides.h:37
virtual const cl::Buffer & cl_buffer() const =0
Interface to be implemented by the child class to return a reference to the OpenCL buffer containing ...
#define ARM_COMPUTE_ERROR_ON_MSG(cond,...)
Definition: Error.h:328

References ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_MSG, ARM_COMPUTE_UNUSED, ICLTensor::cl_buffer(), ITensor::info(), arm_compute::test::validation::info, and IKernel::window().

◆ 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 262 of file ICLKernel.h.

263  {
264  return _config_id;
265  }

Referenced by CLTuner::tune_kernel_dynamic().

◆ 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 132 of file ICLKernel.cpp.

133 {
134  if(_max_workgroup_size == 0)
135  {
136  _max_workgroup_size = CLKernelLibrary::get().max_local_workgroup_size(_kernel);
137  }
138  return _max_workgroup_size;
139 }
static CLKernelLibrary & get()
Access the KernelLibrary singleton.
size_t max_local_workgroup_size(const cl::Kernel &kernel) const
Find the maximum number of local work items in a workgroup can be supported for the kernel.

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

Referenced by arm_compute::enqueue().

◆ get_target()

GPUTarget get_target ( ) const
inline

Get the targeted GPU architecture.

Returns
The targeted GPU architecture.

Definition at line 286 of file ICLKernel.h.

287  {
288  return _target;
289  }

Referenced by CLDepthwiseConvolutionLayer3x3NCHWKernel::configure(), CLGEMMMatrixAccumulateBiasesKernel::configure(), CLGEMMMatrixMultiplyKernel::configure(), CLDirectConvolutionLayerKernel::configure(), and CLTuner::tune_kernel_dynamic().

◆ gws_from_window()

cl::NDRange gws_from_window ( const Window window)
static

Get the global work size given an execution window.

Parameters
[in]windowExecution window
Returns
Global work size of the given execution window

Definition at line 141 of file ICLKernel.cpp.

142 {
143  if((window.x().end() - window.x().start()) == 0 || (window.y().end() - window.y().start()) == 0)
144  {
145  return cl::NullRange;
146  }
147 
148  cl::NDRange gws((window.x().end() - window.x().start()) / window.x().step(),
149  (window.y().end() - window.y().start()) / window.y().step(),
150  (window.z().end() - window.z().start()) / window.z().step());
151 
152  return gws;
153 }
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28
constexpr int step() const
Return the step of the dimension.
Definition: Window.h:102
constexpr const Dimension & z() const
Alias to access the third dimension of the window.
Definition: Window.h:161
constexpr const Dimension & y() const
Alias to access the second dimension of the window.
Definition: Window.h:152
constexpr int end() const
Return the end of the dimension.
Definition: Window.h:97
constexpr int start() const
Return the start of the dimension.
Definition: Window.h:92
constexpr const Dimension & x() const
Alias to access the first dimension of the window.
Definition: Window.h:143

References Window::Dimension::end(), Window::Dimension::start(), IKernel::window(), Window::x(), and Window::y().

Referenced by arm_compute::enqueue().

◆ 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 87 of file ICLKernel.h.

88  {
89  return _kernel;
90  }

Referenced by arm_compute::enqueue().

◆ lws_hint()

cl::NDRange lws_hint ( ) const
inline

Return the Local-Workgroup-Size hint.

Returns
Current lws hint

Definition at line 247 of file ICLKernel.h.

248  {
249  return _lws_hint;
250  }

Referenced by CLLocallyConnectedMatrixMultiplyKernel::configure(), CLReductionOperationKernel::configure(), CLLogits1DMaxShiftExpSumKernel::configure(), ICLSimple2DKernel::run(), ICLSimple3DKernel::run(), CLElementwiseOperationKernel::run(), CLGaussianPyramidHorKernel::run(), CLGradientKernel::run(), CLSobel3x3Kernel::run(), CLDerivativeKernel::run(), CLSobel5x5HorKernel::run(), CLSobel7x7HorKernel::run(), CLHOGOrientationBinningKernel::run(), CLGEMMMatrixAccumulateBiasesKernel::run(), CLGEMMMatrixVectorMultiplyKernel::run(), CLMagnitudePhaseKernel::run(), CLChannelShuffleLayerKernel::run(), CLIntegralImageVertKernel::run(), CLFlattenLayerKernel::run(), CLReverseKernel::run(), CLLocallyConnectedMatrixMultiplyKernel::run(), CLFFTScaleKernel::run(), CLNormalizationLayerKernel::run(), CLWidthConcatenate2TensorsKernel::run(), CLActivationLayerKernel::run(), CLFastCornersKernel::run(), CLComparisonKernel::run(), CLMeanStdDevNormalizationKernel::run(), CLPoolingLayerKernel::run(), CLDepthwiseConvolutionLayer3x3NCHWKernel::run(), CLHarrisScoreKernel::run(), CLFFTDigitReverseKernel::run(), CLReorgLayerKernel::run(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), CLHOGDetectorKernel::run(), CLScaleKernel::run(), CLNormalizePlanarYUVLayerKernel::run(), CLRangeKernel::run(), CLSelectKernel::run(), CLWidthConcatenate4TensorsKernel::run(), CLBatchConcatenateLayerKernel::run(), CLReductionOperationKernel::run(), CLYOLOLayerKernel::run(), CLFFTRadixStageKernel::run(), CLGEMMLowpMatrixMultiplyKernel::run(), CLGEMMLowpMatrixMultiplyNativeKernel::run(), CLDirectConvolutionLayerOutputStageKernel::run(), CLGEMMReshapeLHSMatrixKernel::run(), CLColorConvertKernel::run(), CLDepthwiseIm2ColKernel::run(), CLGEMMLowpMatrixAReductionKernel::run(), CLGEMMLowpMatrixMultiplyReshapedOnlyRHSKernel::run(), CLDeconvolutionReshapeOutputKernel::run(), CLWinogradInputTransformKernel::run(), CLCol2ImKernel::run(), CLGEMMLowpOffsetContributionOutputStageKernel::run(), CLDirectConvolutionLayerKernel::run(), CLBatchNormalizationLayerKernel::run(), CLStridedSliceKernel::run(), CLGaussianPyramidVertKernel::run(), CLGEMMLowpOffsetContributionKernel::run(), CLGEMMLowpMatrixMultiplyReshapedKernel::run(), CLGEMMMatrixMultiplyNativeKernel::run(), CLGEMMMatrixMultiplyKernel::run(), CLWinogradOutputTransformKernel::run(), CLGEMMMatrixMultiplyReshapedOnlyRHSKernel::run(), CLEdgeNonMaxSuppressionKernel::run(), CLHOGBlockNormalizationKernel::run(), CLIm2ColKernel::run(), CLGEMMMatrixMultiplyReshapedKernel::run(), CLSobel5x5VertKernel::run(), CLSobel7x7VertKernel::run(), CLCopyToArrayKernel::run(), CLEdgeTraceKernel::run(), CLLogits1DMaxShiftExpSumKernel::run(), CLConvolutionRectangleKernel::run(), CLLogits1DNormKernel::run(), and ICLKernel::set_lws_hint().

◆ 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 176 of file ICLKernel.h.

177  {
178  return num_arguments_per_array<1>();
179  }

Referenced by CLROIPoolingLayerKernel::configure().

◆ 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 184 of file ICLKernel.h.

185  {
186  return num_arguments_per_tensor<1>();
187  }

Referenced by CLGEMMMatrixVectorMultiplyKernel::configure(), CLReverseKernel::configure(), CLBatchNormalizationLayerKernel::configure(), CLDirectConvolutionLayerOutputStageKernel::configure(), CLDirectConvolutionLayerKernel::configure(), CLSelectKernel::run(), CLWinogradOutputTransformKernel::run(), and CLWeightsReshapeKernel::run().

◆ 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 192 of file ICLKernel.h.

193  {
194  return num_arguments_per_tensor<2>();
195  }

Referenced by CLTableLookupKernel::configure(), CLWarpPerspectiveKernel::configure(), CLWarpAffineKernel::configure(), CLScaleKernel::configure(), CLThresholdKernel::configure(), CLGEMMMatrixVectorMultiplyKernel::configure(), CLHistogramKernel::configure(), CLRemapKernel::configure(), CLMinMaxKernel::configure(), CLPriorBoxLayerKernel::configure(), CLMeanStdDevKernel::configure(), CLL2NormalizeLayerKernel::configure(), CLHarrisScoreKernel::configure(), CLFastCornersKernel::configure(), CLHOGDetectorKernel::configure(), CLAccumulateWeightedKernel::configure(), CLHistogramBorderKernel::configure(), CLAccumulateSquaredKernel::configure(), CLEdgeNonMaxSuppressionKernel::configure(), CLMinMaxLocationKernel::configure(), CLCopyToArrayKernel::configure(), CLEdgeTraceKernel::configure(), CLLKTrackerStage0Kernel::configure(), CLLKTrackerStage1Kernel::configure(), CLGEMMMatrixVectorMultiplyKernel::run(), CLDepthwiseConvolutionLayerReshapeWeightsGenericKernel::run(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), CLReductionOperationKernel::run(), CLGEMMLowpMatrixMultiplyKernel::run(), CLGEMMLowpMatrixMultiplyNativeKernel::run(), CLGEMMLowpMatrixMultiplyReshapedOnlyRHSKernel::run(), CLGEMMMatrixMultiplyKernel::run(), CLGEMMLowpMatrixMultiplyReshapedKernel::run(), CLGEMMMatrixMultiplyNativeKernel::run(), CLGEMMMatrixMultiplyReshapedOnlyRHSKernel::run(), CLWeightsReshapeKernel::run(), CLIm2ColKernel::run(), and CLGEMMMatrixMultiplyReshapedKernel::run().

◆ 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 200 of file ICLKernel.h.

201  {
202  return num_arguments_per_tensor<3>();
203  }

Referenced by CLGEMMMatrixVectorMultiplyKernel::configure(), CLReshapeLayerKernel::configure(), CLDepthConvertLayerKernel::configure(), CLFFTScaleKernel::configure(), CLPixelWiseMultiplicationKernel::configure(), CLFillBorderKernel::configure(), CLL2NormalizeLayerKernel::configure(), CLFFTRadixStageKernel::configure(), CLROIPoolingLayerKernel::configure(), CLBatchNormalizationLayerKernel::configure(), CLDirectConvolutionLayerOutputStageKernel::configure(), CLDirectConvolutionLayerKernel::configure(), CLGEMMMatrixVectorMultiplyKernel::run(), CLDepthwiseConvolutionLayerReshapeWeightsGenericKernel::run(), CLDepthwiseConvolutionLayer3x3NCHWKernel::run(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), CLNormalizePlanarYUVLayerKernel::run(), CLDepthConcatenateLayerKernel::run(), CLBatchConcatenateLayerKernel::run(), CLDirectConvolutionLayerOutputStageKernel::run(), CLGEMMReshapeLHSMatrixKernel::run(), CLGEMMLowpQuantizeDownInt32ToInt16ScaleByFixedPointKernel::run(), CLGEMMLowpQuantizeDownInt32ToUint8ScaleByFloatKernel::run(), CLWinogradInputTransformKernel::run(), CLGEMMLowpQuantizeDownInt32ToUint8ScaleKernel::run(), CLGEMMLowpQuantizeDownInt32ToUint8ScaleByFixedPointKernel::run(), CLBatchNormalizationLayerKernel::run(), CLDirectConvolutionLayerKernel::run(), CLWeightsReshapeKernel::run(), and CLIm2ColKernel::run().

◆ 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 208 of file ICLKernel.h.

209  {
210  return num_arguments_per_tensor<4>();
211  }

Referenced by CLScaleKernel::configure(), CLReverseKernel::configure(), CLWidthConcatenate2TensorsKernel::configure(), CLWidthConcatenate4TensorsKernel::configure(), CLStackLayerKernel::configure(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), and CLWinogradOutputTransformKernel::run().

◆ run()

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

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]windowRegion on which to execute the kernel. (Must be a valid region of the window returned by window()).
[in,out]queueCommand queue on which to enqueue the kernel.

Implemented in CLLogits1DNormKernel, CLLKTrackerStage1Kernel, CLConvolutionRectangleKernel, CLLogits1DMaxShiftExpSumKernel, CLLKTrackerStage0Kernel, CLEdgeTraceKernel, CLComplexPixelWiseMultiplicationKernel, CLGEMMLowpMatrixBReductionKernel, CLLKTrackerFinalizeKernel, CLCopyToArrayKernel, CLSobel5x5VertKernel, CLSobel7x7VertKernel, CLGEMMMatrixMultiplyReshapedKernel, CLIm2ColKernel, CLWeightsReshapeKernel, CLHOGBlockNormalizationKernel, CLMinMaxLocationKernel, CLEdgeNonMaxSuppressionKernel, CLGEMMMatrixMultiplyReshapedOnlyRHSKernel, CLWinogradOutputTransformKernel, CLLKTrackerInitKernel, CLSpaceToBatchLayerKernel, CLFuseBatchNormalizationKernel, CLGEMMLowpMatrixMultiplyReshapedKernel, CLGEMMMatrixMultiplyKernel, CLGEMMMatrixMultiplyNativeKernel, CLGaussianPyramidVertKernel, CLGEMMLowpOffsetContributionKernel, CLHistogramBorderKernel, CLLogits1DShiftExpSumKernel, CLStridedSliceKernel, CLBatchNormalizationLayerKernel, CLDirectConvolutionLayerKernel, CLGEMMLowpOffsetContributionOutputStageKernel, CLWinogradFilterTransformKernel, CLCol2ImKernel, CLGEMMLowpQuantizeDownInt32ToUint8ScaleByFixedPointKernel, CLGEMMLowpQuantizeDownInt32ToUint8ScaleKernel, CLGEMMLowpQuantizeDownInt32ToUint8ScaleByFloatKernel, CLWinogradInputTransformKernel, CLBatchToSpaceLayerKernel, CLDeconvolutionReshapeOutputKernel, CLGEMMLowpQuantizeDownInt32ToInt16ScaleByFixedPointKernel, CLGEMMLowpMatrixMultiplyReshapedOnlyRHSKernel, CLGEMMLowpMatrixAReductionKernel, CLDepthwiseIm2ColKernel, CLROIAlignLayerKernel, CLColorConvertKernel, CLGEMMReshapeLHSMatrixKernel, CLDirectConvolutionLayerOutputStageKernel, CLGEMMLowpMatrixMultiplyNativeKernel, CLStackLayerKernel, CLFFTRadixStageKernel, CLGEMMLowpMatrixMultiplyKernel, CLGEMMReshapeRHSMatrixKernel, CLCropKernel, CLL2NormalizeLayerKernel, CLYOLOLayerKernel, CLPixelWiseMultiplicationKernel, CLBoundingBoxTransformKernel, CLReductionOperationKernel, CLBatchConcatenateLayerKernel, CLDepthConcatenateLayerKernel, CLMeanStdDevKernel, CLScharr3x3Kernel, CLSelectKernel, CLWidthConcatenate4TensorsKernel, CLConvertFullyConnectedWeightsKernel, CLNormalizePlanarYUVLayerKernel, CLRangeKernel, CLScaleKernel, CLChannelCombineKernel, CLDepthwiseConvolutionLayer3x3NHWCKernel, CLHOGDetectorKernel, CLPriorBoxLayerKernel, CLDepthwiseConvolutionLayer3x3NCHWKernel, CLDepthwiseVectorToTensorKernel, CLFFTDigitReverseKernel, CLHarrisScoreKernel, CLHeightConcatenateLayerKernel, CLMinMaxLayerKernel, CLReorgLayerKernel, CLWidthConcatenateLayerKernel, CLActivationLayerKernel, CLComparisonKernel, CLDepthwiseConvolutionLayerReshapeWeightsGenericKernel, CLFastCornersKernel, CLGatherKernel, CLGEMMMatrixAdditionKernel, CLMeanStdDevNormalizationKernel, CLPoolingLayerKernel, CLTileKernel, CLWidthConcatenate2TensorsKernel, CLChannelExtractKernel, CLFFTScaleKernel, CLFillBorderKernel, CLMemsetKernel, CLNormalizationLayerKernel, CLUpsampleLayerKernel, CLDeconvolutionLayerUpsampleKernel, CLComputeAllAnchorsKernel, CLLocallyConnectedMatrixMultiplyKernel, CLCopyKernel, CLDepthToSpaceLayerKernel, CLPermuteKernel, CLQuantizationLayerKernel, CLReverseKernel, CLROIPoolingLayerKernel, CLSpaceToDepthLayerKernel, CLChannelShuffleLayerKernel, CLDepthwiseConvolutionLayerReshapeWeightsKernel, CLFlattenLayerKernel, CLIntegralImageVertKernel, CLReshapeLayerKernel, CLDequantizationLayerKernel, CLFloorKernel, CLGEMMMatrixVectorMultiplyKernel, CLMagnitudePhaseKernel, CLGEMMMatrixAccumulateBiasesKernel, CLAbsoluteDifferenceKernel, CLHOGOrientationBinningKernel, CLSobel5x5HorKernel, CLSobel7x7HorKernel, CLDerivativeKernel, CLSobel3x3Kernel, CLBitwiseAndKernel, CLBitwiseOrKernel, CLBitwiseXorKernel, CLMinMaxKernel, CLRemapKernel, CLGradientKernel, CLHistogramKernel, CLGaussianPyramidHorKernel, CLElementwiseOperationKernel, CLElementWiseUnaryLayerKernel, ICLSimple3DKernel, and ICLSimple2DKernel.

Referenced by CLScheduler::enqueue().

◆ 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_hintLocal-Workgroup-Size to use

Definition at line 237 of file ICLKernel.h.

238  {
239  ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); // lws_hint will be overwritten by configure()
240  _lws_hint = lws_hint;
241  }
cl::NDRange lws_hint() const
Return the Local-Workgroup-Size hint.
Definition: ICLKernel.h:247
#define ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(k)
Definition: Validate.h:940

References ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, and ICLKernel::lws_hint().

Referenced by CLGEMMMatrixMultiplyKernel::configure(), and CLTuner::tune_kernel_dynamic().

◆ set_target() [1/2]

void set_target ( GPUTarget  target)
inline

Set the targeted GPU architecture.

Parameters
[in]targetThe targeted GPU architecture

Definition at line 271 of file ICLKernel.h.

272  {
273  _target = target;
274  }

Referenced by CLDirectConvolutionLayer::configure(), CLSoftmaxLayer::configure(), CLGEMMLowpMatrixMultiplyCore::configure(), CLFullyConnectedLayer::configure(), CLGEMMConvolutionLayer::configure(), and CLDepthwiseConvolutionLayer::configure().

◆ set_target() [2/2]

void set_target ( cl::Device &  device)

Set the targeted GPU architecture according to the CL device.

Parameters
[in]deviceA CL device

Definition at line 127 of file ICLKernel.cpp.

128 {
129  _target = get_target_from_device(device);
130 }
GPUTarget get_target_from_device(const cl::Device &device)
Helper function to get the GPU target from CL device.
Definition: CLHelpers.cpp:131

References arm_compute::get_target_from_device().


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