Compute Library
 20.08
ICLKernel.h
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2016-2020 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 
33 
34 #include <string>
35 
36 namespace arm_compute
37 {
38 template <typename T>
39 class ICLArray;
40 class ICLTensor;
41 class Window;
42 
43 /** Common interface for all the OpenCL kernels */
44 class ICLKernel : public IKernel
45 {
46 private:
47  /** Returns the number of arguments enqueued per array object.
48  *
49  * @return The number of arguments enqueued per array object.
50  */
51  template <unsigned int dimension_size>
52  constexpr static unsigned int num_arguments_per_array()
53  {
54  return num_arguments_per_tensor<dimension_size>();
55  }
56  /** Returns the number of arguments enqueued per tensor object.
57  *
58  * @return The number of arguments enqueued per tensor object.
59  */
60  template <unsigned int dimension_size>
61  constexpr static unsigned int num_arguments_per_tensor()
62  {
63  return 2 + 2 * dimension_size;
64  }
65  using IKernel::configure; //Prevent children from calling IKernel::configure() directly
66 protected:
67  /** Configure the kernel's window and local workgroup size hint.
68  *
69  * @param[in] window The maximum window which will be returned by window()
70  * @param[in] lws_hint (Optional) Local-Workgroup-Size to use.
71  */
72  void configure_internal(const Window &window, cl::NDRange lws_hint = CLKernelLibrary::get().default_ndrange())
73  {
74  _lws_hint = lws_hint;
75  IKernel::configure(window);
76  }
77 
78 public:
79  /** Constructor */
81  : _kernel(nullptr), _target(GPUTarget::MIDGARD), _config_id(arm_compute::default_config_id), _max_workgroup_size(0), _lws_hint()
82  {
83  }
84  /** Returns a reference to the OpenCL kernel of this object.
85  *
86  * @return A reference to the OpenCL kernel of this object.
87  */
88  cl::Kernel &kernel()
89  {
90  return _kernel;
91  }
92  /** Add the passed 1D array's parameters to the object's kernel's arguments starting from the index idx.
93  *
94  * @param[in,out] idx Index at which to start adding the array's arguments. Will be incremented by the number of kernel arguments set.
95  * @param[in] array Array to set as an argument of the object's kernel.
96  * @param[in] strides @ref Strides object containing stride of each dimension in bytes.
97  * @param[in] num_dimensions Number of dimensions of the @p array.
98  * @param[in] window Window the kernel will be executed on.
99  */
100  template <typename T>
101  void add_1D_array_argument(unsigned int &idx, const ICLArray<T> *array, const Strides &strides, unsigned int num_dimensions, const Window &window)
102  {
103  add_array_argument<T, 1>(idx, array, strides, num_dimensions, window);
104  }
105  /** Add the passed 1D tensor's parameters to the object's kernel's arguments starting from the index idx.
106  *
107  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
108  * @param[in] tensor Tensor to set as an argument of the object's kernel.
109  * @param[in] window Window the kernel will be executed on.
110  */
111  void add_1D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window)
112  {
113  add_tensor_argument<1>(idx, tensor, window);
114  }
115  /** Add the passed 1D tensor's parameters to the object's kernel's arguments starting from the index idx if the condition is true.
116  *
117  * @param[in] cond Condition to check
118  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
119  * @param[in] tensor Tensor to set as an argument of the object's kernel.
120  * @param[in] window Window the kernel will be executed on.
121  */
122  void add_1D_tensor_argument_if(bool cond, unsigned int &idx, const ICLTensor *tensor, const Window &window)
123  {
124  if(cond)
125  {
126  add_1D_tensor_argument(idx, tensor, window);
127  }
128  }
129  /** Add the passed 2D tensor's parameters to the object's kernel's arguments starting from the index idx.
130  *
131  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
132  * @param[in] tensor Tensor to set as an argument of the object's kernel.
133  * @param[in] window Window the kernel will be executed on.
134  */
135  void add_2D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window)
136  {
137  add_tensor_argument<2>(idx, tensor, window);
138  }
139  /** Add the passed 2D tensor's parameters to the object's kernel's arguments starting from the index idx if the condition is true.
140  *
141  * @param[in] cond Condition to check
142  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
143  * @param[in] tensor Tensor to set as an argument of the object's kernel.
144  * @param[in] window Window the kernel will be executed on.
145  */
146  void add_2D_tensor_argument_if(bool cond, unsigned int &idx, const ICLTensor *tensor, const Window &window)
147  {
148  if(cond)
149  {
150  add_2D_tensor_argument(idx, tensor, window);
151  }
152  }
153  /** Add the passed 3D tensor's parameters to the object's kernel's arguments starting from the index idx.
154  *
155  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
156  * @param[in] tensor Tensor to set as an argument of the object's kernel.
157  * @param[in] window Window the kernel will be executed on.
158  */
159  void add_3D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window)
160  {
161  add_tensor_argument<3>(idx, tensor, window);
162  }
163  /** Add the passed 4D tensor's parameters to the object's kernel's arguments starting from the index idx.
164  *
165  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
166  * @param[in] tensor Tensor to set as an argument of the object's kernel.
167  * @param[in] window Window the kernel will be executed on.
168  */
169  void add_4D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window)
170  {
171  add_tensor_argument<4>(idx, tensor, window);
172  }
173  /** Returns the number of arguments enqueued per 1D array object.
174  *
175  * @return The number of arguments enqueues per 1D array object.
176  */
177  constexpr static unsigned int num_arguments_per_1D_array()
178  {
179  return num_arguments_per_array<1>();
180  }
181  /** Returns the number of arguments enqueued per 1D tensor object.
182  *
183  * @return The number of arguments enqueues per 1D tensor object.
184  */
185  constexpr static unsigned int num_arguments_per_1D_tensor()
186  {
187  return num_arguments_per_tensor<1>();
188  }
189  /** Returns the number of arguments enqueued per 2D tensor object.
190  *
191  * @return The number of arguments enqueues per 2D tensor object.
192  */
193  constexpr static unsigned int num_arguments_per_2D_tensor()
194  {
195  return num_arguments_per_tensor<2>();
196  }
197  /** Returns the number of arguments enqueued per 3D tensor object.
198  *
199  * @return The number of arguments enqueues per 3D tensor object.
200  */
201  constexpr static unsigned int num_arguments_per_3D_tensor()
202  {
203  return num_arguments_per_tensor<3>();
204  }
205  /** Returns the number of arguments enqueued per 4D tensor object.
206  *
207  * @return The number of arguments enqueues per 4D tensor object.
208  */
209  constexpr static unsigned int num_arguments_per_4D_tensor()
210  {
211  return num_arguments_per_tensor<4>();
212  }
213  /** Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue.
214  *
215  * @note The queue is *not* flushed by this method, and therefore the kernel will not have been executed by the time this method returns.
216  *
217  * @param[in] window Region on which to execute the kernel. (Must be a valid region of the window returned by window()).
218  * @param[in,out] queue Command queue on which to enqueue the kernel.
219  */
220  virtual void run(const Window &window, cl::CommandQueue &queue)
221  {
222  ARM_COMPUTE_UNUSED(window, queue);
223  }
224  /** Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue.
225  *
226  * @note The queue is *not* flushed by this method, and therefore the kernel will not have been executed by the time this method returns.
227  *
228  * @param[in] tensors A vector containing the tensors to operato on.
229  * @param[in] window Region on which to execute the kernel. (Must be a valid region of the window returned by window()).
230  * @param[in,out] queue Command queue on which to enqueue the kernel.
231  */
232  virtual void run_op(ITensorPack &tensors, const Window &window, cl::CommandQueue &queue)
233  {
234  ARM_COMPUTE_UNUSED(tensors, window, queue);
235  }
236  /** Add the passed parameters to the object's kernel's arguments starting from the index idx.
237  *
238  * @param[in,out] idx Index at which to start adding the arguments. Will be incremented by the number of kernel arguments set.
239  * @param[in] value Value to set as an argument of the object's kernel.
240  */
241  template <typename T>
242  void add_argument(unsigned int &idx, T value)
243  {
244  _kernel.setArg(idx++, value);
245  }
246 
247  /** Set the Local-Workgroup-Size hint
248  *
249  * @note This method should be called after the configuration of the kernel
250  *
251  * @param[in] lws_hint Local-Workgroup-Size to use
252  */
253  void set_lws_hint(const cl::NDRange &lws_hint)
254  {
255  ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); // lws_hint will be overwritten by configure()
256  _lws_hint = lws_hint;
257  }
258 
259  /** Return the Local-Workgroup-Size hint
260  *
261  * @return Current lws hint
262  */
263  cl::NDRange lws_hint() const
264  {
265  return _lws_hint;
266  }
267 
268  /** Get the configuration ID
269  *
270  * @note The configuration ID can be used by the caller to distinguish different calls of the same OpenCL kernel
271  * In particular, this method can be used by CLScheduler to keep track of the best LWS for each configuration of the same kernel.
272  * The configuration ID should be provided only for the kernels potentially affected by the LWS geometry
273  *
274  * @note This method should be called after the configuration of the kernel
275  *
276  * @return configuration id string
277  */
278  const std::string &config_id() const
279  {
280  return _config_id;
281  }
282 
283  /** Set the targeted GPU architecture
284  *
285  * @param[in] target The targeted GPU architecture
286  */
287  void set_target(GPUTarget target)
288  {
289  _target = target;
290  }
291 
292  /** Set the targeted GPU architecture according to the CL device
293  *
294  * @param[in] device A CL device
295  */
296  void set_target(cl::Device &device);
297 
298  /** Get the targeted GPU architecture
299  *
300  * @return The targeted GPU architecture.
301  */
303  {
304  return _target;
305  }
306 
307  /** Get the maximum workgroup size for the device the CLKernelLibrary uses.
308  *
309  * @return The maximum workgroup size value.
310  */
311  size_t get_max_workgroup_size();
312  /** Get the global work size given an execution window
313  *
314  * @param[in] window Execution window
315  *
316  * @return Global work size of the given execution window
317  */
318  static cl::NDRange gws_from_window(const Window &window);
319 
320 private:
321  /** Add the passed array's parameters to the object's kernel's arguments starting from the index idx.
322  *
323  * @param[in,out] idx Index at which to start adding the array's arguments. Will be incremented by the number of kernel arguments set.
324  * @param[in] array Array to set as an argument of the object's kernel.
325  * @param[in] strides @ref Strides object containing stride of each dimension in bytes.
326  * @param[in] num_dimensions Number of dimensions of the @p array.
327  * @param[in] window Window the kernel will be executed on.
328  */
329  template <typename T, unsigned int dimension_size>
330  void add_array_argument(unsigned int &idx, const ICLArray<T> *array, const Strides &strides, unsigned int num_dimensions, const Window &window);
331  /** Add the passed tensor's parameters to the object's kernel's arguments starting from the index idx.
332  *
333  * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
334  * @param[in] tensor Tensor to set as an argument of the object's kernel.
335  * @param[in] window Window the kernel will be executed on.
336  */
337  template <unsigned int dimension_size>
338  void add_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window);
339 
340 protected:
341  cl::Kernel _kernel; /**< OpenCL kernel to run */
342  GPUTarget _target; /**< The targeted GPU */
343  std::string _config_id; /**< Configuration ID */
344  size_t _max_workgroup_size; /**< The maximum workgroup size for this kernel */
345 private:
346  cl::NDRange _lws_hint; /**< Local workgroup size hint for the OpenCL kernel */
347 };
348 
349 /** Add the kernel to the command queue with the given window.
350  *
351  * @note Depending on the size of the window, this might translate into several jobs being enqueued.
352  *
353  * @note If kernel->kernel() is empty then the function will return without adding anything to the queue.
354  *
355  * @param[in,out] queue OpenCL command queue.
356  * @param[in] kernel Kernel to enqueue
357  * @param[in] window Window the kernel has to process.
358  * @param[in] lws_hint (Optional) Local workgroup size requested. Default is based on the device target.
359  * @param[in] use_dummy_work_items (Optional) Use dummy work items in order to have two dimensional power of two NDRange. Default is false
360  * Note: it is kernel responsibility to check if the work-item is out-of-range
361  *
362  * @note If any dimension of the lws is greater than the global workgroup size then no lws will be passed.
363  */
364 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);
365 
366 /** Add the passed array's parameters to the object's kernel's arguments starting from the index idx.
367  *
368  * @param[in,out] idx Index at which to start adding the array's arguments. Will be incremented by the number of kernel arguments set.
369  * @param[in] array Array to set as an argument of the object's kernel.
370  * @param[in] strides @ref Strides object containing stride of each dimension in bytes.
371  * @param[in] num_dimensions Number of dimensions of the @p array.
372  * @param[in] window Window the kernel will be executed on.
373  */
374 template <typename T, unsigned int dimension_size>
375 void ICLKernel::add_array_argument(unsigned &idx, const ICLArray<T> *array, const Strides &strides, unsigned int num_dimensions, const Window &window)
376 {
377  ARM_COMPUTE_ERROR_ON(array == nullptr);
378 
379  // Calculate offset to the start of the window
380  unsigned int offset_first_element = 0;
381 
382  for(unsigned int n = 0; n < num_dimensions; ++n)
383  {
384  offset_first_element += window[n].start() * strides[n];
385  }
386 
387  unsigned int idx_start = idx;
388  _kernel.setArg(idx++, array->cl_buffer());
389 
390  for(unsigned int dimension = 0; dimension < dimension_size; dimension++)
391  {
392  _kernel.setArg<cl_uint>(idx++, strides[dimension]);
393  _kernel.setArg<cl_uint>(idx++, strides[dimension] * window[dimension].step());
394  }
395 
396  _kernel.setArg<cl_uint>(idx++, offset_first_element);
397 
398  ARM_COMPUTE_ERROR_ON_MSG_VAR(idx_start + num_arguments_per_array<dimension_size>() != idx,
399  "add_%dD_array_argument() is supposed to add exactly %d arguments to the kernel", dimension_size, num_arguments_per_array<dimension_size>());
400  ARM_COMPUTE_UNUSED(idx_start);
401 }
402 }
403 #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:185
static cl::NDRange gws_from_window(const Window &window)
Get the global work size given an execution window.
Definition: ICLKernel.cpp:141
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's parameters to the object's kernel's arguments starting from the index idx ...
Definition: ICLKernel.h:122
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28
static constexpr unsigned int num_arguments_per_1D_array()
Returns the number of arguments enqueued per 1D array object.
Definition: ICLKernel.h:177
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:220
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 ...
Definition: ICLKernel.h:146
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:39
cl::Kernel & kernel()
Returns a reference to the OpenCL kernel of this object.
Definition: ICLKernel.h:88
cl::NDRange lws_hint() const
Return the Local-Workgroup-Size hint.
Definition: ICLKernel.h:263
void set_lws_hint(const cl::NDRange &lws_hint)
Set the Local-Workgroup-Size hint.
Definition: ICLKernel.h:253
void add_argument(unsigned int &idx, T value)
Add the passed parameters to the object's kernel's arguments starting from the index idx.
Definition: ICLKernel.h:242
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.
Definition: ICLKernel.h:101
#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.
#define ARM_COMPUTE_ERROR_ON_MSG_VAR(cond, msg,...)
Definition: Error.h:457
Common interface for all the OpenCL kernels.
Definition: ICLKernel.h:44
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.
Definition: ICLKernel.h:159
Copyright (c) 2017-2020 Arm Limited.
const std::string & config_id() const
Get the configuration ID.
Definition: ICLKernel.h:278
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:201
#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:302
static constexpr unsigned int num_arguments_per_2D_tensor()
Returns the number of arguments enqueued per 2D tensor object.
Definition: ICLKernel.h:193
static constexpr unsigned int num_arguments_per_4D_tensor()
Returns the number of arguments enqueued per 4D tensor object.
Definition: ICLKernel.h:209
Strides of an item in bytes.
Definition: Strides.h:37
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:135
Interface for OpenCL tensor.
Definition: ICLTensor.h:42
ICLKernel()
Constructor.
Definition: ICLKernel.h:80
GPUTarget
Available GPU Targets.
Definition: GPUTarget.h:34
size_t get_max_workgroup_size()
Get the maximum workgroup size for the device the CLKernelLibrary uses.
Definition: ICLKernel.cpp:132
Tensor packing service.
Definition: ITensorPack.h:36
void set_target(GPUTarget target)
Set the targeted GPU architecture.
Definition: ICLKernel.h:287
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:111
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.
Definition: ICLKernel.h:169
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 ...
Describe a multidimensional execution window.
Definition: Window.h:39
#define ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(k)
Definition: Validate.h:941
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:232