From ac159b94cfe78d256e5b1de0dcbf21ce9153cf2f Mon Sep 17 00:00:00 2001 From: QizhouZhang97 Date: Thu, 29 Oct 2026 07:55:07 +0000 Subject: [PATCH 01/13] ci: run MORI's CI, images and wheels on ROCm 10.0 Move every CI lane from ROCm 7.14 (and 7.2.4) to ROCm 10.0. The base is rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0, which docker/Dockerfile.dev now defaults to together with the 10.0 apt repo, so the separate intranode-test-rocm10 gate is dropped: intranode-test itself runs on ROCm 10 now, on both MI355X and MI300X. The JAX job moves off the from-source jaxlib 0.9.1 build on a ROCm 7.11 base to AMD's rocm/jax 10.0 image (docker/Dockerfile.jax), and the nightly wheels build against ROCm 10.0's RHEL 8 packages for py3.11 and py3.12 -- ROCm 10.0 ships no py3.10 PyTorch image, and 3.10 is end-of-life. setup.py baked hipconfig's version into the package, which on ROCm 10.0 is HIP's 7.15 rather than the release, so every `import mori` warned that it was built against another ROCm than the rocm_sdk it runs on. It now reads rocm-core's rocm_version.h first, and a test checks the baked pattern against the installed rocm_sdk. Co-authored-by: Cursor --- .claude/skills/deploy-mori/SKILL.md | 30 ++++--- .github/workflows/ci.yml | 107 +++--------------------- .github/workflows/ci_cco.yml | 2 +- .github/workflows/codeql.yml | 2 +- .github/workflows/nightly.yml | 45 +++++----- docker/Dockerfile.dev | 27 +++--- docker/Dockerfile.jax | 22 +++++ docker/Dockerfile.jax091-rocm711 | 70 ---------------- docker/Dockerfile.manylinux | 14 ++-- setup.py | 22 +++++ src/application/CMakeLists.txt | 2 +- tests/python/cco/test_rocm_bootstrap.py | 11 +++ tools/bench_ep_asyncll_compare.sh | 6 +- 13 files changed, 135 insertions(+), 225 deletions(-) create mode 100644 docker/Dockerfile.jax delete mode 100644 docker/Dockerfile.jax091-rocm711 diff --git a/.claude/skills/deploy-mori/SKILL.md b/.claude/skills/deploy-mori/SKILL.md index 346c286c15..f002110d1c 100644 --- a/.claude/skills/deploy-mori/SKILL.md +++ b/.claude/skills/deploy-mori/SKILL.md @@ -52,7 +52,7 @@ done ## Step 1: Start the Docker container Use the container name from the user's args if provided; ask if not. -Default image: `rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0` +Default image: `rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0` (use unless the user specifies another) — same base the CI workflows build on. ```bash @@ -427,35 +427,43 @@ mst status -v ## Step 4: Install MORI -### 4.0 — ROCm 7.14 dev toolchain (rocm7.14 base images only) +### 4.0 — ROCm dev toolchain (rocm10.x / rocm7.14 base images only) **Skip on pre-7.14 images** — they already ship a full `/opt/rocm` and must not get a second ROCm stacked on top. -The rocm7.14 pytorch image has **no `/opt/rocm` at all**: ROCm comes from the pip -`rocm-sdk-core` wheel, which ships runtime `.so` files and headers but no CMake -package configs. Without this step the build dies at configure time with +The rocm10.x and rocm7.14 pytorch images have **no `/opt/rocm` at all**: ROCm +comes from the pip `rocm-sdk-core` wheel, which ships runtime `.so` files and +headers but no CMake package configs. Without this step the build dies at +configure time with `does not contain the HIP runtime CMake package ... hip-lang-config.cmake`. +Install the release in the image's tag — mixing them makes every `import mori` +warn that it was built against another ROCm. For ROCm 10.0: + ```bash sudo docker exec $CONTAINER_NAME bash -c ' set -e [ -d /opt/rocm/bin ] && { echo "/opt/rocm already present — skipping"; exit 0; } apt-get install -y --no-install-recommends wget gnupg ca-certificates mkdir -p --mode=0755 /etc/apt/keyrings -wget -qO - https://repo.amd.com/rocm/packages-multi-arch/gpg/rocm.gpg \ +wget -qO - https://stable.repo.amd.com/rocm/gpg/packages.gpg \ | gpg --dearmor > /etc/apt/keyrings/amdrocm.gpg -echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/amdrocm.gpg] https://repo.amd.com/rocm/packages-multi-arch/ubuntu2404 stable main" \ +echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/amdrocm.gpg] https://stable.repo.amd.com/rocm/core/packages/ubuntu2404 stable main" \ > /etc/apt/sources.list.d/amdrocm.list apt-get update -apt-get install -y --no-install-recommends amdrocm-core-dev7.14-gfx942 amdrocm-core-dev7.14-gfx950 -hipconfig --version +apt-get install -y --no-install-recommends amdrocm-core-dev10.0-gfx942 amdrocm-core-dev10.0-gfx950 +cat /opt/rocm/core/.info/version ' ``` +A rocm7.14 image takes `amdrocm-core-dev7.14-` from the +`https://repo.amd.com/rocm/packages-multi-arch/ubuntu2404` repo instead, signed +by `https://repo.amd.com/rocm/packages-multi-arch/gpg/rocm.gpg`. + No environment setup is needed for the rest of the session. CMake resolves ROCm -as `-DROCM_PATH` → `$ROCM_PATH` → `$HIP_PATH` → `/opt/rocm` (where both the -7.2.4 and 7.14 stacks live) and pins the compiler to an absolute +as `-DROCM_PATH` → `$ROCM_PATH` → `$HIP_PATH` → `/opt/rocm` (where the 7.2.4, +7.14 and 10.x stacks all live) and pins the compiler to an absolute `/llvm/bin/amdclang++`; mori's libraries carry `RUNPATH "$ORIGIN:/lib"`, and the JIT builds absolute `/bin/...` paths off the same default. Set `ROCM_PATH` only for ROCm somewhere else (e.g. a diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index a795528708..7ac2f77c6c 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -31,7 +31,7 @@ concurrency: env: IMAGE: rocm/mori:ci - BASE_IMAGE: rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 + BASE_IMAGE: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 CONTAINER: mori_ci_${{ github.run_id }} CT: docker @@ -83,6 +83,12 @@ jobs: $CT exec $CONTAINER \ git config --global --add safe.directory $GITHUB_WORKSPACE + - name: Report ROCm version + run: | + $CT exec $CONTAINER bash -c " + ls -d /opt/rocm/core-* && /opt/rocm/bin/hipconfig --version && echo + " + - name: Build and install mori run: $CT exec $CONTAINER pip install . @@ -454,97 +460,6 @@ jobs: chown -R $(id -u):$(id -g) $GITHUB_WORKSPACE 2>/dev/null || true fi - # Gate for "does mori still build and run on the next ROCm major". Kept to the - # single EP intranode test on one runner: the rest of the intranode suite is - # ROCm-version agnostic, and the ROCm 10 base image is a separate ~30GB pull. - intranode-test-rocm10: - name: mori intranode test (ROCm 10, MI355X_AINIC) - needs: build - runs-on: [self-hosted, MI355X-AINIC-PIT-G17] - env: - # py3.11 rather than the 7.14 lane's py3.12 — this is the tag the ROCm 10 - # build was validated on. - BASE_IMAGE: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 - IMAGE: rocm/mori:ci-rocm10 - CONTAINER: mori_ci_rocm10_${{ github.run_id }} - MORI_RDMA_DEVICES: rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7 - MORI_RDMA_SL: 3 - MORI_RDMA_TC: 104 - # Pinned for the same reason as intranode-test's socket_ifname, see above. - MORI_SOCKET_IFNAME: eno0 - NCCL_SOCKET_IFNAME: eno0 - timeout-minutes: 60 - steps: - - name: Fix workspace permissions - run: | - WS="${GITHUB_WORKSPACE}" - if [ -d "$WS" ]; then - $CT run --rm -v "$WS:$WS" "$BASE_IMAGE" \ - chown -R $(id -u):$(id -g) "$WS" 2>/dev/null || true - fi - - - name: Checkout - uses: actions/checkout@v4 - with: - submodules: true - # A re-run (run_attempt > 1) checks out the PR merged with the current - # main; otherwise it would reuse the merge commit of the first attempt. - ref: ${{ github.event_name == 'pull_request' && github.run_attempt > 1 && format('refs/pull/{0}/merge', github.event.pull_request.number) || '' }} - - # ROCm 10 ships in a different apt repo than 7.x, hence the ROCM_APT_* args. - - name: Build CI image - run: | - $CT build --network=host \ - --build-arg BASE_IMAGE=$BASE_IMAGE \ - --build-arg ROCM_APT_VERSION=10.0 \ - --build-arg ROCM_APT_GPG=https://stable.repo.amd.com/rocm/gpg/packages.gpg \ - --build-arg ROCM_APT_LIST="https://stable.repo.amd.com/rocm/core/packages/ubuntu2404 stable main" \ - -t $IMAGE -f docker/Dockerfile.dev . - - - name: Start container - run: | - $CT rm -f $CONTAINER 2>/dev/null || true - CONTAINER_RUNTIME=$CT ./docker/ci_run.sh --name $CONTAINER \ - -e MORI_RDMA_DEVICES=$MORI_RDMA_DEVICES \ - -e MORI_RDMA_SL=$MORI_RDMA_SL \ - -e MORI_RDMA_TC=$MORI_RDMA_TC \ - -e MORI_SOCKET_IFNAME=$MORI_SOCKET_IFNAME \ - -e NCCL_SOCKET_IFNAME=$NCCL_SOCKET_IFNAME \ - -v $GITHUB_WORKSPACE:$GITHUB_WORKSPACE \ - -w $GITHUB_WORKSPACE \ - $IMAGE sleep infinity - $CT exec $CONTAINER \ - git config --global --add safe.directory $GITHUB_WORKSPACE - - - name: Report ROCm version - run: | - $CT exec $CONTAINER bash -c " - ls -d /opt/rocm/core-* && /opt/rocm/bin/hipconfig --version && echo - " - - - name: Install mori - run: | - $CT exec $CONTAINER bash -c \ - "cd $GITHUB_WORKSPACE && pip install ." - - - name: Verify installation - run: $CT exec $CONTAINER bash -c "cd $GITHUB_WORKSPACE && python tools/verify_install.py" - - - name: MORI-EP (intranode) - run: | - $CT exec -e PYTHONPATH=$GITHUB_WORKSPACE $CONTAINER bash -c " - cd $GITHUB_WORKSPACE && timeout 360 pytest tests/python/ops/test_dispatch_combine_intranode.py -v - " - - - name: Cleanup - if: always() - run: | - $CT rm -f $CONTAINER || true - if $CT image inspect $IMAGE &>/dev/null; then - $CT run --rm -v $GITHUB_WORKSPACE:$GITHUB_WORKSPACE $IMAGE \ - chown -R $(id -u):$(id -g) $GITHUB_WORKSPACE 2>/dev/null || true - fi - jax-intranode-test: name: mori JAX intranode test (${{ matrix.platform }}) needs: build @@ -556,13 +471,11 @@ jobs: - platform: MI355X_AINIC runner: [self-hosted, MI355X-AINIC-PIT-G17] rdma_devices: rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7 - gpu_arch: gfx950 socket_ifname: eno0 # Runner 325_bnxt_108 is decommissioned. # - platform: MI325X_BNXT # runner: [self-hosted, 325_bnxt_108] # rdma_devices: bnxt_re0,bnxt_re1,bnxt_re2,bnxt_re3,bnxt_re4,bnxt_re5,bnxt_re7,bnxt_re8 - # gpu_arch: gfx942 # socket_ifname: enp49s0f1np1 env: # jax.distributed rendezvouses over localhost, but mori.jax.shmem_init_attr @@ -590,13 +503,13 @@ jobs: - name: Build JAX CI image run: | - BUILD_ARGS="--network=host --build-arg GPU_TARGET=${{ matrix.gpu_arch }}" + BUILD_ARGS="--network=host" if [ "$CT" = "podman" ]; then BUILD_ARGS="$BUILD_ARGS --isolation=chroot" fi $CT build $BUILD_ARGS \ -t rocm/mori:jax-ci \ - -f docker/Dockerfile.jax091-rocm711 . + -f docker/Dockerfile.jax . - name: Start container run: | @@ -619,7 +532,7 @@ jobs: $CT exec ${CONTAINER}_jax bash -c \ "cd $GITHUB_WORKSPACE && \ export XLA_FFI_INCLUDE_DIR=\$(python3 -c 'from jax import ffi; print(ffi.include_dir())') && \ - BUILD_UMBP=OFF BUILD_XLA_FFI_OPS=ON pip install --break-system-packages ." + BUILD_UMBP=OFF BUILD_XLA_FFI_OPS=ON pip install ." - name: MORI-EP JAX test run: | diff --git a/.github/workflows/ci_cco.yml b/.github/workflows/ci_cco.yml index 418accc4eb..821f8428f6 100644 --- a/.github/workflows/ci_cco.yml +++ b/.github/workflows/ci_cco.yml @@ -12,7 +12,7 @@ concurrency: env: IMAGE: rocm/mori:ci-cco - BASE_IMAGE: rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 + BASE_IMAGE: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 CONTAINER: mori_cco_ci_${{ github.run_id }} CT: podman # MI355X-AINIC runner is podman rootless (no docker daemon) diff --git a/.github/workflows/codeql.yml b/.github/workflows/codeql.yml index 106c4d5eae..c417193ca0 100644 --- a/.github/workflows/codeql.yml +++ b/.github/workflows/codeql.yml @@ -32,7 +32,7 @@ concurrency: env: IMAGE: rocm/mori:ci - BASE_IMAGE: rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 + BASE_IMAGE: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 CONTAINER: mori_codeql_${{ github.run_id }} CT: docker # falls back to cpp,python on push events (where inputs.languages is empty) diff --git a/.github/workflows/nightly.yml b/.github/workflows/nightly.yml index ad48266c8f..c28108df70 100644 --- a/.github/workflows/nightly.yml +++ b/.github/workflows/nightly.yml @@ -31,16 +31,13 @@ jobs: # Built in docker/Dockerfile.manylinux, so a wheel needs nothing newer # than glibc 2.28 instead of the glibc/libstdc++ of the Ubuntu image that # ships its Python. Each lane builds against the ROCm its test image - # below runs: 7.2.4 for py3.10, 7.14 for py3.12. - - python: "3.10" - rocm: "7.2.4" - rocm_build_args: >- - --build-arg ROCM_VERSION=7.2.4 - --build-arg ROCM_REPO=https://repo.radeon.com/rocm/rhel8/7.2.4/main - --build-arg ROCM_GPG=https://repo.radeon.com/rocm/rocm.gpg.key - --build-arg "ROCM_PACKAGES=rocm-hip-runtime-devel rocm-smi-lib" + # below runs: 10.0 for both, the Dockerfile's default (its header lists + # the args for other ROCm releases). + - python: "3.11" + rocm: "10.0" + rocm_build_args: "" - python: "3.12" - rocm: "7.14" + rocm: "10.0" rocm_build_args: "" steps: # Must precede Checkout: later steps run as root inside the CI container @@ -187,9 +184,9 @@ jobs: # TCP connections from the host itself, so rank 1 spins in "retrying # connection to root" for 1500 attempts instead of failing. - platform: MI355X_AINIC - python: "3.10" + python: "3.11" runner: [self-hosted, MI355X-AINIC-PIT-G17] - base_image: rocm/pytorch:rocm7.2.4_ubuntu22.04_py3.10_pytorch_release_2.8.0 + base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 rdma_devices: rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7 rdma_sl: 3 rdma_tc: 104 @@ -199,7 +196,7 @@ jobs: - platform: MI355X_AINIC python: "3.12" runner: [self-hosted, MI355X-AINIC-PIT-G17] - base_image: rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 + base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.12_pytorch_release_2.13.0 rdma_devices: rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7 rdma_sl: 3 rdma_tc: 104 @@ -208,12 +205,12 @@ jobs: # 8x ConnectX-7 NDR on InfiniBand, so no rdma_sl/rdma_tc: mlx5 applies # both only on the RoCE path. mlx5_8 is deliberately not in # rdma_devices -- it is the 100GbE port behind ens14np0 (the TCP - # rendezvous interface), not one of the 8 rails. py3.12 only: that base + # rendezvous interface), not one of the 8 rails. py3.11 only: that base # image is what ci.yml validates this runner against. - platform: MI300X_MLX - python: "3.12" + python: "3.11" runner: [self-hosted, 300_mlx_p19_29] - base_image: rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 + base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 rdma_devices: mlx5_0,mlx5_1,mlx5_2,mlx5_3,mlx5_4,mlx5_5,mlx5_6,mlx5_7 socket_ifname: ens14np0 timeout_minutes: 90 @@ -434,9 +431,9 @@ jobs: matrix: include: - platform: MI355X_AINIC - python: "3.10" + python: "3.11" runner: [self-hosted, MI355X-AINIC] - base_image: rocm/pytorch:rocm7.2.4_ubuntu22.04_py3.10_pytorch_release_2.8.0 + base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 node1_host: 10.2.80.20 node2_host: 10.2.80.10 node2_port: 22 @@ -450,7 +447,7 @@ jobs: - platform: MI355X_AINIC python: "3.12" runner: [self-hosted, MI355X-AINIC] - base_image: rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 + base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.12_pytorch_release_2.13.0 node1_host: 10.2.80.20 node2_host: 10.2.80.10 node2_port: 22 @@ -469,9 +466,9 @@ jobs: # it drives over ssh (banff-ccs-aus-p19-14). Both addresses are the # ens14np0 100GbE port used for TCP rendezvous. - platform: MI300X_MLX - python: "3.12" + python: "3.11" runner: [self-hosted, 300_mlx_p19_29] - base_image: rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 + base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 node1_host: 10.194.132.110 node2_host: 10.194.132.71 node2_port: 22 @@ -934,10 +931,10 @@ jobs: fail-fast: false matrix: include: - - python: "3.10" - base_image: rocm/pytorch:rocm7.2.4_ubuntu22.04_py3.10_pytorch_release_2.8.0 + - python: "3.11" + base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 - python: "3.12" - base_image: rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 + base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.12_pytorch_release_2.13.0 env: IMAGE: rocm/mori:ci-py${{ matrix.python }} steps: @@ -1017,7 +1014,7 @@ jobs: MORI_RDMA_SL: 3 MORI_RDMA_TC: 104 IMAGE: rocm/mori:ci-flydsl-${{ matrix.flydsl }} - BASE_IMAGE: rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 + BASE_IMAGE: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 CONTAINER: mori_flydsl_compat_${{ github.run_id }}_${{ strategy.job-index }} CT: podman # MI355X-AINIC runner is podman rootless (no docker daemon) timeout-minutes: 90 diff --git a/docker/Dockerfile.dev b/docker/Dockerfile.dev index d15e757ee8..f2879ac76f 100644 --- a/docker/Dockerfile.dev +++ b/docker/Dockerfile.dev @@ -1,21 +1,23 @@ # Tested base images (the apt ROCm install below must match the base's ROCm -# major; mixing them would end up with two ROCm installs stacked): -# rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 (default) -# rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 (ROCM_APT_* below) -ARG BASE_IMAGE=rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 +# version: the base's rocm_sdk wheels warn on every import of a mori built +# against another one, and a different major stacks two ROCm installs): +# rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 (default) +# rocm/pytorch:rocm10.0_ubuntu24.04_py3.12_pytorch_release_2.13.0 +# rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 (ROCM_APT_* below) +ARG BASE_IMAGE=rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 FROM ${BASE_IMAGE} # Space-separated gfx targets to pull ROCm device packages for. The CCO CI # image uses "gfx950 gfx1250". ARG GPU_ARCHS="gfx942 gfx950" # apt ROCm to install when the base image ships no dev stack (see below). -# 10.0 lives in a different repo than 7.x, so the URLs are args too: -# ROCM_APT_VERSION=10.0 -# ROCM_APT_GPG=https://stable.repo.amd.com/rocm/gpg/packages.gpg -# ROCM_APT_LIST="https://stable.repo.amd.com/rocm/core/packages/ubuntu2404 stable main" -ARG ROCM_APT_VERSION=7.14 -ARG ROCM_APT_GPG=https://repo.amd.com/rocm/packages-multi-arch/gpg/rocm.gpg -ARG ROCM_APT_LIST="https://repo.amd.com/rocm/packages-multi-arch/ubuntu2404 stable main" +# 7.14 lives in a different repo than 10.x, so the URLs are args too: +# ROCM_APT_VERSION=7.14 +# ROCM_APT_GPG=https://repo.amd.com/rocm/packages-multi-arch/gpg/rocm.gpg +# ROCM_APT_LIST="https://repo.amd.com/rocm/packages-multi-arch/ubuntu2404 stable main" +ARG ROCM_APT_VERSION=10.0 +ARG ROCM_APT_GPG=https://stable.repo.amd.com/rocm/gpg/packages.gpg +ARG ROCM_APT_LIST="https://stable.repo.amd.com/rocm/core/packages/ubuntu2404 stable main" RUN apt-get update && \ apt-get install -y --no-install-recommends \ @@ -36,8 +38,7 @@ RUN apt-get update && \ # registers /opt/rocm/{bin,lib,include,...} via update-alternatives, the layout # mori expects. # -# Older base images (the 7.2.4 ones the py3.10 nightly lane still uses) already -# ship a full ROCm dev stack, and its apt repo is ubuntu2404-only anyway — so +# Older base images (the 7.2.x ones) already ship a full ROCm dev stack, so # probe for the same file mori's CMakeLists checks and skip if it is there. RUN set -ex; \ if [ -e /opt/rocm/lib/cmake/hip-lang/hip-lang-config.cmake ]; then \ diff --git a/docker/Dockerfile.jax b/docker/Dockerfile.jax new file mode 100644 index 0000000000..7777eadf0d --- /dev/null +++ b/docker/Dockerfile.jax @@ -0,0 +1,22 @@ +# JAX CI image: AMD's ROCm JAX release image plus what mori's build needs. +# +# ROCm comes from the image's pip rocm-sdk wheels, and rocm-sdk-devel carries a +# full dev tree (CMake configs, hsakmt, hipcc), so unlike docker/Dockerfile.dev no +# apt ROCm goes on top. The image's ROCM_PATH names the rocm-sdk-core tree, which +# has no CMake configs, so the devel tree is linked to /opt/rocm -- the root +# mori's build and runtime JIT default to -- and ROCM_PATH follows it. +ARG BASE_IMAGE=rocm/jax:rocm10.0-jax0.10.2-py3.11 +FROM ${BASE_IMAGE} + +RUN apt-get update && \ + apt-get install -y --no-install-recommends \ + libibverbs-dev libnuma-dev libdrm-dev && \ + rm -rf /var/lib/apt/lists/* + +RUN ln -s "$(python3 -m rocm_sdk path --root)" /opt/rocm && \ + test -e /opt/rocm/lib/cmake/hip-lang/hip-lang-config.cmake +ENV ROCM_PATH=/opt/rocm HIP_PATH=/opt/rocm ROCM_HOME=/opt/rocm + +# The amdgpu workaround docker/Dockerfile.dev explains; the symmetric heap +# allocates the same way under JAX. +ENV MALLOC_MMAP_THRESHOLD_=131072 diff --git a/docker/Dockerfile.jax091-rocm711 b/docker/Dockerfile.jax091-rocm711 deleted file mode 100644 index c687b7cbfc..0000000000 --- a/docker/Dockerfile.jax091-rocm711 +++ /dev/null @@ -1,70 +0,0 @@ -############################################################ -# Stage 1: Build jaxlib + ROCm plugin wheels from source -############################################################ -FROM ghcr.io/rocm/jax-base-ubu24.rocm711:latest AS builder - -ARG JAX_BRANCH=rocm-jaxlib-v0.9.1 -ARG BAZEL_VERSION=7.7.0 -ARG GPU_TARGET=gfx942,gfx950 -ARG ROCM_VERSION=7 - -# Install build tools -RUN apt-get update && apt-get install -y --no-install-recommends \ - wget \ - patchelf \ - clang-18 \ - libc6-dev && \ - apt-get clean && rm -rf /var/lib/apt/lists/* - -# Install Bazel -RUN wget -q -O /usr/local/bin/bazel \ - "https://github.com/bazelbuild/bazel/releases/download/${BAZEL_VERSION}/bazel-${BAZEL_VERSION}-linux-x86_64" && \ - chmod +x /usr/local/bin/bazel && \ - bazel --version - -# Clone ROCm JAX repo -RUN git clone --depth 1 --branch ${JAX_BRANCH} \ - https://github.com/ROCmSoftwarePlatform/jax.git /jax - -WORKDIR /jax - -# Build jaxlib + jax-rocm-plugin + jax-rocm-pjrt wheels -RUN ROCM_PATH=$(realpath /opt/rocm) && \ - CLANG_PATH=$(realpath /opt/rocm/llvm/bin/clang) && \ - python3 ./build/build.py build \ - --wheels=jaxlib,jax-rocm-plugin,jax-rocm-pjrt \ - --use_clang=true --clang_path=$CLANG_PATH \ - --rocm_version=${ROCM_VERSION} --rocm_path=$ROCM_PATH \ - --rocm_amdgpu_targets=${GPU_TARGET} - -RUN ls -lah /jax/dist/ - -############################################################ -# Stage 2: Clean runtime image -############################################################ -FROM ghcr.io/rocm/jax-base-ubu24.rocm711:latest - -ARG PYTHON_VERSION=3.12 - -RUN python${PYTHON_VERSION} -m pip install "jax==0.9.1" - -# Install the wheels built in stage 1 -COPY --from=builder /jax/dist/*.whl /tmp/wheels/ -RUN python${PYTHON_VERSION} -m pip install --break-system-packages \ - /tmp/wheels/*.whl && \ - rm -rf /tmp/wheels - -# Common runtime and test dependencies -RUN python${PYTHON_VERSION} -m pip install --break-system-packages \ - numpy scipy cmake pybind11 pytest - -RUN apt-get update && apt-get install -y --no-install-recommends \ - libibverbs-dev && \ - apt-get clean && rm -rf /var/lib/apt/lists/* - -# Sanity check (CPU-only at build time, GPU visible at runtime) -RUN python${PYTHON_VERSION} -c \ - "import jax; print('JAX', jax.__version__); print('Devices:', jax.devices())" \ - || true - -WORKDIR /workspace diff --git a/docker/Dockerfile.manylinux b/docker/Dockerfile.manylinux index f39464ea85..555f789935 100644 --- a/docker/Dockerfile.manylinux +++ b/docker/Dockerfile.manylinux @@ -5,8 +5,12 @@ # libstdc++ instead (manylinux_2_35 on Ubuntu 22.04, _2_39 on 24.04). # tools/build_manylinux_wheel.py is the build entry point. # -# ROCm comes from AMD's RHEL 8 repos. The defaults install the 7.14 multi-arch -# packages; the classic 7.2.x releases need: +# ROCm comes from AMD's RHEL 8 repos. The defaults install the 10.0 packages; +# the 7.14 multi-arch release needs: +# ROCM_VERSION=7.14 +# ROCM_REPO=https://repo.amd.com/rocm/packages-multi-arch/rhel8/x86_64 +# ROCM_GPG=https://repo.amd.com/rocm/packages-multi-arch/gpg/rocm.gpg +# and the classic 7.2.x releases: # ROCM_VERSION=7.2.4 # ROCM_REPO=https://repo.radeon.com/rocm/rhel8/7.2.4/main # ROCM_GPG=https://repo.radeon.com/rocm/rocm.gpg.key @@ -18,9 +22,9 @@ FROM ${BASE_IMAGE} # serve ROCm's own libraries, which mori does not link. The archs mori compiles # for come from MORI_GPU_ARCHS (see setup.py), so gfx1250 builds without them. ARG GPU_ARCHS="gfx942 gfx950" -ARG ROCM_VERSION=7.14 -ARG ROCM_REPO=https://repo.amd.com/rocm/packages-multi-arch/rhel8/x86_64 -ARG ROCM_GPG=https://repo.amd.com/rocm/packages-multi-arch/gpg/rocm.gpg +ARG ROCM_VERSION=10.0 +ARG ROCM_REPO=https://stable.repo.amd.com/rocm/core/packages/rhel8/x86_64 +ARG ROCM_GPG=https://stable.repo.amd.com/rocm/gpg/packages.gpg # Space-separated; {version} expands to ROCM_VERSION, {arch} to each GPU_ARCHS entry. ARG ROCM_PACKAGES="amdrocm-core-devel{version}-{arch}" # Still protobuf 21.12, the 3.21 line UMBP builds against on Ubuntu 24.04, and past diff --git a/setup.py b/setup.py index ca7b4f986b..f56a26d7f8 100644 --- a/setup.py +++ b/setup.py @@ -418,6 +418,9 @@ def _write_cco_build_flags(root_dir: Path, *, sdma: bool) -> None: _ROCM_VERSION_RE = re.compile(r"\b(\d+)\.(\d+)(?:\.\d+)?\b") +_ROCM_VERSION_HEADER_RE = re.compile( + r"^\s*#\s*define\s+ROCM_VERSION_(MAJOR|MINOR|PATCH)\s+(\d+)", re.MULTILINE +) def _detect_rocm_build_version() -> tuple[str | None, str | None]: @@ -428,9 +431,28 @@ def _detect_rocm_build_version() -> tuple[str | None, str | None]: ).resolve() candidates: list[str] = [] + # The ROCm release version, which is what rocm_sdk.__version__ reports. + # hipconfig reports HIP's own version, and that stopped tracking the release + # in ROCm 10.0 (it ships HIP 7.15), so it is only a last resort. + try: + header = (rocm_path / "include" / "rocm-core" / "rocm_version.h").read_text( + encoding="utf-8" + ) + except OSError: + pass + else: + fields = dict(_ROCM_VERSION_HEADER_RE.findall(header)) + if "MAJOR" in fields and "MINOR" in fields: + candidates.append( + ".".join(fields[k] for k in ("MAJOR", "MINOR", "PATCH") if k in fields) + ) + + # .info/ sits at the root of the classic layout and under core/ in the + # multi-arch (7.14+) one. for version_file in ( rocm_path / ".info" / "version", rocm_path / ".info" / "version-dev", + rocm_path / "core" / ".info" / "version", ): try: candidates.append(version_file.read_text(encoding="utf-8")) diff --git a/src/application/CMakeLists.txt b/src/application/CMakeLists.txt index ed0ee1b205..a804bf62d0 100644 --- a/src/application/CMakeLists.txt +++ b/src/application/CMakeLists.txt @@ -10,7 +10,7 @@ if(WITH_MPI) endif() find_package(hsa-runtime64 REQUIRED) -# ROCm 7.14's hsakmtTargets.cmake leaks its build machine into +# The hsakmtTargets.cmake of ROCm 7.14 and 10.0 leaks its build machine into # hsakmt::hsakmt's INTERFACE_LINK_LIBRARIES; undone here and after # find_package(hsakmt) below (both no-ops on 7.2.4, which is clean). # diff --git a/tests/python/cco/test_rocm_bootstrap.py b/tests/python/cco/test_rocm_bootstrap.py index 3fb996734b..65f3a950b3 100644 --- a/tests/python/cco/test_rocm_bootstrap.py +++ b/tests/python/cco/test_rocm_bootstrap.py @@ -129,3 +129,14 @@ def test_baked_version_uses_major_minor_compatibility_pattern(): major, minor, *_ = bootstrap.ROCM_BUILD_VERSION.split(".") assert bootstrap.ROCM_VERSION_PATTERN == f"{major}.{minor}.*" + + +def test_baked_version_matches_the_installed_rocm_sdk(): + # A mismatch only warns at import time, from inside rocm_sdk, so nothing + # else in the suite would notice it. + if bootstrap.ROCM_VERSION_PATTERN is None: + pytest.skip("ROCm build metadata is generated while building the package") + rocm_sdk = pytest.importorskip("rocm_sdk") + + major, minor, *_ = rocm_sdk.__version__.split(".") + assert bootstrap.ROCM_VERSION_PATTERN == f"{major}.{minor}.*" diff --git a/tools/bench_ep_asyncll_compare.sh b/tools/bench_ep_asyncll_compare.sh index c6c76eceb0..019fe9d9d8 100755 --- a/tools/bench_ep_asyncll_compare.sh +++ b/tools/bench_ep_asyncll_compare.sh @@ -148,11 +148,13 @@ PY # Recorded because two numbers from two machines are not comparable without it, # and because two of the three known environment issues change these results # materially. See the known-issues skill. +# .info/ sits under core/ in the multi-arch (7.14+) layout. +ROCM_RELEASE=$(cat /opt/rocm/.info/version /opt/rocm/core/.info/version 2>/dev/null | head -1 || true) { echo "# collected $(date -u +%Y-%m-%dT%H:%M:%SZ)" echo "host: $(hostname)" echo "kernel: $(uname -r)" - echo "rocm: $(cat /opt/rocm/.info/version 2>/dev/null || echo unknown)" + echo "rocm: ${ROCM_RELEASE:-unknown}" echo "git_commit: $(git -C "$REPO_ROOT" rev-parse --short HEAD 2>/dev/null || echo unknown)" echo "git_dirty: $(test -n "$(git -C "$REPO_ROOT" status --porcelain 2>/dev/null)" && echo yes || echo no)" echo "" @@ -194,7 +196,7 @@ print(f'visible_devices: {torch.cuda.device_count()}') fi # Issue 3: on ROCm 7.2.x the VMM path ignores the uncached request, costing # cco combine ~10-15us. The tell is that forcing cached changes nothing. - echo "issue3: ROCm $(cat /opt/rocm/.info/version 2>/dev/null | cut -d. -f1-2) -- 7.0-7.2.x is affected, 7.14+ is not" + echo "issue3: ROCm $(echo "$ROCM_RELEASE" | cut -d. -f1-2) -- 7.0-7.2.x is affected, 7.14+ is not" } > "$ENVFILE" 2>&1 { From 1328cac503fe3df7c03b60ae6b3dbd4024e767fc Mon Sep 17 00:00:00 2001 From: QizhouZhang97 Date: Thu, 29 Oct 2026 08:03:17 +0000 Subject: [PATCH 02/13] ci: keep a ROCm 7.14 intranode gate on MI355X Everything else now runs on ROCm 10.0, so nothing would notice a change that only builds or runs on it. The new intranode-test-rocm714 job builds what intranode-test builds on the 7.14 base and apt ROCm, checks the install, and runs the EP intranode test, on the MI355X_AINIC runner only. Co-authored-by: Cursor --- .github/workflows/ci.yml | 91 ++++++++++++++++++++++++++++++++++++++++ 1 file changed, 91 insertions(+) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 7ac2f77c6c..6877484ecb 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -460,6 +460,97 @@ jobs: chown -R $(id -u):$(id -g) $GITHUB_WORKSPACE 2>/dev/null || true fi + # Gate for "does mori still build and run on ROCm 7.14", the release before + # the 10.0 the other jobs run on. Kept to the EP intranode test on one runner: + # the rest of the intranode suite is ROCm-version agnostic. + intranode-test-rocm714: + name: mori intranode test (ROCm 7.14, MI355X_AINIC) + needs: build + runs-on: [self-hosted, MI355X-AINIC-PIT-G17] + env: + BASE_IMAGE: rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 + IMAGE: rocm/mori:ci-rocm714 + CONTAINER: mori_ci_rocm714_${{ github.run_id }} + MORI_RDMA_DEVICES: rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7 + MORI_RDMA_SL: 3 + MORI_RDMA_TC: 104 + # Pinned for the same reason as intranode-test's socket_ifname, see above. + MORI_SOCKET_IFNAME: eno0 + NCCL_SOCKET_IFNAME: eno0 + timeout-minutes: 60 + steps: + - name: Fix workspace permissions + run: | + WS="${GITHUB_WORKSPACE}" + if [ -d "$WS" ]; then + $CT run --rm -v "$WS:$WS" "$BASE_IMAGE" \ + chown -R $(id -u):$(id -g) "$WS" 2>/dev/null || true + fi + + - name: Checkout + uses: actions/checkout@v4 + with: + submodules: true + # A re-run (run_attempt > 1) checks out the PR merged with the current + # main; otherwise it would reuse the merge commit of the first attempt. + ref: ${{ github.event_name == 'pull_request' && github.run_attempt > 1 && format('refs/pull/{0}/merge', github.event.pull_request.number) || '' }} + + # ROCm 7.14 ships in a different apt repo than 10.x, hence the ROCM_APT_* args. + - name: Build CI image + run: | + $CT build --network=host \ + --build-arg BASE_IMAGE=$BASE_IMAGE \ + --build-arg ROCM_APT_VERSION=7.14 \ + --build-arg ROCM_APT_GPG=https://repo.amd.com/rocm/packages-multi-arch/gpg/rocm.gpg \ + --build-arg ROCM_APT_LIST="https://repo.amd.com/rocm/packages-multi-arch/ubuntu2404 stable main" \ + -t $IMAGE -f docker/Dockerfile.dev . + + - name: Start container + run: | + $CT rm -f $CONTAINER 2>/dev/null || true + CONTAINER_RUNTIME=$CT ./docker/ci_run.sh --name $CONTAINER \ + -e MORI_RDMA_DEVICES=$MORI_RDMA_DEVICES \ + -e MORI_RDMA_SL=$MORI_RDMA_SL \ + -e MORI_RDMA_TC=$MORI_RDMA_TC \ + -e MORI_SOCKET_IFNAME=$MORI_SOCKET_IFNAME \ + -e NCCL_SOCKET_IFNAME=$NCCL_SOCKET_IFNAME \ + -v $GITHUB_WORKSPACE:$GITHUB_WORKSPACE \ + -w $GITHUB_WORKSPACE \ + $IMAGE sleep infinity + $CT exec $CONTAINER \ + git config --global --add safe.directory $GITHUB_WORKSPACE + + - name: Report ROCm version + run: | + $CT exec $CONTAINER bash -c " + ls -d /opt/rocm/core-* && /opt/rocm/bin/hipconfig --version && echo + " + + # The same build as intranode-test's, so a ROCm 10-only API anywhere in it + # fails to compile here. + - name: Install mori + run: | + $CT exec $CONTAINER bash -c \ + "cd $GITHUB_WORKSPACE && BUILD_BENCHMARK=ON BUILD_CCO_SDMA=ON BUILD_EXAMPLES=ON pip install ." + + - name: Verify installation + run: $CT exec $CONTAINER bash -c "cd $GITHUB_WORKSPACE && python tools/verify_install.py" + + - name: MORI-EP (intranode) + run: | + $CT exec -e PYTHONPATH=$GITHUB_WORKSPACE $CONTAINER bash -c " + cd $GITHUB_WORKSPACE && timeout 360 pytest tests/python/ops/test_dispatch_combine_intranode.py -v + " + + - name: Cleanup + if: always() + run: | + $CT rm -f $CONTAINER || true + if $CT image inspect $IMAGE &>/dev/null; then + $CT run --rm -v $GITHUB_WORKSPACE:$GITHUB_WORKSPACE $IMAGE \ + chown -R $(id -u):$(id -g) $GITHUB_WORKSPACE 2>/dev/null || true + fi + jax-intranode-test: name: mori JAX intranode test (${{ matrix.platform }}) needs: build From 65280f5ca69cf508e2f4f8b6f5b99d2d84e2679c Mon Sep 17 00:00:00 2001 From: QizhouZhang97 Date: Fri, 30 Oct 2026 03:24:46 +0000 Subject: [PATCH 03/13] ci: move from ROCm 10.0 to 10.1.0 (torch 2.14) The default base is now rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 with the ROCm 10.1 apt toolchain. 10.1 drops rocm_smi, so this is also the first CI to run mori's amd_smi topology path. The 10.1 images run the system Python, which has no `python`, so docker/Dockerfile.dev adds one. The apt repo URL now takes the base's Ubuntu release ({ubuntu}), which lets the same Dockerfile build the ubuntu22.04 py3.10 image: with 10.1 the nightly wheel lanes are py3.10 and py3.12 again, as they were on 7.2.4/7.14. The rocm/jax 10.1 image ships no rocm-sdk-devel, so the JAX job builds docker/Dockerfile.dev on top of it like the torch images do, and docker/Dockerfile.jax goes away. Co-authored-by: Cursor --- .claude/skills/deploy-mori/SKILL.md | 7 +++--- .github/workflows/ci.yml | 17 ++++++++----- .github/workflows/ci_cco.yml | 2 +- .github/workflows/codeql.yml | 2 +- .github/workflows/nightly.yml | 38 ++++++++++++++--------------- docker/Dockerfile.dev | 23 +++++++++++------ docker/Dockerfile.jax | 22 ----------------- docker/Dockerfile.manylinux | 4 +-- src/application/CMakeLists.txt | 5 ++-- 9 files changed, 56 insertions(+), 64 deletions(-) delete mode 100644 docker/Dockerfile.jax diff --git a/.claude/skills/deploy-mori/SKILL.md b/.claude/skills/deploy-mori/SKILL.md index f002110d1c..b56efb8846 100644 --- a/.claude/skills/deploy-mori/SKILL.md +++ b/.claude/skills/deploy-mori/SKILL.md @@ -52,7 +52,7 @@ done ## Step 1: Start the Docker container Use the container name from the user's args if provided; ask if not. -Default image: `rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0` +Default image: `rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0` (use unless the user specifies another) — same base the CI workflows build on. ```bash @@ -439,7 +439,8 @@ configure time with `does not contain the HIP runtime CMake package ... hip-lang-config.cmake`. Install the release in the image's tag — mixing them makes every `import mori` -warn that it was built against another ROCm. For ROCm 10.0: +warn that it was built against another ROCm. For ROCm 10.1 (an Ubuntu 22.04 +image takes `ubuntu2204` in the repo URL): ```bash sudo docker exec $CONTAINER_NAME bash -c ' @@ -452,7 +453,7 @@ wget -qO - https://stable.repo.amd.com/rocm/gpg/packages.gpg \ echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/amdrocm.gpg] https://stable.repo.amd.com/rocm/core/packages/ubuntu2404 stable main" \ > /etc/apt/sources.list.d/amdrocm.list apt-get update -apt-get install -y --no-install-recommends amdrocm-core-dev10.0-gfx942 amdrocm-core-dev10.0-gfx950 +apt-get install -y --no-install-recommends amdrocm-core-dev10.1-gfx942 amdrocm-core-dev10.1-gfx950 cat /opt/rocm/core/.info/version ' ``` diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 6877484ecb..f2d8ccaaf0 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -31,7 +31,7 @@ concurrency: env: IMAGE: rocm/mori:ci - BASE_IMAGE: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 + BASE_IMAGE: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 CONTAINER: mori_ci_${{ github.run_id }} CT: docker @@ -460,9 +460,9 @@ jobs: chown -R $(id -u):$(id -g) $GITHUB_WORKSPACE 2>/dev/null || true fi - # Gate for "does mori still build and run on ROCm 7.14", the release before - # the 10.0 the other jobs run on. Kept to the EP intranode test on one runner: - # the rest of the intranode suite is ROCm-version agnostic. + # Gate for "does mori still build and run on ROCm 7.14"; every other job runs + # on 10.1. Kept to the EP intranode test on one runner: the rest of the + # intranode suite is ROCm-version agnostic. intranode-test-rocm714: name: mori intranode test (ROCm 7.14, MI355X_AINIC) needs: build @@ -574,6 +574,7 @@ jobs: # intranode-test. MORI_SOCKET_IFNAME: ${{ matrix.socket_ifname }} NCCL_SOCKET_IFNAME: ${{ matrix.socket_ifname }} + JAX_BASE_IMAGE: rocm/jax:rocm10.1.0_ubuntu24.04_py3.12_jax_release_0.10.2 timeout-minutes: 90 steps: - name: Fix workspace permissions @@ -592,6 +593,8 @@ jobs: # main; otherwise it would reuse the merge commit of the first attempt. ref: ${{ github.event_name == 'pull_request' && github.run_attempt > 1 && format('refs/pull/{0}/merge', github.event.pull_request.number) || '' }} + # The JAX image has no ROCm dev stack either, so it takes the same + # Dockerfile as the torch images. - name: Build JAX CI image run: | BUILD_ARGS="--network=host" @@ -599,8 +602,9 @@ jobs: BUILD_ARGS="$BUILD_ARGS --isolation=chroot" fi $CT build $BUILD_ARGS \ + --build-arg BASE_IMAGE=$JAX_BASE_IMAGE \ -t rocm/mori:jax-ci \ - -f docker/Dockerfile.jax . + -f docker/Dockerfile.dev . - name: Start container run: | @@ -620,10 +624,11 @@ jobs: # Resolve the XLA FFI headers from the system jax *outside* pip's # isolated build env (build isolation hides system site-packages, so # CMake can't `import jax` during configure). Pass it via env var. + # The JAX image ships no pytest. $CT exec ${CONTAINER}_jax bash -c \ "cd $GITHUB_WORKSPACE && \ export XLA_FFI_INCLUDE_DIR=\$(python3 -c 'from jax import ffi; print(ffi.include_dir())') && \ - BUILD_UMBP=OFF BUILD_XLA_FFI_OPS=ON pip install ." + BUILD_UMBP=OFF BUILD_XLA_FFI_OPS=ON pip install . && pip install pytest" - name: MORI-EP JAX test run: | diff --git a/.github/workflows/ci_cco.yml b/.github/workflows/ci_cco.yml index 821f8428f6..eca08fa966 100644 --- a/.github/workflows/ci_cco.yml +++ b/.github/workflows/ci_cco.yml @@ -12,7 +12,7 @@ concurrency: env: IMAGE: rocm/mori:ci-cco - BASE_IMAGE: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 + BASE_IMAGE: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 CONTAINER: mori_cco_ci_${{ github.run_id }} CT: podman # MI355X-AINIC runner is podman rootless (no docker daemon) diff --git a/.github/workflows/codeql.yml b/.github/workflows/codeql.yml index c417193ca0..b77595a121 100644 --- a/.github/workflows/codeql.yml +++ b/.github/workflows/codeql.yml @@ -32,7 +32,7 @@ concurrency: env: IMAGE: rocm/mori:ci - BASE_IMAGE: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 + BASE_IMAGE: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 CONTAINER: mori_codeql_${{ github.run_id }} CT: docker # falls back to cpp,python on push events (where inputs.languages is empty) diff --git a/.github/workflows/nightly.yml b/.github/workflows/nightly.yml index c28108df70..722bd26c9a 100644 --- a/.github/workflows/nightly.yml +++ b/.github/workflows/nightly.yml @@ -31,13 +31,13 @@ jobs: # Built in docker/Dockerfile.manylinux, so a wheel needs nothing newer # than glibc 2.28 instead of the glibc/libstdc++ of the Ubuntu image that # ships its Python. Each lane builds against the ROCm its test image - # below runs: 10.0 for both, the Dockerfile's default (its header lists + # below runs: 10.1 for both, the Dockerfile's default (its header lists # the args for other ROCm releases). - - python: "3.11" - rocm: "10.0" + - python: "3.10" + rocm: "10.1" rocm_build_args: "" - python: "3.12" - rocm: "10.0" + rocm: "10.1" rocm_build_args: "" steps: # Must precede Checkout: later steps run as root inside the CI container @@ -184,9 +184,9 @@ jobs: # TCP connections from the host itself, so rank 1 spins in "retrying # connection to root" for 1500 attempts instead of failing. - platform: MI355X_AINIC - python: "3.11" + python: "3.10" runner: [self-hosted, MI355X-AINIC-PIT-G17] - base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.14.0 rdma_devices: rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7 rdma_sl: 3 rdma_tc: 104 @@ -196,7 +196,7 @@ jobs: - platform: MI355X_AINIC python: "3.12" runner: [self-hosted, MI355X-AINIC-PIT-G17] - base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.12_pytorch_release_2.13.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 rdma_devices: rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7 rdma_sl: 3 rdma_tc: 104 @@ -205,12 +205,12 @@ jobs: # 8x ConnectX-7 NDR on InfiniBand, so no rdma_sl/rdma_tc: mlx5 applies # both only on the RoCE path. mlx5_8 is deliberately not in # rdma_devices -- it is the 100GbE port behind ens14np0 (the TCP - # rendezvous interface), not one of the 8 rails. py3.11 only: that base + # rendezvous interface), not one of the 8 rails. py3.12 only: that base # image is what ci.yml validates this runner against. - platform: MI300X_MLX - python: "3.11" + python: "3.12" runner: [self-hosted, 300_mlx_p19_29] - base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 rdma_devices: mlx5_0,mlx5_1,mlx5_2,mlx5_3,mlx5_4,mlx5_5,mlx5_6,mlx5_7 socket_ifname: ens14np0 timeout_minutes: 90 @@ -431,9 +431,9 @@ jobs: matrix: include: - platform: MI355X_AINIC - python: "3.11" + python: "3.10" runner: [self-hosted, MI355X-AINIC] - base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.14.0 node1_host: 10.2.80.20 node2_host: 10.2.80.10 node2_port: 22 @@ -447,7 +447,7 @@ jobs: - platform: MI355X_AINIC python: "3.12" runner: [self-hosted, MI355X-AINIC] - base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.12_pytorch_release_2.13.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 node1_host: 10.2.80.20 node2_host: 10.2.80.10 node2_port: 22 @@ -466,9 +466,9 @@ jobs: # it drives over ssh (banff-ccs-aus-p19-14). Both addresses are the # ens14np0 100GbE port used for TCP rendezvous. - platform: MI300X_MLX - python: "3.11" + python: "3.12" runner: [self-hosted, 300_mlx_p19_29] - base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 node1_host: 10.194.132.110 node2_host: 10.194.132.71 node2_port: 22 @@ -931,10 +931,10 @@ jobs: fail-fast: false matrix: include: - - python: "3.11" - base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 + - python: "3.10" + base_image: rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.14.0 - python: "3.12" - base_image: rocm/pytorch:rocm10.0_ubuntu24.04_py3.12_pytorch_release_2.13.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 env: IMAGE: rocm/mori:ci-py${{ matrix.python }} steps: @@ -1014,7 +1014,7 @@ jobs: MORI_RDMA_SL: 3 MORI_RDMA_TC: 104 IMAGE: rocm/mori:ci-flydsl-${{ matrix.flydsl }} - BASE_IMAGE: rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 + BASE_IMAGE: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 CONTAINER: mori_flydsl_compat_${{ github.run_id }}_${{ strategy.job-index }} CT: podman # MI355X-AINIC runner is podman rootless (no docker daemon) timeout-minutes: 90 diff --git a/docker/Dockerfile.dev b/docker/Dockerfile.dev index f2879ac76f..65c04fcc9f 100644 --- a/docker/Dockerfile.dev +++ b/docker/Dockerfile.dev @@ -1,23 +1,25 @@ # Tested base images (the apt ROCm install below must match the base's ROCm # version: the base's rocm_sdk wheels warn on every import of a mori built # against another one, and a different major stacks two ROCm installs): -# rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 (default) -# rocm/pytorch:rocm10.0_ubuntu24.04_py3.12_pytorch_release_2.13.0 +# rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 (default) +# rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.14.0 +# rocm/jax:rocm10.1.0_ubuntu24.04_py3.12_jax_release_0.10.2 (JAX CI) # rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 (ROCM_APT_* below) -ARG BASE_IMAGE=rocm/pytorch:rocm10.0_ubuntu24.04_py3.11_pytorch_release_2.13.0 +ARG BASE_IMAGE=rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 FROM ${BASE_IMAGE} # Space-separated gfx targets to pull ROCm device packages for. The CCO CI # image uses "gfx950 gfx1250". ARG GPU_ARCHS="gfx942 gfx950" # apt ROCm to install when the base image ships no dev stack (see below). -# 7.14 lives in a different repo than 10.x, so the URLs are args too: +# {ubuntu} in ROCM_APT_LIST expands to the base's release (2404, 2204). 7.14 +# lives in a different repo than 10.x, so the URLs are args too: # ROCM_APT_VERSION=7.14 # ROCM_APT_GPG=https://repo.amd.com/rocm/packages-multi-arch/gpg/rocm.gpg # ROCM_APT_LIST="https://repo.amd.com/rocm/packages-multi-arch/ubuntu2404 stable main" -ARG ROCM_APT_VERSION=10.0 +ARG ROCM_APT_VERSION=10.1 ARG ROCM_APT_GPG=https://stable.repo.amd.com/rocm/gpg/packages.gpg -ARG ROCM_APT_LIST="https://stable.repo.amd.com/rocm/core/packages/ubuntu2404 stable main" +ARG ROCM_APT_LIST="https://stable.repo.amd.com/rocm/core/packages/ubuntu{ubuntu} stable main" RUN apt-get update && \ apt-get install -y --no-install-recommends \ @@ -29,9 +31,13 @@ RUN apt-get update && \ cmake && \ rm -rf /var/lib/apt/lists/* +# The 10.1 torch and jax images run the system Python, which has python3 but no +# python. +RUN command -v python >/dev/null || ln -s "$(command -v python3)" /usr/local/bin/python + # ── apt ROCm dev toolchain (only when the base image lacks one) ── # -# The rocm pytorch images have no /opt/rocm: ROCm comes from the pip +# The rocm pytorch and jax images have no /opt/rocm: ROCm comes from the pip # `rocm-sdk-core` wheel, which ships runtime .so files and headers but no CMake # package configs, so `project(... HIP)` cannot find hip-lang-config.cmake. # The apt ROCm below installs under /opt/rocm/core-${ROCM_APT_VERSION} and @@ -49,7 +55,8 @@ RUN set -ex; \ mkdir -p --mode=0755 /etc/apt/keyrings; \ wget -qO - "${ROCM_APT_GPG}" \ | gpg --dearmor > /etc/apt/keyrings/amdrocm.gpg; \ - echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/amdrocm.gpg] ${ROCM_APT_LIST}" \ + ubuntu="$(. /etc/os-release && echo "${VERSION_ID}" | tr -d .)"; \ + echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/amdrocm.gpg] $(echo "${ROCM_APT_LIST}" | sed "s/{ubuntu}/${ubuntu}/g")" \ > /etc/apt/sources.list.d/amdrocm.list; \ apt-get update; \ for arch in ${GPU_ARCHS}; do pkgs="${pkgs} amdrocm-core-dev${ROCM_APT_VERSION}-${arch}"; done; \ diff --git a/docker/Dockerfile.jax b/docker/Dockerfile.jax deleted file mode 100644 index 7777eadf0d..0000000000 --- a/docker/Dockerfile.jax +++ /dev/null @@ -1,22 +0,0 @@ -# JAX CI image: AMD's ROCm JAX release image plus what mori's build needs. -# -# ROCm comes from the image's pip rocm-sdk wheels, and rocm-sdk-devel carries a -# full dev tree (CMake configs, hsakmt, hipcc), so unlike docker/Dockerfile.dev no -# apt ROCm goes on top. The image's ROCM_PATH names the rocm-sdk-core tree, which -# has no CMake configs, so the devel tree is linked to /opt/rocm -- the root -# mori's build and runtime JIT default to -- and ROCM_PATH follows it. -ARG BASE_IMAGE=rocm/jax:rocm10.0-jax0.10.2-py3.11 -FROM ${BASE_IMAGE} - -RUN apt-get update && \ - apt-get install -y --no-install-recommends \ - libibverbs-dev libnuma-dev libdrm-dev && \ - rm -rf /var/lib/apt/lists/* - -RUN ln -s "$(python3 -m rocm_sdk path --root)" /opt/rocm && \ - test -e /opt/rocm/lib/cmake/hip-lang/hip-lang-config.cmake -ENV ROCM_PATH=/opt/rocm HIP_PATH=/opt/rocm ROCM_HOME=/opt/rocm - -# The amdgpu workaround docker/Dockerfile.dev explains; the symmetric heap -# allocates the same way under JAX. -ENV MALLOC_MMAP_THRESHOLD_=131072 diff --git a/docker/Dockerfile.manylinux b/docker/Dockerfile.manylinux index 555f789935..02a604ad6e 100644 --- a/docker/Dockerfile.manylinux +++ b/docker/Dockerfile.manylinux @@ -5,7 +5,7 @@ # libstdc++ instead (manylinux_2_35 on Ubuntu 22.04, _2_39 on 24.04). # tools/build_manylinux_wheel.py is the build entry point. # -# ROCm comes from AMD's RHEL 8 repos. The defaults install the 10.0 packages; +# ROCm comes from AMD's RHEL 8 repos. The defaults install the 10.1 packages; # the 7.14 multi-arch release needs: # ROCM_VERSION=7.14 # ROCM_REPO=https://repo.amd.com/rocm/packages-multi-arch/rhel8/x86_64 @@ -22,7 +22,7 @@ FROM ${BASE_IMAGE} # serve ROCm's own libraries, which mori does not link. The archs mori compiles # for come from MORI_GPU_ARCHS (see setup.py), so gfx1250 builds without them. ARG GPU_ARCHS="gfx942 gfx950" -ARG ROCM_VERSION=10.0 +ARG ROCM_VERSION=10.1 ARG ROCM_REPO=https://stable.repo.amd.com/rocm/core/packages/rhel8/x86_64 ARG ROCM_GPG=https://stable.repo.amd.com/rocm/gpg/packages.gpg # Space-separated; {version} expands to ROCM_VERSION, {arch} to each GPU_ARCHS entry. diff --git a/src/application/CMakeLists.txt b/src/application/CMakeLists.txt index a804bf62d0..bff1a597be 100644 --- a/src/application/CMakeLists.txt +++ b/src/application/CMakeLists.txt @@ -10,9 +10,10 @@ if(WITH_MPI) endif() find_package(hsa-runtime64 REQUIRED) -# The hsakmtTargets.cmake of ROCm 7.14 and 10.0 leaks its build machine into +# The hsakmtTargets.cmake of ROCm 7.14 and 10.x leaks its build machine into # hsakmt::hsakmt's INTERFACE_LINK_LIBRARIES; undone here and after -# find_package(hsakmt) below (both no-ops on 7.2.4, which is clean). +# find_package(hsakmt) below (both no-ops on 7.2.4, which is clean, and the +# second one on 10.1). # # It links the imported target `numa::numa` but never defines it, so # find_package(hsakmt) would fail at generate time: "The link interface of From bd98cbb6d1189d40587240c4ed22a3e93326b060 Mon Sep 17 00:00:00 2001 From: QizhouZhang97 Date: Fri, 30 Oct 2026 07:02:24 +0000 Subject: [PATCH 04/13] fix(topology): reuse the amd_smi already in the process On ROCm 10.1, where mori reads GPU topology through amd_smi, importing torch first (directly or through FlyDSL) has already loaded the rocm_sdk wheels' libamd_smi.so.27. OpenAmdSmi() then loaded a second copy from /opt/rocm, which binds part of its state to the first one's symbols and reports 0 sockets, so the rank exits with 255. On PR #871 every CCO Python example from 04_flydsl_lsa_put on died this way. The rocm_smi path never hit it: it opens librocm_smi64.so.1 by soname, which reuses a loaded copy. Reuse a libamd_smi already mapped into the process before opening one by path (MORI_AMD_SMI_PATH still wins). amd_smi reference-counts init and shut_down, so mori's pair leaves torch's session alone. Only the 10.1+ branch changes; 7.14 still builds the rocm_smi one. Co-authored-by: Cursor --- src/application/topology/gpu.cpp | 27 +++++++++++++++++++++++++++ 1 file changed, 27 insertions(+) diff --git a/src/application/topology/gpu.cpp b/src/application/topology/gpu.cpp index 38450d3a77..88671eb618 100644 --- a/src/application/topology/gpu.cpp +++ b/src/application/topology/gpu.cpp @@ -22,10 +22,13 @@ #include "mori/application/topology/gpu.hpp" #include +#include #include #include +#include #include +#include #include #include "mori/application/utils/check.hpp" @@ -42,7 +45,31 @@ namespace { #if MORI_USE_AMDSMI // ---- amd_smi path (ROCm 10.1+) ---- +// The libamd_smi already mapped into the process, if any: torch brings its own +// from the rocm_sdk wheels, and a second copy loaded beside it binds part of its +// state to the first one's symbols and then reports 0 sockets. amd_smi counts +// init/shut_down references, so sharing the instance leaves the other user's +// session alone. +void* LoadedAmdSmi() { + std::string path; + dl_iterate_phdr( + [](dl_phdr_info* info, size_t, void* data) { + const char* slash = info->dlpi_name ? strrchr(info->dlpi_name, '/') : nullptr; + if (slash && strncmp(slash + 1, "libamd_smi.so", strlen("libamd_smi.so")) == 0) { + *static_cast(data) = info->dlpi_name; + return 1; + } + return 0; + }, + &path); + return path.empty() ? nullptr : dlopen(path.c_str(), RTLD_NOW | RTLD_NOLOAD); +} + void* OpenAmdSmi() { + const char* explicitPath = std::getenv("MORI_AMD_SMI_PATH"); + if (!(explicitPath && explicitPath[0])) { + if (void* h = LoadedAmdSmi()) return h; + } const char* candidates[] = { std::getenv("MORI_AMD_SMI_PATH"), "libamd_smi.so", "/opt/rocm/lib/libamd_smi.so", std::getenv("MORI_ROCM_SMI_PATH"), "librocm_smi64.so.1", "librocm_smi64.so", From 4806fda7e32ac4cb8cf35117330e3a19d2d06007 Mon Sep 17 00:00:00 2001 From: QizhouZhang97 Date: Fri, 30 Oct 2026 07:02:29 +0000 Subject: [PATCH 05/13] ci: use torch 2.12 on the ROCm 10.1 images The 10.1 lanes ran the pytorch_release_2.14.0 tags. Nothing in mori needs a torch newer than 2.12, the one the 7.14 CI already ran, and torch 2.14's symm_mem.empty() puts its data 0x2400 bytes into the allocation, which ccoMemImport does not accept (examples/cco/python/08_torch_symm_import). Move every 10.1 image to the pytorch_release_2.12.0 tag of the same ROCm and Python. Co-authored-by: Cursor --- .claude/skills/deploy-mori/SKILL.md | 2 +- .github/workflows/ci.yml | 2 +- .github/workflows/ci_cco.yml | 2 +- .github/workflows/codeql.yml | 2 +- .github/workflows/nightly.yml | 18 +++++++++--------- docker/Dockerfile.dev | 6 +++--- 6 files changed, 16 insertions(+), 16 deletions(-) diff --git a/.claude/skills/deploy-mori/SKILL.md b/.claude/skills/deploy-mori/SKILL.md index b56efb8846..56fc02edfd 100644 --- a/.claude/skills/deploy-mori/SKILL.md +++ b/.claude/skills/deploy-mori/SKILL.md @@ -52,7 +52,7 @@ done ## Step 1: Start the Docker container Use the container name from the user's args if provided; ask if not. -Default image: `rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0` +Default image: `rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0` (use unless the user specifies another) — same base the CI workflows build on. ```bash diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index f2d8ccaaf0..8f4851a26c 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -31,7 +31,7 @@ concurrency: env: IMAGE: rocm/mori:ci - BASE_IMAGE: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 + BASE_IMAGE: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 CONTAINER: mori_ci_${{ github.run_id }} CT: docker diff --git a/.github/workflows/ci_cco.yml b/.github/workflows/ci_cco.yml index eca08fa966..abfdf86e8e 100644 --- a/.github/workflows/ci_cco.yml +++ b/.github/workflows/ci_cco.yml @@ -12,7 +12,7 @@ concurrency: env: IMAGE: rocm/mori:ci-cco - BASE_IMAGE: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 + BASE_IMAGE: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 CONTAINER: mori_cco_ci_${{ github.run_id }} CT: podman # MI355X-AINIC runner is podman rootless (no docker daemon) diff --git a/.github/workflows/codeql.yml b/.github/workflows/codeql.yml index b77595a121..91bc63107f 100644 --- a/.github/workflows/codeql.yml +++ b/.github/workflows/codeql.yml @@ -32,7 +32,7 @@ concurrency: env: IMAGE: rocm/mori:ci - BASE_IMAGE: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 + BASE_IMAGE: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 CONTAINER: mori_codeql_${{ github.run_id }} CT: docker # falls back to cpp,python on push events (where inputs.languages is empty) diff --git a/.github/workflows/nightly.yml b/.github/workflows/nightly.yml index 722bd26c9a..9f26dddc43 100644 --- a/.github/workflows/nightly.yml +++ b/.github/workflows/nightly.yml @@ -186,7 +186,7 @@ jobs: - platform: MI355X_AINIC python: "3.10" runner: [self-hosted, MI355X-AINIC-PIT-G17] - base_image: rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.14.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.12.0 rdma_devices: rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7 rdma_sl: 3 rdma_tc: 104 @@ -196,7 +196,7 @@ jobs: - platform: MI355X_AINIC python: "3.12" runner: [self-hosted, MI355X-AINIC-PIT-G17] - base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 rdma_devices: rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7 rdma_sl: 3 rdma_tc: 104 @@ -210,7 +210,7 @@ jobs: - platform: MI300X_MLX python: "3.12" runner: [self-hosted, 300_mlx_p19_29] - base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 rdma_devices: mlx5_0,mlx5_1,mlx5_2,mlx5_3,mlx5_4,mlx5_5,mlx5_6,mlx5_7 socket_ifname: ens14np0 timeout_minutes: 90 @@ -433,7 +433,7 @@ jobs: - platform: MI355X_AINIC python: "3.10" runner: [self-hosted, MI355X-AINIC] - base_image: rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.14.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.12.0 node1_host: 10.2.80.20 node2_host: 10.2.80.10 node2_port: 22 @@ -447,7 +447,7 @@ jobs: - platform: MI355X_AINIC python: "3.12" runner: [self-hosted, MI355X-AINIC] - base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 node1_host: 10.2.80.20 node2_host: 10.2.80.10 node2_port: 22 @@ -468,7 +468,7 @@ jobs: - platform: MI300X_MLX python: "3.12" runner: [self-hosted, 300_mlx_p19_29] - base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 node1_host: 10.194.132.110 node2_host: 10.194.132.71 node2_port: 22 @@ -932,9 +932,9 @@ jobs: matrix: include: - python: "3.10" - base_image: rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.14.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.12.0 - python: "3.12" - base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 + base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 env: IMAGE: rocm/mori:ci-py${{ matrix.python }} steps: @@ -1014,7 +1014,7 @@ jobs: MORI_RDMA_SL: 3 MORI_RDMA_TC: 104 IMAGE: rocm/mori:ci-flydsl-${{ matrix.flydsl }} - BASE_IMAGE: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 + BASE_IMAGE: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 CONTAINER: mori_flydsl_compat_${{ github.run_id }}_${{ strategy.job-index }} CT: podman # MI355X-AINIC runner is podman rootless (no docker daemon) timeout-minutes: 90 diff --git a/docker/Dockerfile.dev b/docker/Dockerfile.dev index 65c04fcc9f..e553002c6c 100644 --- a/docker/Dockerfile.dev +++ b/docker/Dockerfile.dev @@ -1,11 +1,11 @@ # Tested base images (the apt ROCm install below must match the base's ROCm # version: the base's rocm_sdk wheels warn on every import of a mori built # against another one, and a different major stacks two ROCm installs): -# rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 (default) -# rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.14.0 +# rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 (default) +# rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.12.0 # rocm/jax:rocm10.1.0_ubuntu24.04_py3.12_jax_release_0.10.2 (JAX CI) # rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 (ROCM_APT_* below) -ARG BASE_IMAGE=rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.14.0 +ARG BASE_IMAGE=rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 FROM ${BASE_IMAGE} # Space-separated gfx targets to pull ROCm device packages for. The CCO CI From 4a4df30145582aecd64252fe6cc74846a8ef40ed Mon Sep 17 00:00:00 2001 From: QizhouZhang97 Date: Fri, 30 Oct 2026 07:56:18 +0000 Subject: [PATCH 06/13] test(jax): give every worker all GPUs so RCCL can build its topology The JAX EP test pinned each of its 8 workers to one GPU through HIP_VISIBLE_DEVICES. On ROCm 10.1 the RCCL that JAX uses then fails to initialize the communicator for the test's own jax.lax.all_gather on every rank but one ("TopoPreset: Local rank N not found in intra-node graph. Topology misconfiguration."). That rank waits in RCCL init for the others, they wait for it at the next jit, and the job times out (exit 124 on PR #871). A pure-JAX 8-process all_gather hangs the same way; 2 or 4 processes do not. Let every worker see every GPU and hand JAX its own through local_device_ids, and set mori's HIP device explicitly, since mori runs on the calling thread's. Co-authored-by: Cursor --- tests/python/ops/test_dispatch_combine_jax.py | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/tests/python/ops/test_dispatch_combine_jax.py b/tests/python/ops/test_dispatch_combine_jax.py index c282fa2a03..ba7c1fe199 100644 --- a/tests/python/ops/test_dispatch_combine_jax.py +++ b/tests/python/ops/test_dispatch_combine_jax.py @@ -67,20 +67,27 @@ def _jax_worker( ): """Long-lived JAX worker: init distributed + mori, then serve tasks.""" gpus_per_proc = total_gpus // world_size - os.environ["HIP_VISIBLE_DEVICES"] = ",".join( - str(rank * gpus_per_proc + i) for i in range(gpus_per_proc) - ) + local_devices = [rank * gpus_per_proc + i for i in range(gpus_per_proc)] import jax + # Every process sees every GPU and JAX takes its own: RCCL on ROCm 10.1 cannot + # build its intra-node graph ("TopoPreset: Local rank N not found") when + # HIP_VISIBLE_DEVICES leaves each process one GPU, and the first XLA + # collective then hangs. jax.distributed.initialize( coordinator_address=f"localhost:{coordinator_port}", num_processes=world_size, process_id=rank, + local_device_ids=local_devices, ) import mori + from mori.jit.hip_driver import _check, _get_hip_lib + # mori runs on the calling thread's HIP device; with every GPU visible, that + # has to be this process's own. + _check(_get_hip_lib().hipSetDevice(local_devices[0]), "hipSetDevice") mori.jax.shmem_init_attr(rank, world_size) while True: From c1d4ce80a17cbcb7b647924d8ade2bf38aed1a55 Mon Sep 17 00:00:00 2001 From: QizhouZhang97 Date: Mon, 2 Nov 2026 02:16:25 +0000 Subject: [PATCH 07/13] oom report --- CCO_BENCH_OOM_mi355-gpu-51.md | 386 ++++++++++++++++++++++++++++++++++ 1 file changed, 386 insertions(+) create mode 100644 CCO_BENCH_OOM_mi355-gpu-51.md diff --git a/CCO_BENCH_OOM_mi355-gpu-51.md b/CCO_BENCH_OOM_mi355-gpu-51.md new file mode 100644 index 0000000000..d743d808b3 --- /dev/null +++ b/CCO_BENCH_OOM_mi355-gpu-51.md @@ -0,0 +1,386 @@ +# CCO ops-v2 bench 在 mi355-gpu-51 上报 OOM（ROCm 10.1）：交接说明 + +给在 mi355-gpu-51 上接手的 agent。目标：查清 PR #871 的 CCO CI 为什么在这台机器上 +报 `hipMemCreate failed: 2 (out of memory)`，并给出修复方向。下面的结论和数据都来自 +pit2-p03-g17（另一台 MI355X），在那台机器上复现不了。 + +## 1. 背景 + +- PR [ROCm/mori#871](https://github.com/ROCm/mori/pull/871)，分支 `dev/migrate_rocm10`， + head `4a4df301`。它把 CI 从 ROCm 7.14 迁到 ROCm 10.1，镜像是 + `rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0`。 +- 这是 pull_request 事件，CI checkout 的是 merge commit `cc973cdd` + （`4a4df301` 合进 main 的 `f3cebc3a`），不是 PR head。 +- 失败的 job 是 `ci_cco.yml` 的 "CCO unit test (MI355X_AINIC)"，runner 是 + `mi355-gpu-51`。这台 runner 用 rootless podman（workflow 里 `CT: podman`）， + 容器用 `docker/ci_run.sh` 起，默认带 `--privileged`。 +- CI 镜像 tag `rocm/mori:ci-cco` 是共用的：main 的 7.14 构建和 PR 的 10.1 构建都写它。 + +## 2. 现象 + +- 失败的步骤是 "Run ops-v2 dispatch/combine bench (EP8 smoke, both backends)"： + + ``` + cd tests/python/ops/dispatch_combine_v2 && \ + BACKENDS=flydsl,hip COMB_MODE=pull MODES=graph SWEEP=128,512,4096,8192 \ + ITERS=20 WARMUP=10 torchrun --standalone --nproc_per_node=8 bench_ep.py + ``` + +- torchrun 启动约 2.5 秒后，rank 2 在第一次分配上失败： + `[shmem] [error] ccoMemAlloc: hipMemCreate failed: 2 (out of memory)`。调用链是 + `bench_ep.py:308 build('flydsl')`、`flydsl_backend.py:162 SymmArena(comm, regions)`、 + `symm_arena.py:58 comm.alloc_mem(self._total)`、`communicator.py:429 alloc_mem`。 + 这块 arena 是 1.76 GB，GPU 有 288 GB。截图里只看到 rank 2 的 traceback，其他 rank + 是否也失败不清楚。 +- 同一个 job 里前面的 ops-v2 步骤（test、HIP backend、scales、parity、AsyncLL、 + routing snapshot）都通过，它们的 arena 都小于 1 GB。 +- 失败记录（时间都是 UTC）： + + | run | 代码 | 镜像 | bench 步骤 | 结果 | + |---|---|---|---|---| + | [41017647570](https://github.com/ROCm/mori/actions/runs/41017647570/job/123937501153) | PR `4806fda7` | 10.1 | 2026-10-30 07:46:45 至 07:46:49 | 失败 | + | [41023155247](https://github.com/ROCm/mori/actions/runs/41023155247/job/123954966536) | PR `4a4df301` | 10.1 | 2026-10-30 09:47:04 至 09:47:08 | 失败 | + | main 和其他分支，10-30 07:56 之后 | 各自的 | 7.14 | 跑到这一步 26 次，都在 mi355-gpu-51 上 | 26 次全部通过（另有 6 次被取消或没跑到这一步） | + + PR 失败之后紧接着的 main（10-30 08:51，7.14）也通过了，所以这台机器没有一直 + 处在坏状态。问题只在 10.1 镜像跑在 mi355-gpu-51 上时出现。 + +## 3. 在 pit2-p03-g17 上已经确认的（全部正常） + +- 裸 HIP VMM，属性和 `ccoMemAlloc` 一样（uncached、POSIX fd、gpuDirectRDMACapable）。 + 每次 32 GB 反复 alloc/free：`hipMemRelease` 之后显存立即回来；子进程退出后， + `waitpid` 返回时就全部回收。10.1 和 7.14 表现一样。 +- 走 MORI：2 rank，每次 16 GB，`alloc_mem`、`register_window`（P2P 导出和导入 fd）、 + 关闭、free，跑 10 轮。最慢 8 ms 回来，通信器销毁后回到起点。 +- 原样跑 CI 的 bench：PR head `4a4df301` 和 CI 实际测的 merge commit `cc973cdd` 都通过， + 两个后端 8/8 个点都验证通过。每卡峰值约 12 GB，每个 rank 两块 1.76 GB 的 arena。 +- CI 的 8 个 ops-v2 步骤连着跑也都通过。每一步开始时显存都回到空闲水平 + （0.3 到 0.5 GB），步骤之间没有残留进程。AsyncLL sdma 第一次失败在 + `hsaKmtCreateQueueExt`（SDMA 队列不够，同机的 CI job 在用），单独重跑就过了。 +- 用下面附录 A 的 `LD_PRELOAD` 拦截 KFD ioctl： + - 10.1 和 7.14 发给驱动的 VRAM 分配请求完全一样：uncached 是 `flags=0xb5000001`， + pinned 是 `0xb0000001`，没有 contiguous 位（bit 23）。 + - bench 进程里没有任何 SVM 注册（`AMDKFD_IOC_SVM`），加 `HSA_XNACK=1` 也没有。 + - 每个 rank 会有一条 `KFD ioctl nr=0x8 FAILED errno=28`，这是 CREATE_EVENT + 的事件槽用完了，runtime 会回退，可以忽略。 +- 本机环境：kernel `6.8.0-60-generic`，amdgpu-dkms `1:6.16.13.30300100-2303411.22.04` + （带下面说的 93aa919 修复），`amdgpu.noretry=-1`，docker，非 privileged。 + +## 4. 怀疑方向 + +**H1：amdgpu 的 SVM 重叠检查缺陷（`docker/Dockerfile.dev` 里有记录）。** 不带 VA 的 +VRAM 分配（VMM 的 `hipMemCreate` 就是这种）在 `kfd_ioctl_alloc_memory_of_gpu` 里 +会把 `[0, size)` 当成 SVM 重叠检查范围。只要有 SVM 注册的区间低于 `size`，内核就 +返回 EADDRINUSE，HIP 报成 out of memory。镜像里的 python3.12 不是 PIE，brk 堆在 +0.9 到 1 GB 附近，所以只有大于约 1 GB 的分配会撞上，这和"小 arena 通过、1.76 GB +的 bench 失败"对得上。上游修复是 93aa919（"drm/amdkfd: fix vram allocation failure +for a special case"），在 6.18 和 stable 6.17.8、6.12.58、6.6.117、6.1.159 里。 +DKMS 模块会覆盖发行版内核自带的 amdgpu，所以要看 amdgpu-dkms 的版本。 +疑点：pit2 上 10.1 的进程根本没有 SVM 注册，所以如果是 H1，51 上一定有别的东西在 +注册 SVM。 + +**H2：GPU 2 当时真的没有显存了。** 可能是同一个 job 里前面的步骤在 51 上留下了显存 +（比如 RDMA/dma-buf 引用在进程退出后没放），也可能是机器上别的进程。job 结束时 +容器被删，下一个 job 就恢复正常，这也能解释 main 紧接着能过。 + +两者用 errno 就能分开：H1 是 `EADDRINUSE`（98），H2 是 `ENOMEM`（12）。 + +## 5. 在 51 上要做的 + +### 5.0 注意 + +这是共享的 CI runner。动 GPU 前先确认没有 CI job 在跑： +`ps aux | grep '[R]unner.Worker'`。不要覆盖 `rocm/mori:ci-cco` 这个 tag，也不要用 +`mori_cco_ci_*` 的容器名，下面用 `localhost/mori-oom:*` 和 `mori_oom_*`。不要在 CI 的 +workspace（`/shared/amdgpu/home/jiahao_zhou_qle/actions-runner/_work/mori/mori`）里 +操作，另外 clone 一份。 + +### 5.1 不用 GPU 的检查（先做，可能直接定案） + +```bash +uname -r +cat /sys/module/amdgpu/version +dpkg -l | grep -E 'amdgpu-dkms ' || rpm -qa | grep amdgpu-dkms +cat /sys/module/amdgpu/parameters/noretry +# 93aa919 在不在。在的话这一行能找到： +grep -n "Skip a special case that allocates VRAM without VA" /usr/src/amdgpu-*/amd/amdkfd/kfd_chardev.c +# 两次失败时内核有没有报 SVM 冲突（H1 的直接证据）。时间带了 UTC 后缀，按 UTC 解释； +# 需要 persistent journal，如果机器之后重启过又没有，就只能看下面的 dmesg： +journalctl -k --since "2026-10-30 07:40 UTC" --until "2026-10-30 07:50 UTC" | grep -E "already allocated by SVM|amdgpu|kfd" +journalctl -k --since "2026-10-30 09:40 UTC" --until "2026-10-30 09:50 UTC" | grep -E "already allocated by SVM|amdgpu|kfd" +# 没有 journal 的话： +sudo dmesg -T | grep -E "already allocated by SVM" | tail +podman info --format '{{.Host.Security.Rootless}}'; ulimit -l +``` + +如果在那两个时间点看到 `Address: 0x0 already allocated by SVM`，就是 H1。 + +### 5.2 用 10.1 镜像原样跑 bench，并挂 spy + +```bash +git clone https://github.com/ROCm/mori.git mori-oom && cd mori-oom +git fetch origin pull/871/head:pr871 && git checkout pr871 # 4a4df301 +podman build --network=host \ + --build-arg BASE_IMAGE=rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 \ + --build-arg GPU_ARCHS="gfx950 gfx1250" --build-arg WITH_FLYDSL=1 \ + -t localhost/mori-oom:rocm101 -f docker/Dockerfile.dev . +mkdir -p /tmp/oom # 把附录 A 存成 /tmp/oom/kfd_svm_spy.c，附录 B 存成 /tmp/oom/replay_ops_v2.sh +CONTAINER_RUNTIME=podman ./docker/ci_run.sh --name mori_oom_101 \ + -e MORI_RDMA_DEVICES=rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7 \ + -e MORI_RDMA_SL=3 -e MORI_RDMA_TC=104 \ + -v $PWD:$PWD -v /tmp/oom:/w localhost/mori-oom:rocm101 sleep infinity +podman exec mori_oom_101 git config --global --add safe.directory $PWD +podman exec mori_oom_101 bash -c "cd $PWD && BUILD_TESTS=ON BUILD_EXAMPLES=ON BUILD_BENCHMARK=ON \ + MORI_WITH_MPI=ON BUILD_UMBP=OFF BUILD_CCO_SDMA=ON pip install ." +podman exec mori_oom_101 gcc -shared -fPIC -O2 -o /w/kfd_svm_spy.so /w/kfd_svm_spy.c -ldl + +amd-smi metric -m # 先记一下每张卡当前用了多少 +podman exec -e LD_PRELOAD=/w/kfd_svm_spy.so -e SPY_BT=5 mori_oom_101 bash -c \ + "cd $PWD/tests/python/ops/dispatch_combine_v2 && BACKENDS=flydsl,hip COMB_MODE=pull MODES=graph \ + SWEEP=128,512,4096,8192 ITERS=20 WARMUP=10 timeout 900 torchrun --standalone --nproc_per_node=8 bench_ep.py" \ + > /tmp/oom/bench_spy.log 2>&1; echo rc=$? +grep -E '^\[spy' /tmp/oom/bench_spy.log | sed -E 's/\[spy [0-9]+\]/[spy]/' | sort | uniq -c +``` + +pit2 上正常的输出只有这两种行：16 条 +`VRAM alloc va=0 size=1.76 GB flags=0xb5000001 ok, no SVM range below it`，和 8 条 +`KFD ioctl nr=0x8 FAILED errno=28`（可以忽略）。 + +### 5.3 spy 输出怎么看 + +| 失败 rank 上的 spy 输出 | 含义 | 下一步 | +|---|---|---| +| `VRAM alloc ... FAILED errno=98 (Address already in use)`，并且列出了一个低于 size 的 SVM 区间 | H1 | 升级 51 的 amdgpu-dkms 到带 93aa919 的版本。同时看 `SPY_BT` 打出的 backtrace，找出是谁注册的 SVM | +| `SVM op=0 start=...` 行，并且有 `WOULD FAIL on an amdgpu without 93aa919` | 有 SVM 注册（pit2 上没有），分配本身成功了 | 看 backtrace 和 7.14 的对比（5.5） | +| `VRAM alloc ... FAILED errno=12 (Cannot allocate memory)` | H2，显存真的不够 | `amd-smi process`、`amd-smi metric -m`、`podman ps -a`，找是谁占着 GPU 2 | +| `KFD reports only X GB available` | runtime 查到可用显存不够，在进内核前就拒绝了 | 同上 | +| 失败的那次分配没有任何 spy 输出 | 在 HIP/ROCr 里就被拒了，没到内核 | 加 `AMD_LOG_LEVEL=4` 重跑，看 HIP 那边的原因 | + +### 5.4 如果单独跑 bench 能过 + +CI 里 bench 前面还有很多步骤（CCO 单测、SDMA、benchmark、examples、allocator、 +GEMM+AR、ops-v2）。按 `ci_cco.yml` 里 cco-unit-test 的顺序在同一个容器里逐步跑， +同时在 host 上开附录 C 的显存采样，看哪一步之后 GPU 2 的显存不回到空闲水平。 +附录 B 只覆盖最后 8 个 ops-v2 步骤，可以先跑它；不行再把前面的步骤按 workflow +补上。 + +```bash +podman exec mori_oom_101 bash /w/replay_ops_v2.sh $PWD /w/replay_logs > /tmp/oom/replay_marks.log 2>&1 +``` + +### 5.5 对照：7.14 镜像 + +用 main 的 base 再建一个镜像，同样挂 spy 跑 bench，比较两边的 SVM 注册： + +```bash +podman build --network=host \ + --build-arg BASE_IMAGE=rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 \ + --build-arg GPU_ARCHS="gfx950 gfx1250" --build-arg WITH_FLYDSL=1 \ + -t localhost/mori-oom:rocm714 -f docker/Dockerfile.dev . +``` + +然后和 5.2 一样起容器 `mori_oom_714`、装 MORI、跑 bench。7.14 的镜像里 +`/opt/rocm/lib` 不在默认搜索路径上，单独跑的二进制需要 `LD_LIBRARY_PATH=/opt/rocm/lib`。 + +### 5.6 A/B 开关（在 10.1 上能复现之后再用） + +- `CCO_UNCACHED_WINDOW=0`：arena 用 pinned 而不是 uncached（flags 变成 `0xb0000001`）。 +- `CCO_GDR_CAPABLE=0`：分配不带 gpuDirectRDMACapable。 +- `SWEEP=128,512,4096`：最大 token 数从 8192 降到 4096，arena 大约减半。如果这样就能 + 过，说明和分配大小有关，支持 H1。 + +## 附录 A：kfd_svm_spy.c + +```c +// LD_PRELOAD shim: logs KFD SVM registrations, VA-less VRAM allocations and failed +// KFD ioctls (with errno). Flags every VA-less VRAM allocation that an amdgpu +// without 93aa919 would reject with EADDRINUSE: its SVM-overlap check runs over +// [0, size) when va_addr == 0. +// gcc -shared -fPIC -O2 -o kfd_svm_spy.so kfd_svm_spy.c -ldl +// LD_PRELOAD=./kfd_svm_spy.so SPY_BT=3 python ... +#define _GNU_SOURCE +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +struct svm_args { + uint64_t start_addr, size; + uint32_t op, nattr; +}; +struct alloc_args { + uint64_t va_addr, size, handle, mmap_offset; + uint32_t gpu_id, flags; +}; +struct avail_args { + uint64_t available; + uint32_t gpu_id, pad; +}; +#define NR_ALLOC 0x16 +#define NR_SVM 0x20 +#define NR_AVAIL 0x23 +#define VRAM_FLAG 1u + +#define MAX_RANGES 4096 +static uint64_t g_start[MAX_RANGES], g_len[MAX_RANGES]; +static int g_n, g_bt_left = -1, g_other_fail_left = 20; +static pthread_mutex_t g_mu = PTHREAD_MUTEX_INITIALIZER; + +static void Backtrace(void) { + if (g_bt_left < 0) g_bt_left = getenv("SPY_BT") ? atoi(getenv("SPY_BT")) : 0; + if (g_bt_left <= 0) return; + g_bt_left--; + void* frames[24]; + int n = backtrace(frames, 24); + backtrace_symbols_fd(frames + 2, n - 2, 2); +} + +static int LowestSvmBelow(uint64_t size) { + int hit = -1; + for (int i = 0; i < g_n; i++) + if (g_start[i] < size && (hit < 0 || g_start[i] < g_start[hit])) hit = i; + return hit; +} + +int ioctl(int fd, unsigned long req, ...) { + static int (*real)(int, unsigned long, void*); + if (!real) real = (int (*)(int, unsigned long, void*))dlsym(RTLD_NEXT, "ioctl"); + va_list ap; + va_start(ap, req); + void* arg = va_arg(ap, void*); + va_end(ap); + int ret = real(fd, req, arg); + int err = errno; + if (_IOC_TYPE(req) != 'K') return ret; + pthread_mutex_lock(&g_mu); + if (_IOC_NR(req) == NR_SVM) { + struct svm_args* a = (struct svm_args*)arg; + fprintf(stderr, "[spy %d] SVM op=%u start=0x%lx size=0x%lx nattr=%u ret=%d\n", getpid(), a->op, + a->start_addr, a->size, a->nattr, ret); + if (a->op == 0 && ret == 0 && g_n < MAX_RANGES) { + g_start[g_n] = a->start_addr; + g_len[g_n++] = a->size; + Backtrace(); + } + } else if (_IOC_NR(req) == NR_ALLOC) { + struct alloc_args* a = (struct alloc_args*)arg; + if (a->va_addr == 0 && (a->flags & VRAM_FLAG)) { + int hit = LowestSvmBelow(a->size); + if (ret != 0) + fprintf(stderr, + "[spy %d] VRAM alloc va=0 size=%.2f GB flags=0x%08x FAILED errno=%d (%s); lowest " + "SVM range below the size: %s0x%lx\n", + getpid(), a->size / 1073741824.0, a->flags, err, strerror(err), + hit >= 0 ? "" : "none, ", hit >= 0 ? g_start[hit] : 0); + else if (hit >= 0) + fprintf(stderr, + "[spy %d] VRAM alloc va=0 size=%.2f GB flags=0x%08x ok here, WOULD FAIL on an " + "amdgpu without 93aa919: SVM range [0x%lx, +0x%lx) is below the size\n", + getpid(), a->size / 1073741824.0, a->flags, g_start[hit], g_len[hit]); + else if (a->size >= (64ul << 20)) + fprintf(stderr, "[spy %d] VRAM alloc va=0 size=%.2f GB flags=0x%08x ok, no SVM range below it\n", + getpid(), a->size / 1073741824.0, a->flags); + } else if (ret != 0) { + fprintf(stderr, "[spy %d] alloc va=0x%lx size=0x%lx flags=0x%08x FAILED errno=%d (%s)\n", + getpid(), a->va_addr, a->size, a->flags, err, strerror(err)); + } + } else if (_IOC_NR(req) == NR_AVAIL) { + struct avail_args* a = (struct avail_args*)arg; + if (ret == 0 && a->available < (16ul << 30)) + fprintf(stderr, "[spy %d] KFD reports only %.2f GB available on gpu_id %u\n", getpid(), + a->available / 1073741824.0, a->gpu_id); + } else if (ret != 0 && g_other_fail_left > 0) { + g_other_fail_left--; + fprintf(stderr, "[spy %d] KFD ioctl nr=0x%x FAILED errno=%d (%s)\n", getpid(), + (unsigned)_IOC_NR(req), err, strerror(err)); + } + pthread_mutex_unlock(&g_mu); + errno = err; + return ret; +} + +int munmap(void* addr, size_t len) { + static int (*real)(void*, size_t); + if (!real) real = (int (*)(void*, size_t))dlsym(RTLD_NEXT, "munmap"); + int ret = real(addr, len); + int err = errno; + if (ret == 0) { + uint64_t lo = (uint64_t)addr, hi = lo + len; + pthread_mutex_lock(&g_mu); + for (int i = 0; i < g_n; i++) + if (g_start[i] >= lo && g_start[i] + g_len[i] <= hi) { + g_start[i] = g_start[--g_n]; + g_len[i] = g_len[g_n]; + i--; + } + pthread_mutex_unlock(&g_mu); + } + errno = err; + return ret; +} +``` + +## 附录 B：replay_ops_v2.sh（在容器里跑） + +```bash +#!/bin/bash +# Replays ci_cco.yml's ops-v2 steps (dispatch/combine test .. bench) back to back, +# inside the CI container. Usage: replay_ops_v2.sh +W=$1 +LOGS=$2 +D=$W/tests/python/ops/dispatch_combine_v2 +mark() { echo "@@ $(date +%T.%N | cut -c1-10) $*"; } +leftover() { echo "@@ leftover python procs: $(pgrep -c -f '[p]ython|[t]orchrun' || true)"; } +step() { + local name=$1; shift + mark "BEGIN $name" + (cd "$D" && env "$@") > "$LOGS/$name.log" 2>&1 + local rc=$? + mark "END $name rc=$rc" + leftover +} +mkdir -p "$LOGS" +step 24-pytest-intranode bash -c "cd $W && python -m pytest tests/python/ops/dispatch_combine_v2/test_dispatch_combine_v2_intranode.py -q" +step 25-hip-test-op MORI_V2_KERNEL_BACKEND=hip timeout 900 torchrun --standalone --nproc_per_node=8 test_op.py +step 26-hip-scales MORI_V2_KERNEL_BACKEND=hip SCALE_DIM=32 timeout 900 torchrun --standalone --nproc_per_node=8 test_op.py +for disp in fp8 fp4; do + step 27-asym-$disp MORI_V2_KERNEL_BACKEND=hip SCALE_DIM=224 HIDDEN=7168 DISP=$disp timeout 900 torchrun --standalone --nproc_per_node=8 test_asym_dtype.py +done +step 28-parity timeout 900 torchrun --standalone --nproc_per_node=8 test_ep_backend_parity.py +for transport in sdma lsa; do + step 29-asyncll-$transport MORI_ENABLE_SDMA=1 TRANSPORT=$transport timeout 900 torchrun --standalone --nproc_per_node=8 test_asyncll.py +done +step 30-routing-snapshot timeout 900 torchrun --standalone --nproc_per_node=8 test_routing_snapshot.py +step 31-bench BACKENDS=flydsl,hip COMB_MODE=pull MODES=graph SWEEP=128,512,4096,8192 ITERS=20 WARMUP=10 timeout 900 torchrun --standalone --nproc_per_node=8 bench_ep.py +mark DONE +``` + +## 附录 C：host 上的显存采样（每 0.5 秒一行，单位 MB，gpu0 到 gpu7） + +```bash +( while true; do echo "$(date +%T.%N | cut -c1-10) $(amd-smi metric -m --json 2>/dev/null | python3 -c " +import json,sys +d=json.load(sys.stdin); d=d.get('gpu_data',d) if isinstance(d,dict) else d +print(' '.join(str(g['mem_usage']['used_vram']['value']) for g in d))")" >> /tmp/oom/vram.log; sleep 0.5; done ) & +SAMPLER=$! +# ... 跑测试 ... +kill $SAMPLER +``` + +把 `replay_marks.log` 里每一步的 BEGIN/END 时间和 `vram.log` 对上，就能看出每一步 +开始时的显存和峰值。 + +## 附录 D：pit2 上用过的其他测试 + +在 pit2-p03-g17 的 `/var/zqz/zqz_vmm_churn/` 下，需要的话可以拷过来： +`vmm_churn.cpp`（裸 HIP VMM 反复 alloc/free，有同进程和跨进程两种模式）和 +`cco_churn.py`（MORI 的 `alloc_mem`、`register_window`、free 循环，用 mpirun 跑）。 +在 pit2 上这两个都显示释放及时。H2 的情况下可以在 51 上跑来对照。 From e4a4ccf9d4344df5958ec21c17fd8dfe34e11dfe Mon Sep 17 00:00:00 2001 From: qizzhang Date: Mon, 2 Nov 2026 01:09:36 -0600 Subject: [PATCH 08/13] ci: set HSA_USE_SVM=0 so pinned host buffers are not SVM ranges On ROCm 10.1 the CCO ops-v2 bench fails on the MI355X-AINIC runner with "ccoMemAlloc: hipMemCreate failed: 2 (out of memory)" on an empty GPU. The KFD returns EADDRINUSE, not ENOMEM: the runners' amdgpu-dkms (6.14.14) lacks 93aa919, so a VRAM allocation without a VA is checked against the process's SVM ranges over [0, size). The 10.1 HIP pins the bench's 256 KiB idx copies (idx.to(dev), idx.cpu()); 7.14 does not. On a KFD that offers the SVM API, the thunk registers each pinned buffer as an SVM range that outlives the pin, and these come from the brk heap of the non-PIE python, 0.2-1.1 GB up. The first 1.76 GB arena then fails on every rank. MALLOC_MMAP_THRESHOLD_ cannot prevent it: glibc still serves a 256 KiB request from a free heap chunk. HSA_USE_SVM=0 makes the thunk register host memory as userptr, so no SVM range exists. On the runner, with the CI image, the bench fails on all 8 ranks without it and passes with it (both backends, M=8192). Co-authored-by: Cursor --- docker/Dockerfile.dev | 14 ++++++++------ 1 file changed, 8 insertions(+), 6 deletions(-) diff --git a/docker/Dockerfile.dev b/docker/Dockerfile.dev index e553002c6c..dd1cc5a0a5 100644 --- a/docker/Dockerfile.dev +++ b/docker/Dockerfile.dev @@ -68,17 +68,19 @@ RUN set -ex; \ # fail with hipErrorOutOfMemory while VRAM is nearly empty. VMM allocations carry # no VA yet, so the thunk passes va_addr=0, but kfd_ioctl_alloc_memory_of_gpu # still runs its SVM-overlap check over [va_addr, va_addr+size) — with va_addr=0 -# that degenerates into "is anything SVM-registered below ?". A pageable -# host tensor that ROCr registers as SVM for an H2D copy lands on the brk heap at -# a low VA, hits that window, and the kernel returns EADDRINUSE. Pinning glibc's -# mmap threshold (i.e. disabling the dynamic one) keeps such buffers in the mmap -# region, well above any plausible allocation size. +# that degenerates into "is anything SVM-registered below ?". When the KFD +# offers the SVM API, the thunk registers each pageable buffer HIP pins for a +# copy as an SVM range, and the range outlives the pin. python here is not PIE, +# so a tensor copied from the brk heap leaves a range at a low VA and the kernel +# returns EADDRINUSE. glibc's mmap threshold cannot keep tensors off the heap: it +# still serves a larger request from a free heap chunk. HSA_USE_SVM=0 makes the +# thunk register host memory as userptr instead, so no SVM range exists. # # Upstream fix is 93aa919 ("drm/amdkfd: fix vram allocation failure for a special # case"), in 6.18 and stable 6.17.8 / 6.12.58 / 6.6.117 / 6.1.159. Drop this once # the runners' amdgpu-dkms carries it — note the DKMS module overrides the # distro kernel's amdgpu, so bumping the Ubuntu kernel alone does nothing. -ENV MALLOC_MMAP_THRESHOLD_=131072 +ENV HSA_USE_SVM=0 # Declared here, next to their only use: every RUN after an ARG has it in its # cache key, so up top a new FlyDSL version would rebuild the ROCm layers too. From 905b80b506c039df0b3b29c0b3e08eef3351f1de Mon Sep 17 00:00:00 2001 From: qizzhang Date: Wed, 4 Nov 2026 21:48:40 -0600 Subject: [PATCH 09/13] Drop the CCO bench OOM handoff notes Working notes from the investigation added in c1d4ce8; the findings are in e4a4ccf's message and next to HSA_USE_SVM in the workflows. Co-authored-by: Cursor --- CCO_BENCH_OOM_mi355-gpu-51.md | 386 ---------------------------------- 1 file changed, 386 deletions(-) delete mode 100644 CCO_BENCH_OOM_mi355-gpu-51.md diff --git a/CCO_BENCH_OOM_mi355-gpu-51.md b/CCO_BENCH_OOM_mi355-gpu-51.md deleted file mode 100644 index d743d808b3..0000000000 --- a/CCO_BENCH_OOM_mi355-gpu-51.md +++ /dev/null @@ -1,386 +0,0 @@ -# CCO ops-v2 bench 在 mi355-gpu-51 上报 OOM（ROCm 10.1）：交接说明 - -给在 mi355-gpu-51 上接手的 agent。目标：查清 PR #871 的 CCO CI 为什么在这台机器上 -报 `hipMemCreate failed: 2 (out of memory)`，并给出修复方向。下面的结论和数据都来自 -pit2-p03-g17（另一台 MI355X），在那台机器上复现不了。 - -## 1. 背景 - -- PR [ROCm/mori#871](https://github.com/ROCm/mori/pull/871)，分支 `dev/migrate_rocm10`， - head `4a4df301`。它把 CI 从 ROCm 7.14 迁到 ROCm 10.1，镜像是 - `rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0`。 -- 这是 pull_request 事件，CI checkout 的是 merge commit `cc973cdd` - （`4a4df301` 合进 main 的 `f3cebc3a`），不是 PR head。 -- 失败的 job 是 `ci_cco.yml` 的 "CCO unit test (MI355X_AINIC)"，runner 是 - `mi355-gpu-51`。这台 runner 用 rootless podman（workflow 里 `CT: podman`）， - 容器用 `docker/ci_run.sh` 起，默认带 `--privileged`。 -- CI 镜像 tag `rocm/mori:ci-cco` 是共用的：main 的 7.14 构建和 PR 的 10.1 构建都写它。 - -## 2. 现象 - -- 失败的步骤是 "Run ops-v2 dispatch/combine bench (EP8 smoke, both backends)"： - - ``` - cd tests/python/ops/dispatch_combine_v2 && \ - BACKENDS=flydsl,hip COMB_MODE=pull MODES=graph SWEEP=128,512,4096,8192 \ - ITERS=20 WARMUP=10 torchrun --standalone --nproc_per_node=8 bench_ep.py - ``` - -- torchrun 启动约 2.5 秒后，rank 2 在第一次分配上失败： - `[shmem] [error] ccoMemAlloc: hipMemCreate failed: 2 (out of memory)`。调用链是 - `bench_ep.py:308 build('flydsl')`、`flydsl_backend.py:162 SymmArena(comm, regions)`、 - `symm_arena.py:58 comm.alloc_mem(self._total)`、`communicator.py:429 alloc_mem`。 - 这块 arena 是 1.76 GB，GPU 有 288 GB。截图里只看到 rank 2 的 traceback，其他 rank - 是否也失败不清楚。 -- 同一个 job 里前面的 ops-v2 步骤（test、HIP backend、scales、parity、AsyncLL、 - routing snapshot）都通过，它们的 arena 都小于 1 GB。 -- 失败记录（时间都是 UTC）： - - | run | 代码 | 镜像 | bench 步骤 | 结果 | - |---|---|---|---|---| - | [41017647570](https://github.com/ROCm/mori/actions/runs/41017647570/job/123937501153) | PR `4806fda7` | 10.1 | 2026-10-30 07:46:45 至 07:46:49 | 失败 | - | [41023155247](https://github.com/ROCm/mori/actions/runs/41023155247/job/123954966536) | PR `4a4df301` | 10.1 | 2026-10-30 09:47:04 至 09:47:08 | 失败 | - | main 和其他分支，10-30 07:56 之后 | 各自的 | 7.14 | 跑到这一步 26 次，都在 mi355-gpu-51 上 | 26 次全部通过（另有 6 次被取消或没跑到这一步） | - - PR 失败之后紧接着的 main（10-30 08:51，7.14）也通过了，所以这台机器没有一直 - 处在坏状态。问题只在 10.1 镜像跑在 mi355-gpu-51 上时出现。 - -## 3. 在 pit2-p03-g17 上已经确认的（全部正常） - -- 裸 HIP VMM，属性和 `ccoMemAlloc` 一样（uncached、POSIX fd、gpuDirectRDMACapable）。 - 每次 32 GB 反复 alloc/free：`hipMemRelease` 之后显存立即回来；子进程退出后， - `waitpid` 返回时就全部回收。10.1 和 7.14 表现一样。 -- 走 MORI：2 rank，每次 16 GB，`alloc_mem`、`register_window`（P2P 导出和导入 fd）、 - 关闭、free，跑 10 轮。最慢 8 ms 回来，通信器销毁后回到起点。 -- 原样跑 CI 的 bench：PR head `4a4df301` 和 CI 实际测的 merge commit `cc973cdd` 都通过， - 两个后端 8/8 个点都验证通过。每卡峰值约 12 GB，每个 rank 两块 1.76 GB 的 arena。 -- CI 的 8 个 ops-v2 步骤连着跑也都通过。每一步开始时显存都回到空闲水平 - （0.3 到 0.5 GB），步骤之间没有残留进程。AsyncLL sdma 第一次失败在 - `hsaKmtCreateQueueExt`（SDMA 队列不够，同机的 CI job 在用），单独重跑就过了。 -- 用下面附录 A 的 `LD_PRELOAD` 拦截 KFD ioctl： - - 10.1 和 7.14 发给驱动的 VRAM 分配请求完全一样：uncached 是 `flags=0xb5000001`， - pinned 是 `0xb0000001`，没有 contiguous 位（bit 23）。 - - bench 进程里没有任何 SVM 注册（`AMDKFD_IOC_SVM`），加 `HSA_XNACK=1` 也没有。 - - 每个 rank 会有一条 `KFD ioctl nr=0x8 FAILED errno=28`，这是 CREATE_EVENT - 的事件槽用完了，runtime 会回退，可以忽略。 -- 本机环境：kernel `6.8.0-60-generic`，amdgpu-dkms `1:6.16.13.30300100-2303411.22.04` - （带下面说的 93aa919 修复），`amdgpu.noretry=-1`，docker，非 privileged。 - -## 4. 怀疑方向 - -**H1：amdgpu 的 SVM 重叠检查缺陷（`docker/Dockerfile.dev` 里有记录）。** 不带 VA 的 -VRAM 分配（VMM 的 `hipMemCreate` 就是这种）在 `kfd_ioctl_alloc_memory_of_gpu` 里 -会把 `[0, size)` 当成 SVM 重叠检查范围。只要有 SVM 注册的区间低于 `size`，内核就 -返回 EADDRINUSE，HIP 报成 out of memory。镜像里的 python3.12 不是 PIE，brk 堆在 -0.9 到 1 GB 附近，所以只有大于约 1 GB 的分配会撞上，这和"小 arena 通过、1.76 GB -的 bench 失败"对得上。上游修复是 93aa919（"drm/amdkfd: fix vram allocation failure -for a special case"），在 6.18 和 stable 6.17.8、6.12.58、6.6.117、6.1.159 里。 -DKMS 模块会覆盖发行版内核自带的 amdgpu，所以要看 amdgpu-dkms 的版本。 -疑点：pit2 上 10.1 的进程根本没有 SVM 注册，所以如果是 H1，51 上一定有别的东西在 -注册 SVM。 - -**H2：GPU 2 当时真的没有显存了。** 可能是同一个 job 里前面的步骤在 51 上留下了显存 -（比如 RDMA/dma-buf 引用在进程退出后没放），也可能是机器上别的进程。job 结束时 -容器被删，下一个 job 就恢复正常，这也能解释 main 紧接着能过。 - -两者用 errno 就能分开：H1 是 `EADDRINUSE`（98），H2 是 `ENOMEM`（12）。 - -## 5. 在 51 上要做的 - -### 5.0 注意 - -这是共享的 CI runner。动 GPU 前先确认没有 CI job 在跑： -`ps aux | grep '[R]unner.Worker'`。不要覆盖 `rocm/mori:ci-cco` 这个 tag，也不要用 -`mori_cco_ci_*` 的容器名，下面用 `localhost/mori-oom:*` 和 `mori_oom_*`。不要在 CI 的 -workspace（`/shared/amdgpu/home/jiahao_zhou_qle/actions-runner/_work/mori/mori`）里 -操作，另外 clone 一份。 - -### 5.1 不用 GPU 的检查（先做，可能直接定案） - -```bash -uname -r -cat /sys/module/amdgpu/version -dpkg -l | grep -E 'amdgpu-dkms ' || rpm -qa | grep amdgpu-dkms -cat /sys/module/amdgpu/parameters/noretry -# 93aa919 在不在。在的话这一行能找到： -grep -n "Skip a special case that allocates VRAM without VA" /usr/src/amdgpu-*/amd/amdkfd/kfd_chardev.c -# 两次失败时内核有没有报 SVM 冲突（H1 的直接证据）。时间带了 UTC 后缀，按 UTC 解释； -# 需要 persistent journal，如果机器之后重启过又没有，就只能看下面的 dmesg： -journalctl -k --since "2026-10-30 07:40 UTC" --until "2026-10-30 07:50 UTC" | grep -E "already allocated by SVM|amdgpu|kfd" -journalctl -k --since "2026-10-30 09:40 UTC" --until "2026-10-30 09:50 UTC" | grep -E "already allocated by SVM|amdgpu|kfd" -# 没有 journal 的话： -sudo dmesg -T | grep -E "already allocated by SVM" | tail -podman info --format '{{.Host.Security.Rootless}}'; ulimit -l -``` - -如果在那两个时间点看到 `Address: 0x0 already allocated by SVM`，就是 H1。 - -### 5.2 用 10.1 镜像原样跑 bench，并挂 spy - -```bash -git clone https://github.com/ROCm/mori.git mori-oom && cd mori-oom -git fetch origin pull/871/head:pr871 && git checkout pr871 # 4a4df301 -podman build --network=host \ - --build-arg BASE_IMAGE=rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 \ - --build-arg GPU_ARCHS="gfx950 gfx1250" --build-arg WITH_FLYDSL=1 \ - -t localhost/mori-oom:rocm101 -f docker/Dockerfile.dev . -mkdir -p /tmp/oom # 把附录 A 存成 /tmp/oom/kfd_svm_spy.c，附录 B 存成 /tmp/oom/replay_ops_v2.sh -CONTAINER_RUNTIME=podman ./docker/ci_run.sh --name mori_oom_101 \ - -e MORI_RDMA_DEVICES=rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7 \ - -e MORI_RDMA_SL=3 -e MORI_RDMA_TC=104 \ - -v $PWD:$PWD -v /tmp/oom:/w localhost/mori-oom:rocm101 sleep infinity -podman exec mori_oom_101 git config --global --add safe.directory $PWD -podman exec mori_oom_101 bash -c "cd $PWD && BUILD_TESTS=ON BUILD_EXAMPLES=ON BUILD_BENCHMARK=ON \ - MORI_WITH_MPI=ON BUILD_UMBP=OFF BUILD_CCO_SDMA=ON pip install ." -podman exec mori_oom_101 gcc -shared -fPIC -O2 -o /w/kfd_svm_spy.so /w/kfd_svm_spy.c -ldl - -amd-smi metric -m # 先记一下每张卡当前用了多少 -podman exec -e LD_PRELOAD=/w/kfd_svm_spy.so -e SPY_BT=5 mori_oom_101 bash -c \ - "cd $PWD/tests/python/ops/dispatch_combine_v2 && BACKENDS=flydsl,hip COMB_MODE=pull MODES=graph \ - SWEEP=128,512,4096,8192 ITERS=20 WARMUP=10 timeout 900 torchrun --standalone --nproc_per_node=8 bench_ep.py" \ - > /tmp/oom/bench_spy.log 2>&1; echo rc=$? -grep -E '^\[spy' /tmp/oom/bench_spy.log | sed -E 's/\[spy [0-9]+\]/[spy]/' | sort | uniq -c -``` - -pit2 上正常的输出只有这两种行：16 条 -`VRAM alloc va=0 size=1.76 GB flags=0xb5000001 ok, no SVM range below it`，和 8 条 -`KFD ioctl nr=0x8 FAILED errno=28`（可以忽略）。 - -### 5.3 spy 输出怎么看 - -| 失败 rank 上的 spy 输出 | 含义 | 下一步 | -|---|---|---| -| `VRAM alloc ... FAILED errno=98 (Address already in use)`，并且列出了一个低于 size 的 SVM 区间 | H1 | 升级 51 的 amdgpu-dkms 到带 93aa919 的版本。同时看 `SPY_BT` 打出的 backtrace，找出是谁注册的 SVM | -| `SVM op=0 start=...` 行，并且有 `WOULD FAIL on an amdgpu without 93aa919` | 有 SVM 注册（pit2 上没有），分配本身成功了 | 看 backtrace 和 7.14 的对比（5.5） | -| `VRAM alloc ... FAILED errno=12 (Cannot allocate memory)` | H2，显存真的不够 | `amd-smi process`、`amd-smi metric -m`、`podman ps -a`，找是谁占着 GPU 2 | -| `KFD reports only X GB available` | runtime 查到可用显存不够，在进内核前就拒绝了 | 同上 | -| 失败的那次分配没有任何 spy 输出 | 在 HIP/ROCr 里就被拒了，没到内核 | 加 `AMD_LOG_LEVEL=4` 重跑，看 HIP 那边的原因 | - -### 5.4 如果单独跑 bench 能过 - -CI 里 bench 前面还有很多步骤（CCO 单测、SDMA、benchmark、examples、allocator、 -GEMM+AR、ops-v2）。按 `ci_cco.yml` 里 cco-unit-test 的顺序在同一个容器里逐步跑， -同时在 host 上开附录 C 的显存采样，看哪一步之后 GPU 2 的显存不回到空闲水平。 -附录 B 只覆盖最后 8 个 ops-v2 步骤，可以先跑它；不行再把前面的步骤按 workflow -补上。 - -```bash -podman exec mori_oom_101 bash /w/replay_ops_v2.sh $PWD /w/replay_logs > /tmp/oom/replay_marks.log 2>&1 -``` - -### 5.5 对照：7.14 镜像 - -用 main 的 base 再建一个镜像，同样挂 spy 跑 bench，比较两边的 SVM 注册： - -```bash -podman build --network=host \ - --build-arg BASE_IMAGE=rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0 \ - --build-arg GPU_ARCHS="gfx950 gfx1250" --build-arg WITH_FLYDSL=1 \ - -t localhost/mori-oom:rocm714 -f docker/Dockerfile.dev . -``` - -然后和 5.2 一样起容器 `mori_oom_714`、装 MORI、跑 bench。7.14 的镜像里 -`/opt/rocm/lib` 不在默认搜索路径上，单独跑的二进制需要 `LD_LIBRARY_PATH=/opt/rocm/lib`。 - -### 5.6 A/B 开关（在 10.1 上能复现之后再用） - -- `CCO_UNCACHED_WINDOW=0`：arena 用 pinned 而不是 uncached（flags 变成 `0xb0000001`）。 -- `CCO_GDR_CAPABLE=0`：分配不带 gpuDirectRDMACapable。 -- `SWEEP=128,512,4096`：最大 token 数从 8192 降到 4096，arena 大约减半。如果这样就能 - 过，说明和分配大小有关，支持 H1。 - -## 附录 A：kfd_svm_spy.c - -```c -// LD_PRELOAD shim: logs KFD SVM registrations, VA-less VRAM allocations and failed -// KFD ioctls (with errno). Flags every VA-less VRAM allocation that an amdgpu -// without 93aa919 would reject with EADDRINUSE: its SVM-overlap check runs over -// [0, size) when va_addr == 0. -// gcc -shared -fPIC -O2 -o kfd_svm_spy.so kfd_svm_spy.c -ldl -// LD_PRELOAD=./kfd_svm_spy.so SPY_BT=3 python ... -#define _GNU_SOURCE -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -struct svm_args { - uint64_t start_addr, size; - uint32_t op, nattr; -}; -struct alloc_args { - uint64_t va_addr, size, handle, mmap_offset; - uint32_t gpu_id, flags; -}; -struct avail_args { - uint64_t available; - uint32_t gpu_id, pad; -}; -#define NR_ALLOC 0x16 -#define NR_SVM 0x20 -#define NR_AVAIL 0x23 -#define VRAM_FLAG 1u - -#define MAX_RANGES 4096 -static uint64_t g_start[MAX_RANGES], g_len[MAX_RANGES]; -static int g_n, g_bt_left = -1, g_other_fail_left = 20; -static pthread_mutex_t g_mu = PTHREAD_MUTEX_INITIALIZER; - -static void Backtrace(void) { - if (g_bt_left < 0) g_bt_left = getenv("SPY_BT") ? atoi(getenv("SPY_BT")) : 0; - if (g_bt_left <= 0) return; - g_bt_left--; - void* frames[24]; - int n = backtrace(frames, 24); - backtrace_symbols_fd(frames + 2, n - 2, 2); -} - -static int LowestSvmBelow(uint64_t size) { - int hit = -1; - for (int i = 0; i < g_n; i++) - if (g_start[i] < size && (hit < 0 || g_start[i] < g_start[hit])) hit = i; - return hit; -} - -int ioctl(int fd, unsigned long req, ...) { - static int (*real)(int, unsigned long, void*); - if (!real) real = (int (*)(int, unsigned long, void*))dlsym(RTLD_NEXT, "ioctl"); - va_list ap; - va_start(ap, req); - void* arg = va_arg(ap, void*); - va_end(ap); - int ret = real(fd, req, arg); - int err = errno; - if (_IOC_TYPE(req) != 'K') return ret; - pthread_mutex_lock(&g_mu); - if (_IOC_NR(req) == NR_SVM) { - struct svm_args* a = (struct svm_args*)arg; - fprintf(stderr, "[spy %d] SVM op=%u start=0x%lx size=0x%lx nattr=%u ret=%d\n", getpid(), a->op, - a->start_addr, a->size, a->nattr, ret); - if (a->op == 0 && ret == 0 && g_n < MAX_RANGES) { - g_start[g_n] = a->start_addr; - g_len[g_n++] = a->size; - Backtrace(); - } - } else if (_IOC_NR(req) == NR_ALLOC) { - struct alloc_args* a = (struct alloc_args*)arg; - if (a->va_addr == 0 && (a->flags & VRAM_FLAG)) { - int hit = LowestSvmBelow(a->size); - if (ret != 0) - fprintf(stderr, - "[spy %d] VRAM alloc va=0 size=%.2f GB flags=0x%08x FAILED errno=%d (%s); lowest " - "SVM range below the size: %s0x%lx\n", - getpid(), a->size / 1073741824.0, a->flags, err, strerror(err), - hit >= 0 ? "" : "none, ", hit >= 0 ? g_start[hit] : 0); - else if (hit >= 0) - fprintf(stderr, - "[spy %d] VRAM alloc va=0 size=%.2f GB flags=0x%08x ok here, WOULD FAIL on an " - "amdgpu without 93aa919: SVM range [0x%lx, +0x%lx) is below the size\n", - getpid(), a->size / 1073741824.0, a->flags, g_start[hit], g_len[hit]); - else if (a->size >= (64ul << 20)) - fprintf(stderr, "[spy %d] VRAM alloc va=0 size=%.2f GB flags=0x%08x ok, no SVM range below it\n", - getpid(), a->size / 1073741824.0, a->flags); - } else if (ret != 0) { - fprintf(stderr, "[spy %d] alloc va=0x%lx size=0x%lx flags=0x%08x FAILED errno=%d (%s)\n", - getpid(), a->va_addr, a->size, a->flags, err, strerror(err)); - } - } else if (_IOC_NR(req) == NR_AVAIL) { - struct avail_args* a = (struct avail_args*)arg; - if (ret == 0 && a->available < (16ul << 30)) - fprintf(stderr, "[spy %d] KFD reports only %.2f GB available on gpu_id %u\n", getpid(), - a->available / 1073741824.0, a->gpu_id); - } else if (ret != 0 && g_other_fail_left > 0) { - g_other_fail_left--; - fprintf(stderr, "[spy %d] KFD ioctl nr=0x%x FAILED errno=%d (%s)\n", getpid(), - (unsigned)_IOC_NR(req), err, strerror(err)); - } - pthread_mutex_unlock(&g_mu); - errno = err; - return ret; -} - -int munmap(void* addr, size_t len) { - static int (*real)(void*, size_t); - if (!real) real = (int (*)(void*, size_t))dlsym(RTLD_NEXT, "munmap"); - int ret = real(addr, len); - int err = errno; - if (ret == 0) { - uint64_t lo = (uint64_t)addr, hi = lo + len; - pthread_mutex_lock(&g_mu); - for (int i = 0; i < g_n; i++) - if (g_start[i] >= lo && g_start[i] + g_len[i] <= hi) { - g_start[i] = g_start[--g_n]; - g_len[i] = g_len[g_n]; - i--; - } - pthread_mutex_unlock(&g_mu); - } - errno = err; - return ret; -} -``` - -## 附录 B：replay_ops_v2.sh（在容器里跑） - -```bash -#!/bin/bash -# Replays ci_cco.yml's ops-v2 steps (dispatch/combine test .. bench) back to back, -# inside the CI container. Usage: replay_ops_v2.sh -W=$1 -LOGS=$2 -D=$W/tests/python/ops/dispatch_combine_v2 -mark() { echo "@@ $(date +%T.%N | cut -c1-10) $*"; } -leftover() { echo "@@ leftover python procs: $(pgrep -c -f '[p]ython|[t]orchrun' || true)"; } -step() { - local name=$1; shift - mark "BEGIN $name" - (cd "$D" && env "$@") > "$LOGS/$name.log" 2>&1 - local rc=$? - mark "END $name rc=$rc" - leftover -} -mkdir -p "$LOGS" -step 24-pytest-intranode bash -c "cd $W && python -m pytest tests/python/ops/dispatch_combine_v2/test_dispatch_combine_v2_intranode.py -q" -step 25-hip-test-op MORI_V2_KERNEL_BACKEND=hip timeout 900 torchrun --standalone --nproc_per_node=8 test_op.py -step 26-hip-scales MORI_V2_KERNEL_BACKEND=hip SCALE_DIM=32 timeout 900 torchrun --standalone --nproc_per_node=8 test_op.py -for disp in fp8 fp4; do - step 27-asym-$disp MORI_V2_KERNEL_BACKEND=hip SCALE_DIM=224 HIDDEN=7168 DISP=$disp timeout 900 torchrun --standalone --nproc_per_node=8 test_asym_dtype.py -done -step 28-parity timeout 900 torchrun --standalone --nproc_per_node=8 test_ep_backend_parity.py -for transport in sdma lsa; do - step 29-asyncll-$transport MORI_ENABLE_SDMA=1 TRANSPORT=$transport timeout 900 torchrun --standalone --nproc_per_node=8 test_asyncll.py -done -step 30-routing-snapshot timeout 900 torchrun --standalone --nproc_per_node=8 test_routing_snapshot.py -step 31-bench BACKENDS=flydsl,hip COMB_MODE=pull MODES=graph SWEEP=128,512,4096,8192 ITERS=20 WARMUP=10 timeout 900 torchrun --standalone --nproc_per_node=8 bench_ep.py -mark DONE -``` - -## 附录 C：host 上的显存采样（每 0.5 秒一行，单位 MB，gpu0 到 gpu7） - -```bash -( while true; do echo "$(date +%T.%N | cut -c1-10) $(amd-smi metric -m --json 2>/dev/null | python3 -c " -import json,sys -d=json.load(sys.stdin); d=d.get('gpu_data',d) if isinstance(d,dict) else d -print(' '.join(str(g['mem_usage']['used_vram']['value']) for g in d))")" >> /tmp/oom/vram.log; sleep 0.5; done ) & -SAMPLER=$! -# ... 跑测试 ... -kill $SAMPLER -``` - -把 `replay_marks.log` 里每一步的 BEGIN/END 时间和 `vram.log` 对上，就能看出每一步 -开始时的显存和峰值。 - -## 附录 D：pit2 上用过的其他测试 - -在 pit2-p03-g17 的 `/var/zqz/zqz_vmm_churn/` 下，需要的话可以拷过来： -`vmm_churn.cpp`（裸 HIP VMM 反复 alloc/free，有同进程和跨进程两种模式）和 -`cco_churn.py`（MORI 的 `alloc_mem`、`register_window`、free 循环，用 mpirun 跑）。 -在 pit2 上这两个都显示释放及时。H2 的情况下可以在 51 上跑来对照。 From 98efed5b13647800342256657b81cdbcd41956f1 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Wed, 4 Nov 2026 21:48:40 -0600 Subject: [PATCH 10/13] test(cco): drop the rocm_sdk version-match test Co-authored-by: Cursor --- tests/python/cco/test_rocm_bootstrap.py | 11 ----------- 1 file changed, 11 deletions(-) diff --git a/tests/python/cco/test_rocm_bootstrap.py b/tests/python/cco/test_rocm_bootstrap.py index 65f3a950b3..3fb996734b 100644 --- a/tests/python/cco/test_rocm_bootstrap.py +++ b/tests/python/cco/test_rocm_bootstrap.py @@ -129,14 +129,3 @@ def test_baked_version_uses_major_minor_compatibility_pattern(): major, minor, *_ = bootstrap.ROCM_BUILD_VERSION.split(".") assert bootstrap.ROCM_VERSION_PATTERN == f"{major}.{minor}.*" - - -def test_baked_version_matches_the_installed_rocm_sdk(): - # A mismatch only warns at import time, from inside rocm_sdk, so nothing - # else in the suite would notice it. - if bootstrap.ROCM_VERSION_PATTERN is None: - pytest.skip("ROCm build metadata is generated while building the package") - rocm_sdk = pytest.importorskip("rocm_sdk") - - major, minor, *_ = rocm_sdk.__version__.split(".") - assert bootstrap.ROCM_VERSION_PATTERN == f"{major}.{minor}.*" From 558a3cb9ed75c5bc46fb56beb109408b2e17bb2d Mon Sep 17 00:00:00 2001 From: qizzhang Date: Wed, 4 Nov 2026 21:48:41 -0600 Subject: [PATCH 11/13] ci: keep the py3.10 nightly wheel lane on ROCm 7.2.4 With both nightly lanes on 10.1, the next release would ship only 10.1-built wheels and no job would build or test 7.2.x. Keep the py3.10 lane on 7.2.4 as on main: its RHEL 8 build args for Dockerfile.manylinux and its rocm7.2.4 py3.10 test images in test-wheel, test-wheel-internode and test-pypi-wheel. py3.12 moves to 10.1. The wheel artifacts are keyed by Python version, so the two lanes do not collide. Co-authored-by: Cursor --- .github/workflows/nightly.yml | 18 +++++++++++------- 1 file changed, 11 insertions(+), 7 deletions(-) diff --git a/.github/workflows/nightly.yml b/.github/workflows/nightly.yml index 9f26dddc43..4b796c1075 100644 --- a/.github/workflows/nightly.yml +++ b/.github/workflows/nightly.yml @@ -31,11 +31,15 @@ jobs: # Built in docker/Dockerfile.manylinux, so a wheel needs nothing newer # than glibc 2.28 instead of the glibc/libstdc++ of the Ubuntu image that # ships its Python. Each lane builds against the ROCm its test image - # below runs: 10.1 for both, the Dockerfile's default (its header lists - # the args for other ROCm releases). + # below runs: 7.2.4 for py3.10, so a 7.x wheel is still built and tested + # every night, and 10.1, the Dockerfile's default, for py3.12. - python: "3.10" - rocm: "10.1" - rocm_build_args: "" + rocm: "7.2.4" + rocm_build_args: >- + --build-arg ROCM_VERSION=7.2.4 + --build-arg ROCM_REPO=https://repo.radeon.com/rocm/rhel8/7.2.4/main + --build-arg ROCM_GPG=https://repo.radeon.com/rocm/rocm.gpg.key + --build-arg "ROCM_PACKAGES=rocm-hip-runtime-devel rocm-smi-lib" - python: "3.12" rocm: "10.1" rocm_build_args: "" @@ -186,7 +190,7 @@ jobs: - platform: MI355X_AINIC python: "3.10" runner: [self-hosted, MI355X-AINIC-PIT-G17] - base_image: rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.12.0 + base_image: rocm/pytorch:rocm7.2.4_ubuntu22.04_py3.10_pytorch_release_2.8.0 rdma_devices: rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7 rdma_sl: 3 rdma_tc: 104 @@ -433,7 +437,7 @@ jobs: - platform: MI355X_AINIC python: "3.10" runner: [self-hosted, MI355X-AINIC] - base_image: rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.12.0 + base_image: rocm/pytorch:rocm7.2.4_ubuntu22.04_py3.10_pytorch_release_2.8.0 node1_host: 10.2.80.20 node2_host: 10.2.80.10 node2_port: 22 @@ -932,7 +936,7 @@ jobs: matrix: include: - python: "3.10" - base_image: rocm/pytorch:rocm10.1.0_ubuntu22.04_py3.10_pytorch_release_2.12.0 + base_image: rocm/pytorch:rocm7.2.4_ubuntu22.04_py3.10_pytorch_release_2.8.0 - python: "3.12" base_image: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 env: From 5d6ab671b42ebd79aa82274cd4cfb32d5a17efba Mon Sep 17 00:00:00 2001 From: qizzhang Date: Wed, 4 Nov 2026 21:49:08 -0600 Subject: [PATCH 12/13] ci: set HSA_USE_SVM=0 on the MI355X-AINIC jobs, not in the image Only hosts whose amdgpu-dkms lacks 93aa919 need it, and it is not free: with SVM off, libhsakmt pins a pageable buffer again on every copy, which everyone who builds Dockerfile.dev would inherit. Of the runners, only the MI355X-AINIC pair (amdgpu-dkms 6.14.14) lacks the fix. Set it per job instead, on node1's and node2's containers: cco-unit-test and cco-internode-test in ci_cco.yml, internode-test's MI355X_AINIC entry in ci.yml, and test-wheel-internode's MI355X_AINIC entries and flydsl-compat in nightly.yml. Dockerfile.dev drops the workaround block, MALLOC_MMAP_THRESHOLD_ included: it was there for the same bug and does not prevent it. Co-authored-by: Cursor --- .github/workflows/ci.yml | 11 +++++++++++ .github/workflows/ci_cco.yml | 14 ++++++++++++++ .github/workflows/nightly.yml | 14 ++++++++++++++ docker/Dockerfile.dev | 18 ------------------ 4 files changed, 39 insertions(+), 18 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 8f4851a26c..f756110c47 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -769,6 +769,14 @@ jobs: rdma_tc: 104 ct: podman timeout_minutes: 45 + # These hosts' amdgpu-dkms lacks 93aa919 ("drm/amdkfd: fix vram + # allocation failure for a special case"): hipMemCreate (VRAM with no VA + # yet) fails with EADDRINUSE, reported as out of memory, while any SVM + # range sits below its size, and libhsakmt registers every pageable + # buffer HIP pins for a copy as one. HSA_USE_SVM=0 registers them as + # userptr instead. Not in the image: with SVM off every pageable copy is + # pinned again. Drop it once the hosts' amdgpu-dkms carries the fix. + hsa_use_svm: 0 # 8x ConnectX-7 NDR on InfiniBand, so no rdma_sl/rdma_tc: mlx5 applies # both only on the RoCE path. mlx5_8 is deliberately not in # rdma_devices -- it is the 100GbE port behind ens14np0 (the TCP @@ -815,6 +823,7 @@ jobs: MORI_RDMA_SL: ${{ matrix.rdma_sl }} MORI_RDMA_TC: ${{ matrix.rdma_tc }} MORI_SOCKET_IFNAME: ${{ matrix.node1_ifname }} + HSA_USE_SVM: ${{ matrix.hsa_use_svm }} timeout-minutes: ${{ matrix.timeout_minutes }} steps: - name: Fix workspace permissions @@ -859,6 +868,7 @@ jobs: -e MORI_RDMA_DEVICES=$MORI_RDMA_DEVICES \ ${MORI_RDMA_SL:+-e MORI_RDMA_SL=$MORI_RDMA_SL} \ ${MORI_RDMA_TC:+-e MORI_RDMA_TC=$MORI_RDMA_TC} \ + ${HSA_USE_SVM:+-e HSA_USE_SVM=$HSA_USE_SVM} \ -e MORI_SOCKET_IFNAME=$MORI_SOCKET_IFNAME \ -e GLOO_SOCKET_IFNAME=$NODE1_IFNAME \ -e NCCL_SOCKET_IFNAME=$NODE1_IFNAME \ @@ -901,6 +911,7 @@ jobs: -e MORI_RDMA_DEVICES=$MORI_RDMA_DEVICES \ ${MORI_RDMA_SL:+-e MORI_RDMA_SL=$MORI_RDMA_SL} \ ${MORI_RDMA_TC:+-e MORI_RDMA_TC=$MORI_RDMA_TC} \ + ${HSA_USE_SVM:+-e HSA_USE_SVM=$HSA_USE_SVM} \ -e MORI_SOCKET_IFNAME=$MORI_SOCKET_IFNAME \ -e GLOO_SOCKET_IFNAME=$NODE2_IFNAME \ -e NCCL_SOCKET_IFNAME=$NODE2_IFNAME \ diff --git a/.github/workflows/ci_cco.yml b/.github/workflows/ci_cco.yml index abfdf86e8e..fc709466a8 100644 --- a/.github/workflows/ci_cco.yml +++ b/.github/workflows/ci_cco.yml @@ -29,10 +29,19 @@ jobs: rdma_devices: rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7 rdma_sl: 3 rdma_tc: 104 + # These hosts' amdgpu-dkms lacks 93aa919 ("drm/amdkfd: fix vram + # allocation failure for a special case"): hipMemCreate (VRAM with no VA + # yet) fails with EADDRINUSE, reported as out of memory, while any SVM + # range sits below its size, and libhsakmt registers every pageable + # buffer HIP pins for a copy as one. HSA_USE_SVM=0 registers them as + # userptr instead. Not in the image: with SVM off every pageable copy is + # pinned again. Drop it once the hosts' amdgpu-dkms carries the fix. + hsa_use_svm: 0 env: MORI_RDMA_DEVICES: ${{ matrix.rdma_devices }} MORI_RDMA_SL: ${{ matrix.rdma_sl }} MORI_RDMA_TC: ${{ matrix.rdma_tc }} + HSA_USE_SVM: ${{ matrix.hsa_use_svm }} timeout-minutes: 60 steps: - name: Checkout @@ -56,6 +65,7 @@ jobs: -e MORI_RDMA_DEVICES=$MORI_RDMA_DEVICES \ -e MORI_RDMA_SL=$MORI_RDMA_SL \ -e MORI_RDMA_TC=$MORI_RDMA_TC \ + -e HSA_USE_SVM=$HSA_USE_SVM \ -v $GITHUB_WORKSPACE:$GITHUB_WORKSPACE \ -w $GITHUB_WORKSPACE \ $IMAGE sleep infinity @@ -433,6 +443,7 @@ jobs: rdma_devices: rocep105s0,rocep121s0,rocep137s0,rocep153s0,rocep233s0,rocep249s0,rocep25s0,rocep9s0 rdma_sl: 3 rdma_tc: 104 + hsa_use_svm: 0 # the amdgpu-dkms workaround, see cco-unit-test env: # Overrides the workflow-level name, which is keyed by run_id alone: sharing # it would have this job's cleanup delete the other job's container. @@ -447,6 +458,7 @@ jobs: MORI_RDMA_SL: ${{ matrix.rdma_sl }} MORI_RDMA_TC: ${{ matrix.rdma_tc }} MORI_SOCKET_IFNAME: ${{ matrix.node1_ifname }} + HSA_USE_SVM: ${{ matrix.hsa_use_svm }} ENTRY: tests/python/ops/dispatch_combine_v2/test_dispatch_combine_v2_internode.py timeout-minutes: 60 steps: @@ -474,6 +486,7 @@ jobs: -e MORI_SOCKET_IFNAME=$MORI_SOCKET_IFNAME \ -e GLOO_SOCKET_IFNAME=$NODE1_IFNAME \ -e NCCL_SOCKET_IFNAME=$NODE1_IFNAME \ + -e HSA_USE_SVM=$HSA_USE_SVM \ -v $GITHUB_WORKSPACE:$GITHUB_WORKSPACE \ -w $GITHUB_WORKSPACE \ $IMAGE sleep infinity @@ -502,6 +515,7 @@ jobs: -e MORI_SOCKET_IFNAME=$MORI_SOCKET_IFNAME \ -e GLOO_SOCKET_IFNAME=$NODE2_IFNAME \ -e NCCL_SOCKET_IFNAME=$NODE2_IFNAME \ + -e HSA_USE_SVM=$HSA_USE_SVM \ -v $GITHUB_WORKSPACE:$GITHUB_WORKSPACE \ -w $GITHUB_WORKSPACE \ $IMAGE sleep infinity && diff --git a/.github/workflows/nightly.yml b/.github/workflows/nightly.yml index 4b796c1075..7c547eec3d 100644 --- a/.github/workflows/nightly.yml +++ b/.github/workflows/nightly.yml @@ -448,6 +448,14 @@ jobs: rdma_tc: 104 ct: podman timeout_minutes: 45 + # These hosts' amdgpu-dkms lacks 93aa919 ("drm/amdkfd: fix vram + # allocation failure for a special case"): hipMemCreate (VRAM with no VA + # yet) fails with EADDRINUSE, reported as out of memory, while any SVM + # range sits below its size, and libhsakmt registers every pageable + # buffer HIP pins for a copy as one. HSA_USE_SVM=0 registers them as + # userptr instead. Not in the image: with SVM off every pageable copy is + # pinned again. Drop it once the hosts' amdgpu-dkms carries the fix. + hsa_use_svm: 0 - platform: MI355X_AINIC python: "3.12" runner: [self-hosted, MI355X-AINIC] @@ -462,6 +470,7 @@ jobs: rdma_tc: 104 ct: podman timeout_minutes: 45 + hsa_use_svm: 0 # as the py3.10 entry # 8x ConnectX-7 NDR on InfiniBand, so no rdma_sl/rdma_tc: mlx5 applies # both only on the RoCE path. mlx5_8 is deliberately not in # rdma_devices -- it is the 100GbE port behind ens14np0 (the TCP @@ -508,6 +517,7 @@ jobs: MORI_RDMA_SL: ${{ matrix.rdma_sl }} MORI_RDMA_TC: ${{ matrix.rdma_tc }} MORI_SOCKET_IFNAME: ${{ matrix.node1_ifname }} + HSA_USE_SVM: ${{ matrix.hsa_use_svm }} # Perf dashboard capture on node1/rank0 (MI355X_AINIC only for now). MORI_PERF_OUT: ${{ matrix.platform == 'MI355X_AINIC' && format('{0}/_perf/perf.jsonl', github.workspace) || '' }} MORI_PERF_RUN_ID: ${{ github.run_id }} @@ -551,6 +561,7 @@ jobs: -e MORI_SOCKET_IFNAME=$MORI_SOCKET_IFNAME \ -e GLOO_SOCKET_IFNAME=$NODE1_IFNAME \ -e NCCL_SOCKET_IFNAME=$NODE1_IFNAME \ + ${HSA_USE_SVM:+-e HSA_USE_SVM=$HSA_USE_SVM} \ -e MORI_PERF_OUT=$MORI_PERF_OUT \ -e MORI_PERF_RUN_ID=$MORI_PERF_RUN_ID \ -e MORI_PERF_COMMIT=$MORI_PERF_COMMIT \ @@ -594,6 +605,7 @@ jobs: -e MORI_SOCKET_IFNAME=$MORI_SOCKET_IFNAME \ -e GLOO_SOCKET_IFNAME=$NODE2_IFNAME \ -e NCCL_SOCKET_IFNAME=$NODE2_IFNAME \ + ${HSA_USE_SVM:+-e HSA_USE_SVM=$HSA_USE_SVM} \ -e MORI_IO_BENCH_ENGINE_BIN=$GITHUB_WORKSPACE/build_bench/tests/cpp/bench_engine \ -v $GITHUB_WORKSPACE:$GITHUB_WORKSPACE \ -w $GITHUB_WORKSPACE \ @@ -1017,6 +1029,7 @@ jobs: MORI_RDMA_DEVICES: rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7 # as ci_cco.yml MORI_RDMA_SL: 3 MORI_RDMA_TC: 104 + HSA_USE_SVM: 0 # the amdgpu-dkms workaround, see test-wheel-internode IMAGE: rocm/mori:ci-flydsl-${{ matrix.flydsl }} BASE_IMAGE: rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.0 CONTAINER: mori_flydsl_compat_${{ github.run_id }}_${{ strategy.job-index }} @@ -1042,6 +1055,7 @@ jobs: -e MORI_RDMA_DEVICES=$MORI_RDMA_DEVICES \ -e MORI_RDMA_SL=$MORI_RDMA_SL \ -e MORI_RDMA_TC=$MORI_RDMA_TC \ + -e HSA_USE_SVM=$HSA_USE_SVM \ -v $GITHUB_WORKSPACE:$GITHUB_WORKSPACE \ -w $GITHUB_WORKSPACE \ $IMAGE sleep infinity diff --git a/docker/Dockerfile.dev b/docker/Dockerfile.dev index dd1cc5a0a5..384dabb1f4 100644 --- a/docker/Dockerfile.dev +++ b/docker/Dockerfile.dev @@ -64,24 +64,6 @@ RUN set -ex; \ rm -rf /var/lib/apt/lists/*; \ fi -# Works around an amdgpu kernel bug that makes the symmetric heap's hipMemCreate -# fail with hipErrorOutOfMemory while VRAM is nearly empty. VMM allocations carry -# no VA yet, so the thunk passes va_addr=0, but kfd_ioctl_alloc_memory_of_gpu -# still runs its SVM-overlap check over [va_addr, va_addr+size) — with va_addr=0 -# that degenerates into "is anything SVM-registered below ?". When the KFD -# offers the SVM API, the thunk registers each pageable buffer HIP pins for a -# copy as an SVM range, and the range outlives the pin. python here is not PIE, -# so a tensor copied from the brk heap leaves a range at a low VA and the kernel -# returns EADDRINUSE. glibc's mmap threshold cannot keep tensors off the heap: it -# still serves a larger request from a free heap chunk. HSA_USE_SVM=0 makes the -# thunk register host memory as userptr instead, so no SVM range exists. -# -# Upstream fix is 93aa919 ("drm/amdkfd: fix vram allocation failure for a special -# case"), in 6.18 and stable 6.17.8 / 6.12.58 / 6.6.117 / 6.1.159. Drop this once -# the runners' amdgpu-dkms carries it — note the DKMS module overrides the -# distro kernel's amdgpu, so bumping the Ubuntu kernel alone does nothing. -ENV HSA_USE_SVM=0 - # Declared here, next to their only use: every RUN after an ARG has it in its # cache key, so up top a new FlyDSL version would rebuild the ROCm layers too. # 1 installs flydsl, which only the cco Python examples need. From 7062f9c5e7567daadd5abb71b53ae231d4731a19 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Wed, 4 Nov 2026 21:49:08 -0600 Subject: [PATCH 13/13] fix(topology): reuse a loaded libamd_smi only of mori's major version A libamd_smi of another major version than the amdsmi.h mori was built against may lay out its structs differently. Check amdsmi_get_lib_version() against AMDSMI_LIB_VERSION_MAJOR before reusing the copy already in the process; otherwise say so on stderr and open one by path as before. Also note that reuse relies on the first user having initialized amd_smi with AMDSMI_INIT_AMD_GPUS: a later amdsmi_init only takes a reference and ignores its flags. test_topology in the 10.1 image finds all 8 GPUs with rocm_sdk's libamd_smi.so.27 preloaded (reused, one copy mapped) and with a stub reporting major 99 preloaded (not reused, falls back). Co-authored-by: Cursor --- src/application/topology/gpu.cpp | 22 ++++++++++++++++++++-- 1 file changed, 20 insertions(+), 2 deletions(-) diff --git a/src/application/topology/gpu.cpp b/src/application/topology/gpu.cpp index 88671eb618..8c1e19be94 100644 --- a/src/application/topology/gpu.cpp +++ b/src/application/topology/gpu.cpp @@ -49,7 +49,11 @@ namespace { // from the rocm_sdk wheels, and a second copy loaded beside it binds part of its // state to the first one's symbols and then reports 0 sockets. amd_smi counts // init/shut_down references, so sharing the instance leaves the other user's -// session alone. +// session alone. A copy of another major version than the header mori was built +// against may lay out its structs differently, so it is not reused. And a later +// amdsmi_init only takes a reference and ignores its flags: the shared instance +// sees GPUs only if its first user initialized it with AMDSMI_INIT_AMD_GPUS, as +// torch does. void* LoadedAmdSmi() { std::string path; dl_iterate_phdr( @@ -62,7 +66,21 @@ void* LoadedAmdSmi() { return 0; }, &path); - return path.empty() ? nullptr : dlopen(path.c_str(), RTLD_NOW | RTLD_NOLOAD); + if (path.empty()) return nullptr; + void* handle = dlopen(path.c_str(), RTLD_NOW | RTLD_NOLOAD); + if (!handle) return nullptr; + + auto getLibVersion = reinterpret_cast( + dlsym(handle, "amdsmi_get_lib_version")); + amdsmi_version_t version = {}; + if (getLibVersion && getLibVersion(&version) == AMDSMI_STATUS_SUCCESS && + version.major == AMDSMI_LIB_VERSION_MAJOR) { + return handle; + } + fprintf(stderr, "[AMD-SMI] not reusing %s: its major version is %u, mori was built against %d\n", + path.c_str(), version.major, AMDSMI_LIB_VERSION_MAJOR); + dlclose(handle); + return nullptr; } void* OpenAmdSmi() {