#include "arm_compute/core/Error.h"
#include "arm_compute/runtime/CL/CLScheduler.h"
+#include <cerrno>
+#include <fstream>
+#include <iostream>
#include <limits>
#include <string>
using namespace arm_compute;
-CLTuner::CLTuner()
- : real_function(nullptr), _lws_table(), _queue(), _queue_profiler(), _kernel_event()
+namespace
{
-}
+/* Function to be used to intercept kernel enqueues and store their OpenCL Event */
+class Interceptor
+{
+public:
+ explicit Interceptor(CLTuner &tuner);
+
+ /** clEnqueueNDRangeKernel interface
+ *
+ * @param[in] command_queue A valid command-queue. The kernel will be queued for execution on the device associated with command_queue.
+ * @param[in] kernel A valid kernel object. The OpenCL context associated with kernel and command_queue must be the same.
+ * @param[in] work_dim The number of dimensions used to specify the global work-items and work-items in the work-group. work_dim must be greater than zero and less than or equal to CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS.
+ * @param[in] gwo Global-Workgroup-Offset. It can be used to specify an array of work_dim unsigned values that describe the offset used to calculate the global ID of a work-item. If global_work_offset is NULL, the global IDs start at offset (0, 0, ... 0).
+ * @param[in] gws Global-Workgroup-Size. Points to an array of work_dim unsigned values that describe the number of global work-items in work_dim dimensions that will execute the kernel function.
+ * @param[in] lws Local-Workgroup-Size. Points to an array of work_dim unsigned values that describe the number of work-items that make up a work-group
+ * @param[in] num_events_in_wait_list Number of events in the waiting list
+ * @param[in] event_wait_list Event waiting list
+ * @param[in] event OpenCL kernel event
+ *
+ * @return the OpenCL status
+ */
+ cl_int operator()(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,
+ const cl_event *event_wait_list, cl_event *event);
+
+private:
+ CLTuner &_tuner;
+};
-void CLTuner::set_cl_kernel_event(cl_event kernel_event)
+Interceptor::Interceptor(CLTuner &tuner)
+ : _tuner(tuner)
{
- _kernel_event = kernel_event;
}
-void CLTuner::tune_kernel(ICLKernel &kernel)
+cl_int Interceptor::operator()(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,
+ const cl_event *event_wait_list, cl_event *event)
{
- if(real_function == nullptr)
+ ARM_COMPUTE_ERROR_ON_MSG(event != nullptr, "Not supported");
+ ARM_COMPUTE_UNUSED(event);
+ if(_tuner.kernel_event_is_set())
{
- real_function = CLSymbols::get().clEnqueueNDRangeKernel_ptr;
+ // 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.
+ return CL_SUCCESS;
+ }
+ cl_event tmp;
+ cl_int retval = _tuner.real_clEnqueueNDRangeKernel(command_queue, kernel, work_dim, gwo, gws, lws, num_events_in_wait_list, event_wait_list, &tmp);
- // Get the default queue
- _queue = CLScheduler::get().queue();
+ // Set OpenCL event
+ _tuner.set_cl_kernel_event(tmp);
- // Check if we can use the OpenCL timer with the default queue
- cl_command_queue_properties props = _queue.getInfo<CL_QUEUE_PROPERTIES>();
+ return retval;
+}
- if((props & CL_QUEUE_PROFILING_ENABLE) == 0)
- {
- // Set the queue for profiling
- _queue_profiler = cl::CommandQueue(CLScheduler::get().context(), props | CL_QUEUE_PROFILING_ENABLE);
- }
- else
- {
- _queue_profiler = _queue;
- }
- }
+} // namespace
+CLTuner::CLTuner(bool tune_new_kernels)
+ : real_clEnqueueNDRangeKernel(nullptr), _lws_table(), _queue(), _queue_profiler(), _kernel_event(), _tune_new_kernels(tune_new_kernels)
+{
+}
+
+bool CLTuner::kernel_event_is_set() const
+{
+ return _kernel_event() != nullptr;
+}
+void CLTuner::set_cl_kernel_event(cl_event kernel_event)
+{
+ _kernel_event = kernel_event;
+}
+
+void CLTuner::set_tune_new_kernels(bool tune_new_kernels)
+{
+ _tune_new_kernels = tune_new_kernels;
+}
+bool CLTuner::tune_new_kernels() const
+{
+ return _tune_new_kernels;
+}
+
+void CLTuner::tune_kernel(ICLKernel &kernel)
+{
// Get the configuration ID from the kernel
const std::string &config_id = kernel.config_id();
if(p == _lws_table.end())
{
- // Set profiler queue
- CLScheduler::get().set_queue(_queue_profiler);
-
- // Find the optimal LWS for the kernel
- cl::NDRange opt_lws = find_optimal_lws(kernel);
-
- // Insert the optimal LWS in the table
- _lws_table.emplace(config_id, opt_lws);
+ if(_tune_new_kernels)
+ {
+ // Find the optimal LWS for the kernel
+ cl::NDRange opt_lws = find_optimal_lws(kernel);
- // Set Local-Workgroup-Size
- kernel.set_lws_hint(opt_lws);
+ // Insert the optimal LWS in the table
+ add_lws_to_table(config_id, opt_lws);
- // Restore queue
- CLScheduler::get().set_queue(_queue);
+ // Set Local-Workgroup-Size
+ kernel.set_lws_hint(opt_lws);
+ }
}
else
{
}
}
+void CLTuner::add_lws_to_table(const std::string &kernel_id, cl::NDRange optimal_lws)
+{
+ _lws_table.emplace(kernel_id, optimal_lws);
+}
+
cl::NDRange CLTuner::find_optimal_lws(ICLKernel &kernel)
{
+ if(real_clEnqueueNDRangeKernel == nullptr)
+ {
+ real_clEnqueueNDRangeKernel = CLSymbols::get().clEnqueueNDRangeKernel_ptr;
+
+ // Get the default queue
+ _queue = CLScheduler::get().queue();
+
+ // Check if we can use the OpenCL timer with the default queue
+ cl_command_queue_properties props = _queue.getInfo<CL_QUEUE_PROPERTIES>();
+
+ if((props & CL_QUEUE_PROFILING_ENABLE) == 0)
+ {
+ // Set the queue for profiling
+ _queue_profiler = cl::CommandQueue(CLScheduler::get().context(), props | CL_QUEUE_PROFILING_ENABLE);
+ }
+ else
+ {
+ _queue_profiler = _queue;
+ }
+ }
// Start intercepting enqueues:
CLSymbols::get().clEnqueueNDRangeKernel_ptr = Interceptor(*this);
// Run the kernel
kernel.run(kernel.window(), _queue_profiler);
- CLScheduler::get().sync();
+ _queue_profiler.finish();
const cl_ulong start = _kernel_event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
const cl_ulong end = _kernel_event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
const cl_ulong diff = end - start;
+ _kernel_event = nullptr;
min_exec_time = diff;
}
// Run the kernel
kernel.run(kernel.window(), _queue_profiler);
- CLScheduler::get().sync();
+ _queue_profiler.finish();
const cl_ulong start = _kernel_event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
const cl_ulong end = _kernel_event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
const cl_ulong diff = end - start;
+ _kernel_event = nullptr;
// Check the execution time
if(diff < min_exec_time)
}
// Restore real function
- CLSymbols::get().clEnqueueNDRangeKernel_ptr = real_function;
+ CLSymbols::get().clEnqueueNDRangeKernel_ptr = real_clEnqueueNDRangeKernel;
return opt_lws;
}
_lws_table = lws_table;
}
-const std::unordered_map<std::string, cl::NDRange> &CLTuner::export_lws_table()
+const std::unordered_map<std::string, cl::NDRange> &CLTuner::lws_table() const
{
return _lws_table;
}
-Interceptor::Interceptor(CLTuner &tuner)
- : _tuner(tuner)
+void CLTuner::load_from_file(const std::string &filename)
{
+ std::ifstream fs;
+ fs.exceptions(std::ifstream::badbit);
+ fs.open(filename, std::ios::in);
+ if(!fs.is_open())
+ {
+ ARM_COMPUTE_ERROR("Failed to open '%s' (%s [%d])", filename.c_str(), strerror(errno), errno);
+ }
+ std::string line;
+ while(!std::getline(fs, line).fail())
+ {
+ std::istringstream ss(line);
+ std::string token;
+ if(std::getline(ss, token, ';').fail())
+ {
+ ARM_COMPUTE_ERROR("Malformed row '%s' in %s (Should be of the form 'kernel_id;lws[0];lws[1];lws[2]')", ss.str().c_str(), filename.c_str());
+ }
+ std::string kernel_id = token;
+ cl::NDRange lws(1, 1, 1);
+ for(int i = 0; i < 3; i++)
+ {
+ if(std::getline(ss, token, ';').fail())
+ {
+ ARM_COMPUTE_ERROR("Malformed row '%s' in %s (Should be of the form 'kernel_id;lws[0];lws[1];lws[2]')", ss.str().c_str(), filename.c_str());
+ }
+ lws.get()[i] = support::cpp11::stoi(token);
+ }
+
+ // If all dimensions are 0: reset to NullRange (i.e nullptr)
+ if(lws[0] == 0 && lws[1] == 0 && lws[2] == 0)
+ {
+ lws = cl::NullRange;
+ }
+ add_lws_to_table(kernel_id, lws);
+ }
+ fs.close();
}
-cl_int Interceptor::operator()(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,
- const cl_event *event_wait_list, cl_event *event)
+void CLTuner::save_to_file(const std::string &filename) const
{
- ARM_COMPUTE_ERROR_ON_MSG(event != nullptr, "Not supported");
- ARM_COMPUTE_UNUSED(event);
-
- cl_event tmp;
- cl_int retval = _tuner.real_function(command_queue, kernel, work_dim, gwo, gws, lws, num_events_in_wait_list, event_wait_list, &tmp);
-
- // Set OpenCL event
- _tuner.set_cl_kernel_event(tmp);
-
- return retval;
+ std::ofstream fs;
+ fs.exceptions(std::ifstream::failbit | std::ifstream::badbit);
+ fs.open(filename, std::ios::out);
+ for(auto kernel_data : _lws_table)
+ {
+ fs << kernel_data.first << ";" << kernel_data.second[0] << ";" << kernel_data.second[1] << ";" << kernel_data.second[2] << std::endl;
+ }
+ fs.close();
}
projects / platform / upstream / armcl.git / blobdiff