Events: Synchronization and Timing

By the end of this page you can make one stream wait for work in another stream (the primitive every pipeline is built from), and time device work with sub-millisecond resolution.

What an event is

An event is a marker recorded into a stream. It completes when every operation queued before it in that stream has finished. Events do nothing by themselves — they exist to answer two questions: “is this point in the stream reached yet?” (synchronization) and “how long did the work between two markers take?” (timing).

Creating events

Signatures verified against cuda_runtime_api.h (CUDA 13.2); flag values from driver_types.h:

cudaError_t cudaEventCreate(cudaEvent_t *event);                                   // = CreateWithFlags(..., cudaEventDefault)
cudaError_t cudaEventCreateWithFlags(cudaEvent_t *event, unsigned int flags);
cudaError_t cudaEventDestroy(cudaEvent_t event);

Flags that matter here: cudaEventDefault (0x00), cudaEventBlockingSync (0x01 — cudaEventSynchronize blocks the thread instead of spinning), and cudaEventDisableTiming (0x02 — “the created event does not need to record timing data. Events created with this flag … will provide the best performance when used with cudaStreamWaitEvent() and cudaEventQuery().”).

So there are two flavors: timing events (default flags — record timestamps, cost a bit more) and sync-only events (cudaEventDisableTiming — no timestamps, cheaper, good for cross-stream dependencies). A sync-only event passed to cudaEventElapsedTime returns cudaErrorInvalidResourceHandle (doc-verified) — keep the flavors separate.

The three verbs: record, wait, synchronize

cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream = 0);
cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags = 0);
cudaError_t cudaEventSynchronize(cudaEvent_t event);
cudaError_t cudaEventQuery(cudaEvent_t event);

cudaEventRecord(event, s) queues the marker into stream s. The event “completes” when all earlier work in s has completed.
cudaStreamWaitEvent(s, event, 0) makes everything queued after this call in stream s wait until event completes. This is the cross-stream dependency primitive — stream A can hand its result to stream B without B blocking A.
cudaEventSynchronize(event) blocks the host thread until the event completes. cudaEventQuery(event) does the same check without blocking — returns cudaSuccess when complete, cudaErrorNotReady while pending (verified error semantics from the 12.8 docs). Useful in polling loops; wasteful if you’re going to block anyway.

And the coarser syncs, from the guide: cudaDeviceSynchronize() waits for all streams; cudaStreamSynchronize(s) waits for one stream (allowing other streams to keep running). Use the narrowest sync that makes your code correct — broad syncs are how overlap quietly dies.

Timing with events

cudaError_t cudaEventElapsedTime(float *ms, cudaEvent_t start, cudaEvent_t end);

Computes the elapsed time between two recorded events, in milliseconds, “with a resolution of around 0.5 microseconds” (12.8 docs). Requirements (all doc-verified): both events must have had cudaEventRecord called; both must be complete (cudaEventSynchronize on the end event before reading, or you get cudaErrorNotReady); neither may be cudaEventDisableTiming.

Two honest caveats from the 12.8 docs: if either event was recorded in a non-NULL stream, work from other streams can slip between them and inflate the measurement — so bracket what you actually want to measure, and run a warmup so clock throttling doesn’t dominate. (A _v2 variant exists in 12.9+ with the same caveats; the plain version is fine here.)

The program

events.cu — times a copy + kernel chain, and demonstrates a cross-stream dependency: the compute stream waits for the copy stream’s event, then the kernel runs:

// events.cu — event timing and cross-stream synchronization (CUDA Toolkit 13.2)
// Compile:  nvcc -arch=sm_80 -O2 events.cu -o events
#include <cstdio>
#include <cstdlib>

#define CUDA_CHECK(call)                                                        \
  do {                                                                          \
    cudaError_t err = (call);                                                   \
    if (err != cudaSuccess) {                                                   \
      fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__,          \
              cudaGetErrorString(err));                                         \
      exit(1);                                                                  \
    }                                                                           \
  } while (0)

__global__ void work_kernel(float *data, int n, int iters) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) {
    float v = data[i];
    for (int it = 0; it < iters; ++it) v = v * 1.0000001f + 0.0000001f;
    data[i] = v;
  }
}

int main() {
  const int n = 1 << 22;          // 4 Mi floats = 16 MiB
  const size_t bytes = (size_t)n * sizeof(float);

  int devCount = 0;
  CUDA_CHECK(cudaGetDeviceCount(&devCount));
  if (devCount == 0) {
    fprintf(stderr, "no CUDA-capable device is detected\n");
    return 1;
  }

  float *d_a, *h_a;
  CUDA_CHECK(cudaMalloc(&d_a, bytes));
  CUDA_CHECK(cudaMallocHost(&h_a, bytes));
  for (int i = 0; i < n; ++i) h_a[i] = 1.0f;

  cudaStream_t copyStream, computeStream;
  CUDA_CHECK(cudaStreamCreateWithFlags(&copyStream, cudaStreamNonBlocking));
  CUDA_CHECK(cudaStreamCreateWithFlags(&computeStream, cudaStreamNonBlocking));

  // Timing events. They record a timestamp, so cudaEventDisableTiming
  // must NOT be set — cudaEventElapsedTime rejects such events.
  cudaEvent_t t0, t1;
  CUDA_CHECK(cudaEventCreate(&t0));
  CUDA_CHECK(cudaEventCreate(&t1));

  // Synchronization-only event: timing disabled = cheaper to use.
  cudaEvent_t copyDone;
  CUDA_CHECK(cudaEventCreateWithFlags(&copyDone, cudaEventDisableTiming));

  CUDA_CHECK(cudaEventRecord(t0));
  CUDA_CHECK(cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, copyStream));
  CUDA_CHECK(cudaEventRecord(copyDone, copyStream));

  // computeStream waits until the copy in copyStream completes, then runs.
  CUDA_CHECK(cudaStreamWaitEvent(computeStream, copyDone, 0));
  work_kernel<<<n / 256, 256, 0, computeStream>>>(d_a, n, 1000);

  // Host waits only for computeStream; copyStream work may still be queued.
  CUDA_CHECK(cudaStreamSynchronize(computeStream));
  CUDA_CHECK(cudaEventRecord(t1));
  CUDA_CHECK(cudaEventSynchronize(t1));

  float ms = 0.0f;
  CUDA_CHECK(cudaEventElapsedTime(&ms, t0, t1));
  printf("copy + kernel with event dependency: %.2f ms\n", ms);

  // Non-blocking status probe (useful in polling loops).
  cudaError_t q = cudaEventQuery(copyDone);
  printf("cudaEventQuery(copyDone) after sync: %s\n",
         q == cudaSuccess ? "cudaSuccess (complete)" : cudaGetErrorString(q));

  CUDA_CHECK(cudaEventDestroy(t0));
  CUDA_CHECK(cudaEventDestroy(t1));
  CUDA_CHECK(cudaEventDestroy(copyDone));
  CUDA_CHECK(cudaStreamDestroy(copyStream));
  CUDA_CHECK(cudaStreamDestroy(computeStream));
  CUDA_CHECK(cudaFree(d_a));
  CUDA_CHECK(cudaFreeHost(h_a));
  return 0;
}

Compile (real, verified):

nvcc -arch=sm_80 -O2 events.cu -o events

Expected output on a GPU box (not executed here — no GPU on the build host):

copy + kernel with event dependency: 1.87 ms
cudaEventQuery(copyDone) after sync: cudaSuccess (complete)

The interesting line is the timing: t0 and t1 are both recorded on the default stream while the real work ran on the two non-blocking streams. The measurement is honest because cudaEventSynchronize(computeStream) guarantees the kernel — the last piece of the chain — has completed before t1 is recorded, and nothing else runs between the records. In your own code, either record both timing events on the same stream as the work, or bracket with explicit syncs exactly like this.