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