Metadata-Version: 2.4
Name: cute-comm
Version: 0.1.0
Summary: NCCL device-initiated API for the CuTe Python DSL
License: BSD 3-Clause License
        
        Copyright (c) 2026, Bodun Hu
        
        Redistribution and use in source and binary forms, with or without
        modification, are permitted provided that the following conditions are met:
        
        1. Redistributions of source code must retain the above copyright notice, this
           list of conditions and the following disclaimer.
        
        2. Redistributions in binary form must reproduce the above copyright notice,
           this list of conditions and the following disclaimer in the documentation
           and/or other materials provided with the distribution.
        
        3. Neither the name of the copyright holder nor the names of its
           contributors may be used to endorse or promote products derived from
           this software without specific prior written permission.
        
        THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
        AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
        IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
        DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
        FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
        DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
        SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
        CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
        OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
        OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
License-File: LICENSE
Requires-Python: >=3.10
Requires-Dist: cuda-pathfinder>=1.5.0
Requires-Dist: nvidia-cuda-nvcc>=13.2.51
Requires-Dist: nvidia-cutlass-dsl>=4.4.2
Requires-Dist: nvidia-nccl-cu13>=2.29.7
Provides-Extra: dev
Requires-Dist: pytest; extra == 'dev'
Description-Content-Type: text/markdown

# CuTe-comm

NCCL device-initiated API bindings for the [CuTe Python DSL](https://github.com/NVIDIA/cutlass).

Exposes NCCL's GIN (GPU-initiated networking) and LSA (local symmetric access) APIs via `cute.ffi`, with a thin host-side layer for communicator setup.

## Requirements

- Python >= 3.10, CUDA 13, NCCL >= 2.29.7
- GPU with `deviceApiSupport` (check via `query_properties`)

## Installation

```bash
pip install cute-comm
# or from source
uv sync
```

## Usage

```python
import cute_comm.nccl as cn

# Host: bootstrap communicator, register memory
uid = cn.get_unique_id()          # broadcast from rank 0
comm = cn.init_rank(nranks, rank, uid)
dev_comm = cn.create_dev_comm(comm, gin_contexts=1, gin_signals=1, lsa_barriers=1)
win = cn.register_window(comm, buf.data_ptr(), buf.nbytes)

# Device: use inside @cute.jit kernels
@cute.jit
def my_kernel(dev_comm, ctx, win, peer):
    ptr = cn.lsa_ptr(win, offset, peer)
    cn.gin_put_siginc(dev_comm, ctx, peer, dst_win, dst_off, src_win, src_off, nbytes, signal_idx)
    cn.gin_wait_signal(dev_comm, ctx, signal_idx, threshold=1)

# Pass link_options to bundle the bridge PTX (compiled lazily, cached in ~/.cache/cute_comm/)
my_kernel[grid, block, cn.link_options("sm_90a")](dev_comm, ctx, win, peer)

# Cleanup
cn.deregister_window(comm, win)
cn.destroy_dev_comm(comm, dev_comm)
cn.finalize(comm)
cn.destroy(comm)
```

## API

**Host** — `get_unique_id`, `init_rank`, `query_properties`, `create_dev_comm`, `destroy_dev_comm`, `register_window`, `deregister_window`, `finalize`, `destroy`

**Device** (inside `@cute.jit`) — `local_ptr`, `lsa_ptr`, `lsa_barrier`, `gin_put`, `gin_put_siginc`, `gin_put_sigadd`, `gin_read_signal`, `gin_wait_signal`

## License

BSD 3-Clause. See [LICENSE](LICENSE).
