diff options
Diffstat (limited to 'tests/framework/instruments')
-rw-r--r-- | tests/framework/instruments/Instruments.cpp | 1 | ||||
-rw-r--r-- | tests/framework/instruments/Instruments.h | 5 | ||||
-rw-r--r-- | tests/framework/instruments/OpenCLTimer.cpp | 143 | ||||
-rw-r--r-- | tests/framework/instruments/OpenCLTimer.h | 63 |
4 files changed, 212 insertions, 0 deletions
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<CL_KERNEL_FUNCTION_NAME>(); + 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<CL_QUEUE_PROPERTIES>(); + 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_EVENT_COMMAND_EXECUTION_STATUS>(); + cl_ulong queued = kernel.event.getProfilingInfo<CL_PROFILING_COMMAND_QUEUED>(); + cl_ulong submit = kernel.event.getProfilingInfo<CL_PROFILING_COMMAND_SUBMIT>(); + cl_ulong start = kernel.event.getProfilingInfo<CL_PROFILING_COMMAND_START>(); + cl_ulong end = kernel.event.getProfilingInfo<CL_PROFILING_COMMAND_END>(); + + std::list<std::string> 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 <list> + +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<kernel_info> kernels{}; + std::function<decltype(clEnqueueNDRangeKernel)> real_function; +#endif /* ARM_COMPUTE_CL */ +}; +} // namespace framework +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_OPENCL_TIMER */ |