Define min for ssize_t

Appends a printf-formatted string to the output buffer. Unlike with INFO or WARN messages, the caller should terminate lines with '\n' as appropriate.

Format a string representation of a (union ncclSocketAddress *) socket address using getnameinfo() * * Output: "IPv4/IPv6 address<port>" */

Element access via reinterpret_cast

Emit a key and separator. Santize the key. This is only acceptable if the top state is an object Emit a ',' separator of we aren't the first item.

Create a SymmTensor using memory from the pool. Args: shape: Tensor shape (visible element dimensions). dtype: Elemen

Coerce ``x`` to an ``!llvm.ptr`` ir.Value. Args: x: accepted forms — * ``!llvm.ptr`` ir.Value — passthrough. * `

Initializes a new NCCL communicator. Creates a communicator that connects multiple ranks. This is a collective operation: all ranks m

Launch a distributed test script with the appropriate launcher. Args: script: Path to the worker script. num_procs: Number of pro

``cute.ffi`` with ``source=_BC`` injected on every prototype. The DSL pipeline merges the bitcode into each module's ``link-libraries`` attri

Raises if the communicator is not initialized. Args: operation: Name of the operation being attempted (for the error

CUDA device associated with this communicator. Returns a :py:class:`cuda.core.Device` providing additional functionality such as ``to

Write an integer value

Splits this communicator into sub-communicators based on color values. Ranks that pass the same ``color`` value will be part of the same

Auto-selects best allreduce variant based on total tensor size.

Number of elements in the buffer.

Reinterpret a Python-side integer (devcomm_handle / window_handle / raw device pointer) into the proper C type. The Phase 3 ABI passes every pointer-s

Device buffer address.

Construct a ``cute.Tensor`` view over the registered buffer. Canonical input to :meth:`Gin.put`: byte offset (relative to the window)

Returns a new SymmTensor view with the same backing memory but different shape. Handles -1 inference manually without temporary tenso

Compute slot assignments for a2av_token_bf16_mxfp8 from a routing matrix. Tokens are block-distributed: rank r owns global token rows [r * lo

Converts the IP number of a NCCL address to a string (the port part is ignored and no DNS resolution is attempted).

Handles the termination of a RAS socket. We try to do it in stages for established sockets (in READY state). We shut down just the sending side for t

Finalizes the communicator, flushing uncompleted operations and network resources. Typically called before :py:meth:`destroy` to ensure all o

Resolves a stream specification to its raw stream handle (as int). Args: stream: A Stream instance, an integer stream handle, an object

Wrap ncclEpTensorCreate (mirrors helper in upstream ep_test.py).

Release the NCCL pool, window, devcomm, and bootstrapped comm. Tears down in the correct order: dev_comm → window → buffer free → boo

Returns a per-communicator unique ID for use with :py:meth:`grow`. Generates a unique identifier bound to this communicator that can be

Validate that a non-null tensor pointer points at an initialised descriptor. Returns the pointer unchanged when null (absent / optional field). Aborts

Validates that the buffer is on the same device as the communicator. Args: buffer: Resolved buffer to validate. buffe

Compute algorithm and bus bandwidth matching nccl-tests formulas. Returns: (algBw_GB/s, busBw_GB/s)