Compute Library
 23.11
ICLKernel.cpp
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2016-2023 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 #include "src/core/CL/ICLKernel.h"
25 
28 
29 #include "src/core/helpers/Utils.h"
30 
31 #include <cstddef>
32 
33 void arm_compute::enqueue(cl::CommandQueue &queue,
34  ICLKernel &kernel,
35  const Window &window,
36  const cl::NDRange &lws_hint,
37  bool use_dummy_work_items)
38 {
39  if (kernel.kernel()() == nullptr)
40  {
41  return;
42  }
43 
44  for (unsigned int i = 0; i < Coordinates::num_max_dimensions; ++i)
45  {
46  ARM_COMPUTE_ERROR_ON(window[i].step() == 0);
47  // Make sure that dimensions > Z are 1
48  ARM_COMPUTE_ERROR_ON((i >= 3) && ((window[i].end() - window[i].start()) != 1));
49  }
50 
51  cl::NDRange gws = ICLKernel::gws_from_window(window, use_dummy_work_items);
52 
53  // Check for empty NDRange
54  if (gws.dimensions() == 0)
55  {
56  return;
57  }
58 
59  kernel.cache_gws(gws);
60 
61  cl::NDRange valid_lws;
62  if (lws_hint[0] * lws_hint[1] * lws_hint[2] > kernel.get_max_workgroup_size())
63  {
64  valid_lws = cl::NullRange;
65  }
66  else
67  {
68  valid_lws = lws_hint;
69  }
70 
71  cl::NDRange lws = cl::NullRange;
72 
73  if ((valid_lws[0] <= gws[0]) && (valid_lws[1] <= gws[1]) && (valid_lws[2] <= gws[2]))
74  {
75  lws = valid_lws;
76  }
77 
78  if (CLKernelLibrary::get().is_wbsm_supported())
79  {
80  set_wbsm(kernel.kernel(), kernel.wbsm_hint());
81  }
82  queue.enqueueNDRangeKernel(kernel.kernel(), cl::NullRange, gws, lws);
83 }
84 
85 namespace arm_compute
86 {
87 template <unsigned int dimension_size>
88 void ICLKernel::add_tensor_argument(unsigned &idx, const ICLTensor *tensor, const Window &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",
116  dimension_size, num_arguments_per_tensor<dimension_size>());
117  ARM_COMPUTE_UNUSED(idx_start);
118 }
119 
121 {
122  ARM_COMPUTE_ERROR_ON(tensor == nullptr);
123 
124  const ITensorInfo *info = tensor->info();
125  ARM_COMPUTE_ERROR_ON(info == nullptr);
126  const Strides &strides = info->strides_in_bytes();
127 
128  // Tensor poniter
129  _kernel.setArg(idx++, tensor->cl_buffer());
130 
131  // Add stride_y, stride_z
132  _kernel.setArg<cl_uint>(idx++, strides[1]);
133  _kernel.setArg<cl_uint>(idx++, strides[2]);
134 
135  // Tensor dimensions
136  _kernel.setArg<cl_uint>(idx++, info->dimension(0));
137  _kernel.setArg<cl_uint>(idx++, info->dimension(1));
138  _kernel.setArg<cl_uint>(idx++, info->dimension(2));
139 
140  // Offset of first element
141  unsigned int offset_first_element = info->offset_first_element_in_bytes();
142  _kernel.setArg<cl_uint>(idx++, offset_first_element);
143 }
144 
146 {
147  ARM_COMPUTE_ERROR_ON(tensor == nullptr);
148 
149  const ITensorInfo *info = tensor->info();
150  ARM_COMPUTE_ERROR_ON(info == nullptr);
151  const Strides &strides = info->strides_in_bytes();
152 
153  // Tensor poniter
154  _kernel.setArg(idx++, tensor->cl_buffer());
155 
156  // Add stride_y, stride_z and stride_w
157  _kernel.setArg<cl_uint>(idx++, strides[1]);
158  _kernel.setArg<cl_uint>(idx++, strides[2]);
159  _kernel.setArg<cl_uint>(idx++, strides[3]);
160 
161  // Tensor dimensions
162  _kernel.setArg<cl_uint>(idx++, info->dimension(0));
163  _kernel.setArg<cl_uint>(idx++, info->dimension(1));
164  _kernel.setArg<cl_uint>(idx++, info->dimension(2));
165  _kernel.setArg<cl_uint>(idx++, info->dimension(3));
166 
167  // Offset of first element
168  unsigned int offset_first_element = info->offset_first_element_in_bytes();
169  _kernel.setArg<cl_uint>(idx++, offset_first_element);
170 }
171 
172 #ifndef DOXYGEN_SKIP_THIS
173 template void ICLKernel::add_tensor_argument<1>(unsigned &idx, const ICLTensor *tensor, const Window &window);
174 template void ICLKernel::add_tensor_argument<2>(unsigned &idx, const ICLTensor *tensor, const Window &window);
175 template void ICLKernel::add_tensor_argument<3>(unsigned &idx, const ICLTensor *tensor, const Window &window);
176 template void ICLKernel::add_tensor_argument<4>(unsigned &idx, const ICLTensor *tensor, const Window &window);
177 template void ICLKernel::add_tensor_argument<5>(unsigned &idx, const ICLTensor *tensor, const Window &window);
178 #endif /* DOXYGEN_SKIP_THIS */
179 
180 void ICLKernel::set_target(cl::Device &device)
181 {
182  _target = get_target_from_device(device);
183 }
184 
186 {
187  if (_max_workgroup_size == 0)
188  {
189  _max_workgroup_size = CLKernelLibrary::get().max_local_workgroup_size(_kernel);
190  }
191  return _max_workgroup_size;
192 }
193 
194 cl::NDRange ICLKernel::gws_from_window(const Window &window, bool use_dummy_work_items)
195 {
196  if ((window.x().end() - window.x().start()) == 0 || (window.y().end() - window.y().start()) == 0)
197  {
198  return cl::NullRange;
199  }
200 
201  cl::NDRange gws((window.x().end() - window.x().start()) / window.x().step(),
202  (window.y().end() - window.y().start()) / window.y().step(),
203  (window.z().end() - window.z().start()) / window.z().step());
204 
205  if (use_dummy_work_items)
206  {
207  gws.get()[0] = get_next_power_two(gws[0]);
208  gws.get()[1] = get_next_power_two(gws[1]);
209  }
210 
211  return gws;
212 }
213 
214 cl::NDRange ICLKernel::get_cached_gws() const
215 {
216  return _cached_gws;
217 }
218 
219 void ICLKernel::cache_gws(const cl::NDRange &gws)
220 {
221  _cached_gws = gws;
222 }
223 } // namespace arm_compute
arm_compute::Window::Dimension::start
constexpr int start() const
Return the start of the dimension.
Definition: Window.h:96
ICLTensor.h
arm_compute::ICLKernel::get_max_workgroup_size
size_t get_max_workgroup_size()
Get the maximum workgroup size for the device the CLKernelLibrary uses.
Definition: ICLKernel.cpp:185
Helpers.h
ARM_COMPUTE_ERROR_ON_MSG_VAR
#define ARM_COMPUTE_ERROR_ON_MSG_VAR(cond, msg,...)
Definition: Error.h:457
arm_compute::Window::Dimension::step
constexpr int step() const
Return the step of the dimension.
Definition: Window.h:106
arm_compute::ICLKernel::kernel
cl::Kernel & kernel()
Returns a reference to the OpenCL kernel of this object.
Definition: ICLKernel.h:151
arm_compute::ICLKernel::wbsm_hint
cl_int wbsm_hint() const
Return the workgroup batch size modifier hint.
Definition: ICLKernel.h:404
arm_compute::ICLTensor
Interface for OpenCL tensor.
Definition: ICLTensor.h:41
arm_compute::get_next_power_two
unsigned int get_next_power_two(unsigned int x)
Given an integer value, this function returns the next power of two.
Definition: Utils.h:74
arm_compute::set_wbsm
void set_wbsm(cl::Kernel &kernel, cl_int wbsm_hint)
Definition: CLHelpers.cpp:441
arm_compute::ICLKernel::add_3d_tensor_nhw_argument
void add_3d_tensor_nhw_argument(unsigned int &idx, const ICLTensor *tensor)
Add the passed NHW 3D tensor's parameters to the object's kernel's arguments by passing strides,...
Definition: ICLKernel.cpp:120
arm_compute::CLKernelLibrary::get
static CLKernelLibrary & get()
Access the KernelLibrary singleton.
Definition: CLKernelLibrary.cpp:41
arm_compute::Strides
Strides of an item in bytes.
Definition: Strides.h:38
ICLKernel.h
ARM_COMPUTE_ERROR_ON
#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
arm_compute::ICLKernel::gws_from_window
static cl::NDRange gws_from_window(const Window &window, bool use_dummy_work_items)
Get the global work size given an execution window.
Definition: ICLKernel.cpp:194
arm_compute::Window::y
constexpr const Dimension & y() const
Alias to access the second dimension of the window.
Definition: Window.h:167
arm_compute::ICLKernel::set_target
void set_target(GPUTarget target)
Set the targeted GPU architecture.
Definition: ICLKernel.h:428
arm_compute::Window::is_broadcasted
bool is_broadcasted(size_t dimension) const
Return whether a dimension has been broadcasted.
Definition: Window.inl:66
ARM_COMPUTE_UNUSED
#define ARM_COMPUTE_UNUSED(...)
To avoid unused variables warnings.
Definition: Error.h:151
arm_compute::get_target_from_device
GPUTarget get_target_from_device(const cl::Device &device)
Helper function to get the GPU target from CL device.
Definition: CLHelpers.cpp:224
tensor
CLTensor * tensor
Pointer to the auxiliary tensor.
Definition: ClWorkloadRuntime.cpp:67
arm_compute::ICLKernel
Common interface for all the OpenCL kernels.
Definition: ICLKernel.h:67
arm_compute::ICLKernel::cache_gws
void cache_gws(const cl::NDRange &gws)
Cache the latest gws used to enqueue this kernel.
Definition: ICLKernel.cpp:219
arm_compute::IKernel::window
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28
arm_compute::ICLKernel::add_4d_tensor_nhwc_argument
void add_4d_tensor_nhwc_argument(unsigned int &idx, const ICLTensor *tensor)
Add the passed NHWC 4D tensor's parameters to the object's kernel's arguments by passing strides,...
Definition: ICLKernel.cpp:145
arm_compute::Window
Describe a multidimensional execution window.
Definition: Window.h:39
Utils.h
arm_compute
Copyright (c) 2017-2023 Arm Limited.
Definition: introduction.dox:24
arm_compute::mlgo::parser::end
void end(TokenStream &in, bool &valid)
Definition: MLGOParser.cpp:283
arm_compute::CLKernelLibrary::max_local_workgroup_size
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.
Definition: CLKernelLibrary.cpp:118
arm_compute::ITensorInfo
Store the tensor's metadata.
Definition: ITensorInfo.h:44
arm_compute::test::validation::info
ScaleKernelInfo info(interpolation_policy, default_border_mode, PixelValue(), sampling_policy, false)
arm_compute::cpu::step
constexpr int step
Definition: fp32.cpp:35
arm_compute::ICLKernel::get_cached_gws
cl::NDRange get_cached_gws() const
Get the cached gws used to enqueue this kernel.
Definition: ICLKernel.cpp:214
arm_compute::Window::Dimension::end
constexpr int end() const
Return the end of the dimension.
Definition: Window.h:101
arm_compute::Window::x
constexpr const Dimension & x() const
Alias to access the first dimension of the window.
Definition: Window.h:158
arm_compute::Window::z
constexpr const Dimension & z() const
Alias to access the third dimension of the window.
Definition: Window.h:176
arm_compute::enqueue
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:33