[MLIR] Implement tuning - step 4: Use special value to represent defa…

…ult perf_config (#1159) * [MLIR] Add heuristic init request to represent the default perf_config * Refactor bwd with default perf_config * Refactor fwd xdlops with default perf_config * Refactor bwd xdlops with default perf_config * Address review feedbacks
ROCm · Sep 19, 2021 · 502df2b · 502df2b
1 parent 068b6b1
commit 502df2b
Show file tree

Hide file tree

Showing 8 changed files with 65 additions and 72 deletions.
diff --git a/Dockerfile b/Dockerfile
@@ -114,7 +114,7 @@ RUN if [ "$USE_TARGETID" = "OFF" ] ; then echo "MIOpenTensile is not installed."
 
 RUN if [ "$USE_MLIR" = "ON" ]; \
     then cd ~ && \
-    export MLIR_COMMIT=199d667b9d8caaf283436aaa8a48fd20e074f42c && \
+    export MLIR_COMMIT=31579e0c5cf6eb4d7b1db0d349407f8bab547d9b && \
     wget https://github.com/ROCmSoftwarePlatform/llvm-project-mlir/archive/$MLIR_COMMIT.tar.gz && \
     tar -xvzf $MLIR_COMMIT.tar.gz && \
     rm -rf $MLIR_COMMIT.tar.gz && \
@@ -123,4 +123,4 @@ RUN if [ "$USE_MLIR" = "ON" ]; \
     make -j$(nproc) libMLIRMIOpen && \
     $PREFIX/bin/cmake --install . --component libMLIRMIOpen --prefix /opt/rocm && \
     cd ~ && rm -rf llvm-project-mlir-$MLIR_COMMIT; fi
-
+
diff --git a/src/include/miopen/solver.hpp b/src/include/miopen/solver.hpp
@@ -807,6 +807,9 @@ struct PerformanceConvMlirIgemm : Serializable<PerformanceConvMlirIgemm>
     int GemmNPerThread;
     bool use_spare_set;
 
+    /// \ref https://github.com/ROCmSoftwarePlatform/MIOpen/issues/1154
+    static const PerformanceConvMlirIgemm& MlirHeuristicInitRequest();
+
     PerformanceConvMlirIgemm(int, int, int, int, int, int, bool);
 
     PerformanceConvMlirIgemm(int a, int b, int c, int d, int e, int f)
@@ -854,17 +857,19 @@ struct PerformanceConvMlirIgemmXdlops : Serializable<PerformanceConvMlirIgemmXdl
     int GemmMPerBlock; // 2^n[32..128]
     int GemmNPerBlock; // 2^n[8..16]
     int GemmKPerBlock; // 2^n[4..16]
-
     int GemmMPerWave;
     int GemmNPerWave;
-
     int GemmKPACKSize; // 2^[1..4]
 
     // GemmAThreadCopyMoreGemmK is currently a fix value, is untunable
     bool GemmAThreadCopyMoreGemmK;
     bool GemmBThreadCopyMoreGemmKPack;
 
     bool use_spare_set;
+
+    /// \ref https://github.com/ROCmSoftwarePlatform/MIOpen/issues/1154
+    static const PerformanceConvMlirIgemmXdlops& MlirHeuristicInitRequest();
+
     PerformanceConvMlirIgemmXdlops(int, int, int, int, int, int, bool, bool, bool);
 
     PerformanceConvMlirIgemmXdlops();

diff --git a/src/include/miopen/solver/mlir_common.hpp b/src/include/miopen/solver/mlir_common.hpp
@@ -41,10 +41,22 @@ std::string ConstructBuildOptions(const ConvolutionContext& ctx,
                                   bool is_xdlops,
                                   int kernel_id = 0);
 
+template <typename T>
 std::string ConstructBuildOptions(const ConvolutionContext& ctx,
-                                  const std::string& config,
+                                  const T& perf_config,
                                   bool is_xdlops,
-                                  int kernel_id = 0);
+                                  int kernel_id = 0)
+{
+    std::ostringstream options{ConstructBuildOptions(ctx, is_xdlops, kernel_id), std::ios::ate};
+
+    // Library does heuristic initialization when no perf_config
+    // is specified
+    if(!(perf_config == T::MlirHeuristicInitRequest()))
+        options << " --perf_config " + perf_config.ToString();
+
+    return options.str();
+}
+
 } // namespace mlir
 } // namespace solver
 } // namespace miopen

diff --git a/src/solver/conv_mlir_igemm_bwd.cpp b/src/solver/conv_mlir_igemm_bwd.cpp
@@ -57,7 +57,7 @@ bool ConvMlirIgemmBwd::IsApplicable(const ConvolutionContext& ctx) const
 PerformanceConvMlirIgemm ConvMlirIgemmBwd::GetPerformanceConfig(const ConvolutionContext& ctx) const
 {
     std::ignore = ctx;
-    return {};
+    return PerformanceConvMlirIgemm::MlirHeuristicInitRequest();
 }
 
 bool ConvMlirIgemmBwd::IsValidPerformanceConfig(const ConvolutionContext& ctx,
@@ -87,16 +87,8 @@ ConvSolution ConvMlirIgemmBwd::GetSolution(const ConvolutionContext& ctx,
 
         construction_parameters.kernel_name  = mlir::GetKernelName(ctx, false, kernel_id);
         construction_parameters.kernel_file  = construction_parameters.kernel_name + ".mlir";
-
-        if(config == PerformanceConvMlirIgemm())
-            // At this case, do not pass in the invalid perf config and instead make Miir library to
-            // do heuristic initialization
-            construction_parameters.comp_options =
-                mlir::ConstructBuildOptions(ctx, false, kernel_id);
-        else
-            // At this case, Make Miir library to use the valid perf config
-            construction_parameters.comp_options =
-                mlir::ConstructBuildOptions(ctx, config.ToString(), false, kernel_id);
+        construction_parameters.comp_options =
+            mlir::ConstructBuildOptions(ctx, config, false, kernel_id);
 
         size_t local_size  = 0;
         size_t global_size = 0;

diff --git a/src/solver/conv_mlir_igemm_bwd_xdlops.cpp b/src/solver/conv_mlir_igemm_bwd_xdlops.cpp
@@ -61,7 +61,7 @@ PerformanceConvMlirIgemmXdlops
 ConvMlirIgemmBwdXdlops::GetPerformanceConfig(const ConvolutionContext& ctx) const
 {
     std::ignore = ctx;
-    return {};
+    return PerformanceConvMlirIgemmXdlops::MlirHeuristicInitRequest();
 }
 
 bool ConvMlirIgemmBwdXdlops::IsValidPerformanceConfig(
@@ -92,16 +92,8 @@ ConvSolution ConvMlirIgemmBwdXdlops::GetSolution(const ConvolutionContext& ctx,
 
         construction_parameters.kernel_name  = mlir::GetKernelName(ctx, true, kernel_id);
         construction_parameters.kernel_file  = construction_parameters.kernel_name + ".mlir";
-
-        if(config == PerformanceConvMlirIgemmXdlops())
-            // At this case, do not pass in the invalid perf config and instead make Miir library to
-            // do heuristic initialization
-            construction_parameters.comp_options =
-                mlir::ConstructBuildOptions(ctx, true, kernel_id);
-        else
-            // At this case, Make Miir library to use the valid perf config
-            construction_parameters.comp_options =
-                mlir::ConstructBuildOptions(ctx, config.ToString(), true, kernel_id);
+        construction_parameters.comp_options =
+            mlir::ConstructBuildOptions(ctx, config, true, kernel_id);
 
         size_t local_size  = 0;
         size_t global_size = 0;

diff --git a/src/solver/conv_mlir_igemm_fwd.cpp b/src/solver/conv_mlir_igemm_fwd.cpp
@@ -37,6 +37,13 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_MLIR_IGEMM_FWD)
 namespace miopen {
 namespace solver {
 
+const PerformanceConvMlirIgemm& PerformanceConvMlirIgemm::MlirHeuristicInitRequest()
+{
+    static const PerformanceConvMlirIgemm p =
+        PerformanceConvMlirIgemm(-2, -2, -2, -2, -2, -2, false);
+    return p;
+}
+
 PerformanceConvMlirIgemm::PerformanceConvMlirIgemm(int BlockSize_,
                                                    int GemmMPerBlock_,
                                                    int GemmNPerBlock_,
@@ -85,7 +92,10 @@ bool PerformanceConvMlirIgemm::operator==(const PerformanceConvMlirIgemm& other)
 bool PerformanceConvMlirIgemm::IsValid(const ConvolutionContext& ctx) const
 {
 #if MIOPEN_USE_MLIR
-    return MiirIsConfigApplicable(mlir::ConstructBuildOptions(ctx, ToString(), false));
+    if(*this == MlirHeuristicInitRequest())
+        return true;
+
+    return MiirIsConfigApplicable(mlir::ConstructBuildOptions(ctx, *this, false));
 #else
     std::ignore = ctx;
     return false;
@@ -94,6 +104,9 @@ bool PerformanceConvMlirIgemm::IsValid(const ConvolutionContext& ctx) const
 
 bool PerformanceConvMlirIgemm::SetNextValue(const ConvolutionContext& /*config*/)
 {
+    if(*this == MlirHeuristicInitRequest())
+        MIOPEN_THROW("Should not iterate from the heuristic value");
+
     // always search full space, no matter if use_spare_set or not
     do
     {
@@ -125,9 +138,8 @@ std::string PerformanceConvMlirIgemm::ToString() const
 
 PerformanceConvMlirIgemm ConvMlirIgemmFwd::GetPerformanceConfig(const ConvolutionContext& ctx) const
 {
-    // Return the invalid config as the default to signal MLIR do heuristic intialization
     std::ignore = ctx;
-    return {};
+    return PerformanceConvMlirIgemm::MlirHeuristicInitRequest();
 }
 
 bool ConvMlirIgemmFwd::IsValidPerformanceConfig(const ConvolutionContext& ctx,
@@ -168,17 +180,9 @@ ConvSolution ConvMlirIgemmFwd::GetSolution(const ConvolutionContext& ctx,
     ConvSolution result;
     KernelInfo construction_parameters;
 
-    construction_parameters.kernel_name = mlir::GetKernelName(ctx, false);
-    construction_parameters.kernel_file = construction_parameters.kernel_name + ".mlir";
-
-    if(config == PerformanceConvMlirIgemm())
-        // At this case, do not pass in the invalid perf config and instead make Miir library to do
-        // heuristic initialization
-        construction_parameters.comp_options = mlir::ConstructBuildOptions(ctx, false);
-    else
-        // At this case, Make Miir library to use the valid perf config
-        construction_parameters.comp_options =
-            mlir::ConstructBuildOptions(ctx, config.ToString(), false);
+    construction_parameters.kernel_name  = mlir::GetKernelName(ctx, false);
+    construction_parameters.kernel_file  = construction_parameters.kernel_name + ".mlir";
+    construction_parameters.comp_options = mlir::ConstructBuildOptions(ctx, config, false);
 
     size_t local_size  = 0;
     size_t global_size = 0;

diff --git a/src/solver/conv_mlir_igemm_fwd_xdlops.cpp b/src/solver/conv_mlir_igemm_fwd_xdlops.cpp
@@ -38,6 +38,13 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_MLIR_IGEMM_FWD_XDLOPS)
 namespace miopen {
 namespace solver {
 
+const PerformanceConvMlirIgemmXdlops& PerformanceConvMlirIgemmXdlops::MlirHeuristicInitRequest()
+{
+    static const PerformanceConvMlirIgemmXdlops p =
+        PerformanceConvMlirIgemmXdlops(-2, -2, -2, -2, -2, -2, false, false, false);
+    return p;
+}
+
 bool ConvMlirIgemmFwdXdlops::IsApplicable(const ConvolutionContext& ctx) const
 {
 #if MIOPEN_USE_MLIR
@@ -107,8 +114,10 @@ bool PerformanceConvMlirIgemmXdlops::operator==(const PerformanceConvMlirIgemmXd
 bool PerformanceConvMlirIgemmXdlops::IsValid(const ConvolutionContext& ctx) const
 {
 #if MIOPEN_USE_MLIR
-    bool isValid = MiirIsConfigApplicable(mlir::ConstructBuildOptions(ctx, ToString(), true));
-    return isValid;
+    if(*this == MlirHeuristicInitRequest())
+        return true;
+
+    return MiirIsConfigApplicable(mlir::ConstructBuildOptions(ctx, *this, true));
 #else
     std::ignore = ctx;
     return false;
@@ -117,6 +126,9 @@ bool PerformanceConvMlirIgemmXdlops::IsValid(const ConvolutionContext& ctx) cons
 
 bool PerformanceConvMlirIgemmXdlops::SetNextValue(const ConvolutionContext& /*config*/)
 {
+    if(*this == MlirHeuristicInitRequest())
+        MIOPEN_THROW("Should not iterate from the heuristic value");
+
     GemmBThreadCopyMoreGemmKPack = true;
     GemmAThreadCopyMoreGemmK     = true;
     do
@@ -151,7 +163,7 @@ PerformanceConvMlirIgemmXdlops
 ConvMlirIgemmFwdXdlops::GetPerformanceConfig(const ConvolutionContext& ctx) const
 {
     std::ignore = ctx;
-    return {};
+    return PerformanceConvMlirIgemmXdlops::MlirHeuristicInitRequest();
 }
 
 bool ConvMlirIgemmFwdXdlops::IsValidPerformanceConfig(
@@ -176,17 +188,9 @@ ConvSolution ConvMlirIgemmFwdXdlops::GetSolution(const ConvolutionContext& ctx,
     ConvSolution result;
     KernelInfo construction_parameters;
 
-    construction_parameters.kernel_name = mlir::GetKernelName(ctx, true);
-    construction_parameters.kernel_file = construction_parameters.kernel_name + ".mlir";
-
-    if(config == PerformanceConvMlirIgemmXdlops())
-        // At this case, do not pass in the invalid perf config and instead make Miir library to do
-        // heuristic initialization
-        construction_parameters.comp_options = mlir::ConstructBuildOptions(ctx, true);
-    else
-        // At this case, Make Miir library to use the valid perf config
-        construction_parameters.comp_options =
-            mlir::ConstructBuildOptions(ctx, config.ToString(), true);
+    construction_parameters.kernel_name  = mlir::GetKernelName(ctx, true);
+    construction_parameters.kernel_file  = construction_parameters.kernel_name + ".mlir";
+    construction_parameters.comp_options = mlir::ConstructBuildOptions(ctx, config, true);
 
     size_t local_size  = 0;
     size_t global_size = 0;

diff --git a/src/solver/mlir_common.cpp b/src/solver/mlir_common.cpp
@@ -173,22 +173,6 @@ std::string ConstructBuildOptions(const ConvolutionContext& ctx,
     return mlir_handle.str();
 }
 
-std::string ConstructBuildOptions(const ConvolutionContext& ctx,
-                                  const std::string& config,
-                                  bool is_xdlops,
-                                  int kernel_id)
-{
-    std::ostringstream mlir_handle;
-
-    // clang-format off
-    mlir_handle
-        << ConstructBuildOptions(ctx, is_xdlops, kernel_id)
-        << " --perf_config " << config;
-    // clang-format on
-
-    return mlir_handle.str();
-}
-
 } // namespace mlir
 } // namespace solver
 } // namespace miopen