From 3926e491af3aefe9dce2e8f0e11f79b9c184bed2 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Sun, 11 Oct 2026 06:04:14 -0700 Subject: [PATCH 01/18] Mock xio so the API reference builds iris/host/iris.py imports xio unconditionally at module load. The docs CI installs only docs/sphinx/requirements.txt, so xio is never available and every autodoc directive failed with ModuleNotFoundError, leaving the entire API reference empty on the published site. Build goes from 71 warnings to 6, and the reference pages regain their content: host/class 15 signatures, triton/device-functions 13, gluon/device-functions 15, host/tensor-creation 13, host/ccl 5, triton/ops 4. --- .gitignore | 4 ++++ docs/conf.py | 5 +++++ 2 files changed, 9 insertions(+) diff --git a/.gitignore b/.gitignore index 67ebd97bd..4b7306ca7 100644 --- a/.gitignore +++ b/.gitignore @@ -69,3 +69,7 @@ docker/rocm-systems/ .intellikit .github/agents/skills/ docs/benchmark-results/*.png + +# Local Sphinx build +docs/.venv/ +docs/_build/ diff --git a/docs/conf.py b/docs/conf.py index ae01bbb61..098231f0e 100644 --- a/docs/conf.py +++ b/docs/conf.py @@ -96,6 +96,11 @@ "iris.host.distributed.helpers", "iris.host.platform.hip", "tritonblas", + # rocm-xio. iris/host/iris.py imports it unconditionally at module load, so + # without this every autodoc directive fails with ModuleNotFoundError and the + # whole API reference renders empty. The docs CI installs only + # docs/sphinx/requirements.txt, so xio is never present there either. + "xio", ] # Custom mocks that preserve docstrings for Triton Gluon From 03f8fbb6fa5b2da30781d565220349a3b28b6afe Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Sun, 11 Oct 2026 06:08:41 -0700 Subject: [PATCH 02/18] Add the SLURM guide to the navigation getting-started/slurm.md is a 290-line guide linked from index.md three times and from installation.md twice, but it was absent from the sidebar, so Sphinx reported it as belonging to no toctree. Note for anyone editing the nav: sphinx/_toc.yml is GENERATED from sphinx/_toc.yml.in during the build. Both are tracked, so editing _toc.yml appears to work and is silently overwritten on the next build. --- docs/conf.py | 1 + docs/sphinx/_toc.yml.in | 1 + 2 files changed, 2 insertions(+) diff --git a/docs/conf.py b/docs/conf.py index 098231f0e..6cc98ad38 100644 --- a/docs/conf.py +++ b/docs/conf.py @@ -169,6 +169,7 @@ class GluonMock: # Table of contents external_toc_path = "./sphinx/_toc.yml" + # Theme options for AMD ROCm theme html_theme_options = { "flavor": "instinct", diff --git a/docs/sphinx/_toc.yml.in b/docs/sphinx/_toc.yml.in index afe6c23cf..0b42543c9 100644 --- a/docs/sphinx/_toc.yml.in +++ b/docs/sphinx/_toc.yml.in @@ -3,6 +3,7 @@ subtrees: - caption: Getting Started entries: - file: getting-started/installation.md + - file: getting-started/slurm.md - file: reference/examples.md - file: CONTRIBUTING.md - caption: Conceptual From fd05a4a4ed242436d238b19761f8716c846c4bdb Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Sun, 11 Oct 2026 06:12:17 -0700 Subject: [PATCH 03/18] Document the 11 missing public APIs Nine public methods on Iris and two Triton ops had no documentation at all: host put, put_tile, put_tiles, quiet, get_copy_engine_ctx -> new page, reference/host/copy-engine.md host allocate_symmetric, device_barrier, get_device_id, ops -> added to reference/host/class.md triton atomic_sub, quiet -> added to reference/triton/ops.md The first five are the entire host-side SDMA surface, which was undocumented. Also: - get_copy_engine_ctx had no docstring; added one. - atomic_sub's docstring had a line indented 8 spaces inside a 4-space paragraph, which RST parses as a block quote and warns about. - Iris.ops is a property, so it needs autoproperty rather than automethod. Signature counts: class 15 -> 19, triton/ops 4 -> 6, copy-engine 5 new. --- docs/reference/api-reference.md | 1 + docs/reference/host/class.md | 24 ++++++++++++++++ docs/reference/host/copy-engine.md | 46 ++++++++++++++++++++++++++++++ docs/reference/triton/ops.md | 12 ++++++++ docs/sphinx/_toc.yml.in | 1 + iris/host/iris.py | 12 ++++++++ iris/mem/triton/ops.py | 2 +- 7 files changed, 97 insertions(+), 1 deletion(-) create mode 100644 docs/reference/host/copy-engine.md diff --git a/docs/reference/api-reference.md b/docs/reference/api-reference.md index 38ab0ba40..60198630f 100644 --- a/docs/reference/api-reference.md +++ b/docs/reference/api-reference.md @@ -7,6 +7,7 @@ Explore Iris APIs. The host-side API (class, tensor creation, CCL) is identical - [Iris Class](host/class.md) - [Tensor Creation](host/tensor-creation.md) - [Collective Communication (CCL)](host/ccl.md) +- [Copy Engine (SDMA)](host/copy-engine.md) ## Triton Backend diff --git a/docs/reference/host/class.md b/docs/reference/host/class.md index 87a0187bb..1b0128371 100644 --- a/docs/reference/host/class.md +++ b/docs/reference/host/class.md @@ -45,3 +45,27 @@ Broadcast data from a source rank to all ranks. This method automatically detect ```{eval-rst} .. automethod:: iris.host.iris.Iris.broadcast ``` + +## allocate_symmetric + +```{eval-rst} +.. automethod:: iris.host.iris.Iris.allocate_symmetric +``` + +## device_barrier + +```{eval-rst} +.. automethod:: iris.host.iris.Iris.device_barrier +``` + +## get_device_id + +```{eval-rst} +.. automethod:: iris.host.iris.Iris.get_device_id +``` + +## ops + +```{eval-rst} +.. autoproperty:: iris.host.iris.Iris.ops +``` diff --git a/docs/reference/host/copy-engine.md b/docs/reference/host/copy-engine.md new file mode 100644 index 000000000..2b50bffd2 --- /dev/null +++ b/docs/reference/host/copy-engine.md @@ -0,0 +1,46 @@ +# Copy Engine (SDMA) + +Host-initiated transfers that run on the GPU's SDMA copy engines rather than on +compute units. The host enqueues a descriptor and the DMA engine moves the +bytes, so no kernel launch is involved and no CUs are consumed. + +```{note} +These operations require the copy engine to be initialised. See +[`iris.iris()`](class.md) for how to enable it, and the +[Installation guide](../../getting-started/installation.md) for the optional +`rocm-xio` dependency they rely on. +``` + +Unlike the device-side operations, there is no engine choice to make here: a CU +store is an instruction a wave executes, so the host cannot issue one. Anything +host-initiated goes through the copy engine. + +## put + +```{eval-rst} +.. automethod:: iris.host.iris.Iris.put +``` + +## put_tile + +```{eval-rst} +.. automethod:: iris.host.iris.Iris.put_tile +``` + +## put_tiles + +```{eval-rst} +.. automethod:: iris.host.iris.Iris.put_tiles +``` + +## quiet + +```{eval-rst} +.. automethod:: iris.host.iris.Iris.quiet +``` + +## get_copy_engine_ctx + +```{eval-rst} +.. automethod:: iris.host.iris.Iris.get_copy_engine_ctx +``` diff --git a/docs/reference/triton/ops.md b/docs/reference/triton/ops.md index 93be2b31d..529c3fa38 100644 --- a/docs/reference/triton/ops.md +++ b/docs/reference/triton/ops.md @@ -21,3 +21,15 @@ Fused matrix multiplication and collective communication operations accessible v ```{eval-rst} .. automethod:: iris.ops.OpsNamespace.matmul_reduce_scatter ``` + +## atomic_sub + +```{eval-rst} +.. autofunction:: iris.mem.triton.ops.atomic_sub +``` + +## quiet + +```{eval-rst} +.. autofunction:: iris.mem.triton.ops.quiet +``` diff --git a/docs/sphinx/_toc.yml.in b/docs/sphinx/_toc.yml.in index 0b42543c9..d4fec22fa 100644 --- a/docs/sphinx/_toc.yml.in +++ b/docs/sphinx/_toc.yml.in @@ -19,6 +19,7 @@ subtrees: - file: reference/host/class.md - file: reference/host/tensor-creation.md - file: reference/host/ccl.md + - file: reference/host/copy-engine.md - file: reference/triton/overview.md entries: - file: reference/triton/device-functions.md diff --git a/iris/host/iris.py b/iris/host/iris.py index aa1674334..85ac6fcd7 100644 --- a/iris/host/iris.py +++ b/iris/host/iris.py @@ -951,6 +951,18 @@ def get_heap_bases(self): return self.heap_bases def get_copy_engine_ctx(self): + """ + Get the device-side copy engine context for SDMA operations. + + The returned tensor holds the per-rank SDMA queue, read and write pointer, + and doorbell addresses that device-side kernels need in order to enqueue + copy packets. Pass it to ``iris.put`` as ``copy_engine_ctx`` when driving + the copy engine from inside a kernel. + + Returns: + torch.Tensor: The copy engine device context, or ``None`` if the copy + engine was not initialized. + """ return self.copy_engines_device_ctx @staticmethod diff --git a/iris/mem/triton/ops.py b/iris/mem/triton/ops.py index 1d022efac..f50fa6e9d 100644 --- a/iris/mem/triton/ops.py +++ b/iris/mem/triton/ops.py @@ -603,7 +603,7 @@ def atomic_sub( This function performs an atomic subtraction operation by translating the pointer from the `from_rank`'s address space to the `to_rank`'s address space and atomically - subtracting the provided data from the `to_rank` memory location. If the `from_rank` and `to_rank` are the same, + subtracting the provided data from the `to_rank` memory location. If the `from_rank` and `to_rank` are the same, this function performs a local atomic subtraction operation. Args: From 750d680c6617803a7e7c158be87cd8fd1c9792dc Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Sun, 11 Oct 2026 06:17:37 -0700 Subject: [PATCH 04/18] Fix 18 broken images and restore the theme's JavaScript Images. Several conceptual pages reference images from raw HTML blocks, which is how the light/dark variants are selected: Sphinx only tracks images referenced through markdown or image directives, so it copied 4 of the 29 and every raw-HTML reference 404'd on the published site: 11 on conceptual/taxonomy, 6 on conceptual/programming-model, plus the logo on index. A build-finished hook now copies docs/images/ verbatim, which fixes all of them without rewriting the theme-switching markup. JavaScript. Sphinx removed its bundled jQuery in 6.0 and requirements.txt asks only for sphinx>=5.0.0, so CI resolves 9.x. Four rocm-docs-core theme scripts still call $() -- code_word_breaks, renameVersionLinks, rdcMisc and theme_mode_captions -- so every page threw '$ is not defined' and all four silently did nothing. Added sphinxcontrib-jquery. Verified with Playwright over HTTP across all 19 content pages: every page returns 200, renders its heading, nav, breadcrumbs and TOC, and loads all of its images (taxonomy 12/12 and programming-model 7/7, both previously 0). --- docs/conf.py | 32 ++++++++++++++++++++++++++++++++ docs/sphinx/requirements.txt | 2 ++ 2 files changed, 34 insertions(+) diff --git a/docs/conf.py b/docs/conf.py index 6cc98ad38..5339dc798 100644 --- a/docs/conf.py +++ b/docs/conf.py @@ -32,6 +32,11 @@ # ones. extensions = [ "rocm_docs", + # Sphinx removed its bundled jQuery in 6.0, but four rocm-docs-core theme + # scripts (code_word_breaks, renameVersionLinks, rdcMisc, + # theme_mode_captions) still call $(). Without this every page throws + # "$ is not defined" and those scripts silently do nothing. + "sphinxcontrib.jquery", "sphinx.ext.autodoc", "sphinx.ext.napoleon", "sphinx.ext.viewcode", @@ -191,3 +196,30 @@ class GluonMock: "copybutton_hide": copybutton_hide, "copybutton_remove_prompts": copybutton_remove_prompts, } + +# --------------------------------------------------------------------------- +# Copy docs/images/ verbatim into the built site. +# +# Several conceptual pages reference images from raw HTML blocks, e.g. +# +# which is how the light/dark theme variants are selected. Sphinx only tracks +# and copies images referenced through markdown or image directives, so every +# raw-HTML reference 404s on the published site -- 18 of them, across +# conceptual/programming-model, conceptual/taxonomy and index. +# +# Copying the directory wholesale fixes all of them without rewriting the +# theme-switching markup. +# --------------------------------------------------------------------------- +import shutil as _shutil + + +def _copy_images(app, exception): + if exception is not None or app.builder.name != "html": + return + src = os.path.join(os.path.dirname(__file__), "images") + if os.path.isdir(src): + _shutil.copytree(src, os.path.join(app.outdir, "images"), dirs_exist_ok=True) + + +def setup(app): + app.connect("build-finished", _copy_images) diff --git a/docs/sphinx/requirements.txt b/docs/sphinx/requirements.txt index cdd7c0c73..4633e2640 100644 --- a/docs/sphinx/requirements.txt +++ b/docs/sphinx/requirements.txt @@ -4,3 +4,5 @@ sphinx-copybutton>=0.5.0 sphinx-autobuild>=2021.3.14 sphinx_external_toc>=1.0.0 rocm-docs-core +# rocm-docs-core's theme JS uses jQuery, which Sphinx stopped bundling in 6.0. +sphinxcontrib-jquery>=4.1 From 5aa74cade7b248ec01e0f6185cbd0da83cb34fc7 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 12 Oct 2026 13:26:00 -0700 Subject: [PATCH 05/18] Let the docs workflow be triggered by hand This fork exists to preview the docs, and nothing here generates the pushes that the workflow was written to react to. It fires on push to main under a path filter, so once the docs commits had landed there was no remaining event that could rebuild the site -- and the deploy job additionally gated on github.event_name == 'push', so even a manual run would have built an artifact and stopped short of publishing. Adds workflow_dispatch and widens the deploy gate to accept it. The branch condition is unchanged: still main only, so a PR build still cannot publish. Co-Authored-By: Claude --- .github/workflows/docs.yml | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/.github/workflows/docs.yml b/.github/workflows/docs.yml index 902ee6a2e..3fc96f0f4 100644 --- a/.github/workflows/docs.yml +++ b/.github/workflows/docs.yml @@ -1,6 +1,10 @@ name: Build and Deploy Documentation on: + # Rebuild and republish from the Actions tab without needing a commit. This + # fork exists only to preview the docs, so otherwise the sole way to refresh + # the site is to push something to main. + workflow_dispatch: push: branches: [ main ] paths: @@ -74,7 +78,7 @@ jobs: url: ${{ steps.deployment.outputs.page_url }} runs-on: ubuntu-latest needs: build - if: github.ref == 'refs/heads/main' && github.event_name == 'push' + if: github.ref == 'refs/heads/main' && (github.event_name == 'push' || github.event_name == 'workflow_dispatch') permissions: contents: read pages: write From cac054b55955d486565ba3818c62d1ebd65d7ab0 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 12 Oct 2026 13:44:56 -0700 Subject: [PATCH 06/18] Document the Triton DeviceContext iris.mem.triton.context.Context -- exported as iris.DeviceContext and the aggregate you construct inside a @triton.jit kernel -- had no autodoc coverage at all. Twenty-four public methods, zero directives anywhere under docs/, and nothing in the live API index. Nine of them existed nowhere else on the site: all_reduce_atomic all_reduce_spinlock all_reduce_one_shot all_reduce_ring all_reduce_two_shot all_gather gather all_to_all reduce_scatter These are the tile-level collectives -- they run inside the kernel on the tile a program already holds, which is a different surface from the host collectives on reference/host/ccl.md that launch their own kernels. Searching the docs for "all_gather" found only the host one, so the in-kernel form read as missing. The remaining fifteen were reachable only as the module-level free functions on reference/triton/device-functions.md, which take heap_bases and both ranks explicitly. That is the layer underneath; the bound ctx.load(pointer, from_rank) form a reader actually writes was not on the site. The contrast made it stark: Gluon's device context has all fifteen of its methods documented, while Triton -- the full surface, of which Gluon is a subset -- had none. Verified by building on Python 3.13: 24 headings, 24 rendered signatures, no empty sections, and zero autodoc import failures across the whole build. Note the build needs >=3.10; on 3.9 every autodoc import fails at iris/drivers/base.py:49, which uses PEP 604 union syntax at module level. Co-Authored-By: Claude --- docs/reference/triton/device-context.md | 187 ++++++++++++++++++++++++ docs/sphinx/_toc.yml.in | 1 + 2 files changed, 188 insertions(+) create mode 100644 docs/reference/triton/device-context.md diff --git a/docs/reference/triton/device-context.md b/docs/reference/triton/device-context.md new file mode 100644 index 000000000..a788f0b2b --- /dev/null +++ b/docs/reference/triton/device-context.md @@ -0,0 +1,187 @@ +# Device Context + +Device-side methods provided by Iris for remote memory operations, atomics and +collectives. These methods are part of the `DeviceContext` aggregate used inside +`@triton.jit` kernels, and are reached through the context tensor passed in from +the host: + +```python +from iris import DeviceContext + +@triton.jit +def producer(context_tensor, target_buffer, ..., rank: tl.constexpr, world_size: tl.constexpr): + ctx = DeviceContext.initialize(context_tensor, rank, world_size) + ... + ctx.store(target_buffer + offsets, values, to_rank=consumer_rank, mask=mask) + ctx.atomic_cas(flag + pid, 0, 1, to_rank=consumer_rank, sem="release", scope="sys") +``` + +Taken from `examples/06_message_passing/message_passing_device_context.py`. + +```{note} +The collectives below are **tile-level** -- they run inside your kernel, on the +tile a program is already holding, and fuse with whatever compute surrounds +them. They are a different surface from the host collectives in +[Collective Communication Operations](../host/ccl.md), which launch their own +kernels and are called from Python. +``` + +```{note} +Every operation here also exists as a free function -- see +[Device Functions](device-functions.md). The free form is explicit about both +ends of the transfer and the heap: `load(pointer, to_rank, from_rank, heap_bases)`. +The bound form below knows the current rank and carries the heap bases itself, so +the same call is `ctx.load(pointer, from_rank)`. +``` + +## Initialization + +### initialize +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.initialize + :noindex: +``` + +## Memory transfer operations + +### load +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.load + :noindex: +``` + +### store +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.store + :noindex: +``` + +### get +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.get + :noindex: +``` + +### put +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.put + :noindex: +``` + +### copy +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.copy + :noindex: +``` + +## Atomic operations + +### atomic_add +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.atomic_add + :noindex: +``` + +### atomic_sub +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.atomic_sub + :noindex: +``` + +### atomic_cas +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.atomic_cas + :noindex: +``` + +### atomic_xchg +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.atomic_xchg + :noindex: +``` + +### atomic_xor +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.atomic_xor + :noindex: +``` + +### atomic_and +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.atomic_and + :noindex: +``` + +### atomic_or +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.atomic_or + :noindex: +``` + +### atomic_min +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.atomic_min + :noindex: +``` + +### atomic_max +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.atomic_max + :noindex: +``` + +## Collective operations + +### all_reduce_atomic +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.all_reduce_atomic + :noindex: +``` + +### all_reduce_spinlock +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.all_reduce_spinlock + :noindex: +``` + +### all_reduce_one_shot +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.all_reduce_one_shot + :noindex: +``` + +### all_reduce_ring +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.all_reduce_ring + :noindex: +``` + +### all_reduce_two_shot +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.all_reduce_two_shot + :noindex: +``` + +### all_gather +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.all_gather + :noindex: +``` + +### gather +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.gather + :noindex: +``` + +### all_to_all +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.all_to_all + :noindex: +``` + +### reduce_scatter +```{eval-rst} +.. automethod:: iris.mem.triton.context.Context.reduce_scatter + :noindex: +``` diff --git a/docs/sphinx/_toc.yml.in b/docs/sphinx/_toc.yml.in index d4fec22fa..2fc6fbac9 100644 --- a/docs/sphinx/_toc.yml.in +++ b/docs/sphinx/_toc.yml.in @@ -22,6 +22,7 @@ subtrees: - file: reference/host/copy-engine.md - file: reference/triton/overview.md entries: + - file: reference/triton/device-context.md - file: reference/triton/device-functions.md - file: reference/triton/ops.md - file: reference/gluon/overview.md From 088502f3e8e095137417644c2958bb0c7610839b Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 12 Oct 2026 14:07:55 -0700 Subject: [PATCH 07/18] Document the tile types, and fix two rendering faults Tile types. iris/mem/triton/types.py exports Tile, TileView, TensorView, AllReduceConfig and make_tensor_view from the top-level package, and none of them were documented. These are what the DeviceContext collectives operate on, so the new device-context page referred to objects the site never defined. Adds reference/triton/types.md -- 43 rendered signatures. MagicMock in signatures. tl.constexpr and tl.tensor are used as parameter annotations, so autodoc resolved them to mock attributes and printed reprs like "" into the rendered parameter list. Twenty such leaks before this change, and the embedded id changes every build, so the HTML churned too. The mocked modules now carry real classes for those names. Gluon needed the same treatment separately: it annotates with gl.constexpr, which resolves through triton.experimental.gluon .language rather than triton.language. Only the names actually used as annotations are given classes. Adding uint and pointer_type as well tripped a pre-existing circular import between iris/__init__.py and iris.mem.triton.ops, which surfaced as 19 autodoc failures -- neither name appears in a rendered signature, so neither is needed. Collapsed bullet lists. Four docstrings put a bullet list directly beneath a line ending in a colon. RST needs a blank line there, so the dashes folded into the paragraph and rendered as "has the format: - [cur_rank, num_ranks, ...] - If tracing=True ...". Fixed in Iris.broadcast, Iris.put, TileView and Context.initialize, the last as a literal block since it is a tensor layout and smartquotes was also turning its "..." into an ellipsis. Fifty-nine sites match the colon-then-dash pattern repo-wide, but only these four actually render wrong; the rest sit inside Args: blocks, which napoleon handles correctly. Also generalises the "Fused GEMM + CCL Operations" heading to "Fused Communication and Compute Operations" -- GEMM is the compute side today, not the definition of the surface. Verified on Python 3.13 across all nine API pages: 143 signatures rendered, zero mock reprs, zero collapsed lists, zero autodoc import failures. Co-Authored-By: Claude --- docs/conf.py | 28 ++++++++- docs/reference/api-reference.md | 2 +- docs/reference/triton/ops.md | 8 ++- docs/reference/triton/overview.md | 2 +- docs/reference/triton/types.md | 99 +++++++++++++++++++++++++++++++ docs/sphinx/_toc.yml.in | 1 + iris/host/iris.py | 2 + iris/mem/triton/context.py | 8 ++- iris/mem/triton/types.py | 1 + 9 files changed, 143 insertions(+), 8 deletions(-) create mode 100644 docs/reference/triton/types.md diff --git a/docs/conf.py b/docs/conf.py index 5339dc798..ec703b641 100644 --- a/docs/conf.py +++ b/docs/conf.py @@ -125,6 +125,26 @@ def __call__(self, func=None, **kwargs): # Mock triton.language first triton_language_mock = MagicMock() + + +# These tl.* names are used as parameter ANNOTATIONS throughout the device APIs, +# so autodoc resolves them and prints their repr into the rendered signature. A +# bare MagicMock attribute reprs as +# +# which leaks into the parameter list, and the id changes on every build so the +# HTML churns too. Giving each one a real class makes it repr as its own name. +# The list is every `: tl.` annotation appearing in iris/ -- if a new one +# is introduced it will show up as a MagicMock on the rendered page. +def _name_mock_types(mock, module, names): + """Give a mocked module real classes for the names used as annotations.""" + for n in names: + cls = type(n, (), {"__doc__": "Triton %s." % n}) + cls.__module__ = module + setattr(mock, n, cls) + + +_name_mock_types(triton_language_mock, "triton.language", ("constexpr", "tensor")) + sys.modules["triton.language"] = triton_language_mock sys.modules["triton.language.core"] = MagicMock() sys.modules["triton.language.core"]._aggregate = lambda cls: cls # Preserve class @@ -152,7 +172,13 @@ class GluonMock: sys.modules["triton.experimental"] = MagicMock() sys.modules["triton.experimental"].gluon = GluonMock() sys.modules["triton.experimental.gluon"] = GluonMock() -sys.modules["triton.experimental.gluon"].language = MagicMock() +_gluon_language_mock = MagicMock() +# Gluon annotates with gl.constexpr, which resolves through a *different* mock +# module than tl.constexpr -- so it needs the same treatment or it reprs as a +# MagicMock in the rendered signature. +_name_mock_types(_gluon_language_mock, "triton.experimental.gluon.language", ("constexpr", "tensor")) +sys.modules["triton.experimental.gluon"].language = _gluon_language_mock +sys.modules["triton.experimental.gluon.language"] = _gluon_language_mock # Napoleon settings for Google/NumPy docstring parsing napoleon_google_docstring = True diff --git a/docs/reference/api-reference.md b/docs/reference/api-reference.md index 60198630f..54d6d4d76 100644 --- a/docs/reference/api-reference.md +++ b/docs/reference/api-reference.md @@ -13,7 +13,7 @@ Explore Iris APIs. The host-side API (class, tensor creation, CCL) is identical - [Overview](triton/overview.md) - [Device Functions](triton/device-functions.md) -- [Fused GEMM + CCL Operations](triton/ops.md) +- [Fused Communication and Compute Operations](triton/ops.md) ## Gluon Backend (Experimental) diff --git a/docs/reference/triton/ops.md b/docs/reference/triton/ops.md index 529c3fa38..16026e907 100644 --- a/docs/reference/triton/ops.md +++ b/docs/reference/triton/ops.md @@ -1,6 +1,10 @@ -# Fused GEMM + CCL Operations +# Fused Communication and Compute Operations -Fused matrix multiplication and collective communication operations accessible via the `ops` property on the `Iris` instance (e.g. `ctx.ops.matmul_all_reduce(...)`). +Operations that fuse communication with compute in a single kernel, accessible +via the `ops` property on the `Iris` instance (e.g. `ctx.ops.matmul_all_reduce(...)`). +The compute side is currently GEMM, supplied by tritonBLAS; the communication side +is an Iris collective. Fusing them means the collective overlaps the matmul instead +of running as a separate kernel after it. ## matmul_all_reduce ```{eval-rst} diff --git a/docs/reference/triton/overview.md b/docs/reference/triton/overview.md index 46f12f986..d82bcb510 100644 --- a/docs/reference/triton/overview.md +++ b/docs/reference/triton/overview.md @@ -36,5 +36,5 @@ Explore the API by section: - [Tensor Creation](../host/tensor-creation.md) - [Device Functions](device-functions.md) - [Collective Communication (CCL)](../host/ccl.md) -- [Fused GEMM + CCL Operations](ops.md) +- [Fused Communication and Compute Operations](ops.md) diff --git a/docs/reference/triton/types.md b/docs/reference/triton/types.md new file mode 100644 index 000000000..ccfe01303 --- /dev/null +++ b/docs/reference/triton/types.md @@ -0,0 +1,99 @@ +# Tile Types + +The tile abstractions the device APIs operate on. All of these are exported from +the top level, so `iris.Tile`, `iris.TensorView` and so on are the import paths. + +A `TensorView` wraps a pointer plus the tensor's shape and strides, and hands out +tile pointers and masks. A `TileView` carries a tile's coordinates; a `Tile` +carries its coordinates *and* the data. `AllReduceConfig` is the device-side +configuration consumed by the all-reduce methods on +[Device Context](device-context.md). + +## TensorView + +```{eval-rst} +.. autoclass:: iris.mem.triton.types.TensorView + :noindex: +``` + +### TensorView.tile_ptr +```{eval-rst} +.. automethod:: iris.mem.triton.types.TensorView.tile_ptr + :noindex: +``` + +### TensorView.tile_ptr_from_indices +```{eval-rst} +.. automethod:: iris.mem.triton.types.TensorView.tile_ptr_from_indices + :noindex: +``` + +### TensorView.offset_tile_ptr +```{eval-rst} +.. automethod:: iris.mem.triton.types.TensorView.offset_tile_ptr + :noindex: +``` + +### TensorView.offset +```{eval-rst} +.. automethod:: iris.mem.triton.types.TensorView.offset + :noindex: +``` + +### TensorView.number_of_tiles +```{eval-rst} +.. automethod:: iris.mem.triton.types.TensorView.number_of_tiles + :noindex: +``` + +## TileView + +```{eval-rst} +.. autoclass:: iris.mem.triton.types.TileView + :noindex: +``` + +### TileView.layout +```{eval-rst} +.. automethod:: iris.mem.triton.types.TileView.layout + :noindex: +``` + +### TileView.indices +```{eval-rst} +.. automethod:: iris.mem.triton.types.TileView.indices + :noindex: +``` + +## Tile + +```{eval-rst} +.. autoclass:: iris.mem.triton.types.Tile + :noindex: +``` + +### Tile.layout +```{eval-rst} +.. automethod:: iris.mem.triton.types.Tile.layout + :noindex: +``` + +### Tile.indices +```{eval-rst} +.. automethod:: iris.mem.triton.types.Tile.indices + :noindex: +``` + +## AllReduceConfig + +```{eval-rst} +.. autoclass:: iris.mem.triton.types.AllReduceConfig + :noindex: +``` + +## make_tensor_view + +```{eval-rst} +.. autofunction:: iris.mem.triton.types.make_tensor_view + :noindex: +``` diff --git a/docs/sphinx/_toc.yml.in b/docs/sphinx/_toc.yml.in index 2fc6fbac9..49ce51d06 100644 --- a/docs/sphinx/_toc.yml.in +++ b/docs/sphinx/_toc.yml.in @@ -24,6 +24,7 @@ subtrees: entries: - file: reference/triton/device-context.md - file: reference/triton/device-functions.md + - file: reference/triton/types.md - file: reference/triton/ops.md - file: reference/gluon/overview.md entries: diff --git a/iris/host/iris.py b/iris/host/iris.py index 85ac6fcd7..2f7318581 100644 --- a/iris/host/iris.py +++ b/iris/host/iris.py @@ -304,6 +304,7 @@ def broadcast(self, value, src=0): This method automatically detects the type of value and uses the appropriate broadcast mechanism: + - For tensors and arrays: uses efficient PyTorch distributed tensor collectives - For scalars and other objects: uses object broadcast @@ -1016,6 +1017,7 @@ def put( One-sided put operation with optional wait (POLL) and signal (ATOMIC). Supports: + - Simple copy: put(src, to_rank) - Copy + signal: put(src, to_rank, signal_flag=flag) - Wait + copy: put(src, to_rank, wait_flag=flag, wait_value=N) diff --git a/iris/mem/triton/context.py b/iris/mem/triton/context.py index cb6549313..63ee446e9 100644 --- a/iris/mem/triton/context.py +++ b/iris/mem/triton/context.py @@ -98,9 +98,11 @@ def initialize(context_tensor, rank, world_size, tracing: tl.constexpr = False): """ Initialize Context from the encoded context tensor. - The context tensor has the format: - - [cur_rank, num_ranks, heap_base_0, ..., heap_base_N, trace_info...] - - If tracing=True: extracts trace buffer pointers from context_tensor + The context tensor has the format:: + + [cur_rank, num_ranks, heap_base_0, ..., heap_base_N, trace_info...] + + If ``tracing=True``, trace buffer pointers are also extracted from it. Args: context_tensor: Pointer to encoded context data (from Iris.get_device_context()) diff --git a/iris/mem/triton/types.py b/iris/mem/triton/types.py index d04097824..33ea9d795 100644 --- a/iris/mem/triton/types.py +++ b/iris/mem/triton/types.py @@ -109,6 +109,7 @@ class TileView: TileView storing BOTH runtime coordinates AND compile-time block sizes. This class uses the @constexpr_function pattern: + - Stores runtime coordinates (pid_m, pid_n) as tl.tensor (computed from tl.program_id) - Stores compile-time block sizes (block_m, block_n) as tl.constexpr - Constructor is decorated with @constexpr_function to execute at compile-time From 7483336f732dfc96819e5cd254dd5657e105c5f6 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 12 Oct 2026 14:22:18 -0700 Subject: [PATCH 08/18] Stop the tile types page rendering everything twice conf.py sets autodoc_default_options members=True, so .. autoclass already emits every method of the class. The page additionally spelled each one out as its own .. automethod section underneath, so TensorView.tile_ptr and friends appeared once inside the class block and again as a standalone section -- 43 rendered signatures for 36 real objects. Dropped the explicit sections; the autoclass directives carry them. Two more docstring sections were rendering as run-on text. Napoleon recognises "Example:" but not "Example usage:" or "Fields:", so those fell through as plain paragraphs and the code samples collapsed onto one line: Example usage: view = make_tensor_view(ptr, M, N, ...) tile_ptr, mask = ... "Example usage:" -> "Example:" in TileView, Tile, TensorView and Iris.CCL, and "Fields:" -> "Attributes:" in AllReduceConfig, which additionally promotes variant_code and locks_ptr to documented attributes carrying their types rather than one undifferentiated blob. The earlier sweep missed these because it only looked for bullet lists under a colon, and these are indented code under a colon. Verified on 3.13 across all nine API pages: no method rendered twice, no collapsed blocks, no mock reprs, zero autodoc import failures. Co-Authored-By: Claude --- docs/reference/triton/types.md | 54 ---------------------------------- iris/host/iris.py | 2 +- iris/mem/triton/types.py | 8 ++--- 3 files changed, 5 insertions(+), 59 deletions(-) diff --git a/docs/reference/triton/types.md b/docs/reference/triton/types.md index ccfe01303..39c615adf 100644 --- a/docs/reference/triton/types.md +++ b/docs/reference/triton/types.md @@ -16,36 +16,6 @@ configuration consumed by the all-reduce methods on :noindex: ``` -### TensorView.tile_ptr -```{eval-rst} -.. automethod:: iris.mem.triton.types.TensorView.tile_ptr - :noindex: -``` - -### TensorView.tile_ptr_from_indices -```{eval-rst} -.. automethod:: iris.mem.triton.types.TensorView.tile_ptr_from_indices - :noindex: -``` - -### TensorView.offset_tile_ptr -```{eval-rst} -.. automethod:: iris.mem.triton.types.TensorView.offset_tile_ptr - :noindex: -``` - -### TensorView.offset -```{eval-rst} -.. automethod:: iris.mem.triton.types.TensorView.offset - :noindex: -``` - -### TensorView.number_of_tiles -```{eval-rst} -.. automethod:: iris.mem.triton.types.TensorView.number_of_tiles - :noindex: -``` - ## TileView ```{eval-rst} @@ -53,18 +23,6 @@ configuration consumed by the all-reduce methods on :noindex: ``` -### TileView.layout -```{eval-rst} -.. automethod:: iris.mem.triton.types.TileView.layout - :noindex: -``` - -### TileView.indices -```{eval-rst} -.. automethod:: iris.mem.triton.types.TileView.indices - :noindex: -``` - ## Tile ```{eval-rst} @@ -72,18 +30,6 @@ configuration consumed by the all-reduce methods on :noindex: ``` -### Tile.layout -```{eval-rst} -.. automethod:: iris.mem.triton.types.Tile.layout - :noindex: -``` - -### Tile.indices -```{eval-rst} -.. automethod:: iris.mem.triton.types.Tile.indices - :noindex: -``` - ## AllReduceConfig ```{eval-rst} diff --git a/iris/host/iris.py b/iris/host/iris.py index 2f7318581..a24a87fd4 100644 --- a/iris/host/iris.py +++ b/iris/host/iris.py @@ -1516,7 +1516,7 @@ class CCL: Collective Communication Library (CCL) interface for Iris. Provides collective operations that can be called as methods on the Iris instance. - Example usage: + Example: >>> ctx = iris.iris() >>> ctx.ccl.all_to_all(output_tensor, input_tensor) """ diff --git a/iris/mem/triton/types.py b/iris/mem/triton/types.py index 33ea9d795..bb52daea1 100644 --- a/iris/mem/triton/types.py +++ b/iris/mem/triton/types.py @@ -114,7 +114,7 @@ class TileView: - Stores compile-time block sizes (block_m, block_n) as tl.constexpr - Constructor is decorated with @constexpr_function to execute at compile-time - Example usage: + Example: pid = tl.program_id(0) pid_m = pid // num_tiles_n pid_n = pid % num_tiles_n @@ -155,7 +155,7 @@ class Tile: Extends TileView concept to include a data field for storing computed tile data (e.g., GEMM results in registers). - Example usage: + Example: c = tl.dot(a, b) tile = Tile(pid_m, pid_n, BLOCK_M, BLOCK_N, c) ctx.all_reduce_atomic(tile, dst_view) @@ -219,7 +219,7 @@ class TensorView: """ TensorView storing pointer and tensor metadata. - Example usage: + Example: view = make_tensor_view(ptr, M, N, stride_m, stride_n) tile_ptr, mask = view.tile_ptr(tile) """ @@ -288,7 +288,7 @@ class AllReduceConfig: """ Device-side configuration for all_reduce collective. - Fields: + Attributes: variant_code: Integer code for algorithm variant (0=atomic, 1=ring, 2=one_shot, 3=two_shot, 4=spinlock) locks_ptr: Pointer to locks array (ALWAYS required - pass dummy tensor if variant doesn't need locks) """ From 7529821e40025737334a0a2b8777eec24a3774d8 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 12 Oct 2026 14:50:56 -0700 Subject: [PATCH 09/18] Fix the run-on examples and document the tile type fields Two faults on the tile types page, both in iris/mem/triton/types.py. Examples rendered on one line. The Example: sections in TensorView, TileView and Tile hold plain indented code, and RST joins consecutive lines into a single paragraph, so the reader got view = make_tensor_view(ptr, M, N, stride_m, stride_n) tile_ptr, mask = ... Every other example in the codebase uses >>> doctest style, which preserves the breaks. These three now match. Fields had no descriptions. ptr, M, N, stride_m, stride_n on TensorView, and the pid/block fields on TileView and Tile, are annotated but were never described, so undoc-members rendered them as a bare "ptr: tensor" list under __init__. Added Attributes: sections. The wording for TensorView is taken from make_tensor_view's existing Args, which documents the same five values, so the two agree rather than drifting. iris/mem/triton/ops.py put() looked like a fourth instance but is not: its example is already >>> style with a caption line above it, which renders correctly. Verified on 3.13: 150 signatures across nine API pages, no run-on examples, no collapsed lists, no duplicate sections, no mock reprs, zero import failures. The run-on check was potency-tested against both known-bad strings before its zero was believed. Co-Authored-By: Claude --- iris/mem/triton/types.py | 40 ++++++++++++++++++++++++++++++---------- 1 file changed, 30 insertions(+), 10 deletions(-) diff --git a/iris/mem/triton/types.py b/iris/mem/triton/types.py index bb52daea1..4b0081ff8 100644 --- a/iris/mem/triton/types.py +++ b/iris/mem/triton/types.py @@ -114,12 +114,18 @@ class TileView: - Stores compile-time block sizes (block_m, block_n) as tl.constexpr - Constructor is decorated with @constexpr_function to execute at compile-time + Attributes: + pid_m: Tile row index, computed from ``tl.program_id``. + pid_n: Tile column index, computed from ``tl.program_id``. + block_m: Tile height. Compile-time constant. + block_n: Tile width. Compile-time constant. + Example: - pid = tl.program_id(0) - pid_m = pid // num_tiles_n - pid_n = pid % num_tiles_n - tile = TileView(pid_m, pid_n, BLOCK_M, BLOCK_N) - rm, rn, mask = tile.layout(M, N) + >>> pid = tl.program_id(0) + >>> pid_m = pid // num_tiles_n + >>> pid_n = pid % num_tiles_n + >>> tile = TileView(pid_m, pid_n, BLOCK_M, BLOCK_N) + >>> rm, rn, mask = tile.layout(M, N) """ pid_m: tl.tensor @@ -155,10 +161,17 @@ class Tile: Extends TileView concept to include a data field for storing computed tile data (e.g., GEMM results in registers). + Attributes: + pid_m: Tile row index, computed from ``tl.program_id``. + pid_n: Tile column index, computed from ``tl.program_id``. + block_m: Tile height. Compile-time constant. + block_n: Tile width. Compile-time constant. + data: The tile's values, for example a GEMM result held in registers. + Example: - c = tl.dot(a, b) - tile = Tile(pid_m, pid_n, BLOCK_M, BLOCK_N, c) - ctx.all_reduce_atomic(tile, dst_view) + >>> c = tl.dot(a, b) + >>> tile = Tile(pid_m, pid_n, BLOCK_M, BLOCK_N, c) + >>> ctx.all_reduce_atomic(tile, dst_view) """ pid_m: tl.tensor @@ -219,9 +232,16 @@ class TensorView: """ TensorView storing pointer and tensor metadata. + Attributes: + ptr: Pointer to the tensor's data. + M: Number of rows. + N: Number of columns. + stride_m: Stride in the M dimension. + stride_n: Stride in the N dimension. + Example: - view = make_tensor_view(ptr, M, N, stride_m, stride_n) - tile_ptr, mask = view.tile_ptr(tile) + >>> view = make_tensor_view(ptr, M, N, stride_m, stride_n) + >>> tile_ptr, mask = view.tile_ptr(tile) """ ptr: tl.tensor From 87299af81ef5e6de0e608acd5c76a35f876cb567 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 12 Oct 2026 15:00:54 -0700 Subject: [PATCH 10/18] Stop the tile types page listing undocumented members autodoc_default_options sets undoc-members and special-members globally, so every annotated field was rendered a third time as a bare "ptr: tensor" entry with an empty body, underneath an equally empty __init__. The same five values already appear twice above: once in Parameters, once as documented attributes. The "no-" prefix turns a default option off per directive, so the five directives on this page now carry :no-undoc-members: and :no-special-members:. Scoped to this page -- the setting is left alone elsewhere. Co-Authored-By: Claude --- docs/reference/triton/types.md | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/docs/reference/triton/types.md b/docs/reference/triton/types.md index 39c615adf..d1e3d9060 100644 --- a/docs/reference/triton/types.md +++ b/docs/reference/triton/types.md @@ -14,6 +14,8 @@ configuration consumed by the all-reduce methods on ```{eval-rst} .. autoclass:: iris.mem.triton.types.TensorView :noindex: + :no-undoc-members: + :no-special-members: ``` ## TileView @@ -21,6 +23,8 @@ configuration consumed by the all-reduce methods on ```{eval-rst} .. autoclass:: iris.mem.triton.types.TileView :noindex: + :no-undoc-members: + :no-special-members: ``` ## Tile @@ -28,6 +32,8 @@ configuration consumed by the all-reduce methods on ```{eval-rst} .. autoclass:: iris.mem.triton.types.Tile :noindex: + :no-undoc-members: + :no-special-members: ``` ## AllReduceConfig @@ -35,6 +41,8 @@ configuration consumed by the all-reduce methods on ```{eval-rst} .. autoclass:: iris.mem.triton.types.AllReduceConfig :noindex: + :no-undoc-members: + :no-special-members: ``` ## make_tensor_view @@ -42,4 +50,6 @@ configuration consumed by the all-reduce methods on ```{eval-rst} .. autofunction:: iris.mem.triton.types.make_tensor_view :noindex: + :no-undoc-members: + :no-special-members: ``` From 2f73feb82df27891f291c3d0dc7a5b55832e4b70 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 12 Oct 2026 15:05:53 -0700 Subject: [PATCH 11/18] Restore make_tensor_view, emptied by an autoclass-only option The previous commit added :no-undoc-members: and :no-special-members: to all five directives on the page, but make_tensor_view is an autofunction and those options belong to autoclass. Sphinx accepted them without complaint, built cleanly, and emitted the section heading with nothing underneath -- so the function silently vanished from the page while the build still reported success. Removed them from that one directive. The four autoclass directives keep them. Co-Authored-By: Claude --- docs/reference/triton/types.md | 2 -- 1 file changed, 2 deletions(-) diff --git a/docs/reference/triton/types.md b/docs/reference/triton/types.md index d1e3d9060..93615770d 100644 --- a/docs/reference/triton/types.md +++ b/docs/reference/triton/types.md @@ -50,6 +50,4 @@ configuration consumed by the all-reduce methods on ```{eval-rst} .. autofunction:: iris.mem.triton.types.make_tensor_view :noindex: - :no-undoc-members: - :no-special-members: ``` From e381bc78ac4aa6d011080a582cae9e1c8a80cb02 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 12 Oct 2026 15:57:02 -0700 Subject: [PATCH 12/18] Move the shutil import in conf.py to the top of the file ruff E402. The image-copy hook added its own `import shutil as _shutil` next to the function that uses it, which reads fine but is the one lint error in the repo -- and CI runs `ruff check . --fix`, which cannot auto-fix E402, so the job goes red rather than quietly correcting it. Rebuilt after the move: still 29 of 29 images copied into _build/html/images, so the hook is intact and not merely syntactically valid. Co-Authored-By: Claude --- docs/conf.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/conf.py b/docs/conf.py index ec703b641..0a0660de7 100644 --- a/docs/conf.py +++ b/docs/conf.py @@ -11,6 +11,7 @@ # documentation root, use os.path.abspath to make it absolute, like shown here. # import os +import shutil import sys from unittest.mock import MagicMock @@ -236,7 +237,6 @@ class GluonMock: # Copying the directory wholesale fixes all of them without rewriting the # theme-switching markup. # --------------------------------------------------------------------------- -import shutil as _shutil def _copy_images(app, exception): @@ -244,7 +244,7 @@ def _copy_images(app, exception): return src = os.path.join(os.path.dirname(__file__), "images") if os.path.isdir(src): - _shutil.copytree(src, os.path.join(app.outdir, "images"), dirs_exist_ok=True) + shutil.copytree(src, os.path.join(app.outdir, "images"), dirs_exist_ok=True) def setup(app): From f0865b23b56e1e0638650fb24403ca56e6edea67 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 12 Oct 2026 15:58:21 -0700 Subject: [PATCH 13/18] Revert the docs workflow to upstream The workflow_dispatch addition was for previewing on a fork, which gets no pushes of its own. Upstream main does, so the trigger it already has is sufficient there and this does not belong in the change. Co-Authored-By: Claude --- .github/workflows/docs.yml | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/.github/workflows/docs.yml b/.github/workflows/docs.yml index 3fc96f0f4..902ee6a2e 100644 --- a/.github/workflows/docs.yml +++ b/.github/workflows/docs.yml @@ -1,10 +1,6 @@ name: Build and Deploy Documentation on: - # Rebuild and republish from the Actions tab without needing a commit. This - # fork exists only to preview the docs, so otherwise the sole way to refresh - # the site is to push something to main. - workflow_dispatch: push: branches: [ main ] paths: @@ -78,7 +74,7 @@ jobs: url: ${{ steps.deployment.outputs.page_url }} runs-on: ubuntu-latest needs: build - if: github.ref == 'refs/heads/main' && (github.event_name == 'push' || github.event_name == 'workflow_dispatch') + if: github.ref == 'refs/heads/main' && github.event_name == 'push' permissions: contents: read pages: write From f3cf752986cd085604e8e800dcfb742c4ffe1131 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Tue, 13 Oct 2026 08:23:22 -0700 Subject: [PATCH 14/18] Keep the docs build gate out of this PR Swept in by a git add -A during the merge. It is a separate proposal and this PR was deliberately scoped back to docs content after the workflow change was reverted. Kept at /tmp/check_docs_build.py pending a decision on whether it lands here or on its own. --- docs/check_docs_build.py | 241 --------------------------------------- 1 file changed, 241 deletions(-) delete mode 100644 docs/check_docs_build.py diff --git a/docs/check_docs_build.py b/docs/check_docs_build.py deleted file mode 100644 index 8b18351bf..000000000 --- a/docs/check_docs_build.py +++ /dev/null @@ -1,241 +0,0 @@ -#!/usr/bin/env python3 -"""Gate a docs build on things Sphinx will not fail on by itself. - -Run after building: - - python3 -m sphinx -b html -d _build/doctrees . _build/html 2>&1 | tee build.log - python3 check_docs_build.py --log build.log --html _build/html - -Why this exists, concretely. Two separate incidents shipped an API reference -that was partly empty while the build reported success: - -* `iris/host/iris.py` imports `xio` at module load. Docs CI installs only - `docs/sphinx/requirements.txt`, so every autodoc directive raised - ModuleNotFoundError and the reference pages published as headings with no - content. Sphinx logged 71 warnings and exited 0. -* A directive carried `:no-undoc-members:`, which is autoclass-only. On an - autofunction Sphinx logged - `ERROR: An option to autofunction is either unknown or has an invalid value` - and then emitted the section heading with nothing under it. Exit code 0, job - green, the function gone from the page. - -So the log is worth failing on -- but it is not sufficient. A separate class of -defect produces perfectly valid HTML and zero log output: mock objects rendered -into signatures, bullet lists and code examples collapsed into run-on -paragraphs, members documented twice, members rendered with no description. A -build log scan finds none of those. They are only visible in the output. - -Hence two layers: LOG for things Sphinx noticed and discarded, HTML for things -it never noticed. -""" - -import argparse -import re -import sys -from pathlib import Path - -# Warnings that are not this repo's fault and not worth failing on. Anything -# else in the log is fatal. Keep this list short and justify every entry -- -# a long allowlist is how the xio failure stayed invisible for so long. -LOG_ALLOWLIST = [ - # rocm-docs-core looks the project up in an external registry. - r"Current project 'Iris' not found in projects", - # Four external inventories that 404 or serve a malformed header. Network - # dependent, so failing on these would make CI flaky rather than strict. - r"failed to reach any of the inventories", - r"intersphinx inventory .* not (fetchable|readable)", -] - - -def scan_log(text): - out = [] - for line in text.splitlines(): - if not re.search(r"\b(WARNING|ERROR|SEVERE)\b", line): - continue - if any(re.search(p, line) for p in LOG_ALLOWLIST): - continue - out.append(line.strip()[:160]) - return out - - -# -------------------------------------------------------------------------- -# HTML checks. Each returns a list of offending items. -# -------------------------------------------------------------------------- - -def _article(html): - m = re.search(r"]*>(.*?)", html, re.S) - return m.group(1) if m else "" - - -def c_mock(a): - """Mocked modules rendered into a signature, e.g. - tracing ().""" - return re.findall(r"MagicMock name='[^']*'", a) - - -def c_collapsed_list(a): - """A bullet list with no blank line before it, so RST folds the dashes - into the preceding paragraph.""" - return [x for x in re.findall(r"

([^<]{25,}?)

", a) - if re.search(r":\s+[-*]\s", x)] - - -def c_runon_code(a): - """An Example written as bare indented text rather than >>> or a literal - block, so consecutive statements render as one line.""" - return [x for x in re.findall(r"

([^<]{30,}?)

", a) - if len(re.findall(r"=\s*\w+[.(]", x)) >= 2] - - -def c_undocumented(a): - """An API object rendered with an empty body -- present on the page with - no description at all.""" - out = [] - for m in re.finditer(r'
(.*?)
', a, re.S): - kind, blk = m.group(1), m.group(2) - dt = re.search(r'([^<]+)', blk) - for dd in re.findall(r"
(.*?)
", blk, re.S): - if not re.sub(r"<[^>]+>|\s", "", dd): - out.append("%s %s" % (kind, dt.group(1) if dt else "?")) - break - return out - - -def c_empty_section(a): - """A section whose entire body is its own heading -- the directive under - it rendered nothing.""" - out = [] - for m in re.finditer(r'
(.*?)
', a, re.S): - sid, blk = m.group(1), m.group(2) - if ".*?", "", blk, flags=re.S) - if not re.sub(r"<[^>]+>|\s|&#?\w+;", "", rest): - out.append(sid) - return out - - -def c_dup_member(a): - """The same method rendered twice on one page, which happens when a page - spells out automethod sections under an autoclass that already has - members enabled.""" - heads = re.findall(r"

([A-Za-z_]+\.[A-Za-z_]+)tracing (<MagicMock name='mock.constexpr' id='139779281861264'>)

", - "collapsed-list": - "

The context tensor has the format: - [cur_rank, num_ranks, heap_base_0] " - "- If tracing=True: extracts trace buffer pointers

", - "runon-example": - "

view = make_tensor_view(ptr, M, N, stride_m, stride_n) tile_ptr, mask " - "= view.tile_ptr(tile)

", - "undocumented-member": - '
' - 'ptr' - ':' - ' tensor
' - "
", - "empty-section": - '

make_tensor_view' - '#

', - "duplicate-member": - '

TensorView.tile_ptr#

', - "raw-markup": - "

```{eval-rst} .. autoclass:: iris.mem.triton.types.Tile

", -} - -# A correctly rendered section. Every check must stay silent on it, or it is -# flagging normal output and its hits mean nothing. -GOOD = ( - '

TensorView#

' - '
' - 'TensorView
' - "

TensorView storing pointer and tensor metadata.

" - "

Compute tile pointer and mask using the stored dimensions and strides.

" - "
" -) - - -def potency(): - """Every check must fire on real defective markup and stay quiet on good - markup. A check that cannot do both is reported as dead rather than - silently contributing a zero.""" - dead = [] - for name, fn in CHECKS: - if not fn(BAD[name]) or fn(GOOD): - dead.append(name) - return dead - - -def main(): - ap = argparse.ArgumentParser() - ap.add_argument("--log", type=Path, help="sphinx build log") - ap.add_argument("--html", type=Path, help="built html directory") - ap.add_argument("--selftest", action="store_true", - help="run the potency gate and exit") - args = ap.parse_args() - - dead = potency() - print("potency: %d/%d checks can detect their own defect" - % (len(CHECKS) - len(dead), len(CHECKS))) - if dead: - print("FAIL: these checks are inert: %s" % dead) - return 2 - if args.selftest: - return 0 - - failures = 0 - - if args.log: - bad = scan_log(args.log.read_text(errors="replace")) - print("\nlog: %d non-allowlisted warnings/errors" % len(bad)) - for b in bad[:25]: - print(" %s" % b) - failures += len(bad) - - if args.html: - pages = sorted(args.html.rglob("*.html")) - print("\nhtml: %d pages" % len(pages)) - for p in pages: - a = _article(p.read_text(errors="replace")) - if not a: - continue - hits = {n: fn(a) for n, fn in CHECKS} - hits = {n: v for n, v in hits.items() if v} - if hits: - rel = p.relative_to(args.html) - print(" %s" % rel) - for n, v in hits.items(): - print(" %-22s %d e.g. %s" - % (n, len(v), re.sub(r"\s+", " ", str(v[0]))[:70])) - failures += len(v) - - print("\n%s (%d issues)" % ("FAIL" if failures else "PASS", failures)) - return 1 if failures else 0 - - -if __name__ == "__main__": - sys.exit(main()) From 08552000edbd1cb0b88d44416ea464c70d3c785b Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Tue, 13 Oct 2026 08:29:52 -0700 Subject: [PATCH 15/18] Fix two docs breaks that arrived with #568 Quote the return annotation on get_copy_engine_ctx. #568 added `-> torch.Tensor | None`, but torch is in autodoc_mock_imports, so the annotation evaluates mock | None at class-definition time and raises TypeError: unsupported operand type(s) for |: 'Tensor' and 'NoneType'. That killed the import of iris.host.iris and took 45 autodoc directives with it -- the API reference went empty again, by a different route than the xio import did. Quoting defers evaluation; it is the only PEP 604 annotation on a mocked type in iris/, and nothing introspects annotations at runtime. Defer the get_copy_engine_ctx index entry to class.md. #568 documents the method in its new "Optional SDMA Copy Engine" section and this branch already documented it on the dedicated copy-engine page, so Sphinx warned about a duplicate object description. Added :no-index: to this branch's copy rather than touching the section #568 just added. Verified on 3.13: 39 pages, 0 autodoc import failures, 4 warnings (all external intersphinx), and every page rendered in headless Chrome over HTTP -- no dead links, no broken images, no page errors beyond two the rocm-docs-core theme also throws on rocm.github.io/iris. jQuery errors per page: 4 upstream, 0 here. --- docs/reference/host/copy-engine.md | 1 + iris/host/iris.py | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/docs/reference/host/copy-engine.md b/docs/reference/host/copy-engine.md index 2b50bffd2..44387072d 100644 --- a/docs/reference/host/copy-engine.md +++ b/docs/reference/host/copy-engine.md @@ -43,4 +43,5 @@ host-initiated goes through the copy engine. ```{eval-rst} .. automethod:: iris.host.iris.Iris.get_copy_engine_ctx + :no-index: ``` diff --git a/iris/host/iris.py b/iris/host/iris.py index ed4687379..081ab37c5 100644 --- a/iris/host/iris.py +++ b/iris/host/iris.py @@ -984,7 +984,7 @@ def get_heap_bases(self): """ return self.heap_bases - def get_copy_engine_ctx(self) -> torch.Tensor | None: + def get_copy_engine_ctx(self) -> "torch.Tensor | None": """ Get the device-side copy engine context for SDMA operations. From f0dc6c2a78e56e66595e0ef771212475abd9db07 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Tue, 13 Oct 2026 09:01:46 -0700 Subject: [PATCH 16/18] Allow the docs workflow to be dispatched manually The workflow only fires on push to main under a paths filter. A fork gets no pushes of its own, so once the docs commits have landed there is no remaining event that can rebuild the site -- which is why this fork has never produced a run. The deploy job additionally gated on event_name == 'push', so even a manual run would have built an artifact and stopped short of publishing. Adds workflow_dispatch and widens that gate to accept it. The branch condition is unchanged -- still main only, so a PR build still cannot publish. Useful upstream too when the rendered output changes without the sources changing, such as a rocm-docs-core or Sphinx release. --- .github/workflows/docs.yml | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/.github/workflows/docs.yml b/.github/workflows/docs.yml index 902ee6a2e..e42f3fbce 100644 --- a/.github/workflows/docs.yml +++ b/.github/workflows/docs.yml @@ -1,6 +1,11 @@ name: Build and Deploy Documentation on: + # Rebuild and republish without a commit. Needed on a fork, which gets no + # pushes of its own and so can never trigger the paths filter below; also + # useful upstream when the rendered output changes without the sources + # changing, e.g. a rocm-docs-core or Sphinx release. + workflow_dispatch: push: branches: [ main ] paths: @@ -74,7 +79,7 @@ jobs: url: ${{ steps.deployment.outputs.page_url }} runs-on: ubuntu-latest needs: build - if: github.ref == 'refs/heads/main' && github.event_name == 'push' + if: github.ref == 'refs/heads/main' && (github.event_name == 'push' || github.event_name == 'workflow_dispatch') permissions: contents: read pages: write From 04d830f1fb082411dac5908ee0ea64e33097472b Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Tue, 13 Oct 2026 09:11:00 -0700 Subject: [PATCH 17/18] Revert the docs workflow to upstream Keeping this PR to docs content only. --- .github/workflows/docs.yml | 7 +------ 1 file changed, 1 insertion(+), 6 deletions(-) diff --git a/.github/workflows/docs.yml b/.github/workflows/docs.yml index e42f3fbce..902ee6a2e 100644 --- a/.github/workflows/docs.yml +++ b/.github/workflows/docs.yml @@ -1,11 +1,6 @@ name: Build and Deploy Documentation on: - # Rebuild and republish without a commit. Needed on a fork, which gets no - # pushes of its own and so can never trigger the paths filter below; also - # useful upstream when the rendered output changes without the sources - # changing, e.g. a rocm-docs-core or Sphinx release. - workflow_dispatch: push: branches: [ main ] paths: @@ -79,7 +74,7 @@ jobs: url: ${{ steps.deployment.outputs.page_url }} runs-on: ubuntu-latest needs: build - if: github.ref == 'refs/heads/main' && (github.event_name == 'push' || github.event_name == 'workflow_dispatch') + if: github.ref == 'refs/heads/main' && github.event_name == 'push' permissions: contents: read pages: write From 57bbd409202d68cb387cc2aa839dfda4c1743c66 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Thu, 15 Oct 2026 11:13:08 -0700 Subject: [PATCH 18/18] Address review: move the device ops, drop the C++ jargon, stop suppressing anchors atomic_sub and quiet are iris.mem.triton.ops device functions, but sat on reference/triton/ops.md with the iris.ops host namespace because both surfaces are called "ops". Moved to device-functions.md -- atomic_sub with the other atomics, quiet under a new Synchronization heading. ops.md is now only the four fused host ops. "Free function" is C++ for a non-member function and means nothing in Python. Replaced with module-level function vs method, shown rather than described. :noindex: was copied from the Gluon page onto types.md and device-context.md, where nothing else documents those objects, so it suppressed their index anchors for no reason. Removed, 29 lines; the site index goes from 60 to 114 API objects. copy-engine.md keeps its one: #568 documents get_copy_engine_ctx on class.md too, so without it Sphinx warns about a duplicate object description. Checked by removing each one and rebuilding. Built on 3.13: 34 pages, 0 autodoc import failures, 0 duplicate objects. --- docs/reference/triton/device-context.md | 38 +++++------------------ docs/reference/triton/device-functions.md | 10 ++++++ docs/reference/triton/ops.md | 12 ------- docs/reference/triton/types.md | 5 --- 4 files changed, 18 insertions(+), 47 deletions(-) diff --git a/docs/reference/triton/device-context.md b/docs/reference/triton/device-context.md index a788f0b2b..eff6ab1b6 100644 --- a/docs/reference/triton/device-context.md +++ b/docs/reference/triton/device-context.md @@ -26,12 +26,14 @@ them. They are a different surface from the host collectives in kernels and are called from Python. ``` -```{note} -Every operation here also exists as a free function -- see -[Device Functions](device-functions.md). The free form is explicit about both -ends of the transfer and the heap: `load(pointer, to_rank, from_rank, heap_bases)`. -The bound form below knows the current rank and carries the heap bases itself, so -the same call is `ctx.load(pointer, from_rank)`. +Every operation here is also available as a module-level function -- see +[Device Functions](device-functions.md). The module-level form takes the heap +bases and both ranks explicitly; the context already knows the current rank and +holds the heap bases, so its methods take fewer arguments: + +```python +iris.load(pointer, to_rank, from_rank, heap_bases) # module-level function +ctx.load(pointer, from_rank) # method on the context ``` ## Initialization @@ -39,7 +41,6 @@ the same call is `ctx.load(pointer, from_rank)`. ### initialize ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.initialize - :noindex: ``` ## Memory transfer operations @@ -47,31 +48,26 @@ the same call is `ctx.load(pointer, from_rank)`. ### load ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.load - :noindex: ``` ### store ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.store - :noindex: ``` ### get ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.get - :noindex: ``` ### put ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.put - :noindex: ``` ### copy ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.copy - :noindex: ``` ## Atomic operations @@ -79,55 +75,46 @@ the same call is `ctx.load(pointer, from_rank)`. ### atomic_add ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.atomic_add - :noindex: ``` ### atomic_sub ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.atomic_sub - :noindex: ``` ### atomic_cas ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.atomic_cas - :noindex: ``` ### atomic_xchg ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.atomic_xchg - :noindex: ``` ### atomic_xor ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.atomic_xor - :noindex: ``` ### atomic_and ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.atomic_and - :noindex: ``` ### atomic_or ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.atomic_or - :noindex: ``` ### atomic_min ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.atomic_min - :noindex: ``` ### atomic_max ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.atomic_max - :noindex: ``` ## Collective operations @@ -135,53 +122,44 @@ the same call is `ctx.load(pointer, from_rank)`. ### all_reduce_atomic ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.all_reduce_atomic - :noindex: ``` ### all_reduce_spinlock ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.all_reduce_spinlock - :noindex: ``` ### all_reduce_one_shot ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.all_reduce_one_shot - :noindex: ``` ### all_reduce_ring ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.all_reduce_ring - :noindex: ``` ### all_reduce_two_shot ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.all_reduce_two_shot - :noindex: ``` ### all_gather ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.all_gather - :noindex: ``` ### gather ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.gather - :noindex: ``` ### all_to_all ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.all_to_all - :noindex: ``` ### reduce_scatter ```{eval-rst} .. automethod:: iris.mem.triton.context.Context.reduce_scatter - :noindex: ``` diff --git a/docs/reference/triton/device-functions.md b/docs/reference/triton/device-functions.md index c9af17dd3..92ef5d07a 100644 --- a/docs/reference/triton/device-functions.md +++ b/docs/reference/triton/device-functions.md @@ -37,6 +37,11 @@ Device-side functions provided by Iris for remote memory operations and atomics. .. autofunction:: iris.mem.triton.ops.atomic_add ``` +### atomic_sub +```{eval-rst} +.. autofunction:: iris.mem.triton.ops.atomic_sub +``` + ### atomic_cas ```{eval-rst} .. autofunction:: iris.mem.triton.ops.atomic_cas @@ -72,4 +77,9 @@ Device-side functions provided by Iris for remote memory operations and atomics. .. autofunction:: iris.mem.triton.ops.atomic_max ``` +## Synchronization +### quiet +```{eval-rst} +.. autofunction:: iris.mem.triton.ops.quiet +``` diff --git a/docs/reference/triton/ops.md b/docs/reference/triton/ops.md index 16026e907..08b4e9e12 100644 --- a/docs/reference/triton/ops.md +++ b/docs/reference/triton/ops.md @@ -25,15 +25,3 @@ of running as a separate kernel after it. ```{eval-rst} .. automethod:: iris.ops.OpsNamespace.matmul_reduce_scatter ``` - -## atomic_sub - -```{eval-rst} -.. autofunction:: iris.mem.triton.ops.atomic_sub -``` - -## quiet - -```{eval-rst} -.. autofunction:: iris.mem.triton.ops.quiet -``` diff --git a/docs/reference/triton/types.md b/docs/reference/triton/types.md index 93615770d..3042a01d8 100644 --- a/docs/reference/triton/types.md +++ b/docs/reference/triton/types.md @@ -13,7 +13,6 @@ configuration consumed by the all-reduce methods on ```{eval-rst} .. autoclass:: iris.mem.triton.types.TensorView - :noindex: :no-undoc-members: :no-special-members: ``` @@ -22,7 +21,6 @@ configuration consumed by the all-reduce methods on ```{eval-rst} .. autoclass:: iris.mem.triton.types.TileView - :noindex: :no-undoc-members: :no-special-members: ``` @@ -31,7 +29,6 @@ configuration consumed by the all-reduce methods on ```{eval-rst} .. autoclass:: iris.mem.triton.types.Tile - :noindex: :no-undoc-members: :no-special-members: ``` @@ -40,7 +37,6 @@ configuration consumed by the all-reduce methods on ```{eval-rst} .. autoclass:: iris.mem.triton.types.AllReduceConfig - :noindex: :no-undoc-members: :no-special-members: ``` @@ -49,5 +45,4 @@ configuration consumed by the all-reduce methods on ```{eval-rst} .. autofunction:: iris.mem.triton.types.make_tensor_view - :noindex: ```