From 0cf594258444639d11674441bec83a4e808bf4f2 Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Sun, 5 Jul 2026 01:23:20 +0200 Subject: [PATCH 01/11] Add 1.11 and discover MIOpen --- ci/cscs-mi300.yml | 61 +++++++++++++++++++++++++++++++++++++---------- 1 file changed, 49 insertions(+), 12 deletions(-) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index f10c7e4cd..03e2eff26 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -1,11 +1,57 @@ include: - remote: 'https://gitlab.com/cscs-ci/recipes/-/raw/master/templates/v2/.ci-ext.yml' -unit_test: +.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"; + exec "$0" "$@" + ' $JULIA --project -e ' + println("Instantiating project"); + using Pkg; + Pkg.activate(pwd()); + Pkg.instantiate(); + using AMDGPU; + AMDGPU.versioninfo(); + println("Running tests"); + Pkg.test("AMDGPU"; test_args=`--jobs=32`);' + +UnitTest julia 1.11: + extends: .baremetal-runner-beverin-mi300 + variables: + JULIA: /users/lraess/julia_amd/juliaup/julia-1.11.9+0.x64.linux.gnu/bin/julia + UENV: ab517d8a08f537d2 + MIOPEN_PREFIX: /user-environment/linux-zen3/miopen-hip-7.2.3-22x3pq72goxi4xwkknlspqhxbx72owes + 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" + script: *unit_test_script + +UnitTest julia 1.12: extends: .baremetal-runner-beverin-mi300 variables: JULIA: /users/lraess/julia_amd/juliaup/julia-1.12.6+0.x64.linux.gnu/bin/julia - UENV: b2550889de318ab5 + UENV: ab517d8a08f537d2 + MIOPEN_PREFIX: /user-environment/linux-zen3/miopen-hip-7.2.3-22x3pq72goxi4xwkknlspqhxbx72owes SLURM_JOB_NUM_NODES: 1 SLURM_NTASKS_PER_NODE: 1 SLURM_GPUS_PER_TASK: 2 @@ -15,13 +61,4 @@ unit_test: JULIA_AMDGPU_CORE_MUST_LOAD: "1" JULIA_AMDGPU_HIP_MUST_LOAD: "1" JULIA_AMDGPU_DISABLE_ARTIFACTS: "1" - script: - - srun -n 1 --uenv $UENV --view=default $JULIA --project -e ' - println("Instantiating project"); - using Pkg; - Pkg.activate(pwd()); - Pkg.instantiate(); - using AMDGPU; - AMDGPU.versioninfo(); - println("Running tests"); - Pkg.test("AMDGPU"; test_args=`--jobs=32`);' + script: *unit_test_script From e91d50e5069b50c308ff874dfcca2c9476cb860b Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Sun, 5 Jul 2026 01:38:48 +0200 Subject: [PATCH 02/11] rm julia 1.11 --- ci/cscs-mi300.yml | 17 ----------------- 1 file changed, 17 deletions(-) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index 03e2eff26..2ce028f73 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -29,23 +29,6 @@ include: println("Running tests"); Pkg.test("AMDGPU"; test_args=`--jobs=32`);' -UnitTest julia 1.11: - extends: .baremetal-runner-beverin-mi300 - variables: - JULIA: /users/lraess/julia_amd/juliaup/julia-1.11.9+0.x64.linux.gnu/bin/julia - UENV: ab517d8a08f537d2 - MIOPEN_PREFIX: /user-environment/linux-zen3/miopen-hip-7.2.3-22x3pq72goxi4xwkknlspqhxbx72owes - 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" - script: *unit_test_script - UnitTest julia 1.12: extends: .baremetal-runner-beverin-mi300 variables: From 5f2cce066d1f6411b3300b2dcd75c7833b6585b5 Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Sun, 5 Jul 2026 21:56:49 +0200 Subject: [PATCH 03/11] [revert] Add mixed matmul support --- ci/cscs-mi300.yml | 2 +- test/hip_dnn/conv.jl | 14 ++++++++++++++ 2 files changed, 15 insertions(+), 1 deletion(-) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index 2ce028f73..7c258346b 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -27,7 +27,7 @@ include: using AMDGPU; AMDGPU.versioninfo(); println("Running tests"); - Pkg.test("AMDGPU"; test_args=`--jobs=32`);' + Pkg.test("AMDGPU"; test_args=`--jobs=32 hip_dnn`);' UnitTest julia 1.12: extends: .baremetal-runner-beverin-mi300 diff --git a/test/hip_dnn/conv.jl b/test/hip_dnn/conv.jl index 11d76cb2b..aa5b06c6b 100644 --- a/test/hip_dnn/conv.jl +++ b/test/hip_dnn/conv.jl @@ -5,6 +5,11 @@ using AMDGPU.MIOpen @assert AMDGPU.functional(:MIOpen) +function _mark(label) + println("=== ISOLATE: ", label, " ===") + flush(stdout) +end + @testset "Simple Convolution" begin for T in (Float16, Float32), nd in 2:3 ndims = Val{nd}() @@ -18,15 +23,18 @@ using AMDGPU.MIOpen x = AMDGPU.ones(T, (spatial_size..., 1, 1)) w = AMDGPU.ones(T, (spatial_size..., 1, 1)) + _mark("Simple Convolution: T=$T nd=$nd forward") y = MIOpen.convolution(x, w; padding, stride, dilation, groups) @test eltype(y) == T @test size(y) == (target_size..., 1, 1) @test Array(y)[1] == prod(spatial_size) Δ = AMDGPU.ones(T, size(y)) + _mark("Simple Convolution: T=$T nd=$nd weight grad") ∇w = MIOpen.∇convolution_weight(Δ, x, w; padding, stride, dilation, groups) @test size(∇w) == size(w) + _mark("Simple Convolution: T=$T nd=$nd data grad") ∇x = MIOpen.∇convolution_data(Δ, x, w; padding, stride, dilation, groups) @test size(∇x) == size(x) end @@ -40,21 +48,26 @@ end wh2 = rand(Float32, 3, 4, 3, 16) x, w1, w2 = ROCArray.((xh, wh1, wh2)) + _mark("padding/stride/dilation: case 1") y = MIOpen.convolution(x, w1; padding=(0, 0), stride=(1, 1), dilation=(1, 1), groups=1) @test size(y) == (9, 9, 16, 10) + _mark("padding/stride/dilation: case 2") y = MIOpen.convolution(x, w1; padding=(2, 2), stride=(2, 2), dilation=(1, 1), groups=1) @test size(y) == (7, 7, 16, 10) + _mark("padding/stride/dilation: case 3") y = MIOpen.convolution(x, w2; padding=(2, 3), stride=(1, 2), dilation=(1, 1), groups=1) @test size(y) == (12, 7, 16, 10) + _mark("padding/stride/dilation: case 4") y = MIOpen.convolution(x, w1; padding=(2, 3), stride=(1, 2), dilation=(2, 2), groups=1) @test size(y) == (12, 7, 16, 10) # Depthwise convolution. wdh1 = rand(Float32, 2, 2, 1, 3) wd1 = ROCArray(wdh1) + _mark("depthwise convolution") y = MIOpen.convolution(x, wd1; padding=(0, 0), stride=(1, 1), dilation=(1, 1), groups=3) @test size(y) == (9, 9, 3, 10) @@ -62,6 +75,7 @@ end xh = ones(Float32, 10, 10, 4, 10) wdh2 = ones(Float32, 2, 2, 2, 4) x, wd2 = ROCArray.((xh, wdh2)) + _mark("grouped convolution") y = MIOpen.convolution(x, wd2; padding=(0, 0), stride=(1, 1), dilation=(1, 1), groups=2) @test size(y) == (9, 9, 4, 10) end From b4f03694397dc2c308faca2b99928baa95ab67e9 Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Sun, 5 Jul 2026 22:12:15 +0200 Subject: [PATCH 04/11] [revert] locate --- ci/cscs-mi300.yml | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index 7c258346b..559a033f1 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -27,7 +27,7 @@ include: using AMDGPU; AMDGPU.versioninfo(); println("Running tests"); - Pkg.test("AMDGPU"; test_args=`--jobs=32 hip_dnn`);' + Pkg.test("AMDGPU"; test_args=`--jobs=32 hip_dnn/conv`);' UnitTest julia 1.12: extends: .baremetal-runner-beverin-mi300 @@ -35,6 +35,7 @@ UnitTest julia 1.12: 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 + MIOPEN_DEBUG_COMPILE_ONLY: "1" SLURM_JOB_NUM_NODES: 1 SLURM_NTASKS_PER_NODE: 1 SLURM_GPUS_PER_TASK: 2 From 994e23b0c5913f420860594a4bcac64c2093e002 Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Sun, 5 Jul 2026 22:31:55 +0200 Subject: [PATCH 05/11] [revert] debug --- ci/cscs-mi300.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index 559a033f1..cc0eb5dd0 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -35,7 +35,7 @@ UnitTest julia 1.12: 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 - MIOPEN_DEBUG_COMPILE_ONLY: "1" + MIOPEN_DEBUG_FIND_ONLY_SOLVER: ConvHipImplicitGemmFwdXdlops SLURM_JOB_NUM_NODES: 1 SLURM_NTASKS_PER_NODE: 1 SLURM_GPUS_PER_TASK: 2 From 84c4617aae470f28a737e22319d24729d5f1ba2a Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Sun, 5 Jul 2026 22:43:59 +0200 Subject: [PATCH 06/11] [revert] debug --- ci/cscs-mi300.yml | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index cc0eb5dd0..3f139f179 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -25,7 +25,6 @@ include: Pkg.activate(pwd()); Pkg.instantiate(); using AMDGPU; - AMDGPU.versioninfo(); println("Running tests"); Pkg.test("AMDGPU"; test_args=`--jobs=32 hip_dnn/conv`);' @@ -35,7 +34,8 @@ UnitTest julia 1.12: 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 - MIOPEN_DEBUG_FIND_ONLY_SOLVER: ConvHipImplicitGemmFwdXdlops + MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS_NHWC: "0" + MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS: "0" SLURM_JOB_NUM_NODES: 1 SLURM_NTASKS_PER_NODE: 1 SLURM_GPUS_PER_TASK: 2 From e39764b9c1c55b84c23ab5dc9efcc5254dfe1fe5 Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Sun, 5 Jul 2026 22:58:01 +0200 Subject: [PATCH 07/11] [revert] debug --- ci/cscs-mi300.yml | 1 + 1 file changed, 1 insertion(+) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index 3f139f179..5da3e3988 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -36,6 +36,7 @@ UnitTest julia 1.12: MIOPEN_PREFIX: /user-environment/linux-zen3/miopen-hip-7.2.3-22x3pq72goxi4xwkknlspqhxbx72owes 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" SLURM_JOB_NUM_NODES: 1 SLURM_NTASKS_PER_NODE: 1 SLURM_GPUS_PER_TASK: 2 From 6a8650bd30c6ca0abee2f87aafa86c6b13fd99cd Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Sun, 5 Jul 2026 23:09:24 +0200 Subject: [PATCH 08/11] [revert] Debug --- ci/cscs-mi300.yml | 1 + 1 file changed, 1 insertion(+) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index 5da3e3988..24d6e6fa4 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -37,6 +37,7 @@ UnitTest julia 1.12: 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_CONV_IMPLICIT_GEMM_HIP_GROUP_BWD_XDLOPS: "0" SLURM_JOB_NUM_NODES: 1 SLURM_NTASKS_PER_NODE: 1 SLURM_GPUS_PER_TASK: 2 From 0cee75fee18a439ecfcb6896fa0a04d0f258b8fb Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Mon, 6 Jul 2026 00:34:03 +0200 Subject: [PATCH 09/11] Undo debug and skip test for failing arch --- ci/cscs-mi300.yml | 3 +-- test/hip_dnn/conv.jl | 26 +++++++++++--------------- 2 files changed, 12 insertions(+), 17 deletions(-) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index 24d6e6fa4..69b370396 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -26,7 +26,7 @@ include: Pkg.instantiate(); using AMDGPU; println("Running tests"); - Pkg.test("AMDGPU"; test_args=`--jobs=32 hip_dnn/conv`);' + Pkg.test("AMDGPU"; test_args=`--jobs=32`);' UnitTest julia 1.12: extends: .baremetal-runner-beverin-mi300 @@ -37,7 +37,6 @@ UnitTest julia 1.12: 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_CONV_IMPLICIT_GEMM_HIP_GROUP_BWD_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 aa5b06c6b..69149a26a 100644 --- a/test/hip_dnn/conv.jl +++ b/test/hip_dnn/conv.jl @@ -5,10 +5,11 @@ using AMDGPU.MIOpen @assert AMDGPU.functional(:MIOpen) -function _mark(label) - println("=== ISOLATE: ", label, " ===") - flush(stdout) -end +# ConvHipImplicitGemmGroupBwdXdlops segfaults on gfx942 (MI300, MIOpen 7.2.3) +# and ignores its own MIOPEN_DEBUG disable flag. Skip until fixed 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)" @testset "Simple Convolution" begin for T in (Float16, Float32), nd in 2:3 @@ -23,20 +24,21 @@ end x = AMDGPU.ones(T, (spatial_size..., 1, 1)) w = AMDGPU.ones(T, (spatial_size..., 1, 1)) - _mark("Simple Convolution: T=$T nd=$nd forward") y = MIOpen.convolution(x, w; padding, stride, dilation, groups) @test eltype(y) == T @test size(y) == (target_size..., 1, 1) @test Array(y)[1] == prod(spatial_size) Δ = AMDGPU.ones(T, size(y)) - _mark("Simple Convolution: T=$T nd=$nd weight grad") ∇w = MIOpen.∇convolution_weight(Δ, x, w; padding, stride, dilation, groups) @test size(∇w) == size(w) - _mark("Simple Convolution: T=$T nd=$nd data grad") - ∇x = MIOpen.∇convolution_data(Δ, x, w; padding, stride, dilation, groups) - @test size(∇x) == size(x) + if _skip_bwd_data + @test_skip false + else + ∇x = MIOpen.∇convolution_data(Δ, x, w; padding, stride, dilation, groups) + @test size(∇x) == size(x) + end end end @@ -48,26 +50,21 @@ end wh2 = rand(Float32, 3, 4, 3, 16) x, w1, w2 = ROCArray.((xh, wh1, wh2)) - _mark("padding/stride/dilation: case 1") y = MIOpen.convolution(x, w1; padding=(0, 0), stride=(1, 1), dilation=(1, 1), groups=1) @test size(y) == (9, 9, 16, 10) - _mark("padding/stride/dilation: case 2") y = MIOpen.convolution(x, w1; padding=(2, 2), stride=(2, 2), dilation=(1, 1), groups=1) @test size(y) == (7, 7, 16, 10) - _mark("padding/stride/dilation: case 3") y = MIOpen.convolution(x, w2; padding=(2, 3), stride=(1, 2), dilation=(1, 1), groups=1) @test size(y) == (12, 7, 16, 10) - _mark("padding/stride/dilation: case 4") y = MIOpen.convolution(x, w1; padding=(2, 3), stride=(1, 2), dilation=(2, 2), groups=1) @test size(y) == (12, 7, 16, 10) # Depthwise convolution. wdh1 = rand(Float32, 2, 2, 1, 3) wd1 = ROCArray(wdh1) - _mark("depthwise convolution") y = MIOpen.convolution(x, wd1; padding=(0, 0), stride=(1, 1), dilation=(1, 1), groups=3) @test size(y) == (9, 9, 3, 10) @@ -75,7 +72,6 @@ end xh = ones(Float32, 10, 10, 4, 10) wdh2 = ones(Float32, 2, 2, 2, 4) x, wd2 = ROCArray.((xh, wdh2)) - _mark("grouped convolution") y = MIOpen.convolution(x, wd2; padding=(0, 0), stride=(1, 1), dilation=(1, 1), groups=2) @test size(y) == (9, 9, 4, 10) end From 95f4b8417641c35924c9b0df0d523787a2249399 Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Mon, 6 Jul 2026 09:17:36 +0200 Subject: [PATCH 10/11] Fixup 3D --- ci/cscs-mi300.yml | 2 ++ 1 file changed, 2 insertions(+) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index 69b370396..8c518558c 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -37,6 +37,8 @@ UnitTest julia 1.12: 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 From f6b052e3e4ac9947ed7951f68ecfac4468cc427f Mon Sep 17 00:00:00 2001 From: Ludovic Raess Date: Mon, 6 Jul 2026 09:34:16 +0200 Subject: [PATCH 11/11] Improve ref --- ci/cscs-mi300.yml | 2 ++ test/hip_dnn/conv.jl | 3 ++- 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/ci/cscs-mi300.yml b/ci/cscs-mi300.yml index 8c518558c..dc00e2750 100644 --- a/ci/cscs-mi300.yml +++ b/ci/cscs-mi300.yml @@ -34,6 +34,8 @@ UnitTest julia 1.12: 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" diff --git a/test/hip_dnn/conv.jl b/test/hip_dnn/conv.jl index 69149a26a..52515cfd2 100644 --- a/test/hip_dnn/conv.jl +++ b/test/hip_dnn/conv.jl @@ -6,7 +6,8 @@ using AMDGPU.MIOpen @assert AMDGPU.functional(:MIOpen) # ConvHipImplicitGemmGroupBwdXdlops segfaults on gfx942 (MI300, MIOpen 7.2.3) -# and ignores its own MIOPEN_DEBUG disable flag. Skip until fixed upstream. +# and ignores its own MIOPEN_DEBUG disable flag. Skip until fixed upstream: +# https://github.com/ROCm/rocm-libraries/issues/9088 _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)"