Skip to content

[W/A] disable group conv in asm igemm solver #1978

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
junliume merged 4 commits into from
Feb 13, 2023
Merged

[W/A] disable group conv in asm igemm solver #1978

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
246919a
[W/A] disable group conv in asm igemm solver
junliume Feb 11, 2023
bac184b
fix for resolver refactor changes
junliume Feb 11, 2023
e0aa9c0
remove test related to the disabled solver in previous commit
carlushuang Feb 11, 2023
020f112
Comply with MIOpen lifecycle of workarounds
junliume Feb 12, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 4 additions & 0 deletions src/include/miopen/conv/asm_implicit_gemm.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,10 @@
#include <vector>
#include <limits>

/// W/A for issue 1979: igemm solver does not support group conv. See:
/// https://github.com/ROCmSoftwarePlatform/MIOpen/issues/1979
Comment on lines +37 to +38

@atamazov atamazov Feb 12, 2023

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

[Notice] Basically, WORKAROUND_ISSUE_1979 says everything so the comment is not required.

#define WORKAROUND_ISSUE_1979 1

namespace miopen {

namespace solver {
Expand Down
5 changes: 5 additions & 0 deletions src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
5 changes: 5 additions & 0 deletions src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
6 changes: 6 additions & 0 deletions src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
21 changes: 8 additions & 13 deletions test/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)

@atamazov atamazov Feb 12, 2023

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

[Notice] The removal of WORKAROUND_ISSUE_1053 is a by-product which is not mentioned in the description.


# Run the test suite to a depth limit
#limit greater than 2 leads to prolonged testing more than 5hrs per stage.
Expand Down Expand Up @@ -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)
Expand Down Expand Up @@ -1338,11 +1331,12 @@ COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $<TARGET_FILE:test_conv2d> ${MIO
COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $<TARGET_FILE:test_conv2d> ${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} $<TARGET_FILE:test_conv2d> ${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} $<TARGET_FILE:test_conv2d> ${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} $<TARGET_FILE:test_conv2d> ${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} $<TARGET_FILE:test_conv2d> ${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} $<TARGET_FILE:test_conv2d> ${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} $<TARGET_FILE:test_conv2d> ${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} $<TARGET_FILE:test_conv2d> ${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
Expand Down Expand Up @@ -1441,7 +1435,8 @@ COMMAND ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} $<TARGET_FILE:test_conv2d>
COMMAND ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} $<TARGET_FILE:test_conv2d> ${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} $<TARGET_FILE:test_conv2d> ${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} $<TARGET_FILE:test_conv2d> ${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} $<TARGET_FILE:test_conv2d> ${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} $<TARGET_FILE:test_conv2d> ${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
Expand Down