1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
|
//
// Copyright © 2017 Arm Ltd. All rights reserved.
// SPDX-License-Identifier: MIT
//
#include "OpenClTimer.hpp"
#include <string>
#include <sstream>
namespace armnn
{
OpenClTimer::OpenClTimer()
{
}
void OpenClTimer::Start()
{
m_Kernels.clear();
auto interceptor = [this]( 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)
{
cl_int retVal = 0;
// Get the name of the kernel
cl::Kernel retainedKernel(kernel, true);
std::stringstream ss;
ss << retainedKernel.getInfo<CL_KERNEL_FUNCTION_NAME>();
// Embed workgroup sizes into the name
if(gws != nullptr)
{
ss << " GWS[" << gws[0] << "," << gws[1] << "," << gws[2] << "]";
}
if(lws != nullptr)
{
ss << " LWS[" << lws[0] << "," << lws[1] << "," << lws[2] << "]";
}
cl_event customEvent;
// Forward to original OpenCl function
retVal = m_OriginalEnqueueFunction( command_queue,
kernel,
work_dim,
gwo,
gws,
lws,
num_events_in_wait_list,
event_wait_list,
&customEvent);
// Store the Kernel info for later GetMeasurements() call
m_Kernels.emplace_back(ss.str(), customEvent);
return retVal;
};
m_OriginalEnqueueFunction = CLSymbols::get().clEnqueueNDRangeKernel_ptr;
CLSymbols::get().clEnqueueNDRangeKernel_ptr = interceptor;
}
void OpenClTimer::Stop()
{
CLSymbols::get().clEnqueueNDRangeKernel_ptr = m_OriginalEnqueueFunction;
}
std::vector<Measurement> OpenClTimer::GetMeasurements() const
{
std::vector<Measurement> measurements;
cl_command_queue_properties clQueueProperties = CLScheduler::get().queue().getInfo<CL_QUEUE_PROPERTIES>();
int idx = 0;
for (auto& kernel : m_Kernels)
{
std::string name = std::string(this->GetName()) + "/" + std::to_string(idx++) + ": " + kernel.m_Name;
double timeUs = 0.0;
if((clQueueProperties & CL_QUEUE_PROFILING_ENABLE) != 0)
{
// Wait for the event to finish before accessing profile results.
kernel.m_Event.wait();
cl_ulong start = kernel.m_Event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
cl_ulong end = kernel.m_Event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
timeUs = static_cast<double>(end - start) / 1000.0;
}
measurements.emplace_back(name, timeUs, Measurement::Unit::TIME_US);
}
return measurements;
}
} //namespace armnn
|
