Skip to main content

ClikaRT::Stream

class

Header: ClikaRT/compute/stream.h

An execution stream: an ordered queue of device work. Tensors are placed on a stream; operations submitted to it run in order, asynchronously of the host.

Put real work on a stream you create (Stream::create(device)): submission there is pipelined, so the host keeps enqueueing while the device runs, and two such streams overlap independent work. A bare Device placement never runs on a device's default stream: the work follows its inputs and otherwise takes the calling thread's CURRENT lane there, the stream StreamOrDevice(device).resolve() names. The per-device Stream::default_stream(device) dispatches each operation on the calling thread instead: the explicit opt-in to synchronous execution (pass it on purpose, or use SynchronousStreamScope); convenient for a one-off call or a quick script, but it gives up that pipelining, so it is not where a throughput-sensitive workload belongs.

A handle to a device execution stream. Cheap to copy (it is an identifier, not the stream's resources).

A moved-from Stream is an EMPTY handle that names no stream: is_default() is false, query_idle() is true, native_handle() is nullptr, synchronize() completes at once, device() reports no device, it compares equal to another empty handle and to nothing else, a copy of it is empty, and allocate raises InvalidArgument. Assign a live stream to use it again.

Every fallible operation returns its value and raises ClikaRT::Error on failure; opt back into a Result<T> at the call site with CLIKART_TRY(...).

Static member functions​

default_stream()​

static Stream default_stream(Device device)

The per-device default stream; always valid, never fails.

Declared in ClikaRT/compute/stream.h, line 54

create()​

static Stream create(Device device)

Allocate a fresh non-default stream on device (for overlapping work).

Declared in ClikaRT/compute/stream.h, line 57

Member functions​

Stream(Stream)​

Stream(const Stream& other)

Declared in ClikaRT/compute/stream.h, line 59

operator=(Stream)​

Stream& operator=(const Stream& other)

Copy-assign: rebinds this handle to the same underlying stream.

Declared in ClikaRT/compute/stream.h, line 61

Stream(Stream)​

Stream(Stream&& other) noexcept

Move: transfers the handle; other is left empty and stays usable.

Declared in ClikaRT/compute/stream.h, line 63

operator=(Stream)​

Stream& operator=(Stream&& other) noexcept

Move-assign: transfers the handle; other is left empty and stays usable.

Declared in ClikaRT/compute/stream.h, line 65

~Stream()​

~Stream()

Releases this handle (the stream itself outlives its handles).

Declared in ClikaRT/compute/stream.h, line 67

device()​

Device device() const

The device this stream runs on.

Declared in ClikaRT/compute/stream.h, line 70

is_default()​

bool is_default() const

True if this is the device's default stream.

Declared in ClikaRT/compute/stream.h, line 72

query_idle()​

bool query_idle() const

True if the stream currently has no work in flight (a hint, not a barrier).

Declared in ClikaRT/compute/stream.h, line 74

operator==()​

bool operator==(const Stream& other) const

True when both handles name the same lane: a pool lane compares equal to every handle of that lane and unequal to every other lane; a device's default stream equals itself and no other stream.

Declared in ClikaRT/compute/stream.h, line 79

operator!=()​

bool operator!=(const Stream& other) const

Declared in ClikaRT/compute/stream.h, line 80

native_handle()​

void* native_handle() const noexcept

The backend-native stream handle, as an opaque pointer, what you launch a hand-written kernel on. static_cast it to the backend's stream type (cudaStream_t for a CUDA stream; the queue handle for others) inside a custom op's compute():

auto cu = static_cast<cudaStream_t>(stream.native_handle());
my_kernel<<<grid, block, smem, cu>>>(...);

Returns nullptr for the CPU backend (host work has no device stream) and for a stream with no live backend resources. The handle is owned by the runtime; do not destroy it, and keep this Stream alive while using it.

Declared in ClikaRT/compute/stream.h, line 93

synchronize()​

void synchronize()

Block until all work submitted on this stream has completed.

Declared in ClikaRT/compute/stream.h, line 97

allocate()​

Tensor allocate(ClikaRT::Span<const std::int64_t> shape, DataType dtype)

Stream::allocate: defined here (not in stream.h) because a by-value Tensor return needs the complete Tensor type, and stream.h cannot include tensor.h (this header includes stream.h). The throw stays caller-side (inline).

Declared in ClikaRT/compute/stream.h, line 115