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/Framework.cpp | 3 + tests/framework/SConscript | 12 ++- tests/framework/instruments/Instruments.cpp | 1 + tests/framework/instruments/Instruments.h | 5 + tests/framework/instruments/OpenCLTimer.cpp | 143 ++++++++++++++++++++++++++++ tests/framework/instruments/OpenCLTimer.h | 63 ++++++++++++ 6 files changed, 226 insertions(+), 1 deletion(-) create mode 100644 tests/framework/instruments/OpenCLTimer.cpp create mode 100644 tests/framework/instruments/OpenCLTimer.h (limited to 'tests') diff --git a/tests/framework/Framework.cpp b/tests/framework/Framework.cpp index 82ea0fa93d..a5c665c458 100644 --- a/tests/framework/Framework.cpp +++ b/tests/framework/Framework.cpp @@ -50,6 +50,9 @@ Framework::Framework() #ifdef MALI_ENABLED _available_instruments.emplace(InstrumentType::MALI, Instrument::make_instrument); #endif /* MALI_ENABLED */ +#ifdef OPENCL_TIMER_ENABLED + _available_instruments.emplace(InstrumentType::OPENCL_TIMER, Instrument::make_instrument); +#endif /* OPENCL_TIMER_ENABLED */ } std::set Framework::available_instruments() const diff --git a/tests/framework/SConscript b/tests/framework/SConscript index 92cbd552bb..52b8bed10f 100644 --- a/tests/framework/SConscript +++ b/tests/framework/SConscript @@ -28,7 +28,8 @@ Import('vars') # vars is imported from arm_compute: variables = [ BoolVariable("pmu", "Enable PMU counters", False), - BoolVariable("mali", "Enable Mali hardware counters", False) + BoolVariable("mali", "Enable Mali hardware counters", False), + BoolVariable("opencl_timer", "Enable OpenCL timers", False) ] # We need a separate set of Variables for the Help message (Otherwise the global variables will get displayed twice) @@ -62,6 +63,15 @@ if not framework_env['pmu']: else: framework_env.Append(CPPDEFINES = ['PMU_ENABLED']) +if not framework_env['opencl_timer']: + # Remove OpenCLTimer files + files = [f for f in files if "OpenCLTimer" not in os.path.basename(str(f))] +else: + if not framework_env["opencl"]: + print("ERROR: You need opencl=1 to be able to use opencl_timer=1") + Exit(1) + framework_env.Append(CPPDEFINES = ['OPENCL_TIMER_ENABLED']) + if not framework_env['mali']: # Remove MALI files files = [f for f in files if "MaliCounter" not in os.path.basename(str(f))] diff --git a/tests/framework/instruments/Instruments.cpp b/tests/framework/instruments/Instruments.cpp index 1c21f0773d..eca6d6f437 100644 --- a/tests/framework/instruments/Instruments.cpp +++ b/tests/framework/instruments/Instruments.cpp @@ -45,6 +45,7 @@ InstrumentType instrument_type_from_name(const std::string &name) { "pmu_cycles", InstrumentType::PMU_CYCLE_COUNTER }, { "pmu_instructions", InstrumentType::PMU_INSTRUCTION_COUNTER }, { "mali", InstrumentType::MALI }, + { "opencl_timer", InstrumentType::OPENCL_TIMER }, }; try diff --git a/tests/framework/instruments/Instruments.h b/tests/framework/instruments/Instruments.h index df6aa62ec7..ffe7cb681d 100644 --- a/tests/framework/instruments/Instruments.h +++ b/tests/framework/instruments/Instruments.h @@ -25,6 +25,7 @@ #define ARM_COMPUTE_TEST_INSTRUMENTS #include "MaliCounter.h" +#include "OpenCLTimer.h" #include "PMUCounter.h" #include "WallClockTimer.h" @@ -46,6 +47,7 @@ enum class InstrumentType : unsigned int PMU_CYCLE_COUNTER = 0x0201, PMU_INSTRUCTION_COUNTER = 0x0202, MALI = 0x0300, + OPENCL_TIMER = 0x0400, }; InstrumentType instrument_type_from_name(const std::string &name); @@ -77,6 +79,9 @@ inline ::std::stringstream &operator<<(::std::stringstream &stream, InstrumentTy case InstrumentType::MALI: stream << "MALI"; break; + case InstrumentType::OPENCL_TIMER: + stream << "OPENCL_TIMER"; + break; case InstrumentType::ALL: stream << "ALL"; break; 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 diff --git a/tests/framework/instruments/OpenCLTimer.h b/tests/framework/instruments/OpenCLTimer.h new file mode 100644 index 0000000000..d00a2e8bf2 --- /dev/null +++ b/tests/framework/instruments/OpenCLTimer.h @@ -0,0 +1,63 @@ +/* + * 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. + */ +#ifndef ARM_COMPUTE_TEST_OPENCL_TIMER +#define ARM_COMPUTE_TEST_OPENCL_TIMER + +#include "Instrument.h" + +#ifdef ARM_COMPUTE_CL +#include "arm_compute/core/CL/OpenCL.h" +#endif /* ARM_COMPUTE_CL */ + +#include + +namespace arm_compute +{ +namespace test +{ +namespace framework +{ +/** Instrument creating measurements based on the information returned by clGetEventProfilingInfo for each OpenCL kernel executed*/ +class OpenCLTimer : public Instrument +{ +public: + OpenCLTimer(); + std::string id() const override; + void start() override; + void stop() override; + MeasurementsMap measurements() const override; +#ifdef ARM_COMPUTE_CL + struct kernel_info + { + cl::Event event{}; /**< OpenCL event associated to the kernel enqueue */ + std::string name{}; /**< OpenCL Kernel name */ + }; + std::list kernels{}; + std::function real_function; +#endif /* ARM_COMPUTE_CL */ +}; +} // namespace framework +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_OPENCL_TIMER */ -- cgit v1.2.1