2 * Copyright (c) 2017-2018 ARM Limited.
4 * SPDX-License-Identifier: MIT
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:
13 * The above copyright notice and this permission notice shall be included in all
14 * copies or substantial portions of the Software.
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
24 #include "OpenCLTimer.h"
26 #include "../Framework.h"
29 #include "arm_compute/graph/INode.h"
30 #include "arm_compute/runtime/CL/CLScheduler.h"
32 #ifndef ARM_COMPUTE_CL
33 #error "You can't use OpenCLTimer without OpenCL"
34 #endif /* ARM_COMPUTE_CL */
42 std::string OpenCLTimer::id() const
47 OpenCLTimer::OpenCLTimer(ScaleFactor scale_factor)
48 : _kernels(), _real_function(nullptr), _real_graph_function(nullptr), _prefix(), _timer_enabled(false)
50 auto q = CLScheduler::get().queue();
51 cl_command_queue_properties props = q.getInfo<CL_QUEUE_PROPERTIES>();
52 if((props & CL_QUEUE_PROFILING_ENABLE) == 0)
54 CLScheduler::get().set_queue(cl::CommandQueue(CLScheduler::get().context(), props | CL_QUEUE_PROFILING_ENABLE));
59 case ScaleFactor::NONE:
63 case ScaleFactor::TIME_US:
64 _scale_factor = 1000.f;
67 case ScaleFactor::TIME_MS:
68 _scale_factor = 1000000.f;
71 case ScaleFactor::TIME_S:
72 _scale_factor = 1000000000.f;
76 ARM_COMPUTE_ERROR("Invalid scale");
80 void OpenCLTimer::test_start()
82 // Start intercepting enqueues:
83 ARM_COMPUTE_ERROR_ON(_real_function != nullptr);
84 ARM_COMPUTE_ERROR_ON(_real_graph_function != nullptr);
85 _real_function = CLSymbols::get().clEnqueueNDRangeKernel_ptr;
86 _real_graph_function = graph::TaskExecutor::get().execute_function;
87 auto interceptor = [this](
88 cl_command_queue command_queue,
94 cl_uint num_events_in_wait_list,
95 const cl_event * event_wait_list,
98 if(this->_timer_enabled)
100 ARM_COMPUTE_ERROR_ON_MSG(event != nullptr, "Not supported");
101 ARM_COMPUTE_UNUSED(event);
103 OpenCLTimer::kernel_info info;
104 cl::Kernel cpp_kernel(kernel, true);
105 std::stringstream ss;
106 ss << this->_prefix << cpp_kernel.getInfo<CL_KERNEL_FUNCTION_NAME>();
109 ss << " GWS[" << gws[0] << "," << gws[1] << "," << gws[2] << "]";
113 ss << " LWS[" << lws[0] << "," << lws[1] << "," << lws[2] << "]";
115 info.name = ss.str();
117 cl_int retval = this->_real_function(command_queue, kernel, work_dim, gwo, gws, lws, num_events_in_wait_list, event_wait_list, &tmp);
119 this->_kernels.push_back(std::move(info));
124 return this->_real_function(command_queue, kernel, work_dim, gwo, gws, lws, num_events_in_wait_list, event_wait_list, event);
128 // Start intercepting tasks:
129 auto task_interceptor = [this](graph::ExecutionTask & task)
131 if(task.node != nullptr && !task.node->name().empty())
133 this->_prefix = task.node->name() + "/";
139 this->_real_graph_function(task);
143 CLSymbols::get().clEnqueueNDRangeKernel_ptr = interceptor;
144 graph::TaskExecutor::get().execute_function = task_interceptor;
147 void OpenCLTimer::start()
150 _timer_enabled = true;
152 void OpenCLTimer::stop()
154 _timer_enabled = false;
157 void OpenCLTimer::test_stop()
159 // Restore real function
160 CLSymbols::get().clEnqueueNDRangeKernel_ptr = _real_function;
161 graph::TaskExecutor::get().execute_function = _real_graph_function;
162 _real_graph_function = nullptr;
163 _real_function = nullptr;
166 Instrument::MeasurementsMap OpenCLTimer::measurements() const
168 MeasurementsMap measurements;
169 unsigned int kernel_number = 0;
170 for(auto kernel : _kernels)
172 cl_ulong start = kernel.event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
173 cl_ulong end = kernel.event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
175 measurements.emplace(kernel.name + " #" + support::cpp11::to_string(kernel_number++), Measurement((end - start) / _scale_factor, _unit));
180 } // namespace framework
182 } // namespace arm_compute