Streams and events

What they are

A stream is an ordered sequence of operations on an accelerator. Operations issued on the same stream run in order; operations on different streams may run concurrently. An event is a marker recorded onto a stream that other streams can wait on for synchronization.

Every PyTorch accelerator backend (CUDA, ROCm, MPS, XPU) has its own stream/event types but the API surface is similar.

CUDA streams

c10::cuda::CUDAStream (c10/cuda/CUDAStream.h) wraps a cudaStream_t plus the device and a "stream type" (default, current, custom-priority).

PyTorch maintains a per-thread current stream per device:

with torch.cuda.stream(my_stream):
    # all CUDA ops here submit to my_stream
    ...
torch.cuda.current_stream(device=0)  # query

Internally c10::cuda::getCurrentCUDAStream() returns the current stream; ATen kernels submit to it. The c10::cuda::CUDAStreamGuard RAII helper temporarily switches the current stream.

There are special stream values:

• Default stream (legacy) — synchronizes implicitly with all other streams. Compatibility-only; PyTorch uses per-thread default by default.
• Per-thread default stream — separate per-thread streams that don't synchronize with each other; the default since CUDA 7.x.
• Allocator stream — the caching allocator records a "block allocation" against the stream that allocated it, so it knows which stream's work must complete before the block can be reused.

CUDA events

c10::cuda::CUDAEvent (c10/cuda/CUDAEvent.h) wraps a cudaEvent_t.

e = torch.cuda.Event(enable_timing=True)
e.record()                     # mark the current stream
my_other_stream.wait_event(e)  # block other_stream until e completes
elapsed_ms = e.elapsed_time(other_event)

Events are how PyTorch synchronizes between:

• Streams within a process.
• The CPU and the GPU (torch.cuda.synchronize() waits on a synthetic event).
• The DataLoader pin-memory thread and the training loop.
• The autograd engine's per-device backward streams.

Stream-aware allocator

The CUDA caching allocator is stream-aware. When a block is freed:

1. The allocator records an event on the stream that last used the block.
2. The block stays "in flight" until that event completes.
3. New allocations on the same stream can use the block immediately (they execute after the freeing op anyway).
4. New allocations on other streams either wait for the event or pick a different block.

This is why mixing streams without explicit synchronization rarely produces "use after free" GPU bugs in PyTorch — the allocator is doing the bookkeeping for you.

Stream APIs in user code

# Create a side stream for a one-off async op
side = torch.cuda.Stream()
with torch.cuda.stream(side):
    big_async_kernel(x)
torch.cuda.current_stream().wait_stream(side)  # main waits

# Async H2D into a side stream
torch.cuda.synchronize()
buf = torch.empty(N, device="cuda")
with torch.cuda.stream(side):
    buf.copy_(cpu_buf, non_blocking=True)

DataLoader pin-memory + stream

When pin_memory=True, batches are H2D-copied async on a side stream in the main process before being yielded to the training loop. The main loop's first op on the batch synchronizes with that copy event. This is one of the more important latency hides for input-bound jobs.

Other backends

Backend	Stream type	Event type
CUDA	c10::cuda::CUDAStream	c10::cuda::CUDAEvent
ROCm/HIP	c10::hip::HIPStream	c10::hip::HIPEvent
MPS	at::mps::MPSStream (single per device)	at::mps::MPSEvent (MTLSharedEvent)
XPU	c10::xpu::XPUStream	c10::xpu::XPUEvent
MTIA	at::mtia::MTIAStream	at::mtia::MTIAEvent

The torch::ScalarEvent and torch::Stream abstractions in c10/core/Stream.h and c10/core/Event.h give a backend-agnostic interface used by the autograd engine and a few other cross-cutting consumers.

Where to look

File	Purpose
c10/cuda/CUDAStream.h / .cpp	CUDA stream wrapper
c10/cuda/CUDAEvent.h	CUDA event wrapper
c10/cuda/CUDAGuard.h	RAII device + stream guard
c10/cuda/CUDACachingAllocator.cpp	Stream-aware caching allocator
c10/core/Stream.h, Event.h	Backend-agnostic interfaces
aten/src/ATen/mps/MPSStream.mm	MPS stream

Where to read next

• Systems / CUDA backend — usage in CUDA kernels.
• Systems / Autograd — the engine uses streams to overlap backward across devices.