cuda.coop API Reference#
Warning
cuda.coop is an experimental API and is subject to change.
Common API#
The primitive functions below are compiler markers; register is a
host-side configuration function. The installed .pyi files are
authoritative for overload and result typing.
Backend registration#
- cuda.coop.register(
- backend: Literal['numba-cuda-mlir', 'numba_cuda_mlir'],
Register cooperative primitives with the selected compiler backend.
Call this on the host before compiling a kernel, including when
cuda.coopwas imported before the compiler. Repeated calls are safe.- Parameters:
backend – Compiler backend to register.
"numba_cuda_mlir"is an alias for"numba-cuda-mlir".- Raises:
ValueError – The backend name is unsupported.
ImportError – The backend adapter or its required compiler is unavailable or incompatible. Other backend initialization errors propagate unchanged.
Thread groups#
See thread groups and
participation and synchronization for the shared
execution model. ThreadGroup and ThreadHierarchy describe which
threads cooperate. A descriptor’s availability does not imply that every
primitive supports that group.
- cuda.coop.this_thread() ThreadGroup#
Describe the calling thread as a one-thread group.
- Returns:
A symbolic descriptor. Load and Store do not accept one-thread groups. See thread groups.
- Return type:
- cuda.coop.this_warp() ThreadGroup#
Describe the calling thread’s physical warp.
- Returns:
A group of 32 consecutive threads in the block’s linear thread order. Use
group_by()to form smaller logical warps.- Return type:
Notes
Warp Load and Store require a block size divisible by 32. Constructing this descriptor does not turn a partial final warp into a complete group. See participation requirements.
- cuda.coop.this_block() ThreadGroup#
Describe all threads in the calling thread’s CUDA block.
- Returns:
A descriptor whose size comes from the kernel launch’s block dimensions. For multidimensional blocks, thread ranks are linearized with the x coordinate varying fastest.
- Return type:
Notes
The factory takes no size argument and does not synchronize the block. See thread groups and the primitive’s participation requirements.
- cuda.coop.this_cluster() ThreadGroup#
Describe the thread-block cluster containing the calling thread.
- Returns:
A symbolic descriptor. Cluster groups are not Load or Store targets.
- Return type:
Notes
Construction does not create a cluster or enable cluster scheduling. See thread groups.
- cuda.coop.this_grid() ThreadGroup#
Describe all threads in the current kernel grid.
- Returns:
A symbolic descriptor. Grid groups are not Load or Store targets.
- Return type:
Notes
Construction does not request a cooperative launch or synchronize the grid. See thread groups.
- class cuda.coop.ThreadGroup(
- kind: str,
- hierarchy: ThreadHierarchy = <factory>,
- parent: ThreadGroup | None = None,
- mapping: GroupByMapping | None = None,
- source: str = 'explicit',
Describe the participants in a cooperative operation.
Use the
this_*factories to describe groups in the current kernel launch. Constructing a descriptor does not synchronize threads or launch a kernel. Load and Store support blocks, physical warps, and logical warps; see thread groups.Group descriptors can also be constructed in ordinary Python. Runtime rank, size, membership, and synchronization queries are not available in this API layer.
- group_by( ) _ThreadGroupT#
Partition a physical warp by threads or a block by warps.
- Parameters:
count (int) – Positive compile-time number of units in each subgroup. For a warp parent, the unit is one thread; for a block parent, it is one physical warp. Thus
this_warp().group_by(8)describes eight lanes, whilethis_block().group_by(2)describes 64 threads.exhaustive (bool, optional) – Compile-time flag, default
True. An exhaustive partition must divide the parent’s unit count exactly.Falsepermits a remainder outside the complete groups. Each primitive still determines which partitions it supports.
- Returns:
A descriptor for the subgroup containing the calling thread. Nested partitions are unsupported. Load and Store support logical warp widths of 1, 2, 4, 8, 16, or 32; mapped groups of physical warps are not Load or Store targets.
- Return type:
Examples
Descriptors can be inspected without compiling or launching a kernel:
from cuda import coop group = coop.this_warp().group_by(8) assert group.kind == "threads_within_warp" assert group.static_size == 8
- class cuda.coop.ThreadHierarchy#
Describe the launch dimensions known during group planning.
Public construction describes the current kernel launch with unresolved dimensions. Backends supply exact extents from
LaunchFactswhen they resolve a group. Callers cannot pass independent launch dimensions to the public constructor.- cluster_dim#
Blocks along each axis of one cluster, once needed and known.
- grid_dim#
Clusters along each axis of the resolved grid. On a non-cluster launch, each cluster contains one block.
- class cuda.coop.Hierarchy#
Alias for
ThreadHierarchy.
Payloads and temporary storage#
ThreadData constructs the per-thread payload supplied to a primitive;
thread-group descriptors describe the participating threads.
- cuda.coop.ThreadData( ) ThreadDataLike[Any]#
Construct a fixed-size payload owned by the calling thread.
Each thread has its own slots. See per-thread payloads for their relationship to a group tile and the blocked and striped layouts.
- Parameters:
items_per_thread (int) – Positive compile-time number of items owned by each thread. This extent is fixed for the lifetime of the payload and is available inside the kernel as
items.items_per_thread. Pass it as a compile-time kernel argument using the backend’s supported specialization mechanism.dtype (dtype-like, optional) – Optional numeric element type. Leave unspecified for normal use: a supported producer such as
cuda.coop.load()supplies it. Backend-supported typed assignments can also establish the type. Supply it explicitly when the available type information cannot identify the intended representation. All items have the same type; conflicting typed values are an error.alignment (int, optional) – Compile-time minimum storage alignment in bytes, expressed as a positive power of two.
Nonelets the compiler choose. The request applies when payload storage is materialized; it does not assert alignment of a Load source or Store destination.
- Returns:
Writable per-thread payload with indexed reads and writes. Its contents are uninitialized; write every item before reading it. Construction does not synchronize threads.
- Return type:
- class cuda.coop.ThreadDataLike(*args, **kwargs)#
Mutable fixed-size per-thread payload understood by supported backends.
See per-thread payloads for construction, item access, and dtype requirements.
- cuda.coop.TempStorage(
- size_in_bytes: int | None = None,
- *,
- alignment: int | None = None,
- auto_sync: bool | None = False,
- sharing: str = 'shared',
Describe shared scratch for supported cooperative block operations.
Construct the descriptor inside the kernel and pass it as
temp_storageto operations that accept explicit block scratch. See temporary storage for supported operations, allocation lifetime, and synchronization requirements.- Parameters:
size_in_bytes (int, optional) – Positive compile-time capacity in bytes.
Nonelets the compiler determine the capacity from all uses. An explicit capacity must be large enough for those operations; undersized storage is rejected.alignment (int, optional) – Compile-time minimum alignment in bytes, expressed as a positive power of two.
Nonelets the compiler choose. The allocation satisfies both this request and the operations’ alignment needs.auto_sync (bool, optional) – Whether to insert a trailing barrier after each scratch-using call. Defaults to
False;Nonealso disables automatic reuse synchronization. The caller must synchronize before reusing the scratch, including on the next iteration of a loop. PassTrueto request automatic reuse barriers.sharing ({"shared", "exclusive"}, optional) – Compile-time allocation policy, default
"shared". Calls using the same descriptor can reuse one scratch slice."exclusive"gives distinct call sites separate slices, using more shared memory to avoid barriers needed solely for cross-call scratch reuse whenauto_sync=False. Repeated execution of one call site still reuses its slice and requires synchronization before reuse. Omittingtemp_storageleaves layout and reuse barriers to the compiler.
- Returns:
Compiler-recognized scratch descriptor. The storage contents are opaque; keep application data in
cuda.coop.ThreadData()or application-owned arrays.- Return type:
- class cuda.coop.TempStorageLike(*args, **kwargs)#
Explicit scratch descriptor understood by supported backends.
See temporary storage for construction, allocation sharing, and synchronization.
Memory operations#
- cuda.coop.load(
- group: BlockGroup | WarpGroup,
- source: object,
- output: ThreadDataLike[_CommonNumericT],
- /,
- *,
- algorithm: BlockLoadStoreAlgorithm | WarpLoadStoreAlgorithm = 'direct',
- valid_items: ValidItems | None = None,
- oob_default: _CommonNumericT | float | None = None,
- offset: IntegerValue | None = None,
- temp_storage: TempStorageLike | None = None,
Load a group tile from memory into per-thread values.
- Parameters:
group (cuda.coop.ThreadGroup) – Participating threads; see thread groups. Supports blocks and physical or logical warps. Warp loads require an enclosing block size divisible by 32.
source (array) – One-dimensional contiguous source array in device-accessible memory. Its element dtype must match
output; an untypedThreadDatainfers its dtype from this array. The array must contain all elements selected byoffsetandvalid_items.output (cuda.coop.ThreadDataLike) – Writable per-thread payload. Load populates this payload in place. The group’s tile contains
group_size * items_per_threadelements.algorithm (str, optional) – Compile-time load algorithm, default
"direct"."direct"gives each thread consecutive elements (blocked order);"striped"gives neighboring threads neighboring elements at each item index."vectorize"uses vector accesses when possible, and"transpose"uses shared scratch to rearrange striped accesses into blocked order. Both return blocked order. Blocks also support"warp_transpose"and"warp_transpose_timesliced", which return blocked order and require a block size divisible by 32.valid_items (int or integer scalar, optional) – Number of valid elements in the group’s tile, shared by all threads in that group. Supply a value between zero and the tile size, inclusive.
Noneloads the full tile. Slots beyond this valid prefix have unspecified values unlessoob_defaultis given.oob_default (numeric scalar, optional) – Value written to slots beyond
valid_items. Requires an explicitvalid_itemscount. For example, use zero to pad a partial tile. A runtime value must have the payload dtype and be uniform across the group. WithNone, those slots are unspecified, even if initialized before the Load; assign them before reading them.offset (int or integer scalar, optional) – Nonnegative offset in elements from the start of
source, uniform across the group.Nonemeans zero. For block tiles, supply the block’s starting offset explicitly. For Warp tiles, the backend adds(linear_thread_rank // group_size) * tile_sizeautomatically; do not include that within-block group offset a second time.temp_storage (cuda.coop.TempStorageLike, optional) – Scratch descriptor for block transpose-family algorithms.
Noneuses automatic scratch. Direct, striped, and vectorized loads need no shared scratch. Warp loads requireNone.
- Returns:
The call populates
outputin place.- Return type:
None
See also
cub::BlockLoad,cub::WarpLoadC++ block and warp Load primitives.
- cuda.coop.store(
- group: BlockGroup | WarpGroup,
- destination: object,
- value: _CommonNumericT | CommonThreadDataLike[_CommonNumericT],
- /,
- *,
- algorithm: BlockLoadStoreAlgorithm | WarpLoadStoreAlgorithm = 'direct',
- valid_items: ValidItems | None = None,
- offset: IntegerValue | None = None,
- temp_storage: TempStorageLike | None = None,
Store a group tile from per-thread values into memory.
- Parameters:
group (cuda.coop.ThreadGroup) – Participating threads; see thread groups. Supports blocks and physical or logical warps. Warp stores require an enclosing block size divisible by 32.
destination (array) – Writable one-dimensional contiguous array in device-accessible memory, with the same element dtype as
value. It must contain every element selected byoffsetandvalid_items.value (numeric scalar or cuda.coop.ThreadDataLike) – This thread’s value or readable payload. Initialize every item that will be stored. The tile contains
group_size * items_per_threadelements, with one item per thread for a scalar. As in CUB, transpose algorithms may rearrange the payload in place. Do not rely on its contents after Store; copy values before the call if they are needed later.algorithm (str, optional) – Compile-time store algorithm, default
"direct"."direct"expects blocked values;"striped"expects striped values."vectorize"and"transpose"also expect blocked values and use vector accesses or shared-memory rearrangement, respectively. Blocks additionally support"warp_transpose"and"warp_transpose_timesliced", both requiring a block size divisible by 32. See data layouts before pairing different Load and Store algorithms.valid_items (int or integer scalar, optional) – Number of valid elements in the group’s tile, uniform across the group and between zero and the tile size, inclusive.
Nonestores the entire tile. Elements outside the valid prefix are not written.offset (int or integer scalar, optional) – Nonnegative offset in elements from the start of
destination, uniform across the group.Nonemeans zero. Supply each block’s origin explicitly. Warp stores also add the within-block group origin automatically, using the same addressing rule ascuda.coop.load().temp_storage (cuda.coop.TempStorageLike, optional) – Scratch descriptor for block transpose-family algorithms.
Noneuses automatic scratch. Direct, striped, and vectorized stores need no shared scratch. Warp stores requireNone.
- Returns:
Writes to
destinationand may rearrange the input payload.- Return type:
None
See also
cub::BlockStore,cub::WarpStoreC++ block and warp Store primitives.