---
title: "ClikaRT::Stream"
sidebar_label: "Stream"
description: "A handle to a device execution stream. Cheap to copy (it is an identifier, not the stream's resources)."
---

<!-- Generated by tools/doxygen2mdx from Doxygen XML. Do not edit. -->

<span class="api-kind api-kind--class">class</span>

**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](./Stream.md) is an EMPTY handle that names no stream: [`is_default()`](./Stream.md#is_default) is false, [`query_idle()`](./Stream.md#query_idle) is true, [`native_handle()`](./Stream.md#native_handle) is `nullptr`, [`synchronize()`](./Stream.md#synchronize) completes at once, [`device()`](device-namespace/index.md) reports no device, it compares equal to another empty handle and to nothing else, a copy of it is empty, and [`allocate`](./Stream.md#allocate) raises `InvalidArgument`. Assign a live stream to use it again.

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

## Static member functions

### default\_stream() {#default_stream}

<div class="api-signature">

`static` [`Stream`](./Stream.md) **`default_stream`**([`Device`](./Device.md) `device`)

</div>

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

<div class="api-declared">

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

</div>

### create() {#create}

<div class="api-signature">

`static` [`Stream`](./Stream.md) **`create`**([`Device`](./Device.md) `device`)

</div>

Allocate a fresh non-default stream on [`device`](device-namespace/index.md) (for overlapping work).

<div class="api-declared">

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

</div>

### from\_native\_handle() {#from_native_handle}

<div class="api-signature">

`static` [`Stream`](./Stream.md) **`from_native_handle`**([`Device`](./Device.md) `device`, `void*` `native_handle`)

</div>

Adopt a native stream you own (a `cudaStream_t` on CUDA, passed as an opaque pointer) as a stream on [`device`](device-namespace/index.md). Work placed on the returned stream runs on the native stream and is submitted asynchronously of the calling thread, as on a stream from [`create`](./Stream.md#create). The runtime never destroys the native stream. Adopting the same native stream on the same device again returns the same stream.

Scope: an adopted stream is not released before the process ends. The runtime may synchronize it, query it and reclaim memory on it at any point until then, so keep the native stream alive until the process ends.

A null handle raises `InvalidArgument`. On CUDA a null handle is the legacy default stream, which the runtime cannot order against its own streams; torch reports it (`torch.cuda.current_stream().cuda_stream` is 0) until you set a stream. Pass a stream you created (`cudaStreamCreate`, `torch.cuda.Stream()`). A device whose backend cannot adopt a native stream raises `Unsupported`.

<div class="api-declared">

Declared in `ClikaRT/compute/stream.h`, line 78

</div>

## Member functions

### Stream(Stream) {#Stream}

<div class="api-signature">

**`Stream`**(`const Stream&` `other`)

</div>

<div class="api-declared">

Declared in `ClikaRT/compute/stream.h`, line 82

</div>

### operator=(Stream) {#operator-assign}

<div class="api-signature">

[`Stream`](./Stream.md)`&` **`operator=`**(`const` [`Stream`](./Stream.md)`&` `other`)

</div>

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

<div class="api-declared">

Declared in `ClikaRT/compute/stream.h`, line 84

</div>

### Stream(Stream) {#Stream-2}

<div class="api-signature">

**`Stream`**(`Stream&&` `other`) `noexcept`

</div>

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

<div class="api-declared">

Declared in `ClikaRT/compute/stream.h`, line 86

</div>

### operator=(Stream) {#operator-assign-2}

<div class="api-signature">

[`Stream`](./Stream.md)`&` **`operator=`**([`Stream`](./Stream.md)`&&` `other`) `noexcept`

</div>

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

<div class="api-declared">

Declared in `ClikaRT/compute/stream.h`, line 88

</div>

### \~Stream() {#Stream-3}

<div class="api-signature">

**`~Stream`**()

</div>

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

<div class="api-declared">

Declared in `ClikaRT/compute/stream.h`, line 90

</div>

### device() {#device}

<div class="api-signature">

[`Device`](./Device.md) **`device`**() `const`

</div>

The device this stream runs on.

<div class="api-declared">

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

</div>

### is\_default() {#is_default}

<div class="api-signature">

`bool` **`is_default`**() `const`

</div>

True if this is the device's default stream.

<div class="api-declared">

Declared in `ClikaRT/compute/stream.h`, line 95

</div>

### query\_idle() {#query_idle}

<div class="api-signature">

`bool` **`query_idle`**() `const`

</div>

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

<div class="api-declared">

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

</div>

### operator==() {#operator-eq}

<div class="api-signature">

`bool` **`operator==`**(`const` [`Stream`](./Stream.md)`&` `other`) `const`

</div>

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.

<div class="api-declared">

Declared in `ClikaRT/compute/stream.h`, line 102

</div>

### operator\!=() {#operator-ne}

<div class="api-signature">

`bool` **`operator!=`**(`const` [`Stream`](./Stream.md)`&` `other`) `const`

</div>

<div class="api-declared">

Declared in `ClikaRT/compute/stream.h`, line 103

</div>

### native\_handle() {#native_handle}

<div class="api-signature">

`void*` **`native_handle`**() `const noexcept`

</div>

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()`:

```cpp
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`](./Stream.md) alive while using it.

<div class="api-declared">

Declared in `ClikaRT/compute/stream.h`, line 116

</div>

### synchronize() {#synchronize}

<div class="api-signature">

`void` **`synchronize`**()

</div>

Block until all work submitted on this stream has completed.

<div class="api-declared">

Declared in `ClikaRT/compute/stream.h`, line 120

</div>

### wait\_for() {#wait_for}

<div class="api-signature">

`void` **`wait_for`**(`const` [`Tensor`](./Tensor.md)`&` `tensor`)

</div>

Make this stream wait, on the device, for `tensor`'s pending writes: work submitted to this stream after the call runs once the producers of `tensor`'s bytes have finished, with no host wait (a CUDA stream waits on the producer's recorded event; a device's default stream and a CPU stream complete the wait on the calling thread, which is what a synchronous lane is). This stream is also recorded as a reader of those bytes, so a free of `tensor`'s buffer while the work queued here still reads it parks the buffer until this stream has run past that work. The DLPack handoff: a consumer's `stream` argument, adopted with [`from_native_handle`](./Stream.md#from_native_handle), waits here before the capsule is handed over. `tensor` must be defined and live on this stream's device, and this handle must name a stream (`InvalidArgument` otherwise; an empty handle's code reads `FAILED_PRECONDITION`). A tensor with no pending work returns at once.

<div class="api-declared">

Declared in `ClikaRT/compute/stream.h`, line 137

</div>

### allocate() {#allocate}

<div class="api-signature">

[`Tensor`](./Tensor.md) **`allocate`**([`ClikaRT::Span`](./Span.md)`<const std::int64_t>` `shape`, [`DataType`](../index.md#DataType) `dtype`)

</div>

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

<div class="api-declared">

Declared in `ClikaRT/compute/stream.h`, line 155

</div>