41 : real_clEnqueueNDRangeKernel(nullptr),
42 _tuning_params_table(),
45 _tune_new_kernels(tune_new_kernels),
46 _tuning_info(tuning_info)
50 struct CLTuner::IKernelData
52 virtual ~IKernelData() =
default;
53 virtual void do_run(
ICLKernel &kernel, cl::CommandQueue &queue) = 0;
55 struct DefaultKernelData :
public CLTuner::IKernelData
57 DefaultKernelData(ITensorPack &tensors) : _tensors{tensors}
60 ~DefaultKernelData()
override =
default;
61 void do_run(ICLKernel &kernel, cl::CommandQueue &queue)
override
63 const bool inject_memory = !_tensors.empty();
64 inject_memory ? kernel.run_op(_tensors, kernel.window(), queue) : kernel.run(kernel.window(), queue);
68 ITensorPack &_tensors;
73 return _kernel_event() !=
nullptr;
77 _kernel_event = kernel_event;
86 return _tune_new_kernels;
105 void CLTuner::do_tune_kernel_dynamic(
ICLKernel &kernel, IKernelData *data)
112 if (kernel.
config_id() != arm_compute::default_config_id)
114 auto p = _tuning_params_table.find(config_id);
116 if (p == _tuning_params_table.end())
118 if (_tune_new_kernels)
121 CLTuningParams opt_tuning_params = find_optimal_tuning_params(kernel, data);
147 DefaultKernelData data{tensors};
149 do_tune_kernel_dynamic(kernel, &data);
154 _tuning_params_table.emplace(kernel_id, optimal_tuning_params);
160 cl::CommandQueue queue_profiler;
172 cl_command_queue_properties props = default_queue.getInfo<CL_QUEUE_PROPERTIES>();
174 if ((props & CL_QUEUE_PROFILING_ENABLE) == 0)
181 queue_profiler = default_queue;
185 auto interceptor = [
this](cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim,
const size_t *gwo,
186 const size_t *gws,
const size_t *lws, cl_uint num_events_in_wait_list,
187 const cl_event *event_wait_list, cl_event *event)
196 num_events_in_wait_list, event_wait_list, &tmp);
201 if (event !=
nullptr)
212 data->do_run(kernel, queue_profiler);
219 cl::NDRange gws = kernel.get_cached_gws();
222 "[CLTuner] Kernel with config_id '%s' uses %s as the upper-bound for lws search", kernel.config_id().c_str(),
225 queue_profiler.finish();
227 const cl_ulong start = _kernel_event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
228 const cl_ulong
end = _kernel_event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
229 cl_ulong min_exec_time =
end - start;
230 _kernel_event =
nullptr;
232 CLTuningParams opt_tuning_params(cl::NullRange, 0);
236 for (
size_t i = 0; i < tuning_list->size(); ++i)
238 CLTuningParams tuning_test = (*tuning_list)[i];
240 cl::NDRange lws_test = tuning_test.get_lws();
241 auto x = lws_test[0];
242 auto y = lws_test[1];
243 auto z = lws_test[2];
244 const bool invalid_lws = (x * y * z > kernel.get_max_workgroup_size()) || (x == 1 && y == 1 && z == 1);
251 kernel.set_lws_hint(lws_test);
254 cl_int wbsm_test = tuning_test.get_wbsm();
255 kernel.set_wbsm_hint(wbsm_test);
258 to_string(kernel.lws_hint()).c_str(), kernel.wbsm_hint());
261 data->do_run(kernel, queue_profiler);
263 queue_profiler.finish();
265 const cl_ulong start = _kernel_event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
266 const cl_ulong
end = _kernel_event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
267 const cl_ulong diff =
end - start;
268 _kernel_event =
nullptr;
271 if (diff < min_exec_time)
273 min_exec_time = diff;
274 opt_tuning_params.
set_lws(tuning_test.get_lws());
277 opt_tuning_params.
set_wbsm(tuning_test.get_wbsm());
284 return opt_tuning_params;
289 return _tuning_params_table;
294 _tuning_params_table.clear();
301 fs.exceptions(std::ifstream::badbit);
302 fs.open(filename, std::ios::in);
308 bool header_line =
true;
309 while (!std::getline(fs, line).fail())
314 size_t pos_lws = line.find(
"lws");
315 size_t pos_wbsm = line.find(
"wbsm");
317 if (pos_lws != std::string::npos || pos_wbsm != std::string::npos)
320 if (pos_wbsm != std::string::npos)
326 if (std::getline(fs, line).fail())
334 size_t pos = line.find(
";");
335 if (pos == std::string::npos)
339 std::string kernel_id = line.substr(0, pos);
340 line.erase(0, pos + 1);
341 if (!tuning_params.
from_string(_tuning_info, line))
352 if (!_tune_new_kernels || _tuning_params_table.empty() || filename.empty())
357 fs.exceptions(std::ifstream::failbit | std::ifstream::badbit);
358 fs.open(filename, std::ios::out);
359 std::string header_string =
"";
360 header_string +=
"lws";
363 if (!header_string.empty())
365 header_string +=
" ";
367 header_string +=
"wbsm";
369 fs << header_string << std::endl;
370 for (
auto const &kernel_data : _tuning_params_table)
373 fs << kernel_data.first << tun_pams.
to_string(_tuning_info) << std::endl;