From 17667facec3c0b0b42dcc0d0f4bec24b4cfe54ba Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Thu, 27 Aug 2026 12:20:01 -0700 Subject: [PATCH 1/4] Drop the rotated-out internal ROCm apt repo before apt-get update MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The rocm/pytorch base images ship /etc/apt/sources.list.d/rocm.list pinned to a specific internal build on compute-artifactory.amd.com — currently compute-rocm-rel-7.1 20. That build has been rotated out upstream, so the index returns 404, apt-get update exits 100, and because the Dockerfile chains update && install the whole layer fails. With no image, every downstream CI job fails within seconds; the failures on recent PRs are this, not the PR contents. ROCm and torch are already installed in the base image and nothing here installs further rocm-* packages, so the repo is unnecessary. Removing it is also a no-op on images that do not carry the file. Verified against rocm/pytorch:rocm7.1_ubuntu24.04_py3.13_pytorch_release_2.9.1: apt-get update as-is exit 100 rm rocm.list then apt-get update exit 0 rm rocm.list then the full install line exit 0 rocm and torch after removal intact, 2.9.1+rocm7.1.0 Applied to Dockerfile.ccl and apptainer/iris.def as well, which share the same base image family and the same update-and-install chain. Dockerfile.dev builds from ubuntu:24.04 and is unaffected. --- apptainer/iris.def | 2 ++ docker/Dockerfile | 8 +++++++- docker/Dockerfile.ccl | 8 +++++++- 3 files changed, 16 insertions(+), 2 deletions(-) diff --git a/apptainer/iris.def b/apptainer/iris.def index 8842b2b96..bea9e2d67 100644 --- a/apptainer/iris.def +++ b/apptainer/iris.def @@ -13,6 +13,8 @@ From: rocm/pytorch:rocm7.2.1_ubuntu24.04_py3.14_pytorch_2.10.0 export PATH=\"\$ROCM_PATH/bin:\$PATH\" # Install system packages + # see docker/Dockerfile -- base image pins a rotated-out internal ROCm apt repo + rm -f /etc/apt/sources.list.d/rocm.list && \ apt-get update && \ DEBIAN_FRONTEND=noninteractive apt-get install -y \ git wget ninja-build cmake python3-pip python3-dev build-essential jq libdwarf-dev && \ diff --git a/docker/Dockerfile b/docker/Dockerfile index c01e86a8e..04126529a 100644 --- a/docker/Dockerfile +++ b/docker/Dockerfile @@ -20,7 +20,13 @@ ENV OMPI_ALLOW_RUN_AS_ROOT_CONFIRM=1 \ OMPI_ALLOW_RUN_AS_ROOT=1 # Install system packages -RUN apt-get update && \ +# The rocm/pytorch base images ship /etc/apt/sources.list.d/rocm.list pinned to a +# specific internal build (e.g. compute-rocm-rel-7.1 20) on compute-artifactory.amd.com. +# Once that build is rotated out upstream the index 404s, apt-get update returns 100, +# and the whole layer fails -- taking every downstream CI job with it. ROCm and torch +# are already installed in the image, so nothing here needs that repo. +RUN rm -f /etc/apt/sources.list.d/rocm.list && \ + apt-get update && \ DEBIAN_FRONTEND=noninteractive apt-get install -y \ git wget ninja-build cmake python3-pip python3-dev build-essential libdwarf-dev && \ rm -rf /var/lib/apt/lists/* diff --git a/docker/Dockerfile.ccl b/docker/Dockerfile.ccl index 8271c31dd..7a07dd5bd 100644 --- a/docker/Dockerfile.ccl +++ b/docker/Dockerfile.ccl @@ -19,7 +19,13 @@ ENV OMPI_ALLOW_RUN_AS_ROOT_CONFIRM=1 \ OMPI_ALLOW_RUN_AS_ROOT=1 # Install system packages -RUN apt-get update && \ +# The rocm/pytorch base images ship /etc/apt/sources.list.d/rocm.list pinned to a +# specific internal build (e.g. compute-rocm-rel-7.1 20) on compute-artifactory.amd.com. +# Once that build is rotated out upstream the index 404s, apt-get update returns 100, +# and the whole layer fails -- taking every downstream CI job with it. ROCm and torch +# are already installed in the image, so nothing here needs that repo. +RUN rm -f /etc/apt/sources.list.d/rocm.list && \ + apt-get update && \ DEBIAN_FRONTEND=noninteractive apt-get install -y \ git wget ninja-build cmake python3-pip python3-dev build-essential libdwarf-dev && \ rm -rf /var/lib/apt/lists/* From b0876d40460fbf4fdef96797492284f713db836d Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Fri, 28 Aug 2026 16:09:39 -0700 Subject: [PATCH 2/4] Re-run CI Empty commit to re-trigger the workflows. The container build now succeeds (iris-dev-triton-aafec41 present on the runner), and the earlier failures were the artifactory 404 plus the node running out of disk. From 7a32ff9c3cdec13964001082a1a5e53a0b2c341a Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Fri, 28 Aug 2026 18:37:27 -0700 Subject: [PATCH 3/4] Probe each HIP intrinsic separately and fall back to inline asm iris/mem/utils.py imported memrealtime and smid under one try/except. Upstream Triton at the pinned commit (bcbcabdd) exports memrealtime but not smid, so the import failed as a pair and both helpers fell back to stubs that call tl.static_assert(False). Any kernel recording a trace event then failed to compile: CompileTimeAssertionFailure: smid is unavailable in this Triton build at device_utils.get_cu_id() in record_event_start This was masked for months because the CI image is a fixed tag that was never rebuilt; the first clean rebuild surfaced it. Each intrinsic is now probed on its own, so a missing smid no longer disables timestamps, and each fallback emits the instruction the intrinsic would rather than refusing to compile. get_cu_id reads CU_ID from HW_REG_HW_ID exactly as get_xcc_id below it already reads HW_REG_XCC_ID. The CU_ID field is 4 bits, so it identifies the CU within its shader engine rather than globally -- pair it with get_xcc_id. Verified on gfx950: a 256-workgroup launch reports CU_ID 0-8 with XCC_ID 0-7, the latter matching the part's 8 XCDs. Architectures without a fallback report 0 instead of failing to compile, since tracing is diagnostic and losing CU attribution beats breaking every traced kernel. tests/unittests/test_device_context.py: 41 passed on gfx950. --- iris/mem/utils.py | 83 ++++++++++++++++++++++++++++++++++++++++------- 1 file changed, 71 insertions(+), 12 deletions(-) diff --git a/iris/mem/utils.py b/iris/mem/utils.py index d5b771d53..e50cbb626 100644 --- a/iris/mem/utils.py +++ b/iris/mem/utils.py @@ -13,16 +13,29 @@ import triton.language as tl from triton.language.target_info import is_hip_cdna3, is_hip_cdna4 +# Probe each intrinsic separately. They were added to +# ``triton.language.extra.hip`` at different times, so a build can export one +# and not the other -- upstream Triton at the pinned commit ships +# ``memrealtime`` but not ``smid``. Guarding both behind a single try/except +# meant a missing ``smid`` also disabled ``read_realtime``, and the fallback +# was ``static_assert(False)``, which fails compilation of every kernel that +# records a trace event rather than degrading the trace. try: from triton.language.extra.hip import memrealtime as _memrealtime + + _HAS_MEMREALTIME = True +except ImportError: + _HAS_MEMREALTIME = False + +try: from triton.language.extra.hip import smid as _smid - _HAS_HIP_INTRINSICS = True + _HAS_SMID = True except ImportError: - _HAS_HIP_INTRINSICS = False + _HAS_SMID = False -if _HAS_HIP_INTRINSICS: +if _HAS_MEMREALTIME: @triton.jit def read_realtime(): @@ -39,6 +52,34 @@ def read_realtime(): int64: Current timestamp in cycles (100 MHz constant clock) """ return _memrealtime() +else: + + @triton.jit + def read_realtime(): + """ + Read GPU wall clock timestamp on builds without ``tl.extra.hip.memrealtime``. + + Emits the same instruction the intrinsic would. gfx11/gfx12 use a + message rather than ``s_memrealtime``; there is no portable fallback for + those here, so they report 0 and timestamps are simply unavailable. + + Returns: + int64: Timestamp in cycles, or 0 where unsupported + """ + if is_hip_cdna3() or is_hip_cdna4(): + return tl.inline_asm_elementwise( + asm="s_memrealtime $0\n\ts_waitcnt vmcnt(0)", + constraints=("=s"), + args=[], + dtype=tl.int64, + is_pure=False, + pack=1, + ) + else: + return tl.cast(0, tl.int64) + + +if _HAS_SMID: @triton.jit def get_cu_id(): @@ -55,17 +96,35 @@ def get_cu_id(): return _smid() else: - @triton.jit - def read_realtime(): - """Fallback stub when HIP intrinsics are missing.""" - tl.static_assert(False, "memrealtime is unavailable in this Triton build") - return tl.cast(0, tl.int64) - @triton.jit def get_cu_id(): - """Fallback stub when HIP intrinsics are missing.""" - tl.static_assert(False, "smid is unavailable in this Triton build") - return tl.cast(0, tl.int32) + """ + Get compute-unit ID on builds without ``tl.extra.hip.smid``. + + Reads CU_ID out of ``HW_REG_HW_ID`` directly, the same mechanism + ``get_xcc_id`` below uses. The field is 4 bits, so this identifies the + CU within its shader engine rather than globally; pair it with + ``get_xcc_id`` for a fuller picture. Verified on gfx950, where a + 256-workgroup launch reports CU_ID 0-8 alongside XCC_ID 0-7. + + Other architectures report 0 rather than failing to compile: tracing is + diagnostic, and losing CU attribution is preferable to breaking every + traced kernel. + + Returns: + int32: CU ID within the shader engine, or 0 where unsupported + """ + if is_hip_cdna3() or is_hip_cdna4(): + return tl.inline_asm_elementwise( + asm="s_getreg_b32 $0, hwreg(HW_REG_HW_ID, 8, 4)", + constraints=("=s"), + args=[], + dtype=tl.int32, + is_pure=False, + pack=1, + ) + else: + return tl.cast(0, tl.int32) @triton.jit From 889caab6ba6fbe28229833a0e94ec76c90dca545 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Sat, 29 Aug 2026 18:52:03 -0700 Subject: [PATCH 4/4] Keep a communication share of the workgroups in WG specialization The specialized kernels launch num_sms workgroups and split them on pid < GEMM_SMS, giving the remainder to the communication path. Passing gemm_sms equal to the CU count leaves that remainder empty, so every workgroup takes the GEMM branch and the scatter never runs. Validation caught it as C=0.0, while the reported throughput went UP because the communication was simply skipped. The perf job hit this after moving to a 256-CU part, where its hardcoded --gemm_sms 256 matches the CU count exactly. The computed default has the same hole: 2**int(log2(cu_count)) equals cu_count whenever the CU count is a power of two, so it only ever worked on 304-CU parts by accident. Step the default down when it would leave nothing over, and reject an explicit value that does. Measured on 8x MI350X: the CI config now runs at 2904 TFLOPs against a 1440 threshold, and the default at 2530. Co-Authored-By: Claude --- .../workflows/iris-performance-regression-test.yml | 4 +++- .../benchmark.py | 12 ++++++++++++ .../benchmark.py | 12 ++++++++++++ 3 files changed, 27 insertions(+), 1 deletion(-) diff --git a/.github/workflows/iris-performance-regression-test.yml b/.github/workflows/iris-performance-regression-test.yml index c3f1507d0..59014abef 100644 --- a/.github/workflows/iris-performance-regression-test.yml +++ b/.github/workflows/iris-performance-regression-test.yml @@ -27,7 +27,9 @@ jobs: - example_name: "GEMM All-Scatter WG Specialization" example_path: "10_gemm_all_scatter_wg_specialization" tflops_threshold: 1440 # Actual: ~1802 TFLOPs (80% regression threshold) - benchmark_args: "-m 16384 -n 16384 -k 16384 --BLK_M 128 --BLK_N 128 --BLK_K 64 --gsize_m 6 --gemm_sms 256" + # gemm_sms must stay below the CU count -- the remainder is what runs the + # communication path. 256 left zero of it on a 256-CU part. + benchmark_args: "-m 16384 -n 16384 -k 16384 --BLK_M 128 --BLK_N 128 --BLK_K 64 --gsize_m 6 --gemm_sms 192" - example_name: "GEMM All-Scatter" example_path: "07_gemm_all_scatter" diff --git a/examples/10_gemm_all_scatter_wg_specialization/benchmark.py b/examples/10_gemm_all_scatter_wg_specialization/benchmark.py index 910ebdd6f..367677fa2 100755 --- a/examples/10_gemm_all_scatter_wg_specialization/benchmark.py +++ b/examples/10_gemm_all_scatter_wg_specialization/benchmark.py @@ -93,6 +93,18 @@ def _worker(local_rank: int, world_size: int, init_url: str, args: dict): if args["gemm_sms"] is None: # For wg_specialized: use next smaller power of 2 args["gemm_sms"] = 2 ** int(math.log2(cu_count)) if cu_count > 0 else 1 + # The kernel launches num_sms workgroups and gives whatever is left after + # gemm_sms to the communication path. When the CU count is itself a power + # of two (256 on MI350X) that leaves nothing, so step down once. + if args["gemm_sms"] >= args["num_sms"]: + args["gemm_sms"] //= 2 + + if args["gemm_sms"] >= args["num_sms"]: + raise ValueError( + f"gemm_sms ({args['gemm_sms']}) must be less than num_sms ({args['num_sms']}): " + "workgroup specialization needs the remainder for the communication path, " + "and with none of it the all-scatter never runs." + ) # GEMM datatype = torch.float32 diff --git a/examples/22_gemm_one_shot_reduce_scatter_wg_specialization/benchmark.py b/examples/22_gemm_one_shot_reduce_scatter_wg_specialization/benchmark.py index 47728618f..aeda82e80 100644 --- a/examples/22_gemm_one_shot_reduce_scatter_wg_specialization/benchmark.py +++ b/examples/22_gemm_one_shot_reduce_scatter_wg_specialization/benchmark.py @@ -89,6 +89,18 @@ def _worker(local_rank: int, world_size: int, init_url: str, args: dict): if args["gemm_sms"] is None: # Use next smaller power of 2 for GEMM SMs args["gemm_sms"] = 2 ** int(math.log2(cu_count)) if cu_count > 0 else 1 + # The kernel launches num_sms workgroups and gives whatever is left after + # gemm_sms to the communication path. When the CU count is itself a power + # of two (256 on MI350X) that leaves nothing, so step down once. + if args["gemm_sms"] >= args["num_sms"]: + args["gemm_sms"] //= 2 + + if args["gemm_sms"] >= args["num_sms"]: + raise ValueError( + f"gemm_sms ({args['gemm_sms']}) must be less than num_sms ({args['num_sms']}): " + "workgroup specialization needs the remainder for the communication path, " + "and with none of it the reduce-scatter never runs." + ) datatype = torch.float16 if args["datatype"] == "fp16":