The CUDA Handbook ▸ Chapter 6. Streams, Events, and Graphs

6.10 Stream Memory Operations

CUDA events (Section 6.3) let one stream wait until another reaches a point in its work; stream memory operations let a stream wait on, or write, a specific value in memory. cuStreamWaitValue32() and cuStreamWaitValue64() enqueue a wait that blocks the stream until a 32- or 64-bit location satisfies a condition – equal to a value, greater than or equal to it (a cyclic comparison, so a wrapping counter still works), or a bitwise AND or NOR test. cuStreamWriteValue32() and cuStreamWriteValue64() enqueue a write of a value to memory, ordered within the stream like any other operation and, by default, preceded by a system-wide memory fence so that a producer’s earlier writes are visible before it sets the flag. cuStreamBatchMemOp() submits several such operations at once, which the driver dispatches faster than if the same operations were performed one at a time.

The mechanism is the same one the driver uses to track its own progress (Section 2.5): a monotonically increasing value written to a shared location, and a wait that proceeds once the location reaches a threshold. Exposed to the application, it enables fine-grained producer/consumer signaling that never involves the CPU – one stream, or one GPU, waits on a flag that another stream, another GPU, or a network adapter will write. It is the foundation of GPUDirect Async, in which the GPU rings a network adapter’s doorbell and waits on its completion flag directly, keeping the CPU off the critical path of communication.

The memory these operations act on is restricted:

• It must be addressable as a device pointer – device memory, or host memory that was registered with cuMemHostRegister() and whose device pointer was taken from cuMemHostGetDevicePointer(). Ordinary pageable host memory will not do.
• Managed memory (cudaMallocManaged()) is disallowed outright.
• A stream belongs to the device whose context created it, so the location must be a device pointer that stream’s device can reach. There is no rule that it live on that same GPU, but reaching another GPU’s memory follows the usual peer-access rules.

The 64-bit variants require compute capability 7.0 or later, and on Windows the device must be in TCC mode. The ordering these operations establish is invisible to CUDA’s own dependency tracking, which does not know that one stream’s kernel feeds a cuStreamWaitValue32() in another – so the application owns the correctness of the signaling, and misuse can deadlock. Support is reported by the CU_DEVICE_ATTRIBUTE_CAN_USE_STREAM_MEM_OPS device attribute.