Compute Library
 22.08
ICLKernel.h
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2016-2022 Arm Limited.
3  *
4  * SPDX-License-Identifier: MIT
5  *
6  * Permission is hereby granted, free of charge, to any person obtaining a copy
7  * of this software and associated documentation files (the "Software"), to
8  * deal in the Software without restriction, including without limitation the
9  * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
10  * sell copies of the Software, and to permit persons to whom the Software is
11  * furnished to do so, subject to the following conditions:
12  *
13  * The above copyright notice and this permission notice shall be included in all
14  * copies or substantial portions of the Software.
15  *
16  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
19  * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
21  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
22  * SOFTWARE.
23  */
24 #ifndef ARM_COMPUTE_ICLKERNEL_H
25 #define ARM_COMPUTE_ICLKERNEL_H
26 
35 
37 
38 #include <string>
39 
40 #if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
41 namespace arm_compute
42 {
43 namespace experimental
44 {
45 namespace dynamic_fusion
46 {
47 struct TensorBinding;
48 struct ClExecutionDescriptor;
49 } // namespace dynamic_fusion
50 } // namespace experimental
51 } // namespace arm_compute
52 #endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
53 
54 namespace arm_compute
55 {
56 namespace
57 {
58 bool is_same_lws(cl::NDRange lws0, cl::NDRange lws1)
59 {
60  if(lws0.dimensions() != lws1.dimensions())
61  {
62  return false;
63  }
64 
65  for(size_t i = 0; i < lws0.dimensions(); ++i)
66  {
67  if(lws0.get()[i] != lws1.get()[i])
68  {
69  return false;
70  }
71  }
72 
73  return true;
74 }
75 } // namespace
76 template <typename T>
77 class ICLArray;
78 class ICLTensor;
79 class Window;
80 /** Common interface for all the OpenCL kernels */
81 class ICLKernel : public IKernel
82 {
83 private:
84  /** Returns the number of arguments enqueued per array object.
85  *
86  * @return The number of arguments enqueued per array object.
87  */
88  template <unsigned int dimension_size>
89  constexpr static unsigned int num_arguments_per_array()
90  {
91  return num_arguments_per_tensor<dimension_size>();
92  }
93  /** Returns the number of arguments enqueued per tensor object.
94  *
95  * @return The number of arguments enqueued per tensor object.
96  */
97  template <unsigned int dimension_size>
98  constexpr static unsigned int num_arguments_per_tensor()
99  {
100  return 2 + 2 * dimension_size;
101  }
102 
103  cl::NDRange default_lws_tune(const Window &window)
104  {
105  return get_default_lws_for_type(_type, gws_from_window(window));
106  }
107 
108  using IKernel::configure; //Prevent children from calling IKernel::configure() directly
109 protected:
110  /** Configure the kernel's window and local workgroup size hint.
111  *
112  * @param[in] window The maximum window which will be returned by window()
113  * @param[in] lws_hint Local-Workgroup-Size to use.
114  * @param[in] wbsm_hint (Optional) Workgroup-Batch-Size-Modifier to use.
115  */
116  void configure_internal(const Window &window, cl::NDRange lws_hint, cl_int wbsm_hint = 0)
117  {
118  configure_internal(window, CLTuningParams(lws_hint, wbsm_hint));
119  }
120 
121  /** Configure the kernel's window and tuning parameters hints.
122  *
123  * @param[in] window The maximum window which will be returned by window()
124  * @param[in] tuning_params_hint (Optional) Tuning parameters to use.
125  */
126  void configure_internal(const Window &window, CLTuningParams tuning_params_hint = CLTuningParams(CLKernelLibrary::get().default_ndrange(), 0))
127  {
128  _tuning_params_hint = tuning_params_hint;
129 
130  if(is_same_lws(_tuning_params_hint.get_lws(), CLKernelLibrary::get().default_ndrange()))
131  {
132  _tuning_params_hint.set_lws(default_lws_tune(window));
133  }
134 
135  IKernel::configure(window);
136  }
137 
138 public:
139  /** Constructor */
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  }
144  /** Returns a reference to the OpenCL kernel of this object.
145  *
146  * @return A reference to the OpenCL kernel of this object.
147  */
148  cl::Kernel &kernel()
149  {
150  return _kernel;
151  }
152  /** Returns the CL kernel type
153  *
154  * @return The CL kernel type
155  */
157  {
158  return _type;
159  }
160  /** Add the passed 1D array's parameters to the object's kernel's arguments starting from the index idx.
161  *
162  * @param[in,out] idx Index at which to start adding the array's arguments. Will be incremented by the number of kernel arguments set.
163  * @param[in] array Array to set as an argument of the object's kernel.
164  * @param[in] strides @ref Strides object containing stride of each dimension in bytes.
165  * @param[in] num_dimensions Number of dimensions of the @p array.
166  * @param[in] window Window the kernel will be executed on.
167  */
168  template <typename T>
169  void add_1D_array_argument(unsigned int &idx, const ICLArray<T> *array, const Strides &strides, unsigned int num_dimensions, const Window &window)
170  {
171  add_array_argument<T, 1>(idx, array, strides, num_dimensions, window);
172  }
173  /** Add the passed 1D tensor's parameters to the object's kernel's arguments starting from the index idx.
174  *
175  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
176  * @param[in] tensor Tensor to set as an argument of the object's kernel.
177  * @param[in] window Window the kernel will be executed on.
178  */
179  void add_1D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window)
180  {
181  add_tensor_argument<1>(idx, tensor, window);
182  }
183  /** Add the passed 1D tensor's parameters to the object's kernel's arguments starting from the index idx if the condition is true.
184  *
185  * @param[in] cond Condition to check
186  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
187  * @param[in] tensor Tensor to set as an argument of the object's kernel.
188  * @param[in] window Window the kernel will be executed on.
189  */
190  void add_1D_tensor_argument_if(bool cond, unsigned int &idx, const ICLTensor *tensor, const Window &window)
191  {
192  if(cond)
193  {
194  add_1D_tensor_argument(idx, tensor, window);
195  }
196  }
197  /** Add the passed 2D tensor's parameters to the object's kernel's arguments starting from the index idx.
198  *
199  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
200  * @param[in] tensor Tensor to set as an argument of the object's kernel.
201  * @param[in] window Window the kernel will be executed on.
202  */
203  void add_2D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window)
204  {
205  add_tensor_argument<2>(idx, tensor, window);
206  }
207  /** Add the passed 2D tensor's parameters to the object's kernel's arguments starting from the index idx if the condition is true.
208  *
209  * @param[in] cond Condition to check
210  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
211  * @param[in] tensor Tensor to set as an argument of the object's kernel.
212  * @param[in] window Window the kernel will be executed on.
213  */
214  void add_2D_tensor_argument_if(bool cond, unsigned int &idx, const ICLTensor *tensor, const Window &window)
215  {
216  if(cond)
217  {
218  add_2D_tensor_argument(idx, tensor, window);
219  }
220  }
221  /** Add the passed 3D tensor's parameters to the object's kernel's arguments starting from the index idx.
222  *
223  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
224  * @param[in] tensor Tensor to set as an argument of the object's kernel.
225  * @param[in] window Window the kernel will be executed on.
226  */
227  void add_3D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window)
228  {
229  add_tensor_argument<3>(idx, tensor, window);
230  }
231  /** Add the passed 4D tensor's parameters to the object's kernel's arguments starting from the index idx.
232  *
233  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
234  * @param[in] tensor Tensor to set as an argument of the object's kernel.
235  * @param[in] window Window the kernel will be executed on.
236  */
237  void add_4D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window)
238  {
239  add_tensor_argument<4>(idx, tensor, window);
240  }
241  /** Add the passed 5D tensor's parameters to the object's kernel's arguments starting from the index idx.
242  *
243  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
244  * @param[in] tensor Tensor to set as an argument of the object's kernel.
245  * @param[in] window Window the kernel will be executed on.
246  */
247  void add_5D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window)
248  {
249  add_tensor_argument<5>(idx, tensor, window);
250  }
251 
252  /** 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.
253  *
254  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
255  * @param[in] tensor Tensor to set as an argument of the object's kernel.
256  */
257  void add_3d_tensor_nhw_argument(unsigned int &idx, const ICLTensor *tensor);
258 
259  /** Returns the number of arguments enqueued per NHW 3D Tensor object.
260  *
261  * @return The number of arguments enqueued per NHW 3D Tensor object.
262  */
263  constexpr static unsigned int num_arguments_per_3d_tensor_nhw()
264  {
265  constexpr unsigned int no_args_per_3d_tensor_nhw = 7u;
266  return no_args_per_3d_tensor_nhw;
267  }
268 
269  /** 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.
270  *
271  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
272  * @param[in] tensor Tensor to set as an argument of the object's kernel.
273  */
274  void add_4d_tensor_nhwc_argument(unsigned int &idx, const ICLTensor *tensor);
275 
276  /** Returns the number of arguments enqueued per NHWC 4D Tensor object.
277  *
278  * @return The number of arguments enqueued per NHWC 4D Tensor object.
279  */
280  constexpr static unsigned int num_arguments_per_4d_tensor_nhwc()
281  {
282  constexpr unsigned int no_args_per_4d_tensor_nhwc = 9u;
283  return no_args_per_4d_tensor_nhwc;
284  }
285 
286  /** Returns the number of arguments enqueued per 1D array object.
287  *
288  * @return The number of arguments enqueues per 1D array object.
289  */
290  constexpr static unsigned int num_arguments_per_1D_array()
291  {
292  return num_arguments_per_array<1>();
293  }
294  /** Returns the number of arguments enqueued per 1D tensor object.
295  *
296  * @return The number of arguments enqueues per 1D tensor object.
297  */
298  constexpr static unsigned int num_arguments_per_1D_tensor()
299  {
300  return num_arguments_per_tensor<1>();
301  }
302  /** Returns the number of arguments enqueued per 2D tensor object.
303  *
304  * @return The number of arguments enqueues per 2D tensor object.
305  */
306  constexpr static unsigned int num_arguments_per_2D_tensor()
307  {
308  return num_arguments_per_tensor<2>();
309  }
310  /** Returns the number of arguments enqueued per 3D tensor object.
311  *
312  * @return The number of arguments enqueues per 3D tensor object.
313  */
314  constexpr static unsigned int num_arguments_per_3D_tensor()
315  {
316  return num_arguments_per_tensor<3>();
317  }
318  /** Returns the number of arguments enqueued per 4D tensor object.
319  *
320  * @return The number of arguments enqueues per 4D tensor object.
321  */
322  constexpr static unsigned int num_arguments_per_4D_tensor()
323  {
324  return num_arguments_per_tensor<4>();
325  }
326  /** Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue.
327  *
328  * @note The queue is *not* flushed by this method, and therefore the kernel will not have been executed by the time this method returns.
329  *
330  * @param[in] window Region on which to execute the kernel. (Must be a valid region of the window returned by window()).
331  * @param[in,out] queue Command queue on which to enqueue the kernel.
332  */
333  virtual void run(const Window &window, cl::CommandQueue &queue)
334  {
335  ARM_COMPUTE_UNUSED(window, queue);
336  }
337  /** Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue.
338  *
339  * @note The queue is *not* flushed by this method, and therefore the kernel will not have been executed by the time this method returns.
340  *
341  * @param[in] tensors A vector containing the tensors to operato on.
342  * @param[in] window Region on which to execute the kernel. (Must be a valid region of the window returned by window()).
343  * @param[in,out] queue Command queue on which to enqueue the kernel.
344  */
345  virtual void run_op(ITensorPack &tensors, const Window &window, cl::CommandQueue &queue)
346  {
347  ARM_COMPUTE_UNUSED(tensors, window, queue);
348  }
349 
350 #if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
351  /// 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
352  virtual void run_composite_op(ITensorPack &tensors, const Window &window, cl::CommandQueue &queue, const experimental::dynamic_fusion::ClExecutionDescriptor &exec_desc)
353  {
354  ARM_COMPUTE_UNUSED(tensors, window, queue, exec_desc);
355  }
356 #endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
357  /** Add the passed parameters to the object's kernel's arguments starting from the index idx.
358  *
359  * @param[in,out] idx Index at which to start adding the arguments. Will be incremented by the number of kernel arguments set.
360  * @param[in] value Value to set as an argument of the object's kernel.
361  */
362  template <typename T>
363  void add_argument(unsigned int &idx, T value)
364  {
365  _kernel.setArg(idx++, value);
366  }
367 
368  /** Set the Local-Workgroup-Size hint
369  *
370  * @note This method should be called after the configuration of the kernel
371  *
372  * @param[in] lws_hint Local-Workgroup-Size to use
373  */
374  void set_lws_hint(const cl::NDRange &lws_hint)
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  }
379 
380  /** Return the Local-Workgroup-Size hint
381  *
382  * @return Current lws hint
383  */
384  cl::NDRange lws_hint() const
385  {
386  return _tuning_params_hint.get_lws();
387  }
388 
389  /** Set the workgroup batch size modifier hint
390  *
391  * @note This method should be called after the configuration of the kernel
392  *
393  * @param[in] wbsm_hint workgroup batch size modifier value
394  */
395  void set_wbsm_hint(const cl_int &wbsm_hint)
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  }
400 
401  /** Return the workgroup batch size modifier hint
402  *
403  * @return Current wbsm hint
404  */
405  cl_int wbsm_hint() const
406  {
407  return _tuning_params_hint.get_wbsm();
408  }
409 
410  /** Get the configuration ID
411  *
412  * @note The configuration ID can be used by the caller to distinguish different calls of the same OpenCL kernel
413  * In particular, this method can be used by CLScheduler to keep track of the best LWS for each configuration of the same kernel.
414  * The configuration ID should be provided only for the kernels potentially affected by the LWS geometry
415  *
416  * @note This method should be called after the configuration of the kernel
417  *
418  * @return configuration id string
419  */
420  const std::string &config_id() const
421  {
422  return _config_id;
423  }
424 
425  /** Set the targeted GPU architecture
426  *
427  * @param[in] target The targeted GPU architecture
428  */
429  void set_target(GPUTarget target)
430  {
431  _target = target;
432  }
433 
434  /** Set the targeted GPU architecture according to the CL device
435  *
436  * @param[in] device A CL device
437  */
438  void set_target(cl::Device &device);
439 
440  /** Get the targeted GPU architecture
441  *
442  * @return The targeted GPU architecture.
443  */
445  {
446  return _target;
447  }
448 
449  /** Get the maximum workgroup size for the device the CLKernelLibrary uses.
450  *
451  * @return The maximum workgroup size value.
452  */
453  size_t get_max_workgroup_size();
454  /** Get the global work size given an execution window
455  *
456  * @param[in] window Execution window
457  *
458  * @return Global work size of the given execution window
459  */
460  static cl::NDRange gws_from_window(const Window &window);
461 
462 private:
463  /** Add the passed array's parameters to the object's kernel's arguments starting from the index idx.
464  *
465  * @param[in,out] idx Index at which to start adding the array's arguments. Will be incremented by the number of kernel arguments set.
466  * @param[in] array Array to set as an argument of the object's kernel.
467  * @param[in] strides @ref Strides object containing stride of each dimension in bytes.
468  * @param[in] num_dimensions Number of dimensions of the @p array.
469  * @param[in] window Window the kernel will be executed on.
470  */
471  template <typename T, unsigned int dimension_size>
472  void add_array_argument(unsigned int &idx, const ICLArray<T> *array, const Strides &strides, unsigned int num_dimensions, const Window &window);
473  /** Add the passed tensor's parameters to the object's kernel's arguments starting from the index idx.
474  *
475  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
476  * @param[in] tensor Tensor to set as an argument of the object's kernel.
477  * @param[in] window Window the kernel will be executed on.
478  */
479  template <unsigned int dimension_size>
480  void add_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window);
481 
482 protected:
483  cl::Kernel _kernel; /**< OpenCL kernel to run */
484  GPUTarget _target; /**< The targeted GPU */
485  std::string _config_id; /**< Configuration ID */
486  size_t _max_workgroup_size; /**< The maximum workgroup size for this kernel */
487  CLKernelType _type; /**< The CL kernel type */
488 private:
489  CLTuningParams _tuning_params_hint; /**< Tuning parameters hint for the OpenCL kernel */
490 };
491 
492 /** Add the kernel to the command queue with the given window.
493  *
494  * @note Depending on the size of the window, this might translate into several jobs being enqueued.
495  *
496  * @note If kernel->kernel() is empty then the function will return without adding anything to the queue.
497  *
498  * @param[in,out] queue OpenCL command queue.
499  * @param[in] kernel Kernel to enqueue
500  * @param[in] window Window the kernel has to process.
501  * @param[in] lws_hint (Optional) Local workgroup size requested. Default is based on the device target.
502  * @param[in] use_dummy_work_items (Optional) Use dummy work items in order to have two dimensional power of two NDRange. Default is false
503  * Note: it is kernel responsibility to check if the work-item is out-of-range
504  *
505  * @note If any dimension of the lws is greater than the global workgroup size then no lws will be passed.
506  */
507 void enqueue(cl::CommandQueue &queue, ICLKernel &kernel, const Window &window, const cl::NDRange &lws_hint = CLKernelLibrary::get().default_ndrange(), bool use_dummy_work_items = false);
508 
509 /** Add the passed array's parameters to the object's kernel's arguments starting from the index idx.
510  *
511  * @param[in,out] idx Index at which to start adding the array's arguments. Will be incremented by the number of kernel arguments set.
512  * @param[in] array Array to set as an argument of the object's kernel.
513  * @param[in] strides @ref Strides object containing stride of each dimension in bytes.
514  * @param[in] num_dimensions Number of dimensions of the @p array.
515  * @param[in] window Window the kernel will be executed on.
516  */
517 template <typename T, unsigned int dimension_size>
518 void ICLKernel::add_array_argument(unsigned &idx, const ICLArray<T> *array, const Strides &strides, unsigned int num_dimensions, const Window &window)
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 }
545 }
546 #endif /*ARM_COMPUTE_ICLKERNEL_H */
static constexpr unsigned int num_arguments_per_1D_tensor()
Returns the number of arguments enqueued per 1D tensor object.
Definition: ICLKernel.h:298
Common information for all the kernels.
Definition: IKernel.h:33
void add_1D_tensor_argument_if(bool cond, 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:190
static constexpr unsigned int num_arguments_per_1D_array()
Returns the number of arguments enqueued per 1D array object.
Definition: ICLKernel.h:290
virtual void run(const Window &window, cl::CommandQueue &queue)
Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue...
Definition: ICLKernel.h:333
void add_2D_tensor_argument_if(bool cond, 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:214
Unknown CL kernel type.
Definition: CLTypes.h:82
void enqueue(cl::CommandQueue &queue, ICLKernel &kernel, const Window &window, const cl::NDRange &lws_hint=CLKernelLibrary::get().default_ndrange(), bool use_dummy_work_items=false)
Add the kernel to the command queue with the given window.
Definition: ICLKernel.cpp:32
static constexpr unsigned int num_arguments_per_3d_tensor_nhw()
Returns the number of arguments enqueued per NHW 3D Tensor object.
Definition: ICLKernel.h:263
cl::Kernel & kernel()
Returns a reference to the OpenCL kernel of this object.
Definition: ICLKernel.h:148
cl::NDRange lws_hint() const
Return the Local-Workgroup-Size hint.
Definition: ICLKernel.h:384
void set_lws_hint(const cl::NDRange &lws_hint)
Set the Local-Workgroup-Size hint.
Definition: ICLKernel.h:374
cl_int wbsm_hint() const
Return the workgroup batch size modifier hint.
Definition: ICLKernel.h:405
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 inc...
Definition: ICLKernel.h:352
void add_argument(unsigned int &idx, T value)
Add the passed parameters to the object&#39;s kernel&#39;s arguments starting from the index idx...
Definition: ICLKernel.h:363
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&#39;s parameters to the object&#39;s kernel&#39;s arguments starting from the index idx...
Definition: ICLKernel.h:169
cl::NDRange get_default_lws_for_type(CLKernelType kernel_type, cl::NDRange gws)
#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
static CLKernelLibrary & get()
Access the KernelLibrary singleton.
< OpenCL tuner parameters
#define ARM_COMPUTE_ERROR_ON_MSG_VAR(cond, msg,...)
Definition: Error.h:457
Manages all the OpenCL kernels compilation and caching, provides accessors for the OpenCL Context...
Common interface for all the OpenCL kernels.
Definition: ICLKernel.h:81
void add_3D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window)
Add the passed 3D tensor&#39;s parameters to the object&#39;s kernel&#39;s arguments starting from the index idx...
Definition: ICLKernel.h:227
Copyright (c) 2017-2022 Arm Limited.
const std::string & config_id() const
Get the configuration ID.
Definition: ICLKernel.h:420
cl::NDRange default_ndrange() const
Return the default NDRange for the device.
Descriptor containing information required to run a single ClWorkload.
Definition: ClWorkload.h:91
Interface for OpenCL Array.
Definition: ICLArray.h:35
static constexpr unsigned int num_arguments_per_3D_tensor()
Returns the number of arguments enqueued per 3D tensor object.
Definition: ICLKernel.h:314
#define ARM_COMPUTE_UNUSED(...)
To avoid unused variables warnings.
Definition: Error.h:152
GPUTarget get_target() const
Get the targeted GPU architecture.
Definition: ICLKernel.h:444
static constexpr unsigned int num_arguments_per_2D_tensor()
Returns the number of arguments enqueued per 2D tensor object.
Definition: ICLKernel.h:306
static constexpr unsigned int num_arguments_per_4D_tensor()
Returns the number of arguments enqueued per 4D tensor object.
Definition: ICLKernel.h:322
#define ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(k)
Definition: Validate.h:915
Strides of an item in bytes.
Definition: Strides.h:37
static constexpr unsigned int num_arguments_per_4d_tensor_nhwc()
Returns the number of arguments enqueued per NHWC 4D Tensor object.
Definition: ICLKernel.h:280
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
Interface for OpenCL tensor.
Definition: ICLTensor.h:42
ICLKernel()
Constructor.
Definition: ICLKernel.h:140
GPUTarget
Available GPU Targets.
Definition: GPUTarget.h:34
Wrapper to configure the Khronos OpenCL C++ header.
Tensor packing service.
Definition: ITensorPack.h:39
void set_target(GPUTarget target)
Set the targeted GPU architecture.
Definition: ICLKernel.h:429
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
void add_5D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window)
Add the passed 5D tensor&#39;s parameters to the object&#39;s kernel&#39;s arguments starting from the index idx...
Definition: ICLKernel.h:247
void add_4D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window)
Add the passed 4D tensor&#39;s parameters to the object&#39;s kernel&#39;s arguments starting from the index idx...
Definition: ICLKernel.h:237
CLKernelType type() const
Returns the CL kernel type.
Definition: ICLKernel.h:156
virtual const cl::Buffer & cl_buffer() const =0
Interface to be implemented by the child class to return a reference to the OpenCL buffer containing ...
void set_wbsm_hint(const cl_int &wbsm_hint)
Set the workgroup batch size modifier hint.
Definition: ICLKernel.h:395
Describe a multidimensional execution window.
Definition: Window.h:39
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...
Definition: ICLKernel.h:345