**The core rule is local order, global explicitness.** A CUDA stream is an in-order work queue. A kernel, asynchronous copy, memory set, event record, or host function placed later in the same stream observes the required completion order of earlier operations in that stream. Two distinct streams are independent scheduling lanes: the runtime may overlap them, serialize them, or interleave their work according to dependencies and available hardware resources. Concurrency is therefore a permission, not a promise. Multiple streams expose independent work, but simultaneous execution depends on copy engines, kernel resource use, memory bandwidth, occupancy, device architecture, driver/runtime behavior, and contention from other processes. Correct code must produce the same answer whether independent operations overlap or happen serially. | Mechanism | What it orders or waits for | Blocks host thread? | Typical use | Common mistake | |---|---|---:|---|---| | Same-stream enqueue order | Earlier operations in that stream | No | Linear producer/consumer sequence | Assuming it orders another stream | | `cudaStreamWaitEvent` | Consumer work after the wait against the event frontier | No | Cross-stream dependency | Recording the event too early | | `cudaEventQuery` | Reports event completion state | No | Polling/progress engine | Busy-spinning continuously | | `cudaEventSynchronize` | Host waits for one event frontier | Yes | Host needs one result | Using it between every stage | | `cudaStreamQuery` | Reports whether prior stream work is complete | No | Stream readiness check | Treating “not ready” as failure | | `cudaStreamSynchronize` | Host waits for prior work in one stream | Yes | Reuse a stream-owned resource | Draining unrelated work by using device sync instead | | `cudaDeviceSynchronize` | Host waits for prior work across the device context | Yes | Global phase boundary or diagnosis | Placing it in the steady-state loop | | CUDA Graph dependency | Predefined node-to-node edge | No at launch | Repeated stable workflow | Capturing unsupported or unsafe operations | **An event is a completion frontier, not an interrupt.** `cudaEventRecord(event, producer)` enqueues a marker in the producer stream. The event represents work captured in that stream at record time. Operations submitted to the producer after the record do not become part of that recorded frontier. This makes placement exact: record immediately after the last producer operation needed by a consumer. An event object can be recorded more than once. Re-recording replaces its represented state, so careless reuse can make reasoning difficult. A wait already issued against an event uses the captured state applicable when that wait was submitted and is not retroactively changed by a later record. Even so, event pooling should associate reuse with a known generation or buffer slot so humans can audit ownership. ```cpp cudaStream_t upload, compute; cudaEvent_t input_ready; CUDA_CHECK(cudaStreamCreateWithFlags(&upload, cudaStreamNonBlocking)); CUDA_CHECK(cudaStreamCreateWithFlags(&compute, cudaStreamNonBlocking)); CUDA_CHECK(cudaEventCreateWithFlags(&input_ready, cudaEventDisableTiming)); CUDA_CHECK(cudaMemcpyAsync(d_input, h_input, bytes, cudaMemcpyHostToDevice, upload)); CUDA_CHECK(cudaEventRecord(input_ready, upload)); // Enqueues a device dependency; the CPU does not wait here. CUDA_CHECK(cudaStreamWaitEvent(compute, input_ready, 0)); transform<<>>(d_input, d_output, count); CUDA_CHECK(cudaGetLastError()); ``` The consumer stream can accept the wait and subsequent kernel immediately. On the device, the consumer kernel is not eligible to pass the wait until the recorded producer frontier completes. Other streams with no dependency remain eligible. This is fundamentally different from making the CPU wait and then submitting the consumer: the GPU scheduler receives the dependency early and the host remains free to prepare later batches. **Model the workload as a directed acyclic graph.** Nodes represent operations and edges represent true happens-before requirements. For an edge $A\rightarrow B$, ask what data or resource produced by $A$ is consumed, overwritten, freed, or externally observed by $B$. If no such requirement exists, an edge may be unnecessary serialization. A correct schedule is any topological ordering of the graph. Streams partition nodes into ordered lanes; events add edges between lanes. The design objective is not “maximum streams,” but a graph whose critical path is short and whose independent nodes are visible to the runtime. For a path with operation durations $t_i$, an idealized lower bound is $$T_{critical}=\max_{p\in\mathcal{P}}\sum_{i\in p}t_i$$ where $\mathcal{P}$ is the set of dependency paths. Real execution adds launch overhead, contention, imperfect engine overlap, and synchronization latency. Extra edges can only maintain or lengthen the critical path; they cannot create useful parallelism. **Choose event boundaries from buffer ownership.** A robust pipeline defines when each buffer slot is writable by upload, readable by compute, writable by compute, readable by download, and reusable by the host. Each transition gets a clear owner and, only when it crosses streams or host/device domains, an explicit completion mechanism. For triple-buffered processing, slot $k$ might have `input_ready[k]`, `compute_done[k]`, and `download_done[k]`. Before uploading the next generation into slot $k$, the upload stream waits for the prior generation’s completion event. That edge prevents overwrite while allowing other slots to progress. ```cpp struct Slot { void* h_in; void* h_out; void* d_in; void* d_out; cudaEvent_t reusable; cudaEvent_t input_ready; cudaEvent_t output_ready; }; for (std::size_t batch = 0; batch < batches; ++batch) { Slot& s = slots[batch % depth]; // GPU-side protection against reusing this slot too early. CUDA_CHECK(cudaStreamWaitEvent(upload, s.reusable, 0)); prepare_input(s.h_in, batch); CUDA_CHECK(cudaMemcpyAsync(s.d_in, s.h_in, bytes, cudaMemcpyHostToDevice, upload)); CUDA_CHECK(cudaEventRecord(s.input_ready, upload)); CUDA_CHECK(cudaStreamWaitEvent(compute, s.input_ready, 0)); model<<>>(s.d_in, s.d_out); CUDA_CHECK(cudaGetLastError()); CUDA_CHECK(cudaEventRecord(s.output_ready, compute)); CUDA_CHECK(cudaStreamWaitEvent(download, s.output_ready, 0)); CUDA_CHECK(cudaMemcpyAsync(s.h_out, s.d_out, bytes, cudaMemcpyDeviceToHost, download)); CUDA_CHECK(cudaEventRecord(s.reusable, download)); } ``` The host must not modify a pinned input buffer while an asynchronous host-to-device transfer reads it, and it must not consume a pinned output buffer until the device-to-host transfer finishes. If host preparation or consumption touches the same slot, use event query/synchronize at the slot boundary or a higher-level completion queue. A device-side event wait alone does not tell the CPU that host memory is safe. **Separate host waiting from device ordering.** `cudaStreamWaitEvent` orders GPU work and normally returns to the host after enqueueing. `cudaEventSynchronize` and `cudaStreamSynchronize` make the host wait. Choose based on who consumes the result: - GPU consumer in another stream: enqueue `cudaStreamWaitEvent` in that stream. - CPU consumer of one milestone: wait or query the corresponding event. - CPU needs every prior operation in one stream: synchronize that stream. - CPU requires a global quiescent point: synchronize the device, deliberately. - CPU has other useful work: query periodically, use a completion thread, or integrate a bounded progress mechanism rather than immediately blocking. Polling should include useful work, backoff, or integration with an event loop. A tight `cudaEventQuery` spin consumes a CPU core and can increase system contention. Blocking-event flags may reduce CPU spinning for long waits, but latency and scheduling tradeoffs should be measured for the application. **Default-stream semantics can insert invisible edges.** Code that omits a stream argument uses a default stream, but its synchronization behavior depends on build and stream configuration. Under legacy default-stream semantics, work in the legacy stream synchronizes with blocking streams in the same context. An accidental default-stream kernel or blocking copy can therefore split otherwise independent work into serialized phases. Non-blocking streams created with `cudaStreamNonBlocking` do not participate in legacy default-stream synchronization. Per-thread default-stream mode gives each host thread an independent implicit stream, enabled consistently through the compiler option or API macro. Mixed compilation units, libraries, or assumptions can make behavior surprising, so production code should document its default-stream policy and pass explicit streams through asynchronous APIs. Do not infer that “non-blocking stream” means host calls can never block. That flag describes interaction with the legacy default stream. Allocation, pageable-memory staging, module loading, context initialization, resource pressure, and some APIs can still introduce host-side latency or synchronization. **Asynchronous copies need the right memory and hardware path.** `cudaMemcpyAsync` provides stream ordering, but host/device transfer overlap generally requires pinned page-locked host memory. With pageable memory, the runtime may stage or behave synchronously, preserving correctness while defeating the intended pipeline. Pinned memory is a limited system resource. Pool and reuse it; do not pin arbitrary large regions indefinitely. NUMA placement also matters on multi-socket hosts. The buffer should be allocated and initialized near the CPU and I/O path that will use it, then benchmarked under realistic topology and load. For $N$ equal batches with separate upload, compute, and download engines, an optimistic pipeline model is $$T_N\approx T_{fill}+(N-1)\max(t_{H2D},t_K,t_{D2H})+T_{drain}$$ rather than $N(t_{H2D}+t_K+t_{D2H})$. This bound assumes enough buffers, independent engines, no bandwidth bottleneck, and kernels that leave resources for concurrent work. If compute saturates memory bandwidth, overlapping a copy may slow both operations and provide little net benefit. A useful measured overlap efficiency is $$\eta_{overlap}=\frac{T_{serial}-T_{async}}{T_{serial}-T_{ideal}}$$ where $T_{serial}$ is a correct serialized baseline, $T_{async}$ is measured asynchronous time, and $T_{ideal}$ is the pipeline lower bound for the same workload. Values near one indicate good realization of available overlap; values below zero indicate that orchestration overhead or contention made the asynchronous version worse. **Event timing and event dependency are related but different uses.** Default events can carry timestamps for elapsed-time measurement. Place start and stop events in a controlled stream around the operation set being measured, ensure the stop event has completed, then call `cudaEventElapsedTime`. Timing across unrelated streams can include scheduler gaps and intervening work, so define the interval precisely. Dependency-only events should often be created with `cudaEventDisableTiming`. This can reduce event-management overhead and communicates intent. Timing-disabled events cannot be used with elapsed-time calculation. Do not use host wall-clock timing around asynchronous launches without a completion boundary; it measures submission latency, not GPU execution. Warm up contexts, modules, allocators, libraries, and clocks before collecting steady-state timing. Report distribution, not one sample. GPU clocks, thermal state, power limits, contention, memory residency, and initialization can dominate small kernels. Use profiler timelines to confirm which engines actually overlap and where waits occur. **Stream priorities are scheduling hints, not dependency edges.** A higher-priority stream can influence selection of pending kernels, but priority does not prove order, preempt already running work in a general way, or guarantee transfer priority. Correctness must still come from same-stream ordering or explicit synchronization. Use priorities for latency policy only after the dependency graph is correct. Too many streams can increase launch overhead, event traffic, memory footprint, and scheduling noise without exposing more concurrency. Start with lanes that map to real roles—upload, compute, download, communication, or independent request classes—then add streams only when profiling shows queued independent work and available hardware capacity. **Resource lifetime is part of synchronization correctness.** A stream or event handle, device allocation, pinned host buffer, library handle, descriptor, workspace, graph executable, or IPC object must remain valid while queued work can reference it. Scope exit on the host does not imply GPU completion. Associate each asynchronous resource with an ownership protocol. Options include slot completion events, stream-ordered allocation/free, reference-counted request objects released by a completion thread, or framework futures that ultimately map to a GPU event. Avoid a destructor that calls `cudaDeviceSynchronize`; it hides a global barrier in an innocent-looking operation. Libraries often accept a stream on a handle or call. Ensure the library handle, workspace, and descriptors are not concurrently mutated from another host thread. If a library internally uses additional streams, follow its documented completion contract rather than assuming the passed stream captures all work. **Errors are asynchronous too.** A successful kernel launch call means the launch was accepted, not that execution succeeded. Check immediate launch configuration errors with `cudaGetLastError` or an equivalent wrapper, and check completion APIs for delayed execution faults. An error reported at a later event or stream synchronization may originate from an earlier operation. ```cpp #define CUDA_CHECK(expr) do { \ cudaError_t status_ = (expr); \ if (status_ != cudaSuccess) { \ throw std::runtime_error(cudaGetErrorString(status_)); \ } \ } while (0) work<<>>(args...); CUDA_CHECK(cudaGetLastError()); // launch-side validation CUDA_CHECK(cudaEventRecord(done, stream)); // Later, at the request's completion boundary: CUDA_CHECK(cudaEventSynchronize(done)); // execution-side validation ``` For diagnosis, temporarily forcing blocking launches can localize a failing operation, but it changes timing and removes concurrency. Treat that mode as a debugging aid, not a production fix. Tools such as compute sanitizers and timeline profilers provide stronger evidence for races, invalid accesses, and unexpected serialization. **Memory visibility follows documented synchronization and API rules.** An event dependency orders captured work before later consumer work in the waiting stream. It is not a replacement for correct intra-kernel synchronization, atomics, memory scopes, or external-memory protocols. Threads within one kernel need CUDA’s device programming primitives; interprocess, graphics, network, or external-device sharing may require external semaphores and the relevant memory-consistency contract. A data race can remain even when the GPU eventually becomes idle. If two streams write the same allocation without an ordering edge, later device synchronization only waits for both; it does not define which value wins. Establish ownership or order before conflicting accesses. **Build fork/join patterns explicitly.** A common workflow starts in one stream, fans out independent kernels, and joins before a reduction: ```cpp cudaEvent_t input_ready, left_done, right_done; preprocess<<>>(input); CUDA_CHECK(cudaEventRecord(input_ready, root)); CUDA_CHECK(cudaStreamWaitEvent(left, input_ready, 0)); left_branch<<>>(input, a); CUDA_CHECK(cudaEventRecord(left_done, left)); CUDA_CHECK(cudaStreamWaitEvent(right, input_ready, 0)); right_branch<<>>(input, b); CUDA_CHECK(cudaEventRecord(right_done, right)); CUDA_CHECK(cudaStreamWaitEvent(root, left_done, 0)); CUDA_CHECK(cudaStreamWaitEvent(root, right_done, 0)); combine<<>>(a, b, output); ``` The join stream waits for both branch frontiers. No host synchronization is required to express the graph. Recording one event after both waits and the combine provides a single request-completion token for downstream work. **Avoid these synchronization anti-patterns.** They often produce correct test output while destroying scalability or retaining races: 1. Calling `cudaDeviceSynchronize` after every kernel. This converts asynchronous submission into lockstep execution and masks missing edges. 2. Launching producer and consumer in different streams with no event because “they usually run in order.” Stream scheduling order is not a contract. 3. Recording the producer event before the final write needed by the consumer. 4. Waiting on the host and then launching the consumer when a stream wait would express the dependency. 5. Reusing one event across in-flight buffer generations without a clear record/wait ownership rule. 6. Modifying pinned input memory before its upload completes or reading output memory before download completion. 7. Accidentally using the legacy default stream between independent blocking streams. 8. Assuming stream priority creates ordering or preemption guarantees. 9. Measuring launch calls instead of completed GPU work. 10. Freeing buffers, workspaces, events, or handles while queued operations still reference them. 11. Capturing a graph while issuing prohibited synchronization or mixing captured and uncaptured dependencies. 12. Optimizing concurrency before validating results against a deterministic serialized reference. **CUDA Graphs are the next step for stable repeated DAGs.** Streams and event waits dynamically describe a graph through enqueue calls. When the same workflow repeats, graph definition or stream capture can create the dependency structure once, instantiate an executable graph, and launch it repeatedly with reduced CPU submission overhead. Graph nodes can represent kernels, memory operations, event record/wait operations, host functions, and other supported work. Edges constrain scheduling. Once all dependencies of a node are satisfied, CUDA may schedule it according to available resources. Graphs improve submission efficiency; they do not create hardware concurrency that resource usage prevents. Stream capture can incorporate cross-stream dependencies when events are recorded and waited within the same capture graph. Forked capture streams must be joined back to the origin stream before ending capture. Synchronizing or querying captured work is invalid because captured operations have not been submitted for execution. Libraries used during capture must explicitly support capture. Choose streams/events when topology or work changes frequently, graph construction would dominate, or interactive control is important. Choose graphs when launch overhead is significant and the operation topology repeats. Many systems use both: events manage request lifetimes and external stages, while each request’s stable compute core is a graph launch. **Tune from evidence, not stream count.** Establish a correct serialized baseline, then profile the asynchronous version. Inspect: - CPU submission gaps and launch rate; - copy-engine and compute-engine utilization; - event wait duration and where consumers become ready; - pageable versus pinned transfers; - kernels that monopolize registers, shared memory, or occupancy; - DRAM and interconnect bandwidth saturation; - implicit default-stream barriers; - allocation or library calls that synchronize; - buffer depth and backpressure; - tail latency as well as throughput; - correctness under randomized delays and different device loads. Pipeline depth should cover latency without creating unbounded queueing. By Little’s law, an approximate number of in-flight batches is $$L\approx\lambda W$$ where $\lambda$ is target throughput and $W$ is end-to-end latency. Round up and add only justified slack. Excess depth raises memory use and tail latency and can make cancellation or error recovery harder. Backpressure belongs at admission or slot acquisition. If every slot is in flight, wait for or poll the oldest relevant completion rather than globally draining the device. This keeps bounded memory and preserves progress in unrelated streams. **Use a verification workflow that separates correctness from performance.** First run a single stream and one buffer. Compare outputs against a CPU or known-good reference. Add explicit events and multiple buffers while retaining deterministic checks. Randomize stream submission timing, batch sizes, and host delays. Test under sanitizers. Only then remove redundant edges, increase depth, and profile overlap. ```flowchart Define every operation, buffer, handle, and external side effect → Draw producer-to-consumer, overwrite, reuse, and host-observation edges → Put naturally sequential operations in the same stream → Add an event immediately after each cross-stream producer frontier → Enqueue cudaStreamWaitEvent before each consumer or resource reuse → Choose host query/event wait/stream wait only where the CPU consumes a result → Set and document legacy, per-thread, or explicit non-blocking stream policy → Use pinned pooled buffers for transfers intended to overlap → Check launch errors immediately and execution errors at completion boundaries → Validate against a serialized reference and randomized schedules → Profile engine overlap, wait gaps, launch overhead, bandwidth, and occupancy → Remove only edges proven redundant by the ownership model → Bound in-flight slots and apply backpressure at admission → Capture a CUDA Graph when a stable DAG repeats and launch overhead matters → Revalidate on supported devices, toolkit versions, libraries, and production load ``` **A practical review checklist keeps dependencies auditable.** For every event, record its producer stream, exact record point, consumer streams, buffer generation, timing flag, and reuse condition. For every host wait, state which CPU access requires it. For every device-wide synchronization, document why a narrower event or stream boundary is insufficient. For every asynchronous copy, document host memory type, lifetime, ownership during transfer, device destination, and expected engine overlap. For every external library, record stream and capture support. For every graph, define update rules, invalidation behavior, and fallback path when topology changes. The strongest design uses a dependency-graph and minimal-synchronization lens. Correctness comes from explicit happens-before edges and disciplined resource ownership; performance comes from withholding edges that are not required. Streams expose ordered lanes, events connect only true dependencies, host waits occur only at host consumption boundaries, and graphs compress stable launch sequences. The result is asynchronous GPU execution that remains correct under different scheduling decisions while preserving the overlap the hardware can actually deliver.