diff --git a/.claude/skills/deploy-mori/SKILL.md b/.claude/skills/deploy-mori/SKILL.md index 346c286c15..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:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.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 @@ -427,35 +427,44 @@ 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.1 (an Ubuntu 22.04 +image takes `ubuntu2204` in the repo URL): + ```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.1-gfx942 amdrocm-core-dev10.1-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..f756110c47 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.1.0_ubuntu24.04_py3.12_pytorch_release_2.12.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,19 +460,17 @@ 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) + # 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 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 }} + 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 @@ -491,14 +495,14 @@ 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) || '' }} - # ROCm 10 ships in a different apt repo than 7.x, hence the ROCM_APT_* args. + # 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=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" \ + --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 @@ -522,10 +526,12 @@ jobs: 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 && pip install ." + "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" @@ -556,13 +562,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 @@ -570,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 @@ -588,15 +593,18 @@ 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 --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 \ + --build-arg BASE_IMAGE=$JAX_BASE_IMAGE \ -t rocm/mori:jax-ci \ - -f docker/Dockerfile.jax091-rocm711 . + -f docker/Dockerfile.dev . - name: Start container run: | @@ -616,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 --break-system-packages ." + BUILD_UMBP=OFF BUILD_XLA_FFI_OPS=ON pip install . && pip install pytest" - name: MORI-EP JAX test run: | @@ -760,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 @@ -806,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 @@ -850,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 \ @@ -892,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 418accc4eb..fc709466a8 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.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) @@ -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/codeql.yml b/.github/workflows/codeql.yml index 106c4d5eae..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:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.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 ad48266c8f..7c547eec3d 100644 --- a/.github/workflows/nightly.yml +++ b/.github/workflows/nightly.yml @@ -31,7 +31,8 @@ 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. + # 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: "7.2.4" rocm_build_args: >- @@ -40,7 +41,7 @@ jobs: --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: "7.14" + rocm: "10.1" rocm_build_args: "" steps: # Must precede Checkout: later steps run as root inside the CI container @@ -199,7 +200,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.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 @@ -213,7 +214,7 @@ jobs: - platform: MI300X_MLX python: "3.12" 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.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 @@ -447,10 +448,18 @@ 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] - base_image: rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.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 @@ -461,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 @@ -471,7 +481,7 @@ jobs: - platform: MI300X_MLX python: "3.12" 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.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 @@ -507,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 }} @@ -550,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 \ @@ -593,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 \ @@ -937,7 +950,7 @@ jobs: - python: "3.10" base_image: rocm/pytorch:rocm7.2.4_ubuntu22.04_py3.10_pytorch_release_2.8.0 - python: "3.12" - base_image: rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.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: @@ -1016,8 +1029,9 @@ 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:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.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 @@ -1041,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 d15e757ee8..384dabb1f4 100644 --- a/docker/Dockerfile.dev +++ b/docker/Dockerfile.dev @@ -1,21 +1,25 @@ # 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.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.12.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" +# {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.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/ubuntu{ubuntu} stable main" RUN apt-get update && \ apt-get install -y --no-install-recommends \ @@ -27,17 +31,20 @@ 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 # 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 \ @@ -48,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; \ @@ -56,22 +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 ?". 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. -# -# 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 - # 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. 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..02a604ad6e 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.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 +# 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.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. 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..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) -# ROCm 7.14's hsakmtTargets.cmake 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 diff --git a/src/application/topology/gpu.cpp b/src/application/topology/gpu.cpp index 38450d3a77..8c1e19be94 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,49 @@ 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. 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( + [](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); + 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() { + 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", 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: 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 {