Compute Library
 21.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...
 
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...
 
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_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 68 of file ICLKernel.h.

Constructor & Destructor Documentation

◆ ICLKernel()

ICLKernel ( )
inline

Constructor.

Definition at line 127 of file ICLKernel.h.

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

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

References IKernel::window().

157  {
158  add_array_argument<T, 1>(idx, array, strides, num_dimensions, window);
159  }
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 166 of file ICLKernel.h.

References IKernel::window().

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

167  {
168  add_tensor_argument<1>(idx, tensor, window);
169  }
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 177 of file ICLKernel.h.

References ICLKernel::add_1D_tensor_argument().

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

178  {
179  if(cond)
180  {
181  add_1D_tensor_argument(idx, tensor, window);
182  }
183  }
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:166

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

References IKernel::window().

Referenced by ICLKernel::add_2D_tensor_argument_if(), ICLSimple2DKernel::run(), CLBitwiseKernel::run(), CLRemapKernel::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(), ClGemmMatrixMultiplyKernel::run_op(), ClGemmLowpMatrixAReductionKernel::run_op(), ClGemmMatrixMultiplyReshapedOnlyRhsKernel::run_op(), ClIm2ColKernel::run_op(), ClGemmMatrixMultiplyReshapedKernel::run_op(), and ClGemmLowpMatrixBReductionKernel::run_op().

191  {
192  add_tensor_argument<2>(idx, tensor, window);
193  }
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 201 of file ICLKernel.h.

References ICLKernel::add_2D_tensor_argument().

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

202  {
203  if(cond)
204  {
205  add_2D_tensor_argument(idx, tensor, window);
206  }
207  }
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:190

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

References IKernel::window().

Referenced by ICLSimple3DKernel::run(), CLInstanceNormalizationLayerKernel::run(), CLDepthToSpaceLayerKernel::run(), CLSelectKernel::run(), CLSpaceToDepthLayerKernel::run(), CLDeconvolutionLayerUpsampleKernel::run(), CLMaxUnpoolingLayerKernel::run(), CLFFTScaleKernel::run(), CLComparisonKernel::run(), CLMinMaxLayerKernel::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(), ClFillKernel::run_op(), ClElementWiseUnaryKernel::run_op(), ClActivationKernel::run_op(), ClQuantizeKernel::run_op(), ClDepthConcatenateKernel::run_op(), ClBatchConcatenateKernel::run_op(), ClCropKernel::run_op(), ClGemmReshapeLhsMatrixKernel::run_op(), ClWinogradInputTransformKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleByFixedPointKernel::run_op(), ClCastKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleByFloatKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleKernel::run_op(), ClDirectConv2dKernel::run_op(), ClGemmReshapeRhsMatrixKernel::run_op(), ClCol2ImKernel::run_op(), ClGemmLowpOffsetContributionKernel::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().

215  {
216  add_tensor_argument<3>(idx, tensor, window);
217  }
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 224 of file ICLKernel.h.

References IKernel::window().

Referenced by CLChannelShuffleLayerKernel::run(), CLInstanceNormalizationLayerKernel::run(), CLDepthToSpaceLayerKernel::run(), CLReverseKernel::run(), CLSpaceToDepthLayerKernel::run(), CLRemapKernel::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(), ClScaleKernel::run_op(), ClWidthConcatenateKernel::run_op(), ClWidthConcatenate4TensorsKernel::run_op(), ClPermuteKernel::run_op(), ClWinogradInputTransformKernel::run_op(), ClDirectConv2dKernel::run_op(), CLStridedSliceKernel::run_op(), and ClCol2ImKernel::run_op().

225  {
226  add_tensor_argument<4>(idx, tensor, window);
227  }
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 297 of file ICLKernel.h.

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

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

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

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

453 {
454  ARM_COMPUTE_ERROR_ON(array == nullptr);
455 
456  // Calculate offset to the start of the window
457  unsigned int offset_first_element = 0;
458 
459  for(unsigned int n = 0; n < num_dimensions; ++n)
460  {
461  offset_first_element += window[n].start() * strides[n];
462  }
463 
464  unsigned int idx_start = idx;
465  _kernel.setArg(idx++, array->cl_buffer());
466 
467  for(unsigned int dimension = 0; dimension < dimension_size; dimension++)
468  {
469  _kernel.setArg<cl_uint>(idx++, strides[dimension]);
470  _kernel.setArg<cl_uint>(idx++, strides[dimension] * window[dimension].step());
471  }
472 
473  _kernel.setArg<cl_uint>(idx++, offset_first_element);
474 
475  ARM_COMPUTE_ERROR_ON_MSG_VAR(idx_start + num_arguments_per_array<dimension_size>() != idx,
476  "add_%dD_array_argument() is supposed to add exactly %d arguments to the kernel", dimension_size, num_arguments_per_array<dimension_size>());
477  ARM_COMPUTE_UNUSED(idx_start);
478 }
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++, 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 354 of file ICLKernel.h.

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

355  {
356  return _config_id;
357  }

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

GPUTarget get_target ( ) const
inline

◆ 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::Dimension::step(), Window::x(), Window::y(), and Window::z().

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

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

136  {
137  return _kernel;
138  }

◆ lws_hint()

cl::NDRange lws_hint ( ) const
inline

Return the Local-Workgroup-Size hint.

Returns
Current lws hint

Definition at line 318 of file ICLKernel.h.

References CLTuningParams::get_lws().

Referenced by ClLogits1DMaxShiftExpSumKernel::configure(), CLArgMinMaxLayerKernel::configure(), ICLKernel::get_target(), ICLSimple2DKernel::run(), ICLSimple3DKernel::run(), CLBitwiseKernel::run(), CLChannelShuffleLayerKernel::run(), CLInstanceNormalizationLayerKernel::run(), CLReverseKernel::run(), CLSelectKernel::run(), CLSpaceToDepthLayerKernel::run(), CLDepthToSpaceLayerKernel::run(), CLRemapKernel::run(), CLDeconvolutionLayerUpsampleKernel::run(), CLFFTScaleKernel::run(), CLComputeAllAnchorsKernel::run(), CLMaxUnpoolingLayerKernel::run(), CLQLSTMLayerNormalizationKernel::run(), CLROIPoolingLayerKernel::run(), CLGatherKernel::run(), CLComparisonKernel::run(), CLMinMaxLayerKernel::run(), CLNormalizationLayerKernel::run(), CLFFTDigitReverseKernel::run(), CLTileKernel::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(), ClReshapeKernel::run_op(), ClTransposeKernel::run_op(), ClDequantizeKernel::run_op(), ClFloorKernel::run_op(), ClFillKernel::run_op(), ClCopyKernel::run_op(), ClElementWiseUnaryKernel::run_op(), ClWidthConcatenate2TensorsKernel::run_op(), ClWidthConcatenateKernel::run_op(), ClHeightConcatenateKernel::run_op(), ClScaleKernel::run_op(), ClActivationKernel::run_op(), ClQuantizeKernel::run_op(), ClPermuteKernel::run_op(), ClWidthConcatenate4TensorsKernel::run_op(), ClBatchConcatenateKernel::run_op(), ClDepthConcatenateKernel::run_op(), ClCropKernel::run_op(), ClConvertFullyConnectedWeightsKernel::run_op(), ClGemmReshapeLhsMatrixKernel::run_op(), ClGemmLowpMatrixMultiplyNativeKernel::run_op(), ClWinogradInputTransformKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleByFixedPointKernel::run_op(), ClCastKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleKernel::run_op(), ClGemmLowpQuantizeDownInt32ScaleByFloatKernel::run_op(), ClGemmMatrixMultiplyNativeKernel::run_op(), ClGemmMatrixMultiplyKernel::run_op(), ClDirectConv2dKernel::run_op(), CLStridedSliceKernel::run_op(), ClGemmReshapeRhsMatrixKernel::run_op(), ClCol2ImKernel::run_op(), ClGemmLowpOffsetContributionKernel::run_op(), ClGemmLowpOffsetContributionOutputStageKernel::run_op(), ClLogits1DMaxShiftExpSumKernel::run_op(), ClMulKernel::run_op(), ClGemmLowpMatrixAReductionKernel::run_op(), CLFillBorderKernel::run_op(), ClGemmMatrixMultiplyReshapedOnlyRhsKernel::run_op(), ClIm2ColKernel::run_op(), ClGemmMatrixMultiplyReshapedKernel::run_op(), ClLogits1DNormKernel::run_op(), ClComplexMulKernel::run_op(), and ClGemmLowpMatrixBReductionKernel::run_op().

319  {
320  return _tuning_params_hint.get_lws();
321  }
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 232 of file ICLKernel.h.

233  {
234  return num_arguments_per_array<1>();
235  }

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

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

241  {
242  return num_arguments_per_tensor<1>();
243  }

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

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

249  {
250  return num_arguments_per_tensor<2>();
251  }

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

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

265  {
266  return num_arguments_per_tensor<4>();
267  }

◆ 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, CLMinMaxLayerKernel, CLNormalizationLayerKernel, CLROIPoolingLayerKernel, CLQLSTMLayerNormalizationKernel, CLFFTScaleKernel, CLComputeAllAnchorsKernel, CLMaxUnpoolingLayerKernel, CLDeconvolutionLayerUpsampleKernel, CLRemapKernel, CLDepthToSpaceLayerKernel, CLReverseKernel, CLSelectKernel, CLSpaceToDepthLayerKernel, CLChannelShuffleLayerKernel, CLInstanceNormalizationLayerKernel, CLBitwiseKernel, ICLSimple3DKernel, and ICLSimple2DKernel.

Definition at line 275 of file ICLKernel.h.

References ARM_COMPUTE_UNUSED.

Referenced by CLScheduler::init().

276  {
277  ARM_COMPUTE_UNUSED(window, queue);
278  }
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, ClGemmLowpOffsetContributionOutputStageKernel, ClCol2ImKernel, ClGemmLowpOffsetContributionKernel, CLStridedSliceKernel, ClGemmLowpMatrixMultiplyReshapedKernel, ClGemmReshapeRhsMatrixKernel, ClDirectConv2dKernel, ClGemmMatrixMultiplyKernel, ClGemmMatrixMultiplyNativeKernel, ClGemmLowpQuantizeDownInt32ScaleByFloatKernel, ClGemmLowpQuantizeDownInt32ScaleKernel, ClWinogradOutputTransformKernel, ClGemmLowpQuantizeDownInt32ScaleByFixedPointKernel, ClWinogradFilterTransformKernel, ClWinogradInputTransformKernel, ClGemmLowpMatrixMultiplyNativeKernel, ClGemmReshapeLhsMatrixKernel, ClCropKernel, ClPermuteKernel, ClQuantizeKernel, ClPool2dKernel, ClScaleKernel, ClCopyKernel, ClElementWiseUnaryKernel, ClFillKernel, ClDequantizeKernel, ClFloorKernel, ClReshapeKernel, and ClTransposeKernel.

Definition at line 287 of file ICLKernel.h.

References ARM_COMPUTE_UNUSED.

Referenced by CLScheduler::init().

288  {
289  ARM_COMPUTE_UNUSED(tensors, window, queue);
290  }
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 308 of file ICLKernel.h.

References ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, and CLTuningParams::set_lws().

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

309  {
310  ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); // lws_hint will be overwritten by configure()
311  _tuning_params_hint.set_lws(lws_hint);
312  }
cl::NDRange lws_hint() const
Return the Local-Workgroup-Size hint.
Definition: ICLKernel.h:318
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 363 of file ICLKernel.h.

364  {
365  _target = target;
366  }

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

References ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, and CLTuningParams::set_wbsm().

Referenced by CLTuner::tune_kernel_dynamic().

330  {
331  ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); // wbsm_hint will be overwritten by configure()
332  _tuning_params_hint.set_wbsm(wbsm_hint);
333  }
cl_int wbsm_hint() const
Return the workgroup batch size modifier hint.
Definition: ICLKernel.h:339
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 143 of file ICLKernel.h.

144  {
145  return _type;
146  }

◆ wbsm_hint()

cl_int wbsm_hint ( ) const
inline

Return the workgroup batch size modifier hint.

Returns
Current wbsm hint

Definition at line 339 of file ICLKernel.h.

References CLTuningParams::get_wbsm().

Referenced by arm_compute::enqueue().

340  {
341  return _tuning_params_hint.get_wbsm();
342  }

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