From 6530353d8bed4f6dd96f2bda8231feb9c9e122d2 Mon Sep 17 00:00:00 2001
From: atamazov <artem.tamazov@gmail.com>
Date: Tue, 7 Nov 2023 19:46:20 +0300
Subject: [PATCH] workaround_issue_2492(03) [debug] Disable
 MIOPEN_DEBUG_WORKAROUND_ISSUE_2493 during driver warm-up.

---
 driver/conv_driver.hpp                   |  5 +++++
 src/execution_context.cpp                |  8 ++++++++
 src/include/miopen/execution_context.hpp | 10 ++++++++++
 src/solver/conv_winoRxS.cpp              |  2 +-
 4 files changed, 24 insertions(+), 1 deletion(-)

diff --git a/driver/conv_driver.hpp b/driver/conv_driver.hpp
index 5d627ba1a5..90ebf2c437 100644
--- a/driver/conv_driver.hpp
+++ b/driver/conv_driver.hpp
@@ -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>
@@ -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;
@@ -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
diff --git a/src/execution_context.cpp b/src/execution_context.cpp
index fb64a10c92..9593aa9fa1 100644
--- a/src/execution_context.cpp
+++ b/src/execution_context.cpp
@@ -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())
diff --git a/src/include/miopen/execution_context.hpp b/src/include/miopen/execution_context.hpp
index a308338eb0..eac09a3ff6 100644
--- a/src/include/miopen/execution_context.hpp
+++ b/src/include/miopen/execution_context.hpp
@@ -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
diff --git a/src/solver/conv_winoRxS.cpp b/src/solver/conv_winoRxS.cpp
index 439a2e2d14..4f05b5053e 100644
--- a/src/solver/conv_winoRxS.cpp
+++ b/src/solver/conv_winoRxS.cpp
@@ -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;