From 75cf64e6c1db0f7e962e4ad185b3769c4f0f213c Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Mon, 17 Aug 2026 16:22:03 +0200 Subject: [PATCH 1/3] CSCS MI300 CI: select the ROCm provider explicitly The MI300 job set JULIA_AMDGPU_DISABLE_ARTIFACTS=1 to pick up the uenv's ROCm. That variable no longer exists after the discovery refactor, so the job silently resolved to the TheRock artifacts instead, and every test died with hipErrorOutOfMemory out of hipStreamCreateWithPriority (3167 occurrences; no hipMalloc ever failed). The .rocm-miopen merge and MIOPEN_PREFIX symlinking the script does were inert as a result. Select the provider through ROCm_Runtime's "local" preference instead, set before Pkg.instantiate() since it is a compile-time preference that also gates artifact resolution, and dev the two workspace packages the way .buildkite/pipeline.yml now does. The resolved provider is asserted against the requested one afterwards, so a future change of default fails loudly instead of quietly swapping the ROCm underneath the job. The script is parameterised by ROCM_LOCAL so both providers can be run. Only the uenv variant is enabled: this is the project's only gfx942 CI, so covering the default artifact path here is worthwhile, but TheRock ships a ROCm 7.14 userspace that cannot create a HIP stream at all against the 6.12.12 amdgpu driver on this machine, so that job is left commented out until the mismatch is resolved. Verified on beverin (2x MI300A, gfx942, Julia 1.12.6) at c644fc5: uenv ROCm, full suite, --jobs=32: 16162 pass, 0 OOM (was 3167 OOM) A/B in one allocation, only the preference differing: local=true HIP 7.2.53211 hipStreamCreateWithPriority OK local=false HIP 7.14.60850 hipErrorOutOfMemory Co-Authored-By: Claude Opus 5 --- ci/cscs-mi300.yml | 102 +++++++++++++++++++++++++++++++--------------- 1 file changed, 69 insertions(+), 33 deletions(-) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index dc00e2750..2243f9049 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -1,53 +1,89 @@ include: - remote: 'https://gitlab.com/cscs-ci/recipes/-/raw/master/templates/v2/.ci-ext.yml' +# Slurm/Julia settings shared by both ROCm provider variants. +.common_variables: &common_variables + JULIA: /users/lraess/julia_amd/juliaup/julia-1.12.6+0.x64.linux.gnu/bin/julia + UENV: ab517d8a08f537d2 + MIOPEN_PREFIX: /user-environment/linux-zen3/miopen-hip-7.2.3-22x3pq72goxi4xwkknlspqhxbx72owes + # Workaround for MIOpen CK Group-Xdlops solvers segfaulting on gfx942 (MI300). + # Remove once fixed upstream: https://github.com/ROCm/rocm-libraries/issues/9088 + MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS_NHWC: "0" + MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS: "0" + MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS: "0" + MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS: "0" + MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS: "0" + SLURM_JOB_NUM_NODES: 1 + SLURM_NTASKS_PER_NODE: 1 + SLURM_GPUS_PER_TASK: 2 + SLURM_TIMELIMIT: "00:45:00" + JULIA_NUM_THREADS: 4 + JULIA_DEPOT_PATH: "${CI_PROJECT_DIR}/.julia" + JULIA_AMDGPU_CORE_MUST_LOAD: "1" + JULIA_AMDGPU_HIP_MUST_LOAD: "1" + +# AMDGPU.jl takes ROCm either from the uenv ($ROCM_LOCAL == "true": the merged +# ROCM_PATH below, resolved through ROCm_Runtime_Discovery) or from the TheRock +# artifacts shipped by ROCm_Runtime ($ROCM_LOCAL == "false"). Which one is used +# is decided by ROCm_Runtime's "local" preference; it is a compile-time +# preference that also gates artifact resolution, so it has to be set before +# Pkg.instantiate(). It is set explicitly in both variants rather than relying +# on the default, and the resolved provider is asserted afterwards, so that a +# change of default cannot silently swap the ROCm underneath a job. .unit_test_script: &unit_test_script - srun -n 1 --uenv $UENV --view=default bash -c ' set -euo pipefail; shopt -s nullglob; - MERGED_ROCM="$CI_PROJECT_DIR/.rocm-miopen"; - rm -rf "$MERGED_ROCM"; - mkdir -p "$MERGED_ROCM/lib"; - for entry in "$ROCM_PATH"/*; do - name=$(basename "$entry"); - [ "$name" = "lib" ] && continue; - ln -s "$entry" "$MERGED_ROCM/$name"; - done; - ln -s "$ROCM_PATH"/lib/* "$MERGED_ROCM/lib/"; - for libdir in "$MIOPEN_PREFIX/lib" "$MIOPEN_PREFIX/lib64"; do - ln -s "$libdir"/libMIOpen* "$MERGED_ROCM/lib/" 2>/dev/null || true; - done; - export ROCM_PATH="$MERGED_ROCM"; + if [ "$ROCM_LOCAL" = "true" ]; then + MERGED_ROCM="$CI_PROJECT_DIR/.rocm-miopen"; + rm -rf "$MERGED_ROCM"; + mkdir -p "$MERGED_ROCM/lib"; + for entry in "$ROCM_PATH"/*; do + name=$(basename "$entry"); + [ "$name" = "lib" ] && continue; + ln -s "$entry" "$MERGED_ROCM/$name"; + done; + ln -s "$ROCM_PATH"/lib/* "$MERGED_ROCM/lib/"; + for libdir in "$MIOPEN_PREFIX/lib" "$MIOPEN_PREFIX/lib64"; do + ln -s "$libdir"/libMIOpen* "$MERGED_ROCM/lib/" 2>/dev/null || true; + done; + export ROCM_PATH="$MERGED_ROCM"; + fi; exec "$0" "$@" ' $JULIA --project -e ' println("Instantiating project"); using Pkg; Pkg.activate(pwd()); + Pkg.develop([PackageSpec(path="ROCm_Runtime"), PackageSpec(path="ROCm_Runtime_Discovery")]); + using Preferences; + want_local = ENV["ROCM_LOCAL"] == "true"; + set_preferences!(Base.UUID("3129f4d2-de71-4ff3-9833-76037e3ea355"), "local" => string(want_local); force=true); Pkg.instantiate(); using AMDGPU; + AMDGPU.versioninfo(); + AMDGPU.local_rocm == want_local || error("requested local_rocm=$want_local, but AMDGPU resolved to local_rocm=$(AMDGPU.local_rocm)"); println("Running tests"); Pkg.test("AMDGPU"; test_args=`--jobs=32`);' -UnitTest julia 1.12: +UnitTest julia 1.12 (system ROCm): extends: .baremetal-runner-beverin-mi300 variables: - JULIA: /users/lraess/julia_amd/juliaup/julia-1.12.6+0.x64.linux.gnu/bin/julia - UENV: ab517d8a08f537d2 - MIOPEN_PREFIX: /user-environment/linux-zen3/miopen-hip-7.2.3-22x3pq72goxi4xwkknlspqhxbx72owes - # Workaround for MIOpen CK Group-Xdlops solvers segfaulting on gfx942 (MI300). - # Remove once fixed upstream: https://github.com/ROCm/rocm-libraries/issues/9088 - MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS_NHWC: "0" - MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS: "0" - MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS: "0" - MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS: "0" - MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS: "0" - SLURM_JOB_NUM_NODES: 1 - SLURM_NTASKS_PER_NODE: 1 - SLURM_GPUS_PER_TASK: 2 - SLURM_TIMELIMIT: "00:45:00" - JULIA_NUM_THREADS: 4 - JULIA_DEPOT_PATH: "${CI_PROJECT_DIR}/.julia" - JULIA_AMDGPU_CORE_MUST_LOAD: "1" - JULIA_AMDGPU_HIP_MUST_LOAD: "1" - JULIA_AMDGPU_DISABLE_ARTIFACTS: "1" + <<: *common_variables + ROCM_LOCAL: "true" script: *unit_test_script + +# Coverage of the artifact path -- what AMDGPU.jl uses by default -- is worth +# having here, since this is the only gfx942 CI in the project. It is disabled +# for now because it cannot get off the ground on this machine: TheRock ships a +# ROCm 7.14 userspace, and against the 6.12.12 amdgpu driver here every +# hipStreamCreateWithPriority fails with hipErrorOutOfMemory before any test +# body runs. Re-enable (with allow_failure until it is green) once that is +# resolved; the ROCM_LOCAL switch above already handles both variants. +# +# UnitTest julia 1.12 (ROCm artifacts): +# extends: .baremetal-runner-beverin-mi300 +# allow_failure: true +# variables: +# <<: *common_variables +# ROCM_LOCAL: "false" +# script: *unit_test_script From d47c99c0965ec0d0769a703db1cff144f648fb64 Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Mon, 17 Aug 2026 17:25:26 +0200 Subject: [PATCH 2/3] CSCS MI300 CI: correct the reason the artifact job is disabled The comment blamed a ROCm 7.14 userspace / 6.12.12 amdgpu driver mismatch. That was wrong: the driver side is fine (rocminfo enumerates all agents, hipMalloc succeeds). What actually fails is ROCclr's blit-kernel build on first stream creation, which cannot resolve __amd_streamOpsIncrement and __amd_streamOpsDecrement, and is then reported as hipErrorOutOfMemory -- which is what made this look like memory exhaustion for so long. Reported upstream as https://github.com/ROCm/TheRock/issues/7426, so the comment now just points there rather than restating the analysis. Co-Authored-By: Claude Opus 5 --- ci/cscs-mi300.yml | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index 2243f9049..9c52d66b6 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -73,12 +73,12 @@ UnitTest julia 1.12 (system ROCm): script: *unit_test_script # Coverage of the artifact path -- what AMDGPU.jl uses by default -- is worth -# having here, since this is the only gfx942 CI in the project. It is disabled -# for now because it cannot get off the ground on this machine: TheRock ships a -# ROCm 7.14 userspace, and against the 6.12.12 amdgpu driver here every -# hipStreamCreateWithPriority fails with hipErrorOutOfMemory before any test -# body runs. Re-enable (with allow_failure until it is green) once that is -# resolved; the ROCM_LOCAL switch above already handles both variants. +# having here, since this is the only gfx942 CI in the project. Disabled for +# now: with the ROCm 7.14 TheRock bundle, every HIP application fails at its +# first stream creation on gfx942, reported as hipErrorOutOfMemory even though +# nothing runs out of memory. See https://github.com/ROCm/TheRock/issues/7426. +# Re-enable, with allow_failure until it is green, once that is fixed; the +# ROCM_LOCAL switch above already handles both variants. # # UnitTest julia 1.12 (ROCm artifacts): # extends: .baremetal-runner-beverin-mi300 From 4fafd6cbf278f9b02c00234a4c0409d125f821da Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Mon, 17 Aug 2026 21:18:41 +0200 Subject: [PATCH 3/3] Fixup --- ci/cscs-mi300.yml | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index 9c52d66b6..9ce0b2957 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -72,8 +72,7 @@ UnitTest julia 1.12 (system ROCm): ROCM_LOCAL: "true" script: *unit_test_script -# Coverage of the artifact path -- what AMDGPU.jl uses by default -- is worth -# having here, since this is the only gfx942 CI in the project. Disabled for +# Coverage of the artifact path, what AMDGPU.jl uses by default. Disabled for # now: with the ROCm 7.14 TheRock bundle, every HIP application fails at its # first stream creation on gfx942, reported as hipErrorOutOfMemory even though # nothing runs out of memory. See https://github.com/ROCm/TheRock/issues/7426.