From fe35c54c91b082a02ff6276a0cc3665571cc409d Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sun, 13 Sep 2026 15:53:44 -0400 Subject: [PATCH 1/8] rccl-tests: extract shared gin_sdma_harness.py (AICOMRCCL-2221) Move MPI launcher assembly, connectivity-gate retry, and Broadcast debug-log capture into gin_sdma_harness.py so AllGather, AllToAll, and Broadcast GIN-SDMA pytest harnesses share one implementation and stop drifting. Co-authored-by: Cursor --- projects/rccl-tests/test/gin_sdma_harness.py | 356 ++++++++++++++++++ projects/rccl-tests/test/test_AllGather.py | 149 ++++---- projects/rccl-tests/test/test_AllToAll.py | 182 +++++---- projects/rccl-tests/test/test_Broadcast.py | 371 +++---------------- 4 files changed, 554 insertions(+), 504 deletions(-) create mode 100644 projects/rccl-tests/test/gin_sdma_harness.py diff --git a/projects/rccl-tests/test/gin_sdma_harness.py b/projects/rccl-tests/test/gin_sdma_harness.py new file mode 100644 index 000000000000..b6159858b080 --- /dev/null +++ b/projects/rccl-tests/test/gin_sdma_harness.py @@ -0,0 +1,356 @@ +################################################################################# +# Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell cop- +# ies of the Software, and to permit persons to whom the Software is furnished +# to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in all +# copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IM- +# PLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS +# FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR +# COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER +# IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNE- +# CTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. +################################################################################ + +"""Shared helpers for GIN-SDMA rccl-tests pytest harnesses (AllGather, AllToAll, +Broadcast). MPI launcher assembly, connectivity-gate retry, and (for Broadcast) +NCCL debug-file capture live here so the per-collective test modules stay thin.""" + +import glob +import os +import re +import shlex +import shutil +import signal +import subprocess +import tempfile +import time + +import pytest + +MiB = 1024 * 1024 +GiB = 1024 * MiB + +# Intermittent gfx950 cuMem-VMM connectivity-gate abort (not a data error). +CONN_GATE_RE = re.compile( + r"LSA signal connectivity gate failed|unhandled system error", re.I +) +DATA_FAIL_RE = re.compile( + r"#wrong\s*=\s*[1-9]|mismatch|check.*fail|Out of bounds values\s*:\s*[1-9]", + re.I, +) + + +def detect_ngpus(): + if os.environ.get("ROCR_VISIBLE_DEVICES") is not None: + return len(os.environ["ROCR_VISIBLE_DEVICES"].split(",")) + if os.environ.get("HIP_VISIBLE_DEVICES") is not None: + return len(os.environ["HIP_VISIBLE_DEVICES"].split(",")) + try: + out = subprocess.check_output( + 'rocminfo | grep "Device Type:.\\s*.GPU" | wc -l', shell=True + ) + return int(out) + except Exception: + return 0 + + +def env_int(name, default): + """Parse an integer env var; bad values fall back so collection stays usable.""" + raw = os.environ.get(name) + if raw is None or raw == "": + return default + try: + return int(raw) + except ValueError: + return default + + +def mpi_launch_prefix(request, launcher, np, mpi_opts): + """Assemble [launcher, -np, np, ...mpi_opts, (-host hostfile)?].""" + hostfile = request.config.getoption("--hostfile") + launch = [launcher, "-np", str(np)] + list(mpi_opts) + if hostfile: + launch += ["-host", hostfile] + return launch + + +def gin_env_xflags(kv_list): + """Turn ["K=V", ...] into mpirun -x flags.""" + flags = [] + for kv in kv_list: + flags += ["-x", kv] + return flags + + +def launch_mpi_shell(cmd, timeout_s, hang_msg): + """Run a shell MPI command in a new session; kill the process group on timeout.""" + print(cmd) + proc = subprocess.Popen( + cmd, + shell=True, + universal_newlines=True, + stdout=subprocess.PIPE, + stderr=subprocess.STDOUT, + start_new_session=True, + ) + try: + out, _ = proc.communicate(timeout=timeout_s) + except subprocess.TimeoutExpired: + try: + os.killpg(os.getpgid(proc.pid), signal.SIGKILL) + except ProcessLookupError: + pass + out, _ = proc.communicate() + pytest.fail(hang_msg.format(out=(out or "")[-2000:])) + print(out) + return proc.returncode, out + + +def run_with_conn_gate_retry( + launch_fn, conn_retries, settle_s=3, is_data_failure=None +): + """Retry launch_fn on connectivity-gate aborts; never retry genuine data failures. + + launch_fn must return (returncode, stdout).""" + if is_data_failure is None: + is_data_failure = lambda out: bool(DATA_FAIL_RE.search(out or "")) + + rc, out = 1, "" + for attempt in range(1, max(1, conn_retries) + 1): + rc, out = launch_fn() + if rc == 0: + return rc, out + if is_data_failure(out): + return rc, out + if CONN_GATE_RE.search(out or "") and attempt < conn_retries: + print( + "=== connectivity-gate abort (attempt {}/{}); " + "re-launching after settle ===".format(attempt, conn_retries) + ) + time.sleep(settle_s) + continue + return rc, out + return rc, out + + +def read_nccl_debug_logs(debug_dir): + """Merged per-rank NCCL debug output written under debug_dir.""" + chunks = [] + for log_path in sorted(glob.glob(os.path.join(debug_dir, "nccl-debug.*"))): + try: + with open(log_path, errors="replace") as fh: + chunks.append(fh.read()) + except OSError: + pass + return "\n".join(chunks) + + +# --- Broadcast-specific parsing / assertions --------------------------------- + +PLUGIN_RE = re.compile(r"gin-anvil-sdma", re.I) + +_NUM = r"[-+]?(?:\d+\.?\d*|\.\d+)(?:[eE][-+]?\d+)?" +_WRONG = r"(?:{}|N/A)".format(_NUM) +BCAST_ROW_RE = re.compile( + r"^\s*(\d+)\s+(\d+)\s+\S+\s+\S+\s+(\d+)" + r"\s+{n}\s+{n}\s+{n}\s+({w})" + r"\s+{n}\s+{n}\s+{n}\s+({w})".format(n=_NUM, w=_WRONG) +) +BCAST_OOB_RE = re.compile(r"Out of bounds values\s*:\s*(\d+)") +BCAST_TIER_RE = re.compile(r"^#\[bcast-tier\]\s+(\S+)", re.M) + + +def bcast_tiers(out): + """Set of -D 3 tier names the run reported.""" + return set(BCAST_TIER_RE.findall(out or "")) + + +def bcast_rows(out): + """Measured rows as (size, root, wrong_out_of_place, wrong_in_place).""" + rows = [] + for line in (out or "").splitlines(): + m = BCAST_ROW_RE.match(line) + if m: + rows.append((int(m.group(1)), int(m.group(3)), m.group(4), m.group(5))) + return rows + + +def bcast_wrong_count(field): + """Numeric #wrong, or None when the column is N/A because checks were off.""" + try: + return float(field) + except ValueError: + return None + + +def bcast_data_failed(out): + """True when the run produced demonstrably wrong results. + + Deliberately not a text match on "Wrong": rccl-tests prints a "#wrong" column + header on every run that reaches the results table.""" + for _, _, oop, ip in bcast_rows(out): + if any(bcast_wrong_count(f) not in (None, 0.0) for f in (oop, ip)): + return True + m = BCAST_OOB_RE.search(out or "") + return bool(m and m.group(1) != "0") + + +def assert_bcast_ok(rc, out, debug, ctx, expect_tier=None): + """Assert the run bound GIN-SDMA, measured something, and checked clean.""" + tail = (out or "")[-2000:] + assert rc == 0, "Broadcast {} failed (exit {}). Output tail:\n{}".format( + ctx, rc, tail + ) + assert PLUGIN_RE.search(debug or "") or PLUGIN_RE.search(out or ""), ( + "Broadcast {} never bound the GIN Anvil-SDMA backend, so the run proves " + "nothing about the GIN path. Output tail:\n{}".format(ctx, tail) + ) + rows = bcast_rows(out) + assert rows, "Broadcast {} produced no measured rows. Output tail:\n{}".format( + ctx, tail + ) + unchecked = [ + (s, r) + for s, r, oop, ip in rows + if bcast_wrong_count(oop) is None or bcast_wrong_count(ip) is None + ] + assert not unchecked, ( + "Broadcast {} reported #wrong as N/A at (size, root) {}, so the data check " + "never ran and the result is unverified.".format(ctx, unchecked) + ) + bad = [ + (s, r, oop, ip) + for s, r, oop, ip in rows + if bcast_wrong_count(oop) != 0.0 or bcast_wrong_count(ip) != 0.0 + ] + assert ( + not bad + ), "Broadcast {} reported nonzero #wrong (size, root, oop, ip): {}".format( + ctx, bad + ) + m = BCAST_OOB_RE.search(out or "") + assert m and m.group(1) == "0", "Broadcast {} out-of-bounds count is {}".format( + ctx, m.group(1) if m else "absent" + ) + if expect_tier is not None: + allowed = ( + {expect_tier} if isinstance(expect_tier, str) else set(expect_tier) + ) + tiers = bcast_tiers(out) + assert tiers, ( + "Broadcast {} reported no #[bcast-tier] line, so which tier ran is " + "unverified. Is RCCL_TESTS_BCAST_TIER_LOG set for this launch, and is " + "the binary new enough to emit it? Output tail:\n{}".format(ctx, tail) + ) + assert len(tiers) == 1 and tiers <= allowed, ( + "Broadcast {} ran tier(s) {} but expected exactly one of {}. The tier " + "gate moved, or the tier under test fell through to another " + "kernel.".format(ctx, sorted(tiers), sorted(allowed)) + ) + + +def launch_bcast_gin_sdma( + request, + msg_bytes, + dtype, + *, + exe, + np, + launcher, + mpi_opts, + xenv, + ctas, + gin_type, + timeout_s, + force_flat_gin=False, + force_sag_gin=False, +): + """Launch broadcast_perf -D 3; returns (returncode, stdout, debug_text).""" + size = str(int(msg_bytes)) + debug_dir = tempfile.mkdtemp(prefix="bcast-gin-dbg-") + gin_env = [ + "NCCL_GIN_ENABLE=1", + "NCCL_GIN_TYPE={}".format(gin_type), + "NCCL_GIN_ANVIL_SDMA_THRESHOLD=0", + "NCCL_GIN_ANVIL_SDMA_THRESHOLD_BROADCAST=0", + "NCCL_DEBUG=INFO", + "NCCL_DEBUG_SUBSYS=INIT,NET", + "NCCL_DEBUG_FILE={}/nccl-debug.%h.%p.log".format(debug_dir), + "RCCL_TESTS_BCAST_TIER_LOG=1", + ] + if force_flat_gin: + gin_env += [ + "NCCL_GIN_ANVIL_BCAST_SCATTER_AG_MIN_BYTES=0", + "NCCL_GIN_ANVIL_BCAST_RING_MIN_BYTES=0", + ] + elif force_sag_gin: + gin_env += ["NCCL_GIN_ANVIL_BCAST_RING_MIN_BYTES=0"] + gin_env += list(xenv) + + args = ( + mpi_launch_prefix(request, launcher, np, mpi_opts) + + gin_env_xflags(gin_env) + + [ + exe, + "-b", + size, + "-e", + size, + "-f", + "2", + "-g", + "1", + "-R", + "2", + "-D", + "3", + "-A", + "1", + "-V", + ctas, + "-d", + dtype, + "-c", + "1", + "-w", + "1", + "-n", + "3", + ] + ) + cmd = " ".join(shlex.quote(a) for a in args) + hang_msg = ( + "Broadcast GIN-SDMA HANG: no completion within {}s at msg={} bytes " + "({} MiB), dtype={}. Output tail:\n{{}}".format( + timeout_s, size, msg_bytes // MiB, dtype + ) + ) + try: + rc, out = launch_mpi_shell(cmd, timeout_s, hang_msg) + debug = read_nccl_debug_logs(debug_dir) + finally: + shutil.rmtree(debug_dir, ignore_errors=True) + return rc, out, debug + + +def run_bcast_with_conn_gate_retry(launch_fn, conn_retries, settle_s=3): + """Connectivity-gate retry for Broadcast launches that return debug text too.""" + state = {"debug": ""} + + def attempt(): + rc, out, debug = launch_fn() + state["debug"] = debug + return rc, out + + rc, out = run_with_conn_gate_retry( + attempt, conn_retries, settle_s=settle_s, is_data_failure=bcast_data_failed + ) + return rc, out, state["debug"] diff --git a/projects/rccl-tests/test/test_AllGather.py b/projects/rccl-tests/test/test_AllGather.py index 7f958d33cd34..98896811f4f8 100644 --- a/projects/rccl-tests/test/test_AllGather.py +++ b/projects/rccl-tests/test/test_AllGather.py @@ -22,14 +22,22 @@ import os import re import shlex -import signal import subprocess -import time import itertools import math import pytest +from gin_sdma_harness import ( + GiB, + MiB, + env_int, + gin_env_xflags, + launch_mpi_shell, + mpi_launch_prefix, + run_with_conn_gate_retry, +) + ngpus = 0 if os.environ.get('ROCR_VISIBLE_DEVICES') is not None: ngpus = len(os.environ['ROCR_VISIBLE_DEVICES'].split(",")) @@ -126,16 +134,13 @@ def test_AllGatherSingleProcess(nthreads, ngpus_single, byte_range, op, step_fac # 256/512 MiB per-rank (2/4 segments) and 2 GiB total all pass with #wrong=0, # no hang (busbw ~421-427 GB/s). -MiB = 1024 * 1024 -GiB = 1024 * MiB - _ag_enabled = os.environ.get("RCCL_TESTS_GIN_SDMA_AG", "") not in ("", "0", "false", "False") -AG_NP = int(os.environ.get("RCCL_TESTS_AG_NP", "0")) or ngpus +AG_NP = env_int("RCCL_TESTS_AG_NP", 0) or ngpus AG_LAUNCHER = os.environ.get("RCCL_TESTS_MPI_LAUNCHER", "mpirun") AG_CTAS = os.environ.get("RCCL_TESTS_AG_CTAS", "8") -AG_TIMEOUT_S = int(os.environ.get("RCCL_TESTS_AG_TIMEOUT_S", "900")) -AG_CONN_RETRIES = int(os.environ.get("RCCL_TESTS_AG_CONN_RETRIES", "5")) +AG_TIMEOUT_S = env_int("RCCL_TESTS_AG_TIMEOUT_S", 900) +AG_CONN_RETRIES = env_int("RCCL_TESTS_AG_CONN_RETRIES", 5) AG_MPI_OPTS = shlex.split(os.environ.get("RCCL_TESTS_MPI_OPTS", "")) AG_XENV = shlex.split(os.environ.get("RCCL_TESTS_AG_XENV", "")) AG_EXE = os.environ.get( @@ -147,98 +152,70 @@ def test_AllGatherSingleProcess(nthreads, ngpus_single, byte_range, op, step_fac "a GIN-SDMA-capable (e.g. 8x MI355X) node to enable.") -# Intermittent gfx950 cuMem-VMM connectivity-gate abort (not a data error): the -# LSA signal connectivity self-test fails to map a peer aperture and the run -# aborts before any collective work. Re-launched after a short settle, exactly -# as the shell harness does. A genuine wrong-element data check is never retried. -_CONN_GATE_RE = re.compile(r"LSA signal connectivity gate failed|unhandled system error", re.I) -_DATA_FAIL_RE = re.compile( - r"#wrong\s*=\s*[1-9]|mismatch|check.*fail|Out of bounds values\s*:\s*[1-9]", re.I) +# Intermittent gfx950 cuMem-VMM connectivity-gate abort: see gin_sdma_harness.py. _AG_DEVTIME_TIER_RE = re.compile(r"#\[ag-devtime\].*tier\s+(LSA|SDMA)", re.I) def _launch_ag_gin_sdma(request, total_bytes, dtype, force_sdma_tier=True, device_timing=False): - """Launch all_gather_perf -D 3 once at a fixed TOTAL gathered size (per-rank - chunk = total/NP). When force_sdma_tier is True (default), threshold=0 pins the - SDMA tier for every size. Returns (returncode, stdout). A timeout (hang) fails - the test immediately.""" + """Launch all_gather_perf -D 3 once at a fixed TOTAL gathered size.""" size = str(int(total_bytes)) - gin_env = [] - for kv in ["NCCL_GIN_ENABLE=1", "NCCL_GIN_TYPE=6"] + AG_XENV: - gin_env += ["-x", kv] + gin_kv = ["NCCL_GIN_ENABLE=1", "NCCL_GIN_TYPE=6"] + AG_XENV if force_sdma_tier: - for kv in ["NCCL_GIN_ANVIL_SDMA_THRESHOLD=0", - "NCCL_GIN_ANVIL_SDMA_THRESHOLD_ALLGATHER=0"]: - gin_env += ["-x", kv] - - hostfile = request.config.getoption("--hostfile") - launch = [AG_LAUNCHER, "-np", str(AG_NP)] + AG_MPI_OPTS - if hostfile: - launch += ["-host", hostfile] + gin_kv += [ + "NCCL_GIN_ANVIL_SDMA_THRESHOLD=0", + "NCCL_GIN_ANVIL_SDMA_THRESHOLD_ALLGATHER=0", + ] - args = launch + gin_env + [ - AG_EXE, - "-b", size, "-e", size, - "-f", "2", - "-g", "1", - "-R", "2", - "-D", "3", # GinHybridAllGatherKernel - "-A", "1", # emit algo/proto/nchannels columns - "-V", AG_CTAS, - "-d", dtype, - "-c", "1", # data check: nonzero exit on any wrong element - "-w", "1", - "-n", "3", - ] + args = ( + mpi_launch_prefix(request, AG_LAUNCHER, AG_NP, AG_MPI_OPTS) + + gin_env_xflags(gin_kv) + + [ + AG_EXE, + "-b", + size, + "-e", + size, + "-f", + "2", + "-g", + "1", + "-R", + "2", + "-D", + "3", + "-A", + "1", + "-V", + AG_CTAS, + "-d", + dtype, + "-c", + "1", + "-w", + "1", + "-n", + "3", + ] + ) if device_timing: - args += ["-B", "1"] # augment stdout with #[ag-devtime] tier line + args += ["-B", "1"] cmd = " ".join(shlex.quote(a) for a in args) - print(cmd) - # New session so a hang can be killed as a whole process group -- otherwise a - # wedged launcher leaves orphaned ranks pinning the GPUs (and exhausting SDMA - # queues for later tests). - proc = subprocess.Popen(cmd, shell=True, universal_newlines=True, - stdout=subprocess.PIPE, stderr=subprocess.STDOUT, - start_new_session=True) - try: - out, _ = proc.communicate(timeout=AG_TIMEOUT_S) - except subprocess.TimeoutExpired: - try: - os.killpg(os.getpgid(proc.pid), signal.SIGKILL) - except ProcessLookupError: - pass - out, _ = proc.communicate() - pytest.fail( - "AllGather GIN-SDMA HANG: no completion within {}s at total={} bytes " - "({} MiB/rank), dtype={}. Output tail:\n{}".format( - AG_TIMEOUT_S, size, total_bytes // AG_NP // MiB, dtype, (out or "")[-2000:])) - print(out) - return proc.returncode, out + hang_msg = ( + "AllGather GIN-SDMA HANG: no completion within {}s at total={} bytes " + "({} MiB/rank), dtype={}. Output tail:\n{{}}".format( + AG_TIMEOUT_S, size, total_bytes // AG_NP // MiB, dtype + ) + ) + return launch_mpi_shell(cmd, AG_TIMEOUT_S, hang_msg) def _run_ag_gin_sdma(request, total_bytes, dtype, **launch_kw): - """Launch with connectivity-gate retry. Returns (returncode, stdout) of the - first run that either succeeds, hits a genuine data-check failure, or fails - for a non-connectivity-gate reason; otherwise the last attempt.""" if AG_NP < 2: pytest.skip("need >= 2 ranks/GPUs for GIN-SDMA AllGather") - - rc, out = 1, "" - for attempt in range(1, max(1, AG_CONN_RETRIES) + 1): - rc, out = _launch_ag_gin_sdma(request, total_bytes, dtype, **launch_kw) - if rc == 0: - return rc, out - # A real wrong-element mismatch must fail now, never retried. - if _DATA_FAIL_RE.search(out or ""): - return rc, out - # Intermittent gfx950 cuMem-VMM connectivity-gate abort: settle + retry. - if _CONN_GATE_RE.search(out or "") and attempt < AG_CONN_RETRIES: - print("=== connectivity-gate abort (attempt {}/{}); re-launching after settle ===".format( - attempt, AG_CONN_RETRIES)) - time.sleep(3) - continue - return rc, out - return rc, out + return run_with_conn_gate_retry( + lambda: _launch_ag_gin_sdma(request, total_bytes, dtype, **launch_kw), + AG_CONN_RETRIES, + ) # (items 1 + 2) Multi-segment loop and the 1 GiB single-copy boundary. per_rank diff --git a/projects/rccl-tests/test/test_AllToAll.py b/projects/rccl-tests/test/test_AllToAll.py index 094a119d5d76..73afa06f292a 100644 --- a/projects/rccl-tests/test/test_AllToAll.py +++ b/projects/rccl-tests/test/test_AllToAll.py @@ -46,9 +46,9 @@ # RCCL_TESTS_MPI_LAUNCHER launcher binary (default: mpirun) # RCCL_TESTS_MPI_OPTS extra launcher opts (e.g. --allow-run-as-root -mca ...) # RCCL_TESTS_A2A_XENV extra "-x K=V" env the backend needs on this cluster -# (e.g. HSA_FORCE_FINE_GRAIN_PCIE=1 NCCL_CUMEM_ENABLE=1) # RCCL_TESTS_A2A_EXE path to alltoall_perf (default: ../build/alltoall_perf) # RCCL_TESTS_A2A_TIMEOUT_S per-run hang timeout in seconds (default: 900) +# RCCL_TESTS_A2A_CONN_RETRIES connectivity-gate retries (default: 3) # # Verified on 8x MI355X (ROCm 7.13, NCCL_GIN_TYPE=6, force1ch): 256 MiB/peer and # 2 GiB/peer int32 and the 2 GiB-total guard all pass with #wrong=0, no hang @@ -56,128 +56,120 @@ import os import shlex -import signal -import subprocess import pytest -MiB = 1024 * 1024 -GiB = 1024 * MiB +from gin_sdma_harness import ( + GiB, + MiB, + detect_ngpus, + env_int, + gin_env_xflags, + launch_mpi_shell, + mpi_launch_prefix, + run_with_conn_gate_retry, +) path = os.path.dirname(os.path.abspath(__file__)) -# Path to alltoall_perf. Override with RCCL_TESTS_A2A_EXE when the binary is not -# at the default build/ location (e.g. an installed/container path). executable = os.environ.get( - "RCCL_TESTS_A2A_EXE", os.path.join(path, "..", "build", "alltoall_perf")) + "RCCL_TESTS_A2A_EXE", os.path.join(path, "..", "build", "alltoall_perf") +) -# Opt-in: only run where the GIN-SDMA backend + hardware are available. -_enabled = os.environ.get("RCCL_TESTS_GIN_SDMA_A2A", "") not in ("", "0", "false", "False") +_enabled = os.environ.get("RCCL_TESTS_GIN_SDMA_A2A", "") not in ( + "", + "0", + "false", + "False", +) - -def _detect_ngpus(): - if os.environ.get("ROCR_VISIBLE_DEVICES") is not None: - return len(os.environ["ROCR_VISIBLE_DEVICES"].split(",")) - if os.environ.get("HIP_VISIBLE_DEVICES") is not None: - return len(os.environ["HIP_VISIBLE_DEVICES"].split(",")) - try: - out = subprocess.check_output( - 'rocminfo | grep "Device Type:.\\s*.GPU" | wc -l', shell=True) - return int(out) - except Exception: - return 0 - - -NP = int(os.environ.get("RCCL_TESTS_A2A_NP", "0")) or _detect_ngpus() +NP = env_int("RCCL_TESTS_A2A_NP", 0) or detect_ngpus() LAUNCHER = os.environ.get("RCCL_TESTS_MPI_LAUNCHER", "mpirun") CTAS = os.environ.get("RCCL_TESTS_A2A_CTAS", "16") -# Generous per-call timeout; a real hang spins forever, so a finite cap is what -# converts item (3) into a detectable failure. Scaled up for the multi-GiB case. -TIMEOUT_S = int(os.environ.get("RCCL_TESTS_A2A_TIMEOUT_S", "900")) - -# Extra launcher options (e.g. "--allow-run-as-root -mca pml ob1 ...") and any -# deployment-specific env the GIN-SDMA backend needs beyond the essentials below -# (e.g. HSA_FORCE_FINE_GRAIN_PCIE=1 NCCL_CUMEM_ENABLE=1 RCCL_ENABLE_INTRANET=1). -# GIN enablement itself is site-specific, so it is injected here rather than -# hard-coded, keeping the test portable across clusters. +TIMEOUT_S = env_int("RCCL_TESTS_A2A_TIMEOUT_S", 900) +CONN_RETRIES = env_int("RCCL_TESTS_A2A_CONN_RETRIES", 3) MPI_OPTS = shlex.split(os.environ.get("RCCL_TESTS_MPI_OPTS", "")) XENV = shlex.split(os.environ.get("RCCL_TESTS_A2A_XENV", "")) pytestmark = pytest.mark.skipif( not _enabled, reason="GIN-SDMA AllToAll tests are opt-in; set RCCL_TESTS_GIN_SDMA_A2A=1 on " - "a GIN-SDMA-capable (e.g. 8x MI355X) node to enable.") + "a GIN-SDMA-capable (e.g. 8x MI355X) node to enable.", +) + + +def _launch_a2a(request, total_bytes, dtype): + """Launch alltoall_perf -D 3 once at a fixed total (per-rank) size.""" + size = str(int(total_bytes)) + gin_env = gin_env_xflags( + [ + "NCCL_GIN_ENABLE=1", + "NCCL_GIN_TYPE=6", + "NCCL_GIN_ANVIL_SDMA_THRESHOLD=0", + "NCCL_GIN_ANVIL_SDMA_THRESHOLD_ALLTOALL=0", + ] + + XENV + ) + + args = ( + mpi_launch_prefix(request, LAUNCHER, NP, MPI_OPTS) + + gin_env + + [ + executable, + "-b", + size, + "-e", + size, + "-f", + "2", + "-g", + "1", + "-R", + "2", + "-D", + "3", + "-V", + CTAS, + "-d", + dtype, + "-c", + "1", + "-w", + "1", + "-n", + "3", + ] + ) + cmd = " ".join(shlex.quote(a) for a in args) + hang_msg = ( + "AllToAll GIN-SDMA HANG: no completion within {}s at total={} bytes " + "({} MiB/peer), dtype={}. Output tail:\n{{}}".format( + TIMEOUT_S, size, total_bytes // NP // MiB, dtype + ) + ) + return launch_mpi_shell(cmd, TIMEOUT_S, hang_msg) def _run_a2a(request, total_bytes, dtype): - """Run one all-to-all at a fixed total (per-rank) size, forcing the GIN-SDMA - tier. Returns (returncode, stdout). A timeout (hang) fails the test.""" + """Run one all-to-all with connectivity-gate retry on gfx950 VMM aborts.""" if NP < 2: pytest.skip("need >= 2 ranks/GPUs for AllToAll") + return run_with_conn_gate_retry( + lambda: _launch_a2a(request, total_bytes, dtype), CONN_RETRIES + ) - size = str(int(total_bytes)) - # Essentials to exercise the GIN-SDMA put path; force the SDMA (large) tier - # for every size. Deployment-specific env comes from RCCL_TESTS_A2A_XENV. - gin_env = [] - for kv in ["NCCL_GIN_ENABLE=1", "NCCL_GIN_TYPE=6", - "NCCL_GIN_ANVIL_SDMA_THRESHOLD=0", - "NCCL_GIN_ANVIL_SDMA_THRESHOLD_ALLTOALL=0"] + XENV: - gin_env += ["-x", kv] - - hostfile = request.config.getoption("--hostfile") - launch = [LAUNCHER, "-np", str(NP)] + MPI_OPTS - if hostfile: - launch += ["-host", hostfile] - - args = launch + gin_env + [ - executable, - "-b", size, "-e", size, - "-f", "2", - "-g", "1", - "-R", "2", - "-D", "3", - "-V", CTAS, - "-d", dtype, - "-c", "1", # data check: nonzero exit on any wrong element - "-w", "1", - "-n", "3", - ] - cmd = " ".join(shlex.quote(a) for a in args) - print(cmd) - # Run in a new session so a hang can be killed as a whole process group -- - # otherwise a wedged launcher leaves orphaned ranks pinning the GPUs (and - # exhausting SDMA queues for later tests). - proc = subprocess.Popen(cmd, shell=True, universal_newlines=True, - stdout=subprocess.PIPE, stderr=subprocess.STDOUT, - start_new_session=True) - try: - out, _ = proc.communicate(timeout=TIMEOUT_S) - except subprocess.TimeoutExpired: - try: - os.killpg(os.getpgid(proc.pid), signal.SIGKILL) - except ProcessLookupError: - pass - out, _ = proc.communicate() - pytest.fail( - "AllToAll GIN-SDMA HANG: no completion within {}s at total={} bytes " - "({} MiB/peer), dtype={}. Output tail:\n{}".format( - TIMEOUT_S, size, total_bytes // NP // MiB, dtype, (out or "")[-2000:])) - print(out) - return proc.returncode, out - - -# (item 1 + 2) Multi-segment loop and the 1 GiB truncation boundary. per_peer is -# the transfer to each peer; total per-rank bytes = per_peer * NP. -@pytest.mark.parametrize("per_peer_mib", [256, 2048]) # 256 MiB (2 seg), 2 GiB (16 seg) + +@pytest.mark.parametrize("per_peer_mib", [256, 2048]) @pytest.mark.parametrize("dtype", ["int32", "int64", "uint8"]) def test_AllToAllGinSdmaLargeSegmented(request, per_peer_mib, dtype): total = per_peer_mib * MiB * NP rc, _ = _run_a2a(request, total, dtype) - assert rc == 0, "AllToAll data check failed (nonzero exit) at {} MiB/peer, dtype={}".format( - per_peer_mib, dtype) + assert rc == 0, ( + "AllToAll data check failed (nonzero exit) at {} MiB/peer, dtype={}".format( + per_peer_mib, dtype + ) + ) -# (item 3) The original 2 GiB-total hang repro, as a completion guard. At 8 ranks -# this is 256 MiB/peer -- the size that hung as a single unclamped put. def test_AllToAllGinSdma2GiBTotalHangGuard(request): rc, _ = _run_a2a(request, 2 * GiB, "int32") assert rc == 0, "AllToAll 2 GiB-total data check failed (nonzero exit)" diff --git a/projects/rccl-tests/test/test_Broadcast.py b/projects/rccl-tests/test/test_Broadcast.py index 6ed5657663a4..28639052cad3 100644 --- a/projects/rccl-tests/test/test_Broadcast.py +++ b/projects/rccl-tests/test/test_Broadcast.py @@ -20,19 +20,25 @@ ################################################################################ import os -import re -import glob import shlex -import shutil -import signal import subprocess import itertools import math -import tempfile -import time import pytest +from gin_sdma_harness import ( + GiB, + MiB, + assert_bcast_ok, + bcast_data_failed, + bcast_rows, + bcast_tiers, + env_int, + launch_bcast_gin_sdma, + run_bcast_with_conn_gate_retry, +) + ngpus = 0 if os.environ.get("ROCR_VISIBLE_DEVICES") is not None: ngpus = len(os.environ["ROCR_VISIBLE_DEVICES"].split(",")) @@ -132,9 +138,6 @@ def test_BroadcastSingleProcess( # is 6 on develop but 5 on the NCCL 2.30.7 line, so this must be settable # rather than baked in -- a wrong value silently loads no GIN plugin. -MiB = 1024 * 1024 -GiB = 1024 * MiB - _bcast_enabled = os.environ.get("RCCL_TESTS_GIN_SDMA_BCAST", "") not in ( "", "0", @@ -142,28 +145,12 @@ def test_BroadcastSingleProcess( "False", ) - -def _env_int(name, default): - """Parse an integer env var; bad values fall back so collection stays usable.""" - raw = os.environ.get(name) - if raw is None or raw == "": - return default - try: - return int(raw) - except ValueError: - return default - - -BCAST_NP = _env_int("RCCL_TESTS_BCAST_NP", 0) or ngpus +BCAST_NP = env_int("RCCL_TESTS_BCAST_NP", 0) or ngpus BCAST_LAUNCHER = os.environ.get("RCCL_TESTS_MPI_LAUNCHER", "mpirun") BCAST_CTAS = os.environ.get("RCCL_TESTS_BCAST_CTAS", "8") BCAST_GIN_TYPE = os.environ.get("RCCL_TESTS_BCAST_GIN_TYPE", "6") -# Manual / SUT defaults. GIN CI (run-gin-ci.sh) overrides both so -# HW_CASES * retries * TIMEOUT_S stays under GIN_PYTEST_TIMEOUT; otherwise -# GNU timeout kills pytest and leaves the mpirun session (start_new_session) -# orphaned because killpg only runs in this TimeoutExpired handler. -BCAST_TIMEOUT_S = _env_int("RCCL_TESTS_BCAST_TIMEOUT_S", 900) -BCAST_CONN_RETRIES = _env_int("RCCL_TESTS_BCAST_CONN_RETRIES", 5) +BCAST_TIMEOUT_S = env_int("RCCL_TESTS_BCAST_TIMEOUT_S", 900) +BCAST_CONN_RETRIES = env_int("RCCL_TESTS_BCAST_CONN_RETRIES", 5) BCAST_MPI_OPTS = shlex.split(os.environ.get("RCCL_TESTS_MPI_OPTS", "")) BCAST_XENV = shlex.split(os.environ.get("RCCL_TESTS_BCAST_XENV", "")) BCAST_EXE = os.environ.get( @@ -176,267 +163,18 @@ def _env_int(name, default): "a GIN-SDMA-capable (e.g. 8x MI355X) node to enable.", ) -_CONN_GATE_RE = re.compile( - r"LSA signal connectivity gate failed|unhandled system error", re.I -) - -# The GIN Anvil-SDMA backend has to actually bind for any of this to mean -# something: a run that falls back to another transport still exits 0 and still -# prints a full results table. NCCL_DEBUG_SUBSYS=INIT,NET below is what puts the -# bind line in the captured output. -_PLUGIN_RE = re.compile(r"gin-anvil-sdma", re.I) - -# One measured rccl-tests results row: -# size count type redop root | time algbw busbw #wrong | time algbw busbw #wrong -# out-of-place columns first, then in-place. Timings come from getFloatStr, which -# falls back to scientific notation when a value will not fit its width, so the -# numeric fields have to admit exponents. The row is not end-anchored: an algo / -# proto / nchannels group and a timestamp may follow the in-place columns, and -# concurrent rank output can splice itself onto the end of the line. -_NUM = r"[-+]?(?:\d+\.?\d*|\.\d+)(?:[eE][-+]?\d+)?" -_WRONG = r"(?:{}|N/A)".format(_NUM) -_ROW_RE = re.compile( - r"^\s*(\d+)\s+(\d+)\s+\S+\s+\S+\s+(\d+)" - r"\s+{n}\s+{n}\s+{n}\s+({w})" - r"\s+{n}\s+{n}\s+{n}\s+({w})".format(n=_NUM, w=_WRONG) -) -_OOB_RE = re.compile(r"Out of bounds values\s*:\s*(\d+)") - -# "#[bcast-tier] ", emitted once per distinct tier by BroadcastRunColl when -# RCCL_TESTS_BCAST_TIER_LOG is set. Named tiers are the contract between the -D 3 -# dispatch and the tier assertions below; see bcastReportTier in broadcast.cu. -_TIER_RE = re.compile(r"^#\[bcast-tier\]\s+(\S+)", re.M) - - -def _bcast_tiers(out): - """Set of -D 3 tier names the run reported.""" - return set(_TIER_RE.findall(out or "")) - - -def _bcast_rows(out): - """Measured rows as (size, root, wrong_out_of_place, wrong_in_place).""" - rows = [] - for line in (out or "").splitlines(): - m = _ROW_RE.match(line) - if m: - rows.append((int(m.group(1)), int(m.group(3)), m.group(4), m.group(5))) - return rows - - -def _wrong_count(field): - """Numeric #wrong, or None when the column is N/A because checks were off.""" - try: - return float(field) - except ValueError: - return None - - -def _data_failed(out): - """True when the run produced demonstrably wrong results. - - Deliberately not a text match on "Wrong": rccl-tests prints a "#wrong" column - header on every run that reaches the results table, so a bare pattern matched - every healthy run and made the connectivity retry below unreachable. N/A is - not a failure here -- it means checking was off, which _assert_bcast_ok - rejects separately rather than burning retries on. - """ - for _, _, oop, ip in _bcast_rows(out): - if any(_wrong_count(f) not in (None, 0.0) for f in (oop, ip)): - return True - m = _OOB_RE.search(out or "") - return bool(m and m.group(1) != "0") - - -def _read_debug_logs(debug_dir): - """Merged per-rank NCCL debug output written under debug_dir.""" - chunks = [] - for log_path in sorted(glob.glob(os.path.join(debug_dir, "nccl-debug.*"))): - try: - with open(log_path, errors="replace") as fh: - chunks.append(fh.read()) - except OSError: - pass - return "\n".join(chunks) - - -def _assert_bcast_ok(rc, out, debug, ctx, expect_tier=None): - """Assert the run bound GIN-SDMA, measured something, and checked clean. - - Exit status alone is not enough: it stays 0 for a run that binds a different - backend, and for one that produces no measured rows at all. - - Nor is the data check enough to say which tier ran. Every -D 3 tier produces - a correct broadcast, so #wrong stays 0 when a tier is deleted and its - messages fall through to the next one -- the test keeps passing while the - kernel it named is no longer reached. expect_tier pins the tier that has to - have run, so removing it fails here instead of silently going uncovered. - """ - tail = (out or "")[-2000:] - assert rc == 0, "Broadcast {} failed (exit {}). Output tail:\n{}".format( - ctx, rc, tail - ) - assert _PLUGIN_RE.search(debug or "") or _PLUGIN_RE.search(out or ""), ( - "Broadcast {} never bound the GIN Anvil-SDMA backend, so the run proves " - "nothing about the GIN path. Output tail:\n{}".format(ctx, tail) - ) - rows = _bcast_rows(out) - assert rows, "Broadcast {} produced no measured rows. Output tail:\n{}".format( - ctx, tail - ) - unchecked = [ - (s, r) - for s, r, oop, ip in rows - if _wrong_count(oop) is None or _wrong_count(ip) is None - ] - assert not unchecked, ( - "Broadcast {} reported #wrong as N/A at (size, root) {}, so the data check " - "never ran and the result is unverified.".format(ctx, unchecked) - ) - bad = [ - (s, r, oop, ip) - for s, r, oop, ip in rows - if _wrong_count(oop) != 0.0 or _wrong_count(ip) != 0.0 - ] - assert ( - not bad - ), "Broadcast {} reported nonzero #wrong (size, root, oop, ip): {}".format(ctx, bad) - m = _OOB_RE.search(out or "") - assert m and m.group(1) == "0", "Broadcast {} out-of-bounds count is {}".format( - ctx, m.group(1) if m else "absent" - ) - if expect_tier is not None: - # A string pins one tier; a collection allows a set of equivalent ones. - # The ring tier needs the latter: it reports ring-table where an - # edge-disjoint decomposition exists and ring-multi where none does - # (N=4 and N=6), and BCAST_NP follows whatever GPUs the host has. - # - # The ring also requires the LSA team to cover the world, so a multi-node - # launch takes scatter+allgather instead and fails here. That is the - # intent: such a run is not exercising the ring the caller asked for. - allowed = ( - {expect_tier} if isinstance(expect_tier, str) else set(expect_tier) - ) - tiers = _bcast_tiers(out) - assert tiers, ( - "Broadcast {} reported no #[bcast-tier] line, so which tier ran is " - "unverified. Is RCCL_TESTS_BCAST_TIER_LOG set for this launch, and is " - "the binary new enough to emit it? Output tail:\n{}".format(ctx, tail) - ) - # Exactly one tier, and an allowed one. These launches pin a single size, - # so a second tier means some roots took a different kernel; merely - # intersecting `allowed` would let that through. - assert len(tiers) == 1 and tiers <= allowed, ( - "Broadcast {} ran tier(s) {} but expected exactly one of {}. The tier " - "gate moved, or the tier under test fell through to another " - "kernel.".format(ctx, sorted(tiers), sorted(allowed)) - ) - -def _launch_bcast_gin_sdma( - request, msg_bytes, dtype, *, force_flat_gin=False, force_sag_gin=False -): - """Launch broadcast_perf -D 3 at a fixed message size. When force_flat_gin, - disable the SAG/ring large tiers so the root flat gin.put() path is used. - When force_sag_gin, disable only the ring tier so scatter+allgather runs. - - Returns (returncode, stdout, debug_text); debug_text is the merged per-rank - NCCL debug log, kept off stdout so the results table stays parseable.""" - size = str(int(msg_bytes)) - debug_dir = tempfile.mkdtemp(prefix="bcast-gin-dbg-") - gin_env = [ - "NCCL_GIN_ENABLE=1", - "NCCL_GIN_TYPE={}".format(BCAST_GIN_TYPE), - "NCCL_GIN_ANVIL_SDMA_THRESHOLD=0", - "NCCL_GIN_ANVIL_SDMA_THRESHOLD_BROADCAST=0", - # Debug output is what lets _PLUGIN_RE confirm which backend actually - # bound. It goes to per-rank files rather than stdout because ranks - # interleave: on stdout an INFO line splices itself into the middle of a - # results row, which corrupts the very table we need to parse. - "NCCL_DEBUG=INFO", - "NCCL_DEBUG_SUBSYS=INIT,NET", - "NCCL_DEBUG_FILE={}/nccl-debug.%h.%p.log".format(debug_dir), - # Makes BroadcastRunColl name the tier it launched, which is what lets - # _assert_bcast_ok tell the -D 3 tiers apart. One line from the main - # proc, so the results table stays parseable. - "RCCL_TESTS_BCAST_TIER_LOG=1", - ] - if force_flat_gin: - gin_env += [ - "NCCL_GIN_ANVIL_BCAST_SCATTER_AG_MIN_BYTES=0", - "NCCL_GIN_ANVIL_BCAST_RING_MIN_BYTES=0", - ] - elif force_sag_gin: - # Ring is checked before SAG; disable ring only so SAG is exercised. - gin_env += ["NCCL_GIN_ANVIL_BCAST_RING_MIN_BYTES=0"] - gin_env += BCAST_XENV - gin_env = sum([["-x", kv] for kv in gin_env], []) - - hostfile = request.config.getoption("--hostfile") - launch = [BCAST_LAUNCHER, "-np", str(BCAST_NP)] + BCAST_MPI_OPTS - if hostfile: - launch += ["-host", hostfile] - - args = ( - launch - + gin_env - + [ - BCAST_EXE, - "-b", - size, - "-e", - size, - "-f", - "2", - "-g", - "1", - "-R", - "2", - "-D", - "3", # GinHybridBroadcastKernel - "-A", - "1", - "-V", - BCAST_CTAS, - "-d", - dtype, - "-c", - "1", - "-w", - "1", - "-n", - "3", - ] - ) - cmd = " ".join(shlex.quote(a) for a in args) - print(cmd) - try: - proc = subprocess.Popen( - cmd, - shell=True, - universal_newlines=True, - stdout=subprocess.PIPE, - stderr=subprocess.STDOUT, - start_new_session=True, - ) - try: - out, _ = proc.communicate(timeout=BCAST_TIMEOUT_S) - except subprocess.TimeoutExpired: - try: - os.killpg(os.getpgid(proc.pid), signal.SIGKILL) - except ProcessLookupError: - pass - out, _ = proc.communicate() - pytest.fail( - "Broadcast GIN-SDMA HANG: no completion within {}s at msg={} bytes " - "({} MiB), dtype={}. Output tail:\n{}".format( - BCAST_TIMEOUT_S, size, msg_bytes // MiB, dtype, (out or "")[-2000:] - ) - ) - debug = _read_debug_logs(debug_dir) - finally: - shutil.rmtree(debug_dir, ignore_errors=True) - print(out) - return proc.returncode, out, debug +def _bcast_launch_kwargs(): + return { + "exe": BCAST_EXE, + "np": BCAST_NP, + "launcher": BCAST_LAUNCHER, + "mpi_opts": BCAST_MPI_OPTS, + "xenv": BCAST_XENV, + "ctas": BCAST_CTAS, + "gin_type": BCAST_GIN_TYPE, + "timeout_s": BCAST_TIMEOUT_S, + } def _run_bcast_gin_sdma( @@ -444,30 +182,17 @@ def _run_bcast_gin_sdma( ): if BCAST_NP < 2: pytest.skip("need >= 2 ranks/GPUs for GIN-SDMA Broadcast") - - rc, out, debug = 1, "", "" - for attempt in range(1, max(1, BCAST_CONN_RETRIES) + 1): - rc, out, debug = _launch_bcast_gin_sdma( + return run_bcast_with_conn_gate_retry( + lambda: launch_bcast_gin_sdma( request, msg_bytes, dtype, force_flat_gin=force_flat_gin, force_sag_gin=force_sag_gin, - ) - if rc == 0: - return rc, out, debug - if _data_failed(out): - return rc, out, debug - if _CONN_GATE_RE.search(out or "") and attempt < BCAST_CONN_RETRIES: - print( - "=== connectivity-gate abort (attempt {}/{}); re-launching after settle ===".format( - attempt, BCAST_CONN_RETRIES - ) - ) - time.sleep(3) - continue - return rc, out, debug - return rc, out, debug + **_bcast_launch_kwargs(), + ), + BCAST_CONN_RETRIES, + ) @_bcast_skip @@ -478,7 +203,7 @@ def test_BroadcastGinSdmaLargeSegmented(request, msg_mib, dtype): rc, out, debug = _run_bcast_gin_sdma( request, msg_mib * MiB, dtype, force_flat_gin=True ) - _assert_bcast_ok( + assert_bcast_ok( rc, out, debug, @@ -493,7 +218,7 @@ def test_BroadcastGinSdmaScatterAllgather(request): rc, out, debug = _run_bcast_gin_sdma( request, 256 * MiB, "int32", force_sag_gin=True ) - _assert_bcast_ok( + assert_bcast_ok( rc, out, debug, "SAG 256 MiB", expect_tier="scatter-allgather" ) @@ -504,7 +229,7 @@ def test_BroadcastGinSdma2GiBHangGuard(request): rc, out, debug = _run_bcast_gin_sdma( request, 2 * GiB, "int32", force_flat_gin=False ) - _assert_bcast_ok( + assert_bcast_ok( rc, out, debug, @@ -519,7 +244,7 @@ def test_BroadcastGinSdma4GiBHangGuard(request): rc, out, debug = _run_bcast_gin_sdma( request, 4 * GiB, "int32", force_flat_gin=False ) - _assert_bcast_ok( + assert_bcast_ok( rc, out, debug, @@ -539,7 +264,7 @@ def test_BroadcastGinSdmaRowRegexAcceptsScientificNotation(): " 1.23e+02 1.00e+02 1.00e+02 0" " 1.23e+02 1.00e+02 1.00e+02 0" ) - rows = _bcast_rows(line) + rows = bcast_rows(line) assert rows == [(134217728, 0, "0", "0")] @@ -548,7 +273,7 @@ def test_BroadcastGinSdmaDataFailedIgnoresColumnHeader(): "# size count type redop root" " time algbw busbw #wrong" ) - assert not _data_failed(header) + assert not bcast_data_failed(header) def test_BroadcastGinSdmaDataFailedDetectsNonzeroWrong(): @@ -557,7 +282,7 @@ def test_BroadcastGinSdmaDataFailedDetectsNonzeroWrong(): " 12.34 1.00 1.00 1" " 12.34 1.00 1.00 0" ) - assert _data_failed(line) + assert bcast_data_failed(line) def test_BroadcastGinSdmaDataFailedTreatsNaAsUncheckedNotFailed(): @@ -566,7 +291,7 @@ def test_BroadcastGinSdmaDataFailedTreatsNaAsUncheckedNotFailed(): " 12.34 1.00 1.00 N/A" " 12.34 1.00 1.00 0" ) - assert not _data_failed(line) + assert not bcast_data_failed(line) def _clean_bcast_output(tier): @@ -587,16 +312,16 @@ def _clean_bcast_output(tier): # expands its args unquoted, so a filter with a space in it is not available. def test_BroadcastGinSdmaTierGuardLineIsNotAResultsRow(): """The tier line shares stdout with the table it must not corrupt.""" - assert _bcast_rows("#[bcast-tier] scatter-allgather") == [] - assert _bcast_tiers(_clean_bcast_output("ring-table")) == {"ring-table"} - assert _bcast_rows(_clean_bcast_output("ring-table")) + assert bcast_rows("#[bcast-tier] scatter-allgather") == [] + assert bcast_tiers(_clean_bcast_output("ring-table")) == {"ring-table"} + assert bcast_rows(_clean_bcast_output("ring-table")) def test_BroadcastGinSdmaTierGuardAcceptsExpectedTier(): out = _clean_bcast_output("scatter-allgather") - _assert_bcast_ok(0, out, "gin-anvil-sdma", "guard", expect_tier="scatter-allgather") + assert_bcast_ok(0, out, "gin-anvil-sdma", "guard", expect_tier="scatter-allgather") # Ring passes against either decomposition outcome. - _assert_bcast_ok( + assert_bcast_ok( 0, _clean_bcast_output("ring-multi"), "gin-anvil-sdma", @@ -614,9 +339,9 @@ def test_BroadcastGinSdmaTierGuardRejectsFallThrough(): would go green while covering nothing. """ fell_through = _clean_bcast_output("flat-gin") - _assert_bcast_ok(0, fell_through, "gin-anvil-sdma", "guard", expect_tier="flat-gin") + assert_bcast_ok(0, fell_through, "gin-anvil-sdma", "guard", expect_tier="flat-gin") with pytest.raises(AssertionError, match="expected exactly one of"): - _assert_bcast_ok( + assert_bcast_ok( 0, fell_through, "gin-anvil-sdma", "guard", expect_tier="scatter-allgather" ) @@ -625,7 +350,7 @@ def test_BroadcastGinSdmaTierGuardRejectsMissingTierLine(): """A binary predating the tier log must fail loudly, not pass vacuously.""" no_tier = "\n".join(_clean_bcast_output("flat-gin").splitlines()[1:]) with pytest.raises(AssertionError, match="no .*bcast-tier.* line"): - _assert_bcast_ok( + assert_bcast_ok( 0, no_tier, "gin-anvil-sdma", "guard", expect_tier="flat-gin" ) From a7fabe15f3ff086e65d2364f24218a67d751159a Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sun, 13 Sep 2026 15:53:49 -0400 Subject: [PATCH 2/8] rccl-tests: GIN device-API CI and harness cleanup (AICOMRCCL-2221) Build device-api CI with ENABLE_DEVICE_API + ENABLE_ROCSHMEM_GIN via cmake, dedupe rocshmem RDC link blocks, move GPU AllGather gtest to gpu-only label, add AllGather policy header missing STATUS, and link rocSHMEM into FixturesDebug for GIN unit tests. Co-authored-by: Cursor --- projects/rccl-tests/src/CMakeLists.txt | 27 ++++---- projects/rccl/test/CMakeLists.txt | 34 +++++++++- projects/rccl/tools/ci/lib/device-api.sbatch | 71 +++++++++++++++----- projects/rccl/tools/ci/run-device-api-ci.sh | 19 ++++-- 4 files changed, 113 insertions(+), 38 deletions(-) diff --git a/projects/rccl-tests/src/CMakeLists.txt b/projects/rccl-tests/src/CMakeLists.txt index 3f671325cce2..cfd0e4f3facc 100644 --- a/projects/rccl-tests/src/CMakeLists.txt +++ b/projects/rccl-tests/src/CMakeLists.txt @@ -146,25 +146,26 @@ add_rccl_test(all_gather) add_rccl_test(all_reduce) add_rccl_test(alltoall) if (ENABLE_ROCSHMEM_GIN AND TARGET roc::rocshmem) - target_compile_definitions(alltoall_perf PRIVATE ENABLE_ROCSHMEM_GIN) - target_compile_options(alltoall_perf PRIVATE -fgpu-rdc $<$:-ggdb -O0>) - # --whole-archive: librccl.so (ENABLE_ROCSHMEM_GIN) leaves rocshmem host - # symbols unresolved (--allow-shlib-undefined). The executable must pull - # in ALL rocshmem symbols (not just those it references directly) so that - # -rdynamic exports them for librccl.so to resolve at runtime. - target_link_libraries(alltoall_perf PRIVATE - -Wl,--whole-archive roc::rocshmem -Wl,--no-whole-archive) - target_link_options(alltoall_perf PRIVATE -fgpu-rdc --hip-link -rdynamic) + function(_rccl_tests_link_rocshmem_gin_rdc target) + target_compile_definitions(${target} PRIVATE ENABLE_ROCSHMEM_GIN) + target_compile_options(${target} PRIVATE -fgpu-rdc $<$:-ggdb -O0>) + # --whole-archive: librccl.so (ENABLE_ROCSHMEM_GIN) leaves rocshmem host + # symbols unresolved (--allow-shlib-undefined). The executable must pull + # in ALL rocshmem symbols (not just those it references directly) so that + # -rdynamic exports them for librccl.so to resolve at runtime. + target_link_libraries(${target} PRIVATE + -Wl,--whole-archive roc::rocshmem -Wl,--no-whole-archive) + target_link_options(${target} PRIVATE -fgpu-rdc --hip-link -rdynamic) + endfunction() + + _rccl_tests_link_rocshmem_gin_rdc(alltoall_perf) # GIN-SDMA hybrid AllGather (all_gather_perf -D 3): same device-RDC link as # alltoall_perf so the GinHybridAllGatherKernel resolves the GIN device backend. if (ENABLE_DEVICE_API) target_compile_definitions(all_gather_perf PRIVATE ENABLE_DEVICE_API) endif() - target_compile_options(all_gather_perf PRIVATE -fgpu-rdc $<$:-ggdb -O0>) - target_link_libraries(all_gather_perf PRIVATE - -Wl,--whole-archive roc::rocshmem -Wl,--no-whole-archive) - target_link_options(all_gather_perf PRIVATE -fgpu-rdc --hip-link -rdynamic) + _rccl_tests_link_rocshmem_gin_rdc(all_gather_perf) endif() if (ENABLE_ROCSHMEM_GIN AND TARGET roc::rocshmem AND RCCL_SOURCE_DIR) add_executable(gin_rocshmem_gda_factory_test ${CMAKE_CURRENT_SOURCE_DIR}/gin_rocshmem_gda_factory_test.cu) diff --git a/projects/rccl/test/CMakeLists.txt b/projects/rccl/test/CMakeLists.txt index c300cde31c99..d156fa9bfa43 100644 --- a/projects/rccl/test/CMakeLists.txt +++ b/projects/rccl/test/CMakeLists.txt @@ -546,7 +546,8 @@ if(BUILD_TESTS) # direct-LSA gather addressing (the pointer-arithmetic core of # AllGatherLsaDirect) on real hardware, self-contained (no librccl / GIN # comm). Requires a visible GPU at run time (skips cleanly otherwise), so it - # carries a "gpu" label alongside "unit". + # carries the "gpu" label only so `ctest -L unit` on CPU nodes does not + # count an all-skip binary as a passing unit test. add_executable(rccl-UnitTestsGinSdmaAllGatherGpu gin/gin_sdma_allgather_gpu_test.cpp) set_source_files_properties(gin/gin_sdma_allgather_gpu_test.cpp PROPERTIES LANGUAGE CXX) target_include_directories(rccl-UnitTestsGinSdmaAllGatherGpu PRIVATE @@ -569,7 +570,10 @@ if(BUILD_TESTS) RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/test" CXX_STANDARD 17 CXX_STANDARD_REQUIRED ON) add_test(NAME rccl-UnitTestsGinSdmaAllGatherGpu COMMAND rccl-UnitTestsGinSdmaAllGatherGpu) - set_tests_properties(rccl-UnitTestsGinSdmaAllGatherGpu PROPERTIES LABELS "unit;gpu") + set_tests_properties(rccl-UnitTestsGinSdmaAllGatherGpu PROPERTIES LABELS "gpu") + else() + message(STATUS "Skipping GIN-SDMA AllGather policy/gpu unit tests: " + "${_gin_sdma_ag_hdr_dir}/gin_sdma_allgather_policy.h not found") endif() # GIN Anvil-SDMA hybrid Broadcast (-D 3) policy host unit test (no GPU). @@ -677,6 +681,29 @@ if(BUILD_TESTS) -Wl,--export-dynamic-symbol=ncclNetPlugin_v10 -Wl,--export-dynamic-symbol=ncclNetPlugin_v11 -Wl,--export-dynamic-symbol=ncclCollNetPlugin_v11) + + # Transitive nccl_device/gin headers pull sdma/anvil_device.hpp from rocSHMEM. + if(ENABLE_ROCSHMEM_GIN) + if(ROCSHMEM_SOURCE_DIR) + target_include_directories(rccl-UnitTestsFixturesDebug PRIVATE + "${ROCSHMEM_SOURCE_DIR}/src" + "${ROCSHMEM_SOURCE_DIR}/include") + else() + target_include_directories(rccl-UnitTestsFixturesDebug PRIVATE + ${CMAKE_SOURCE_DIR}/../rocshmem/src + ${CMAKE_SOURCE_DIR}/../rocshmem/include) + endif() + if(ROCSHMEM_INCLUDE_DIR) + target_include_directories(rccl-UnitTestsFixturesDebug PRIVATE ${ROCSHMEM_INCLUDE_DIR}) + elseif(ROCSHMEM_INSTALL_DIR) + target_include_directories(rccl-UnitTestsFixturesDebug PRIVATE "${ROCSHMEM_INSTALL_DIR}/include") + endif() + if(ROCSHMEM_BUILD_DIR) + target_include_directories(rccl-UnitTestsFixturesDebug PRIVATE ${ROCSHMEM_BUILD_DIR}) + elseif(ROCSHMEM_SOURCE_DIR AND EXISTS "${ROCSHMEM_SOURCE_DIR}/build/include") + target_include_directories(rccl-UnitTestsFixturesDebug PRIVATE "${ROCSHMEM_SOURCE_DIR}/build/include") + endif() + endif() endif() # rccl-UnitTestsAltRsmi: Compiles alt_rsmi.cc directly (not via librccl.so) with @@ -1039,7 +1066,8 @@ if(BUILD_TESTS) # into the Fixtures binary so the executable satisfies those symbols when # the dynamic loader brings in librccl.so. if(ENABLE_ROCSHMEM_GIN AND GIN_ANVIL_UNIT_TESTS AND NOT ENABLE_ROCSHMEM - AND ROCSHMEM_INSTALL_DIR AND "${test_executable}" STREQUAL "rccl-UnitTestsFixtures") + AND ROCSHMEM_INSTALL_DIR AND ("${test_executable}" STREQUAL "rccl-UnitTestsFixtures" + OR "${test_executable}" STREQUAL "rccl-UnitTestsFixturesDebug")) set(_ROCSHMEM_FIXTURES_OFFLOAD_FLAGS "") foreach(_g IN LISTS GPU_TARGETS) if(_g) diff --git a/projects/rccl/tools/ci/lib/device-api.sbatch b/projects/rccl/tools/ci/lib/device-api.sbatch index 81549ecd1883..7f9bb7f648a6 100755 --- a/projects/rccl/tools/ci/lib/device-api.sbatch +++ b/projects/rccl/tools/ci/lib/device-api.sbatch @@ -15,9 +15,10 @@ # On one 8-GPU node: # 1. consume ROCm (provisioned on the head node) via .ci-out/rocm.env # 2. lib/build-ompi.sh build OpenMPI (--with-rocm) -> ompi.env -# 3. projects/rccl/install.sh build RCCL (in-tree, release) -# 4. rccl-tests install.sh build rccl-tests (ENABLE_DEVICE_API=1) -# 5. run-device-api-ci.sh run the device-API benchmark suite +# 3. lib/build-rocshmem.sh build rocSHMEM -> rocshmem.env +# 4. projects/rccl/install.sh build RCCL (in-tree, release, --rocshmem-gin) +# 5. rccl-tests (cmake) build rccl-tests (ENABLE_DEVICE_API + ENABLE_ROCSHMEM_GIN) +# 6. run-device-api-ci.sh run the device-API benchmark suite # # Required env (from the workflow via `sbatch --export=ALL,...`): # RCCL_DEVICE_API_WORKDIR GHA workspace (rocm-systems root) @@ -108,28 +109,64 @@ echo "==> verified hipcc = ${resolved_hipcc}" # shellcheck source=/dev/null source "${ci_dir}/lib/ensure-python-yaml.sh" -# 3) RCCL via its own install.sh (-l = local gpu arch only). -echo "==> [3/5] build RCCL (targets=${amdgpu_targets})" -( cd "${workdir}/projects/rccl" \ +# 3) rocSHMEM (in-tree, IPC+SDMA) -> rocshmem.env (ROCSHMEM_INSTALL_DIR) +echo "==> [3/6] build rocSHMEM" +GPU_TARGETS="${amdgpu_targets}" "${ci_dir}/lib/build-rocshmem.sh" +set -a +# shellcheck source=/dev/null +source "${workdir}/.ci-out/rocshmem.env" +set +a + +# 4) RCCL with GIN / rocSHMEM (device-API + GIN-SDMA -D 3 paths need rocshmem in the executable). +echo "==> [4/6] build RCCL (GIN, targets=${amdgpu_targets})" +( + cd "${workdir}/projects/rccl" \ && MPI_PATH="${MPI_HOME}" ./install.sh -lj "${build_jobs}" \ - --amdgpu_targets="${amdgpu_targets}" ) + --amdgpu_targets="${amdgpu_targets}" \ + --rocshmem-gin --no_clean \ + --cmake-options "-DROCSHMEM_INSTALL_DIR=${ROCSHMEM_INSTALL_DIR}" +) rccl_prefix="${workdir}/projects/rccl/build/release" +rccl_cfg="$(ls "${rccl_prefix}"/lib*/cmake/rccl/*[Cc]onfig.cmake 2>/dev/null | head -n1)" +if [[ -z "${rccl_cfg}" ]]; then + echo "ERROR: no RCCL cmake config under ${rccl_prefix}/lib*/cmake/rccl" >&2 + ls -laR "${rccl_prefix}"/lib*/cmake 2>/dev/null >&2 || true + exit 1 +fi +rccl_cmake_dir="$(dirname "${rccl_cfg}")" +echo " in-tree RCCL cmake config: ${rccl_cmake_dir}" -# 4) rccl-tests: built in-tree from projects/rccl-tests (must match -# RCCL_TESTS_DIR in run-device-api-ci.sh). install.sh wipes its own build/. -echo "==> [4/5] build rccl-tests (ENABLE_DEVICE_API=1)" -rccl_tests_dir="${workdir}/projects/rccl-tests" -( cd "${rccl_tests_dir}" \ - && ENABLE_DEVICE_API=1 ./install.sh \ - --mpi --mpi_home="${MPI_HOME}" --rccl_home="${rccl_prefix}" \ - --gpu_targets="${amdgpu_targets}" ) +# 5) rccl-tests via cmake so all_gather_perf -D 3 links rocshmem (install.sh/Makefile +# path does not enable ENABLE_ROCSHMEM_GIN). RCCL_TESTS_BIN_DIR is consumed by +# run-device-api-ci.sh. +echo "==> [5/6] build rccl-tests (ENABLE_DEVICE_API + ENABLE_ROCSHMEM_GIN)" +rccl_tests_src="${workdir}/projects/rccl-tests" +rccl_tests_build="${workdir}/.ci-out/rccl-tests-build" +rm -rf "${rccl_tests_build}"; mkdir -p "${rccl_tests_build}" +( + cd "${rccl_tests_build}" + cmake -DCMAKE_BUILD_TYPE=Release -DUSE_MPI=ON \ + -DENABLE_DEVICE_API=ON -DENABLE_ROCSHMEM_GIN=ON \ + -DROCSHMEM_INSTALL_DIR="${ROCSHMEM_INSTALL_DIR}" \ + -DRCCL_DIR="${rccl_cmake_dir}" \ + -DCMAKE_PREFIX_PATH="${rccl_prefix};${MPI_HOME}" \ + -DGPU_TARGETS="${amdgpu_targets}" \ + "${rccl_tests_src}" + make -j"${build_jobs}" -k all_gather_perf alltoall_perf all_reduce_perf || true +) +if [[ ! -x "${rccl_tests_build}/all_reduce_perf" ]]; then + echo "ERROR: all_reduce_perf not built under ${rccl_tests_build}" >&2 + ls -la "${rccl_tests_build}" 2>/dev/null >&2 || true + exit 1 +fi +export RCCL_TESTS_BIN_DIR="${rccl_tests_build}" } > "${build_out}" 2> "${build_err}" echo "==> build phase logs: ${build_out} / ${build_err}" # --- test phase -> test-%j.{out,err} --- { -# 5) Benchmarks. -echo "==> [5/5] run benchmarks" +# 6) Benchmarks. +echo "==> [6/6] run benchmarks" # GIN/RoCE QP creation pins memory; a low memlock limit makes ibv_create_qp fail with ENOMEM. Raise it for the mpirun ranks and log it. ulimit -l unlimited 2>/dev/null || echo "WARN: could not raise RLIMIT_MEMLOCK (ulimit -l)" echo "==> RLIMIT_MEMLOCK (ulimit -l) = $(ulimit -l)" diff --git a/projects/rccl/tools/ci/run-device-api-ci.sh b/projects/rccl/tools/ci/run-device-api-ci.sh index aae310f913d8..837480545795 100755 --- a/projects/rccl/tools/ci/run-device-api-ci.sh +++ b/projects/rccl/tools/ci/run-device-api-ci.sh @@ -72,14 +72,23 @@ fi cd "${RCCL_TESTS_DIR}" export PATH="${MPI_HOME}/bin:${ROCM_PATH}/bin:${PATH}" -export LD_LIBRARY_PATH="${RCCL_LIB_DIR}:${MPI_HOME}/lib:${ROCM_PATH}/lib:${LD_LIBRARY_PATH:-}" +if [[ -n "${ROCSHMEM_INSTALL_DIR:-}" ]]; then + export LD_LIBRARY_PATH="${RCCL_LIB_DIR}:${ROCSHMEM_INSTALL_DIR}/lib:${MPI_HOME}/lib:${ROCM_PATH}/lib:${LD_LIBRARY_PATH:-}" +else + export LD_LIBRARY_PATH="${RCCL_LIB_DIR}:${MPI_HOME}/lib:${ROCM_PATH}/lib:${LD_LIBRARY_PATH:-}" +fi -PERF_DIR=build +if [[ -n "${RCCL_TESTS_BIN_DIR:-}" ]]; then + PERF_DIR="${RCCL_TESTS_BIN_DIR}" +else + PERF_DIR="${RCCL_TESTS_DIR}/build" +fi if [[ ! -d "${PERF_DIR}" || ! -f "${PERF_DIR}/all_reduce_perf" ]]; then echo "rccl-tests perf binaries not found under ${PERF_DIR}" - ls -la + ls -la "${PERF_DIR}" 2>/dev/null || ls -la "${RCCL_TESTS_DIR}" exit 1 fi +echo "==> rccl-tests perf dir = ${PERF_DIR}" PARSER="${script_dir}/lib/parse_device_api_config.py" [[ -f "${CONFIG}" ]] || { echo "ERROR: test-matrix config not found: ${CONFIG}" >&2; exit 1; } @@ -136,7 +145,7 @@ run_bench() { set +e timeout --kill-after="${BENCH_KILL_AFTER}" "${BENCH_TIMEOUT}" \ mpirun -np "${NP}" ${env_flags} -x LD_LIBRARY_PATH \ - "./${PERF_DIR}/${bin}" ${BENCH_ARGS} ${extra_args} + "${PERF_DIR}/${bin}" ${BENCH_ARGS} ${extra_args} local rc=$? set -e if [[ ${rc} -ne 0 ]]; then @@ -175,7 +184,7 @@ echo "All device-api benchmark runs succeeded." # Requires ENABLE_DEVICE_API=ON rccl-tests build (alltoall_perf with -B flags). run_devtime_smoke() { local pytest_dir="${RCCL_TESTS_DIR}/test" - local perf_bin="${RCCL_TESTS_DIR}/${PERF_DIR}/alltoall_perf" + local perf_bin="${PERF_DIR}/alltoall_perf" if [[ ! -x "${perf_bin}" ]]; then echo "WARN: skip devtime smoke: ${perf_bin} not found" return 0 From 264938bdf366aacea5470420bf0d284107ed3ccb Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sun, 13 Sep 2026 16:43:18 -0400 Subject: [PATCH 3/8] rccl: do not include Anvil-SDMA headers in every DeviceLinker TU. ENABLE_ROCSHMEM_GIN was defaulting NCCL_GIN_ANVIL_SDMA_ENABLE to 1, so non-GIN fat objects pulled rocSHMEM anvil_device.hpp and failed under RCCL_POISON_HIP_ATOMICS. Keep SDMA opt-in on GIN kernel TUs and rccl-tests RDC binaries. Co-authored-by: Cursor --- projects/rccl-tests/CMakeLists.txt | 9 ++++----- projects/rccl-tests/src/CMakeLists.txt | 4 ++-- .../src/include/nccl_device/gin/gin_device_common.h | 11 ++++++----- projects/rccl/test/CMakeLists.txt | 7 +++---- projects/rccl/test/host/CMakeLists.txt | 4 ---- 5 files changed, 15 insertions(+), 20 deletions(-) diff --git a/projects/rccl-tests/CMakeLists.txt b/projects/rccl-tests/CMakeLists.txt index 7b54fced4abd..fcb97dc7307d 100644 --- a/projects/rccl-tests/CMakeLists.txt +++ b/projects/rccl-tests/CMakeLists.txt @@ -238,12 +238,11 @@ if (ENABLE_ROCSHMEM_GIN) if (NOT ROCSHMEM_SOURCE_DIR AND EXISTS "${_mono_root}/projects/rocshmem/CMakeLists.txt") set(ROCSHMEM_SOURCE_DIR "${_mono_root}/projects/rocshmem") endif() - # Installed rocshmem headers (rocshmem_config.h) are needed globally: - # RCCL's public nccl_device.h → gin_anvil_sdma.h → anvil_device.hpp - # → log.hpp → rocshmem/rocshmem_config.h. Must come before source - # includes so the generated config is found first. + # Installed rocshmem headers (rocshmem_config.h) for GIN-SDMA perf binaries + # that set NCCL_GIN_ANVIL_SDMA_ENABLE=1. Must come before source includes + # so the generated config is found first. include_directories(BEFORE ${ROCSHMEM_INSTALL_DIR}/include) - # Internal rocshmem headers (sdma/anvil_device.hpp) needed globally. + # Internal rocshmem headers (sdma/anvil_device.hpp) for those same TUs. if (ROCSHMEM_SOURCE_DIR) include_directories(${ROCSHMEM_SOURCE_DIR}/src) endif() diff --git a/projects/rccl-tests/src/CMakeLists.txt b/projects/rccl-tests/src/CMakeLists.txt index cfd0e4f3facc..730412f01106 100644 --- a/projects/rccl-tests/src/CMakeLists.txt +++ b/projects/rccl-tests/src/CMakeLists.txt @@ -147,7 +147,7 @@ add_rccl_test(all_reduce) add_rccl_test(alltoall) if (ENABLE_ROCSHMEM_GIN AND TARGET roc::rocshmem) function(_rccl_tests_link_rocshmem_gin_rdc target) - target_compile_definitions(${target} PRIVATE ENABLE_ROCSHMEM_GIN) + target_compile_definitions(${target} PRIVATE ENABLE_ROCSHMEM_GIN NCCL_GIN_ANVIL_SDMA_ENABLE=1) target_compile_options(${target} PRIVATE -fgpu-rdc $<$:-ggdb -O0>) # --whole-archive: librccl.so (ENABLE_ROCSHMEM_GIN) leaves rocshmem host # symbols unresolved (--allow-shlib-undefined). The executable must pull @@ -190,7 +190,7 @@ if (ENABLE_ROCSHMEM_GIN AND TARGET roc::rocshmem) # GIN-SDMA hybrid Broadcast (broadcast_perf -D 3): same device-RDC link as # alltoall_perf/all_gather_perf so the Gin*BroadcastKernel variants resolve # the GIN device backend (SDMA puts + LSA/LL peer stores). - target_compile_definitions(broadcast_perf PRIVATE ENABLE_ROCSHMEM_GIN) + target_compile_definitions(broadcast_perf PRIVATE ENABLE_ROCSHMEM_GIN NCCL_GIN_ANVIL_SDMA_ENABLE=1) target_compile_options(broadcast_perf PRIVATE -fgpu-rdc $<$:-ggdb -O0>) target_link_libraries(broadcast_perf PRIVATE -Wl,--whole-archive roc::rocshmem -Wl,--no-whole-archive) diff --git a/projects/rccl/src/include/nccl_device/gin/gin_device_common.h b/projects/rccl/src/include/nccl_device/gin/gin_device_common.h index 41b48ee310b0..8f548bf1cbc4 100644 --- a/projects/rccl/src/include/nccl_device/gin/gin_device_common.h +++ b/projects/rccl/src/include/nccl_device/gin/gin_device_common.h @@ -39,7 +39,8 @@ #endif #endif -// GIN rocshmem device templates (GDA, SDMA) gate on ENABLE_ROCSHMEM_GIN. +// GIN rocSHMEM GDA device templates. ENABLE_ROCSHMEM_GIN is library-wide; +// Anvil-SDMA does not follow this default (opt-in per TU below). #ifndef NCCL_GIN_ROCSHMEM_GDA_ENABLE #if defined(__HIP_PLATFORM_AMD__) && defined(ENABLE_ROCSHMEM_GIN) #define NCCL_GIN_ROCSHMEM_GDA_ENABLE 1 @@ -49,12 +50,12 @@ #endif #ifndef NCCL_GIN_ANVIL_SDMA_ENABLE -#if defined(__HIP_PLATFORM_AMD__) && defined(ENABLE_ROCSHMEM_GIN) -#define NCCL_GIN_ANVIL_SDMA_ENABLE 1 -#else +// Opt-in only. ENABLE_ROCSHMEM_GIN is a library-wide host flag; if this +// defaults to 1, every DeviceLinker TU that includes gin__funcs.h pulls +// rocSHMEM anvil_device.hpp and hits RCCL_POISON_HIP_ATOMICS. GIN-SDMA +// kernels pass -DNCCL_GIN_ANVIL_SDMA_ENABLE=1 on their own TUs. #define NCCL_GIN_ANVIL_SDMA_ENABLE 0 #endif -#endif // Test seam for the system-scope fence issued by the HIP GIN Put/PutValue // templates. Expands to the real fence in production (zero cost); unit tests may diff --git a/projects/rccl/test/CMakeLists.txt b/projects/rccl/test/CMakeLists.txt index d156fa9bfa43..094dc19a5d36 100644 --- a/projects/rccl/test/CMakeLists.txt +++ b/projects/rccl/test/CMakeLists.txt @@ -960,8 +960,8 @@ if(BUILD_TESTS) target_include_directories(${test_executable} PRIVATE ${RCCL_COMMON_INCLUDE_DIRS}) target_compile_definitions(${test_executable} PRIVATE ${RCCL_COMMON_COMPILE_DEFS}) - # RCCL_COMMON_COMPILE_DEFS inherits ENABLE_ROCSHMEM_GIN from rccl's COMPILE_DEFINITIONS, - # which defaults NCCL_GIN_ANVIL_SDMA_ENABLE to 1 and pulls in rocSHMEM's sdma headers. + # Belt-and-suspenders: Anvil-SDMA device headers are opt-in per TU. Keep the + # explicit 0 on non-stub tests so ENABLE_ROCSHMEM_GIN cannot leak them back in. if(ENABLE_ROCSHMEM_GIN AND NOT "${test_executable}" IN_LIST RCCL_ANVIL_SDMA_STUB_TARGETS) target_compile_definitions(${test_executable} PRIVATE NCCL_GIN_ANVIL_SDMA_ENABLE=0) endif() @@ -1094,8 +1094,7 @@ if(BUILD_TESTS) endforeach() # Suite H: shadow rocSHMEM sdma headers with no-op test stubs (higher-priority include path). - # Gated on ENABLE_ROCSHMEM_GIN because that is what now turns on NCCL_GIN_ANVIL_SDMA_ENABLE, - # and these are the only two test targets that define it. + # These two targets opt into NCCL_GIN_ANVIL_SDMA_ENABLE=1; other tests keep it at 0. if(ENABLE_ROCSHMEM_GIN) if(ROCSHMEM_SOURCE_DIR) target_include_directories(rccl-UnitTestsFixtures PRIVATE ${ROCSHMEM_SOURCE_DIR}/src) diff --git a/projects/rccl/test/host/CMakeLists.txt b/projects/rccl/test/host/CMakeLists.txt index bf00b3489d1b..05c429a5ee08 100644 --- a/projects/rccl/test/host/CMakeLists.txt +++ b/projects/rccl/test/host/CMakeLists.txt @@ -49,8 +49,6 @@ if(BUILD_TESTS) target_compile_definitions(rccl-UnitTestsMicro PRIVATE ${RCCL_COMMON_COMPILE_DEFS}) # Mirrors the opt-out in test/CMakeLists.txt, whose loop does not reach this target. - # Inheriting ENABLE_ROCSHMEM_GIN defaults NCCL_GIN_ANVIL_SDMA_ENABLE to 1, which pulls - # in rocSHMEM's sdma device headers. if(ENABLE_ROCSHMEM_GIN) target_compile_definitions(rccl-UnitTestsMicro PRIVATE NCCL_GIN_ANVIL_SDMA_ENABLE=0) endif() @@ -364,8 +362,6 @@ if(BUILD_TESTS) ENQUEUE_CC_PATH="${PROJECT_BINARY_DIR}/hipify/src/enqueue.cc") # Mirrors the opt-out in test/CMakeLists.txt, whose loop does not reach this target. - # Inheriting ENABLE_ROCSHMEM_GIN defaults NCCL_GIN_ANVIL_SDMA_ENABLE to 1, which pulls - # in rocSHMEM's sdma device headers. if(ENABLE_ROCSHMEM_GIN) target_compile_definitions(rccl-UnitTestsMicroEnqueue PRIVATE NCCL_GIN_ANVIL_SDMA_ENABLE=0) endif() From 6c1e471a5968659c9b543a81971d9144941c394d Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sun, 13 Sep 2026 16:58:41 -0400 Subject: [PATCH 4/8] rccl: address GIN harness and device-api review Co-authored-by: Cursor --- projects/rccl-tests/src/CMakeLists.txt | 6 +-- projects/rccl-tests/test/gin_sdma_harness.py | 5 +- projects/rccl-tests/test/test_AllGather.py | 11 ++-- projects/rccl-tests/test/test_AllToAll.py | 7 +-- projects/rccl-tests/test/test_Broadcast.py | 52 ++++++++++++++----- projects/rccl/test/CMakeLists.txt | 47 ++++++++--------- .../rccl/tools/ci/lib/device-api-tests.json | 6 +++ projects/rccl/tools/ci/lib/device-api.sbatch | 22 +++++--- projects/rccl/tools/ci/run-device-api-ci.sh | 21 ++++---- 9 files changed, 105 insertions(+), 72 deletions(-) diff --git a/projects/rccl-tests/src/CMakeLists.txt b/projects/rccl-tests/src/CMakeLists.txt index 730412f01106..ae11b9df1912 100644 --- a/projects/rccl-tests/src/CMakeLists.txt +++ b/projects/rccl-tests/src/CMakeLists.txt @@ -190,11 +190,7 @@ if (ENABLE_ROCSHMEM_GIN AND TARGET roc::rocshmem) # GIN-SDMA hybrid Broadcast (broadcast_perf -D 3): same device-RDC link as # alltoall_perf/all_gather_perf so the Gin*BroadcastKernel variants resolve # the GIN device backend (SDMA puts + LSA/LL peer stores). - target_compile_definitions(broadcast_perf PRIVATE ENABLE_ROCSHMEM_GIN NCCL_GIN_ANVIL_SDMA_ENABLE=1) - target_compile_options(broadcast_perf PRIVATE -fgpu-rdc $<$:-ggdb -O0>) - target_link_libraries(broadcast_perf PRIVATE - -Wl,--whole-archive roc::rocshmem -Wl,--no-whole-archive) - target_link_options(broadcast_perf PRIVATE -fgpu-rdc --hip-link -rdynamic) + _rccl_tests_link_rocshmem_gin_rdc(broadcast_perf) endif() add_rccl_test(gather) add_rccl_test(hypercube) diff --git a/projects/rccl-tests/test/gin_sdma_harness.py b/projects/rccl-tests/test/gin_sdma_harness.py index b6159858b080..4690398c7d70 100644 --- a/projects/rccl-tests/test/gin_sdma_harness.py +++ b/projects/rccl-tests/test/gin_sdma_harness.py @@ -109,7 +109,7 @@ def launch_mpi_shell(cmd, timeout_s, hang_msg): except ProcessLookupError: pass out, _ = proc.communicate() - pytest.fail(hang_msg.format(out=(out or "")[-2000:])) + pytest.fail(hang_msg.format((out or "")[-2000:])) print(out) return proc.returncode, out @@ -164,7 +164,10 @@ def read_nccl_debug_logs(debug_dir): r"\s+{n}\s+{n}\s+{n}\s+({w})" r"\s+{n}\s+{n}\s+{n}\s+({w})".format(n=_NUM, w=_WRONG) ) +# Deliberately not end-anchored: rccl-tests may append an algo/proto/nchannels +# group and a timestamp, and concurrent rank output can splice onto the line. BCAST_OOB_RE = re.compile(r"Out of bounds values\s*:\s*(\d+)") +# Matches the tier emitted by bcastReportTier() in src/broadcast.cu. BCAST_TIER_RE = re.compile(r"^#\[bcast-tier\]\s+(\S+)", re.M) diff --git a/projects/rccl-tests/test/test_AllGather.py b/projects/rccl-tests/test/test_AllGather.py index 98896811f4f8..86b455bd20a4 100644 --- a/projects/rccl-tests/test/test_AllGather.py +++ b/projects/rccl-tests/test/test_AllGather.py @@ -28,9 +28,10 @@ import pytest -from gin_sdma_harness import ( +from .gin_sdma_harness import ( GiB, MiB, + detect_ngpus, env_int, gin_env_xflags, launch_mpi_shell, @@ -38,13 +39,7 @@ run_with_conn_gate_retry, ) -ngpus = 0 -if os.environ.get('ROCR_VISIBLE_DEVICES') is not None: - ngpus = len(os.environ['ROCR_VISIBLE_DEVICES'].split(",")) -elif os.environ.get('HIP_VISIBLE_DEVICES') is not None: - ngpus = len(os.environ['HIP_VISIBLE_DEVICES'].split(",")) -else: - ngpus = int(subprocess.check_output("rocminfo | grep \"Device Type:.\s*.GPU\" | wc -l",shell=True)) +ngpus = max(1, detect_ngpus()) log_ngpus = int(math.log2(ngpus)) nthreads = ["1"] diff --git a/projects/rccl-tests/test/test_AllToAll.py b/projects/rccl-tests/test/test_AllToAll.py index 73afa06f292a..4216837640d2 100644 --- a/projects/rccl-tests/test/test_AllToAll.py +++ b/projects/rccl-tests/test/test_AllToAll.py @@ -47,8 +47,9 @@ # RCCL_TESTS_MPI_OPTS extra launcher opts (e.g. --allow-run-as-root -mca ...) # RCCL_TESTS_A2A_XENV extra "-x K=V" env the backend needs on this cluster # RCCL_TESTS_A2A_EXE path to alltoall_perf (default: ../build/alltoall_perf) +# RCCL_TESTS_A2A_CTAS device CTA count (-V) (default: 16) # RCCL_TESTS_A2A_TIMEOUT_S per-run hang timeout in seconds (default: 900) -# RCCL_TESTS_A2A_CONN_RETRIES connectivity-gate retries (default: 3) +# RCCL_TESTS_A2A_CONN_RETRIES connectivity-gate retries (default: 5) # # Verified on 8x MI355X (ROCm 7.13, NCCL_GIN_TYPE=6, force1ch): 256 MiB/peer and # 2 GiB/peer int32 and the 2 GiB-total guard all pass with #wrong=0, no hang @@ -59,7 +60,7 @@ import pytest -from gin_sdma_harness import ( +from .gin_sdma_harness import ( GiB, MiB, detect_ngpus, @@ -86,7 +87,7 @@ LAUNCHER = os.environ.get("RCCL_TESTS_MPI_LAUNCHER", "mpirun") CTAS = os.environ.get("RCCL_TESTS_A2A_CTAS", "16") TIMEOUT_S = env_int("RCCL_TESTS_A2A_TIMEOUT_S", 900) -CONN_RETRIES = env_int("RCCL_TESTS_A2A_CONN_RETRIES", 3) +CONN_RETRIES = env_int("RCCL_TESTS_A2A_CONN_RETRIES", 5) MPI_OPTS = shlex.split(os.environ.get("RCCL_TESTS_MPI_OPTS", "")) XENV = shlex.split(os.environ.get("RCCL_TESTS_A2A_XENV", "")) diff --git a/projects/rccl-tests/test/test_Broadcast.py b/projects/rccl-tests/test/test_Broadcast.py index 28639052cad3..13f6e2851a54 100644 --- a/projects/rccl-tests/test/test_Broadcast.py +++ b/projects/rccl-tests/test/test_Broadcast.py @@ -27,29 +27,21 @@ import pytest -from gin_sdma_harness import ( +from .gin_sdma_harness import ( GiB, MiB, assert_bcast_ok, bcast_data_failed, bcast_rows, bcast_tiers, + detect_ngpus, env_int, launch_bcast_gin_sdma, run_bcast_with_conn_gate_retry, + run_with_conn_gate_retry, ) -ngpus = 0 -if os.environ.get("ROCR_VISIBLE_DEVICES") is not None: - ngpus = len(os.environ["ROCR_VISIBLE_DEVICES"].split(",")) -elif os.environ.get("HIP_VISIBLE_DEVICES") is not None: - ngpus = len(os.environ["HIP_VISIBLE_DEVICES"].split(",")) -else: - ngpus = int( - subprocess.check_output( - 'rocminfo | grep "Device Type:.\\s*.GPU" | wc -l', shell=True - ) - ) +ngpus = max(1, detect_ngpus()) log_ngpus = int(math.log2(ngpus)) nthreads = ["1"] @@ -294,6 +286,42 @@ def test_BroadcastGinSdmaDataFailedTreatsNaAsUncheckedNotFailed(): assert not bcast_data_failed(line) +def test_BroadcastGinSdmaConnRetryReturnsCleanPassWithoutRetry(): + calls = [] + + def launch(): + calls.append(None) + return 0, "clean" + + assert run_with_conn_gate_retry(launch, 5, settle_s=0) == (0, "clean") + assert len(calls) == 1 + + +def test_BroadcastGinSdmaConnRetryExhaustsConnectivityAborts(): + calls = [] + + def launch(): + calls.append(None) + return 1, "LSA signal connectivity gate failed" + + assert run_with_conn_gate_retry(launch, 3, settle_s=0) == ( + 1, + "LSA signal connectivity gate failed", + ) + assert len(calls) == 3 + + +def test_BroadcastGinSdmaConnRetryNeverRetriesDataFailure(): + calls = [] + + def launch(): + calls.append(None) + return 1, "#wrong = 1" + + assert run_with_conn_gate_retry(launch, 5, settle_s=0) == (1, "#wrong = 1") + assert len(calls) == 1 + + def _clean_bcast_output(tier): """A minimal healthy -D 3 run reporting `tier`, for the guards below.""" return "\n".join( diff --git a/projects/rccl/test/CMakeLists.txt b/projects/rccl/test/CMakeLists.txt index 094dc19a5d36..1e6d9b2d828c 100644 --- a/projects/rccl/test/CMakeLists.txt +++ b/projects/rccl/test/CMakeLists.txt @@ -418,21 +418,34 @@ if(BUILD_TESTS) target_compile_definitions(rccl-UnitTestsFixtures PRIVATE ENABLE_ROCSHMEM) endif() - if(ENABLE_ROCSHMEM_GIN OR GIN_ANVIL_UNIT_TESTS) + function(_rccl_test_add_rocshmem_includes target) + set(_rocshmem_test_includes "") if(ROCSHMEM_INCLUDE_DIR) - target_include_directories(rccl-UnitTestsFixtures PRIVATE ${ROCSHMEM_INCLUDE_DIR}) + list(APPEND _rocshmem_test_includes "${ROCSHMEM_INCLUDE_DIR}") elseif(ROCSHMEM_INSTALL_DIR) - target_include_directories(rccl-UnitTestsFixtures PRIVATE - "${ROCSHMEM_INSTALL_DIR}/include") + list(APPEND _rocshmem_test_includes "${ROCSHMEM_INSTALL_DIR}/include") endif() if(ROCSHMEM_SOURCE_DIR) - target_include_directories(rccl-UnitTestsFixtures PRIVATE + list(APPEND _rocshmem_test_includes "${ROCSHMEM_SOURCE_DIR}/src" "${ROCSHMEM_SOURCE_DIR}/include") + elseif(EXISTS "${CMAKE_SOURCE_DIR}/../rocshmem") + list(APPEND _rocshmem_test_includes + "${CMAKE_SOURCE_DIR}/../rocshmem/src" + "${CMAKE_SOURCE_DIR}/../rocshmem/include") endif() if(ROCSHMEM_BUILD_DIR) - target_include_directories(rccl-UnitTestsFixtures PRIVATE ${ROCSHMEM_BUILD_DIR}) + list(APPEND _rocshmem_test_includes "${ROCSHMEM_BUILD_DIR}") + elseif(ROCSHMEM_SOURCE_DIR AND EXISTS "${ROCSHMEM_SOURCE_DIR}/build/include") + list(APPEND _rocshmem_test_includes "${ROCSHMEM_SOURCE_DIR}/build/include") + endif() + if(_rocshmem_test_includes) + target_include_directories(${target} PRIVATE ${_rocshmem_test_includes}) endif() + endfunction() + + if(ENABLE_ROCSHMEM_GIN OR GIN_ANVIL_UNIT_TESTS) + _rccl_test_add_rocshmem_includes(rccl-UnitTestsFixtures) endif() if(ENABLE_ROCSHMEM_GIN AND ROCM_VERSION VERSION_GREATER_EQUAL "60400") @@ -682,27 +695,9 @@ if(BUILD_TESTS) -Wl,--export-dynamic-symbol=ncclNetPlugin_v11 -Wl,--export-dynamic-symbol=ncclCollNetPlugin_v11) - # Transitive nccl_device/gin headers pull sdma/anvil_device.hpp from rocSHMEM. + # GIN-enabled Debug fixtures need the same rocSHMEM header layout as fixtures. if(ENABLE_ROCSHMEM_GIN) - if(ROCSHMEM_SOURCE_DIR) - target_include_directories(rccl-UnitTestsFixturesDebug PRIVATE - "${ROCSHMEM_SOURCE_DIR}/src" - "${ROCSHMEM_SOURCE_DIR}/include") - else() - target_include_directories(rccl-UnitTestsFixturesDebug PRIVATE - ${CMAKE_SOURCE_DIR}/../rocshmem/src - ${CMAKE_SOURCE_DIR}/../rocshmem/include) - endif() - if(ROCSHMEM_INCLUDE_DIR) - target_include_directories(rccl-UnitTestsFixturesDebug PRIVATE ${ROCSHMEM_INCLUDE_DIR}) - elseif(ROCSHMEM_INSTALL_DIR) - target_include_directories(rccl-UnitTestsFixturesDebug PRIVATE "${ROCSHMEM_INSTALL_DIR}/include") - endif() - if(ROCSHMEM_BUILD_DIR) - target_include_directories(rccl-UnitTestsFixturesDebug PRIVATE ${ROCSHMEM_BUILD_DIR}) - elseif(ROCSHMEM_SOURCE_DIR AND EXISTS "${ROCSHMEM_SOURCE_DIR}/build/include") - target_include_directories(rccl-UnitTestsFixturesDebug PRIVATE "${ROCSHMEM_SOURCE_DIR}/build/include") - endif() + _rccl_test_add_rocshmem_includes(rccl-UnitTestsFixturesDebug) endif() endif() diff --git a/projects/rccl/tools/ci/lib/device-api-tests.json b/projects/rccl/tools/ci/lib/device-api-tests.json index 721fcc14b9d3..3b24d2efc3f0 100644 --- a/projects/rccl/tools/ci/lib/device-api-tests.json +++ b/projects/rccl/tools/ci/lib/device-api-tests.json @@ -41,6 +41,12 @@ "args": "-R 2 -D 3", "env": ["NCCL_DMABUF_ENABLE=1", "NCCL_GIN_TYPE=2", "HSA_NO_SCRATCH_RECLAIM=1", "NCCL_ENV_PLUGIN=none", "RCCL_ENABLE_INTRANET=1"] }, + { + "name": "gin-d3-allgather", + "bins": ["all_gather_perf"], + "args": "-R 2 -D 3", + "env": ["NCCL_DMABUF_ENABLE=1", "NCCL_GIN_TYPE=6", "HSA_NO_SCRATCH_RECLAIM=1", "NCCL_ENV_PLUGIN=none", "RCCL_ENABLE_INTRANET=1"] + }, { "name": "gin-d4", "bins": ["alltoall_perf"], diff --git a/projects/rccl/tools/ci/lib/device-api.sbatch b/projects/rccl/tools/ci/lib/device-api.sbatch index 7f9bb7f648a6..e92e7f8aef49 100755 --- a/projects/rccl/tools/ci/lib/device-api.sbatch +++ b/projects/rccl/tools/ci/lib/device-api.sbatch @@ -63,7 +63,7 @@ build_jobs="${RCCL_BUILD_JOBS}" { # 1) ROCm: provisioned on the head node from TheRock CI artifacts; consume the # workspace tree via rocm.env. -echo "==> [1/5] use ROCm (provisioned from TheRock artifacts on head node)" +echo "==> [1/6] use ROCm (provisioned from TheRock artifacts on head node)" rocm_env="${workdir}/.ci-out/rocm.env" if [[ ! -f "${rocm_env}" ]]; then echo "ERROR: ${rocm_env} not found. The workflow must run the 'Provision ROCm" >&2 @@ -84,7 +84,7 @@ echo " hipcc: $("${ROCM_PATH}/bin/hipcc" --version 2>/dev/null | head -n1)" # 2) OpenMPI: built here against ROCM_PATH into the per-job workspace cache # (no shared cache), then ompi.env (MPI_HOME) is sourced below. -echo "==> [2/5] build OpenMPI" +echo "==> [2/6] build OpenMPI" "${ci_dir}/lib/build-ompi.sh" set -a # shellcheck source=/dev/null # runtime fragment written by build-ompi.sh @@ -152,13 +152,19 @@ rm -rf "${rccl_tests_build}"; mkdir -p "${rccl_tests_build}" -DCMAKE_PREFIX_PATH="${rccl_prefix};${MPI_HOME}" \ -DGPU_TARGETS="${amdgpu_targets}" \ "${rccl_tests_src}" - make -j"${build_jobs}" -k all_gather_perf alltoall_perf all_reduce_perf || true + make -j"${build_jobs}" -k || true ) -if [[ ! -x "${rccl_tests_build}/all_reduce_perf" ]]; then - echo "ERROR: all_reduce_perf not built under ${rccl_tests_build}" >&2 - ls -la "${rccl_tests_build}" 2>/dev/null >&2 || true - exit 1 -fi +required_perf_bins=( + all_gather_perf all_reduce_perf alltoall_perf broadcast_perf + gather_perf reduce_perf scatter_perf sendrecv_perf +) +for perf_bin in "${required_perf_bins[@]}"; do + if [[ ! -x "${rccl_tests_build}/${perf_bin}" ]]; then + echo "ERROR: ${perf_bin} not built under ${rccl_tests_build}" >&2 + ls -la "${rccl_tests_build}" 2>/dev/null >&2 || true + exit 1 + fi +done export RCCL_TESTS_BIN_DIR="${rccl_tests_build}" } > "${build_out}" 2> "${build_err}" echo "==> build phase logs: ${build_out} / ${build_err}" diff --git a/projects/rccl/tools/ci/run-device-api-ci.sh b/projects/rccl/tools/ci/run-device-api-ci.sh index 837480545795..8ad9e300baf8 100755 --- a/projects/rccl/tools/ci/run-device-api-ci.sh +++ b/projects/rccl/tools/ci/run-device-api-ci.sh @@ -4,9 +4,9 @@ # rocJenkins "device-api" testCommand. # # Consumes ROCM_PATH (rocm.env), MPI_HOME (ompi.env), the in-tree RCCL build at -# projects/rccl/build/release, and rccl-tests/build. ROCM_PATH / MPI_HOME come -# from the environment (device-api.sbatch exports them) or, when run standalone, -# from $WORKDIR/.ci-out/{rocm,ompi}.env. +# projects/rccl/build/release, and rccl-tests/build. Build paths come from the +# environment (device-api.sbatch exports them) or, when run standalone, from +# $WORKDIR/.ci-out/{rocm,ompi,rocshmem}.env. # # Each bench is wrapped in `timeout` so a hung mpirun/driver can't wedge the job; # failures are collected and surfaced at the end (exit non-zero iff any failed). @@ -17,7 +17,9 @@ # debug_env to every run. # # Environment overrides: -# ROCM_PATH / MPI_HOME Else read from .ci-out/{rocm,ompi}.env +# ROCM_PATH / MPI_HOME / ROCSHMEM_INSTALL_DIR +# Else read from matching .ci-out/*.env fragments +# RCCL_TESTS_BIN_DIR rccl-tests build directory (default: rccl-tests/build) # NP MPI ranks per run (default: 8) # BENCH_ARGS Common bench args (default: from JSON bench_args) # BENCH_TIMEOUT Per-bench wall-clock cap (default: 600s) @@ -41,11 +43,12 @@ RCCL_LIB_DIR="${WORKDIR}/projects/rccl/build/release" RCCL_TESTS_DIR="${WORKDIR}/projects/rccl-tests" CONFIG="${CONFIG:-${script_dir}/lib/device-api-tests.json}" -# Prefer the build stages' env fragments over any ambient ROCM_PATH / MPI_HOME. -# shellcheck source=/dev/null # runtime fragment written by the provision step -[[ -f "${WORKDIR}/.ci-out/rocm.env" ]] && source "${WORKDIR}/.ci-out/rocm.env" -# shellcheck source=/dev/null # runtime fragment written by build-ompi.sh -[[ -f "${WORKDIR}/.ci-out/ompi.env" ]] && source "${WORKDIR}/.ci-out/ompi.env" +# Prefer the build stages' env fragments over ambient build paths. +for frag in rocm ompi rocshmem; do + env_file="${WORKDIR}/.ci-out/${frag}.env" + # shellcheck source=/dev/null # runtime fragment written by the build stages + [[ -f "${env_file}" ]] && source "${env_file}" +done : "${ROCM_PATH:?run-device-api-ci.sh: ROCM_PATH unset (provisioned via rocm.env / sbatch)}" : "${MPI_HOME:?run-device-api-ci.sh: MPI_HOME unset (run build-ompi.sh / via sbatch)}" From f822f0261aa71157b7e4f135b861d376f6d83ebb Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sun, 13 Sep 2026 23:20:33 -0400 Subject: [PATCH 5/8] rccl: RDC-link device-api matrix binaries; harness dedupe Device-api CI builds one --rocshmem-gin librccl.so, which leaves GIN rocSHMEM host symbols to the executable. Every binary the matrix launches now whole-archive RDC-links rocSHMEM, not only the GIN collectives. all_reduce_perf still does not build under ENABLE_ROCSHMEM_GIN (same as gin.sbatch), so it is not RDC-linked, not required after `make -k`, and the LSA -D 1/-D 2 suites that launched it are dropped until it builds. The runner's perf-dir sentinel moves to alltoall_perf. Harness: extract gin_perf_argv()/gin_hang_msg() so the three collectives share one argv vector and the AllToAll -A 1 gap is explicit; make detect_ngpus() fail when rocminfo is missing instead of reporting zero GPUs; pin CONN_GATE_RE with a non-gate failure case. Co-authored-by: Cursor --- projects/rccl-tests/src/CMakeLists.txt | 63 +++++----- projects/rccl-tests/test/gin_sdma_harness.py | 111 +++++++++++------- projects/rccl-tests/test/test_AllGather.py | 53 +++------ projects/rccl-tests/test/test_AllToAll.py | 41 ++----- projects/rccl-tests/test/test_Broadcast.py | 42 ++++++- .../rccl/tools/ci/lib/device-api-tests.json | 14 +-- projects/rccl/tools/ci/lib/device-api.sbatch | 9 +- projects/rccl/tools/ci/run-device-api-ci.sh | 11 +- 8 files changed, 190 insertions(+), 154 deletions(-) diff --git a/projects/rccl-tests/src/CMakeLists.txt b/projects/rccl-tests/src/CMakeLists.txt index ae11b9df1912..1b238dcce638 100644 --- a/projects/rccl-tests/src/CMakeLists.txt +++ b/projects/rccl-tests/src/CMakeLists.txt @@ -145,28 +145,6 @@ endif() add_rccl_test(all_gather) add_rccl_test(all_reduce) add_rccl_test(alltoall) -if (ENABLE_ROCSHMEM_GIN AND TARGET roc::rocshmem) - function(_rccl_tests_link_rocshmem_gin_rdc target) - target_compile_definitions(${target} PRIVATE ENABLE_ROCSHMEM_GIN NCCL_GIN_ANVIL_SDMA_ENABLE=1) - target_compile_options(${target} PRIVATE -fgpu-rdc $<$:-ggdb -O0>) - # --whole-archive: librccl.so (ENABLE_ROCSHMEM_GIN) leaves rocshmem host - # symbols unresolved (--allow-shlib-undefined). The executable must pull - # in ALL rocshmem symbols (not just those it references directly) so that - # -rdynamic exports them for librccl.so to resolve at runtime. - target_link_libraries(${target} PRIVATE - -Wl,--whole-archive roc::rocshmem -Wl,--no-whole-archive) - target_link_options(${target} PRIVATE -fgpu-rdc --hip-link -rdynamic) - endfunction() - - _rccl_tests_link_rocshmem_gin_rdc(alltoall_perf) - - # GIN-SDMA hybrid AllGather (all_gather_perf -D 3): same device-RDC link as - # alltoall_perf so the GinHybridAllGatherKernel resolves the GIN device backend. - if (ENABLE_DEVICE_API) - target_compile_definitions(all_gather_perf PRIVATE ENABLE_DEVICE_API) - endif() - _rccl_tests_link_rocshmem_gin_rdc(all_gather_perf) -endif() if (ENABLE_ROCSHMEM_GIN AND TARGET roc::rocshmem AND RCCL_SOURCE_DIR) add_executable(gin_rocshmem_gda_factory_test ${CMAKE_CURRENT_SOURCE_DIR}/gin_rocshmem_gda_factory_test.cu) set_source_files_properties(${CMAKE_CURRENT_SOURCE_DIR}/gin_rocshmem_gda_factory_test.cu PROPERTIES LANGUAGE CXX) @@ -186,12 +164,6 @@ if (ENABLE_ROCSHMEM_GIN AND TARGET roc::rocshmem AND RCCL_SOURCE_DIR) endif() add_rccl_test(alltoallv) add_rccl_test(broadcast) -if (ENABLE_ROCSHMEM_GIN AND TARGET roc::rocshmem) - # GIN-SDMA hybrid Broadcast (broadcast_perf -D 3): same device-RDC link as - # alltoall_perf/all_gather_perf so the Gin*BroadcastKernel variants resolve - # the GIN device backend (SDMA puts + LSA/LL peer stores). - _rccl_tests_link_rocshmem_gin_rdc(broadcast_perf) -endif() add_rccl_test(gather) add_rccl_test(hypercube) add_rccl_test(reduce_scatter) @@ -199,3 +171,38 @@ add_rccl_test(reduce) add_rccl_test(scatter) add_rccl_test(sendrecv) add_rccl_test(all_reduce_bias) +if (ENABLE_ROCSHMEM_GIN AND TARGET roc::rocshmem) + function(_rccl_tests_link_rocshmem_gin_rdc target) + target_compile_definitions(${target} PRIVATE ENABLE_ROCSHMEM_GIN NCCL_GIN_ANVIL_SDMA_ENABLE=1) + target_compile_options(${target} PRIVATE -fgpu-rdc $<$:-ggdb -O0>) + # --whole-archive: librccl.so (ENABLE_ROCSHMEM_GIN) leaves rocshmem host + # symbols unresolved (--allow-shlib-undefined). The executable must pull + # in ALL rocshmem symbols (not just those it references directly) so that + # -rdynamic exports them for librccl.so to resolve at runtime. + target_link_libraries(${target} PRIVATE + -Wl,--whole-archive roc::rocshmem -Wl,--no-whole-archive) + target_link_options(${target} PRIVATE -fgpu-rdc --hip-link -rdynamic) + endfunction() + + # GIN-SDMA hybrid AllGather (all_gather_perf -D 3) needs ENABLE_DEVICE_API + # so the GinHybridAllGatherKernel is compiled into this binary. + if (ENABLE_DEVICE_API) + target_compile_definitions(all_gather_perf PRIVATE ENABLE_DEVICE_API) + endif() + + # Device-api CI builds one --rocshmem-gin librccl.so and runs the matrix + # against it. Every launched binary must RDC-link rocSHMEM so GIN host + # symbols resolve from the executable. all_reduce_perf is omitted: it + # still does not build under ENABLE_ROCSHMEM_GIN (gin.sbatch uses + # `make -k` for the same reason). + foreach(_gin_rdc_bin IN ITEMS + all_gather_perf + alltoall_perf + broadcast_perf + gather_perf + reduce_perf + scatter_perf + sendrecv_perf) + _rccl_tests_link_rocshmem_gin_rdc(${_gin_rdc_bin}) + endforeach() +endif() diff --git a/projects/rccl-tests/test/gin_sdma_harness.py b/projects/rccl-tests/test/gin_sdma_harness.py index 4690398c7d70..5ad34f9457ba 100644 --- a/projects/rccl-tests/test/gin_sdma_harness.py +++ b/projects/rccl-tests/test/gin_sdma_harness.py @@ -49,17 +49,29 @@ def detect_ngpus(): + """Return the visible GPU count. + + Prefer ROCR_VISIBLE_DEVICES / HIP_VISIBLE_DEVICES. Otherwise require + rocminfo on PATH: a missing binary used to look like zero GPUs because + `rocminfo | grep | wc -l` still exits 0, and callers that clamp with + max(1, ...) then silently collect a one-GPU matrix. + """ if os.environ.get("ROCR_VISIBLE_DEVICES") is not None: return len(os.environ["ROCR_VISIBLE_DEVICES"].split(",")) if os.environ.get("HIP_VISIBLE_DEVICES") is not None: return len(os.environ["HIP_VISIBLE_DEVICES"].split(",")) - try: - out = subprocess.check_output( - 'rocminfo | grep "Device Type:.\\s*.GPU" | wc -l', shell=True + rocminfo = shutil.which("rocminfo") + if rocminfo is None: + raise RuntimeError( + "rocminfo not on PATH; cannot detect GPU count. Set " + "ROCR_VISIBLE_DEVICES, HIP_VISIBLE_DEVICES, or RCCL_TESTS_*_NP." ) - return int(out) - except Exception: - return 0 + out = subprocess.check_output( + [rocminfo], universal_newlines=True, stderr=subprocess.DEVNULL + ) + return sum( + 1 for line in out.splitlines() if re.search(r"Device Type:\s*GPU", line) + ) def env_int(name, default): @@ -90,6 +102,54 @@ def gin_env_xflags(kv_list): return flags +def gin_perf_argv(exe, size, dtype, ctas, extra=None): + """Fixed-size GIN -D 3 perf argv. + + `extra` is inserted after `-D 3` and before `-V`, so the AllToAll vs + AllGather/Broadcast `-A 1` gap is an explicit extra rather than a + copy-paste drift. AllToAll omits `-A 1`; the others pass extra=["-A", "1"]. + """ + args = [ + exe, + "-b", + str(size), + "-e", + str(size), + "-f", + "2", + "-g", + "1", + "-R", + "2", + "-D", + "3", + ] + args += list(extra or []) + args += [ + "-V", + str(ctas), + "-d", + dtype, + "-c", + "1", + "-w", + "1", + "-n", + "3", + ] + return args + + +def gin_hang_msg(collective, timeout_s, size, dtype, size_desc): + """Timeout diagnostic; size_desc is e.g. '256 MiB/rank' or '256 MiB'.""" + return ( + "{} GIN-SDMA HANG: no completion within {}s at {} bytes " + "({}), dtype={}. Output tail:\n{{}}".format( + collective, timeout_s, size, size_desc, dtype + ) + ) + + def launch_mpi_shell(cmd, timeout_s, hang_msg): """Run a shell MPI command in a new session; kill the process group on timeout.""" print(cmd) @@ -159,13 +219,13 @@ def read_nccl_debug_logs(debug_dir): _NUM = r"[-+]?(?:\d+\.?\d*|\.\d+)(?:[eE][-+]?\d+)?" _WRONG = r"(?:{}|N/A)".format(_NUM) +# Deliberately not end-anchored: rccl-tests may append an algo/proto/nchannels +# group and a timestamp, and concurrent rank output can splice onto the line. BCAST_ROW_RE = re.compile( r"^\s*(\d+)\s+(\d+)\s+\S+\s+\S+\s+(\d+)" r"\s+{n}\s+{n}\s+{n}\s+({w})" r"\s+{n}\s+{n}\s+{n}\s+({w})".format(n=_NUM, w=_WRONG) ) -# Deliberately not end-anchored: rccl-tests may append an algo/proto/nchannels -# group and a timestamp, and concurrent rank output can splice onto the line. BCAST_OOB_RE = re.compile(r"Out of bounds values\s*:\s*(\d+)") # Matches the tier emitted by bcastReportTier() in src/broadcast.cu. BCAST_TIER_RE = re.compile(r"^#\[bcast-tier\]\s+(\S+)", re.M) @@ -301,40 +361,11 @@ def launch_bcast_gin_sdma( args = ( mpi_launch_prefix(request, launcher, np, mpi_opts) + gin_env_xflags(gin_env) - + [ - exe, - "-b", - size, - "-e", - size, - "-f", - "2", - "-g", - "1", - "-R", - "2", - "-D", - "3", - "-A", - "1", - "-V", - ctas, - "-d", - dtype, - "-c", - "1", - "-w", - "1", - "-n", - "3", - ] + + gin_perf_argv(exe, size, dtype, ctas, extra=["-A", "1"]) ) cmd = " ".join(shlex.quote(a) for a in args) - hang_msg = ( - "Broadcast GIN-SDMA HANG: no completion within {}s at msg={} bytes " - "({} MiB), dtype={}. Output tail:\n{{}}".format( - timeout_s, size, msg_bytes // MiB, dtype - ) + hang_msg = gin_hang_msg( + "Broadcast", timeout_s, size, dtype, "{} MiB".format(msg_bytes // MiB) ) try: rc, out = launch_mpi_shell(cmd, timeout_s, hang_msg) diff --git a/projects/rccl-tests/test/test_AllGather.py b/projects/rccl-tests/test/test_AllGather.py index 86b455bd20a4..2067dbdcd34e 100644 --- a/projects/rccl-tests/test/test_AllGather.py +++ b/projects/rccl-tests/test/test_AllGather.py @@ -34,12 +34,18 @@ detect_ngpus, env_int, gin_env_xflags, + gin_hang_msg, + gin_perf_argv, launch_mpi_shell, mpi_launch_prefix, run_with_conn_gate_retry, ) -ngpus = max(1, detect_ngpus()) +try: + _detected_ngpus = detect_ngpus() +except RuntimeError: + _detected_ngpus = 0 +ngpus = max(1, _detected_ngpus) log_ngpus = int(math.log2(ngpus)) nthreads = ["1"] @@ -131,7 +137,9 @@ def test_AllGatherSingleProcess(nthreads, ngpus_single, byte_range, op, step_fac _ag_enabled = os.environ.get("RCCL_TESTS_GIN_SDMA_AG", "") not in ("", "0", "false", "False") -AG_NP = env_int("RCCL_TESTS_AG_NP", 0) or ngpus +AG_NP = env_int("RCCL_TESTS_AG_NP", 0) or ( + detect_ngpus() if _ag_enabled else _detected_ngpus +) AG_LAUNCHER = os.environ.get("RCCL_TESTS_MPI_LAUNCHER", "mpirun") AG_CTAS = os.environ.get("RCCL_TESTS_AG_CTAS", "8") AG_TIMEOUT_S = env_int("RCCL_TESTS_AG_TIMEOUT_S", 900) @@ -147,7 +155,6 @@ def test_AllGatherSingleProcess(nthreads, ngpus_single, byte_range, op, step_fac "a GIN-SDMA-capable (e.g. 8x MI355X) node to enable.") -# Intermittent gfx950 cuMem-VMM connectivity-gate abort: see gin_sdma_harness.py. _AG_DEVTIME_TIER_RE = re.compile(r"#\[ag-devtime\].*tier\s+(LSA|SDMA)", re.I) @@ -161,45 +168,21 @@ def _launch_ag_gin_sdma(request, total_bytes, dtype, force_sdma_tier=True, devic "NCCL_GIN_ANVIL_SDMA_THRESHOLD_ALLGATHER=0", ] + extra = ["-A", "1"] args = ( mpi_launch_prefix(request, AG_LAUNCHER, AG_NP, AG_MPI_OPTS) + gin_env_xflags(gin_kv) - + [ - AG_EXE, - "-b", - size, - "-e", - size, - "-f", - "2", - "-g", - "1", - "-R", - "2", - "-D", - "3", - "-A", - "1", - "-V", - AG_CTAS, - "-d", - dtype, - "-c", - "1", - "-w", - "1", - "-n", - "3", - ] + + gin_perf_argv(AG_EXE, size, dtype, AG_CTAS, extra=extra) ) if device_timing: args += ["-B", "1"] cmd = " ".join(shlex.quote(a) for a in args) - hang_msg = ( - "AllGather GIN-SDMA HANG: no completion within {}s at total={} bytes " - "({} MiB/rank), dtype={}. Output tail:\n{{}}".format( - AG_TIMEOUT_S, size, total_bytes // AG_NP // MiB, dtype - ) + hang_msg = gin_hang_msg( + "AllGather", + AG_TIMEOUT_S, + size, + dtype, + "{} MiB/rank".format(total_bytes // AG_NP // MiB), ) return launch_mpi_shell(cmd, AG_TIMEOUT_S, hang_msg) diff --git a/projects/rccl-tests/test/test_AllToAll.py b/projects/rccl-tests/test/test_AllToAll.py index 4216837640d2..d79485282d13 100644 --- a/projects/rccl-tests/test/test_AllToAll.py +++ b/projects/rccl-tests/test/test_AllToAll.py @@ -66,6 +66,8 @@ detect_ngpus, env_int, gin_env_xflags, + gin_hang_msg, + gin_perf_argv, launch_mpi_shell, mpi_launch_prefix, run_with_conn_gate_retry, @@ -83,7 +85,7 @@ "False", ) -NP = env_int("RCCL_TESTS_A2A_NP", 0) or detect_ngpus() +NP = env_int("RCCL_TESTS_A2A_NP", 0) or (detect_ngpus() if _enabled else 0) LAUNCHER = os.environ.get("RCCL_TESTS_MPI_LAUNCHER", "mpirun") CTAS = os.environ.get("RCCL_TESTS_A2A_CTAS", "16") TIMEOUT_S = env_int("RCCL_TESTS_A2A_TIMEOUT_S", 900) @@ -114,38 +116,15 @@ def _launch_a2a(request, total_bytes, dtype): args = ( mpi_launch_prefix(request, LAUNCHER, NP, MPI_OPTS) + gin_env - + [ - executable, - "-b", - size, - "-e", - size, - "-f", - "2", - "-g", - "1", - "-R", - "2", - "-D", - "3", - "-V", - CTAS, - "-d", - dtype, - "-c", - "1", - "-w", - "1", - "-n", - "3", - ] + + gin_perf_argv(executable, size, dtype, CTAS) ) cmd = " ".join(shlex.quote(a) for a in args) - hang_msg = ( - "AllToAll GIN-SDMA HANG: no completion within {}s at total={} bytes " - "({} MiB/peer), dtype={}. Output tail:\n{{}}".format( - TIMEOUT_S, size, total_bytes // NP // MiB, dtype - ) + hang_msg = gin_hang_msg( + "AllToAll", + TIMEOUT_S, + size, + dtype, + "{} MiB/peer".format(total_bytes // NP // MiB), ) return launch_mpi_shell(cmd, TIMEOUT_S, hang_msg) diff --git a/projects/rccl-tests/test/test_Broadcast.py b/projects/rccl-tests/test/test_Broadcast.py index 13f6e2851a54..48ed878421e6 100644 --- a/projects/rccl-tests/test/test_Broadcast.py +++ b/projects/rccl-tests/test/test_Broadcast.py @@ -36,12 +36,17 @@ bcast_tiers, detect_ngpus, env_int, + gin_perf_argv, launch_bcast_gin_sdma, run_bcast_with_conn_gate_retry, run_with_conn_gate_retry, ) -ngpus = max(1, detect_ngpus()) +try: + _detected_ngpus = detect_ngpus() +except RuntimeError: + _detected_ngpus = 0 +ngpus = max(1, _detected_ngpus) log_ngpus = int(math.log2(ngpus)) nthreads = ["1"] @@ -137,7 +142,9 @@ def test_BroadcastSingleProcess( "False", ) -BCAST_NP = env_int("RCCL_TESTS_BCAST_NP", 0) or ngpus +BCAST_NP = env_int("RCCL_TESTS_BCAST_NP", 0) or ( + detect_ngpus() if _bcast_enabled else _detected_ngpus +) BCAST_LAUNCHER = os.environ.get("RCCL_TESTS_MPI_LAUNCHER", "mpirun") BCAST_CTAS = os.environ.get("RCCL_TESTS_BCAST_CTAS", "8") BCAST_GIN_TYPE = os.environ.get("RCCL_TESTS_BCAST_GIN_TYPE", "6") @@ -322,6 +329,37 @@ def launch(): assert len(calls) == 1 +def test_BroadcastGinSdmaConnRetryDoesNotRetryNonGateError(): + calls = [] + + def launch(): + calls.append(None) + return 1, "some other error" + + assert run_with_conn_gate_retry(launch, 5, settle_s=0) == ( + 1, + "some other error", + ) + assert len(calls) == 1 + + +def test_BroadcastGinSdmaDetectNgpusRequiresRocminfo(monkeypatch): + monkeypatch.delenv("ROCR_VISIBLE_DEVICES", raising=False) + monkeypatch.delenv("HIP_VISIBLE_DEVICES", raising=False) + from . import gin_sdma_harness as harness + + monkeypatch.setattr(harness.shutil, "which", lambda _cmd: None) + with pytest.raises(RuntimeError, match="rocminfo"): + harness.detect_ngpus() + + +def test_BroadcastGinSdmaPerfArgvMakesAverageCountGapExplicit(): + ag = gin_perf_argv("all_gather_perf", 8, "int32", "8", extra=["-A", "1"]) + a2a = gin_perf_argv("alltoall_perf", 8, "int32", "16") + assert ag[ag.index("-D") + 2 : ag.index("-V")] == ["-A", "1"] + assert a2a[a2a.index("-D") + 2 : a2a.index("-V")] == [] + + def _clean_bcast_output(tier): """A minimal healthy -D 3 run reporting `tier`, for the guards below.""" return "\n".join( diff --git a/projects/rccl/tools/ci/lib/device-api-tests.json b/projects/rccl/tools/ci/lib/device-api-tests.json index 3b24d2efc3f0..0f13baa57788 100644 --- a/projects/rccl/tools/ci/lib/device-api-tests.json +++ b/projects/rccl/tools/ci/lib/device-api-tests.json @@ -10,7 +10,9 @@ " {LOGDIR} in debug_env is replaced by the runner with the per-job debug dir,", " so NCCL_DEBUG_FILE writes per-rank logs there instead of flooding stdout.", "suites: [ { name, bins:[...], args, env:[...] (optional) } ]", - "Switched from YAML to JSON because PyYAML is not installed on the CI hosts." + "Switched from YAML to JSON because PyYAML is not installed on the CI hosts.", + "LSA all_reduce_perf (-D 1/-D 2) is omitted: that binary still does not build", + "under ENABLE_ROCSHMEM_GIN (same as gin.sbatch). Re-add those suites when it does." ], "bench_args": "-b 8 -e 1G -f 2 -g 1", @@ -25,16 +27,6 @@ "bins": ["broadcast_perf", "reduce_perf", "sendrecv_perf", "scatter_perf", "gather_perf"], "args": "-R 2" }, - { - "name": "lsa-d1", - "bins": ["all_reduce_perf"], - "args": "-R 2 -D 1" - }, - { - "name": "lsa-d2", - "bins": ["all_reduce_perf"], - "args": "-R 2 -D 2" - }, { "name": "gin-d3", "bins": ["alltoall_perf"], diff --git a/projects/rccl/tools/ci/lib/device-api.sbatch b/projects/rccl/tools/ci/lib/device-api.sbatch index e92e7f8aef49..c0223481fda4 100755 --- a/projects/rccl/tools/ci/lib/device-api.sbatch +++ b/projects/rccl/tools/ci/lib/device-api.sbatch @@ -136,8 +136,9 @@ fi rccl_cmake_dir="$(dirname "${rccl_cfg}")" echo " in-tree RCCL cmake config: ${rccl_cmake_dir}" -# 5) rccl-tests via cmake so all_gather_perf -D 3 links rocshmem (install.sh/Makefile -# path does not enable ENABLE_ROCSHMEM_GIN). RCCL_TESTS_BIN_DIR is consumed by +# 5) rccl-tests via cmake so every device-api matrix binary RDC-links rocshmem +# (install.sh/Makefile path does not enable ENABLE_ROCSHMEM_GIN). Needed because +# this job's librccl.so is --rocshmem-gin. RCCL_TESTS_BIN_DIR is consumed by # run-device-api-ci.sh. echo "==> [5/6] build rccl-tests (ENABLE_DEVICE_API + ENABLE_ROCSHMEM_GIN)" rccl_tests_src="${workdir}/projects/rccl-tests" @@ -152,10 +153,12 @@ rm -rf "${rccl_tests_build}"; mkdir -p "${rccl_tests_build}" -DCMAKE_PREFIX_PATH="${rccl_prefix};${MPI_HOME}" \ -DGPU_TARGETS="${amdgpu_targets}" \ "${rccl_tests_src}" + # all_reduce_perf still fails under ENABLE_ROCSHMEM_GIN; -k keeps the + # matrix binaries. Do not require all_reduce_perf (see gin.sbatch). make -j"${build_jobs}" -k || true ) required_perf_bins=( - all_gather_perf all_reduce_perf alltoall_perf broadcast_perf + all_gather_perf alltoall_perf broadcast_perf gather_perf reduce_perf scatter_perf sendrecv_perf ) for perf_bin in "${required_perf_bins[@]}"; do diff --git a/projects/rccl/tools/ci/run-device-api-ci.sh b/projects/rccl/tools/ci/run-device-api-ci.sh index 8ad9e300baf8..1675860a20d0 100755 --- a/projects/rccl/tools/ci/run-device-api-ci.sh +++ b/projects/rccl/tools/ci/run-device-api-ci.sh @@ -1,7 +1,8 @@ #!/usr/bin/env bash -# Run the RCCL device-API benchmark suite (symmetric memory / LSA device kernels -# / GIN proxy) against the freshly built rccl-tests tree. Mirrors the legacy -# rocJenkins "device-api" testCommand. +# Run the RCCL device-API benchmark suite (symmetric memory / GIN proxy) +# against the freshly built rccl-tests tree. Mirrors the legacy rocJenkins +# "device-api" testCommand. LSA all_reduce_perf (-D 1/-D 2) is not launched: +# that binary still does not build under ENABLE_ROCSHMEM_GIN. # # Consumes ROCM_PATH (rocm.env), MPI_HOME (ompi.env), the in-tree RCCL build at # projects/rccl/build/release, and rccl-tests/build. Build paths come from the @@ -86,7 +87,9 @@ if [[ -n "${RCCL_TESTS_BIN_DIR:-}" ]]; then else PERF_DIR="${RCCL_TESTS_DIR}/build" fi -if [[ ! -d "${PERF_DIR}" || ! -f "${PERF_DIR}/all_reduce_perf" ]]; then +# all_reduce_perf is not a valid sentinel here: it still fails to build under +# ENABLE_ROCSHMEM_GIN. alltoall_perf is the binary this GIN job always needs. +if [[ ! -d "${PERF_DIR}" || ! -x "${PERF_DIR}/alltoall_perf" ]]; then echo "rccl-tests perf binaries not found under ${PERF_DIR}" ls -la "${PERF_DIR}" 2>/dev/null || ls -la "${RCCL_TESTS_DIR}" exit 1 From 97377b539e89c7b56eb643a54d8f2cde554fea44 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sun, 13 Sep 2026 23:34:50 -0400 Subject: [PATCH 6/8] rccl: isolate Anvil headers from atomics poison Parse rocSHMEM's external anvil_device.hpp before RCCL poisons the __hip_atomic_* builtins, while retaining the poison for all following RCCL code. Share the GIN-SDMA backend definitions between regular and DeviceLinker compilation so both pipelines opt in consistently. Restore NCCL_GIN_ANVIL_SDMA_ENABLE=1 on stub-backed tests and let Debug fixtures link rocSHMEM without requiring GIN_ANVIL_UNIT_TESTS. Co-authored-by: Cursor --- projects/rccl/cmake/DeviceLinker.cmake | 24 +++++++++---------- projects/rccl/src/CMakeLists.txt | 17 +++++++++++++ .../rccl/src/include/poison_hip_atomics.h | 9 +++++++ projects/rccl/test/CMakeLists.txt | 19 +++++++++------ 4 files changed, 50 insertions(+), 19 deletions(-) diff --git a/projects/rccl/cmake/DeviceLinker.cmake b/projects/rccl/cmake/DeviceLinker.cmake index f562a993959f..a792c6c2ba48 100644 --- a/projects/rccl/cmake/DeviceLinker.cmake +++ b/projects/rccl/cmake/DeviceLinker.cmake @@ -853,10 +853,16 @@ add_custom_command( # =========================================================================== # gin_alltoall_sdma.cu.cpp: GIN-SDMA alltoall kernel. # -# The defines go here because this file is filtered out of the rccl target, so -# per-source properties never apply. Leaving only SDMA on keeps ncclGinCallImpl -# on its single backend branch. GDA must stay off, as this object gets no QP bitcode. +# The files are filtered out of the rccl target, so convert their shared +# per-source definitions from src/CMakeLists.txt to flags for these custom +# commands. Leaving only SDMA on keeps ncclGinCallImpl on its single backend +# branch. GDA must stay off, as these objects get no QP bitcode. # =========================================================================== +set(_gin_sdma_def_flags "") +foreach(_def IN LISTS RCCL_GIN_SDMA_COMPILE_DEFINITIONS) + list(APPEND _gin_sdma_def_flags "-D${_def}") +endforeach() + set(GIN_ALLTOALL_SDMA_FAT_OBJ "") if(ENABLE_ROCSHMEM_GIN) set(GIN_ALLTOALL_SDMA_FAT_OBJ "${DEVICE_BUILD_DIR}/gin_alltoall_sdma.o") @@ -866,9 +872,7 @@ if(ENABLE_ROCSHMEM_GIN) -x hip ${DL_OFFLOAD_ARCH_FLAGS} ${DL_HIP_COMPILER_FLAGS} -DRCCL_DEVICE_LINKER - -DNCCL_GIN_ANVIL_SDMA_ENABLE=1 - -DNCCL_GIN_PROXY_ENABLE=0 - -DNCCL_GIN_ROCSHMEM_GDA_ENABLE=0 + ${_gin_sdma_def_flags} ${_link_def_flags} ${_host_inc_flags} ${DL_OPT_FLAGS} @@ -887,9 +891,7 @@ endif() # =========================================================================== # gin_all_reduce_sdma.cu.cpp: GIN-SDMA allreduce kernel. # -# The defines go here because this file is filtered out of the rccl target, so -# per-source properties never apply. Leaving only SDMA on keeps ncclGinCallImpl -# on its single backend branch. GDA must stay off, as this object gets no QP bitcode. +# Uses the same `_gin_sdma_def_flags` as alltoall above. # =========================================================================== set(GIN_ALLREDUCE_SDMA_FAT_OBJ "") if(ENABLE_ROCSHMEM_GIN) @@ -900,9 +902,7 @@ if(ENABLE_ROCSHMEM_GIN) -x hip ${DL_OFFLOAD_ARCH_FLAGS} ${DL_HIP_COMPILER_FLAGS} -DRCCL_DEVICE_LINKER - -DNCCL_GIN_ANVIL_SDMA_ENABLE=1 - -DNCCL_GIN_PROXY_ENABLE=0 - -DNCCL_GIN_ROCSHMEM_GDA_ENABLE=0 + ${_gin_sdma_def_flags} ${_link_def_flags} ${_host_inc_flags} ${DL_OPT_FLAGS} diff --git a/projects/rccl/src/CMakeLists.txt b/projects/rccl/src/CMakeLists.txt index 01d97c1668b1..dce6ceddca18 100644 --- a/projects/rccl/src/CMakeLists.txt +++ b/projects/rccl/src/CMakeLists.txt @@ -844,6 +844,23 @@ if(ENABLE_DEVICE_LINKER) message(STATUS "Device Linker: filtered HIP_SOURCES (removed device .cpp under gensrc/ + subdirs, common.cu, onerank.cu, collectives.cc, dda_all_reduce_ipc.cu.cpp, dda_reduce_scatter_ipc.cu.cpp, dda_all_gather_ipc.cu.cpp, dda_alltoall_ipc.cu.cpp, dda_all_reduce_fabric.cu.cpp, dda_all_reduce_fabric_ll.cu.cpp, dda_reduce_scatter_fabric.cu.cpp, dda_all_gather_fabric.cu.cpp, dda_all_gather_fabric_ll.cu.cpp, dda_alltoall_fabric.cu.cpp)") endif() +# Keep the backend selection in one list for both compilation pipelines. +# DeviceLinker.cmake converts this list to -D flags because the sources are +# filtered out of the rccl target there; a normal build consumes the source +# properties directly. +if(ENABLE_ROCSHMEM_GIN) + set(RCCL_GIN_SDMA_COMPILE_DEFINITIONS + NCCL_GIN_ANVIL_SDMA_ENABLE=1 + NCCL_GIN_PROXY_ENABLE=0 + NCCL_GIN_ROCSHMEM_GDA_ENABLE=0) + foreach(_gin_sdma_source IN ITEMS + "${HIPIFY_DIR}/src/algorithms/gin/sdma/gin_alltoall_sdma.cu.cpp" + "${HIPIFY_DIR}/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu.cpp") + set_source_files_properties("${_gin_sdma_source}" PROPERTIES + COMPILE_DEFINITIONS "${RCCL_GIN_SDMA_COMPILE_DEFINITIONS}") + endforeach() +endif() + # Create an initial git_version.cpp file (that will be updated with latest git version) #================================================================================================== # Create initial empty file at configure time diff --git a/projects/rccl/src/include/poison_hip_atomics.h b/projects/rccl/src/include/poison_hip_atomics.h index 22c92beafd18..844766bf59a7 100644 --- a/projects/rccl/src/include/poison_hip_atomics.h +++ b/projects/rccl/src/include/poison_hip_atomics.h @@ -32,6 +32,15 @@ #include #endif +// The Anvil-SDMA backend deliberately calls rocSHMEM device code. Parse that +// external header before poisoning the HIP builtins so its implementation is +// outside this RCCL-only check; the include guard keeps it from being parsed +// again when gin_anvil_sdma.h includes it. RCCL code following that include is +// still checked. Only the two explicitly opted-in GIN-SDMA TUs take this arm. +#if defined(NCCL_GIN_ANVIL_SDMA_ENABLE) && NCCL_GIN_ANVIL_SDMA_ENABLE +#include "sdma/anvil_device.hpp" +#endif + // clang-format off // clang-format has no notion of a line-continued #pragma: it joins the whole // directive onto one ~300-column line. Keep it off across this block so the diff --git a/projects/rccl/test/CMakeLists.txt b/projects/rccl/test/CMakeLists.txt index 1e6d9b2d828c..1b6a1d5a902c 100644 --- a/projects/rccl/test/CMakeLists.txt +++ b/projects/rccl/test/CMakeLists.txt @@ -955,10 +955,15 @@ if(BUILD_TESTS) target_include_directories(${test_executable} PRIVATE ${RCCL_COMMON_INCLUDE_DIRS}) target_compile_definitions(${test_executable} PRIVATE ${RCCL_COMMON_COMPILE_DEFS}) - # Belt-and-suspenders: Anvil-SDMA device headers are opt-in per TU. Keep the - # explicit 0 on non-stub tests so ENABLE_ROCSHMEM_GIN cannot leak them back in. - if(ENABLE_ROCSHMEM_GIN AND NOT "${test_executable}" IN_LIST RCCL_ANVIL_SDMA_STUB_TARGETS) - target_compile_definitions(${test_executable} PRIVATE NCCL_GIN_ANVIL_SDMA_ENABLE=0) + # Belt-and-suspenders: Anvil-SDMA device headers are opt-in. Stub-backed + # tests need an explicit 1 now that gin_device_common.h defaults to 0; + # everything else stays explicitly opted out. + if(ENABLE_ROCSHMEM_GIN) + if("${test_executable}" IN_LIST RCCL_ANVIL_SDMA_STUB_TARGETS) + target_compile_definitions(${test_executable} PRIVATE NCCL_GIN_ANVIL_SDMA_ENABLE=1) + else() + target_compile_definitions(${test_executable} PRIVATE NCCL_GIN_ANVIL_SDMA_ENABLE=0) + endif() endif() if(ENABLE_DEVICE_LINKER) # Mirror the rccl library define so device-linker-specific test expectations compile in. @@ -1060,9 +1065,9 @@ if(BUILD_TESTS) # rocshmem::envvar::log_flags undefined at load time. Link host rocSHMEM # into the Fixtures binary so the executable satisfies those symbols when # the dynamic loader brings in librccl.so. - if(ENABLE_ROCSHMEM_GIN AND GIN_ANVIL_UNIT_TESTS AND NOT ENABLE_ROCSHMEM - AND ROCSHMEM_INSTALL_DIR AND ("${test_executable}" STREQUAL "rccl-UnitTestsFixtures" - OR "${test_executable}" STREQUAL "rccl-UnitTestsFixturesDebug")) + if(ENABLE_ROCSHMEM_GIN AND NOT ENABLE_ROCSHMEM AND ROCSHMEM_INSTALL_DIR + AND (("${test_executable}" STREQUAL "rccl-UnitTestsFixtures" AND GIN_ANVIL_UNIT_TESTS) + OR "${test_executable}" STREQUAL "rccl-UnitTestsFixturesDebug")) set(_ROCSHMEM_FIXTURES_OFFLOAD_FLAGS "") foreach(_g IN LISTS GPU_TARGETS) if(_g) From 097f4da38c3b8dd505039efd626cdfb511e805d6 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sun, 13 Sep 2026 23:35:08 -0400 Subject: [PATCH 7/8] rccl: add external Anvil poison-boundary test Compile a guarded stand-in for rocSHMEM's anvil_device.hpp with a __hip_atomic_load call and verify that the explicitly opted-in external header is parsed before RCCL activates its atomics poison. Co-authored-by: Cursor --- .../src/include/test_poison_hip_atomics.py | 38 ++++++++++++++++++- 1 file changed, 37 insertions(+), 1 deletion(-) diff --git a/projects/rccl/src/include/test_poison_hip_atomics.py b/projects/rccl/src/include/test_poison_hip_atomics.py index 8c946090b6cc..dcdf5390d40d 100644 --- a/projects/rccl/src/include/test_poison_hip_atomics.py +++ b/projects/rccl/src/include/test_poison_hip_atomics.py @@ -129,7 +129,7 @@ def _probe_src(call): """) -def _compile(src, *, poison, offload): +def _compile(src, *, poison, offload, extra_args=()): """Return (returncode, combined stdout+stderr).""" cmd = [CXX, "-x", "hip", "-nogpulib", "-fsyntax-only"] if offload == "host": @@ -138,6 +138,7 @@ def _compile(src, *, poison, offload): cmd += ["--offload-device-only", "--offload-arch=gfx942"] else: raise ValueError(offload) + cmd += list(extra_args) if poison: cmd += [f"--include={POISON_H}"] with tempfile.NamedTemporaryFile("w", suffix=".cpp", delete=False) as f: @@ -197,6 +198,41 @@ def test_hip_compiler_available(self): @unittest.skipUnless(CXX, "amdclang++/hipcc not found (set ROCM_PATH or HIPCXX)") class PoisonHipAtomicsTest(unittest.TestCase): + def _compile_with_external_anvil_stub(self, src): + """Compile with a stand-in for rocSHMEM's guarded Anvil header.""" + with tempfile.TemporaryDirectory() as tmp: + sdma_dir = os.path.join(tmp, "sdma") + os.makedirs(sdma_dir) + with open(os.path.join(sdma_dir, "anvil_device.hpp"), "w") as f: + f.write(textwrap.dedent("""\ + #ifndef TEST_ANVIL_DEVICE_HPP_ + #define TEST_ANVIL_DEVICE_HPP_ + __device__ __forceinline__ unsigned int external_load(unsigned int* p) { + return __hip_atomic_load( + p, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } + #endif + """)) + return _compile( + src, + poison=True, + offload="device", + extra_args=( + "-DNCCL_GIN_ANVIL_SDMA_ENABLE=1", + f"-I{tmp}", + ), + ) + + def test_external_anvil_header_parsed_before_poison(self): + src = textwrap.dedent("""\ + #include "sdma/anvil_device.hpp" + __global__ void k(unsigned int *p, unsigned int *out) { + *out = external_load(p); + } + """) + rc, out = self._compile_with_external_anvil_stub(src) + self.assertEqual(rc, 0, msg=out) + def test_hip_atomic_load_rejected_on_device(self): rc, out = _compile( _probe_src(_call_expr("__hip_atomic_load")), From c0b537c3fd7008038410c3fe3a653bf7db771603 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sun, 13 Sep 2026 23:35:20 -0400 Subject: [PATCH 8/8] rccl: add post-Anvil atomics poison test After parsing the external Anvil stub, spell __hip_atomic_load directly in the RCCL probe and require the compiler to reject it. This ensures the external header exception does not disable the poison for following RCCL code. Co-authored-by: Cursor --- .../rccl/src/include/test_poison_hip_atomics.py | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/projects/rccl/src/include/test_poison_hip_atomics.py b/projects/rccl/src/include/test_poison_hip_atomics.py index dcdf5390d40d..571dcb1a7685 100644 --- a/projects/rccl/src/include/test_poison_hip_atomics.py +++ b/projects/rccl/src/include/test_poison_hip_atomics.py @@ -233,6 +233,19 @@ def test_external_anvil_header_parsed_before_poison(self): rc, out = self._compile_with_external_anvil_stub(src) self.assertEqual(rc, 0, msg=out) + def test_poison_still_rejects_rccl_code_after_external_anvil_header(self): + src = textwrap.dedent("""\ + #include "sdma/anvil_device.hpp" + __global__ void k(unsigned int *p, unsigned int *out) { + *out = __hip_atomic_load( + p, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } + """) + rc, out = self._compile_with_external_anvil_stub(src) + self.assertNotEqual(rc, 0, msg=out) + self.assertIn("poisoned identifier", out, msg=out) + self.assertIn("__hip_atomic_load", out, msg=out) + def test_hip_atomic_load_rejected_on_device(self): rc, out = _compile( _probe_src(_call_expr("__hip_atomic_load")),