Compute Library
 22.08
CLTuner.cpp
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2017-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  */
26 
27 #include "arm_compute/core/Error.h"
29 #include "src/core/CL/ICLKernel.h"
30 #include "support/StringSupport.h"
31 #if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
33 #endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
34 
35 #include <cerrno>
36 #include <fstream>
37 #include <limits>
38 
39 namespace arm_compute
40 {
41 CLTuner::CLTuner(bool tune_new_kernels, CLTuningInfo tuning_info)
42  : real_clEnqueueNDRangeKernel(nullptr), _tuning_params_table(), _lws_table(), _kernel_event(), _tune_new_kernels(tune_new_kernels), _tuning_info(tuning_info)
43 {
44 }
45 
46 struct CLTuner::IKernelData
47 {
48  virtual ~IKernelData() = default;
49  virtual void do_run(ICLKernel &kernel, cl::CommandQueue &queue) = 0;
50 };
51 struct DefaultKernelData : public CLTuner::IKernelData
52 {
53  DefaultKernelData(ITensorPack &tensors)
54  : _tensors{ tensors }
55  {
56  }
57  ~DefaultKernelData() override = default;
58  void do_run(ICLKernel &kernel, cl::CommandQueue &queue) override
59  {
60  const bool inject_memory = !_tensors.empty();
61  inject_memory ? kernel.run_op(_tensors, kernel.window(), queue) : kernel.run(kernel.window(), queue);
62  }
63 
64 private:
65  ITensorPack &_tensors;
66 };
67 
68 #if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
69 struct CompositeKernelData : public CLTuner::IKernelData
70 {
71  CompositeKernelData(ITensorPack &tensors, const experimental::dynamic_fusion::ClExecutionDescriptor &exec_desc)
72  : _tensors{ tensors }, _exec_desc{ exec_desc }
73  {
74  }
75  ~CompositeKernelData() override = default;
76  void do_run(ICLKernel &kernel, cl::CommandQueue &queue) override
77  {
78  // ClCompositeKernel is purely stateless, and thus always requires memory injection
79  kernel.run_composite_op(_tensors, kernel.window(), queue, _exec_desc);
80  }
81 
82 private:
83  ITensorPack &_tensors;
85 };
86 #endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
87 
89 {
90  return _kernel_event() != nullptr;
91 }
92 void CLTuner::set_cl_kernel_event(cl_event kernel_event)
93 {
94  _kernel_event = kernel_event;
95 }
96 
98 {
99  _tune_new_kernels = tune_new_kernels;
100 }
102 {
103  return _tune_new_kernels;
104 }
105 
107 {
108  _tuning_info.tuner_mode = mode;
109 }
110 
112 {
113  ARM_COMPUTE_UNUSED(kernel);
114 }
115 
117 {
119  tune_kernel_dynamic(kernel, pack);
120 }
121 
122 void CLTuner::do_tune_kernel_dynamic(ICLKernel &kernel, IKernelData *data)
123 {
124  // Get the configuration ID from the kernel and append GPU target name and number of available compute units
125  const std::string config_id = kernel.config_id() + "_" + string_from_target(kernel.get_target()) + "_MP" + support::cpp11::to_string(CLKernelLibrary::get().get_num_compute_units());
126 
127  // Check if we need to find the Optimal LWS. If the kernel's config_id is equal to default_config_id, the kernel does not require to be tuned
128  if(kernel.config_id() != arm_compute::default_config_id)
129  {
130  auto p = _tuning_params_table.find(config_id);
131 
132  if(p == _tuning_params_table.end())
133  {
134  if(_tune_new_kernels)
135  {
136  // Find the optimal LWS for the kernel
137  CLTuningParams opt_tuning_params = find_optimal_tuning_params(kernel, data);
138 
139  // Insert the optimal LWS in the table
140  add_tuning_params(config_id, opt_tuning_params);
141 
142  // Set Local-Workgroup-Size
143  kernel.set_lws_hint(opt_tuning_params.get_lws());
144  if(_tuning_info.tune_wbsm)
145  {
146  kernel.set_wbsm_hint(opt_tuning_params.get_wbsm());
147  }
148  }
149  }
150  else
151  {
152  // Set Local-Workgroup-Size
153  kernel.set_lws_hint(p->second.get_lws());
154  if(_tuning_info.tune_wbsm)
155  {
156  kernel.set_wbsm_hint(p->second.get_wbsm());
157  }
158  }
159  }
160 }
162 {
163  DefaultKernelData data{ tensors };
164 
165  do_tune_kernel_dynamic(kernel, &data);
166 }
167 
168 #if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
170 {
171  CompositeKernelData data{ tensors, exec_desc };
172 
173  do_tune_kernel_dynamic(kernel, &data);
174 }
175 #endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
176 
177 void CLTuner::add_tuning_params(const std::string &kernel_id, CLTuningParams optimal_tuning_params)
178 {
179  _tuning_params_table.emplace(kernel_id, optimal_tuning_params);
180 }
181 
182 CLTuningParams CLTuner::find_optimal_tuning_params(ICLKernel &kernel, IKernelData *data)
183 {
184  // Profiling queue
185  cl::CommandQueue queue_profiler;
186 
187  // Extract real OpenCL function to intercept
188  if(real_clEnqueueNDRangeKernel == nullptr)
189  {
191  }
192 
193  // Get the default queue
194  cl::CommandQueue default_queue = CLScheduler::get().queue();
195 
196  // Check if we can use the OpenCL timer with the default queue
197  cl_command_queue_properties props = default_queue.getInfo<CL_QUEUE_PROPERTIES>();
198 
199  if((props & CL_QUEUE_PROFILING_ENABLE) == 0)
200  {
201  // Set the queue for profiling
202  queue_profiler = cl::CommandQueue(CLScheduler::get().context(), props | CL_QUEUE_PROFILING_ENABLE);
203  }
204  else
205  {
206  queue_profiler = default_queue;
207  }
208 
209  // Start intercepting enqueues:
210  auto interceptor = [this](cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *gwo, const size_t *gws, const size_t *lws, cl_uint num_events_in_wait_list,
211  const cl_event * event_wait_list, cl_event * event)
212  {
213  if(this->kernel_event_is_set())
214  {
215  // If the event is already set it means the kernel enqueue is sliced: given that we only time the first slice we can save time by skipping the other enqueues.
216  return CL_SUCCESS;
217  }
218  cl_event tmp;
219  cl_int retval = this->real_clEnqueueNDRangeKernel(command_queue, kernel, work_dim, gwo, gws, lws, num_events_in_wait_list, event_wait_list, &tmp);
220 
221  // Set OpenCL event
222  this->set_cl_kernel_event(tmp);
223 
224  if(event != nullptr)
225  {
226  //return cl_event from the intercepted call
227  clRetainEvent(tmp);
228  *event = tmp;
229  }
230  return retval;
231  };
233 
234  cl::NDRange gws = ICLKernel::gws_from_window(kernel.window());
235 
236  // Run the kernel with default lws to be used as baseline
237  data->do_run(kernel, queue_profiler);
238 
239  queue_profiler.finish();
240 
241  const cl_ulong start = _kernel_event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
242  const cl_ulong end = _kernel_event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
243  cl_ulong min_exec_time = end - start;
244  _kernel_event = nullptr;
245 
246  CLTuningParams opt_tuning_params(cl::NullRange, 0);
247 
248  // Construct the list of tuning parameters values to be tested based on the tuner mode.
249  auto tuning_list = cl_tuner::get_tuning_parameters_list(_tuning_info, gws);
250  for(size_t i = 0; i < tuning_list->size(); ++i)
251  {
252  CLTuningParams tuning_test = (*tuning_list)[i];
253  // Setting the lws
254  cl::NDRange lws_test = tuning_test.get_lws();
255  auto x = lws_test[0];
256  auto y = lws_test[1];
257  auto z = lws_test[2];
258  const bool invalid_lws = (x * y * z > kernel.get_max_workgroup_size()) || (x == 1 && y == 1 && z == 1);
259 
260  if(invalid_lws)
261  {
262  continue;
263  }
264 
265  kernel.set_lws_hint(lws_test);
266  if(_tuning_info.tune_wbsm && CLKernelLibrary::get().is_wbsm_supported())
267  {
268  cl_int wbsm_test = tuning_test.get_wbsm();
269  kernel.set_wbsm_hint(wbsm_test);
270  }
271 
272  // Run the kernel
273  data->do_run(kernel, queue_profiler);
274 
275  queue_profiler.finish();
276 
277  const cl_ulong start = _kernel_event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
278  const cl_ulong end = _kernel_event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
279  const cl_ulong diff = end - start;
280  _kernel_event = nullptr;
281 
282  // Check the execution time
283  if(diff < min_exec_time)
284  {
285  min_exec_time = diff;
286  opt_tuning_params.set_lws(tuning_test.get_lws());
287  if(_tuning_info.tune_wbsm)
288  {
289  opt_tuning_params.set_wbsm(tuning_test.get_wbsm());
290  }
291  }
292  }
293 
294  // Restore real function
296  return opt_tuning_params;
297 }
298 
299 const std::unordered_map<std::string, CLTuningParams> &CLTuner::tuning_params_table() const
300 {
301  return _tuning_params_table;
302 }
303 
304 void CLTuner::import_tuning_params(const std::unordered_map<std::string, CLTuningParams> &tuning_params_table)
305 {
306  _tuning_params_table.clear();
307  _tuning_params_table = tuning_params_table;
308 }
309 
310 void CLTuner::load_from_file(const std::string &filename)
311 {
312  std::ifstream fs;
313  fs.exceptions(std::ifstream::badbit);
314  fs.open(filename, std::ios::in);
315  if(!fs.is_open())
316  {
317  ARM_COMPUTE_ERROR_VAR("Failed to open '%s' (%s [%d])", filename.c_str(), strerror(errno), errno);
318  }
319  std::string line;
320  bool header_line = true;
321  while(!std::getline(fs, line).fail())
322  {
323  if(header_line)
324  {
325  header_line = false;
326  size_t pos_lws = line.find("lws");
327  size_t pos_wbsm = line.find("wbsm");
328  _tuning_info.tune_wbsm = false;
329  if(pos_lws != std::string::npos || pos_wbsm != std::string::npos)
330  {
331  // The file has in the first line the parameters it has been tuned on
332  if(pos_wbsm != std::string::npos)
333  {
334  _tuning_info.tune_wbsm = true;
335  }
336  // Once the line with the tuning parameter is read we can
337  // read the next one to start collecting the values
338  if(std::getline(fs, line).fail())
339  {
340  break;
341  }
342  }
343  }
344 
345  CLTuningParams tuning_params;
346  size_t pos = line.find(";");
347  if(pos == std::string::npos)
348  {
349  ARM_COMPUTE_ERROR_VAR("Malformed row '%s' in %s", line.c_str(), filename.c_str());
350  }
351  std::string kernel_id = line.substr(0, pos);
352  line.erase(0, pos + 1);
353  if(!tuning_params.from_string(_tuning_info, line))
354  {
355  ARM_COMPUTE_ERROR_VAR("Malformed row '%s' in %s", line.c_str(), filename.c_str());
356  }
357  add_tuning_params(kernel_id, tuning_params);
358  }
359  fs.close();
360 }
361 
362 bool CLTuner::save_to_file(const std::string &filename) const
363 {
364  if(!_tune_new_kernels || _tuning_params_table.empty() || filename.empty())
365  {
366  return false;
367  }
368  std::ofstream fs;
369  fs.exceptions(std::ifstream::failbit | std::ifstream::badbit);
370  fs.open(filename, std::ios::out);
371  std::string header_string = "";
372  header_string += "lws";
373  if(_tuning_info.tune_wbsm)
374  {
375  if(!header_string.empty())
376  {
377  header_string += " ";
378  }
379  header_string += "wbsm";
380  }
381  fs << header_string << std::endl;
382  for(auto const &kernel_data : _tuning_params_table)
383  {
384  CLTuningParams tun_pams(kernel_data.second);
385  fs << kernel_data.first << tun_pams.to_string(_tuning_info) << std::endl;
386  }
387  fs.close();
388  return true;
389 }
390 } // namespace arm_compute
static cl::NDRange gws_from_window(const Window &window)
Get the global work size given an execution window.
Definition: ICLKernel.cpp:193
void set_tuner_mode(CLTunerMode mode)
Set OpenCL tuner mode.
Definition: CLTuner.cpp:106
const Window & window() const
The maximum window the kernel can be executed on.
Definition: IKernel.cpp:28
void set_cl_kernel_event(cl_event kernel_event)
Set the OpenCL kernel event.
Definition: CLTuner.cpp:92
cl::NDRange get_lws() const
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
const std::unordered_map< std::string, CLTuningParams > & tuning_params_table() const
Give read access to the tuning params table.
Definition: CLTuner.cpp:299
static CLScheduler & get()
Access the scheduler singleton.
void import_tuning_params(const std::unordered_map< std::string, CLTuningParams > &tuning_params_table)
Import tuning parameters table.
Definition: CLTuner.cpp:304
void tune_kernel_static(ICLKernel &kernel) override
Tune OpenCL kernel statically.
Definition: CLTuner.cpp:111
bool empty() const
Checks if pack is empty.
Definition: ITensorPack.cpp:80
void tune_kernel_dynamic(ICLKernel &kernel) override
Tune OpenCL kernel dynamically.
Definition: CLTuner.cpp:116
void set_lws_hint(const cl::NDRange &lws_hint)
Set the Local-Workgroup-Size hint.
Definition: ICLKernel.h:374
std::string to_string(T &&value)
Convert integer and float values to string.
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 load_from_file(const std::string &filename)
Load the tuning parameters table from file.
Definition: CLTuner.cpp:310
#define ARM_COMPUTE_ERROR_VAR(msg,...)
Print the given message then throw an std::runtime_error.
Definition: Error.h:346
static CLKernelLibrary & get()
Access the KernelLibrary singleton.
< OpenCL tuner parameters
const std::string & string_from_target(GPUTarget target)
Translates a given gpu device target to string.
Definition: GPUTarget.cpp:151
cl_int clRetainEvent(cl_event event)
Definition: OpenCL.cpp:907
Common interface for all the OpenCL kernels.
Definition: ICLKernel.h:81
Copyright (c) 2017-2022 Arm Limited.
CLTuner(bool tune_new_kernels=true, CLTuningInfo tuning_info=CLTuningInfo())
Constructor.
Definition: CLTuner.cpp:41
std::function< decltype(clEnqueueNDRangeKernel)> real_clEnqueueNDRangeKernel
clEnqueueNDRangeKernel symbol
Definition: CLTuner.h:106
cl_uint get_num_compute_units()
Return the maximum number of compute units in the device.
void set_lws(cl::NDRange lws)
const std::string & config_id() const
Get the configuration ID.
Definition: ICLKernel.h:420
bool kernel_event_is_set() const
Is the kernel_event set ?
Definition: CLTuner.cpp:88
Interface to enqueue OpenCL kernels and get/set the OpenCL CommandQueue and ICLTuner.
std::function< decltype(clEnqueueNDRangeKernel)> clEnqueueNDRangeKernel_ptr
Definition: OpenCL.h:99
Descriptor containing information required to run a single ClWorkload.
Definition: ClWorkload.h:91
#define ARM_COMPUTE_UNUSED(...)
To avoid unused variables warnings.
Definition: Error.h:152
bool save_to_file(const std::string &filename) const
Save the content of the tuning parameters table to file.
Definition: CLTuner.cpp:362
void add_tuning_params(const std::string &kernel_id, CLTuningParams optimal_tuning_params)
Manually add tuning parameters for a kernel.
Definition: CLTuner.cpp:177
void set_tune_new_kernels(bool tune_new_kernels)
Setter for tune_new_kernels option.
Definition: CLTuner.cpp:97
GPUTarget get_target() const
Get the targeted GPU architecture.
Definition: ICLKernel.h:444
void set_wbsm(cl_int wbsm)
void end(TokenStream &in, bool &valid)
Definition: MLGOParser.cpp:290
cl::CommandQueue & queue()
Accessor for the associated CL command queue.
Definition: CLScheduler.cpp:43
std::string to_string(CLTuningInfo tuning_info)
CLTunerMode
< OpenCL tuner modes
Definition: CLTunerTypes.h:35
Tensor packing service.
Definition: ITensorPack.h:39
std::unique_ptr< ICLTuningParametersList > get_tuning_parameters_list(CLTuningInfo tuning_info, const cl::NDRange &gws)
Construct an ICLTuningParametersList object for the given tuner mode and gws configuration.
void set_wbsm_hint(const cl_int &wbsm_hint)
Set the workgroup batch size modifier hint.
Definition: ICLKernel.h:395
bool tune_new_kernels() const
Tune kernels that are not in the tuning parameters table.
Definition: CLTuner.cpp:101
bool from_string(CLTuningInfo tuning_info, std::string tuning_params_string)
static CLSymbols & get()
Get the static instance of CLSymbols.
Definition: OpenCL.cpp:45
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