Skip broken MIOpen XDLOPS conv solvers on gfx90a too, and move the workaround into the test

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

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

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
 

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)

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

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+)"