Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 3 additions & 1 deletion docs/python/compute/index.rst
Original file line number Diff line number Diff line change
Expand Up @@ -52,7 +52,9 @@ API conventions

* **Output semantics** — Algorithms write results into a user-provided array or iterator
rather than returning them. This keeps memory ownership explicit and lifetimes under
your control.
your control. As with any asynchronous CUDA launch, ``cuda.compute`` does not keep
the arrays you pass alive: keep your own references until the work submitted on the
stream has completed.
Comment on lines +55 to +57

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🩺 Stability & Availability | 🟡 Minor | ⚡ Quick win

important: Include captured arrays in the lifetime rule. The phrase “arrays you pass” does not cover an array captured by an operator, as in test_caching.py. The stateful operator packs that array’s device pointer into runtime state. If the caller releases the array before asynchronous work completes, GPU work can use a freed allocation. Tell callers to retain captured arrays and arrays behind iterators until stream work completes. (raw.githubusercontent.com) As per path instructions, documentation changes must “focus on technical accuracy.”

Source: Path instructions


* **Operators** — Many algorithms accept an ``op`` parameter. This can be a built-in
:class:`OpKind <cuda.compute.op.OpKind>` value or a
Expand Down
6 changes: 3 additions & 3 deletions python/cuda_cccl/cuda/compute/_bindings.pyi
Original file line number Diff line number Diff line change
Expand Up @@ -129,9 +129,9 @@ class Value:

class Pointer:
def __init__(self, arg): ...
def rebind(self, ptr: int | ctypes.c_void_p, owner: Any) -> None: ...
def rebind(self, ptr: int | ctypes.c_void_p) -> None: ...

def make_pointer_object(ptr: int | ctypes.c_void_p, owner: Any) -> Pointer: ...
def make_pointer_object(ptr: int | ctypes.c_void_p) -> Pointer: ...

class IteratorState(Buffer):
def __init__(self, arg): ...
Expand Down Expand Up @@ -170,7 +170,7 @@ class Iterator:
def as_bytes(self) -> bytes: ...
@property
def is_ptr_kind(self) -> bool: ...
def bind_pointer_state(self, ptr: int | ctypes.c_void_p, owner: Any) -> None: ...
def bind_pointer_state(self, ptr: int | ctypes.c_void_p) -> None: ...
@property
def host_advance_fn(self): ...
@host_advance_fn.setter
Expand Down
34 changes: 16 additions & 18 deletions python/cuda_cccl/cuda/compute/_bindings_impl.pyx
Original file line number Diff line number Diff line change
Expand Up @@ -609,33 +609,31 @@ cdef class Pointer(StateBase):
)
self.set_state(ptr, ref)

def rebind(self, ptr, owner):
"""Update the ptr and ref in place"""
def rebind(self, ptr):
"""Update the ptr in place from a raw address.
Comment thread
oleksandr-pavlyk marked this conversation as resolved.
Resolved

Clears ``ref``: a raw address has no Python object to keep alive, and
recording one here (e.g. the device array the address came from) would
keep it alive for as long as this Pointer does, which for the Pointer
cached on an Iterator is for as long as the algorithm's wrapper cache
entry lives.
"""
if isinstance(ptr, int):
self.ptr = int_as_ptr(ptr)
elif isinstance(ptr, ctypes.c_void_p):
self.ptr = int_as_ptr(ptr.value)
else:
raise TypeError(
"First argument must be an integer, or ctypes.c_void_p, "
"Argument must be an integer, or ctypes.c_void_p, "
f"got {type(ptr)}"
)
self.ref = owner
self.ref = None


def make_pointer_object(ptr, owner):
def make_pointer_object(ptr):
"""Make a Pointer from a raw address."""
cdef Pointer res = Pointer(0)

if isinstance(ptr, int):
res.ptr = int_as_ptr(ptr)
elif isinstance(ptr, ctypes.c_void_p):
res.ptr = int_as_ptr(ptr.value)
else:
raise TypeError(
"First argument must be an integer, or ctypes.c_void_p, "
f"got {type(ptr)}"
)
res.ref = owner
res.rebind(ptr)
return res


Expand Down Expand Up @@ -840,15 +838,15 @@ cdef class Iterator:
self.is_ptr_kind = (it_kind == cccl_iterator_kind_t.POINTER)
self._cached_ptr_obj = None

def bind_pointer_state(self, ptr, owner):
def bind_pointer_state(self, ptr):
"""Set state from a raw pointer, reusing a cached Pointer instead
of allocating one each call. Only valid when is_ptr_kind is True.
"""
cdef Pointer cached = self._cached_ptr_obj
if cached is None:
cached = Pointer(0)
self._cached_ptr_obj = cached
cached.rebind(ptr, owner)
cached.rebind(ptr)
Comment thread
oleksandr-pavlyk marked this conversation as resolved.
Resolved
self.state = cached

@property
Expand Down
141 changes: 121 additions & 20 deletions python/cuda_cccl/cuda/compute/_caching.py
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,7 @@
from __future__ import annotations

import functools
import hashlib
import threading
import types
import weakref
Expand All @@ -19,7 +20,13 @@
except ImportError:
from cuda.core.experimental._utils.cuda_utils import CUDAError

from ._utils.protocols import get_dtype, get_shape, is_device_array
from ._utils.protocols import (
get_data_pointer,
get_dtype,
get_shape,
is_contiguous,
is_device_array,
)
from .struct import _Struct

try:
Expand Down Expand Up @@ -499,6 +506,32 @@ def build_privately():
)


def _proxy_device_arrays(args: tuple, kwargs: dict) -> tuple[tuple, dict]:
"""Replace device-array arguments with dtype-only ``ProxyArray`` placeholders.

A factory builds an algorithm from its arguments' *types*; the arrays' data
is bound at ``__call__``. Building from proxies therefore produces the same
algorithm, but the wrapper, which lives in the process-wide wrapper cache,
is never constructed from (and so can never retain) a caller's allocation,
and the single-phase and ahead-of-time (``ProxyArray``) build paths are
identical. Non-contiguous arrays are left as is so the build keeps raising
the same error it does today. Objects that merely contain arrays (iterators,
``DoubleBuffer``) are not device arrays and pass through unchanged.
"""
from ._proxy import ProxyArray

def proxy(value):
if (
is_device_array(value)
and not isinstance(value, ProxyArray)
and is_contiguous(value)
):
return ProxyArray(get_dtype(value))
return value

return tuple(proxy(a) for a in args), {k: proxy(v) for k, v in kwargs.items()}


class _CacheWithRegisteredKeyFunctions:
"""
Decorator to cache the result of the decorated function.
Expand Down Expand Up @@ -573,8 +606,9 @@ def inner(*args, **kwargs):
)
_thread_local.factory_device_info = target
try:
build_args, build_kwargs = _proxy_device_arrays(args, kwargs)
with target_cc(target_cc_arg):
result = func(*args, **kwargs)
result = func(*build_args, **build_kwargs)
finally:
_thread_local.factory_device_info = previous_device_info
cache[cache_key] = result
Expand Down Expand Up @@ -604,19 +638,65 @@ def register(self, type_: type, key_function: Callable[[Any], Hashable]) -> None
cache_with_registered_key_functions = _CacheWithRegisteredKeyFunctions()


def _make_hashable(value):
"""Hash a global/closure variable for caching purposes."""
def _function_key(func, in_progress: tuple):
"""Key a captured function by content, unless its key is already being
built further up the walk (the operator refers to itself, directly or via
another function). Such a back-reference is keyed by the position of that
function in ``in_progress``; identical walks produce identical positions,
and descending again would never terminate.
"""
for depth, active in enumerate(in_progress):
if active is func:
return ("recursive", depth)
return CachableFunction(func, _in_progress=in_progress)


def _make_hashable(value, keep_alive: list, in_progress: tuple = ()):
"""Hash a global/closure variable for caching purposes.

Values that cannot be keyed by content fall back to ``id()``. Such a value
is also appended to ``keep_alive`` so the caller can keep it alive. This
prevents cache key collisions: once the object is freed, Python may give
its address to an unrelated object, and two different operators would
then produce equal keys.

``in_progress`` holds the functions whose keys are being built up the
walk; see :func:`_function_key`.
"""
# duck-type check for a JIT backend's dispatcher:
if hasattr(value, "py_func") and callable(value.py_func):
return CachableFunction(value.py_func)
return _function_key(value.py_func, in_progress)
elif isinstance(value, types.FunctionType):

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

A module-level operator such as def add(x, y): return np.add(x, y) now raises RecursionError while constructing its cache key.

co_names includes the attribute name add; the existing globals lookup resolves that name to the operator itself, and the new function branch recursively hashes it forever.

# A captured plain Python function (e.g. a helper called from the
# operator) is keyed by its own content, recursively.
return _function_key(value, in_progress)
Comment thread
coderabbitai[bot] marked this conversation as resolved.
Resolved
elif is_device_array(value):
if len(in_progress) > 1:
# A device array captured by a *nested* function (a helper the
# operator calls) is not threaded through the runtime state; the
# compiler bakes its pointer, shape and strides into the device
# code as constants. Key it by exactly those, so two operators
# whose helpers capture different arrays never share a kernel.
cai = value.__cuda_array_interface__
return (
"device_array.baked",
get_dtype(value),
tuple(cai["shape"]),
cai.get("strides"),
get_data_pointer(value),
)
# The stateful op machinery threads the array's pointer and
# shape through the *runtime* state buffer. Thus, the key
# includes only dtype and rank, not pointer or shape.
return (get_dtype(value), len(get_shape(value)))
elif isinstance(value, np.ndarray):
# A captured host array is frozen into the compiled operator as a
# constant, so arrays with different contents are different operators.
digest = hashlib.sha256(np.ascontiguousarray(value).tobytes()).digest()
return ("numpy.ndarray", value.dtype, value.shape, digest)
elif isinstance(value, (np.number, np.bool_)):
return ("numpy.scalar", value.dtype.str, value.tobytes())
elif isinstance(value, (bool, int, float)):
elif isinstance(value, (bool, int, float, str, bytes)):
# Python scalars are immutable values; key them by type and value so
# equal-valued scalars share a cache entry. Without this they fall
# through to ``id(value)`` below, and a fresh (non-interned) ``int``/
Expand All @@ -625,12 +705,19 @@ def _make_hashable(value):
# collisions between like-named scalar subclasses from other modules).
return ("python.scalar", _type_fqn(value), value)
elif isinstance(value, (list, tuple)):
return tuple(_make_hashable(v) for v in value)
return tuple(_make_hashable(v, keep_alive, in_progress) for v in value)
elif isinstance(value, dict):
return tuple(
sorted((_make_hashable(k), _make_hashable(v)) for k, v in value.items())
sorted(
(
_make_hashable(k, keep_alive, in_progress),
_make_hashable(v, keep_alive, in_progress),
)
for k, v in value.items()
)
)
else:
keep_alive.append(value)
return id(value)


Expand Down Expand Up @@ -662,7 +749,7 @@ def clear_all_caches():
_clear_wrapper_caches()
_process_wide_build_results_cache.clear()
# Auxiliary caches registered process-wide (_jit._infer_return_type, the
# compiled-operator memo _jit._compile_op_impl, and the NVRTC memos in
# compiled-operator memo _jit.compile_op, and the NVRTC memos in
# _cpp_compile) must be cleared too, so builds after a clear really are
# cold. Factory entries' cache_clear is idempotent with
# _clear_wrapper_caches above. Snapshot the registry first: a lazy import
Expand All @@ -674,25 +761,31 @@ def clear_all_caches():

class CachableFunction:
"""
A type that wraps a function and provides custom comparison
(__eq__) and hash (__hash__) implementations.

The purpose of this class is to enable caching and comparison of
functions based on their bytecode, constants, and closures, while
ignoring other attributes such as their names or docstrings.
A cache key for a function, with comparison (__eq__) and hash (__hash__)
based on its bytecode, constants, and closures, while ignoring other
attributes such as its docstring.

It does not hold a reference to the function it was built from. It does
hold a reference to any captured object that could only be keyed by its
``id()`` (see :func:`_make_hashable`), so that address stays valid for
the lifetime of the key.
"""

# TODO: eventually, move this class to _jit.py as it only
# has to do with caching of Python callables that will be
# JIT compiled.
def __init__(self, func):
self._func = func
def __init__(self, func, _in_progress: tuple = ()):
self._name = func.__name__

# Functions whose keys are being built, outermost first, so a
# reference back to one of them does not recurse forever.
in_progress = (*_in_progress, func)
closure = func.__closure__ if func.__closure__ is not None else []
keep_alive: list = []
contents = []
# Make closure contents hashable
for cell in closure:
contents.append(_make_hashable(cell.cell_contents))
contents.append(_make_hashable(cell.cell_contents, keep_alive, in_progress))
self._identity = (
func.__name__,
func.__code__.co_code,
Expand All @@ -706,10 +799,13 @@ def __init__(self, func):
# use the hash of the name itself. Assumes numba
# known how to interpret the dotted name at JIT
# time.
_make_hashable(func.__globals__.get(name, name))
_make_hashable(
func.__globals__.get(name, name), keep_alive, in_progress
)
for name in func.__code__.co_names
),
)
self._keep_alive = tuple(keep_alive)

def __eq__(self, other):
return self._identity == other._identity
Expand All @@ -718,7 +814,12 @@ def __hash__(self):
return hash(self._identity)

def __repr__(self):
return str(self._func)
# Iterators derive device symbol names from str() of their cache key
# (see IteratorBase._get_uid), so the repr must distinguish functions
# that compare unequal: two same-named operators with different bodies
# must not share a symbol. Hash the identity rather than printing it.
digest = hashlib.sha256(repr(self._identity).encode()).hexdigest()[:16]
return f"<CachableFunction {self._name} {digest}>"


# Register keyers for built-in types
Expand Down
27 changes: 19 additions & 8 deletions python/cuda_cccl/cuda/compute/_cccl_interop.py
Original file line number Diff line number Diff line change
Expand Up @@ -97,18 +97,22 @@ def _is_well_known_op(op: OpKind) -> bool:


def _device_array_to_cccl_iter(array: DeviceArrayLike) -> Iterator:
from ._proxy import ProxyArray
from ._proxy import PROXY_PLACEHOLDER_ADDRESS, ProxyArray

if not is_contiguous(array):
raise ValueError("Non-contiguous arrays are not supported.")
dtype = get_dtype(array)

info = _type_info_from_dtype(dtype)
state_info = _type_info_from_dtype(np.intp)
# A ProxyArray has no GPU allocation: leave the pointer NULL for build-time
# (ahead-of-time) compilation. The real pointer is bound at __call__ via
# set_cccl_iterator_state().
state = None if isinstance(array, ProxyArray) else get_data_pointer(array)
# A ProxyArray has no GPU allocation: bind a placeholder address for the
# build (see PROXY_PLACEHOLDER_ADDRESS for why it is not NULL). The real
# pointer is bound at __call__ via set_cccl_iterator_state().
state = (
PROXY_PLACEHOLDER_ADDRESS
if isinstance(array, ProxyArray)
else get_data_pointer(array)
)
return Iterator(
state_info.alignment,
IteratorKind.POINTER,
Expand Down Expand Up @@ -228,15 +232,22 @@ def get_value_type(


def set_cccl_iterator_state(cccl_it: Iterator, input_it):
"""Bind the per-call state of ``input_it`` to the cached ``cccl_it``.

Only the pointer value (or state bytes) is bound; ``input_it`` itself is
not kept alive by ``cccl_it``. The caller holds ``input_it`` for the
duration of the call, which is all the launch needs, whereas ``cccl_it``
Comment on lines +237 to +239

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🩺 Stability & Availability | 🟠 Major | ⚡ Quick win

🔎 Supported by static analysis

🏁 Script executed:

#!/bin/bash
set -euo pipefail
fd -i 'device_array.py|protocols.py'
rg -n -C 4 'def (__del__|close|free|sync)\b|cudaFreeAsync|cudaFree|set_cccl_iterator_state\(' python/cuda_cccl

Repository: NVIDIA/cccl

Length of output: 20385


🏁 Script executed:

set -euo pipefail
printf '%s\n' '--- changed hunk and helper ---'
nl -ba python/cuda_cccl/cuda/compute/_cccl_interop.py | sed -n '210,255p'
printf '%s\n' '--- array protocols ---'
nl -ba python/cuda_cccl/cuda/compute/_utils/protocols.py | sed -n '1,260p'
printf '%s\n' '--- test device array ---'
nl -ba python/cuda_cccl/tests/_utils/device_array.py | sed -n '1,260p'
printf '%s\n' '--- relevant launch call sites ---'
nl -ba python/cuda_cccl/cuda/compute/algorithms/_transform.py | sed -n '60,105p'
nl -ba python/cuda_cccl/cuda/compute/algorithms/_reduce.py | sed -n '145,205p'
printf '%s\n' '--- stream and lifetime references ---'
rg -n -F --glob 'python/cuda_cccl/**' -e 'synchronize' -e 'lifetime' -e 'keep alive' -e 'cudaFree' -e 'cudaMalloc' -e '__cuda_array_interface__' -e 'compute(' python/cuda_cccl | sed -n '1,260p'
printf '%s\n' '--- changed files from merge base ---'
git diff --stat dcee23c7a0e4c5a2036d12c6572c7ac5db89bd89 3decf2d554be2e85130fe71c388a18daefa0d2fb -- python/cuda_cccl/cuda/compute/_cccl_interop.py python/cuda_cccl/tests

Repository: NVIDIA/cccl

Length of output: 35536


🏁 Script executed:

set -euo pipefail
printf '%s\n' '--- binding iterator ownership and compute implementation ---'
nl -ba python/cuda_cccl/cuda/compute/_bindings_impl.pyx | sed -n '220,270p'
nl -ba python/cuda_cccl/cuda/compute/_bindings_impl.pyx | sed -n '590,635p'
nl -ba python/cuda_cccl/cuda/compute/_bindings_impl.pyx | sed -n '1120,1175p'
printf '%s\n' '--- pointer iterator and public typing ---'
nl -ba python/cuda_cccl/cuda/compute/iterators/_pointer.py | sed -n '1,120p'
nl -ba python/cuda_cccl/cuda/compute/typing.py | sed -n '1,100p'
printf '%s\n' '--- docs and tests mentioning async lifetime ---'
rg -n -i -F --glob 'python/cuda_cccl/**' --glob 'docs/**' --glob '*.md' -e 'keep.*alive' -e 'lifetime' -e 'asynchronous' -e 'async' -e 'synchroniz' -e '__cuda_array_interface__' | sed -n '1,320p'
printf '%s\n' '--- PR test diff ---'
git diff --unified=20 dcee23c7a0e4c5a2036d12c6572c7ac5db89bd89 3decf2d554be2e85130fe71c388a18daefa0d2fb -- python/cuda_cccl/tests/compute/test_caching.py | sed -n '1,280p'

Repository: NVIDIA/cccl

Length of output: 41770


🏁 Script executed:

set -euo pipefail
printf '%s\n' '--- CCCL compute declarations and docs ---'
rg -n -C 5 -F --glob '*' -e 'cccl_device_reduce' -e 'cccl_device_transform' -e 'cccl_device_unique_by_key' -e 'stream' python/cuda_cccl/cuda/compute/_bindings_impl.pyx python/cuda_cccl/cuda/compute/_bindings.pyi python/cuda_cccl | sed -n '1,360p'
printf '%s\n' '--- public algorithm entrypoints and docstrings ---'
rg -n -C 8 -F --glob 'python/cuda_cccl/cuda/compute/**' -e 'def reduce_into' -e 'def unary_transform' -e 'def unique_by_key' -e 'stream=None' | sed -n '1,360p'
printf '%s\n' '--- array-interface lifetime wording in repository ---'
rg -n -i -C 5 --glob '*.py' --glob '*.rst' --glob '*.md' -e 'producer.*alive' -e 'consumer.*alive' -e 'keep.*reference' -e 'keep.*alive' -e 'until.*sync' -e 'until.*complete' -e 'lifetime' python docs README.md | sed -n '1,260p'

Repository: NVIDIA/cccl

Length of output: 41826


🌐 Web query:

CUDA Array Interface specification object lifetime producer consumer asynchronous operations

💡 Result:

<source_evidence>
<source kind="open_page">
<title>CUDA Array Interface (Version 3) — Numba CUDA</title>
<location>https://nvidia.github.io/numba-cuda/user/cuda_array_interface.html</location>
<excerpt>citeturn1view0 [wordlim: 200] Crawled: today; Content type: text/html; Source: open({&quot;ref_id&quot;:&quot;turn0search0&quot;,&quot;lineno&quot;:null}); Total lines: 361
L0: Back to top
L1: 
L2: [Input: Search the docs ...] `Ctrl`+`K`
L3: 
L4: Table of Contents
L5:   * 【0†User guide】
L6:     * 【1†Installation】
L7:     * 【2†Writing CUDA Kernels】
L8:     * 【3†Memory management】
L9:     * 【4†Writing Device Functions】
L10:     * 【5†Global Variables and Captured Values】
L11:     * 【6†Supported Python features in CUDA Python】
L12:     * 【7†CUDA Fast Math】
L13:     * 【8†Supported Atomic Operations】
L14:     * 【9†Cooperative Groups】
L15:     * 【10†Random Number Generation】
L16:     * 【11†Device management】
L17:     * 【12†Examples】
L18:     * 【13†Debugging Numba CUDA Programs with Visual Studio Code and CUDA GDB】
L19:     * 【14†Debugging CUDA Python with the the CUDA Simulator】
L20:     * 【15†GPU Reduction】
L21:     * 【16†CUDA Ufuncs and Generalized Ufuncs】
L22:     * 【17†Sharing CUDA Memory】
L23:     * 【18†CUDA Array Interface (Version 3)】
L24:     * 【19†External Memory Management (EMM) Plugin interface】
L25:     * 【20†CUDA Bindings】
L26:     * 【21†Calling foreign functions from Python kernels】
L27:     * 【22†Compiling Python functions for use with other languages】
L28:     * 【23†CUDA device call conventions】
L29:     * 【24†On-disk Kernel Caching】
L30:     * 【25†CUDA Minor Version Compatibility】
L31:     * 【26†CUDA Frequently Asked Questions】
L32:   * 【27†Reference documentation】
L33:     * 【28†CUDA Host API】
L34:     * 【29†CUDA Kernel API】
L35:     * 【30†Cache Hints for Memory Operations】
L36:     * 【31†CUDA-Specific Types】
L37:     * 【32†Memory Management】
L38:     * 【33†Libdevice functions】
L39:     * 【34†Environment Variables】
L40: # CUDA Array Interface (Version 3)【35†#】
L41: 
L42: The CUDA Array Interface (or CAI) is created for interoperability between different implementations of CUDA array-like objects in various projects. The idea is borrowed from the 【36†NumPy array interface†docs.scipy.org】.
L43: 
L44: Note
L45: 
L46: Currently, we only define the Python-side interface. In the future, we may add a C-side interface for efficient exchange of the information in compiled code.
L47: ## Python Interface Specification【37†#】
L48: 
L49: Note
L50: 
L51: Experimental feature. Specification may change.
L52: 
L53: The `__cuda_array_interface__` attribute returns a dictionary (`dict`) that must contain the following entries:
L54: 
L55:   * shape: `(integer, ...)`
L56: 
L57: A tuple of `int` (or `long`) representing the size of each dimension.
L58: 
L59:   * typestr: `str`
L60: 
L61: The type string. This has the same definition as `typestr` in the 【36†NumPy array interface†docs.scipy.org】.
L62: 
L63:   * data: `(integer, boolean)`
L64: The data is a 2-tuple. The first element is the data pointer as a Python `int` (or `long`). The data must be device-accessible. For zero-size arrays, use `0` here. The second element is the read-only flag as a Python `bool`.
L65: Because the user of the interface may or may not be in the same context, the most common case is to use `cuPointerGetAttribute` with `CU_POINTER_ATTRIBUTE_DEVICE_POINTER` in the CUDA driver API (or the equivalent CUDA Runtime API) to retrieve a device pointer that is usable in the currently active context.
L66: 
L67:   * version: `integer`
L68: 
L69: An integer for the version of the interface being exported. The current version is 3.
L70: 
L71: The following are optional entries:
L72: 
L73:   * strides: `None` or `(integer, ...)`
L74: If strides is not given, or it is `None`, the array is in C-contiguous layout. Otherwise, a tuple of `int` (or `long`) is explicitly given for representing the number of bytes to skip to access the next element at each dimension.
L75: 
L76:   * descr
L77: 
L78: This is for describing more complicated types. This follows the same specification as in the 【36†NumPy array interface†docs.scipy.org】.
L79: 
L80:   * mask: `None` or objec…[truncated]</excerpt>
</source>
<source kind="search">
<title>CUDA Array Interface (Version 3) — Numba CUDA</title>
<location>https://nvidia.github.io/numba-cuda/user/cuda_array_interface.html</location>
<excerpt>citeturn0search0 [wordlim: 200] Crawled: 3 months ago; The CAI specification does not prescribe the exact mechanism by which these options are set, or related options that Producers or Consumers might provide to allow the user further control over synchronization behavior. ... Obtaining the value of the `__cuda_array_interface__` property of any object has no effect on the lifetime of the object from which it was created. ... Numba provides two mechanisms for creating device arrays from objects exporting the CUDA Array Interface.
</excerpt>
</source>
<source kind="search">
<title>Numba: docs/source/cuda/cuda_array_interface.rst | Fossies</title>
<location>https://fossies.org/linux/numba/docs/source/cuda/cuda_array_interface.rst</location>
<excerpt>citeturn0search1 [wordlim: 200] Crawled: 3 weeks ago; This follows the same specification as in the NumPy array interface. ... Numba takes no steps to maintain the lifetime of an object from which the interface is exported - it is the user&#39;s responsibility to ensure that the underlying object is kept alive for the duration that the exported interface might be used. ... Numba provides two mechanisms for creating device arrays from objects exporting the CUDA Array Interface.
</excerpt>
</source>
<source kind="search">
<title>Numba Documentation</title>
<location>https://numba.readthedocs.io/_/downloads/en/0.53.1/pdf/</location>
<excerpt>citeturn0search12 [wordlim: 200] Published: 5.6 years ago; Obtaining the value of the `__cuda_array_interface__` property of any object has no effect on the lifetime of the object from which it was created. ... Numba takes no steps to maintain the lifetime of an object from which the interface is exported - it is the user’s responsibility to ensure that the underlying object is kept alive for the duration that the exported interface might be used.
</excerpt>
</source>
<source kind="search">
<title>User guide — Numba CUDA</title>
<location>https://nvidia.github.io/numba-cuda/user/index.html</location>
<excerpt>citeturn0search2 [wordlim: 200] Crawled: 3 months ago;   * CUDA Array Interface (Version 3)    * Python Interface Specification ...       * Lifetime management in Numba
</excerpt>
</source>
<omitted_sources count="28">Additional source records were omitted from this bounded result.</omitted_sources>
</source_evidence>

Citations:

- 1: https://nvidia.github.io/numba-cuda/user/cuda_array_interface.html
- 2: https://nvidia.github.io/numba-cuda/user/cuda_array_interface.html
- 3: https://nvidia.github.io/numba-cuda/user/cuda_array_interface.html
- 4: https://nvidia.github.io/numba-cuda/user/cuda_array_interface.html

🌐 Web query:

official CCCL DeviceReduce stream asynchronous API documentation

💡 Result:

<source_evidence>
<source kind="search">
<title>Streams — CUDA Core Compute Libraries</title>
<location>https://nvidia.github.io/cccl/unstable/libcudacxx/runtime/stream.html</location>
<excerpt>citeturn0search0 [wordlim: 200] Crawled: 3 weeks ago; It is passed as an argument to all asynchronous operations like kernel launch, memory copy and allocations. ... Availability: CCCL 2.2.0 / CUDA 12.3 ...        // Pass to a stream-ordered API

# Streams#

Stream is conceptually a queue of operations for a specific device. It is passed as an argument to all asynchronous operations like kernel launch, memory copy and allocations.

## `cuda::stream_ref`#

`cuda::stream_ref` is a non-owning wrapper around a `cudaStream_t`. It prevents unsafe implicit constructions from `nullptr` or integer literals and provides convenient helpers for:

  * `sync()`: wait for the recorded work to complete

  * `is_done()`: non-blocking completion query

  * comparison operators against other `cuda::stream_ref` or `cudaStream_t`

Availability: CCCL 2.2.0 / CUDA 12.3

Example:
    
    #include &lt;cuda/stream&gt;
    
    cudaStream_t stream;
    cudaStreamCreate(&amp;stream);
    cuda::stream_ref ref{stream};
    
    ref.sync();            // synchronizes the stream via cudaStreamSynchronize
    assert(ref.is_done()); // verifies that the stream has finished all operations via cudaStreamQuery
    
    // compare against other stream_ref or cudaStream_t
    assert(ref == stream);
    assert(ref != cuda::invalid_stream);
    
    cudaStreamDestroy(stream);

## `cuda::stream`#

`cuda::stream` is an owning wrapper around a `cudaStream_t` that manages the lifetime of the underlying CUDA stream. It derives from `cuda::stream_ref`, provides all of its functionality, and can be used anywhere a `cuda::stream_ref` is expected. It can be constructed for a specific `cuda::device_ref`, moved (but not copied), and converted from or to a `cudaStream_t` via `from_native_handle`/`release()`. Constructing a new `cuda::stream` always creates a non-blocking stream; see non-blocking stream creation for CUDA Runtime interop details.

Availability: CCCL 3.1.0 / CUDA 13.1
    
    #include &lt;cuda/stream&gt;
    #include &lt;cuda/devices&gt;
    
    int main() {
      {
        // Create a stream on a specific device
        cuda::stream s{cuda::devices[0]};
    
       // Pass to a stream-ordered API
    
        // Synchronize the stream
        s.sync();
      } // Stream is automatically destroyed here
    }</excerpt>
</source>
<source kind="search">
<title>cub::DeviceReduce — CUDA Core Compute Libraries</title>
<location>https://nvidia.github.io/cccl/unstable/cub/api/structcub_1_1DeviceReduce.html</location>
<excerpt>citeturn0search2 [wordlim: 200] Crawled: 4 days ago; &gt;       __host__ __device__ constexpr auto operator()(cuda::compute_capability cc) const -&gt; cub::ReducePolicy ... &gt;     const auto error = cub::DeviceReduce::Reduce( ... The problem size is read in stream order by the reduction kernels. ... See Two-Phase API (explicit temporary storage management) for usage guidance.

# cub::DeviceReduce#

struct DeviceReduce#
    

DeviceReduce provides device-wide, parallel operations for computing a reduction across a sequence of data items residing within device-accessible memory.

## Overview#

A reduction (or fold) uses a binary combining operator to compute a single aggregate from a sequence of input elements.

## Usage Considerations#

&gt;   * Dynamic parallelism. DeviceReduce methods can be called within kernel code on devices in which CUDA dynamic parallelism is supported.
&gt; 
&gt;   * Determinism. The default reproducibility guarantee is `run_to_run`. A different guarantee can be requested through the execution environment with `cuda::execution::require`. See Determinism in CUB for the supported guarantees.

## Performance#

The work-complexity of reduction, reduce-by-key, and run-length encode as a function of input size is linear, resulting in performance throughput that plateaus with problem sizes large enough to saturate the GPU.
## Deferred problem sizes#

`Reduce`, `Sum`, `Min`, `Max`, and `TransformReduce` allow specifying the problem size from a value that resides in device memory through a single-value `cuda::args::deferred` argument. The deferred source may be a device pointer, a one-element span, or a random-access fancy iterator whose element is a non-`bool` 32- or 64-bit integer.

The problem size is read in stream order by the reduction kernels. Work that produces the count in the same stream is ordered automatically; a producer in another stream requires an event or an equivalent dependency. The source and its value must remain accessible and unchanged until all reduction kernels complete. The value must be nonnegative and `[d_in, d_in + num_items)` must be accessible. A temporary-storage query does not dereference the deferred source.

Deferred reductions are CUDA Graph capturable. The pointed-to count may change between graph replays without updating or recapturing the graph. Compile-time and runtime bounds are accepted as caller preconditions, but do not currently change temporary storage, grid dimensions, or pass selection. When those conditions are not met, the call falls back to `run_to_run` rather than failing.

Public Static Functions

template&lt;typename InputIteratorT, typename OutputIteratorT, typename ReductionOpT, typename T, typename NumItemsT&gt;
static inline cudaError_t Reduce(

    void *d_temp_storage,
    size_t &amp;temp_storage_bytes,
    InputIteratorT d_in,
    OutputIteratorT d_out,
    NumItemsT num_items,
    ReductionOpT reduction_op,
    T init,
    cudaStream_t stream = nullptr
)#
    

Computes a device-wide reduction using the specified binary `reduction_op` functor and initial value `init`.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

  * Does not support binary reduction operators that are non-commutative.

  * Provides “run-to-run” determinism for pseudo-associative reduction (e.g., addition of floating point types) on the same GPU device. However, results for pseudo-associative reduction may be inconsistent from one device to a another device of a different compute-capability because CUB can employ different tile-sizing for different architectures.

  * The range `[d_in, d_in + num_items)` shall not overlap `d_out`.

  * Temporary storage for this operation. If `d_temp_storage` is `nullptr`, the required size is written to `temp_storage_bytes` without dereferencing iterators or launching kernels. Otherwise, `d_temp_storage` must point to a device-accessible allocation of at least `temp_storage_bytes` bytes. No special alignment is requi…[truncated]</excerpt>
</source>
<source kind="search">
<title>Lower-level API — CUDA Core Compute Libraries</title>
<location>https://nvidia.github.io/cccl/unstable/cudax/stf/lower_level_api.html</location>
<excerpt>citeturn0search1 [wordlim: 200] Crawled: 3 weeks ago; # Lower-level API# ... This object provides access to the local description of the data associated with the task and a CUDA stream that can be used to submit work asynchronously.
        }
    
        auto lX = ctx.logical_data(X);
        auto lY = ctx.logical_data(Y);
    
        double alpha = 3.14;
    
        /* Compute Y = Y + alpha X */
        auto t = ctx.task(lX.read(), lY.rw());
        t.start();
        slice&lt;double&gt; sX = t.get&lt;0&gt;();
        slice&lt;double&gt; sY = t.get&lt;1&gt;();
        axpy&lt;&lt;&lt;16, 128, 0, t.get_stream()&gt;&gt;&gt;(sX.size(), alpha, sX.data_handle(), sY.data_handle());
        t.end();
    
        ctx.sync();
    }
    

The `ctx.task()` call returns a task object. This object provides access to the local description of the data associated with the task and a CUDA stream that can be used to submit work asynchronously. The beginning of the task body and its end are delimited by the `.start()` and `.end()` calls. Failing to call either of these methods or calling them more than once or in the wrong order results in undefined behavior.

Asynchrony is achieved by using the CUDA stream, which provides a mechanism to submit work on the execution place (here, implicitly the current CUDA device). </excerpt>
</source>
<source kind="search">
<title>CUDA Runtime interactions — CUDA Core Compute Libraries</title>
<location>https://nvidia.github.io/cccl/unstable/libcudacxx/runtime/cudart_interactions.html</location>
<excerpt>citeturn0search3 [wordlim: 200] Crawled: last week;       // Use default_stream with cccl-runtime APIs. ... In the CUDA Runtime API, the blocking/non-blocking stream creation flag controls synchronization with the CUDA default (NULL) stream.
Instead, it can be wrapped into `cuda::stream_ref` when needed for interop.

Note

When wrapping the NULL stream, the current device must be set explicitly first. CUDA binds the NULL stream to the active device, so the wrapper must be created after selecting the correct device.

## Example: wrapping the default stream#
    
    #include &lt;cuda/stream&gt;
    
    void use_default_stream(int device_id) {
      cudaSetDevice(device_id);
    
      cuda::stream_ref default_stream{cudaStreamPerThread};
      // Use default_stream with cccl-runtime APIs.
    }
    

The above applies to Driver API interop cases as well, where the current context must be managed by the user rather than the current device setting.

## Non-blocking stream creation#

Constructing a new `cuda::stream` always creates a stream with CUDA Runtime non-blocking behavior. This is the behavior of CCCL Runtime-created streams; wrapping or taking ownership of an existing `cudaStream_t` preserves the behavior of that handle.

In the CUDA Runtime API, the blocking/non-blocking stream creation flag controls synchronization with the CUDA default (NULL) stream. Because CCCL Runtime treats the default stream as an interop case rather than a first-class object, as described in default stream interop, `cuda::stream` does not expose a blocking/non-blocking construction option.

New Runtime code should express ordering between explicit streams directly, for example by making one `cuda::stream_ref` wait on another. Code that needs legacy CUDA Runtime implicit stream semantics should wrap the relevant CUDA Runtime stream handle in `cuda::stream_ref` (or take ownership with `cuda::stream::from_native_handle`); operations submitted through the wrapper use the same native handle and preserve that handle’s CUDA Runtime semantics, including any default-stream synchronization semantics.</excerpt>
</source>
<source kind="search">
<title>CUDASTF — CUDA Core Compute Libraries</title>
<location>https://nvidia.github.io/cccl/unstable/cudax/stf.html</location>
<excerpt>citeturn0search4 [wordlim: 200] Crawled: yesterday; This illustrates how a task in CUDASTF encapsulates a piece of work that is asynchronous with respect to CUDA stream semantics. ... ### Lower-level task API#

  * `ctx.submit()` initiates the submission of all asynchronous tasks within the sequence

  * `ctx.finalize()` awaits the conclusion of all outstanding asynchronous operations in the context, automatically invoking `ctx.submit()` if not previously called by user code

Usually, creating the task and invoking `ctx.finalize()` is sufficient. However, manually calling `ctx.submit()` can be beneficial in at least two situations. First, it allows for executing additional unrelated work on the CPU (or another GPU) between submission and synchronization. Second, when it’s necessary for two contexts to run concurrently, using the sequence `ctx1.submit(); ctx2.submit(); ctx1.finalize(); ctx2.finalize();` achieves this goal (whereas calling `ctx1.finalize(); ctx2.finalize();` without the `submit()` calls would wait for the completion of the first task before starting the second).

To wait for the completion of all pending operations (tasks, transfers, …), an asynchronous fence mechanism is available :
    
    cudaStream_t stream = ctx.fence();
    cudaStreamSynchronize(stream);
    

Another synchronization mechanism is the `wait` method of the context object. It is typically used in combination with the `reduce()` access mode for dynamic control flow. `auto val = ctx.wait(ld)` is a blocking call that returns the content of the `ld` logical data. The type of the returned value is defined by the `owning_container_of&lt;interface&gt;` trait class where `interface` is the data interface of the logical data. For a logical data with a regular data interface, the `wait` method therefore cannot be called unless that interface overloads this trait class.

This mechanism is illustrated in the dot product example of the Reduce access mode section.

`ctx.wait` can also be called on a token. Since a token carries no content, `ctx.wait(token)` returns `void` and simply blocks the host until the work the token depends on has completed (see the token section for an example).

</excerpt>
</source>
<omitted_sources count="11">Additional source records were omitted from this bounded result.</omitted_sources>
</source_evidence>

Citations:

- 1: https://nvidia.github.io/cccl/unstable/cub/api/structcub_1_1DeviceReduce.html?utm_source=openai
- 2: https://nvidia.github.io/cccl/unstable/cub/api/structcub_1_1DeviceReduce.html?utm_source=openai
- 3: https://nvidia.github.io/cccl/unstable/libcudacxx/runtime/stream.html?utm_source=openai

important: Correct the input lifetime contract.

set_cccl_iterator_state binds only a raw pointer, and the CCCL call can leave work pending on the supplied stream. The CUDA Array Interface does not extend the producer's lifetime. Therefore, keeping input_it alive only until the Python call returns can allow its allocation to be released while queued kernels still use it. Require callers to keep the input and its allocation alive until the submitted work completes, or implement per-call stream-ordered retention. Do not retain it in the process-wide cached iterator.

belongs to an algorithm object that lives in the process-wide wrapper
cache.
"""
if cccl_it.is_ptr_kind:
ptr = get_data_pointer(input_it)
cccl_it.bind_pointer_state(ptr, input_it)
Comment thread
oleksandr-pavlyk marked this conversation as resolved.
Resolved
cccl_it.bind_pointer_state(get_data_pointer(input_it))
else:
state_ = input_it.state
if isinstance(state_, (IteratorState, Pointer)):
cccl_it.state = state_
else:
cccl_it.state = make_pointer_object(state_, input_it)
Comment thread
oleksandr-pavlyk marked this conversation as resolved.
Resolved
cccl_it.state = make_pointer_object(state_)


@functools.lru_cache()
Expand Down
Loading
Loading