Skip to content
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

[HOTFIX] WORKAROUND_ISSUE_2038: Disable FP16 and BF16 in ConvHipImplicitGemmV4R1Fwd for MI100/200 and all new targets. #2041

Closed
wants to merge 5 commits into from
Closed

[HOTFIX] WORKAROUND_ISSUE_2038: Disable FP16 and BF16 in ConvHipImplicitGemmV4R1Fwd for MI100/200 and all new targets. #2041

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
bdbaf2e
wa-issue-2038(01) Disable BF16 in ConvHipImplicitGemmV4R1Fwd for MI10…
atamazov Mar 23, 2023
e84170c
wa-issue-2038(02) [tests] Do not smoke test ConvHipImplicitGemmV4R1Fw…
atamazov Mar 23, 2023
9ee84e0
Disable FP16 in ConvHipImplicitGemmV4R1Fwd for MI100/200 and all new …
junliume Mar 24, 2023
9538f39
wa-issue-2038(03) Disable also FP16.
atamazov Mar 24, 2023
1364cf2
Merge branch 'wa-issue-2038' of https://github.com/atamazov/MIOpen in…
atamazov Mar 24, 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
16 changes: 14 additions & 2 deletions src/solver/conv_hip_implicit_gemm_fwd_v4r1.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,8 @@

#include <cstddef>

#define WORKAROUND_ISSUE_2038 1

MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1)
MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R1)

Expand Down Expand Up @@ -61,10 +63,20 @@ bool ConvHipImplicitGemmV4R1Fwd::IsApplicable(const ConvolutionContext& ctx,
return false;
if(!problem.IsLayoutDefault())
return false;
if(ctx.GetStream().GetDeviceName() == "gfx90a" &&
problem.conv_problem.IsGfx90aFp16altRequired())
const auto device_name = ctx.GetStream().GetDeviceName();
if(device_name == "gfx90a" && problem.conv_problem.IsGfx90aFp16altRequired())
return false;

#if WORKAROUND_ISSUE_2038
if(!miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1{}))
if(problem.IsBfp16() || problem.IsFp16())
// Explicitly enable all currently known GPUs except xDLOPs ones, which also means that
// all new GPUs will be disabled by default.
if(!(device_name == "gfx803" || device_name == "gfx900" || device_name == "gfx906" ||
StartsWith(device_name, "gfx103")))
return false;
#endif

std::size_t n = problem.batch_sz;
std::size_t k = problem.n_outputs / problem.group_counts;
std::size_t c = problem.n_inputs / problem.group_counts;
Expand Down
17 changes: 16 additions & 1 deletion test/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2069,18 +2069,33 @@ add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV1R1 GFX103X_ENABLED TEST

# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1=1 is necessary due to WORKAROUND_iGemm_936 in Jenkinsfile,
# which disables ConvHipImplicitGemmV4R1Fwd, but we still want to check that the solver is not broken.
add_custom_test(smoke_solver_ConvHipImplicitGemmV4R1 GFX103X_ENABLED HALF_ENABLED BF16_ENABLED TEST_TUNING
add_custom_test(smoke_solver_ConvHipImplicitGemmV4R1Fwd_fp32 GFX103X_ENABLED TEST_TUNING
COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5
MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1=1
MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0
MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1Fwd $<TARGET_FILE:test_conv2d>
${TEST_CONV_VERBOSE_F} --input 256 32 27 27 --weights 128 32 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
)

add_custom_test(smoke_solver_ConvHipImplicitGemmV4R1WrW GFX103X_ENABLED HALF_ENABLED BF16_ENABLED TEST_TUNING
COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5
MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0
MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1WrW $<TARGET_FILE:test_conv2d>
${TEST_CONV_VERBOSE_W} --input 64 64 55 55 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
)

# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1=1 is necessary due to WORKAROUND_iGemm_936 in Jenkinsfile,
# which disables ConvHipImplicitGemmV4R1Fwd, but we still want to check that the solver is not broken.
# smoke_solver_ConvHipImplicitGemmV4R1Fwd is split to FP16+BF16 and FP32 tests because of
# WORKAROUND_ISSUE_2038, which disables GFX908 and GFX90A for FP16 and BF16.
add_custom_test(smoke_solver_ConvHipImplicitGemmV4R1Fwd_fp16_bf16 GFX908_DISABLED GFX90A_DISABLED GFX103X_ENABLED FLOAT_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING
COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5
MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1=1
MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0
MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1Fwd $<TARGET_FILE:test_conv2d>
${TEST_CONV_VERBOSE_F} --input 256 32 27 27 --weights 128 32 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
)

# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1=1 is necessary due to WORKAROUND_SWDEV_229277_227616_229195,
# which disables ConvHipImplicitGemmBwdDataV4R1, but we still want to check that the solver is not broken.
add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV4R1 GFX103X_ENABLED TEST_TUNING
Expand Down