Timing a kernel with clock_gettime on the host gives the wrong answer, because a launch is asynchronous and returns before the GPU has done anything. Events are markers recorded in a stream that the device timestamps as it reaches them.
cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start); kernel<<<blocks, threads>>>(d_a); cudaEventRecord(stop); cudaEventSynchronize(stop); // wait until the device reaches the marker float ms = 0.0f; cudaEventElapsedTime(&ms, start, stop); printf("%.3f ms\n", ms);
cudaEventElapsedTime returns milliseconds as a float, with resolution around half a microsecond. The cudaEventSynchronize is what makes the number valid, since the timestamp does not exist until the device has actually passed the marker.
Events are recorded into a stream, so they measure work in that stream only. Timing a specific stream while others run concurrently works as expected, which is the main advantage over cudaDeviceSynchronize plus a host timer.
Two things skew results. The first launch in a program includes CUDA context creation, which can take a hundred milliseconds or more, so a warm-up launch before timing is standard. The second is that a single run of a short kernel is dominated by launch overhead of a few microseconds, so timing a loop of many launches and dividing is more reliable.
cudaEventQuery(ev) tests without blocking, returning cudaSuccess when the event has been reached and cudaErrorNotReady otherwise. That is useful for overlapping host work with device work rather than sitting in a synchronise.
For anything beyond a single number, nsys gives a timeline of kernels and transfers, and ncu gives per-kernel counters. Events answer “how long did this take”; the profilers answer “and why”.