From 235f79e64bcee3fb90b843ac4c4ab8e16a275321 Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Wed, 23 Sep 2026 21:40:48 +0200 Subject: [PATCH 1/3] Skip MIOpen XDLOPS conv solvers on gfx90a too, set workaround in test Co-Authored-By: Claude Opus 5.5 --- ci/cscs-mi300.yml | 14 -------------- test/hip_dnn/conv.jl | 24 +++++++++++++++--------- 2 files changed, 15 insertions(+), 23 deletions(-) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index c3d653d3d..f5b719cc2 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -35,13 +35,6 @@ UnitTest julia 1.12: JULIA: /users/lraess/julia_amd/julia_amd-1.12 UENV: 56ebb909a6164680 MIOPEN_PREFIX: /user-environment/linux-zen3/miopen-hip-7.2.3-kxzfucnamke4rlezapmrmrdrkmri325u - # 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 @@ -67,13 +60,6 @@ UnitTest julia 1.13: JULIA: /users/lraess/julia_amd/julia_amd-1.13 UENV: 56ebb909a6164680 MIOPEN_PREFIX: /user-environment/linux-zen3/miopen-hip-7.2.3-kxzfucnamke4rlezapmrmrdrkmri325u - # 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 diff --git a/test/hip_dnn/conv.jl b/test/hip_dnn/conv.jl index ce44f70b1..baca1f6fc 100644 --- a/test/hip_dnn/conv.jl +++ b/test/hip_dnn/conv.jl @@ -5,16 +5,22 @@ using AMDGPU.MIOpen @assert AMDGPU.functional(:MIOpen) -# ConvHipImplicitGemmGroupBwdXdlops segfaults in MIOpen's CK grouped-conv invoker -# setup. Not gfx942-specific (reproduces on gfx90a too) and not fixed by rebuilding -# MIOpen with matching gfx942 code objects; we gate on gfx942 only because that is -# where CI runs. The solver's disable flag is inverted upstream: -# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_GROUP_BWD_XDLOPS=1 disables it, =0 does not, -# so we skip rather than depend on that. Skip until fixed upstream: +# MIOpen's CK grouped-conv XDLOPS solvers segfault on gfx90a/gfx942: # https://github.com/ROCm/rocm-libraries/issues/9088 +# Disable fwd/wrw ones (read once, so before the first convolution) and skip +# bwd-data, whose disable flag is inverted upstream. _arch_str = first(split(AMDGPU.HIP.gcn_arch(AMDGPU.device()), ':')) -_skip_bwd_data = _arch_str == "gfx942" -_skip_bwd_data && @info "Skipping convolution backward-data tests (MIOpen bug on gfx942)" +_xdlops_bug = _arch_str in ("gfx90a", "gfx942") +if _xdlops_bug + for var in ("MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS_NHWC", + "MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS", + "MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS", + "MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS", + "MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS") + get!(ENV, var, "0") + end + @info "Skipping convolution backward-data tests (MIOpen bug on $_arch_str)" +end @testset "Simple Convolution" begin for T in (Float16, Float32), nd in 2:3 @@ -38,7 +44,7 @@ _skip_bwd_data && @info "Skipping convolution backward-data tests (MIOpen bug on ∇w = MIOpen.∇convolution_weight(Δ, x, w; padding, stride, dilation, groups) @test size(∇w) == size(w) - if _skip_bwd_data + if _xdlops_bug @test_skip false else ∇x = MIOpen.∇convolution_data(Δ, x, w; padding, stride, dilation, groups) From 27d447954c069265cd699af97052ad6c5cb9a713 Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Wed, 23 Sep 2026 22:00:37 +0200 Subject: [PATCH 2/3] Add CSCS MI250 CI pipeline --- ci/cscs-mi250.yml | 80 +++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 80 insertions(+) create mode 100644 ci/cscs-mi250.yml diff --git a/ci/cscs-mi250.yml b/ci/cscs-mi250.yml new file mode 100644 index 000000000..467247c21 --- /dev/null +++ b/ci/cscs-mi250.yml @@ -0,0 +1,80 @@ +include: + - remote: 'https://gitlab.com/cscs-ci/recipes/-/raw/master/templates/v2/.ci-ext.yml' + +.unit_test_script: &unit_test_script + - srun -n 1 --uenv $UENV --view=default bash -c ' + set -euo pipefail; + ulimit -c 0; + 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"; + exec "$0" "$@" + ' $JULIA --project -e ' + println("Instantiating project"); + using Pkg; + Pkg.activate(pwd()); + Pkg.instantiate(); + using AMDGPU; + println("Running tests"); + Pkg.test("AMDGPU"; test_args=`--jobs=32`);' + +UnitTest julia 1.12: + extends: .baremetal-runner-beverin-mi200 + variables: + JULIA: /users/lraess/julia_amd/julia_amd-1.12 + UENV: 56ebb909a6164680 + MIOPEN_PREFIX: /user-environment/linux-zen3/miopen-hip-7.2.3-kxzfucnamke4rlezapmrmrdrkmri325u + SLURM_JOB_NUM_NODES: 1 + SLURM_NTASKS_PER_NODE: 1 + SLURM_GPUS_PER_TASK: 2 + SLURM_TIMELIMIT: "01:00:00" + JULIA_NUM_THREADS: 4 + JULIA_DEPOT_PATH: "${CI_PROJECT_DIR}/.julia" # overrides default ~/.julia + JULIA_AMDGPU_CORE_MUST_LOAD: "1" + JULIA_AMDGPU_HIP_MUST_LOAD: "1" + JULIA_AMDGPU_DISABLE_ARTIFACTS: "1" + # Trace layer: prints the exception behind an opaque rocfft_status_failure. + ROCFFT_LAYER: "1" + ROCFFT_LOG_TRACE_PATH: "${CI_PROJECT_DIR}/rocfft-trace.log" + script: *unit_test_script + artifacts: + when: always + expire_in: 1 week + paths: + - rocfft-trace.log + +UnitTest julia 1.13: + extends: .baremetal-runner-beverin-mi200 + variables: + JULIA: /users/lraess/julia_amd/julia_amd-1.13 + UENV: 56ebb909a6164680 + MIOPEN_PREFIX: /user-environment/linux-zen3/miopen-hip-7.2.3-kxzfucnamke4rlezapmrmrdrkmri325u + SLURM_JOB_NUM_NODES: 1 + SLURM_NTASKS_PER_NODE: 1 + SLURM_GPUS_PER_TASK: 2 + SLURM_TIMELIMIT: "01:00:00" + JULIA_NUM_THREADS: 4 + JULIA_DEPOT_PATH: "${CI_PROJECT_DIR}/.julia" # overrides default ~/.julia + JULIA_AMDGPU_CORE_MUST_LOAD: "1" + JULIA_AMDGPU_HIP_MUST_LOAD: "1" + JULIA_AMDGPU_DISABLE_ARTIFACTS: "1" + # Trace layer: prints the exception behind an opaque rocfft_status_failure. + ROCFFT_LAYER: "1" + ROCFFT_LOG_TRACE_PATH: "${CI_PROJECT_DIR}/rocfft-trace.log" + script: *unit_test_script + artifacts: + when: always + expire_in: 1 week + paths: + - rocfft-trace.log From ac8071f8d5a5cb195fb0c0a93b6628bb7fcb4a6a Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Wed, 23 Sep 2026 22:23:44 +0200 Subject: [PATCH 3/3] Fix MI250 test failures: hex arch parse in WMMA tests, bounds-check-free codegen kernel Co-Authored-By: Claude Opus 5.5 --- test/core/codegen.jl | 4 +++- test/wmma_rdna3_tests.jl | 6 +++--- test/wmma_rdna4_tests.jl | 4 ++-- 3 files changed, 8 insertions(+), 6 deletions(-) diff --git a/test/core/codegen.jl b/test/core/codegen.jl index 64c078f37..24c72164d 100644 --- a/test/core/codegen.jl +++ b/test/core/codegen.jl @@ -87,7 +87,9 @@ end j = (workgroupIdx().y - 1) * workgroupDim().y + workitemIdx().y k = (workgroupIdx().z - 1) * workgroupDim().z + workitemIdx().z n = i + j + k - n <= length(A) && (@inbounds A[n] = n) + # not `A[n] = n`: under --check-bounds=yes its exception path emits + # private-memory buffer_loads on gfx90a + n <= length(A) && unsafe_store!(pointer(A), Float32(n), n) return end diff --git a/test/wmma_rdna3_tests.jl b/test/wmma_rdna3_tests.jl index bab48728c..c18918404 100644 --- a/test/wmma_rdna3_tests.jl +++ b/test/wmma_rdna3_tests.jl @@ -8,9 +8,9 @@ AMDGPU.allowscalar(false) # Only run WMMA_RDNA3 tests on RDNA3+ (gfx1100+) _arch_str = first(split(AMDGPU.HIP.gcn_arch(AMDGPU.device()), ':')) -gfx = parse(Int, _arch_str[4:end]) -is_rdna3 = 1100 ≤ gfx < 1200 -is_rdna4 = 1200 ≤ gfx < 1300 +gfx = parse(Int, _arch_str[4:end]; base=16) # hex: gfx90a +is_rdna3 = 0x1100 ≤ gfx < 0x1200 +is_rdna4 = 0x1200 ≤ gfx < 0x1300 if !is_rdna3 && !is_rdna4 @info "Skipping WMMA_RDNA3 tests (requires RDNA3+)" else diff --git a/test/wmma_rdna4_tests.jl b/test/wmma_rdna4_tests.jl index 92f5e7040..b8b7a28b7 100644 --- a/test/wmma_rdna4_tests.jl +++ b/test/wmma_rdna4_tests.jl @@ -8,8 +8,8 @@ AMDGPU.allowscalar(false) # Only run WMMA_RDNA4 tests on RDNA4+ (gfx1200+) _arch_str = first(split(AMDGPU.HIP.gcn_arch(AMDGPU.device()), ':')) -gfx = parse(Int, _arch_str[4:end]) -is_rdna4 = 1200 <= gfx < 1300 +gfx = parse(Int, _arch_str[4:end]; base=16) # hex: gfx90a +is_rdna4 = 0x1200 <= gfx < 0x1300 if !is_rdna4 @info "Skipping WMMA_RDNA4 tests (requires RDNA4+ / gfx1200+)"