# Events (CUDA)
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 [[streams-cuda|stream]] that the device timestamps as it reaches them.

```c
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".