Compute Library
 22.08
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...
 
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...
 
virtual void run_composite_op (ITensorPack &tensors, const Window &window, cl::CommandQueue &queue, const experimental::dynamic_fusion::ClExecutionDescriptor &exec_desc)
 The execution is carried out through run_op method. But the run_op method needs to be extended to include ClExecutionDescriptor as now LWS GWS tuning will be separated from the IKernel. 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...
 
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 81 of file ICLKernel.h.

Constructor & Destructor Documentation

◆ ICLKernel()

ICLKernel ( )
inline

Constructor.

Definition at line 140 of file ICLKernel.h.

References arm_compute::MIDGARD, and arm_compute::UNKNOWN.

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

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

170  {
171  add_array_argument<T, 1>(idx, array, strides, num_dimensions, window);
172  }
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 179 of file ICLKernel.h.

Referenced by CLReverseKernel::run(), CLSelectKernel::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(), ClGemmLowpQuantizeDownInt32ScaleByFixedPointKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleByFloatKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleKernel::run_op(), ClDirectConv2dKernel::run_op(), and ClDirectConv3dKernel::run_op().

180  {
181  add_tensor_argument<1>(idx, tensor, window);
182  }
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 190 of file ICLKernel.h.

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

191  {
192  if(cond)
193  {
194  add_1D_tensor_argument(idx, tensor, window);
195  }
196  }
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:179

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

Referenced by 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().

204  {
205  add_tensor_argument<2>(idx, tensor, window);
206  }
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 214 of file ICLKernel.h.

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

215  {
216  if(cond)
217  {
218  add_2D_tensor_argument(idx, tensor, window);
219  }
220  }
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:203

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

Referenced by ICLSimple3DKernel::run(), CLInstanceNormalizationLayerKernel::run(), CLDepthToSpaceLayerKernel::run(), CLSelectKernel::run(), CLSpaceToDepthLayerKernel::run(), CLDeconvolutionLayerUpsampleKernel::run(), CLMaxUnpoolingLayerKernel::run(), CLFFTScaleKernel::run(), CLComparisonKernel::run(), CLNormalizationLayerKernel::run(), CLROIPoolingLayerKernel::run(), CLFFTDigitReverseKernel::run(), CLReorgLayerKernel::run(), CLNormalizePlanarYUVLayerKernel::run(), CLReductionOperationKernel::run(), CLPadLayerKernel::run(), CLFFTRadixStageKernel::run(), CLFillBorderKernel::run(), CLL2NormalizeLayerKernel::run(), CLArgMinMaxLayerKernel::run(), CLROIAlignLayerKernel::run(), CLDeconvolutionReshapeOutputKernel::run(), CLBatchToSpaceLayerKernel::run(), CLBatchNormalizationLayerKernel::run(), CLFuseBatchNormalizationKernel::run(), CLSpaceToBatchLayerKernel::run(), CLComputeMeanVariance::run(), ClElementwiseKernel::run_op(), ClDequantizeKernel::run_op(), ClFloorKernel::run_op(), ClReshapeKernel::run_op(), ClCopyKernel::run_op(), ClElementWiseUnaryKernel::run_op(), ClFillKernel::run_op(), ClActivationKernel::run_op(), ClQuantizeKernel::run_op(), ClDepthConcatenateKernel::run_op(), ClBatchConcatenateKernel::run_op(), ClCropKernel::run_op(), ClWinogradInputTransformKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleByFixedPointKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleByFloatKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleKernel::run_op(), ClCastKernel::run_op(), ClGemmLowpOffsetContributionKernel::run_op(), ClCol2ImKernel::run_op(), ClDirectConv2dKernel::run_op(), ClGemmLowpOffsetContributionOutputStageKernel::run_op(), ClLogits1DMaxShiftExpSumKernel::run_op(), ClMulKernel::run_op(), ClGemmLowpMatrixAReductionKernel::run_op(), CLFillBorderKernel::run_op(), ClIm2ColKernel::run_op(), ClComplexMulKernel::run_op(), ClLogits1DNormKernel::run_op(), and ClGemmLowpMatrixBReductionKernel::run_op().

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

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

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 ClGemmReshapeLhsMatrixKernel::run_op(), ClGemmMatrixMultiplyReshapedOnlyRhsMMULKernel::run_op(), and ClGemmReshapeRhsMatrixKernel::run_op().

120 {
121  ARM_COMPUTE_ERROR_ON(tensor == nullptr);
122 
123  const ITensorInfo *info = tensor->info();
124  ARM_COMPUTE_ERROR_ON(info == nullptr);
125  const Strides &strides = info->strides_in_bytes();
126 
127  // Tensor poniter
128  _kernel.setArg(idx++, tensor->cl_buffer());
129 
130  // Add stride_y, stride_z
131  _kernel.setArg<cl_uint>(idx++, strides[1]);
132  _kernel.setArg<cl_uint>(idx++, strides[2]);
133 
134  // Tensor dimensions
135  _kernel.setArg<cl_uint>(idx++, info->dimension(0));
136  _kernel.setArg<cl_uint>(idx++, info->dimension(1));
137  _kernel.setArg<cl_uint>(idx++, info->dimension(2));
138 
139  // Offset of first element
140  unsigned int offset_first_element = info->offset_first_element_in_bytes();
141  _kernel.setArg<cl_uint>(idx++, offset_first_element);
142 }
#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
ScaleKernelInfo info(interpolation_policy, default_border_mode, PixelValue(), sampling_policy, false)

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

Referenced by CLChannelShuffleLayerKernel::run(), CLInstanceNormalizationLayerKernel::run(), CLDepthToSpaceLayerKernel::run(), CLReverseKernel::run(), CLSpaceToDepthLayerKernel::run(), CLGatherKernel::run(), CLTileKernel::run(), CLReductionOperationKernel::run(), CLStackLayerKernel::run(), CLDepthwiseConvolutionLayerNativeKernel::run(), CLArgMinMaxLayerKernel::run(), CLBatchToSpaceLayerKernel::run(), CLSpaceToBatchLayerKernel::run(), CLComputeMeanVariance::run(), ClWidthConcatenate2TensorsKernel::run_op(), ClHeightConcatenateKernel::run_op(), ClWidthConcatenateKernel::run_op(), ClPermuteKernel::run_op(), ClWidthConcatenate4TensorsKernel::run_op(), ClWinogradInputTransformKernel::run_op(), CLStridedSliceKernel::run_op(), ClCol2ImKernel::run_op(), and ClDirectConv3dKernel::run_op().

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

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

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(), and IKernel::window().

Referenced by CLDepthwiseConvolutionLayerNativeKernel::run(), ClScaleKernel::run_op(), and ClDirectConv2dKernel::run_op().

145 {
146  ARM_COMPUTE_ERROR_ON(tensor == nullptr);
147 
148  const ITensorInfo *info = tensor->info();
149  ARM_COMPUTE_ERROR_ON(info == nullptr);
150  const Strides &strides = info->strides_in_bytes();
151 
152  // Tensor poniter
153  _kernel.setArg(idx++, tensor->cl_buffer());
154 
155  // Add stride_y, stride_z and stride_w
156  _kernel.setArg<cl_uint>(idx++, strides[1]);
157  _kernel.setArg<cl_uint>(idx++, strides[2]);
158  _kernel.setArg<cl_uint>(idx++, strides[3]);
159 
160  // Tensor dimensions
161  _kernel.setArg<cl_uint>(idx++, info->dimension(0));
162  _kernel.setArg<cl_uint>(idx++, info->dimension(1));
163  _kernel.setArg<cl_uint>(idx++, info->dimension(2));
164  _kernel.setArg<cl_uint>(idx++, info->dimension(3));
165 
166  // Offset of first element
167  unsigned int offset_first_element = info->offset_first_element_in_bytes();
168  _kernel.setArg<cl_uint>(idx++, offset_first_element);
169 }
#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
ScaleKernelInfo info(interpolation_policy, default_border_mode, PixelValue(), sampling_policy, false)

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

248  {
249  add_tensor_argument<5>(idx, tensor, window);
250  }
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 363 of file ICLKernel.h.

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

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

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

519 {
520  ARM_COMPUTE_ERROR_ON(array == nullptr);
521 
522  // Calculate offset to the start of the window
523  unsigned int offset_first_element = 0;
524 
525  for(unsigned int n = 0; n < num_dimensions; ++n)
526  {
527  offset_first_element += window[n].start() * strides[n];
528  }
529 
530  unsigned int idx_start = idx;
531  _kernel.setArg(idx++, array->cl_buffer());
532 
533  for(unsigned int dimension = 0; dimension < dimension_size; dimension++)
534  {
535  _kernel.setArg<cl_uint>(idx++, strides[dimension]);
536  _kernel.setArg<cl_uint>(idx++, strides[dimension] * window[dimension].step());
537  }
538 
539  _kernel.setArg<cl_uint>(idx++, offset_first_element);
540 
541  ARM_COMPUTE_ERROR_ON_MSG_VAR(idx_start + num_arguments_per_array<dimension_size>() != idx,
542  "add_%dD_array_argument() is supposed to add exactly %d arguments to the kernel", dimension_size, num_arguments_per_array<dimension_size>());
543  ARM_COMPUTE_UNUSED(idx_start);
544 }
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(), arm_compute::test::validation::n, ITensorInfo::num_dimensions(), ITensorInfo::offset_first_element_in_bytes(), and ITensorInfo::strides_in_bytes().

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++, window.is_broadcasted(d) ? 0 : strides[d]);
109  _kernel.setArg<cl_uint>(idx++, window.is_broadcasted(d) ? 0 : (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 420 of file ICLKernel.h.

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

421  {
422  return _config_id;
423  }

◆ 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().

185 {
186  if(_max_workgroup_size == 0)
187  {
188  _max_workgroup_size = CLKernelLibrary::get().max_local_workgroup_size(_kernel);
189  }
190  return _max_workgroup_size;
191 }
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()

GPUTarget get_target ( ) const
inline

Get the targeted GPU architecture.

Returns
The targeted GPU architecture.

Definition at line 444 of file ICLKernel.h.

References arm_compute::enqueue(), and CLKernelLibrary::get().

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

445  {
446  return _target;
447  }

◆ 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 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().

194 {
195  if((window.x().end() - window.x().start()) == 0 || (window.y().end() - window.y().start()) == 0)
196  {
197  return cl::NullRange;
198  }
199 
200  cl::NDRange gws((window.x().end() - window.x().start()) / window.x().step(),
201  (window.y().end() - window.y().start()) / window.y().step(),
202  (window.z().end() - window.z().start()) / window.z().step());
203 
204  return gws;
205 }
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:107
constexpr const Dimension & z() const
Alias to access the third dimension of the window.
Definition: Window.h:177
constexpr const Dimension & y() const
Alias to access the second dimension of the window.
Definition: Window.h:168
constexpr int end() const
Return the end of the dimension.
Definition: Window.h:102
constexpr int start() const
Return the start of the dimension.
Definition: Window.h:97
constexpr const Dimension & x() const
Alias to access the first dimension of the window.
Definition: Window.h:159

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

Referenced by arm_compute::enqueue().

149  {
150  return _kernel;
151  }

◆ lws_hint()

cl::NDRange lws_hint ( ) const
inline

Return the Local-Workgroup-Size hint.

Returns
Current lws hint

Definition at line 384 of file ICLKernel.h.

Referenced by ClLogits1DMaxShiftExpSumKernel::configure(), CLArgMinMaxLayerKernel::configure(), ICLSimple2DKernel::run(), ICLSimple3DKernel::run(), CLBitwiseKernel::run(), CLChannelShuffleLayerKernel::run(), CLInstanceNormalizationLayerKernel::run(), CLReverseKernel::run(), CLSelectKernel::run(), CLSpaceToDepthLayerKernel::run(), CLDepthToSpaceLayerKernel::run(), CLDeconvolutionLayerUpsampleKernel::run(), CLFFTScaleKernel::run(), CLComputeAllAnchorsKernel::run(), CLMaxUnpoolingLayerKernel::run(), CLQLSTMLayerNormalizationKernel::run(), CLROIPoolingLayerKernel::run(), CLGatherKernel::run(), CLComparisonKernel::run(), CLNormalizationLayerKernel::run(), CLTileKernel::run(), CLFFTDigitReverseKernel::run(), CLMeanStdDevNormalizationKernel::run(), CLRangeKernel::run(), CLReorgLayerKernel::run(), CLNormalizePlanarYUVLayerKernel::run(), CLPadLayerKernel::run(), CLPriorBoxLayerKernel::run(), CLFFTRadixStageKernel::run(), CLFillBorderKernel::run(), CLL2NormalizeLayerKernel::run(), CLBoundingBoxTransformKernel::run(), CLStackLayerKernel::run(), CLArgMinMaxLayerKernel::run(), CLDepthwiseConvolutionLayerNativeKernel::run(), CLROIAlignLayerKernel::run(), CLDeconvolutionReshapeOutputKernel::run(), CLBatchToSpaceLayerKernel::run(), CLBatchNormalizationLayerKernel::run(), CLFuseBatchNormalizationKernel::run(), CLSpaceToBatchLayerKernel::run(), CLComputeMeanVariance::run(), ClElementwiseKernel::run_op(), ClTransposeKernel::run_op(), ClDequantizeKernel::run_op(), ClReshapeKernel::run_op(), ClFloorKernel::run_op(), ClCopyKernel::run_op(), ClElementWiseUnaryKernel::run_op(), ClFillKernel::run_op(), ClWidthConcatenate2TensorsKernel::run_op(), ClScaleKernel::run_op(), ClHeightConcatenateKernel::run_op(), ClActivationKernel::run_op(), ClWidthConcatenateKernel::run_op(), ClQuantizeKernel::run_op(), ClWidthConcatenate4TensorsKernel::run_op(), ClPermuteKernel::run_op(), ClBatchConcatenateKernel::run_op(), ClDepthConcatenateKernel::run_op(), ClConvertFullyConnectedWeightsKernel::run_op(), ClCropKernel::run_op(), ClGemmLowpMatrixMultiplyNativeKernel::run_op(), ClGemmReshapeLhsMatrixKernel::run_op(), ClWinogradInputTransformKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleByFixedPointKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleKernel::run_op(), ClCastKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleByFloatKernel::run_op(), ClGemmMatrixMultiplyNativeKernel::run_op(), CLStridedSliceKernel::run_op(), ClGemmReshapeRhsMatrixKernel::run_op(), ClGemmLowpOffsetContributionKernel::run_op(), ClCol2ImKernel::run_op(), ClDirectConv2dKernel::run_op(), ClGemmLowpOffsetContributionOutputStageKernel::run_op(), ClDirectConv3dKernel::run_op(), ClLogits1DMaxShiftExpSumKernel::run_op(), ClMulKernel::run_op(), ClGemmLowpMatrixAReductionKernel::run_op(), CLFillBorderKernel::run_op(), ClGemmMatrixMultiplyReshapedOnlyRhsKernel::run_op(), ClIm2ColKernel::run_op(), ClGemmMatrixMultiplyReshapedKernel::run_op(), ClComplexMulKernel::run_op(), ClLogits1DNormKernel::run_op(), and ClGemmLowpMatrixBReductionKernel::run_op().

385  {
386  return _tuning_params_hint.get_lws();
387  }
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 290 of file ICLKernel.h.

291  {
292  return num_arguments_per_array<1>();
293  }

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

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

299  {
300  return num_arguments_per_tensor<1>();
301  }

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

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

307  {
308  return num_arguments_per_tensor<2>();
309  }

◆ num_arguments_per_3D_tensor()

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

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

264  {
265  constexpr unsigned int no_args_per_3d_tensor_nhw = 7u;
266  return no_args_per_3d_tensor_nhw;
267  }

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

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

323  {
324  return num_arguments_per_tensor<4>();
325  }

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

Referenced by ClScaleKernel::configure().

281  {
282  constexpr unsigned int no_args_per_4d_tensor_nhwc = 9u;
283  return no_args_per_4d_tensor_nhwc;
284  }

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

References ARM_COMPUTE_UNUSED.

Referenced by CLTuner::CLTuner(), and CLScheduler::init().

334  {
335  ARM_COMPUTE_UNUSED(window, queue);
336  }
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_composite_op()

virtual void run_composite_op ( ITensorPack tensors,
const Window window,
cl::CommandQueue &  queue,
const experimental::dynamic_fusion::ClExecutionDescriptor exec_desc 
)
inlinevirtual

The execution is carried out through run_op method. But the run_op method needs to be extended to include ClExecutionDescriptor as now LWS GWS tuning will be separated from the IKernel.

Reimplemented in ClCompositeKernel.

Definition at line 352 of file ICLKernel.h.

References ARM_COMPUTE_UNUSED.

Referenced by CLTuner::CLTuner(), and CLScheduler::init().

353  {
354  ARM_COMPUTE_UNUSED(tensors, window, queue, exec_desc);
355  }
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 ClGemmLowpMatrixBReductionKernel, ClComplexMulKernel, ClGemmMatrixMultiplyReshapedKernel, ClIm2ColKernel, ClGemmMatrixMultiplyReshapedOnlyRhsKernel, CLFillBorderKernel, ClWeightsReshapeKernel, ClGemmLowpMatrixMultiplyReshapedOnlyRhsKernel, ClGemmLowpMatrixAReductionKernel, ClMulKernel, ClDirectConv3dKernel, ClDirectConv2dKernel, ClGemmLowpMatrixMultiplyReshapedOnlyRhsMMULKernel, ClGemmLowpOffsetContributionOutputStageKernel, ClCol2ImKernel, ClGemmLowpOffsetContributionKernel, CLStridedSliceKernel, ClGemmLowpMatrixMultiplyReshapedKernel, ClGemmReshapeRhsMatrixKernel, ClGemmMatrixMultiplyReshapedOnlyRhsMMULKernel, ClGemmMatrixMultiplyNativeKernel, ClGemmLowpQuantizeDownInt32ScaleByFloatKernel, ClGemmLowpQuantizeDownInt32ScaleKernel, ClWinogradOutputTransformKernel, ClGemmLowpQuantizeDownInt32ScaleByFixedPointKernel, ClWinogradFilterTransformKernel, ClWinogradInputTransformKernel, ClGemmLowpMatrixMultiplyNativeKernel, ClGemmReshapeLhsMatrixKernel, ClCropKernel, ClPermuteKernel, ClQuantizeKernel, ClPool2dKernel, ClPool3dKernel, ClScaleKernel, ClCopyKernel, ClElementWiseUnaryKernel, ClFillKernel, ClDequantizeKernel, ClFloorKernel, ClReshapeKernel, and ClTransposeKernel.

Definition at line 345 of file ICLKernel.h.

References ARM_COMPUTE_UNUSED.

Referenced by CLTuner::CLTuner(), and CLScheduler::init().

346  {
347  ARM_COMPUTE_UNUSED(tensors, window, queue);
348  }
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 374 of file ICLKernel.h.

References ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL.

Referenced by CLTuner::tune_kernel_dynamic().

375  {
376  ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); // lws_hint will be overwritten by configure()
377  _tuning_params_hint.set_lws(lws_hint);
378  }
cl::NDRange lws_hint() const
Return the Local-Workgroup-Size hint.
Definition: ICLKernel.h:384
void set_lws(cl::NDRange lws)
#define ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(k)
Definition: Validate.h:915

◆ set_target() [1/2]

void set_target ( GPUTarget  target)
inline

Set the targeted GPU architecture.

Parameters
[in]targetThe targeted GPU architecture

Definition at line 429 of file ICLKernel.h.

430  {
431  _target = target;
432  }

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

References arm_compute::get_target_from_device().

180 {
181  _target = get_target_from_device(device);
182 }
GPUTarget get_target_from_device(const cl::Device &device)
Helper function to get the GPU target from CL device.
Definition: CLHelpers.cpp:223

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

References ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL.

Referenced by CLTuner::tune_kernel_dynamic().

396  {
397  ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); // wbsm_hint will be overwritten by configure()
398  _tuning_params_hint.set_wbsm(wbsm_hint);
399  }
cl_int wbsm_hint() const
Return the workgroup batch size modifier hint.
Definition: ICLKernel.h:405
void set_wbsm(cl_int wbsm)
#define ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(k)
Definition: Validate.h:915

◆ type()

CLKernelType type ( ) const
inline

Returns the CL kernel type.

Returns
The CL kernel type

Definition at line 156 of file ICLKernel.h.

157  {
158  return _type;
159  }

◆ wbsm_hint()

cl_int wbsm_hint ( ) const
inline

Return the workgroup batch size modifier hint.

Returns
Current wbsm hint

Definition at line 405 of file ICLKernel.h.

Referenced by arm_compute::enqueue().

406  {
407  return _tuning_params_hint.get_wbsm();
408  }

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