100 DAYS OF GPU Course preview
Day 52Module 6
in-technical-review

Events and cross-stream dependencies

Here are two programs. Both submit one producer kernel and four consumer kernels across five streams. Both record one fork event with cudaEventRecord and issue four cudaStreamWaitEvent calls against it.

Every call in both returns cudaSuccess.

One of them has a real dependency between the producer and the consumers. The other has none at all, and the only difference is which line the record sits on. The compiler, runtime, and sanitizer report no error.

The timeline shows which version has the dependency.

By the end of this page you will have built a fork-join diamond across four streams, made its shape appear in an Nsight Systems capture, and seen how moving one line turns the diamond into a race.

An event is a snapshot, not a flag

Day 51 gave each stream its own ordered queue and showed two queues overlapping. What it did not give you is an edge between queues, and a fork-join needs five of them: four from the producer to the branches, and four more from the branches to the join (the producer edge is one event waited on four times).

A CUDA event is that edge, and the CUDA 12.6 runtime reference defines both halves precisely. Of cudaEventRecord: "Captures in event the contents of stream at the time of this call." Of cudaStreamWaitEvent: "Makes all future work submitted to stream wait for all work captured in event." (https://docs.nvidia.com/cuda/archive/12.6.3/cuda-runtime-api/group__CUDART__EVENT.html and https://docs.nvidia.com/cuda/archive/12.6.3/cuda-runtime-api/group__CUDART__STREAM.html , both checked 2026-09-01.) Read the first quote again: the contents of the stream at the time of this call. An event is a snapshot of work already queued, not a flag that someone sets later.

The wait costs the host nothing. cudaStreamWaitEvent returns as soon as the wait is queued, and the stream reference says: "The synchronization will be performed efficiently on the device when applicable." Compare that with the tools you had before this page: a cudaDeviceSynchronize or a cudaEventSynchronize parks the host thread until the device catches up, which adds a host wait to the timeline as day 41 measured. A cross-stream event edge keeps the host free to keep queueing, which is the whole reason launch overhead can hide behind running kernels.

Diagram: the diamond, and the diamond with one line moved. Two bands, each five stream rows on one time axis, kernels as bars. Band 1, correct order: one grindFma bar on stream 0, then four grindFma bars on streams 1 to 4 all opening at the first bar's right edge, then one sumFour bar on stream 0 opening at the latest branch's edge. Caption "1 event recorded, 8 waits, 4 branches side by side." Band 2, record moved above the producer launch: the four branch bars now open at time zero, under the producer bar, and only the join bar still waits. Caption "Same 9 launches, 0 real dependencies on the fork side." Alt text: "A fork-join diamond needs one recorded event and eight waits. Recording the event one line early makes the four branch kernels start at time zero, overlapping the producer they were meant to follow."

The intuition that fails: treating the event like a boolean

If you come from CPU threading, an event sounds like a condition variable: wait on it anywhere, and you block until someone signals. CUDA events do not work that way, and the runtime reference is blunt about the empty case: "Before the first call to cudaEventRecord(), an event represents an empty set of work, so for example cudaEventQuery() would return cudaSuccess." A wait on an empty or too-early snapshot is satisfied immediately. It is not an error, it is a no-op, and the fork half of your diamond turns into five launches with no ordering among them.

The same page settles reuse: cudaEventRecord "can be called multiple times on the same event and will overwrite the previously captured state", and waits "use the most recently captured state at the time of the API call".

One event per dependency edge, re-recorded every iteration, is fine as long as the host issues record before wait each time around. The order of your host lines is the dependency graph. Nothing else is.

One program, three submissions

Full program in code/day52-events/events_forkjoin.cu. It stands or falls on three choices.

Dependency events do not time; timing events do not order. The fork and branch events are created with cudaEventDisableTiming, which the event reference calls out as the fast path: events with this flag "will provide the best performance when used with cudaStreamWaitEvent() and cudaEventQuery()". The stopwatch events inside the timing helper keep the default flags, because cudaEventElapsedTime refuses timestamps that were never taken. The program proves that refusal rather than asserting it.

Every kernel is small enough to leave room for other work. Each grindFma launch has 8 blocks. Day 51's rule applies: two kernels overlap only when the device has room for blocks from both grids. A diamond whose branches queue for free SMs would serialise whatever the events say.

Correctness gates run before any clock. The diamond and a serial version of the same nine launches are both checked against a CPU reference that mirrors the kernels' fmaf chain step for step, and a mismatch is a real failing exit, not a warning.

Here is the diamond. Every call returns as soon as its work is queued; the two event calls carry all of the ordering:

// The diamond. Every call here returns as soon as the work is queued; the
// ordering lives in the two event calls, not in the call order.
static void enqueueDiamond(const Diamond& d) {
    nvtxRangePushA("upstream");
    grindFma<<<kGrindBlocks, kThreadsPerBlock, 0, d.s[0]>>>(
        d.d_in, d.d_mid, kA, 1.0f, kElems, kIters);
    CUDA_CHECK(cudaEventRecord(d.fork, d.s[0]));
    nvtxRangePop();

    nvtxRangePushA("fork");
    for (int j = 0; j < kBranches; ++j) {
        CUDA_CHECK(cudaStreamWaitEvent(d.s[1 + j], d.fork, 0));
        grindFma<<<kGrindBlocks, kThreadsPerBlock, 0, d.s[1 + j]>>>(
            d.d_mid, d.d_branch[j], kA, kBranchB[j], kElems, kIters);
        CUDA_CHECK(cudaEventRecord(d.branchDone[j], d.s[1 + j]));
    }
    nvtxRangePop();

    nvtxRangePushA("join");
    for (int j = 0; j < kBranches; ++j) {
        CUDA_CHECK(cudaStreamWaitEvent(d.s[0], d.branchDone[j], 0));
    }
    const int blocks =
        static_cast<int>((kElems + kThreadsPerBlock - 1) / kThreadsPerBlock);
    sumFour<<<blocks, kThreadsPerBlock, 0, d.s[0]>>>(
        d.d_branch[0], d.d_branch[1], d.d_branch[2], d.d_branch[3], d.d_out,
        kElems);
    nvtxRangePop();
}

And here is the trap, as one moved line rather than a paragraph of theory. The record now runs while stream 0 is empty, so the snapshot contains nothing, every wait is honoured against nothing, and the branches race the producer that feeds them:

// The trap: one line moved. The record now runs before the upstream launch,
// so the event captures s[0] while it is empty. Every cudaStreamWaitEvent
// below is honoured, immediately, against that empty snapshot, and the
// branches race the kernel that feeds them. No API returns an error.
static void enqueueBroken(const Diamond& d, cudaEvent_t forkTooEarly) {
    nvtxRangePushA("broken-order");
    CUDA_CHECK(cudaEventRecord(forkTooEarly, d.s[0]));  // captures nothing
    grindFma<<<kGrindBlocks, kThreadsPerBlock, 0, d.s[0]>>>(
        d.d_in, d.d_mid, kA, 1.0f, kElems, kIters);
    for (int j = 0; j < kBranches; ++j) {
        CUDA_CHECK(cudaStreamWaitEvent(d.s[1 + j], forkTooEarly, 0));
        grindFma<<<kGrindBlocks, kThreadsPerBlock, 0, d.s[1 + j]>>>(
            d.d_mid, d.d_branch[j], kA, kBranchB[j], kElems, kIters);
        CUDA_CHECK(cudaEventRecord(d.branchDone[j], d.s[1 + j]));
    }
    for (int j = 0; j < kBranches; ++j) {
        CUDA_CHECK(cudaStreamWaitEvent(d.s[0], d.branchDone[j], 0));
    }
    const int blocks =
        static_cast<int>((kElems + kThreadsPerBlock - 1) / kThreadsPerBlock);
    sumFour<<<blocks, kThreadsPerBlock, 0, d.s[0]>>>(
        d.d_branch[0], d.d_branch[1], d.d_branch[2], d.d_branch[3], d.d_out,
        kElems);
    nvtxRangePop();
}

The broken run reads a zeroed intermediate buffer, so a branch that starts too early computes from zeros and the program can count the damage. One caveat: that count is an observation, never a gate. A race that happens to produce the right answer is still a race, which is the lesson day 27 paid for in full.

Results

Re-verified 2026-09-02 on the same Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88). Diamond and serial correctness reproduced; diamond held at 4.826 ms while serial moved from 8.366 to 8.305 ms, leaving the ratio effectively unchanged at 1.72x.

The race count changed, as expected. Both transcripts and the fresh CUDA 13 report are listed in front matter. The timeline capture and its two CSV exports ship at content/profiles/cuda-events-dependencies/, so the exercise is doable without a GPU.

Measurement Value
diamond, ms per run 4.826
serial, ms per run 8.305
serial / diamond 1.72x
broken order, elements wrong of 1,048,576 984,128
cudaEventElapsedTime on the fork event cudaErrorInvalidResourceHandle

What the timeline should show

  1. The diamond beats the serial version by 2.0x to 3.0x. The five grind kernels are the same size, so the serial shape costs about six grind units (producer, four branches, join) and the diamond about two and a fraction. Below 1.5x the branches did not overlap and the 40-block residency argument above is wrong somewhere.
  2. The broken ordering produces a nonzero mismatch count on this run, and the capture shows the four branch bars opening before the producer bar closes. The count itself I refuse to predict: it is a race, and a different schedule gives a different number. If it comes back zero while the timeline still shows the overlap, the observation column is wrong, not the timeline.
  3. cudaEventElapsedTime on the disable-timing fork event returns cudaErrorInvalidResourceHandle, which the program gates on. The runtime reference commits to this in writing; anything else is a documentation bug worth reporting.

What the run said

  1. The band died; the floor and the shape held. Serial over diamond measured 1.72x, above the 1.5 falsification line but below the predicted 2.0 to 3.0. The prediction's arithmetic was the error: it priced the diamond at "two and a fraction" grind units, but the diamond's critical path is three whole units deep (producer, branch layer, join), so even a perfect card caps the ratio at 2.0, and 1.73 is that ceiling minus launch and event overhead. The four branch bars do sit abreast between the two single bars in the capture.
  2. Held. The broken ordering came back 984,128 elements wrong of 1,048,576 on this run, with the branch bars opening before the producer bar closes. The count is a race artifact; expect a different number on a different run.
  3. Held exactly. cudaEventElapsedTime on the disable-timing fork event returned cudaErrorInvalidResourceHandle, as the runtime reference commits to in writing.

The shape to keep is the diamond itself: four bars abreast between two single bars. On any card, whatever the milliseconds, source order created none of it. The two event calls did.

Run it yourself

Use a CUDA GPU with Nsight Systems installed.

nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o events_forkjoin \
    events_forkjoin.cu
nsys profile -t cuda,nvtx -o day52-forkjoin ./events_forkjoin

There is no Compiler Explorer embed: the CPU reference alone runs about 2.7 billion fused multiply-adds against the 20 second cap, and the artifact that matters is the timeline, which no embed can draw. Without any GPU, read /setup/learn-cuda-without-a-gpu and open the shipped report; the diamond, the serial block and the broken-order NVTX range are all findable in Nsight Systems by name.

Exercise

Rebuild the fork on two streams instead of four: two branch streams, each running two of the four branch kernels back to back, with the join still waiting on every branch buffer. Keep the program's harness passing, then report your serial-to-diamond ratio next to the four-stream one and explain the change using critical-path arithmetic, not adjectives.

Time: 30 to 45 minutes. Submit: your modified events_forkjoin.cu, both ratios, and two sentences: one deriving each ratio from grind units on the critical path.

Check: harness. The program checks both shapes against the CPU reference and prints the first mismatching index with got and want values on failure; a correct two-stream fork still passes with 0 mismatches, because the dependency edges changed and the dataflow did not.

On a card with at least 24 SMs it also requires fork-join to beat serial by 1.3x and exits nonzero otherwise; below 24 SMs it prints the ratio and says it is not gated. The broken-order count at the end is reported either way and passing it is meaningless: it is a race.

Hint 1

Nothing about correctness cares how many streams there are. Count what sits between the producer's last element and the join's first element, in units of one grind kernel, for four streams and then for two.

Hint 2

With four branch streams the branch phase costs one grind unit, because all four run abreast. Two kernels back to back on one stream cost two units whatever the other stream does. How many events do you need per stream now, and which of the four branchDone records can you drop?

Solution

Each branch stream waits once on the fork event, launches its two branch kernels in order (the stream itself orders them, no event needed between them), and records one done event after the second. The join waits on two events instead of four. Correctness passes untouched because every branch still reads the finished d_mid and the join still waits for every buffer it sums.

The arithmetic: four streams put producer, one branch unit and the join on the critical path, about 2 grind units plus the cheap join, against 6 for serial, near 2.7x. Two streams make the branch phase two units, about 3 total, near 2x. The ratio the harness prints should fall but stay above the 1.3x gate.

The dependency edges you declare set the critical path. The timeline shows the edges the program created, not the ones you meant to create. Fewer streams do not change correctness, but they make the path longer.

Pitfalls

Your waits do nothing and no call ever fails. The record ran after the wait, or never ran, so the wait saw an empty snapshot: "Before the first call to cudaEventRecord(), an event represents an empty set of work." Fix the host-side line order; then confirm on the timeline, because this bug is invisible everywhere else. This page's broken-order range is what it looks like.

You reused one event for two in-flight iterations. A wait latches "the most recently captured state at the time of the API call", so re-recording an event each loop iteration is safe only while record always precedes the waits that pair with it. Two iterations queued ahead of the device need two events, or a per-iteration record-then-wait order like this program's.

cudaEventElapsedTime returns cudaErrorInvalidResourceHandle on events that exist. They were created with cudaEventDisableTiming, or one of them was never recorded; the 12.6 reference lists both causes. Keep two kinds of events: default-flag pairs for the stopwatch, disable-timing for every dependency edge.

Your timing helper reports numbers that are too good. The stopwatch events here are recorded on the default stream, which orders against streams from cudaStreamCreate. Create your streams with cudaStreamNonBlocking and that ordering disappears: the stop event can complete while branch work is still running, and the "measurement" shrinks. Record the stop on a stream that provably follows the work, or sync first.

You ordered the device by parking the host. cudaEventSynchronize after the producer does create the dependency, and it also blocks the host thread that day 41 taught you to keep feeding the card. If the consumer of the ordering is device work, use cudaStreamWaitEvent and leave the host out of it.

Go deeper

Next

Day 53 puts copies on the timeline: pinned against pageable host memory, and cudaMemcpyAsync joining the streams this page ordered. The diamond you built here comes back on day 56, where CUDA graphs capture the whole dependency shape once and replay it without paying launch overhead per kernel, and day 48's launch-counting habit tells you whether that capture was worth it.