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
6 changes: 3 additions & 3 deletions .github/scripts/run_new_examples.sh
Original file line number Diff line number Diff line change
Expand Up @@ -24,11 +24,11 @@ fi
if [ "$INSTALL_METHOD" = "git" ]; then
REPO=${GITHUB_REPOSITORY:-"ROCm/iris"}
SHA=${GITHUB_SHA:-"HEAD"}
INSTALL_CMD="pip install git+https://github.com/${REPO}.git@${SHA}"
INSTALL_CMD="pip install 'iris[sdma] @ git+https://github.com/${REPO}.git@${SHA}'"
elif [ "$INSTALL_METHOD" = "editable" ]; then
INSTALL_CMD="pip install -e ."
INSTALL_CMD="pip install -e '.[sdma]'"
elif [ "$INSTALL_METHOD" = "install" ]; then
INSTALL_CMD="pip install ."
INSTALL_CMD="pip install '.[sdma]'"
else
echo "[ERROR] Invalid install_method: $INSTALL_METHOD"
exit 1
Expand Down
2 changes: 1 addition & 1 deletion .github/scripts/run_perf_benchmark.sh
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,7 @@ echo "[PERF-BENCHMARK] Using GPUs: $GPU_DEVICES"
set -e

cd /iris_workspace
pip install -e .
pip install -e '.[sdma]'
torchrun --rdzv-backend=c10d --rdzv-endpoint=localhost:0 --nnodes=1 --nproc_per_node=8 examples/${EXAMPLE_PATH}/benchmark.py \
--benchmark \
--validate \
Expand Down
12 changes: 6 additions & 6 deletions .github/scripts/run_tests.sh
Original file line number Diff line number Diff line change
Expand Up @@ -8,9 +8,9 @@
# num_ranks: number of GPU ranks (1, 2, 4, or 8)
# gpu_devices: comma-separated GPU device IDs (optional, if not provided will use allocator)
# install_method: pip install method - "git", "editable", or "install" (optional, default: "editable")
# - "git": pip install git+https://github.com/${{ github.repository }}.git@${{ github.sha }}
# - "editable": pip install -e .
# - "install": pip install .
# - "git": pip install 'iris[sdma] @ git+https://github.com/${{ github.repository }}.git@${{ github.sha }}'
# - "editable": pip install -e '.[sdma]'
# - "install": pip install '.[sdma]'
#
# CONTAINER_IMAGE, if set, selects the container image (e.g. the nightly Triton
# image) instead of the one container_build.sh produced.
Expand Down Expand Up @@ -68,11 +68,11 @@ if [ "$INSTALL_METHOD" = "git" ]; then
# For git install, we need the repository and SHA from environment or use defaults
REPO=${GITHUB_REPOSITORY:-"ROCm/iris"}
SHA=${GITHUB_SHA:-"HEAD"}
INSTALL_CMD="pip install git+https://github.com/${REPO}.git@${SHA}"
INSTALL_CMD="pip install 'iris[sdma] @ git+https://github.com/${REPO}.git@${SHA}'"
elif [ "$INSTALL_METHOD" = "editable" ]; then
INSTALL_CMD="pip install -e ."
INSTALL_CMD="pip install -e '.[sdma]'"
elif [ "$INSTALL_METHOD" = "install" ]; then
INSTALL_CMD="pip install ."
INSTALL_CMD="pip install '.[sdma]'"
fi

# Run tests in container
Expand Down
2 changes: 1 addition & 1 deletion .github/workflows/copilot-setup-steps.yml
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,7 @@ jobs:
/opt/venv/bin/python -m venv --system-site-packages .venv
source .venv/bin/activate
python -m pip install --upgrade pip
python -m pip install -e . --no-cache-dir
python -m pip install -e ".[sdma]" --no-cache-dir

- name: Install native build dependencies
run: |
Expand Down
6 changes: 6 additions & 0 deletions README.md
Original file line number Diff line number Diff line change
Expand Up @@ -203,6 +203,12 @@ For a quick installation directly from the repository:
pip install git+https://github.com/ROCm/iris.git
```

The SDMA copy engine needs the optional `sdma` extra ([rocm-xio](https://github.com/ROCm/rocm-xio)):

```shell
pip install "iris[sdma] @ git+https://github.com/ROCm/iris.git"
```

### Docker Compose (Recommended for Development)

The recommended way to get started is using Docker Compose, which provides a development environment with the Iris directory mounted inside the container. This allows you to make changes to the code outside the container and see them reflected inside.
Expand Down
20 changes: 20 additions & 0 deletions docs/getting-started/installation.md
Original file line number Diff line number Diff line change
Expand Up @@ -94,6 +94,26 @@ pip install -e .
```


## Optional: SDMA copy engine

Shader load/store, atomics, and collectives work with the base install. The SDMA
copy engine (`use_copy_engine=True` in device code, and host-initiated `put`,
`put_tile`, `put_tiles`, and `quiet`) needs [rocm-xio](https://github.com/ROCm/rocm-xio),
which is an optional extra:

```shell
pip install "iris[sdma] @ git+https://github.com/ROCm/iris.git"
# or, from a checkout
pip install -e ".[sdma]"
```

To add SDMA to an existing Iris install, rerun the install with the extra. If an
older rocm-xio is already installed, run `pip uninstall rocm-xio` first: every
rocm-xio commit has the same version number, so pip keeps the installed one.

Installing rocm-xio does not turn SDMA on; see [Iris Class](../reference/host/class.md)
for enabling it.

## Selecting the AMD runtime for fabric communication

When multiple ROCm versions are installed, set `LD_LIBRARY_PATH` before starting
Expand Down
23 changes: 23 additions & 0 deletions docs/reference/host/class.md
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,29 @@ Prefer using the convenience factory over calling the constructor directly:
.. autofunction:: iris.host.iris.iris
```

## Optional SDMA Copy Engine

The `engines` parameter defaults to `()`. In this mode, Iris skips SDMA
initialization and queue creation, and `get_copy_engine_ctx()` returns `None`.
Regular Triton and Gluon memory operations and collectives work without SDMA.

Pass `engines=("sdma",)` to initialize both host- and device-initiated SDMA queues.
This needs the `sdma` extra ([Installation](../../getting-started/installation.md)):

```python
ctx = iris.iris(heap_size=1 << 30, engines=("sdma",))
copy_engine_ctx = ctx.get_copy_engine_ctx()
```

This is required for the host methods `put`, `put_tile`, `put_tiles`, and `quiet`,
and for device operations using `use_copy_engine=True`. Host SDMA methods raise
`RuntimeError` if the copy engine is disabled. Pass the returned context to
device operations through their `copy_engine_ctx` argument.

```{eval-rst}
.. automethod:: iris.host.iris.Iris.get_copy_engine_ctx
```

## Core Methods

```{eval-rst}
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -174,7 +174,7 @@ def _worker(local_rank: int, world_size: int, init_url: str, args: dict):
)

# Main benchmark logic
shmem = iris.iris(args["heap_size"])
shmem = iris.iris(args["heap_size"], engines=("sdma",))
dtype = torch_dtype_from_str(args["datatype"])
cur_rank = shmem.get_rank()
world_size = shmem.get_num_ranks()
Expand Down
2 changes: 1 addition & 1 deletion examples/06_message_passing/message_passing_put.py
Original file line number Diff line number Diff line change
Expand Up @@ -157,7 +157,7 @@ def _worker(local_rank: int, world_size: int, init_url: str, args: dict):
)

# Main benchmark logic
shmem = iris.iris(args["heap_size"])
shmem = iris.iris(args["heap_size"], engines=("sdma",) if args["use_copy_engine"] else ())
dtype = torch_dtype_from_str(args["datatype"])
cur_rank = shmem.get_rank()
world_size = shmem.get_num_ranks()
Expand Down
22 changes: 21 additions & 1 deletion iris/device/sdma_utils.py
Original file line number Diff line number Diff line change
Expand Up @@ -11,9 +11,29 @@

import triton
import triton.language as tl
from xio import sdma_ep
from iris.mem.utils import wait_cnt

sdma_ep = None


def bind_sdma_ep():
"""Import rocm-xio and bind ``sdma_ep`` for the SDMA device functions."""
global sdma_ep
if sdma_ep is None:
try:
from xio import sdma_ep as imported
except ModuleNotFoundError as exc:
if exc.name != "xio":
raise
raise ModuleNotFoundError(
"rocm-xio is required for the SDMA copy engine. "
'Install it with `pip install iris[sdma]`, or drop "sdma" from engines '
"to use shader load/store only.",
name="xio",
) from exc
sdma_ep = imported
return sdma_ep
Comment thread
drprajap marked this conversation as resolved.


@triton.jit
def wrap_into_ring(index: tl.uint64):
Expand Down
85 changes: 66 additions & 19 deletions iris/host/iris.py
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,6 @@
count_devices,
)

from xio import sdma_ep
from iris.host.memory.symmetric_heap import SymmetricHeap
import numpy as np
from typing import Any
Expand All @@ -72,6 +71,8 @@
from iris.host.memory import tensors as tensor_creation
from iris.host.platform.utils import is_simulation_env

_SUPPORTED_ENGINES = frozenset({"sdma"})


class Iris:
"""
Expand All @@ -83,6 +84,9 @@ class Iris:
Args:
heap_size (int): Size of the symmetric heap in bytes. Default: 1GB (2^30)
allocator_type (str): Type of allocator to use. Options: "torch" (default), "vmem"
engines (str or tuple of str, optional): Engine name or names to set up. Use
``("sdma",)`` for host- and device-initiated SDMA operations (needs
``pip install iris[sdma]``). Defaults to ``()`` (none).

Example:
>>> ctx = iris.iris(heap_size=2**31) # 2GB heap with torch allocator
Expand All @@ -91,9 +95,25 @@ class Iris:

>>> # Use VMem allocator for memory oversubscription
>>> ctx = iris.iris(heap_size=2**31, allocator_type="vmem")

"""

def __init__(self, heap_size=1 << 30, allocator_type="torch"):
def __init__(self, heap_size=1 << 30, allocator_type="torch", engines=()):
if engines is None:
engines = ()
elif isinstance(engines, str):
engines = (engines,)
else:
engines = tuple(engines)
unsupported = set(engines) - _SUPPORTED_ENGINES
if unsupported:
raise ValueError(f"Unsupported engines: {sorted(unsupported)}. Supported: {sorted(_SUPPORTED_ENGINES)}.")
self._sdma_ep = None
if "sdma" in engines:
from iris.device.sdma_utils import bind_sdma_ep

self._sdma_ep = bind_sdma_ep()

# Initialize distributed environment
comm, cur_rank, num_ranks = init_distributed()
num_gpus = count_devices()
Expand Down Expand Up @@ -137,7 +157,27 @@ def __init__(self, heap_size=1 << 30, allocator_type="torch"):

distributed_barrier()

# initialize copy engines
self.copy_engines_device_ctx = None
if self._sdma_ep is not None:
self._init_copy_engines(num_gpus, num_ranks, cur_rank)

# Initialize CCL interface
self.ccl = self.CCL(self)

# Lazy initialization for ops interface
self._ops = None

# Device-side barrier state, keyed by process group (None = all ranks).
self._device_barrier_state: dict[Any, torch.Tensor] = {}

# Initialize tracing
self.tracing = Tracing(self)

# Pre-build the device context tensor (rebuilt when tracing is enabled)
self._build_device_context()

def _init_copy_engines(self, num_gpus, num_ranks, cur_rank):
sdma_ep = self._sdma_ep
sdma_ep.init()

context_size = sdma_ep.QUEUE_DEVICE_CTX_SIZE
Expand Down Expand Up @@ -167,20 +207,13 @@ def __init__(self, heap_size=1 << 30, allocator_type="torch"):
self.copy_engines_device_ctx[local_rank][3] = handle.doorbell
self.copy_engines_device_ctx[local_rank][4] = handle.cached_wptr
self.copy_engines_device_ctx[local_rank][5] = handle.committed_wptr
# Initialize CCL interface
self.ccl = self.CCL(self)

# Lazy initialization for ops interface
self._ops = None

# Device-side barrier state, keyed by process group (None = all ranks).
self._device_barrier_state: dict[Any, torch.Tensor] = {}

# Initialize tracing
self.tracing = Tracing(self)

# Pre-build the device context tensor (rebuilt when tracing is enabled)
self._build_device_context()
def _require_copy_engine(self):
if self._sdma_ep is None:
raise RuntimeError(
'SDMA is disabled; initialize Iris with engines=("sdma",) (needs `pip install iris[sdma]`).'
)
return self._sdma_ep

def __del__(self):
"""Cleanup resources on deletion."""
Expand Down Expand Up @@ -950,7 +983,8 @@ def get_heap_bases(self):
"""
return self.heap_bases

def get_copy_engine_ctx(self):
def get_copy_engine_ctx(self) -> torch.Tensor | None:
"""Return the device copy-engine context, or None when copy engines are disabled."""
return self.copy_engines_device_ctx

@staticmethod
Expand Down Expand Up @@ -1003,6 +1037,8 @@ def put(
"""
One-sided put operation with optional wait (POLL) and signal (ATOMIC).

Requires initialization with ``engines=("sdma",)``.

Supports:
- Simple copy: put(src, to_rank)
- Copy + signal: put(src, to_rank, signal_flag=flag)
Expand Down Expand Up @@ -1035,6 +1071,7 @@ def put(
... wait_flag=batch_ready, wait_value=256,
... signal_flag=transfer_done, signal_value=1)
"""
sdma_ep = self._require_copy_engine()
if to_tensor is None:
to_tensor = from_tensor

Expand Down Expand Up @@ -1094,6 +1131,8 @@ def put_tile(
"""
2D tile transfer with optional wait/signal (sub-window copy).

Requires initialization with ``engines=("sdma",)``.

Low-level API - caller provides pre-translated pointers for performance.

Args:
Expand Down Expand Up @@ -1125,6 +1164,7 @@ def put_tile(
>>> shmem.put_tile(tile, to_rank=1, to_ptr=dst_ptr, to_stride=dst_stride,
... wait_flag=wait_ptr, wait_value=256, signal_flag=signal_ptr)
"""
sdma_ep = self._require_copy_engine()
from_rank = self.get_rank()

wait_ptr, wait_bits = self._flag_pointer_and_bits(wait_flag, default_bits=32)
Expand Down Expand Up @@ -1176,6 +1216,8 @@ def put_tiles(
"""
Batched 2D tile transfer with optional shared wait/signal.

Requires initialization with ``engines=("sdma",)``.

Args:
tiles: Sequence of pre-configured sdma_ep.Tile objects
to_rank: Destination rank
Expand All @@ -1188,6 +1230,7 @@ def put_tiles(
async_op: If True, don't wait for completion
channel: SDMA channel to use
"""
sdma_ep = self._require_copy_engine()
from_rank = self.get_rank()

if len(tiles) != len(to_ptrs) or len(tiles) != len(to_strides):
Expand Down Expand Up @@ -1232,6 +1275,8 @@ def quiet(self, to_rank: int = None, channel: int = 0):
"""
Wait for all outstanding SDMA operations to complete.

Requires initialization with ``engines=("sdma",)``.

Args:
to_rank: If specified, wait only for ops to this rank.
If None, wait for ops to all ranks.
Expand All @@ -1242,6 +1287,7 @@ def quiet(self, to_rank: int = None, channel: int = 0):
>>> shmem.quiet(to_rank=1) # Wait for completion
>>> shmem.quiet() # Wait for all ranks
"""
sdma_ep = self._require_copy_engine()
src_rank = self.get_rank()
if to_rank is not None:
sdma_ep.quiet(src_rank, to_rank, channel)
Expand Down Expand Up @@ -1698,14 +1744,15 @@ def reduce_scatter(self, output_tensor, input_tensor, op=None, group=None, async
)


def iris(heap_size=1 << 30, allocator_type="torch"):
def iris(heap_size=1 << 30, allocator_type="torch", engines=()):
"""
Create and return an Iris instance with the specified heap size.

Args:
heap_size (int): Size of the heap in bytes. Defaults to 1GB.
allocator_type (str): Type of allocator to use. Options: "torch" (default), "vmem".
Can be overridden with IRIS_ALLOCATOR environment variable.
engines (str or tuple of str, optional): See :class:`Iris`.

Returns:
Iris: An initialized Iris instance.
Expand All @@ -1719,4 +1766,4 @@ def iris(heap_size=1 << 30, allocator_type="torch"):
>>> iris_ctx = iris.iris(2**30, allocator_type="vmem")
>>> tensor = iris_ctx.zeros(1024, 1024)
"""
return Iris(heap_size, allocator_type)
return Iris(heap_size, allocator_type, engines=engines)
Loading
Loading