Add API reference for iris.ccl, iris.ops, and iris.x modules

[Copilot]

Three core API modules (iris.ccl, iris.ops, iris.x) were undocumented in the reference section, making discovery difficult for users.

Changes

Collective Communication Library (iris.ccl)

• Host-side tensor-level collectives: all_reduce, all_gather, all_to_all, reduce_scatter
• Configuration via Config class with algorithm selection and tuning parameters
• ReduceOp enum for reduction operations

Fused GEMM+CCL Operations (iris.ops)

• Computation-communication overlap primitives: matmul_all_reduce, all_gather_matmul, matmul_all_gather, matmul_reduce_scatter
• Workspace management via FusedConfig and FusedWorkspace
• Accessible through shmem.ops namespace or standalone

Device-Side Tile-Level Primitives (iris.x)

• Fine-grained control in Triton kernels via TileView, TensorView, DeviceContext
• Five all-reduce algorithms: atomic, ring, two-shot, one-shot, spinlock
• Algorithm selection via AllReduceConfig
• Helper functions: tile_layout, tile_ptr, offset_ptr

Documentation Structure

docs/reference/
├── ccl/
│   ├── overview.md      # Usage patterns, async operations
│   ├── operations.md    # Autodoc for CCL methods
│   └── config.md        # Config and ReduceOp classes
├── ops/
│   ├── overview.md      # Fused operations, workspace management
│   ├── operations.md    # Autodoc for ops methods
│   └── config.md        # FusedConfig and FusedWorkspace
└── x/
    ├── overview.md      # Tile-level API, algorithm selection
    ├── core.md          # Core abstractions and helpers
    └── operations.md    # Device-side collectives

Each section includes usage examples and follows the existing Triton/Gluon API reference patterns.

Original prompt

---

This section details on the original issue you should resolve

<issue_title>[Documentation]: Add API reference for iris.ccl, iris.ops and iris.x</issue_title>
<issue_description>### Description of errors

• https://github.com/ROCm/iris/tree/main/docs/reference
• https://github.com/ROCm/iris/blob/main/docs/reference/api-reference.md

Attach any links, screenshots, or additional evidence you think will be helpful.

No response</issue_description>

Comments on the Issue (you are @copilot in this section)

• Fixes [Documentation]: Add API reference for iris.ccl, iris.ops and iris.x #344

---

✨ Let Copilot coding agent set things up for you — coding agent works faster and does higher quality work when set up for your repo.

[Copilot]

Pull request overview

This pull request adds comprehensive API reference documentation for three core Iris modules (iris.ccl, iris.ops, and iris.x) that were previously undocumented. The documentation follows the established pattern from existing Triton and Gluon reference sections.

Changes:

• Added structured documentation for collective communication operations (CCL), fused GEMM+CCL operations (ops), and device-side tile-level primitives (x)
• Created overview pages with usage patterns and examples for each module
• Added autodoc-based API reference pages for operations, configuration classes, and core abstractions
• Updated the main API reference navigation to include links to the new sections

Reviewed changes

Copilot reviewed 10 out of 10 changed files in this pull request and generated 6 comments.

Show a summary per file

File	Description
docs/reference/api-reference.md	Added navigation entries for the three new modules in the main API reference index
docs/reference/ccl/overview.md	Overview of collective communication library with usage patterns and examples
docs/reference/ccl/operations.md	Autodoc references for CCL collective operations (all_reduce, all_gather, all_to_all, reduce_scatter)
docs/reference/ccl/config.md	Autodoc references for CCL configuration classes (Config, ReduceOp)
docs/reference/ops/overview.md	Overview of fused GEMM+CCL operations with usage patterns and workspace management examples
docs/reference/ops/operations.md	Autodoc references for fused operations (matmul_all_reduce, all_gather_matmul, matmul_all_gather, matmul_reduce_scatter) and OpsNamespace
docs/reference/ops/config.md	Autodoc references for ops configuration classes (FusedConfig, FusedWorkspace)
docs/reference/x/overview.md	Overview of device-side tile-level primitives with algorithm selection and usage patterns
docs/reference/x/core.md	Autodoc references for core abstractions (TileView, TensorView, DeviceContext, AllReduceConfig) and helper functions
docs/reference/x/operations.md	Autodoc references for device-side collective operations (all-reduce variants, all_gather, all_to_all, reduce_scatter, gather)

[Copilot]

The new documentation sections for iris.ccl, iris.ops, and iris.x are not included in the Sphinx table of contents configuration file (docs/sphinx/_toc.yml). This means these pages will not be included in the generated documentation navigation and may not be properly built.

You need to add entries to docs/sphinx/_toc.yml under the "Reference" section, similar to how triton and gluon are structured. The entries should follow this pattern after the gluon section.

[Copilot]

The documentation shows DeviceContext methods being called (e.g., ctx.all_reduce(tile, src_view, dst_view)), but the actual API uses standalone functions where ctx is passed as the last parameter (e.g., iris.x.all_reduce_atomic(tile, src_view, dst_view, ctx)).

Based on the actual implementation in iris/x/core.py lines 466-489, DeviceContext does not have collective methods. The correct usage is to call the standalone functions from iris.x module, passing ctx as a parameter.

The examples throughout this file (lines 30, 55-58, 90-92) need to be corrected to match the actual API.

[Copilot]

The AllReduceConfig initialization examples show string arguments like AllReduceConfig("ring"), but the actual implementation in iris/x/core.py (lines 452-462) expects an integer variant_code (0-4) and a locks_ptr parameter.

The correct usage should be:

• AllReduceConfig(1, locks_ptr) for ring (not AllReduceConfig("ring"))
• AllReduceConfig(4, locks_ptr) for spinlock (not AllReduceConfig("spinlock", locks_ptr))

The mapping is documented in core.py lines 441-446:

• 0 = atomic
• 1 = ring
• 2 = one_shot
• 3 = two_shot
• 4 = spinlock

All examples using AllReduceConfig need to use integer codes, not strings.

Suggested change

	# Use ring algorithm
	config = iris.x.AllReduceConfig("ring")
	ctx.all_reduce(tile, src_view, dst_view, config=config)
	# Use spinlock algorithm with locks
	config = iris.x.AllReduceConfig("spinlock", locks_ptr)
	# Use ring algorithm (variant_code = 1)
	config = iris.x.AllReduceConfig(1, locks_ptr)
	ctx.all_reduce(tile, src_view, dst_view, config=config)
	# Use spinlock algorithm with locks (variant_code = 4)
	config = iris.x.AllReduceConfig(4, locks_ptr)

[Copilot]

The code example uses an undefined variable heap_size. It should either define it (e.g., heap_size = 2**30) or use a literal value like in the first example (line 14: iris.iris(heap_size=2**30)).

[Copilot]

The code example uses an undefined variable heap_size. It should either define it (e.g., heap_size = 2**30) or use a literal value like in the first example (line 14: iris.iris(heap_size=2**30)).

[Copilot]

The kernel signature is missing tl.constexpr annotations for M, N parameters. TensorView requires dimensions and strides to be constexpr (as documented in iris/x/core.py lines 264-265 and 269-272).

The parameters M and N should be annotated as:

def my_kernel(input_ptr, output_ptr, M: tl.constexpr, N: tl.constexpr,
              rank: tl.constexpr, world_size: tl.constexpr, heap_bases,
              BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr):

Additionally, stride_m and stride_n should either be passed as constexpr parameters or computed from constexpr values. See tests/x/test_all_to_all.py lines 21-27 for the correct pattern.

Suggested change

	def my_kernel(input_ptr, output_ptr, M, N,
	rank, world_size, heap_bases,
	def my_kernel(input_ptr, output_ptr, M: tl.constexpr, N: tl.constexpr,
	rank: tl.constexpr, world_size: tl.constexpr, heap_bases,