DLPack Tensor Exchange#

ovstage exchanges attribute values as DLPack tensors. Each tensor describes a buffer’s element type, shape, memory device, and strides (the distances between elements).

The snippets come from separate examples and tests; do not run them as consecutive steps in one program. stage / stage_ and query / query_ are already created. Python snippets use NumPy as np. The C tabs call the C API from C++. Test snippets use assertions to check calls, waitOp() to wait and check asynchronous errors, and seal(N) to advance the watermark and wait.

Passing Attribute Values#

Pass tensors to ovstage.Stage.write_attribute() and read them through ovstage.ReadGroup or ovstage.MapGroup. They associate tensor values with prims. On writes, set is_array explicitly. Shape and tensor count do not determine whether an attribute is array-valued.

Value per prim	`is_array`	Write tensors
Scalar, vector, or matrix	`False`	One tensor containing the values for all selected prims.
Variable-length array	`True`	One tensor per data row, or one packed tensor. See Writing Attributes for array layout rules.

The ovstage.DLTensor reference describes data, device, dtype, shape, strides, and byte_offset. Within dtype, code identifies the scalar type, bits its width, and lanes the component count.

For a complete write-and-read example in Python and C, see Application Flow.

The snippets below select /World/A, /World/B, and /World/C, then write temperatures 1.0, 2.0, and 3.0 in that order. is_array=False means each prim receives one fixed-size value. The stage and its path dictionary already exist.

Python

Select the three prims and resolve the temperature attribute:

attr = paths.intern_token("temperature")
prim_paths = paths.create_path_list_from_strings(["/World/A", "/World/B", "/World/C"])
query = stage.query_from_path_list(prim_paths)

Write one value for each prim in query order:

stage.write_attribute(
    query,
    attr,
    ordinal=1,
    tensors=np.array([1.0, 2.0, 3.0], np.float32),
    is_array=False,
).wait()

C

Select the three prims and resolve the temperature attribute:

ovx_string_t attrName = str("temperature");
ASSERT_EQ(path_dictionary_create_tokens_from_strings(dict_, &attrName, 1, &attr_).status,
          OVX_API_SUCCESS);

const ovx_string_t paths[] = { str("/World/A"), str("/World/B"), str("/World/C") };
ASSERT_EQ(path_dictionary_create_path_list_from_strings(dict_, paths, 3, &pathList_).status,
          OVX_API_SUCCESS);
ASSERT_EQ(ovstage_query_from_path_list(stage_, pathList_, &query_), OVSTAGE_OK);

Write one value for each prim in query order:

float values[] = { 1.0f, 2.0f, 3.0f };
int64_t shape[] = { 3 };
int64_t strides[] = { 1 };
DLTensor tensor{};
tensor.data = values;
tensor.device = { kDLCPU, 0 };
tensor.ndim = 1;
tensor.dtype = { kDLFloat, 32, 1 };
tensor.shape = shape;
tensor.strides = strides;

ovstage_write_data_t write{};
write.tensors = &tensor;
write.tensor_count = 1;
write.is_array = false;

ovx_string_or_token_t attrArg{ attr_, {} };
ASSERT_TRUE(waitOp(stage_,
    ovstage_write_attribute(stage_, query_, attrArg, /*ordinal*/ 1, write, OVSTAGE_PRIM_MODE_UPSERT),
    "write_attribute"));

Scalar, Vector, and Matrix Layouts#

Reads and maps return compact scalar-lane tensors (dtype.lanes == 1). For fixed-size attributes, the leading axis counts transported data rows. Each per-row array tensor instead counts array elements on its leading axis. In both cases, trailing axes follow the stored component width and semantic. The original input shape is not stored. Returned tensors have explicit compact row-major strides in scalar elements, for example [16, 4, 1] for [N, 4, 4]. Python array() and dlpack() preserve these typed shapes. Write inputs can still use null strides for compact data.

Logical prims select those data rows one-to-one when index_map is NULL. If index_map is present, logical prim i selects row index_map[i] and N may differ from count: it can be smaller when rows are shared, or larger when a query touches only part of a transported bucket.

Copy-in writes may use either the canonical form or compact convenience dimensions. Convenience dimensions are folded into dtype.lanes and are not preserved as schema.

Without index_map, the leading dimension must equal the logical element count. A flat shape = [N * L], lanes = 1 tensor is not a convenience encoding of N rows of width L; use shape = [N, L], lanes = 1 or the canonical shape = [N], lanes = L form.

Value per data row	Canonical write	Accepted convenience write	Raw read/map / Python DLPack export
Scalar	`shape = [N]`, `lanes = 1`	Same	`[N]`, `lanes = 1` / `(N,)`
`point3f`	`shape = [N]`, `lanes = 3`	`shape = [N, 3]`, `lanes = 1`	`[N, 3]`, `lanes = 1` / `(N, 3)`
`matrix4d`	`shape = [N]`, `lanes = 16`	`shape = [N, 4, 4]`, `lanes = 1` (or `[N, 4]` with 4 lanes)	`[N, 4, 4]`, `lanes = 1` / `(N, 4, 4)`

Important

Map creation dtypes and canonical write inputs encode the complete component width in dtype.lanes. Square MATRIX values return two component axes for sides 2, 3 or 4; generic tuples retain one axis, including 12-component tuples. Use array().reshape(-1) when flat scalar indexing is needed.

Forwarding typed tensors to writes#

Writes keep their canonical input contract. To forward a descriptor known to come from an OVStage read or map, use ovstage_tensor_to_write() from ovstage/ovstage_api/ovstage_tensor_utils.h or Python typed_tensor_to_write(). These adapters fold compact component dimensions into lanes, preserve outer row strides, device, pointer and byte offset, and copy metadata only. They do not infer types for arbitrary scalar arrays. Preserve the attribute semantic, array kind and prim/data index mapping.

A returned descriptor still borrows native payload. For a rewrite to pinned read storage, copy the payload before releasing the read group, then adapt the owned descriptor. Keep all input payload and metadata alive until the write completes, or through the availability-event synchronization when supplied (see Completion, Availability, and Buffer Lifetime). A map’s borrowed payload stops being valid at accepted unmap submission. Synchronously rejected unmap leaves the map usable.

Warning

Array attributes use different rules. An (N, 3) tensor with one lane supplies 3N scalar array elements, not N three-component elements. To write vector array elements, set dtype.lanes = 3 explicitly. Python’s ovstage.make_dltensor() can describe a compatible compact producer buffer this way without copying it. See Writing Attributes.

For non-NumPy DLPack producers, ovstage.make_dltensor() combines only complete trailing dimensions of a compact, byte-aligned buffer. The base type and bit width must remain unchanged; supported lane counts are 1–255. If all source dimensions are combined, the result is shape=(1,), ndim=1. Explicit shape and stride overrides must match that result: shape=[1], ndim=1, strides=[1] in this case. The function’s API reference describes these restrictions and the separate NumPy path.

These snippets write both scalar and three-component values for three prims. paths is the stage’s path dictionary. C uses the temperature and velocity tokens. Both representations put the complete component count in dtype.lanes.

Python

# A fixed-size column stacks one row per prim in ONE tensor: a scalar
# is a 1-lane dtype; a float3 is the SAME write with a 3-lane dtype.
# The shape stays [prim_count]; the tuple width lives in dtype.lanes.
temperature = paths.intern_token("temperature")
velocity = paths.intern_token("velocity")
stage.write_attribute(
    query, temperature, ordinal=1,
    tensors=np.array([18.5, 19.5, 20.5], np.float32), is_array=False,
).wait()
stage.write_attribute(
    query, velocity, ordinal=1,
    tensors=make_dltensor(
        np.array([1, 0, 0, 0, 1, 0, 0, 0, 1], np.float32),
        dtype=DLDataType(code=DLDataTypeCode.kDLFloat, bits=32, lanes=3),
        shape=[3]
    ),
    is_array=False,
).wait()
stage.advance_write_floor(ordinal=1).wait()

C

// A fixed-size column stacks one row per prim in ONE tensor: a scalar is a
// 1-lane dtype; a float3 is the SAME write with a 3-lane dtype. The shape
// stays {prim_count}; the tuple width lives in dtype.lanes.
float temperatures[] = { 18.5f, 19.5f, 20.5f };
float velocities[] = { 1, 0, 0, 0, 1, 0, 0, 0, 1 }; // 3 prims x 3 lanes
int64_t rows[] = { 3 };
int64_t strides[] = { 1 };

DLTensor tempTensor{};
tempTensor.data = temperatures;
tempTensor.device = { kDLCPU, 0 };
tempTensor.ndim = 1;
tempTensor.dtype = { kDLFloat, 32, 1 };
tempTensor.shape = rows;
tempTensor.strides = strides;

DLTensor velTensor{};
velTensor.data = velocities;
velTensor.device = { kDLCPU, 0 };
velTensor.ndim = 1;
velTensor.dtype = { kDLFloat, 32, 3 };
velTensor.shape = rows;
velTensor.strides = strides;

ovstage_write_data_t scalarWrite{};
scalarWrite.tensors = &tempTensor;
scalarWrite.tensor_count = 1;
scalarWrite.is_array = false;
ASSERT_TRUE(waitOp(stage_,
    ovstage_write_attribute(stage_, query_, { temperature, {} }, 1, scalarWrite, OVSTAGE_PRIM_MODE_UPSERT),
    "write scalar"));

ovstage_write_data_t vecWrite{};
vecWrite.tensors = &velTensor;
vecWrite.tensor_count = 1;
vecWrite.is_array = false;
ASSERT_TRUE(waitOp(stage_,
    ovstage_write_attribute(stage_, query_, { velocity, {} }, 1, vecWrite, OVSTAGE_PRIM_MODE_UPSERT),
    "write vec3"));
seal(1);

Buffer Ownership and Lifetime#

For C writes, provide exactly one of tensors or managed_tensors. Caller-owned tensor data must remain valid until the operation completes. With a placement availability event, retain it through device availability instead; see Completion, Availability, and Buffer Lifetime. Managed tensors transfer ownership to storage, which calls their deleter when no longer needed. See ovstage_write_data_t for the ownership contract.
Read results borrow storage. Keep each returned group alive while using its tensors or any array views made from them. Releasing the read handle does not release those groups.
Copy values that must outlive their result group. A DLPack consumer array does not keep the owning group alive for you. Outstanding groups can block overlapping writes.

The watermark controls which changes may be read; it does not extend the lifetime of returned buffers. See Reading Attributes for cleanup.

Copy values before releasing their group if you need them afterwards. Here group is an already-fetched result containing three float32 scalars; the C code has checked its layout and allocated float out[3]. Both copies survive release of the group. Release the read handle separately when finished fetching.

Python

values = np.array(group.array(0))
stage.release_group(group)

C

std::memcpy(out, group.data.tensors[0].data, sizeof(out));

ovstage_release_group(stage_, &group);

Experimental Device Placement#

Pass a ovstage.Placement as placement= to Stage.read_attributes(), Stage.write_attribute(), Stage.write_attributes(), or Stage.map_attribute(). In C, include <ovstage/ovstage_api/experimental/ovstage_placement_extension.h> and use the corresponding ovstage_*_on_device wrapper. The extension is experimental: its API, ABI, and semantics may change. An unavailable or incompatible extension reports NOT_FOUND; ordinary calls remain available.

device_type selects CPU (kDLCPU) or CUDA (kDLCUDA), and device_id selects the CUDA device. One placement applies to the whole operation. Omitting placement preserves existing behavior: writes follow the source tensor’s device, and reads/maps return CPU buffers. A zero device type within a placement leaves device selection unspecified with those same defaults.

Placement selects a device’s copy of an attribute; it does not pin that attribute to one device permanently. Writes and map commits produce values on the selected device. Reads use the copy already current there, or materialize one if needed. A read of an already-current copy avoids migration. Inspect each returned DLTensor.device and use a consumer for that device. group.array() is CPU-only; CUDA buffers require a CUDA-capable DLPack consumer. Keep each group alive through its consumer’s device work before releasing it or committing a map.

For a complete producer-to-consumer sequence in C and Python, see Device Placement (Experimental).

Completion, Availability, and Buffer Lifetime#

Placement’s stream carries destination work; a write’s cuda_stream and cuda_event describe the independent source handoff. availability_event is a caller-owned CUDA event that ovstage records for destination availability. Without it, device work is drained before the operation reports completion. With it, an operation can complete while its device work is still running:

Write or map commit: enqueue, wait for the operation, and check its status; then wait on the availability event or submit downstream work on the placement stream. Keep write source buffers alive and unchanged until device work finishes.
Read: fetch a group successfully; only then wait on the availability event or submit its consumer on the placement stream before dereferencing the payload. The event covers prior writes as well as any transfer needed by the read.
Map: the event covers each commit, not the staging buffers returned by setup. Finish producing mapped values and provide the source handoff needed by a CUDA commit; stop accessing buffers when unmap submission is accepted.

Waiting on the event or stream before these boundaries can observe an old or unrecorded event. Stream 2 denotes the worker thread’s per-thread default stream, which may differ from the caller’s; use the event to order the caller’s work in that case. Stream/event handles remain caller-owned, must be valid for the recording device, and must survive their use. The placement descriptor itself is copied during enqueue.

Python Operation.wait() releases its references to write sources. When using an availability event, keep your own source references until that event has completed. Also wait on outstanding write events before an attached renderer performs its own hierarchy propagation.

An event can be recorded for a failing operation: its completion does not prove success. If CUDA refuses to record it, the event remains untouched and a write or map commit may report OP_FAILED after committing its effects. Do not wait on that event as proof of availability. After observing the operation failure, synchronize the specific placement stream, or the recording device when using a default/per-thread stream, before releasing write sources. See Placement Availability Failures and ovstage.Placement for the full rules.

Current Placement Limits#

Placement does not guarantee asynchronous transfers in every case. A write staged between devices drains its staging buffer before returning completion. Ragged writes use the default stream rather than the requested placement stream; their availability event is already signalled when the operation completes. CUDA maps support fixed-size attributes only; array-valued CUDA maps return NOT_SUPPORTED. Map commit scatter uses the mapped device’s default stream and is drained before operation completion. An availability event does not make that scatter asynchronous with respect to completion, although submission still enqueues the operation. Staged host rows with negative or zero/broadcast strides are not supported. Hierarchy outputs follow the selected hierarchy computation model; placement does not change that model.

CPU/CUDA Data and Synchronization#

Writes accept CPU or CUDA source tensors. Reads and maps return CPU buffers unless a device is selected through Experimental Device Placement. A DLPack view avoids a copy of the returned buffer; transfers between devices can still be necessary.

For CUDA input, ovstage.Stage.write_attribute() accepts cuda_stream and cuda_event to synchronize before consuming the data. You can use both:

Parameter	Meaning
`cuda_stream`	`None` or `0`: no stream synchronization. `1`: CUDA default stream. Otherwise pass the integer CUDA stream handle. Synchronization waits for queued work on that stream to finish before consuming the input.
`cuda_event`	`None` or `0`: no event. Otherwise pass the integer CUDA event handle to wait for.

Leave both unset for CPU data or CUDA data whose producer has already completed. For writes with several tensors, do not mix CPU and CUDA tensors or different CUDA devices. In C, these parameters correspond to cuda_sync.stream and cuda_sync.wait_event in ovstage_cuda_sync_t; use {0, 0} for no additional synchronization.

These examples finish producing three CUDA values before submitting the write, so no additional synchronization is needed. The GPU buffer remains alive until the write completes. Python uses Warp as wp; C uses the CUDA runtime and a checkCuda() helper that reports CUDA errors. The query selects three prims, and gpu_samples / attribute names the destination attribute. See the CUDA tensor-write example for setup.

Python

# Finish producing GPU data, then retain it until the write completes.
device_values = wp.array(np.array([5.0, 6.0, 7.0], np.float32), dtype=wp.float32, device="cuda:0")
wp.synchronize()
write = stage.write_attribute(
    query, gpu_samples, ordinal=8, tensors=device_values, is_array=False
)
# Keep device_values alive and unchanged while the write is in progress.
write.wait()
del device_values  # Safe to release our reference after completion.

C

float values[] = {5.0f, 6.0f, 7.0f};
float* deviceValues = nullptr;
checkCuda(cudaMalloc(reinterpret_cast<void**>(&deviceValues), sizeof(values)));
checkCuda(cudaMemcpyAsync(deviceValues, values, sizeof(values), cudaMemcpyHostToDevice));
// Finish producing the GPU data before submitting it to ovstage.
checkCuda(cudaDeviceSynchronize());

int64_t shape[] = {3};
int64_t strides[] = {1};
DLTensor tensor{};
tensor.data = deviceValues;
tensor.device = {kDLCUDA, 0};
tensor.ndim = 1;
tensor.dtype = {kDLFloat, 32, 1};
tensor.shape = shape;
tensor.strides = strides;
ovstage_write_data_t write{};
write.tensors = &tensor;
write.tensor_count = 1;
write.is_array = false;
write.cuda_sync = {}; // Producer already synchronized: no additional wait.
waitOp(stage, ovstage_write_attribute(stage, query, {attribute, {}}, 8,
                                     write, OVSTAGE_PRIM_MODE_UPSERT), "CUDA write");
// Keep the GPU allocation alive until the write completes.
checkCuda(cudaFree(deviceValues));

Selecting Values and Prims#

By default, a write assigns the first input value to the first prim in the query, the second value to the second prim, and so on.

Use index_map to change which input value each prim receives. For example, [1, 0] gives the second input value to the first prim and the first value to the second prim.

Use mask to update only some prims in the query. It contains one bit per prim addressed by the write: a set bit writes that prim; an unset bit leaves its existing value unchanged. The input still needs a value for every prim addressed by the write, including those skipped by the mask.

In C, supply an array of uint64_t; in Python, supply a list of integers. Each integer masks 64 prims. The least significant bit of the first integer controls the first prim, the next bit controls the second, and so on. The second integer starts at the 65th prim. For four prims [A, B, C, D], mask=[0b1001] writes only A and D.

A write can use index_map or mask, but not both. count specifies how many prims the write addresses, starting at the beginning of the query. In C, set it explicitly when using either option. Python can infer it from the length of index_map; with mask, set it explicitly. Supply at least (count + 63) // 64 mask integers. Bits beyond count are ignored.

See Writing Attributes for index bounds and packed array rules. Packed data must divide evenly across the inferred row count, with a whole number of dtype elements per row; otherwise the write is rejected. See ovstage.Stage.write_attribute() for the Python parameters and ovstage_write_data_t for the C descriptor.

The query below selects A, B, C, and D, initially holding 10, 20, 30, and 40. The first write uses index_map to assign 222, 333, and 111 to A, B, and C. The second uses mask to change only D to 44. In C, tensor(values, shape) builds a compact CPU float32 tensor; three and four are shapes {3} and {4}. readFour() copies the four current values into the verification buffer float values[4].

Python

# index_map chooses source rows; mask chooses target prims.
# These are separate writes because the two options cannot be combined.
stage.write_attribute(
    query, heat, ordinal=2, tensors=np.array([111.0, 222.0, 333.0], np.float32),
    is_array=False, index_map=[1, 2, 0],  # A<-row1(222), B<-row2(333), C<-row0(111)
).wait()
stage.advance_write_floor(ordinal=2).wait()
assert _values(stage, query, heat, 2, order) == [222.0, 333.0, 111.0, 40.0]

stage.write_attribute(
    query, heat, ordinal=3, tensors=np.array([0.0, 0.0, 0.0, 44.0], np.float32),
    is_array=False, mask=[0b1000], count=4,  # bit 3 set -> only D (row 3)
).wait()
stage.advance_write_floor(ordinal=3).wait()

C

// index_map chooses source rows; mask chooses target prims.
// These are separate writes because the two options cannot be combined.
float gathered[] = { 111.0f, 222.0f, 333.0f };
const uint32_t indexMap[] = { 1, 2, 0 }; // A<-row1(222), B<-row2(333), C<-row0(111)
DLTensor gatheredTensor = tensor(gathered, three);
ovstage_write_data_t gatherWrite{};
gatherWrite.tensors = &gatheredTensor;
gatherWrite.tensor_count = 1;
gatherWrite.is_array = false;
gatherWrite.count = 3;
gatherWrite.index_map = indexMap;
ASSERT_TRUE(waitOp(stage_,
    ovstage_write_attribute(stage_, query_, { heat_, {} }, 2, gatherWrite, OVSTAGE_PRIM_MODE_UPSERT), "gather"));
seal(2);

readFour(2, values);
EXPECT_FLOAT_EQ(values[0], 222.0f);
EXPECT_FLOAT_EQ(values[1], 333.0f);
EXPECT_FLOAT_EQ(values[2], 111.0f);
EXPECT_FLOAT_EQ(values[3], 40.0f);

float masked[] = { 0.0f, 0.0f, 0.0f, 44.0f };
const uint64_t maskBits[] = { 0b1000ull }; // bit 3 set -> only D (row 3)
DLTensor maskedTensor = tensor(masked, four);
ovstage_write_data_t maskWrite{};
maskWrite.tensors = &maskedTensor;
maskWrite.tensor_count = 1;
maskWrite.is_array = false;
maskWrite.count = 4;
maskWrite.mask = maskBits;
ASSERT_TRUE(waitOp(stage_,
    ovstage_write_attribute(stage_, query_, { heat_, {} }, 3, maskWrite, OVSTAGE_PRIM_MODE_UPSERT), "mask"));
seal(3);

Writing through Mapped Buffers#

Mapping without placement gives you writable CPU staging buffers. The experimental Experimental Device Placement also supports CUDA map buffers for fixed-size attributes. They need not share memory with the final attribute storage; unmap may copy or scatter the values. The call order is:

Call ovstage.Stage.map_attribute() with an ordinal and attribute. Creating an attribute requires its dtype.
Iterate ovstage.Map.groups() and write each buffer’s values.
Call ovstage.Map.unmap_group() for each group, then ovstage.Map.unmap() to finish remaining writes and release the map. A with block calls the final unmap for you. Leave cuda_event and cuda_stream unset for these CPU buffers.
Call ovstage.Stage.advance_write_floor() and wait before reading the changes.

Ready means that staging buffers are available. It does not create visible prims or attributes, change their schemas or parents, or record changes or ordinal ownership. Each successful unmap commits its group’s values and structural changes. Final unmap commits only groups that remain, then closes the session; earlier successful commits remain if a later group fails.

Unmap rechecks the current schema, prim-creation policy, ordinal admission, and read overlap. A conflicting dtype or authored semantic requires deleting the attribute first and mapping at a later ordinal, or setting attribute_mode=AttributeMode.FORCE with an explicit dtype to replace the schema when each group commits. An omitted semantic adopts an unambiguous compatible existing one; specify dtype and semantic when targets have multiple same-name identities. Changes to unrelated prims do not by themselves invalidate the map. Finish unmapping before advancing the watermark to or beyond the map’s ordinal.

In Python, groups fetched through a Map retain that map. Keep the group or its DLPack producer alive while using a view. Stop accessing the group and all previously obtained views as soon as its group commit or final unmap submission is accepted. A synchronously rejected submission leaves it usable. Views already returned cannot check whether they are still valid.

Note

Prefer mapping.fetch_next(), mapping.groups(), and mapping.unmap() on the Python Map object. The lower-level Stage methods also accept the map’s integer ID: stage.fetch_map_next(mapping.handle) returns a group that does not retain the Map, and stage.unmap_attribute(mapping.handle) does not mark the Python Map or its groups as unusable. Passing the Map object instead of its integer ID preserves that lifetime tracking.

A successful transform unmap makes derived transforms need updating; use the computation sequence in Transforms. A pending or failed map does not supply a committed parent. Children reconnect when that parent is later committed successfully. For failed-map cleanup, see Error Handling and Diagnostics.

The snippets below map an existing scalar attribute on two prims, write 10 and 20, and finish the map before advancing the watermark to 2. Mapped buffers are write-only and uninitialized: fill every selected element. The index calculations associate each buffer row with its prim. Python’s with block performs the final unmap; C calls it explicitly.

Python

fill = {0: 10.0, 1: 20.0}  # by list-relative prim index (A, B)
with stage.map_attribute(query, existing, ordinal=2) as session:
    session.wait()
    for group in session.groups():
        buffer = group.array(0)
        for local in range(group.prim_count):
            row = group.data_row_index(local) if group.has_data_index_map else local
            buffer[row] = fill[group.prim_index(local)]
        session.unmap_group(group).wait()  # streaming commit of this group
stage.advance_write_floor(ordinal=2).wait()

C

ovstage_map_desc_t existingDesc{};
existingDesc.attribute = { existing, {} };
existingDesc.prim_mode = OVSTAGE_PRIM_MODE_UPSERT;
ovstage_map_handle_t map = OVSTAGE_INVALID_MAP_HANDLE;
ASSERT_TRUE(waitOp(stage_,
    ovstage_map_attribute(stage_, query_, &existingDesc, 2, nullptr, 0, &map), "map existing"));
for (;;)
{
    ovstage_map_group_t mapGroup{};
    const ovstage_api_status_t fetched = ovstage_fetch_map_next(stage_, map, OVSTAGE_TIMEOUT_INFINITE, &mapGroup);
    if (fetched == OVSTAGE_ERROR_END_OF_ITERATION)
        break;
    ASSERT_EQ(fetched, OVSTAGE_OK);
    float* buffer = static_cast<float*>(mapGroup.data.tensors[0].data);
    for (uint32_t local = 0; local < mapGroup.prims.count; ++local)
    {
        const size_t listIndex = mapGroup.prims.index_map ? mapGroup.prims.index_map[local]
                                                          : mapGroup.prims.offset + local;
        const uint32_t row = mapGroup.data.index_map ? mapGroup.data.index_map[local] : local;
        buffer[row] = (listIndex == 0) ? 10.0f : 20.0f;
    }
    ASSERT_TRUE(waitOp(stage_, ovstage_unmap_group(stage_, map, &mapGroup, ovstage_cuda_sync_t{}), "unmap_group"));
}
ASSERT_TRUE(waitOp(stage_, ovstage_unmap_attribute(stage_, map, ovstage_cuda_sync_t{}), "unmap_attribute"));
seal(2);