diff --git a/doc/src/find_and_immediate.md b/doc/src/find_and_immediate.md index 0c6d10cced..a866e34ecd 100644 --- a/doc/src/find_and_immediate.md +++ b/doc/src/find_and_immediate.md @@ -167,10 +167,10 @@ MIOpen provides a set of Find modes which are used to accelerate the Find calls. - `NORMAL`, or `1`: Normal Find: This is the full Find mode call, which will benchmark all the solvers and return a list. - `FAST`, or `2`: Fast Find: Checks the [Find-Db](https://rocmsoftwareplatform.github.io/MIOpen/doc/html/finddb.html) for an entry. If there is a Find-Db hit, use that entry. If there is a miss, utilize the Immediate mode fallback. If Start-up times are expected to be faster, but worse GPU performance. - `HYBRID`, or `3`, or unset `MIOPEN_FIND_MODE`: Hybrid Find: Checks the [Find-Db](https://rocmsoftwareplatform.github.io/MIOpen/doc/html/finddb.html) for an entry. If there is a Find-Db hit, use that entry. If there is a miss, use the existing Find machinery. Slower start-up times than Fast Find, but no GPU performance drop. -- `FAST_HYBRID`, or `4`: Fast Hybrid Find: Checks the [Find-Db](https://rocmsoftwareplatform.github.io/MIOpen/doc/html/finddb.html) for an entry. If there is a Find-Db hit, uses that entry. If there is a miss, uses the existing Find machinery with skipping slow-compiling kernels. Faster start-up times than Hybrid Find, but GPU performance is a bit worse. -- `DYNAMIC_HYBRID`, or `5`: Dynamic Hybrid Find: This mode is similar to Fast Hybrid, but in case of Find-db miss, skips all non-dynamic kernels, thus saving compilation time. Versus FAST_HYBRID, we expect similar start-up times but better GPU performance. Use with caution, this mode is experimental for now. +- `4`: This value is reserved and should not be used. +- `DYNAMIC_HYBRID`, or `5`: Dynamic Hybrid Find: Checks the [Find-Db](https://rocmsoftwareplatform.github.io/MIOpen/doc/html/finddb.html) for an entry. If there is a Find-Db hit, uses that entry. If there is a miss, uses the existing Find machinery with skipping non-dynamic kernels. Faster start-up times than Hybrid Find, but GPU performance may be a bit worse. - As of MIOpen 2.7, the default mode is set to `HYBRID` mode as default. To run the full `NORMAL` Find mode, set the environment as: + Currently, the default Find mode is `DYNAMIC_HYBRID`. To run the full `NORMAL` Find mode, set the environment as: ``` export MIOPEN_FIND_MODE=NORMAL ``` diff --git a/src/convolution.cpp b/src/convolution.cpp index c0fa29acf6..f1b4485027 100644 --- a/src/convolution.cpp +++ b/src/convolution.cpp @@ -358,8 +358,6 @@ std::size_t ConvolutionDescriptor::ForwardGetWorkSpaceSize(Handle& handle, GetForwardSolutions(handle, wDesc, xDesc, yDesc, 1, &count, &sol, &fallback); if(count < 1 || (findMode.IsHybrid(ctx) && fallback)) { - ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage = - findMode.IsFastHybrid(ctx); ctx.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); break; // Fall down to Normal Find. } @@ -438,8 +436,6 @@ ConvolutionDescriptor::BackwardDataGetWorkSpaceSize(Handle& handle, GetBackwardSolutions(handle, dyDesc, wDesc, dxDesc, 1, &count, &sol, &fallback); if(count < 1 || (findMode.IsHybrid(ctx) && fallback)) { - ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage = - findMode.IsFastHybrid(ctx); ctx.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); break; // Fall down to Normal Find. } @@ -730,8 +726,6 @@ ConvolutionDescriptor::BackwardWeightsGetWorkSpaceSize(Handle& handle, GetWrwSolutions(handle, dyDesc, xDesc, dwDesc, 1, &count, &sol, &fallback); if(count < 1 || (findMode.IsHybrid(ctx) && fallback)) { - ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage = - findMode.IsFastHybrid(ctx); ctx.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); break; // Fall down to Normal Find. } diff --git a/src/find_controls.cpp b/src/find_controls.cpp index f4e2458f76..a12e27689e 100644 --- a/src/find_controls.cpp +++ b/src/find_controls.cpp @@ -165,7 +165,7 @@ const char* ToCString(const FindMode::Values mode) case FindMode::Values::Normal: return "NORMAL"; case FindMode::Values::Fast: return "FAST"; case FindMode::Values::Hybrid: return "HYBRID"; - case FindMode::Values::FastHybrid: return "FAST_HYBRID"; + case FindMode::Values::DeprecatedFastHybrid: break; case FindMode::Values::DynamicHybrid: return "DYNAMIC_HYBRID"; case FindMode::Values::End_: break; } @@ -191,8 +191,6 @@ FindMode::Values GetFindModeValueImpl2() return FindMode::Values::Fast; else if(str == "HYBRID") return FindMode::Values::Hybrid; - else if(str == "FAST_HYBRID") - return FindMode::Values::FastHybrid; else if(str == "DYNAMIC_HYBRID") return FindMode::Values::DynamicHybrid; else @@ -232,9 +230,6 @@ static_assert(miopenConvolutionFindModeFast == static_assert(miopenConvolutionFindModeHybrid == static_cast(FindMode::Values::Hybrid), "API is not in sync with the implementation."); -static_assert(miopenConvolutionFindModeFastHybrid == - static_cast(FindMode::Values::FastHybrid), - "API is not in sync with the implementation."); static_assert(miopenConvolutionFindModeDynamicHybrid == static_cast(FindMode::Values::DynamicHybrid), "API is not in sync with the implementation."); diff --git a/src/include/miopen/execution_context.hpp b/src/include/miopen/execution_context.hpp index 50ba586304..565f53414d 100644 --- a/src/include/miopen/execution_context.hpp +++ b/src/include/miopen/execution_context.hpp @@ -79,9 +79,8 @@ struct ExecutionContext // to optimize the getWorkspaceSize() calls for speed. This specific optimization is correct // because Solvers shall be written so that the required workspace size does not depend on the // performance config. - bool disable_perfdb_access = false; - bool skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage = false; - bool use_dynamic_solutions_only = false; + bool disable_perfdb_access = false; + bool use_dynamic_solutions_only = false; inline Handle& GetStream() const { return *stream; } inline void SetStream(Handle* stream_) { stream = stream_; } @@ -275,25 +274,18 @@ struct ExecutionContext class AutoUseFastDynamicSolutions { - bool prev_skip_slow_; bool prev_use_dynamic_; ExecutionContext* const ctx; public: AutoUseFastDynamicSolutions(ExecutionContext& ctx_) : ctx(&ctx_) { - prev_skip_slow_ = ctx->skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage; prev_use_dynamic_ = ctx->use_dynamic_solutions_only; - ctx->skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage = true; - ctx->use_dynamic_solutions_only = true; + ctx->use_dynamic_solutions_only = true; } - ~AutoUseFastDynamicSolutions() - { - ctx->skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage = prev_skip_slow_; - ctx->use_dynamic_solutions_only = prev_use_dynamic_; - } + ~AutoUseFastDynamicSolutions() { ctx->use_dynamic_solutions_only = prev_use_dynamic_; } }; bool IsHipKernelsEnabled(const TargetProperties& target); diff --git a/src/include/miopen/find_controls.hpp b/src/include/miopen/find_controls.hpp index b27239af3c..5bd6c7edcb 100644 --- a/src/include/miopen/find_controls.hpp +++ b/src/include/miopen/find_controls.hpp @@ -114,7 +114,7 @@ class FindMode Normal = Begin_, Fast, Hybrid, - FastHybrid, + DeprecatedFastHybrid, DynamicHybrid, End_, Default_ = DynamicHybrid, @@ -148,15 +148,7 @@ class FindMode template bool IsHybrid(const Context& context) const { - return (value == Values::Hybrid || value == Values::FastHybrid || - value == Values::DynamicHybrid) && - IsEnabled(context); - } - - template - bool IsFastHybrid(const Context& context) const - { - return value == Values::FastHybrid && IsEnabled(context); + return (value == Values::Hybrid || value == Values::DynamicHybrid) && IsEnabled(context); } template diff --git a/src/include/miopen/miopen_internal.h b/src/include/miopen/miopen_internal.h index ccc8ecd997..93ead3340b 100644 --- a/src/include/miopen/miopen_internal.h +++ b/src/include/miopen/miopen_internal.h @@ -60,20 +60,17 @@ extern "C" { * use the existing Find machinery. Slower start-up times than Fast Find, but no GPU performance * drop. * - * * Fast Hybrid: Checks the Find-db for an entry. If there is a hit, use that entry. If there is a - * miss, uses the existing Find machinery with skipping slow-compiling kernels. Faster start-up - * times than Hybrid Find, but GPU performance is a bit worse. - * - * * Dynamic Hybrid: This mode is similar to Fast Hybrid, but in case of Find-db miss skips all - * non-dynamic kernels, thus saving compilation time. Versus Fast Hybrid, we expect similar start-up - * times but better GPU performance. + * * Dynamic Hybrid: Checks the Find-db for an entry. If there is a hit, uses that entry. If there + * is a miss, uses the existing Find machinery with skipping non-dynamic kernels, thus saving + * compilation time.slow-compiling kernels. Faster start-up times than Hybrid Find, but GPU + * performance may be a bit worse. */ typedef enum { miopenConvolutionFindModeNormal = 1, /*!< Normal mode */ miopenConvolutionFindModeFast = 2, /*!< Fast mode */ miopenConvolutionFindModeHybrid = 3, /*!< Hybrid mode */ - miopenConvolutionFindModeFastHybrid = 4, /*!< Fast Hybrid mode */ + miopenConvolutionFindModeReserved_4 = 4, /*!< Reserved - do not use */ miopenConvolutionFindModeDynamicHybrid = 5, /*!< Dynamic Hybrid mode */ miopenConvolutionFindModeDefault = miopenConvolutionFindModeDynamicHybrid, /*!< Default setting */ diff --git a/src/ocl/convolutionocl.cpp b/src/ocl/convolutionocl.cpp index 88367b8fbc..fe95fc509f 100755 --- a/src/ocl/convolutionocl.cpp +++ b/src/ocl/convolutionocl.cpp @@ -181,8 +181,6 @@ ConvolutionDescriptor::FindDataDirectSolutions(Handle& handle, const auto dir = isForward ? conv::Direction::Forward : conv::Direction::BackwardData; auto ctx = ConvolutionContext{xDesc, wDesc, yDesc, *this, dir}; - ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage = - findMode.IsFastHybrid(ctx); ctx.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); ctx.do_search = exhaustiveSearch; ctx.save_srch_req = true; @@ -220,8 +218,6 @@ ConvolutionDescriptor::FindDataImplicitGemmSolutions(Handle& handle, const auto dir = isForward ? conv::Direction::Forward : conv::Direction::BackwardData; auto ctx = ConvolutionContext{xDesc, wDesc, yDesc, *this, dir}; - ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage = - findMode.IsFastHybrid(ctx); ctx.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); ctx.do_search = exhaustiveSearch; ctx.save_srch_req = true; @@ -494,8 +490,6 @@ void ConvolutionDescriptor::FindConvFwdAlgorithm(Handle& handle, ConvolutionUserBuffers bufs(workSpace, workSpaceSize); bufs.SetFwd(x, w, y); ctx.SetBufs(bufs); - ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage = - findMode.IsFastHybrid(ctx); ctx.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); perf_db = UserFindDbRecord::TryLoad(handle, problem, [&](DbRecord& record) { DirConvFindCore(handle, @@ -1130,15 +1124,12 @@ void ConvolutionDescriptor::FindConvBwdDataAlgorithm(Handle& handle, }(); perf_db = UserFindDbRecord::TryLoad(handle, problem, [&](DbRecord& record) { - const auto network_config = problem.BuildConfKey(); - const auto invoke_ctx = conv::DataInvokeParams{InvokeType::Evaluate, + const auto network_config = problem.BuildConfKey(); + const auto invoke_ctx = conv::DataInvokeParams{InvokeType::Evaluate, {dyDesc, dy, wDesc, w, dxDesc, dx}, workSpace, workSpaceSize, this->attribute.gfx90aFp16alt.GetBwd()}; - - ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage = - findMode.IsFastHybrid(ctx); ctx.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); // Find solutions @@ -1494,8 +1485,6 @@ void ConvolutionDescriptor::FindConvBwdWeightsAlgorithm(Handle& handle, perf_db = UserFindDbRecord::TryLoad(handle, problem, [&](DbRecord& record) { ConvolutionUserBuffers bufs(workSpace, workSpaceSize); bufs.SetWrW(x, dw, dy); - ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage = - findMode.IsFastHybrid(ctx); ctx.use_dynamic_solutions_only = findMode.IsDynamicHybrid(ctx); ctx.do_search = exhaustiveSearch; ctx.SetStream(&handle); diff --git a/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp b/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp index 6834c520c9..d059481fd3 100644 --- a/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp +++ b/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp @@ -638,8 +638,6 @@ bool ConvHipImplicitGemmBwdDataV1R1::IsApplicable(const ConvolutionContext& ctx) { if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1{})) return false; - if(ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage) - return false; if(!ctx.use_hip_kernels) return false; if(!ctx.IsLayoutDefault()) diff --git a/src/solver/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp index 4840d3f36b..6b10310404 100644 --- a/src/solver/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp @@ -748,9 +748,6 @@ bool ConvHipImplicitGemmBwdDataV1R1Xdlops::IsApplicable(const ConvolutionContext if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1_XDLOPS{})) return false; - if(ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage) - return false; - if(!IsComposableKernelSupportedHardware(ctx)) return false; diff --git a/src/solver/conv_hip_implicit_gemm_bwd_v4r1.cpp b/src/solver/conv_hip_implicit_gemm_bwd_v4r1.cpp index ba09272a26..297a8cd5c0 100644 --- a/src/solver/conv_hip_implicit_gemm_bwd_v4r1.cpp +++ b/src/solver/conv_hip_implicit_gemm_bwd_v4r1.cpp @@ -745,9 +745,6 @@ bool ConvHipImplicitGemmBwdDataV4R1::IsApplicable(const ConvolutionContext& ctx) if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1{})) return false; - if(ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage) - return false; - if(!IsComposableKernelSupportedHardware(ctx)) return false; diff --git a/src/solver/conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp index 2ea0c60a5e..00b64bf7a3 100644 --- a/src/solver/conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp @@ -824,8 +824,6 @@ bool ConvHipImplicitGemmBwdDataV4R1Xdlops::IsApplicable(const ConvolutionContext if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS{})) return false; #endif - if(ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage) - return false; if(!IsComposableKernelSupportedHardware(ctx)) return false; if(!ctx.direction.IsBackwardData()) diff --git a/src/solver/conv_hip_implicit_gemm_fwd_v4r1.cpp b/src/solver/conv_hip_implicit_gemm_fwd_v4r1.cpp index aeced97a10..e9c68a68e9 100644 --- a/src/solver/conv_hip_implicit_gemm_fwd_v4r1.cpp +++ b/src/solver/conv_hip_implicit_gemm_fwd_v4r1.cpp @@ -44,8 +44,6 @@ bool ConvHipImplicitGemmV4R1Fwd::IsApplicable(const ConvolutionContext& ctx) con { if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1{})) return false; - if(ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage) - return false; if(!IsComposableKernelSupportedHardware(ctx)) return false; if(!ctx.direction.IsForward()) @@ -82,8 +80,6 @@ bool ConvHipImplicitGemmV4R1WrW::IsApplicable(const ConvolutionContext& ctx) con { if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R1{})) return false; - if(ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage) - return false; if(!IsComposableKernelSupportedHardware(ctx)) return false; if(!ctx.direction.IsBackwardWrW()) diff --git a/src/solver/conv_hip_implicit_gemm_fwd_v4r4.cpp b/src/solver/conv_hip_implicit_gemm_fwd_v4r4.cpp index 52b7c598e8..44a29a86f9 100644 --- a/src/solver/conv_hip_implicit_gemm_fwd_v4r4.cpp +++ b/src/solver/conv_hip_implicit_gemm_fwd_v4r4.cpp @@ -586,8 +586,6 @@ bool ConvHipImplicitGemmV4R4Fwd::IsApplicable(const ConvolutionContext& ctx) con { if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4{})) return false; - if(ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage) - return false; if(!ctx.use_hip_kernels) return false; if(!ctx.IsLayoutDefault()) diff --git a/src/solver/conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp index 963bcd8138..5dd367a22b 100644 --- a/src/solver/conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp @@ -967,9 +967,6 @@ bool ConvHipImplicitGemmForwardV4R4Xdlops::IsApplicable(const ConvolutionContext if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4_XDLOPS{})) return false; - if(ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage) - return false; - if(!ctx.use_hip_kernels) return false; diff --git a/src/solver/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp b/src/solver/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp index 93e4c080db..916410f05e 100644 --- a/src/solver/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp +++ b/src/solver/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp @@ -1043,9 +1043,6 @@ bool ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm::IsApplicable( if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4_PADDED_GEMM_XDLOPS{})) return false; - if(ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage) - return false; - if(!ctx.use_hip_kernels) return false; diff --git a/src/solver/conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp index 928729edea..e81a5ff293 100644 --- a/src/solver/conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp @@ -997,9 +997,6 @@ bool ConvHipImplicitGemmForwardV4R5Xdlops::IsApplicable(const ConvolutionContext if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R5_XDLOPS{})) return false; - if(ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage) - return false; - if(!ctx.use_hip_kernels) return false; diff --git a/src/solver/conv_hip_implicit_gemm_wrw_v4r4.cpp b/src/solver/conv_hip_implicit_gemm_wrw_v4r4.cpp index 9ee4e1d516..36be4fd96a 100644 --- a/src/solver/conv_hip_implicit_gemm_wrw_v4r4.cpp +++ b/src/solver/conv_hip_implicit_gemm_wrw_v4r4.cpp @@ -590,8 +590,6 @@ bool ConvHipImplicitGemmV4R4WrW::IsApplicable(const ConvolutionContext& ctx) con { if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4{})) return false; - if(ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage) - return false; if(!ctx.use_hip_kernels) return false; if(!ctx.IsLayoutDefault()) diff --git a/src/solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp index 36ba5aa0be..24a6d7a687 100644 --- a/src/solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp @@ -1036,9 +1036,6 @@ bool ConvHipImplicitGemmWrwV4R4Xdlops::IsApplicable(const ConvolutionContext& ct if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4_XDLOPS{})) return false; - if(ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage) - return false; - if(!ctx.use_hip_kernels) return false; diff --git a/src/solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp b/src/solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp index 6136014412..edddc01fa9 100644 --- a/src/solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp +++ b/src/solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp @@ -1108,9 +1108,6 @@ bool ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm::IsApplicable(const Convolutio if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4_PADDED_GEMM_XDLOPS{})) return false; - if(ctx.skip_solutions_that_take_long_time_to_build_and_have_narrow_coverage) - return false; - if(!IsComposableKernelSupportedHardware(ctx)) return false;