diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index dc00e2750..9ce0b2957 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -1,53 +1,88 @@ 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. 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 +# allow_failure: true +# variables: +# <<: *common_variables +# ROCM_LOCAL: "false" +# script: *unit_test_script