Compute Library
 21.05
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...
 
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 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.

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

References IKernel::window().

◆ add_1D_tensor_argument()

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

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

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

Definition at line 124 of file ICLKernel.h.

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

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(), CLGEMMLowpQuantizeDownInt32ScaleByFixedPointKernel::run(), CLRangeKernel::run(), CLNormalizePlanarYUVLayerKernel::run(), CLReductionOperationKernel::run(), CLBoundingBoxTransformKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleByFloatKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleKernel::run(), CLDeconvolutionReshapeOutputKernel::run(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), CLBatchToSpaceLayerKernel::run(), CLBatchNormalizationLayerKernel::run(), CLFuseBatchNormalizationKernel::run(), CLDepthwiseConvolutionLayer3x3NCHWKernel::run(), CLWeightsReshapeKernel::run(), CLWinogradOutputTransformKernel::run(), CLDepthwiseConvolutionLayerNativeKernel::run(), and ClDirectConvolutionKernel::run_op().

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

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's parameters to the object's kernel's arguments starting from the index idx.
Definition: ICLKernel.h:124

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

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

◆ add_2D_tensor_argument()

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

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

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

Definition at line 148 of file ICLKernel.h.

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

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(), CLGEMMLowpMatrixMultiplyNativeKernel::run(), CLArgMinMaxLayerKernel::run(), CLROIAlignLayerKernel::run(), CLGEMMMatrixMultiplyKernel::run(), CLGEMMLowpMatrixMultiplyReshapedKernel::run(), CLGEMMMatrixMultiplyNativeKernel::run(), CLWeightsReshapeKernel::run(), CLIm2ColKernel::run(), CLGEMMLowpMatrixAReductionKernel::run(), CLGEMMLowpMatrixMultiplyReshapedOnlyRHSKernel::run(), CLGEMMMatrixMultiplyReshapedOnlyRHSKernel::run(), CLGEMMLowpMatrixBReductionKernel::run(), ClTransposeKernel::run_op(), ClScaleKernel::run_op(), and ClConvertFullyConnectedWeightsKernel::run_op().

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

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's parameters to the object's kernel's arguments starting from the index idx.
Definition: ICLKernel.h:148

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

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

◆ add_3D_tensor_argument()

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

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

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

Definition at line 172 of file ICLKernel.h.

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

References IKernel::window().

Referenced by ICLSimple3DKernel::run(), CLInstanceNormalizationLayerKernel::run(), CLDepthToSpaceLayerKernel::run(), CLSelectKernel::run(), CLSpaceToDepthLayerKernel::run(), CLDeconvolutionLayerUpsampleKernel::run(), CLMaxUnpoolingLayerKernel::run(), CLFFTScaleKernel::run(), CLMinMaxLayerKernel::run(), CLNormalizationLayerKernel::run(), CLROIPoolingLayerKernel::run(), CLComparisonKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleByFixedPointKernel::run(), CLFFTDigitReverseKernel::run(), CLReorgLayerKernel::run(), CLNormalizePlanarYUVLayerKernel::run(), CLReductionOperationKernel::run(), CLPadLayerKernel::run(), CLFillBorderKernel::run(), CLFFTRadixStageKernel::run(), CLL2NormalizeLayerKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleByFloatKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleKernel::run(), CLArgMinMaxLayerKernel::run(), CLGEMMReshapeLHSMatrixKernel::run(), CLCol2ImKernel::run(), CLDeconvolutionReshapeOutputKernel::run(), CLROIAlignLayerKernel::run(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), CLBatchToSpaceLayerKernel::run(), CLGEMMLowpOffsetContributionKernel::run(), CLWinogradInputTransformKernel::run(), CLBatchNormalizationLayerKernel::run(), CLWinogradFilterTransformKernel::run(), CLFuseBatchNormalizationKernel::run(), CLDepthwiseConvolutionLayer3x3NCHWKernel::run(), CLWeightsReshapeKernel::run(), CLSpaceToBatchLayerKernel::run(), CLDepthwiseConvolutionLayerNativeKernel::run(), CLComputeMeanVariance::run(), CLGEMMLowpOffsetContributionOutputStageKernel::run(), CLIm2ColKernel::run(), CLGEMMLowpMatrixAReductionKernel::run(), CLGEMMReshapeRHSMatrixKernel::run(), CLGEMMLowpMatrixBReductionKernel::run(), ClElementwiseKernel::run_op(), ClFloorKernel::run_op(), ClReshapeKernel::run_op(), ClCopyKernel::run_op(), ClElementWiseUnaryKernel::run_op(), ClDequantizationKernel::run_op(), ClActivationKernel::run_op(), ClPoolingKernel::run_op(), ClQuantizationKernel::run_op(), ClFillKernel::run_op(), ClDepthConcatenateKernel::run_op(), ClBatchConcatenateKernel::run_op(), ClMulKernel::run_op(), ClDirectConvolutionKernel::run_op(), ClCropKernel::run_op(), CLFillBorderKernel::run_op(), ClLogits1DMaxShiftExpSumKernel::run_op(), ClComplexMulKernel::run_op(), and ClLogits1DNormKernel::run_op().

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

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

References IKernel::window().

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

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

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

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

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

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

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

◆ add_tensor_argument()

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

Definition at line 88 of file ICLKernel.cpp.

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)

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

◆ config_id()

const std::string& config_id ( ) const
inline

Get the configuration ID.

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

Definition at line 312 of file ICLKernel.h.

313  {
314  return _config_id;
315  }

Referenced by CLTuner::tune_kernel_dynamic().

◆ get_max_workgroup_size()

size_t get_max_workgroup_size ( )

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

Returns
The maximum workgroup size value.

Definition at line 131 of file ICLKernel.cpp.

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.

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

Referenced by arm_compute::enqueue().

◆ get_target()

GPUTarget get_target ( ) const
inline

Get the targeted GPU architecture.

Returns
The targeted GPU architecture.

Definition at line 336 of file ICLKernel.h.

337  {
338  return _target;
339  }

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

◆ gws_from_window()

cl::NDRange gws_from_window ( const Window window)
static

Get the global work size given an execution window.

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

Definition at line 140 of file ICLKernel.cpp.

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

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

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

102  {
103  return _kernel;
104  }

Referenced by arm_compute::enqueue().

◆ lws_hint()

cl::NDRange lws_hint ( ) const
inline

Return the Local-Workgroup-Size hint.

Returns
Current lws hint

Definition at line 276 of file ICLKernel.h.

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

References CLTuningParams::get_lws().

Referenced by ClLogits1DMaxShiftExpSumKernel::configure(), CLArgMinMaxLayerKernel::configure(), ICLSimple2DKernel::run(), ICLSimple3DKernel::run(), CLBitwiseKernel::run(), CLRemapKernel::run(), CLInstanceNormalizationLayerKernel::run(), CLChannelShuffleLayerKernel::run(), CLDepthToSpaceLayerKernel::run(), CLReverseKernel::run(), CLSelectKernel::run(), CLSpaceToDepthLayerKernel::run(), CLDeconvolutionLayerUpsampleKernel::run(), CLComputeAllAnchorsKernel::run(), CLFFTScaleKernel::run(), CLMaxUnpoolingLayerKernel::run(), CLQLSTMLayerNormalizationKernel::run(), CLMinMaxLayerKernel::run(), CLGatherKernel::run(), CLNormalizationLayerKernel::run(), CLComparisonKernel::run(), CLROIPoolingLayerKernel::run(), CLTileKernel::run(), CLFFTDigitReverseKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleByFixedPointKernel::run(), CLMeanStdDevNormalizationKernel::run(), CLRangeKernel::run(), CLReorgLayerKernel::run(), CLNormalizePlanarYUVLayerKernel::run(), CLPadLayerKernel::run(), CLPriorBoxLayerKernel::run(), CLFFTRadixStageKernel::run(), CLFillBorderKernel::run(), CLL2NormalizeLayerKernel::run(), CLBoundingBoxTransformKernel::run(), CLGEMMLowpQuantizeDownInt32ScaleKernel::run(), CLStackLayerKernel::run(), CLArgMinMaxLayerKernel::run(), CLGEMMLowpMatrixMultiplyNativeKernel::run(), CLGEMMReshapeLHSMatrixKernel::run(), CLCol2ImKernel::run(), CLDepthwiseConvolutionLayer3x3NHWCKernel::run(), CLROIAlignLayerKernel::run(), CLDeconvolutionReshapeOutputKernel::run(), CLBatchToSpaceLayerKernel::run(), CLWinogradInputTransformKernel::run(), CLGEMMLowpOffsetContributionKernel::run(), CLBatchNormalizationLayerKernel::run(), CLGEMMMatrixMultiplyKernel::run(), CLWinogradFilterTransformKernel::run(), CLFuseBatchNormalizationKernel::run(), CLGEMMLowpMatrixMultiplyReshapedKernel::run(), CLDepthwiseConvolutionLayer3x3NCHWKernel::run(), CLGEMMMatrixMultiplyNativeKernel::run(), CLWeightsReshapeKernel::run(), CLSpaceToBatchLayerKernel::run(), CLWinogradOutputTransformKernel::run(), CLDepthwiseConvolutionLayerNativeKernel::run(), CLComputeMeanVariance::run(), CLGEMMLowpOffsetContributionOutputStageKernel::run(), CLIm2ColKernel::run(), CLGEMMLowpMatrixAReductionKernel::run(), CLGEMMReshapeRHSMatrixKernel::run(), CLGEMMLowpMatrixMultiplyReshapedOnlyRHSKernel::run(), CLGEMMMatrixMultiplyReshapedOnlyRHSKernel::run(), CLGEMMLowpMatrixBReductionKernel::run(), ClElementwiseKernel::run_op(), ClTransposeKernel::run_op(), ClFloorKernel::run_op(), ClReshapeKernel::run_op(), ClCopyKernel::run_op(), ClElementWiseUnaryKernel::run_op(), ClDequantizationKernel::run_op(), ClWidthConcatenate2TensorsKernel::run_op(), ClHeightConcatenateKernel::run_op(), ClWidthConcatenateKernel::run_op(), ClActivationKernel::run_op(), ClPoolingKernel::run_op(), ClQuantizationKernel::run_op(), ClBatchConcatenateKernel::run_op(), ClDepthConcatenateKernel::run_op(), ClFillKernel::run_op(), ClWidthConcatenate4TensorsKernel::run_op(), ClPermuteKernel::run_op(), ClConvertFullyConnectedWeightsKernel::run_op(), ClScaleKernel::run_op(), CLStridedSliceKernel::run_op(), ClMulKernel::run_op(), ClDirectConvolutionKernel::run_op(), ClCropKernel::run_op(), CLFillBorderKernel::run_op(), ClLogits1DMaxShiftExpSumKernel::run_op(), ClComplexMulKernel::run_op(), ClLogits1DNormKernel::run_op(), and ICLKernel::set_lws_hint().

◆ num_arguments_per_1D_array()

static constexpr unsigned int num_arguments_per_1D_array ( )
inlinestatic

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

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

Definition at line 190 of file ICLKernel.h.

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.

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

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

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

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

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

◆ 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 CLGEMMLowpMatrixBReductionKernel, CLGEMMMatrixMultiplyReshapedKernel, CLGEMMMatrixMultiplyReshapedOnlyRHSKernel, CLGEMMLowpMatrixMultiplyReshapedOnlyRHSKernel, CLGEMMReshapeRHSMatrixKernel, CLGEMMLowpMatrixAReductionKernel, CLIm2ColKernel, CLGEMMLowpOffsetContributionOutputStageKernel, CLComputeMeanVariance, CLDepthwiseConvolutionLayerNativeKernel, CLWinogradOutputTransformKernel, CLSpaceToBatchLayerKernel, CLWeightsReshapeKernel, CLDepthwiseConvolutionLayer3x3NCHWKernel, CLGEMMMatrixMultiplyNativeKernel, CLGEMMLowpMatrixMultiplyReshapedKernel, CLFuseBatchNormalizationKernel, CLGEMMMatrixMultiplyKernel, CLWinogradFilterTransformKernel, CLBatchNormalizationLayerKernel, CLGEMMLowpOffsetContributionKernel, CLWinogradInputTransformKernel, CLBatchToSpaceLayerKernel, CLDeconvolutionReshapeOutputKernel, CLDepthwiseConvolutionLayer3x3NHWCKernel, CLROIAlignLayerKernel, CLCol2ImKernel, CLGEMMReshapeLHSMatrixKernel, CLArgMinMaxLayerKernel, CLGEMMLowpMatrixMultiplyNativeKernel, CLStackLayerKernel, CLGEMMLowpQuantizeDownInt32ScaleByFloatKernel, CLGEMMLowpQuantizeDownInt32ScaleKernel, CLBoundingBoxTransformKernel, CLL2NormalizeLayerKernel, CLFFTRadixStageKernel, CLFillBorderKernel, CLPadLayerKernel, CLPriorBoxLayerKernel, CLNormalizePlanarYUVLayerKernel, CLReductionOperationKernel, CLRangeKernel, CLReorgLayerKernel, CLMeanStdDevNormalizationKernel, CLFFTDigitReverseKernel, CLGEMMLowpQuantizeDownInt32ScaleByFixedPointKernel, CLTileKernel, CLComparisonKernel, CLGatherKernel, CLMinMaxLayerKernel, CLNormalizationLayerKernel, CLROIPoolingLayerKernel, CLQLSTMLayerNormalizationKernel, CLFFTScaleKernel, CLComputeAllAnchorsKernel, CLMaxUnpoolingLayerKernel, CLDeconvolutionLayerUpsampleKernel, CLDepthToSpaceLayerKernel, CLReverseKernel, CLSelectKernel, CLSpaceToDepthLayerKernel, CLChannelShuffleLayerKernel, CLInstanceNormalizationLayerKernel, CLRemapKernel, CLBitwiseKernel, ICLSimple3DKernel, and ICLSimple2DKernel.

Definition at line 233 of file ICLKernel.h.

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

References ARM_COMPUTE_UNUSED, and IKernel::window().

◆ 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 ClComplexMulKernel, CLFillBorderKernel, ClCropKernel, ClDirectConvolutionKernel, ClMulKernel, CLStridedSliceKernel, ClScaleKernel, ClPermuteKernel, ClFillKernel, ClQuantizationKernel, ClPoolingKernel, ClDequantizationKernel, ClCopyKernel, ClElementWiseUnaryKernel, ClFloorKernel, ClReshapeKernel, and ClTransposeKernel.

Definition at line 245 of file ICLKernel.h.

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

References ARM_COMPUTE_UNUSED, and IKernel::window().

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

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:915

References ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, ICLKernel::lws_hint(), and CLTuningParams::set_lws().

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

◆ set_target() [1/2]

void set_target ( GPUTarget  target)
inline

Set the targeted GPU architecture.

Parameters
[in]targetThe targeted GPU architecture

Definition at line 321 of file ICLKernel.h.

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.

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

References arm_compute::get_target_from_device().

◆ set_wbsm_hint()

void set_wbsm_hint ( const cl_int &  wbsm_hint)
inline

Set the workgroup batch size modifier hint.

Note
This method should be called after the configuration of the kernel
Parameters
[in]wbsm_hintworkgroup batch size modifier value

Definition at line 287 of file ICLKernel.h.

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:915

References ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL, CLTuningParams::set_wbsm(), and ICLKernel::wbsm_hint().

Referenced by CLTuner::tune_kernel_dynamic().

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

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

References CLTuningParams::get_wbsm().

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


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