Compute Library
 21.02
ICLKernel Class Reference

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

Constructor & Destructor Documentation

◆ ICLKernel()

ICLKernel ( )
inline

Constructor.

Definition at line 93 of file ICLKernel.h.

94  : _kernel(nullptr), _target(GPUTarget::MIDGARD), _config_id(arm_compute::default_config_id), _max_workgroup_size(0), _tuning_params_hint()
95  {
96  }

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

References IKernel::window().

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

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

References IKernel::window().

Referenced by ICLKernel::add_1D_tensor_argument_if(), CLReverseKernel::run(), CLSelectKernel::run(), CLComputeAllAnchorsKernel::run(), CLQLSTMLayerNormalizationKernel::run(), CLMinMaxLayerKernel::run(), CLGatherKernel::run(), CLFFTDigitReverseKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleByFixedPointKernel::run(), CLRangeKernel::run(), CLNormalizePlanarYUVLayerKernel::run(), CLReductionOperationKernel::run(), CLBoundingBoxTransformKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleByFloatKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleKernel::run(), CLDeconvolutionReshapeOutputKernel::run(), CLBatchToSpaceLayerKernel::run(), CLDepthwiseConvolutionLayer3x3NCHWKernel::run(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), CLBatchNormalizationLayerKernel::run(), CLFuseBatchNormalizationKernel::run(), CLDirectConvolutionLayerKernel::run(), CLWeightsReshapeKernel::run(), CLWinogradOutputTransformKernel::run(), and CLDepthwiseConvolutionLayerNativeKernel::run().

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

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

References ICLKernel::add_1D_tensor_argument().

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

136  {
137  if(cond)
138  {
139  add_1D_tensor_argument(idx, tensor, window);
140  }
141  }
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&#39;s parameters to the object&#39;s kernel&#39;s arguments starting from the index idx...
Definition: ICLKernel.h:124

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

References IKernel::window().

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

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

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

References ICLKernel::add_2D_tensor_argument().

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

160  {
161  if(cond)
162  {
163  add_2D_tensor_argument(idx, tensor, window);
164  }
165  }
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&#39;s parameters to the object&#39;s kernel&#39;s arguments starting from the index idx...
Definition: ICLKernel.h:148

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

References IKernel::window().

Referenced by ICLSimple3DKernel::run(), CLDequantizationLayerKernel::run(), CLDepthwiseConvolutionLayerReshapeWeightsKernel::run(), CLDepthToSpaceLayerKernel::run(), CLSelectKernel::run(), CLSpaceToDepthLayerKernel::run(), CLDeconvolutionLayerUpsampleKernel::run(), CLFFTScaleKernel::run(), CLMaxUnpoolingLayerKernel::run(), CLQuantizationLayerKernel::run(), CLMinMaxLayerKernel::run(), CLNormalizationLayerKernel::run(), CLComparisonKernel::run(), CLFFTDigitReverseKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleByFixedPointKernel::run(), CLReorgLayerKernel::run(), CLROIPoolingLayerKernel::run(), CLNormalizePlanarYUVLayerKernel::run(), CLPadLayerKernel::run(), CLReductionOperationKernel::run(), CLFFTRadixStageKernel::run(), CLFillBorderKernel::run(), CLL2NormalizeLayerKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleByFloatKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleKernel::run(), CLLogits1DMaxShiftExpSumKernel::run(), CLArgMinMaxLayerKernel::run(), CLGEMMReshapeLHSMatrixKernel::run(), CLCol2ImKernel::run(), CLDeconvolutionReshapeOutputKernel::run(), CLROIAlignLayerKernel::run(), CLBatchToSpaceLayerKernel::run(), CLDepthwiseConvolutionLayer3x3NCHWKernel::run(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), CLGEMMLowpOffsetContributionKernel::run(), CLWinogradInputTransformKernel::run(), CLBatchNormalizationLayerKernel::run(), CLWinogradFilterTransformKernel::run(), CLFuseBatchNormalizationKernel::run(), CLDirectConvolutionLayerKernel::run(), CLSpaceToBatchLayerKernel::run(), CLWeightsReshapeKernel::run(), CLDepthwiseConvolutionLayerNativeKernel::run(), CLGEMMLowpOffsetContributionOutputStageKernel::run(), CLIm2ColKernel::run(), CLGEMMLowpMatrixAReductionKernel::run(), CLGEMMReshapeRHSMatrixKernel::run(), CLLogits1DNormKernel::run(), CLGEMMLowpMatrixBReductionKernel::run(), ClElementwiseKernel::run_op(), ClFloorKernel::run_op(), ClReshapeKernel::run_op(), ClCopyKernel::run_op(), ClElementWiseUnaryKernel::run_op(), ClActivationKernel::run_op(), ClPoolingKernel::run_op(), ClFillKernel::run_op(), ClDepthConcatenateKernel::run_op(), ClBatchConcatenateKernel::run_op(), CLFillBorderKernel::run_op(), ClCropKernel::run_op(), CLPixelWiseMultiplicationKernel::run_op(), and CLComplexPixelWiseMultiplicationKernel::run_op().

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

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

References IKernel::window().

Referenced by CLChannelShuffleLayerKernel::run(), CLDepthToSpaceLayerKernel::run(), CLReverseKernel::run(), CLSpaceToDepthLayerKernel::run(), CLScaleKernel::run(), CLGatherKernel::run(), CLTileKernel::run(), CLInstanceNormalizationLayerKernel::run(), CLReductionOperationKernel::run(), CLStackLayerKernel::run(), CLArgMinMaxLayerKernel::run(), CLCol2ImKernel::run(), CLBatchToSpaceLayerKernel::run(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), CLWinogradFilterTransformKernel::run(), CLSpaceToBatchLayerKernel::run(), CLWinogradOutputTransformKernel::run(), CLDepthwiseConvolutionLayerNativeKernel::run(), ClWidthConcatenate2TensorsKernel::run_op(), ClHeightConcatenateKernel::run_op(), ClWidthConcatenateKernel::run_op(), ClPoolingKernel::run_op(), ClWidthConcatenate4TensorsKernel::run_op(), CLStridedSliceKernel::run_op(), and ClPermuteKernel::run_op().

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

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

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

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

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

References ARM_COMPUTE_ERROR_ON, ARM_COMPUTE_ERROR_ON_MSG_VAR, ARM_COMPUTE_UNUSED, and ICLArray< T >::cl_buffer().

410 {
411  ARM_COMPUTE_ERROR_ON(array == nullptr);
412 
413  // Calculate offset to the start of the window
414  unsigned int offset_first_element = 0;
415 
416  for(unsigned int n = 0; n < num_dimensions; ++n)
417  {
418  offset_first_element += window[n].start() * strides[n];
419  }
420 
421  unsigned int idx_start = idx;
422  _kernel.setArg(idx++, array->cl_buffer());
423 
424  for(unsigned int dimension = 0; dimension < dimension_size; dimension++)
425  {
426  _kernel.setArg<cl_uint>(idx++, strides[dimension]);
427  _kernel.setArg<cl_uint>(idx++, strides[dimension] * window[dimension].step());
428  }
429 
430  _kernel.setArg<cl_uint>(idx++, offset_first_element);
431 
432  ARM_COMPUTE_ERROR_ON_MSG_VAR(idx_start + num_arguments_per_array<dimension_size>() != idx,
433  "add_%dD_array_argument() is supposed to add exactly %d arguments to the kernel", dimension_size, num_arguments_per_array<dimension_size>());
434  ARM_COMPUTE_UNUSED(idx_start);
435 }
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:466
#define ARM_COMPUTE_ERROR_ON_MSG_VAR(cond, msg,...)
Definition: Error.h:457
#define ARM_COMPUTE_UNUSED(...)
To avoid unused variables warnings.
Definition: Error.h:152

◆ 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(), ITensorInfo::num_dimensions(), ITensorInfo::offset_first_element_in_bytes(), ITensorInfo::strides_in_bytes(), and IKernel::window().

89 {
90  ARM_COMPUTE_ERROR_ON(tensor == nullptr);
91 
92  const ITensorInfo *info = tensor->info();
93  const Strides &strides = info->strides_in_bytes();
94 
95  // Calculate offset to the start of the window
96  unsigned int offset_first_element = info->offset_first_element_in_bytes();
97 
98  for(unsigned int n = 0; n < info->num_dimensions(); ++n)
99  {
100  offset_first_element += (window.is_broadcasted(n) ? 0 : window[n].start()) * strides[n];
101  }
102 
103  unsigned int idx_start = idx;
104  _kernel.setArg(idx++, tensor->cl_buffer());
105 
106  for(unsigned int d = 0; d < dimension_size; ++d)
107  {
108  _kernel.setArg<cl_uint>(idx++, strides[d]);
109  _kernel.setArg<cl_uint>(idx++, strides[d] * window[d].step());
110  }
111 
112  _kernel.setArg<cl_uint>(idx++, offset_first_element);
113 
114  ARM_COMPUTE_ERROR_ON_MSG_VAR(idx_start + num_arguments_per_tensor<dimension_size>() != idx,
115  "add_%dD_tensor_argument() is supposed to add exactly %d arguments to the kernel", dimension_size, num_arguments_per_tensor<dimension_size>());
116  ARM_COMPUTE_UNUSED(idx_start);
117 }
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:466
#define ARM_COMPUTE_ERROR_ON_MSG_VAR(cond, msg,...)
Definition: Error.h:457
#define ARM_COMPUTE_UNUSED(...)
To avoid unused variables warnings.
Definition: Error.h:152
bool is_broadcasted(size_t dimension) const
Return whether a dimension has been broadcasted.
Definition: Window.inl:62
ScaleKernelInfo info(interpolation_policy, default_border_mode, PixelValue(), sampling_policy, false)

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

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

313  {
314  return _config_id;
315  }

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

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

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

132 {
133  if(_max_workgroup_size == 0)
134  {
135  _max_workgroup_size = CLKernelLibrary::get().max_local_workgroup_size(_kernel);
136  }
137  return _max_workgroup_size;
138 }
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...

◆ get_target()

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

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

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

141 {
142  if((window.x().end() - window.x().start()) == 0 || (window.y().end() - window.y().start()) == 0)
143  {
144  return cl::NullRange;
145  }
146 
147  cl::NDRange gws((window.x().end() - window.x().start()) / window.x().step(),
148  (window.y().end() - window.y().start()) / window.y().step(),
149  (window.z().end() - window.z().start()) / window.z().step());
150 
151  return gws;
152 }
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:104
constexpr const Dimension & z() const
Alias to access the third dimension of the window.
Definition: Window.h:163
constexpr const Dimension & y() const
Alias to access the second dimension of the window.
Definition: Window.h:154
constexpr int end() const
Return the end of the dimension.
Definition: Window.h:99
constexpr int start() const
Return the start of the dimension.
Definition: Window.h:94
constexpr const Dimension & x() const
Alias to access the first dimension of the window.
Definition: Window.h:145

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

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

102  {
103  return _kernel;
104  }

◆ lws_hint()

cl::NDRange lws_hint ( ) const
inline

Return the Local-Workgroup-Size hint.

Returns
Current lws hint

Definition at line 276 of file ICLKernel.h.

References CLTuningParams::get_lws().

Referenced by CLLogits1DMaxShiftExpSumKernel::configure(), CLReductionOperationKernel::configure(), CLArgMinMaxLayerKernel::configure(), ICLKernel::get_target(), ICLSimple2DKernel::run(), ICLSimple3DKernel::run(), CLBitwiseKernel::run(), CLGaussianPyramidHorKernel::run(), CLMinMaxKernel::run(), CLGradientKernel::run(), CLAbsoluteDifferenceKernel::run(), CLRemapKernel::run(), CLHOGOrientationBinningKernel::run(), CLSobel3x3Kernel::run(), CLDequantizationLayerKernel::run(), CLDerivativeKernel::run(), CLLKTrackerInitKernel::run(), CLSobel7x7HorKernel::run(), CLSobel5x5HorKernel::run(), CLChannelShuffleLayerKernel::run(), CLDepthwiseConvolutionLayerReshapeWeightsKernel::run(), CLReverseKernel::run(), CLSelectKernel::run(), CLDepthToSpaceLayerKernel::run(), CLSpaceToDepthLayerKernel::run(), CLDeconvolutionLayerUpsampleKernel::run(), CLMaxUnpoolingLayerKernel::run(), CLFFTScaleKernel::run(), CLComputeAllAnchorsKernel::run(), CLMagnitudePhaseKernel::run(), CLQLSTMLayerNormalizationKernel::run(), CLQuantizationLayerKernel::run(), CLScaleKernel::run(), CLIntegralImageVertKernel::run(), CLMinMaxLayerKernel::run(), CLComparisonKernel::run(), CLNormalizationLayerKernel::run(), CLGatherKernel::run(), CLFastCornersKernel::run(), CLFFTDigitReverseKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleByFixedPointKernel::run(), CLTileKernel::run(), CLMeanStdDevNormalizationKernel::run(), CLInstanceNormalizationLayerKernel::run(), CLConvertFullyConnectedWeightsKernel::run(), CLRangeKernel::run(), CLReorgLayerKernel::run(), CLROIPoolingLayerKernel::run(), CLMeanStdDevKernel::run(), CLNormalizePlanarYUVLayerKernel::run(), CLPadLayerKernel::run(), CLPriorBoxLayerKernel::run(), CLScharr3x3Kernel::run(), CLChannelExtractKernel::run(), CLHarrisScoreKernel::run(), CLReductionOperationKernel::run(), CLHOGDetectorKernel::run(), CLFFTRadixStageKernel::run(), CLFillBorderKernel::run(), CLL2NormalizeLayerKernel::run(), CLBoundingBoxTransformKernel::run(), CLChannelCombineKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleKernel::run(), CLLogits1DMaxShiftExpSumKernel::run(), CLStackLayerKernel::run(), CLLKTrackerFinalizeKernel::run(), CLGEMMLowpMatrixMultiplyNativeKernel::run(), CLArgMinMaxLayerKernel::run(), CLGEMMReshapeLHSMatrixKernel::run(), CLCol2ImKernel::run(), CLDeconvolutionReshapeOutputKernel::run(), CLROIAlignLayerKernel::run(), CLBatchToSpaceLayerKernel::run(), CLGaussianPyramidVertKernel::run(), CLDepthwiseConvolutionLayer3x3NCHWKernel::run(), CLGEMMLowpOffsetContributionKernel::run(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), CLWinogradInputTransformKernel::run(), CLBatchNormalizationLayerKernel::run(), CLWinogradFilterTransformKernel::run(), CLGEMMMatrixMultiplyKernel::run(), CLFuseBatchNormalizationKernel::run(), CLGEMMLowpMatrixMultiplyReshapedKernel::run(), CLDirectConvolutionLayerKernel::run(), CLColorConvertKernel::run(), CLGEMMMatrixMultiplyNativeKernel::run(), CLSpaceToBatchLayerKernel::run(), CLWeightsReshapeKernel::run(), CLHOGBlockNormalizationKernel::run(), CLEdgeNonMaxSuppressionKernel::run(), CLWinogradOutputTransformKernel::run(), CLMinMaxLocationKernel::run(), CLDepthwiseConvolutionLayerNativeKernel::run(), CLGEMMLowpOffsetContributionOutputStageKernel::run(), CLIm2ColKernel::run(), CLGEMMLowpMatrixAReductionKernel::run(), CLCopyToArrayKernel::run(), CLSobel7x7VertKernel::run(), CLSobel5x5VertKernel::run(), CLGEMMReshapeRHSMatrixKernel::run(), CLGEMMLowpMatrixMultiplyReshapedOnlyRHSKernel::run(), CLLKTrackerStage0Kernel::run(), CLLogits1DNormKernel::run(), CLGEMMMatrixMultiplyReshapedOnlyRHSKernel::run(), CLEdgeTraceKernel::run(), CLGEMMLowpMatrixBReductionKernel::run(), CLLKTrackerStage1Kernel::run(), CLConvolutionRectangleKernel::run(), ClElementwiseKernel::run_op(), ClFloorKernel::run_op(), ClReshapeKernel::run_op(), ClElementWiseUnaryKernel::run_op(), ClCopyKernel::run_op(), ClWidthConcatenate2TensorsKernel::run_op(), ClHeightConcatenateKernel::run_op(), ClActivationKernel::run_op(), ClWidthConcatenateKernel::run_op(), ClPoolingKernel::run_op(), ClDepthConcatenateKernel::run_op(), ClBatchConcatenateKernel::run_op(), ClFillKernel::run_op(), ClWidthConcatenate4TensorsKernel::run_op(), CLStridedSliceKernel::run_op(), ClPermuteKernel::run_op(), CLFillBorderKernel::run_op(), ClCropKernel::run_op(), CLPixelWiseMultiplicationKernel::run_op(), CLComplexPixelWiseMultiplicationKernel::run_op(), and arm_compute::test::validation::TEST_CASE().

277  {
278  return _tuning_params_hint.get_lws();
279  }
cl::NDRange get_lws() const

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

Referenced by CLROIPoolingLayerKernel::configure().

191  {
192  return num_arguments_per_array<1>();
193  }

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

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

199  {
200  return num_arguments_per_tensor<1>();
201  }

◆ num_arguments_per_2D_tensor()

static constexpr unsigned int num_arguments_per_2D_tensor ( )
inlinestatic

◆ num_arguments_per_3D_tensor()

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

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

223  {
224  return num_arguments_per_tensor<4>();
225  }

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

Reimplemented in CLConvolutionRectangleKernel, CLLKTrackerStage1Kernel, CLGEMMLowpMatrixBReductionKernel, CLEdgeTraceKernel, CLGEMMMatrixMultiplyReshapedKernel, CLGEMMMatrixMultiplyReshapedOnlyRHSKernel, CLLogits1DNormKernel, CLLKTrackerStage0Kernel, CLGEMMLowpMatrixMultiplyReshapedOnlyRHSKernel, CLGEMMReshapeRHSMatrixKernel, CLSobel5x5VertKernel, CLSobel7x7VertKernel, CLCopyToArrayKernel, CLGEMMLowpMatrixAReductionKernel, CLIm2ColKernel, CLGEMMLowpOffsetContributionOutputStageKernel, CLDepthwiseConvolutionLayerNativeKernel, CLMinMaxLocationKernel, CLWinogradOutputTransformKernel, CLEdgeNonMaxSuppressionKernel, CLHOGBlockNormalizationKernel, CLSpaceToBatchLayerKernel, CLWeightsReshapeKernel, CLColorConvertKernel, CLGEMMMatrixMultiplyNativeKernel, CLDirectConvolutionLayerKernel, CLGEMMLowpMatrixMultiplyReshapedKernel, CLFuseBatchNormalizationKernel, CLGEMMMatrixMultiplyKernel, CLWinogradFilterTransformKernel, CLBatchNormalizationLayerKernel, CLDepthwiseConvolutionLayer3x3NHWCKernel, CLGEMMLowpOffsetContributionKernel, CLWinogradInputTransformKernel, CLDepthwiseConvolutionLayer3x3NCHWKernel, CLGaussianPyramidVertKernel, CLHistogramBorderKernel, CLBatchToSpaceLayerKernel, CLROIAlignLayerKernel, CLDeconvolutionReshapeOutputKernel, CLCol2ImKernel, CLGEMMReshapeLHSMatrixKernel, CLArgMinMaxLayerKernel, CLGEMMLowpMatrixMultiplyNativeKernel, CLLKTrackerFinalizeKernel, CLLogits1DMaxShiftExpSumKernel, CLStackLayerKernel, CLGEMMLowpQuantizeDownInt32ScaleByFloatKernel, CLGEMMLowpQuantizeDownInt32ScaleKernel, CLChannelCombineKernel, CLBoundingBoxTransformKernel, CLL2NormalizeLayerKernel, CLFFTRadixStageKernel, CLFillBorderKernel, CLHOGDetectorKernel, CLReductionOperationKernel, CLHarrisScoreKernel, CLChannelExtractKernel, CLPadLayerKernel, CLPriorBoxLayerKernel, CLScharr3x3Kernel, CLMeanStdDevKernel, CLNormalizePlanarYUVLayerKernel, CLConvertFullyConnectedWeightsKernel, CLRangeKernel, CLReorgLayerKernel, CLROIPoolingLayerKernel, CLInstanceNormalizationLayerKernel, CLMeanStdDevNormalizationKernel, CLFastCornersKernel, CLFFTDigitReverseKernel, CLGEMMLowpQuantizeDownInt32ScaleByFixedPointKernel, CLTileKernel, CLComparisonKernel, CLGatherKernel, CLIntegralImageVertKernel, CLMinMaxLayerKernel, CLNormalizationLayerKernel, CLMagnitudePhaseKernel, CLQLSTMLayerNormalizationKernel, CLQuantizationLayerKernel, CLScaleKernel, CLFFTScaleKernel, CLComputeAllAnchorsKernel, CLMaxUnpoolingLayerKernel, CLDeconvolutionLayerUpsampleKernel, CLDepthToSpaceLayerKernel, CLReverseKernel, CLSelectKernel, CLSpaceToDepthLayerKernel, CLChannelShuffleLayerKernel, CLDepthwiseConvolutionLayerReshapeWeightsKernel, CLLKTrackerInitKernel, CLSobel5x5HorKernel, CLSobel7x7HorKernel, CLDequantizationLayerKernel, CLDerivativeKernel, CLSobel3x3Kernel, CLAbsoluteDifferenceKernel, CLHOGOrientationBinningKernel, CLRemapKernel, CLGradientKernel, CLMinMaxKernel, CLHistogramKernel, CLBitwiseKernel, CLGaussianPyramidHorKernel, ICLSimple3DKernel, and ICLSimple2DKernel.

Definition at line 233 of file ICLKernel.h.

References ARM_COMPUTE_UNUSED.

Referenced by CLScheduler::init().

234  {
235  ARM_COMPUTE_UNUSED(window, queue);
236  }
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28
#define ARM_COMPUTE_UNUSED(...)
To avoid unused variables warnings.
Definition: Error.h:152

◆ 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]tensorsA vector containing the tensors to operato on.
[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.

Reimplemented in CLComplexPixelWiseMultiplicationKernel, CLPixelWiseMultiplicationKernel, CLFillBorderKernel, ClCropKernel, ClPermuteKernel, CLStridedSliceKernel, ClFillKernel, ClPoolingKernel, ClCopyKernel, ClElementWiseUnaryKernel, ClFloorKernel, and ClReshapeKernel.

Definition at line 245 of file ICLKernel.h.

References ARM_COMPUTE_UNUSED.

Referenced by CLScheduler::init().

246  {
247  ARM_COMPUTE_UNUSED(tensors, window, queue);
248  }
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28
#define ARM_COMPUTE_UNUSED(...)
To avoid unused variables warnings.
Definition: Error.h:152

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

References ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, and CLTuningParams::set_lws().

Referenced by CLGEMMMatrixMultiplyKernel::configure(), arm_compute::test::validation::TEST_CASE(), and CLTuner::tune_kernel_dynamic().

267  {
268  ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); // lws_hint will be overwritten by configure()
269  _tuning_params_hint.set_lws(lws_hint);
270  }
cl::NDRange lws_hint() const
Return the Local-Workgroup-Size hint.
Definition: ICLKernel.h:276
void set_lws(cl::NDRange lws)
#define ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(k)
Definition: Validate.h:941

◆ set_target() [1/2]

void set_target ( GPUTarget  target)
inline

Set the targeted GPU architecture.

Parameters
[in]targetThe targeted GPU architecture

Definition at line 321 of file ICLKernel.h.

Referenced by arm_compute::test::validation::TEST_CASE().

322  {
323  _target = target;
324  }

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

References arm_compute::get_target_from_device().

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

◆ 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_hintworkgroup batch size modifier value

Definition at line 287 of file ICLKernel.h.

References ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, and CLTuningParams::set_wbsm().

Referenced by CLTuner::tune_kernel_dynamic().

288  {
289  ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); // wbsm_hint will be overwritten by configure()
290  _tuning_params_hint.set_wbsm(wbsm_hint);
291  }
cl_int wbsm_hint() const
Return the workgroup batch size modifier hint.
Definition: ICLKernel.h:297
void set_wbsm(cl_int wbsm)
#define ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(k)
Definition: Validate.h:941

◆ wbsm_hint()

cl_int wbsm_hint ( ) const
inline

Return the workgroup batch size modifier hint.

Returns
Current wbsm hint

Definition at line 297 of file ICLKernel.h.

References CLTuningParams::get_wbsm().

298  {
299  return _tuning_params_hint.get_wbsm();
300  }

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