The samples time host-side intervals with the C++ standard library’s
std::chrono::steady_clock, a monotonic, high-resolution
clock. An example usage is as follows:
float
TimeNULLKernelLaunches(int cIterations = 1000000 )
{
std::chrono::steady_clock::time_point start, stop;
start = std::chrono::steady_clock::now();
for ( int i = 0; i < cIterations; i++ ) {
NullKernel<<<1,1>>>();
}
cudaDeviceSynchronize();
stop = std::chrono::steady_clock::now();
return 1e6*std::chrono::duration(stop - start).count() /
(float) cIterations;
}
This function times the specified number of kernel launches and
returns the microseconds per launch. A
std::chrono::steady_clock::time_point holds a
high-resolution timestamp; since it is a single point in time, two of
them are needed to compute an interval.
std::chrono::steady_clock::now() takes a snapshot of the
current time. Subtracting two time points and wrapping the difference in
std::chrono::duration<double> gives the elapsed
seconds; .count() extracts it as a double. The
clock is monotonic – it never runs backward when the system clock is
adjusted – and its resolution is very fine (nanoseconds of
representation, tens of nanoseconds in practice), so elapsed times are
kept in double precision.
We may use CUDA events when measuring performance in isolation on the CUDA-capable GPU, such as when measuring the device memory bandwidth of a kernel. Using CUDA events for timing is a two-edged sword: they are less affected by spurious system-level events, such as network traffic, but that sometimes can lead to overly optimistic timing results.