From 02f6520a9c6438b649d90673877594a9b5270fe5 Mon Sep 17 00:00:00 2001 From: carlushuang Date: Mon, 23 Oct 2023 04:07:48 -0400 Subject: [PATCH 1/3] fix fwd trans conv get_workspace API --- src/convolution_api.cpp | 12 ++---------- 1 file changed, 2 insertions(+), 10 deletions(-) diff --git a/src/convolution_api.cpp b/src/convolution_api.cpp index 555b90f926..fb18011594 100644 --- a/src/convolution_api.cpp +++ b/src/convolution_api.cpp @@ -56,16 +56,8 @@ static inline auto MakeFwdCtxAndProblem(miopenHandle_t handle, const auto direction = (conv.mode != miopenTranspose) ? Direction::Forward : Direction::BackwardData; - auto problem = (conv.mode != miopenTranspose) ? ProblemDescription{miopen::deref(xDesc), - miopen::deref(wDesc), - miopen::deref(yDesc), - conv, - direction} - : ProblemDescription{miopen::deref(yDesc), - miopen::deref(wDesc), - miopen::deref(xDesc), - conv, - direction}; + auto problem = ProblemDescription{ + miopen::deref(xDesc), miopen::deref(wDesc), miopen::deref(yDesc), conv, direction}; auto ctx = ExecutionContext{&miopen::deref(handle)}; problem.SetupFloats(ctx); From 08162003d1395a33f7a2d6caf8dcc734229e8e9d Mon Sep 17 00:00:00 2001 From: carlushuang Date: Mon, 23 Oct 2023 06:03:44 -0400 Subject: [PATCH 2/3] disable bwd_gemm solver --- src/solver/gemm_bwd.cpp | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/src/solver/gemm_bwd.cpp b/src/solver/gemm_bwd.cpp index df7d08304b..b00f345b92 100644 --- a/src/solver/gemm_bwd.cpp +++ b/src/solver/gemm_bwd.cpp @@ -43,6 +43,7 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_CONV_PRECISE_ROCBLAS_TIMING) #define WORKAROUND_MIOPENGEMM_ISSUE_59 1 +#define WORKAROUND_MIOPENGEMM_ISSUE_2474 1 // copy from convolution.cpp // Workaround for issue 1430. @@ -246,6 +247,9 @@ bool GemmBwd1x1_stride2::IsApplicable(const ExecutionContext& context, const conv::ProblemDescription& problem) const { #if MIOPEN_USE_GEMM +#if WORKAROUND_MIOPENGEMM_ISSUE_2474 + return false; +#endif if(!GemmBwdBase::IsApplicable(context, problem)) return false; @@ -482,7 +486,9 @@ bool GemmBwd1x1_stride1::IsApplicableBeforeWorkaround(const ExecutionContext& co bool GemmBwd1x1_stride1::IsApplicable(const ExecutionContext& context, const conv::ProblemDescription& problem) const { -#if MIOPEN_USE_GEMM && (!MIOPEN_USE_MIOPENGEMM || !WORKAROUND_MIOPENGEMM_ISSUE_59) +#if WORKAROUND_MIOPENGEMM_ISSUE_2474 + return false; +#elif MIOPEN_USE_GEMM && (!MIOPEN_USE_MIOPENGEMM || !WORKAROUND_MIOPENGEMM_ISSUE_59) return IsApplicableBeforeWorkaround(context, problem); #else std::ignore = context; @@ -683,6 +689,9 @@ bool GemmBwdRest::IsApplicable(const ExecutionContext& context, const conv::ProblemDescription& problem) const { #if MIOPEN_USE_GEMM +#if WORKAROUND_MIOPENGEMM_ISSUE_2474 + return false; +#endif if(!GemmBwdBase::IsApplicable(context, problem)) return false; From 81ce1bec6b01bb87c50d549dd6e87797708b30fd Mon Sep 17 00:00:00 2001 From: carlushuang Date: Tue, 24 Oct 2023 01:24:04 -0400 Subject: [PATCH 3/3] Revert "disable bwd_gemm solver" This reverts commit 08162003d1395a33f7a2d6caf8dcc734229e8e9d. --- src/solver/gemm_bwd.cpp | 11 +---------- 1 file changed, 1 insertion(+), 10 deletions(-) diff --git a/src/solver/gemm_bwd.cpp b/src/solver/gemm_bwd.cpp index b00f345b92..df7d08304b 100644 --- a/src/solver/gemm_bwd.cpp +++ b/src/solver/gemm_bwd.cpp @@ -43,7 +43,6 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_CONV_PRECISE_ROCBLAS_TIMING) #define WORKAROUND_MIOPENGEMM_ISSUE_59 1 -#define WORKAROUND_MIOPENGEMM_ISSUE_2474 1 // copy from convolution.cpp // Workaround for issue 1430. @@ -247,9 +246,6 @@ bool GemmBwd1x1_stride2::IsApplicable(const ExecutionContext& context, const conv::ProblemDescription& problem) const { #if MIOPEN_USE_GEMM -#if WORKAROUND_MIOPENGEMM_ISSUE_2474 - return false; -#endif if(!GemmBwdBase::IsApplicable(context, problem)) return false; @@ -486,9 +482,7 @@ bool GemmBwd1x1_stride1::IsApplicableBeforeWorkaround(const ExecutionContext& co bool GemmBwd1x1_stride1::IsApplicable(const ExecutionContext& context, const conv::ProblemDescription& problem) const { -#if WORKAROUND_MIOPENGEMM_ISSUE_2474 - return false; -#elif MIOPEN_USE_GEMM && (!MIOPEN_USE_MIOPENGEMM || !WORKAROUND_MIOPENGEMM_ISSUE_59) +#if MIOPEN_USE_GEMM && (!MIOPEN_USE_MIOPENGEMM || !WORKAROUND_MIOPENGEMM_ISSUE_59) return IsApplicableBeforeWorkaround(context, problem); #else std::ignore = context; @@ -689,9 +683,6 @@ bool GemmBwdRest::IsApplicable(const ExecutionContext& context, const conv::ProblemDescription& problem) const { #if MIOPEN_USE_GEMM -#if WORKAROUND_MIOPENGEMM_ISSUE_2474 - return false; -#endif if(!GemmBwdBase::IsApplicable(context, problem)) return false;