31 #if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION) 33 #endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION) 42 : real_clEnqueueNDRangeKernel(nullptr), _tuning_params_table(), _lws_table(), _kernel_event(), _tune_new_kernels(tune_new_kernels), _tuning_info(tuning_info)
46 struct CLTuner::IKernelData
48 virtual ~IKernelData() =
default;
49 virtual void do_run(
ICLKernel &kernel, cl::CommandQueue &queue) = 0;
51 struct DefaultKernelData :
public CLTuner::IKernelData
57 ~DefaultKernelData()
override =
default;
58 void do_run(
ICLKernel &kernel, cl::CommandQueue &queue)
override 60 const bool inject_memory = !_tensors.
empty();
68 #if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION) 69 struct CompositeKernelData :
public CLTuner::IKernelData
72 : _tensors{ tensors }, _exec_desc{ exec_desc }
75 ~CompositeKernelData()
override =
default;
76 void do_run(
ICLKernel &kernel, cl::CommandQueue &queue)
override 86 #endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION) 90 return _kernel_event() !=
nullptr;
94 _kernel_event = kernel_event;
103 return _tune_new_kernels;
108 _tuning_info.tuner_mode =
mode;
122 void CLTuner::do_tune_kernel_dynamic(
ICLKernel &kernel, IKernelData *data)
128 if(kernel.
config_id() != arm_compute::default_config_id)
130 auto p = _tuning_params_table.find(config_id);
132 if(p == _tuning_params_table.end())
134 if(_tune_new_kernels)
137 CLTuningParams opt_tuning_params = find_optimal_tuning_params(kernel, data);
144 if(_tuning_info.tune_wbsm)
154 if(_tuning_info.tune_wbsm)
163 DefaultKernelData data{ tensors };
165 do_tune_kernel_dynamic(kernel, &data);
168 #if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION) 171 CompositeKernelData data{ tensors, exec_desc };
173 do_tune_kernel_dynamic(kernel, &data);
175 #endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION) 179 _tuning_params_table.emplace(kernel_id, optimal_tuning_params);
185 cl::CommandQueue queue_profiler;
197 cl_command_queue_properties props = default_queue.getInfo<CL_QUEUE_PROPERTIES>();
199 if((props & CL_QUEUE_PROFILING_ENABLE) == 0)
202 queue_profiler = cl::CommandQueue(
CLScheduler::get().context(), props | CL_QUEUE_PROFILING_ENABLE);
206 queue_profiler = default_queue;
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)
219 cl_int retval = this->
real_clEnqueueNDRangeKernel(command_queue, kernel, work_dim, gwo, gws, lws, num_events_in_wait_list, event_wait_list, &tmp);
237 data->do_run(kernel, queue_profiler);
239 queue_profiler.finish();
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;
250 for(
size_t i = 0; i < tuning_list->size(); ++i)
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);
265 kernel.set_lws_hint(lws_test);
268 cl_int wbsm_test = tuning_test.
get_wbsm();
269 kernel.set_wbsm_hint(wbsm_test);
273 data->do_run(kernel, queue_profiler);
275 queue_profiler.finish();
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;
283 if(diff < min_exec_time)
285 min_exec_time = diff;
287 if(_tuning_info.tune_wbsm)
296 return opt_tuning_params;
301 return _tuning_params_table;
306 _tuning_params_table.clear();
313 fs.exceptions(std::ifstream::badbit);
314 fs.open(filename, std::ios::in);
320 bool header_line =
true;
321 while(!std::getline(fs, line).fail())
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)
332 if(pos_wbsm != std::string::npos)
334 _tuning_info.tune_wbsm =
true;
338 if(std::getline(fs, line).fail())
346 size_t pos = line.find(
";");
347 if(pos == std::string::npos)
351 std::string kernel_id = line.substr(0, pos);
352 line.erase(0, pos + 1);
364 if(!_tune_new_kernels || _tuning_params_table.empty() || filename.empty())
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)
375 if(!header_string.empty())
377 header_string +=
" ";
379 header_string +=
"wbsm";
381 fs << header_string << std::endl;
382 for(
auto const &kernel_data : _tuning_params_table)
385 fs << kernel_data.first << tun_pams.
to_string(_tuning_info) << std::endl;
static cl::NDRange gws_from_window(const Window &window)
Get the global work size given an execution window.
void set_tuner_mode(CLTunerMode mode)
Set OpenCL tuner mode.
const Window & window() const
The maximum window the kernel can be executed on.
void set_cl_kernel_event(cl_event kernel_event)
Set the OpenCL kernel event.
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...
const std::unordered_map< std::string, CLTuningParams > & tuning_params_table() const
Give read access to the tuning params table.
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.
void tune_kernel_static(ICLKernel &kernel) override
Tune OpenCL kernel statically.
bool empty() const
Checks if pack is empty.
void tune_kernel_dynamic(ICLKernel &kernel) override
Tune OpenCL kernel dynamically.
void set_lws_hint(const cl::NDRange &lws_hint)
Set the Local-Workgroup-Size hint.
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...
void load_from_file(const std::string &filename)
Load the tuning parameters table from file.
#define ARM_COMPUTE_ERROR_VAR(msg,...)
Print the given message then throw an std::runtime_error.
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.
cl_int clRetainEvent(cl_event event)
Common interface for all the OpenCL kernels.
Copyright (c) 2017-2022 Arm Limited.
CLTuner(bool tune_new_kernels=true, CLTuningInfo tuning_info=CLTuningInfo())
Constructor.
std::function< decltype(clEnqueueNDRangeKernel)> real_clEnqueueNDRangeKernel
clEnqueueNDRangeKernel symbol
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.
bool kernel_event_is_set() const
Is the kernel_event set ?
Interface to enqueue OpenCL kernels and get/set the OpenCL CommandQueue and ICLTuner.
std::function< decltype(clEnqueueNDRangeKernel)> clEnqueueNDRangeKernel_ptr
Descriptor containing information required to run a single ClWorkload.
#define ARM_COMPUTE_UNUSED(...)
To avoid unused variables warnings.
bool save_to_file(const std::string &filename) const
Save the content of the tuning parameters table to file.
void add_tuning_params(const std::string &kernel_id, CLTuningParams optimal_tuning_params)
Manually add tuning parameters for a kernel.
void set_tune_new_kernels(bool tune_new_kernels)
Setter for tune_new_kernels option.
GPUTarget get_target() const
Get the targeted GPU architecture.
void set_wbsm(cl_int wbsm)
void end(TokenStream &in, bool &valid)
cl::CommandQueue & queue()
Accessor for the associated CL command queue.
std::string to_string(CLTuningInfo tuning_info)
CLTunerMode
< OpenCL tuner modes
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.
bool tune_new_kernels() const
Tune kernels that are not in the tuning parameters table.
bool from_string(CLTuningInfo tuning_info, std::string tuning_params_string)
static CLSymbols & get()
Get the static instance of CLSymbols.
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...