Blas auto-tuning for GEMMs (#1668)

src/targets/gpu/CMakeLists.txt

 # ####################################################################################
 # The MIT License (MIT)
 #
-# Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -245,10 +245,14 @@ else()
 endif()
 
 # Check miopen find mode api
+
 include(CheckLibraryExists)
 get_target_property(MIOPEN_LOCATION MIOpen LOCATION)
+get_target_property(ROCBLAS_LOCATION roc::rocblas LOCATION)
 check_library_exists(MIOpen "miopenHiddenSetConvolutionFindMode" "${MIOPEN_LOCATION}" HAS_FIND_MODE_API)
 check_library_exists(MIOpen "miopenFindSolutions" "${MIOPEN_LOCATION}" HAS_FIND_2_API)
+# Beta API for automated GEMM tuning
+check_library_exists(roc::rocblas "rocblas_gemm_ex_get_solutions" "${ROCBLAS_LOCATION}" HAS_ROCBLAS_TUNING_BETA_FEATURE_API)
 
 set(MIGRAPHX_USE_FIND_2_API "${HAS_FIND_2_API}" CACHE BOOL "")
 
@@ -271,6 +275,13 @@ else()
     message(STATUS "MIOpen does not have find mode api")
 endif()
 
+if(HAS_ROCBLAS_TUNING_BETA_FEATURE_API)
+    target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_USE_ROCBLAS_TUNING_API -DROCBLAS_BETA_FEATURES_API -DROCBLAS_NO_DEPRECATED_WARNINGS)
+    message(STATUS "MIGraphx is using Beta API of rocBLAS")
+else()
+    message(STATUS "rocBLAS does not have User Tuning Beta API")
+endif()
+
 target_link_libraries(migraphx_gpu PUBLIC migraphx MIOpen roc::rocblas)
 target_link_libraries(migraphx_gpu PRIVATE migraphx_device migraphx_kernels)
 if(MIGRAPHX_USE_COMPOSABLEKERNEL)

src/targets/gpu/include/migraphx/gpu/gemm.hpp

 /*
  * The MIT License (MIT)
  *
- * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -40,9 +40,8 @@ inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
 struct context;
-
-void blas_shape(const shape& s);
 shape transpose_batch(const shape& s, unsigned trans_batch);
+void blas_shape(const shape& s);
 
 template <class Op>
 struct rocblas_gemm
@@ -52,6 +51,7 @@ struct rocblas_gemm
     float beta           = 0;
     bool compute_fp32    = false;
     unsigned trans_batch = 0;
+    int32_t solution_idx = 0;
 
     template <class Self, class F>
     static auto reflect(Self& self, F f)
@@ -60,7 +60,8 @@ struct rocblas_gemm
                          pack(f(self.alpha, "alpha"),
                               f(self.beta, "beta"),
                               f(self.compute_fp32, "compute_fp32"),
-                              f(self.trans_batch, "trans_batch")));
+                              f(self.trans_batch, "trans_batch"),
+                              f(self.solution_idx, "solution_idx")));
     }
 
     std::string name() const
@@ -76,6 +77,8 @@ struct rocblas_gemm
     {
         std::vector<shape> in_shapes(inputs);
         in_shapes.pop_back();
+        // When input shapes are A, B, C the GEMM equation is  C  =  α AB+ β C   where α, β are
+        // scalars
         check_shapes{in_shapes, *this}.has(2, 3);
         blas_shape(inputs[0]);
         blas_shape(inputs[1]);
@@ -111,11 +114,12 @@ struct rocblas_gemm
     {
         if(this->name() == "gpu::gemm")
         {
-            gemm(ctx, output_shape, args, alpha, beta, compute_fp32);
+            gemm_compute(ctx, output_shape, args, alpha, beta, compute_fp32, solution_idx);
         }
         else
         {
-            gemm(ctx, output_shape, args, int32_t(alpha), int32_t(beta), compute_fp32);
+            gemm_compute(
+                ctx, output_shape, args, int32_t(alpha), int32_t(beta), compute_fp32, solution_idx);
         }
         return args.back();
     }
@@ -124,6 +128,33 @@ struct rocblas_gemm
     {
         return shapes.size() - 1;
     }
+
+    void finalize(context& ctx, const shape& output_shape, const std::vector<shape>& input_shapes)
+    {
+#ifdef MIGRAPHX_USE_ROCBLAS_TUNING_API
+        if(enabled(MIGRAPHX_ENABLE_GEMM_TUNING{}) or ctx.get_exhaustive_tune_flag())
+        {
+            if(this->name() == "gpu::gemm")
+            {
+                solution_idx = gemm_finalize(
+                    ctx, output_shape, input_shapes, alpha, beta, compute_fp32, solution_idx);
+            }
+            else
+            {
+                solution_idx = gemm_finalize(ctx,
+                                             output_shape,
+                                             input_shapes,
+                                             int32_t(alpha),
+                                             int32_t(beta),
+                                             compute_fp32,
+                                             solution_idx);
+            }
+        }
+#else
+        // suppress compiler warnings
+        (void)ctx, (void)output_shape, (void)input_shapes;
+#endif
+    }
 };
 
 } // namespace gpu

src/targets/gpu/include/migraphx/gpu/gemm_impl.hpp

 /*
  * The MIT License (MIT)
  *
- * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -24,27 +24,65 @@
 #ifndef MIGRAPHX_GUARD_RTGLIB_GEMM_IMPL_HPP
 #define MIGRAPHX_GUARD_RTGLIB_GEMM_IMPL_HPP
 
+#include <iterator>
 #include <migraphx/shape.hpp>
 #include <migraphx/argument.hpp>
 #include <migraphx/gpu/context.hpp>
 
+// Set this environment variable to "true" to perform GEMM tuning even when the
+// --exhaustive-tune option isn't set.  Can be used to skip slow convolution tuning.
+MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_ENABLE_GEMM_TUNING);
+
+using milliseconds = std::chrono::duration<double, std::milli>;
+using microseconds = std::chrono::duration<double, std::micro>;
+
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-void gemm(context& ctx,
+/**
+ * @brief Templated implementations of the compute() and finalize() methods of the Gemm operator.
+ *        For each function there are overloads using either float or int32_t for the arguments
+ * alpha and beta.
+ *
+ * @param ctx .
+ * @param output_shape .
+ * @param args .
+ * @param alpha .
+ * @param beta .
+ * @param compute_fp32 .
+ */
+void gemm_compute(context& ctx,
                   const shape& output_shape,
                   const std::vector<argument>& args,
                   float alpha,
                   float beta,
-          bool compute_fp32);
-void gemm(context& ctx,
+                  bool compute_fp32,
+                  int32_t solution_idx);
+
+void gemm_compute(context& ctx,
                   const shape& output_shape,
                   const std::vector<argument>& args,
                   int32_t alpha,
                   int32_t beta,
+                  bool compute_fp32,
+                  int32_t solution_idx);
+
+int32_t gemm_finalize(context& ctx,
+                      const shape& output_shape,
+                      const std::vector<shape>& input_shapes,
+                      float alpha,
+                      float beta,
                       bool compute_fp32);
 
+int32_t gemm_finalize(context& ctx,
+                      const shape& output_shape,
+                      const std::vector<shape>& input_shapes,
+                      int32_t alpha,
+                      int32_t beta,
+                      bool compute_fp32,
+                      int32_t solution_idx);
+
 } // namespace gpu
 } // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/targets/gpu/include/migraphx/gpu/rocblas.hpp

 /*
  * The MIT License (MIT)
  *
- * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal

test/gpu/gemm_tune.cpp

+/*
+ * The MIT License (MIT)
+ *
+ * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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 <iostream>
+#include <vector>
+#include <migraphx/gpu/gemm.hpp>
+#include <hip/hip_runtime_api.h>
+#include <migraphx/gpu/target.hpp>
+#include <migraphx/verify.hpp>
+#include <test.hpp>
+#include <migraphx/make_op.hpp>
+#include <migraphx/iterator_for.hpp>
+
+// includes needed for run_lowering
+#include <migraphx/gpu/lowering.hpp>
+#include <migraphx/auto_contiguous.hpp>
+#include <migraphx/instruction.hpp>
+#include <migraphx/pass_manager.hpp>
+
+// Abbreviated lowering; we don't need the usual cleanup passes for this test
+void run_lowering(migraphx::program& p, bool offload_copy = false)
+{
+    auto ctx = migraphx::gpu::context{};
+    migraphx::run_passes(
+        *p.get_main_module(),
+        {migraphx::auto_contiguous{}, migraphx::gpu::lowering{&ctx, offload_copy}});
+}
+
+/**
+ * Tests the automatic GEMM tuning feature.  In the finalize() method of the gemm op,
+ * rocBLAS API functions are called to quickly benchmark all the GEMM solutions
+ * available in the currently installed rocBLAS library and choose the index of the fastest.
+ */
+TEST_CASE(gemm_tune_with_rocblas)
+{
+    migraphx::program p;
+    auto* mm = p.get_main_module();
+
+    migraphx::shape sa{migraphx::shape::float_type, {4, 2}};
+    migraphx::shape sb{migraphx::shape::float_type, {2, 3}};
+    auto a = mm->add_parameter("a", sa);
+    auto b = mm->add_parameter("b", sb);
+
+    migraphx::operation dot_op = migraphx::make_op("dot");
+    mm->add_instruction(dot_op, a, b);
+
+    // lowering adds gemm implementation for dot operator
+    run_lowering(p);
+
+    migraphx::target gpu_t = migraphx::gpu::target{};
+    migraphx::compile_options options;
+    options.exhaustive_tune = true;
+    p.compile(gpu_t, options);
+
+    migraphx::value solution_idx(0);
+    for(auto ins : iterator_for(*p.get_main_module()))
+    {
+        if(ins->name() == "gpu::gemm")
+        {
+            auto gemm_op = migraphx::get_operation(ins);
+
+            // tuned solution index is not deterministic, but anything other than 0
+            // (default, invalid, or not available) is good.
+            // gemm_op.to_value().debug_print();
+            solution_idx = gemm_op.to_value()["solution_idx"];
+            break;
+        }
+    }
+#ifdef MIGRAPHX_USE_ROCBLAS_TUNING_API
+    EXPECT(0 != solution_idx.to<std::size_t>());
+#else
+    EXPECT(0 == solution_idx.to<std::size_t>());
+#endif
+}
+
+// GEMM tuning of a strided-batch matrix; invokes rocblas_gemm_strided_batched_ex
+TEST_CASE(gemm_tune_strided)
+{
+    migraphx::program p;
+    auto* mm = p.get_main_module();
+
+    migraphx::shape sa{migraphx::shape::float_type, {4, 2, 2}};
+    migraphx::shape sb{migraphx::shape::float_type, {4, 2, 2}};
+    migraphx::shape s_output{migraphx::shape::float_type, {4, 2, 2}};
+    auto a      = mm->add_parameter("a", sa);
+    auto b      = mm->add_parameter("b", sb);
+    auto output = mm->add_parameter("out", s_output);
+
+    auto gemm_oper = migraphx::make_op("gpu::gemm", {{"beta", 2}});
+    mm->add_instruction(gemm_oper, a, b, output);
+
+    migraphx::target gpu_t = migraphx::gpu::target{};
+    migraphx::compile_options options;
+    options.exhaustive_tune = true;
+    p.compile(gpu_t, options);
+
+    migraphx::value solution_idx(0);
+    for(auto ins : iterator_for(*p.get_main_module()))
+    {
+        if(ins->name() == "gpu::gemm")
+        {
+            auto gemm_op = migraphx::get_operation(ins);
+            auto gemmv   = gemm_op.to_value();
+
+            // tuned solution index is not deterministic, but anything other than 0
+            // (default, invalid, or not available) is good.
+            solution_idx = gemm_op.to_value()["solution_idx"];
+            break;
+        }
+    }
+#ifdef MIGRAPHX_USE_ROCBLAS_TUNING_API
+    EXPECT(0 != solution_idx.to<std::size_t>());
+#else
+    EXPECT(0 == solution_idx.to<std::size_t>());
+#endif
+}
+
+// GEMM tuning of a strided-batch matrix; created by lowering
+TEST_CASE(gemm_tune_strided_lowered)
+{
+    migraphx::program p;
+    auto* mm = p.get_main_module();
+
+    // At time of writing this test, gemm_impl considers a shape is strided if it has
+    // at least three dimensions and the 3rd-to-last is nonzero, invoking
+    // rocblas_gemm_strided_batched_ex. Also, DOT operator requires all dimensions except the last
+    // two to be equal.
+    migraphx::shape sa{migraphx::shape::float_type, {4, 2, 5}};
+    migraphx::shape sb{migraphx::shape::float_type, {4, 5, 3}};
+    auto a = mm->add_parameter("a", sa);
+    auto b = mm->add_parameter("b", sb);
+
+    migraphx::operation dot_op = migraphx::make_op("dot");
+    mm->add_instruction(dot_op, a, b);
+
+    // lowering adds gemm implementation for dot operator
+    run_lowering(p);
+
+    migraphx::target gpu_t = migraphx::gpu::target{};
+    migraphx::compile_options options;
+    options.exhaustive_tune = true;
+    p.compile(gpu_t, options);
+
+    migraphx::value solution_idx(0);
+    for(auto ins : iterator_for(*p.get_main_module()))
+    {
+        if(ins->name() == "gpu::gemm")
+        {
+            auto gemm_op = migraphx::get_operation(ins);
+
+            // tuned solution index is not deterministic, but anything other than 0
+            // (default, invalid, or not available) is good.
+            solution_idx = gemm_op.to_value()["solution_idx"];
+            break;
+        }
+    }
+#ifdef MIGRAPHX_USE_ROCBLAS_TUNING_API
+    EXPECT(0 != solution_idx.to<std::size_t>());
+#else
+    EXPECT(0 == solution_idx.to<std::size_t>());
+#endif
+}
+
+TEST_CASE(gemm_tune_invalid_sol_index)
+{
+    migraphx::program p;
+    auto* mm = p.get_main_module();
+
+    migraphx::shape sa{migraphx::shape::float_type, {4, 2}};
+    migraphx::shape sb{migraphx::shape::float_type, {2, 3}};
+    migraphx::shape s_output{migraphx::shape::float_type, {4, 3}};
+    auto a      = mm->add_parameter("a", sa);
+    auto b      = mm->add_parameter("b", sb);
+    auto output = mm->add_parameter("out", s_output);
+
+    auto gemm_oper = migraphx::make_op("gpu::gemm", {{"solution_idx", 987654321}});
+    mm->add_instruction(gemm_oper, a, b, output);
+
+    migraphx::target gpu_t = migraphx::gpu::target{};
+    migraphx::compile_options options;
+    options.exhaustive_tune = true;
+    p.compile(gpu_t, options);
+
+    migraphx::value solution_idx(0);
+    for(auto ins : iterator_for(*p.get_main_module()))
+    {
+        if(ins->name() == "gpu::gemm")
+        {
+            auto gemm_op = migraphx::get_operation(ins);
+            auto gemmv   = gemm_op.to_value();
+
+            // given invalid starting index, should return default 0
+            solution_idx = gemm_op.to_value()["solution_idx"];
+            break;
+        }
+    }
+#ifdef MIGRAPHX_USE_ROCBLAS_TUNING_API
+    EXPECT(0 == solution_idx.to<std::size_t>());
+#else
+    EXPECT(0 != solution_idx.to<std::size_t>());
+#endif
+}
+
+int main(int argc, const char* argv[]) { test::run(argc, argv); }

test/verify/gemm_2args_mm_8.cpp

+/*
+ * The MIT License (MIT)
+ *
+ * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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 "verify_program.hpp"
+#include <migraphx/program.hpp>
+#include <migraphx/generate.hpp>
+#include <migraphx/make_op.hpp>
+
+struct gemm_2args_mm_8 : verify_program<gemm_2args_mm_8>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        auto* mm = p.get_main_module();
+        migraphx::shape a_shape{migraphx::shape::float_type, {2, 128, 32}, {4096, 1, 128}};
+        migraphx::shape b_shape{migraphx::shape::float_type, {32, 32}};
+        auto a  = mm->add_parameter("a", a_shape);
+        auto b  = mm->add_parameter("b", b_shape);
+        auto bb = mm->add_instruction(
+            migraphx::make_op("multibroadcast", {{"out_lens", {2, 32, 32}}}), b);
+
+        mm->add_instruction(migraphx::make_op("dot"), a, bb);
+
+        return p;
+    }
+};