[tests] Added smoke tests for conv solvers (ConvAsmBwdWrW1x1 and more). Some fixes.

CMakeLists.txt

@@ -29,7 +29,7 @@ macro(set_var_to_condition var)
     if(${ARGN})
         set(${var} TRUE)
     else()
-        set(${name} FALSE)
+        set(${var} FALSE)
     endif()
 endmacro()
 

src/CMakeLists.txt

@@ -112,6 +112,7 @@ set( MIOpen_Source
     find_db.cpp
     fused_api.cpp
     fusion.cpp
+    generic_search.cpp
     handle_api.cpp
     invoker_cache.cpp
     kernel_build_params.cpp

src/generic_search.cpp

@@ -0,0 +1,44 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2022 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+
+#include <miopen/env.hpp>
+#include <miopen/generic_search.hpp>
+
+#include <cstddef>
+#include <limits>
+
+namespace miopen {
+namespace solver {
+
+MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_TUNING_ITERATIONS_MAX)
+
+std::size_t GetTuningIterationsMax()
+{
+    return Value(MIOPEN_DEBUG_TUNING_ITERATIONS_MAX{}, std::numeric_limits<std::size_t>::max());
+}
+
+} // namespace solver
+} // namespace miopen

src/include/miopen/generic_search.hpp

@@ -38,6 +38,7 @@
 #include <miopen/timer.hpp>
 #include <miopen/type_traits.hpp>
 
+#include <algorithm>
 #include <vector>
 #include <cstdlib>
 #include <limits>
@@ -211,30 +212,6 @@ class HeartBeat
     }
 };
 
-inline void InitRandomly(std::vector<float>& vec, const double offset, const double factor)
-{
-    float* p = vec.data();
-    for(unsigned long i = 0; i < vec.size(); ++i)
-        *p++ = static_cast<float>(
-            (rand() * (1.0 / RAND_MAX) + offset) * // NOLINT (concurrency-mt-unsafe)
-            factor);
-}
-
-inline void InitRandomly(std::vector<float>& vec)
-{
-    float* p = vec.data();
-    for(unsigned long i = 0; i < vec.size(); ++i)
-        *p++ = static_cast<float>(rand() * (1.0 / RAND_MAX)); // NOLINT (concurrency-mt-unsafe)
-}
-
-inline size_t divide_round_plus_inf(const size_t x, const unsigned y)
-{
-    assert(y > 0);
-    if(x % y != 0)
-        return x / y + 1;
-    return x / y;
-}
-
 /// Solver member function requirements:
 /// * GetDefaultPerformanceConfig shall be implemented.
 ///   - Its return type shall be suitable for instantiation of the ComputedContainer.
@@ -320,6 +297,8 @@ auto GenericSearch(const Solver s,
     return GenericSearch(s, ctx, invoke_ctx);
 }
 
+std::size_t GetTuningIterationsMax();
+
 template <class Solver, class Context>
 auto GenericSearch(const Solver s, const Context& context_, const AnyInvokeParams& invoke_ctx_)
     -> decltype(s.GetDefaultPerformanceConfig(context_))
@@ -344,8 +323,10 @@ auto GenericSearch(const Solver s, const Context& context_, const AnyInvokeParam
     auto& profile_h = context.GetStream();
     AutoEnableProfiling enableProfiling{profile_h};
 
-    auto all_configs       = GetAllConfigs(s, context);
-    const int n_runs_total = std::distance(all_configs.begin(), all_configs.end());
+    auto all_configs = GetAllConfigs(s, context);
+    const std::size_t n_runs_total =
+        std::min(static_cast<std::size_t>(std::distance(all_configs.begin(), all_configs.end())),
+                 GetTuningIterationsMax());
 
     bool is_passed  = false; // left false only if all iterations failed.
     float best_time = std::numeric_limits<float>::max();
@@ -357,15 +338,19 @@ auto GenericSearch(const Solver s, const Context& context_, const AnyInvokeParam
     if(!miopen::IsCacheDisabled()) // Otherwise precompilation is useless.
     {
         std::vector<KernelInfo> kernels;
+        size_t n_current = 0;
         for(const auto& current_config : all_configs)
         {
+            if(n_current >= n_runs_total)
+                break;
             ConvSolution current_solution = s.GetSolution(context, current_config);
             for(auto&& kernel : current_solution.construction_params)
             {
                 if(profile_h.HasProgram(kernel.kernel_file, kernel.comp_options))
                     continue;
                 kernels.push_back(kernel);
             }
+            ++n_current;
         }
         std::ignore = PrecompileKernels(profile_h, kernels);
     }
@@ -375,6 +360,9 @@ auto GenericSearch(const Solver s, const Context& context_, const AnyInvokeParam
         size_t n_current = 0;
         for(const auto& current_config : all_configs)
         {
+            if(n_current >= n_runs_total)
+                break;
+
             float elapsed_time = 0.0f;
             int ret            = 0;
             MIOPEN_LOG_I2('#' << n_current << '/' << n_failed << '/' << n_runs_total << ' '

src/solver/conv_asm_1x1u_stride2.cpp

@@ -99,6 +99,14 @@ inline static bool IsLinear(const int v)
     return L <= v && v <= H;
 }
 
+static inline size_t divide_round_plus_inf(const size_t x, const unsigned y)
+{
+    assert(x >= 0 && y > 0);
+    if(x % y != 0)
+        return x / y + 1;
+    return x / y;
+}
+
 enum class MemLayout : int
 {
     NCHW = 0,
@@ -497,9 +505,6 @@ bool ConvAsm1x1UV2::IsApplicable(const ConvolutionContext& ctx,
         return false;
     }
 
-    if(name == "gfx90a" && problem.conv_problem.IsGfx90aFp16altRequired())
-        return false;
-
     const auto elements_in_dword = 4 / GetTypeSize(problem.in_data_type);
     // clang-format off
     const auto img_hw = problem.out_height * problem.out_width;

src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp

@@ -1165,7 +1165,7 @@ ConvSolution ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::GetSolution(
 
                 if(workSpace == nullptr || workSpaceSize < required_workspace_size)
                     MIOPEN_THROW("Not enough workspace has been provided for "
-                                 "ConvAsmImplicitGemmGTCDynamicWrwXdlops with fp16 and atomic "
+                                 "ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC with fp16 and atomic "
                                  "add.");
                 auto trans_input_buf =
                     trans_input_size == 0

src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp

@@ -644,19 +644,16 @@ bool ConvHipImplicitGemmBwdDataV1R1::IsApplicable(const ConvolutionContext& ctx,
     if(!problem.Is2d() && !(problem.Is3d() && problem.IsFp32()))
         return false;
 
-    // TBD Renable fp16 once the root cause of
-    // fp16 failures, observed in CI, is resolved.
     if(!(problem.IsFp32() || problem.IsBfp16()))
         return false;
     if(problem.group_counts != 1)
         return false;
+
 #if WORKAROUND_ISSUE_309
     if(problem.IsBfp16())
-        return false;
+        if(!miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1{}))
+            return false;
 #endif
-    if(ctx.GetStream().GetDeviceName() == "gfx90a" &&
-       problem.conv_problem.IsGfx90aFp16altRequired())
-        return false;
 
     const auto k = ProblemInterpreter::GetOutputChannelK(problem);
     if(k % GetEPackLength(ctx, problem, false) != 0)