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 aa41373b94..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,6 +897,11 @@ 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")) 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..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,6 +835,11 @@ 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")) 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..59fe6157ec 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,15 @@ 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 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")) return false; diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index d0c030deb7..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,11 +1331,12 @@ 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 -) +# 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 +#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 +1435,8 @@ 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 +# 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 ) add_custom_test(test_conv_igemm_dynamic_xdlops_nhwc_bwd SKIP_UNLESS_ALL HALF_ENABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON