Workaround for issue #2492 part 2 (improvement) (#2510)

* workaround_issue_2492(01) Disable ConvBinWinoRxS when granularity loss is > 0.995 (performance drops 200 times)

* workaround_issue_2492(02) Allow disabing the W/A by setting MIOPEN_DEBUG_WORKAROUND_ISSUE_2493=0 in the env.

* workaround_issue_2492(03) [debug] Disable MIOPEN_DEBUG_WORKAROUND_ISSUE_2493 during driver warm-up.

* workaround_issue_2492(04) [quality] Make the compuation of max granularity loss more clear.

* workaround_issue_2492_01(02) [debug] Log granularity loss when ConvBinWinogradRxSf2x3* solver is skipped.

* workaround_issue_2492_01(03) [tests] test_db_sync: Disable WORKAROUND_ISSUE_2493 via environment. Support reading legacy fdb (WORKAROUND_ISSUE_1987). Allow FDB testing on gfx1030 (SKIP_KDB_PDB_TESTING). Add W/A for ConvOclDirectFwdFused on gfx1030. Print number of failures per testing thread.

* workaround_issue_2492_01(04) Remove leftovers from gfx1030 testing

* workaround_issue_2492_01(05) More gfx1030 leftovers removed

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

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())

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

src/solver/conv_winoRxS.cpp

@@ -682,10 +682,15 @@ 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)
+        constexpr double max_perf_drop_due_to_granularity = 200; // Times.
+        const auto gl = ShaderModel(ctx, problem, Winodata, Winofilter).GetGranularityLoss();
+        if(gl > (1.0 - 1.0 / max_perf_drop_due_to_granularity))
+        {
+            MIOPEN_LOG_I("granularity_loss =" << gl);
             return false;
+        }
     }
 #endif
 

test/gtest/db_sync.cpp

@@ -34,14 +34,22 @@
 #include <miopen/find_db.hpp>
 #include <miopen/tensor.hpp>
 #include <miopen/conv/problem_description.hpp>
+#include <miopen/conv_algo_name.hpp>
 #include <miopen/solver_id.hpp>
 #include <miopen/any_solver.hpp>
 #include <miopen/mt_queue.hpp>
 
+#include <cstdlib>
 #include <regex>
 #include <exception>
 #include <unordered_set>
 
+#define WORKAROUND_ISSUE_2493 1
+
+#define WORKAROUND_ISSUE_1987 0      // Allows testing FDB on gfx1030 (legacy fdb).
+#define SKIP_KDB_PDB_TESTING 0       // Allows testing FDB on gfx1030.
+#define SKIP_CONVOCLDIRECTFWDFUSED 0 // Allows testing FDB on gfx1030 (legacy fdb).
+
 struct KDBKey
 {
     std::string program_file;
@@ -64,6 +72,50 @@ struct std::hash<KDBKey>
     }
 };
 
+#if WORKAROUND_ISSUE_2493
+static void SetEnvironmentVariable(const std::string& name, const std::string& value)
+{
+#ifdef _WIN32
+    const auto ret = _putenv_s(env_var.c_str(), value.c_str());
+#else
+    const auto ret = setenv(name.c_str(), value.c_str(), 1);
+#endif
+    ASSERT_TRUE(ret == 0);
+}
+#endif // WORKAROUND_ISSUE_2493
+
+#if WORKAROUND_ISSUE_1987
+/// \todo Copied from src/db_record.cpp
+/// Transform find-db (v.1.0) ID:VALUES to the current format.
+/// Implementation is intentionally straightforward.
+/// Do not include the 1st value from VALUES (solver name) into transformed VALUES.
+/// Ignore FdbKCache_Key pair (last two values).
+/// Append id (algorithm) to VALUES.
+/// Use solver name as ID.
+static bool TransformFindDbItem10to20(std::string& id, std::string& values)
+{
+    MIOPEN_LOG_T("Legacy find-db item: " << id << ':' << values);
+    std::size_t pos = values.find(',');
+    if(pos == std::string::npos)
+        return false;
+    const auto solver = values.substr(0, pos);
+
+    const auto time_workspace_pos = pos + 1;
+    pos                           = values.find(',', time_workspace_pos);
+    if(pos == std::string::npos)
+        return false;
+    pos = values.find(',', pos + 1);
+    if(pos == std::string::npos)
+        return false;
+    const auto time_workspace = values.substr(time_workspace_pos, pos - time_workspace_pos);
+
+    values = time_workspace + ',' + id;
+    id     = solver;
+    MIOPEN_LOG_T("Transformed find-db item: " << id << ':' << values);
+    return true;
+}
+#endif
+
 namespace miopen {
 conv::Direction GetDirectionFromString(const std::string& direction)
 {
@@ -254,8 +306,19 @@ void ParseFDBbVal(const std::string& val, std::vector<FDBVal>& fdb_vals)
     {
         const auto id_size = id_val.find(':');
         ASSERT_TRUE(id_size != std::string::npos) << "Ill formed value: " << id_val;
-        auto id        = id_val.substr(0, id_size);
-        auto values    = id_val.substr(id_size + 1);
+        auto id     = id_val.substr(0, id_size);
+        auto values = id_val.substr(id_size + 1);
+#if WORKAROUND_ISSUE_1987
+        /// \todo Copied from src/db_record.cpp
+        /// Detect legacy find-db item (v.1.0 ID:VALUES) and transform it to the current format.
+        /// For now, *only* legacy find-db record use convolution algorithm as ID, so if ID is
+        /// a valid algorithm, then we can safely assume that the item is in legacy format.
+        if(IsValidConvolutionDirAlgo(id))
+        {
+            ASSERT_TRUE(TransformFindDbItem10to20(id, values))
+                << "Ill-formed legacy find-db item: " << values;
+        }
+#endif
         const auto tmp = FDBVal{id, values};
         fdb_vals.emplace_back(tmp);
     }
@@ -369,17 +432,20 @@ void SetupPaths(boost::filesystem::path& fdb_file_path,
         << "Db file does not exist" << fdb_file_path;
     ASSERT_TRUE(boost::filesystem::exists(pdb_file_path))
         << "Db file does not exist" << pdb_file_path;
-    ASSERT_TRUE(boost::filesystem::exists(kdb_file_path))
+    ASSERT_TRUE(SKIP_KDB_PDB_TESTING || boost::filesystem::exists(kdb_file_path))
         << "Db file does not exist" << kdb_file_path;
 }
 
 TEST(DBSync, KDBTargetID)
 {
     boost::filesystem::path fdb_file_path, pdb_file_path, kdb_file_path;
+#if WORKAROUND_ISSUE_2493
+    SetEnvironmentVariable("MIOPEN_DEBUG_WORKAROUND_ISSUE_2493", "0");
+#endif
     SetupPaths(fdb_file_path, pdb_file_path, kdb_file_path, get_handle());
     std::ignore = fdb_file_path;
     std::ignore = pdb_file_path;
-    EXPECT_FALSE(miopen::CheckKDBForTargetID(kdb_file_path));
+    EXPECT_FALSE(!SKIP_KDB_PDB_TESTING && miopen::CheckKDBForTargetID(kdb_file_path));
 }
 
 bool LogBuildMessage()
@@ -539,6 +605,7 @@ void CheckFDBEntry(size_t thread_index,
     SetupPaths(fdb_file_path, pdb_file_path, kdb_file_path, _ctx.GetStream());
     std::unordered_set<KDBKey> checked_kdbs;
     const auto data_size = data.size();
+    auto failures        = 0;
     for(auto kidx = thread_index; kidx < data_size; kidx += total_threads)
     {
         const auto& kinder = data.at(kidx);
@@ -550,6 +617,7 @@ void CheckFDBEntry(size_t thread_index,
         problem.Serialize(ss);
         // moment of truth
         EXPECT_TRUE(ss.str() == kinder.first)
+            << '[' << (++failures) << "] " //
             << "Failed to parse FDB key:" << kidx << ":Parsed Key: " << ss.str();
 
         std::vector<miopen::FDBVal> fdb_vals;
@@ -562,6 +630,23 @@ void CheckFDBEntry(size_t thread_index,
         for(const auto& val : fdb_vals)
         {
             miopen::solver::Id id{val.solver_id};
+            EXPECT_TRUE(id.IsValid())
+                << '[' << (++failures) << "] " //
+                << "Solver " << id.Value() << "/" << id.ToString() << ", val.solver_id "
+                << val.solver_id << ", val.vals " << val.vals;
+
+#if SKIP_CONVOCLDIRECTFWDFUSED
+            /// \todo Workaround: solv.IsApplicable() asserts with ConvOclDirectFwdFused
+            /// on gfx1030. AnySolver instance is empty (nullptr) due to some unknown reason.
+            if(val.solver_id == "ConvOclDirectFwdFused")
+            {
+                MIOPEN_LOG_I("Skipping: val.solver_id " << val.solver_id << ", val.vals "
+                                                        << val.vals);
+                ++fdb_idx;
+                continue;
+            }
+#endif
+
             const auto solv = id.GetSolver();
             // Skip MLIR
             if(miopen::StartsWith(id.ToString(), "ConvMlir"))
@@ -570,24 +655,27 @@ void CheckFDBEntry(size_t thread_index,
                 ++fdb_idx;
                 continue;
             }
-            EXPECT_TRUE(solv.IsApplicable(ctx, problem))
+            EXPECT_TRUE(solv.IsApplicable(ctx, problem)) //
+                << '[' << (++failures) << "] "           //
                 << "Solver is not applicable fdb-key:" << kinder.first
                 << " Solver: " << id.ToString();
             miopen::solver::ConvSolution sol;
             if(solv.IsTunable())
             {
                 const auto pdb_entry_exists = pdb_vals.find(val.solver_id) != pdb_vals.end();
                 // TODO: Print the SQL query
-                EXPECT_TRUE(pdb_entry_exists)
+                EXPECT_TRUE(SKIP_KDB_PDB_TESTING || pdb_entry_exists)
+                    << '[' << (++failures) << "] " //
                     << "PDB entry does not exist for tunable fdb-key:" << kinder.first << ": solver"
                     << val.solver_id << " pdb-select-query: " << pdb_select_query;
                 auto db               = miopen::GetDb(ctx);
                 std::string pdb_entry = "";
-                if(pdb_entry_exists)
+                if(!SKIP_KDB_PDB_TESTING && pdb_entry_exists)
                 {
                     pdb_entry = pdb_vals.at(val.solver_id);
                     bool res  = solv.TestPerfCfgParams(ctx, problem, pdb_vals.at(val.solver_id));
                     EXPECT_TRUE(res)
+                        << '[' << (++failures) << "] " //
                         << "Invalid perf config found fdb-key:" << kinder.first
                         << " Solver: " << solv.GetSolverDbId() << ":" << pdb_vals.at(val.solver_id)
                         << " pdb-select-query: " << pdb_select_query;
@@ -603,9 +691,10 @@ void CheckFDBEntry(size_t thread_index,
                 }
                 // TODO Generate the Select query for pdb
                 EXPECT_TRUE(sol.Succeeded())
+                    << '[' << (++failures) << "] " //
                     << "Invalid solution fdb-key:" << kinder.first << " Solver: " << id.ToString()
                     << " pdb-val:" << pdb_entry;
-                if(fdb_idx == 0)
+                if(!SKIP_KDB_PDB_TESTING && fdb_idx == 0)
                 {
                     for(const auto& kern : sol.construction_params)
                     {
@@ -630,6 +719,7 @@ void CheckFDBEntry(size_t thread_index,
                             found = checked_kdbs.count(KDBKey{program_file, compile_options}) > 0;
                         if(!found)
                             EXPECT_TRUE(found)
+                                << '[' << (++failures) << "] " //
                                 << "KDB entry not found for  fdb-key:" << kinder.first
                                 << " Solver: " << id.ToString() << " pdb-val:" << pdb_entry
                                 << " filename: " << program_file << " compile args: "
@@ -642,6 +732,7 @@ void CheckFDBEntry(size_t thread_index,
             }
             else
                 EXPECT_TRUE(pdb_vals.find(val.solver_id) == pdb_vals.end())
+                    << '[' << (++failures) << "] " //
                     << "Non-Tunable solver found in PDB" << solv.GetSolverDbId();
             ++fdb_idx;
         }
@@ -688,8 +779,10 @@ void StaticFDBSync(const std::string& arch, const size_t num_cu)
     SetupPaths(fdb_file_path, pdb_file_path, kdb_file_path, handle);
     std::cout << "Handle CU count: " << handle.GetMaxComputeUnits()
               << " Parameter Value: " << num_cu << std::endl;
+#if !SKIP_KDB_PDB_TESTING
     // Warmup the kdb cache
     miopen::CheckKDBObjects(kdb_file_path, "", "");
+#endif
     const auto& find_db = miopen::ReadonlyRamDb::GetCached(fdb_file_path.string(), true);
     // assert that find_db.cache is not empty, since that indicates the file was not readable
     ASSERT_TRUE(!find_db.GetCacheMap().empty()) << "Find DB does not have any entries";
@@ -736,5 +829,6 @@ TEST_P(DBSync, StaticFDBSync)
 INSTANTIATE_TEST_SUITE_P(DBSyncSuite,
                          DBSync,
                          testing::Values(std::make_pair("gfx90a", 104),
+                                         std::make_pair("gfx1030", 36),
                                          std::make_pair("gfx90a", 110),
                                          std::make_pair("gfx908", 120)));