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

[Copilot]

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

• Explore the repository structure and existing documentation pattern
• Create API reference documentation for iris.ccl
  • Create overview.md for iris.ccl
  • Create operations.md for iris.ccl operations (all_reduce, all_gather, all_to_all, reduce_scatter)
  • Create config.md for iris.ccl configuration
• Create API reference documentation for iris.ops
  • Create overview.md for iris.ops
  • Create operations.md for iris.ops operations (matmul_all_reduce, all_gather_matmul, matmul_all_gather, matmul_reduce_scatter)
  • Create config.md for iris.ops configuration
• Create API reference documentation for iris.x
  • Create overview.md for iris.x
  • Create core.md for iris.x core abstractions (TileView, TensorView, DeviceContext, etc.)
  • Create operations.md for iris.x device-side operations
• Update docs/reference/api-reference.md to include links to new documentation
• Update docs/sphinx/_toc.yml to include new documentation in table of contents
• Test documentation builds successfully
• Address review comments
  • Fix iris.x examples to use correct API
  • Update AllReduceConfig to use integer codes
  • Add missing tl.constexpr annotations
  • Fix undefined heap_size variable
  • Update _toc.yml (restored after merge)
  • Use 'ctx' instead of 'shmem' in ccl documentation

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)

[mawad-amd]

@copilot apply changes based on the comments in this thread

[Copilot]

Pull request overview

Copilot reviewed 10 out of 10 changed files in this pull request and generated 1 comment.

[Copilot]

Pull request overview

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

[Copilot]

The section header labels iris.x.Tile as “deprecated”, but the implementation doesn’t mark it deprecated and several iris.x collectives (e.g., all_reduce_atomic, all_gather, reduce_scatter) require a Tile because they read tile.data. Please remove or reword the “deprecated” claim, and clarify when to use TileView vs Tile (TileView for coordinates/shape-only, Tile when carrying computed tile data).

[Copilot]

In this example, tile is a TileView, but iris.x.all_reduce_atomic expects a Tile with a populated tile.data (the local tile’s computed values). As written, the example can’t work because TileView has no data field. Update the snippet to compute a tile result (e.g., load/compute) and construct iris.x.Tile(pid_m, pid_n, BLOCK_M, BLOCK_N, data) before calling all_reduce_atomic (or use an op that accepts TileView).

Suggested change

	tile = iris.x.TileView(pid_m, pid_n, BLOCK_M, BLOCK_N)
	src_view = iris.x.make_tensor_view(input_ptr, M, N, stride_m, stride_n)
	dst_view = iris.x.make_tensor_view(output_ptr, M, N, stride_m, stride_n)
	ctx = iris.DeviceContext.initialize(context_tensor, rank, world_size)
	src_view = iris.x.make_tensor_view(input_ptr, M, N, stride_m, stride_n)
	dst_view = iris.x.make_tensor_view(output_ptr, M, N, stride_m, stride_n)
	ctx = iris.DeviceContext.initialize(context_tensor, rank, world_size)
	# Compute the local tile's data (e.g., load from input_ptr)
	offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)[:, None]
	offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)[None, :]
	ptrs = input_ptr + offs_m * stride_m + offs_n * stride_n
	data = tl.load(ptrs)
	# Construct a Tile with populated data for the collective
	tile = iris.x.Tile(pid_m, pid_n, BLOCK_M, BLOCK_N, data)

[Copilot]

This usage-pattern snippet calls iris.x.all_to_all(tile, dst_view, dst_view, N_per_rank, ctx) but N_per_rank is not defined/passed in the kernel signature, and all_to_all requires a src_view argument (signature: all_to_all(tile, src_view, dst_view, N_per_rank, ctx)). The example should define N_per_rank: tl.constexpr in the kernel params and create src_view (likely from an input pointer) rather than passing dst_view twice.

Suggested change

	BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr):
	pid_m = tl.program_id(0)
	pid_n = tl.program_id(1)
	tile = iris.x.TileView(pid_m, pid_n, BLOCK_M, BLOCK_N)
	dst_view = iris.x.make_tensor_view(output_ptr, M, N, stride_m, stride_n)
	ctx = iris.DeviceContext.initialize(context_tensor, rank, world_size)
	# Call collectives directly
	iris.x.all_reduce_atomic(tile, dst_view, ctx)
	iris.x.all_gather(tile, dst_view, dim=0, ctx=ctx)
	iris.x.all_to_all(tile, dst_view, dst_view, N_per_rank, ctx)
	N_per_rank: tl.constexpr,
	BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr):
	pid_m = tl.program_id(0)
	pid_n = tl.program_id(1)
	tile = iris.x.TileView(pid_m, pid_n, BLOCK_M, BLOCK_N)
	src_view = iris.x.make_tensor_view(input_ptr, M, N, stride_m, stride_n)
	dst_view = iris.x.make_tensor_view(output_ptr, M, N, stride_m, stride_n)
	ctx = iris.DeviceContext.initialize(context_tensor, rank, world_size)
	# Call collectives directly
	iris.x.all_reduce_atomic(tile, dst_view, ctx)
	iris.x.all_gather(tile, dst_view, dim=0, ctx=ctx)
	iris.x.all_to_all(tile, src_view, dst_view, N_per_rank, ctx)

[Copilot]

In this algorithm-selection example, config = iris.x.AllReduceConfig(1, dummy_locks) is created but never used, and the subsequent call iris.x.all_reduce_ring(tile, dst_view, ctx) doesn’t match the actual all_reduce_ring signature (it requires src_view and dst_view). Either remove AllReduceConfig from this example and show the correct call form for all_reduce_ring(tile, src_view, dst_view, ctx), or add a real API that consumes AllReduceConfig and document that instead.

[Copilot]

This snippet has multiple undefined/mismatched pieces: tile_id = pid_m * num_tiles_n + pid_n references num_tiles_n which isn’t defined in the example, and iris.x.all_reduce_spinlock takes the locks tensor directly (no AllReduceConfig parameter), so the config = iris.x.AllReduceConfig(4, locks_ptr) line is misleading unless there is an API that consumes it. Consider removing AllReduceConfig here, define num_tiles_n (e.g., via tl.cdiv(N, BLOCK_N) or dst_view helpers), and show the correct all_reduce_spinlock(tile, dst_view, locks_ptr, ctx) call with the required locks layout.