Skip to content

Commit

Permalink
workaround_issue_2492(03) [debug] Disable MIOPEN_DEBUG_WORKAROUND_ISS…
Browse files Browse the repository at this point in the history
…UE_2493 during driver warm-up.
  • Loading branch information
atamazov committed Nov 7, 2023
1 parent 35806a4 commit 6530353
Show file tree
Hide file tree
Showing 4 changed files with 24 additions and 1 deletion.
5 changes: 5 additions & 0 deletions driver/conv_driver.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -57,6 +57,7 @@
#include <miopen/convolution.hpp>
#include <miopen/solver.hpp>
#include <miopen/find_controls.hpp>
#include <miopen/execution_context.hpp>
#include "random.hpp"
#include <numeric>
#include <sstream>
Expand Down Expand Up @@ -100,8 +101,10 @@ struct AutoMiopenWarmupMode
{
debug_logging_quiet_prev = miopen::debug::LoggingQuiet;
debug_find_enforce_disable_prev = miopen::debug::FindEnforceDisable;
debug_is_warmup_ongoing_prev = miopen::debug::IsWarmupOngoing;
miopen::debug::LoggingQuiet = true;
miopen::debug::FindEnforceDisable = true;
miopen::debug::IsWarmupOngoing = true;
}
AutoMiopenWarmupMode(const AutoMiopenWarmupMode&) = delete;
AutoMiopenWarmupMode(AutoMiopenWarmupMode&&) = delete;
Expand All @@ -111,11 +114,13 @@ struct AutoMiopenWarmupMode
{
miopen::debug::LoggingQuiet = debug_logging_quiet_prev;
miopen::debug::FindEnforceDisable = debug_find_enforce_disable_prev;
miopen::debug::IsWarmupOngoing = debug_is_warmup_ongoing_prev;
}

private:
bool debug_logging_quiet_prev;
bool debug_find_enforce_disable_prev;
bool debug_is_warmup_ongoing_prev;
};

struct AutoPrepareForGpuReference
Expand Down
8 changes: 8 additions & 0 deletions src/execution_context.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,14 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_AMD_ROCM_PRECOMPILED_BINARIES)
MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_AMD_ROCM_METADATA_ENFORCE)
MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_AMD_ROCM_METADATA_PREFER_OLDER)

namespace miopen {
namespace debug {

bool IsWarmupOngoing = false; // NOLINT (cppcoreguidelines-avoid-non-const-global-variables)

} // namespace debug
} // namespace miopen

static std::ostream& operator<<(std::ostream& os, const rocm_meta_version& rmv)
{
switch(rmv.getValue())
Expand Down
10 changes: 10 additions & 0 deletions src/include/miopen/execution_context.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -65,6 +65,16 @@ class rocm_meta_version

namespace miopen {

namespace debug {

/// Inform the library that some warm-up (e.g. the one implemented in the driver)
/// is in progress. The library can use this, for example, to disable some
/// workarounds that would affect warm-up otherwise.
/// WARNING: This switch is not intended for use in multi-threaded applications.
extern bool IsWarmupOngoing; // NOLINT (cppcoreguidelines-avoid-non-const-global-variables)

} // namespace debug

struct ExecutionContext
{
// Solution-specific
Expand Down
2 changes: 1 addition & 1 deletion src/solver/conv_winoRxS.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -682,7 +682,7 @@ static bool IsApplicableBase(const ExecutionContext& ctx, const ProblemDescripti
// clang-format on

#if WORKAROUND_ISSUE_2493
if(!miopen::IsDisabled(MIOPEN_DEBUG_WORKAROUND_ISSUE_2493{}))
if(!miopen::IsDisabled(MIOPEN_DEBUG_WORKAROUND_ISSUE_2493{}) && !miopen::debug::IsWarmupOngoing)
{
if(ShaderModel(ctx, problem, Winodata, Winofilter).GetGranularityLoss() > 0.995)
return false;
Expand Down

0 comments on commit 6530353

Please sign in to comment.