From 6e4334962913a4b4d88d3d2b9d76ef2b7850a56c Mon Sep 17 00:00:00 2001 From: Anthony Barbier Date: Thu, 9 Nov 2017 15:52:00 +0000 Subject: COMPMID-663 Adding instrument to return timestamps of clGetEventProfilingInfo Change-Id: I1037054615593205f07e25fb9b16fecd13407c2c Reviewed-on: http://mpd-gerrit.cambridge.arm.com/95142 Tested-by: Kaizen Reviewed-by: Georgios Pinitas --- tests/framework/instruments/OpenCLTimer.cpp | 143 ++++++++++++++++++++++++++++ 1 file changed, 143 insertions(+) create mode 100644 tests/framework/instruments/OpenCLTimer.cpp (limited to 'tests/framework/instruments/OpenCLTimer.cpp') diff --git a/tests/framework/instruments/OpenCLTimer.cpp b/tests/framework/instruments/OpenCLTimer.cpp new file mode 100644 index 0000000000..819ecdb743 --- /dev/null +++ b/tests/framework/instruments/OpenCLTimer.cpp @@ -0,0 +1,143 @@ +/* + * Copyright (c) 2017 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "OpenCLTimer.h" + +#include "../Framework.h" +#include "../Utils.h" + +#include "arm_compute/runtime/CL/CLScheduler.h" + +#ifndef ARM_COMPUTE_CL +#error "You can't use OpenCLTimer without OpenCL" +#endif /* ARM_COMPUTE_CL */ + +namespace arm_compute +{ +namespace test +{ +namespace framework +{ +std::string OpenCLTimer::id() const +{ + return "OpenCLTimer"; +} + +/* Function to be used to intercept kernel enqueues and store their OpenCL Event */ +class Interceptor +{ +public: + explicit Interceptor(OpenCLTimer &timer) + : _timer(timer) + { + } + + 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) + { + ARM_COMPUTE_ERROR_ON_MSG(event != nullptr, "Not supported"); + ARM_COMPUTE_UNUSED(event); + + OpenCLTimer::kernel_info info; + cl::Kernel cpp_kernel(kernel, true); + std::stringstream ss; + ss << cpp_kernel.getInfo(); + if(gws != nullptr) + { + ss << " GWS[" << gws[0] << "," << gws[1] << "," << gws[2] << "]"; + } + if(lws != nullptr) + { + ss << " LWS[" << lws[0] << "," << lws[1] << "," << lws[2] << "]"; + } + info.name = ss.str(); + cl_event tmp; + cl_int retval = _timer.real_function(command_queue, kernel, work_dim, gwo, gws, lws, num_events_in_wait_list, event_wait_list, &tmp); + info.event = tmp; + _timer.kernels.push_back(std::move(info)); + return retval; + } + +private: + OpenCLTimer &_timer; +}; + +OpenCLTimer::OpenCLTimer() + : real_function(CLSymbols::get().clEnqueueNDRangeKernel_ptr) +{ + auto q = CLScheduler::get().queue(); + cl_command_queue_properties props = q.getInfo(); + if((props & CL_QUEUE_PROFILING_ENABLE) == 0) + { + CLScheduler::get().set_queue(cl::CommandQueue(CLScheduler::get().context(), props | CL_QUEUE_PROFILING_ENABLE)); + } +} + +void OpenCLTimer::start() +{ + kernels.clear(); + // Start intercepting enqueues: + CLSymbols::get().clEnqueueNDRangeKernel_ptr = Interceptor(*this); +} + +void OpenCLTimer::stop() +{ + // Restore real function + CLSymbols::get().clEnqueueNDRangeKernel_ptr = real_function; +} + +Instrument::MeasurementsMap OpenCLTimer::measurements() const +{ + MeasurementsMap measurements; + unsigned int kernel_number = 0; + for(auto kernel : kernels) + { + //cl_int status = kernel.event.getInfo((); + cl_ulong queued = kernel.event.getProfilingInfo(); + cl_ulong submit = kernel.event.getProfilingInfo(); + cl_ulong start = kernel.event.getProfilingInfo(); + cl_ulong end = kernel.event.getProfilingInfo(); + + std::list raw_data = + { + "queued", support::cpp11::to_string(queued), + "submit", support::cpp11::to_string(submit), + "start", support::cpp11::to_string(start), + "end", support::cpp11::to_string(end), + }; + measurements.emplace(kernel.name + " #" + support::cpp11::to_string(kernel_number++), Measurement(end - start, "ns", raw_data)); + } + + return measurements; +} +} // namespace framework +} // namespace test +} // namespace arm_compute -- cgit v1.2.1