One atomic change to the API contract. Only fields relevant to ``kind`` need to be populated; the others default to None / empty. The replaye

Raised when an argument provided to an NCCL4Py API is invalid. Used for argument validation errors that the Python layer detects before forwa

Symmetric memory pool allocator for UB-X communication primitives. Manages a pool of symmetric memory accessible by all ranks in a process group.

Resolves user-provided buffer specifications to raw pointer, count, dtype, and device. This class handles various buffer input formats and extrac

Result for a single (size, backend) benchmark point.

Minimal mock allocator for SymmTensor tests without distributed setup.

A minimal mock that replicates SymmAllocator's pool management logic without requiring symmetric memory or distributed setup. This tests the

Benchmark configuration matching nccl-tests CLI interface.

Pointer wrapper for ``ncclCoopAny``. Construct via :func:`cta` / :func:`warp` / :func:`thread` / :func:`lanes` / :func:`warp_span`. Methods/p

NCCL data type, mirroring :c:type:`ncclDataType_t`. Used as the ``dtype`` of buffer specs and as the ``datatype`` argument of NCCL collective

NCCL unique identifier for communicator initialization. A UniqueId is used to coordinate communicator initialization across multiple ranks. A

``struct ncclGinBarrierHandle { ncclGinSignal_t signal0; ncclDevResourceHandle_t unused; }`` (src/include/nccl_device/impl/gin_barrier__types.h).

``struct ncclTeam { int nRanks, rank, stride; }`` (src/include/nccl_device/core.h).

NCCL communicator for collective and point-to-point operations. A communicator represents a group of ranks that can perform collective operat

Context manager that temporarily switches the current CUDA device. On enter, sets device as the current device if it differs from the previou

Version, CUDA build variant, and loaded-path info for a shared library.

NCCL device communicator requirements configuration. Provides configuration for device communicator creation, allowing fine-tuning of resourc

GIN transport type, mirroring :c:type:`ncclGinType_t`. Reported by :py:attr:`Communicator.gin_type` and :py:attr:`Communicator.railed_gin_typ

``struct ncclLsaBarrierHandle { ncclDevResourceHandle_t bufHandle; int nBarriers; }`` (src/include/nccl_device/impl/lsa_barrier__types.h).

``struct ncclMultimemHandle { void* mcBasePtr; }`` (src/include/nccl_device/impl/core__types.h).

Hybrid barrier — LSA inner stage + GIN outer stage. Constructed via :func:`hybrid_session`. Only ``sync`` is supported.

NCCL user-defined custom reduction operator. Created by :py:meth:`Communicator.create_pre_mul_sum`. The PreMulSum operator performs ``output

NCCL device communicator resource for device-side operations. Wraps ``ncclDevComm_t`` and manages its lifecycle. Created by :py:meth:`Communi

Owning pair of a cudaMalloc'd device buffer and its ``nccl_ep.Tensor``.

Wraps a ``ncclGin_C*`` (pointer to a stack-alloca'd ``ncclGin_C`` struct). Instances must be produced through :meth:`DevComm.gin`.

GIN (network) barrier session — inter-node. Constructed via :func:`gin_session`. Only ``sync`` is supported.

Result of an NCCL group simulation. Returned by :py:func:`group_end` when called with ``simulate=True``.

Per-step routing handle for dispatch/combine. Construct via :py:meth:`Group.create_handle`; release with :meth:`destroy`.

LSA (Locally Shared Address) barrier session — intra-node, peer-access based. Constructed via :func:`lsa_session`.

Raised when a checkpoint shim call returns a non-success status.

Raised when checkpoint shim symbols are unavailable in this process.

NCCL-backed device memory resource for Buffer allocations. Uses NCCL's memory allocation functions to allocate device memory that is optimize

NCCL-allocated symmetric memory pool with a registered window + devcomm. Drop-in replacement for the (pool buffer, multicast_ptr) pair that `

NCCL registered buffer handle for zero-copy optimized communication. Registers a user buffer with the communicator to enable performance opti

NCCL registered window handle for Remote Memory Access (RMA) operations. Registers a memory window with the communicator for one-sided communicat

Custom tensor subclass that uses symmetric memory pool backing. For blocked formats (e.g. 'mxfp8'), the allocation in the pool is larger than

Version information for NCCL4Py and the NCCL library. For the full stack version (including ``libnccl_ep.so``), use :py:func:`nccl.get_versio

Aggregate version snapshot of the NCCL stack.

``struct ncclCoopAny`` — 16-byte aligned-to-ptr storage + vtable ptr (src/include/nccl_device/coop.h).

A set of consecutive warps that the user has also supplied with a unique id from [0..15]. It is an error for two different warp spans with the same id

Pass `ncclGinAllContexts(comm)` to a barrier in place of an `ncclGin` to expand the fence across every GIN context on the comm.

``struct ncclGin_C`` (src/include/nccl_device/gin.h). The bitfield triple ``{nConnections:8, connectionId:8, _ginBackend:8}`` is represented

Used as completion actions for ncclGinSession::put

EP communication algorithm, mirroring :c:type:`ncclEpAlgorithm_t`. Set on :py:attr:`GroupConfig.algorithm` before calling :py:meth:`Group.cre

Allocator hooks for :py:attr:`GroupConfig.alloc`. Mirrors :c:struct:`ncclEpAllocConfig_t`. Leaving every field at 0 selects NCCL EP's default