From 246919a67bab8013dacaf17af9b15f6cc6c8f9cd Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Fri, 10 Feb 2023 22:50:08 -0800 Subject: [PATCH 1/4] [W/A] disable group conv in asm igemm solver --- src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp | 3 +++ src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp | 3 +++ src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp | 4 ++++ 3 files changed, 10 insertions(+) diff --git a/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp b/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp index aa41373b94..c32e4865ba 100644 --- a/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp +++ b/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp @@ -897,6 +897,9 @@ bool ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::IsApplicable( if(problem.conv_problem.GetConv().attribute.deterministic) return false; + if(ctx.group_counts > 1) + return false; + const auto device_name = ctx.GetStream().GetDeviceName(); if((device_name != "gfx908") && (device_name != "gfx90a")) return false; diff --git a/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp b/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp index e58aee4def..7c6a5ec3ec 100644 --- a/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp +++ b/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp @@ -835,6 +835,9 @@ bool ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::IsApplicable( if(problem.conv_problem.GetConv().attribute.deterministic) return false; + if(ctx.group_counts > 1) + return false; + const auto device_name = ctx.GetStream().GetDeviceName(); if((device_name != "gfx908") && (device_name != "gfx90a")) return false; diff --git a/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp b/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp index 18a33d3be1..dba4139b74 100644 --- a/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp +++ b/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp @@ -819,9 +819,13 @@ bool ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::IsApplicable( { if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS_NHWC{})) return false; + if(problem.conv_problem.GetConv().attribute.deterministic) return false; + if(ctx.group_counts > 1) + return false; + const auto device_name = ctx.GetStream().GetDeviceName(); if((device_name != "gfx908") && (device_name != "gfx90a")) return false; From bac184b7dded2ad594cd8de1eb24fb013ed1a4ab Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Fri, 10 Feb 2023 23:24:34 -0800 Subject: [PATCH 2/4] fix for resolver refactor changes --- src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp | 2 +- src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp | 2 +- src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp b/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp index c32e4865ba..ce534e0ae1 100644 --- a/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp +++ b/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp @@ -897,7 +897,7 @@ bool ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::IsApplicable( if(problem.conv_problem.GetConv().attribute.deterministic) return false; - if(ctx.group_counts > 1) + if(problem.group_counts > 1) return false; const auto device_name = ctx.GetStream().GetDeviceName(); diff --git a/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp b/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp index 7c6a5ec3ec..c577823ff9 100644 --- a/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp +++ b/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp @@ -835,7 +835,7 @@ bool ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::IsApplicable( if(problem.conv_problem.GetConv().attribute.deterministic) return false; - if(ctx.group_counts > 1) + if(problem.group_counts > 1) return false; const auto device_name = ctx.GetStream().GetDeviceName(); diff --git a/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp b/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp index dba4139b74..2c98151963 100644 --- a/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp +++ b/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp @@ -823,7 +823,7 @@ bool ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::IsApplicable( if(problem.conv_problem.GetConv().attribute.deterministic) return false; - if(ctx.group_counts > 1) + if(problem.group_counts > 1) return false; const auto device_name = ctx.GetStream().GetDeviceName(); From e0aa9c0591617819ad37c5f1d5ddd7023cb8f2aa Mon Sep 17 00:00:00 2001 From: carlushuang Date: Sat, 11 Feb 2023 06:59:13 -0800 Subject: [PATCH 3/4] remove test related to the disabled solver in previous commit --- test/CMakeLists.txt | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index d0c030deb7..1aaf7c412d 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1338,11 +1338,11 @@ COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIO COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-weights ) -add_custom_test(test_conv_igemm_dynamic_xdlops_bwd_group SKIP_UNLESS_ALL HALF_ENABLED FLOAT_DISABLED GFX90A_DISABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON -COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 32 28 28 --weights 16 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 --group-count 2 --disable-forward --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 17 17 --weights 64 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 --group-count 4 --disable-forward --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 8 128 56 56 --weights 128 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 8 --disable-forward --disable-backward-weights -) +#add_custom_test(test_conv_igemm_dynamic_xdlops_bwd_group SKIP_UNLESS_ALL HALF_ENABLED FLOAT_DISABLED GFX90A_DISABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON +#COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 32 28 28 --weights 16 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 --group-count 2 --disable-forward --disable-backward-weights +#COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 17 17 --weights 64 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 --group-count 4 --disable-forward --disable-backward-weights +#COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 8 128 56 56 --weights 128 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 8 --disable-forward --disable-backward-weights +#) add_custom_test(test_conv_igemm_dynamic_xdlops_bwd_float SKIP_UNLESS_ALL HALF_DISABLED FLOAT_ENABLED GFX90A_DISABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 512 128 128 --weights 12 512 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights @@ -1441,7 +1441,7 @@ COMMAND ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} $ COMMAND ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 192 17 17 --weights 224 192 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights COMMAND ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 3 17 17 --weights 64 3 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights COMMAND ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 19 19 --weights 510 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 3 224 224 --weights 63 1 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 3 --disable-backward-data --disable-backward-weights +#COMMAND ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 3 224 224 --weights 63 1 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 3 --disable-backward-data --disable-backward-weights ) add_custom_test(test_conv_igemm_dynamic_xdlops_nhwc_bwd SKIP_UNLESS_ALL HALF_ENABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON From 020f112731d2ef067e9505414f36273946e68f26 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Sat, 11 Feb 2023 19:05:54 -0800 Subject: [PATCH 4/4] Comply with MIOpen lifecycle of workarounds --- src/include/miopen/conv/asm_implicit_gemm.hpp | 4 ++++ src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp | 2 ++ src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp | 2 ++ src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp | 2 ++ test/CMakeLists.txt | 9 ++------- 5 files changed, 12 insertions(+), 7 deletions(-) diff --git a/src/include/miopen/conv/asm_implicit_gemm.hpp b/src/include/miopen/conv/asm_implicit_gemm.hpp index 7e72cd309e..8d6b9c5abd 100644 --- a/src/include/miopen/conv/asm_implicit_gemm.hpp +++ b/src/include/miopen/conv/asm_implicit_gemm.hpp @@ -34,6 +34,10 @@ #include #include +/// W/A for issue 1979: igemm solver does not support group conv. See: +/// https://github.com/ROCmSoftwarePlatform/MIOpen/issues/1979 +#define WORKAROUND_ISSUE_1979 1 + namespace miopen { namespace solver { diff --git a/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp b/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp index ce534e0ae1..592219e348 100644 --- a/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp +++ b/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp @@ -897,8 +897,10 @@ bool ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::IsApplicable( if(problem.conv_problem.GetConv().attribute.deterministic) return false; +#if WORKAROUND_ISSUE_1979 if(problem.group_counts > 1) return false; +#endif const auto device_name = ctx.GetStream().GetDeviceName(); if((device_name != "gfx908") && (device_name != "gfx90a")) diff --git a/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp b/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp index c577823ff9..54b92a2f2e 100644 --- a/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp +++ b/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp @@ -835,8 +835,10 @@ bool ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::IsApplicable( if(problem.conv_problem.GetConv().attribute.deterministic) return false; +#if WORKAROUND_ISSUE_1979 if(problem.group_counts > 1) return false; +#endif const auto device_name = ctx.GetStream().GetDeviceName(); if((device_name != "gfx908") && (device_name != "gfx90a")) diff --git a/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp b/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp index 2c98151963..59fe6157ec 100644 --- a/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp +++ b/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp @@ -823,8 +823,10 @@ bool ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::IsApplicable( if(problem.conv_problem.GetConv().attribute.deterministic) return false; +#if WORKAROUND_ISSUE_1979 if(problem.group_counts > 1) return false; +#endif const auto device_name = ctx.GetStream().GetDeviceName(); if((device_name != "gfx908") && (device_name != "gfx90a")) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 1aaf7c412d..0263ecee40 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -48,7 +48,6 @@ set_var_to_condition(MIOPEN_TEST_WITH_MIOPENDRIVER_DEFAULT MIOPEN_BUILD_DRIVER) option( MIOPEN_TEST_WITH_MIOPENDRIVER "Use MIOpenDriver in tests" ${MIOPEN_TEST_WITH_MIOPENDRIVER_DEFAULT}) option( WORKAROUND_ISSUE_936 "" ON) -option( WORKAROUND_ISSUE_1053 "" OFF) # TODO: Remove this W/A after ~6 months (in January 2023) # Run the test suite to a depth limit #limit greater than 2 leads to prolonged testing more than 5hrs per stage. @@ -226,12 +225,6 @@ if (MIOPEN_NO_GPU) test_pooling3d test_perfdb) endif() -if(MIOPEN_TEST_GFX103X OR MIOPEN_TEST_GFX110X) - if(WORKAROUND_ISSUE_1053 AND MIOPEN_TEST_ALL) - list(APPEND SKIP_TESTS test_lrn_test) - endif() -endif() - #TODO Code Quality WORKAROUND ROCm 5.1 update if(MIOPEN_BACKEND_OPENCL AND MIOPEN_TEST_ALL) if(MIOPEN_TEST_GFX900 OR MIOPEN_TEST_GFX906) @@ -1338,6 +1331,7 @@ COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIO COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-weights ) +# TODO: disabled for WORKAROUND_ISSUE_1979 #add_custom_test(test_conv_igemm_dynamic_xdlops_bwd_group SKIP_UNLESS_ALL HALF_ENABLED FLOAT_DISABLED GFX90A_DISABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON #COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 32 28 28 --weights 16 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 --group-count 2 --disable-forward --disable-backward-weights #COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 17 17 --weights 64 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 --group-count 4 --disable-forward --disable-backward-weights @@ -1441,6 +1435,7 @@ COMMAND ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} $ COMMAND ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 192 17 17 --weights 224 192 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights COMMAND ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 3 17 17 --weights 64 3 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights COMMAND ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 19 19 --weights 510 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights +# TODO: disabled for WORKAROUND_ISSUE_1979 #COMMAND ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 3 224 224 --weights 63 1 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 3 --disable-backward-data --disable-backward-weights )