recreated branch by hand merge with develop branch

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
@@ -215,10 +215,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_BETA_FEATURES_API)
 
 set(MIGRAPHX_USE_FIND_2_API "${HAS_FIND_2_API}" CACHE BOOL "")
 
@@ -236,6 +240,13 @@ else()
     message(STATUS "MIOpen does not have find mode api")
 endif()
 
+if(HAS_ROCBLAS_BETA_FEATURES_API)
+    target_compile_definitions(migraphx_gpu PUBLIC -DROCBLAS_BETA_FEATURES_API -DROCBLAS_NO_DEPRECATED_WARNINGS)
+    message(STATUS "MIGraphx is using Beta API of rocBLAS")
+else()
+    message(STATUS "rocBLAS does not have Beta API")
+endif()
+
 target_link_libraries(migraphx_gpu PUBLIC migraphx MIOpen roc::rocblas)
 target_link_libraries(migraphx_gpu PRIVATE migraphx_device migraphx_kernels)
 

src/targets/gpu/gemm_impl.cpp

 /*
  * 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
@@ -21,10 +21,24 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
  * THE SOFTWARE.
  */
+
+/**
+ * Contains a templated struct implementation that wraps several rocBLAS API calls
+ * used by the GEMM operator.  These are accessed by methods declared in gemm_impl.hpp
+ *
+ */
+
 #include <rocblas/rocblas.h>
 #include <migraphx/gpu/gemm_impl.hpp>
-#include <migraphx/reduce_dims.hpp>
-#include <migraphx/permutation.hpp>
+#include <migraphx/time.hpp>
+
+using microseconds = std::chrono::duration<double, std::micro>;
+
+#if ROCBLAS_VERSION_MAJOR > 2 || (ROCBLAS_VERSION_MAJOR == 2 && ROCBLAS_VERSION_MINOR >= 38)
+using flag_type = rocblas_gemm_flags;
+#else
+using flag_type = int;
+#endif
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
@@ -81,198 +95,515 @@ shape transpose_batch(const shape& s, unsigned trans_batch)
     return shape::from_permutation(s.type(), s.lens(), perm);
 }
 
-template <class R, class... Ts, class... Us>
-R rocblas_invoke(R (*f)(Ts...), Us... xs)
+/**
+ * Returns results of rocblas_status_success, rocblas_status_perf_degraded,
+ * or rocblas_status_invalid_value.  Caller
+ * is expected to check for invalid index.  Any other result causes an exception.
+ *
+ */
+template <class F, class Pack, class... Ts>
+auto rocblas_invoke(F f, Pack p, Ts... xs)
 {
-    if constexpr(sizeof...(Ts) == sizeof...(Us))
-        return f(xs...);
+    return p([=](auto... ws) {
+        auto status = f(ws..., xs...);
+        if(status != rocblas_status_success and status != rocblas_status_invalid_value)
+        {
+            if(status == rocblas_status_perf_degraded)
+            {
+                std::cerr << "WARNING: degraded perf. in rocBLAS call" << std::endl;
+            }
             else
-        return f(xs..., nullptr, nullptr);
+                MIGRAPHX_THROW("rocblas_invoke: rocBLAS call failed with status " +
+                               std::to_string(status));
+        }
+        return status;
+    });
 }
 
-static bool is_transposed(const shape& s)
-{
-    if(not s.transposed())
-        return false;
-    return s.strides().back() != 1;
-}
+static bool is_transposed(const shape& s) { return s.transposed() and s.strides().back() != 1; }
 
 static rocblas_int get_batch_stride(const argument& a)
 {
+    // This value is not needed for non-strided inputs
+    if(a.get_shape().strides().size() < 3)
+        return 0;
+    else
         return a.get_shape().strides()[a.get_shape().strides().size() - 3];
 }
 
-template <class T>
-void gemm_impl(context& ctx,
-               const shape& output_shape,
-               const std::vector<argument>& args,
-               T alpha,
-               T beta,
-               bool int8_x4_format,
-               bool compute_fp32)
+/**
+ * Wrapper for multiple rocBLAS calls.  The constructor creates parameters for
+ * these calls based on data shapes and other values contained in the associated
+ * instruction and operation.
+ *
+ * The template parameter T is not the type of the input data but of the weighting
+ * coefficients alpha and beta (these are float in rocBLAS internals)
+ */
+template <typename T>
+struct gemm_impl
 {
-    const bool is_3inputs = (args.size() == 4);
+    gemm_impl(const shape& output_shape,
+              const std::vector<shape>& input_shapes,
+              T alpha_param,
+              T beta_param,
+              bool int8_x4_format,
+              bool compute_fp32_flag)
+        : alpha(alpha_param),
+          beta(beta_param),
+          is_3inputs(input_shapes.size() == 4),
+          compute_fp32(compute_fp32_flag)
+    {
         if(not is_3inputs)
         {
             beta = 0;
         }
 
-    bool transa     = is_transposed(args[0].get_shape());
-    bool transb     = is_transposed(args[1].get_shape());
+        // Create lambdas that will cast alpha, beta to the output shape type
+        // and retain the values being pointed to
+        output_shape.visit_type([&](auto as) {
+            auto alpha_r = as(alpha);
+            auto beta_r  = as(beta);
+            if(compute_fp32)
+            {
+                get_alpha = [=] { return &alpha; };
+                get_beta  = [=] { return &beta; };
+            }
+            else
+            {
+                get_alpha = [=] { return &alpha_r; };
+                get_beta  = [=] { return &beta_r; };
+            }
+        });
+
+        transa     = is_transposed(input_shapes[0]);
+        transb     = is_transposed(input_shapes[1]);
         auto n_dim = output_shape.lens().size();
-    auto dim_1      = n_dim - 1;
         auto dim_0 = n_dim - 2;
-    rocblas_int lda = args[0].get_shape().strides()[transa ? dim_1 : dim_0];
-    rocblas_int ldb = args[1].get_shape().strides()[transb ? dim_1 : dim_0];
-    rocblas_int ldc = args[2].get_shape().strides()[dim_0];
-    rocblas_int ldd = is_3inputs ? args[3].get_shape().strides()[dim_0] : ldc;
+        auto dim_1 = n_dim - 1;
+        // Leading dimensions of matrices
+        lda = input_shapes[0].strides()[transa ? dim_1 : dim_0];
+        ldb = input_shapes[1].strides()[transb ? dim_1 : dim_0];
+        ldc = input_shapes[2].strides()[dim_0];
+        ldd = is_3inputs ? input_shapes[3].strides()[dim_0] : ldc;
 
-    rocblas_datatype arg_type = get_type(args[0].get_shape().type());
-    auto output_type          = arg_type;
+        arg_type    = get_type(input_shapes[0].type());
+        output_type = arg_type;
         if(output_type == rocblas_datatype_i8_r)
         {
             output_type = rocblas_datatype_i32_r;
         }
-    auto compute_type = output_type;
+        compute_type = output_type;
         if(compute_fp32)
         {
             if(arg_type == rocblas_datatype_f16_r)
                 compute_type = rocblas_datatype_f32_r;
         }
 
-#if ROCBLAS_VERSION_MAJOR >= 2 && ROCBLAS_VERSION_MINOR >= 38
-    rocblas_gemm_flags flag =
-        int8_x4_format ? rocblas_gemm_flags_pack_int8x4 : rocblas_gemm_flags_none;
-#else
-    (void)int8_x4_format;
-    int flag = 0;
-#endif
-
-    auto a_lens = args[0].get_shape().lens();
-    auto b_lens = args[1].get_shape().lens();
-    output_shape.visit_type([&](auto as) {
-        auto alpha_r = as(alpha);
-        auto beta_r  = as(beta);
-
-        // use void pointer to select different data type if using fp32 mode
-        void* alpha_v = &alpha_r;
-        void* beta_v  = &beta_r;
+        int8_flag = int8_x4_format ? rocblas_gemm_flags_pack_int8x4 : rocblas_gemm_flags_none;
 
-        if(compute_fp32)
-        {
-            alpha_v = &alpha;
-            beta_v  = &beta;
-        }
+        auto a_lens = input_shapes[0].lens();
+        auto b_lens = input_shapes[1].lens();
 
         auto out_lens = output_shape.lens();
-        rocblas_int m   = out_lens[dim_0];
-        rocblas_int n   = out_lens[dim_1];
-        rocblas_int k   = args[0].get_shape().lens()[dim_1];
-        auto to_pointer = [&](auto&& arg) { return as.from(arg.data()); };
-        if(args[0].get_shape().type() == shape::int8_type and (k % 4) != 0 and int8_x4_format)
+        m             = out_lens[dim_0];
+        n             = out_lens[dim_1];
+        k             = input_shapes[0].lens()[dim_1];
+        if(input_shapes[0].type() == shape::int8_type and (k % 4) != 0 and int8_x4_format)
         {
-            MIGRAPHX_THROW("ROCBLAS_GEMM: k size of int8 type input must be mutlple of 4!");
+            MIGRAPHX_THROW("ROCBLAS_GEMM: k size of int8 type input must be multiple of 4!");
         }
 
-        auto num_matrices = std::accumulate(
+        a_stride     = get_batch_stride(input_shapes[0]);
+        b_stride     = get_batch_stride(input_shapes[1]);
+        c_stride     = get_batch_stride(input_shapes[2]);
+        d_stride     = is_3inputs ? get_batch_stride(input_shapes[3]) : c_stride;
+        num_matrices = std::accumulate(
             out_lens.rbegin() + 2, out_lens.rend(), std::size_t{1}, std::multiplies<std::size_t>());
-        if(num_matrices == 1 or (num_matrices > 1 and get_batch_stride(args[1]) == 0))
+
+        if(num_matrices == 1 or (num_matrices > 1 and b_stride == 0))
         {
             // If the batch dimension of B is broadcasted, then we can
             // multiply m by the batch_size and use rocblas_gemm_ex
             // instead of rocblas_gemm_strided_batched_ex.
             m *= num_matrices;
+            strided_batched = false;
+        }
+    }
+
+    void run(context& ctx, const std::vector<argument>& input_args, int32_t solution_idx = 0) const
+    {
+        if(strided_batched)
+        {
+            auto common_args = create_strided_batched_args_common(ctx, input_args);
+            rocblas_invoke(&rocblas_gemm_strided_batched_ex,
+                           common_args,
+                           rocblas_gemm_algo_standard,
+                           solution_idx,
+                           int8_flag);
+        }
+        else
+        {
+            auto common_args = create_gemm_ex_args_common(ctx, input_args);
+            rocblas_invoke(
+                &rocblas_gemm_ex, common_args, rocblas_gemm_algo_standard, solution_idx, int8_flag);
+        }
+    }
+
+#ifdef ROCBLAS_BETA_FEATURES_API
+    auto validate(context& ctx, const std::vector<shape>& input_shapes, int32_t solution_idx) const
+    {
+        std::vector<argument> input_args;
+        std::transform(input_shapes.begin(),
+                       input_shapes.end(),
+                       std::back_inserter(input_args),
+                       [](const shape& x) { return to_gpu(generate_argument(x)); });
+
+        return validate(ctx, input_args, solution_idx);
+    }
+
+    /**
+     * Checks a particular solution for validity by running it with the flag
+     * rocblas_gemm_flags_check_solution_index (could be invalid if this model was
+     * tuned with a different rocBLAS version)
+     *
+     * @return Returns either solution_idx if valid, or else the default value 0
+     * if not.  The default does not mean list index 0, but tells the picker
+     * to choose a solution.
+     */
+    int32_t
+    validate(context& ctx, const std::vector<argument>& input_args, int32_t solution_idx) const
+    {
+        rocblas_status_ check_valid(rocblas_status_success);
+
+        if(strided_batched)
+        {
+            auto common_args = create_strided_batched_args_common(ctx, input_args);
+            check_valid      = rocblas_invoke(&rocblas_gemm_strided_batched_ex,
+                                         common_args,
+                                         rocblas_gemm_algo_solution_index,
+                                         solution_idx,
+                                         rocblas_gemm_flags_check_solution_index);
+        }
+        else
+        {
+            auto common_args = create_gemm_ex_args_common(ctx, input_args);
+            check_valid      = rocblas_invoke(&rocblas_gemm_ex,
+                                         common_args,
+                                         rocblas_gemm_algo_standard,
+                                         solution_idx,
+                                         rocblas_gemm_flags_check_solution_index);
+        }
 
-            // the rocblas_gemm API handles inputs and output matrices as
-            // column-major format. When doing a C = A * B, we actually do
-            // C^T = (B^T) * (A^T). That is the reason we input args[1] as
-            // A and args[0] as B in calling the rocblas_gemm.
-            rocblas_invoke(&rocblas_gemm_ex,
-                           ctx.get_stream().get_rocblas(),
+        if(check_valid == rocblas_status_invalid_value)
+        {
+            std::cerr << "WARNING:  tuned solution is invalid; reverting to default" << std::endl;
+            return 0;
+        }
+        return solution_idx;
+    }
+#endif
+
+    /**
+     * Helper method to create that subset of a long rocBLAS argument list that is common
+     * to multiple "...strided_batched..." calls.
+     *
+     * The rocblas_gemm API handles inputs and output matrices as
+     *  column-major format. When doing a C = A * B, we actually do
+     *  C^T = (B^T) * (A^T). That is the reason we input args[1] as
+     *   A and args[0] as B in calling the rocblas_gemm.
+     *
+     */
+    auto create_strided_batched_args_common(context& ctx, const std::vector<argument>& args) const
+    {
+        return pack(ctx.get_stream().get_rocblas(),
                     transb ? rocblas_operation_transpose : rocblas_operation_none,
                     transa ? rocblas_operation_transpose : rocblas_operation_none,
                     n,
                     m,
                     k,
-                           alpha_v,
-                           to_pointer(args.at(1)),
+                    get_alpha(),
+                    args[1].data(),
                     arg_type,
                     ldb,
-                           to_pointer(args.at(0)),
+                    b_stride,
+                    args[0].data(),
                     arg_type,
                     lda,
-                           beta_v,
-                           to_pointer(args[2]),
+                    a_stride,
+                    get_beta(),
+                    args[2].data(),
                     output_type,
                     ldc,
-                           is_3inputs ? to_pointer(args[3]) : to_pointer(args[2]),
+                    c_stride,
+                    is_3inputs ? args[3].data() : args[2].data(),
                     output_type,
                     ldd,
-                           compute_type,
-                           rocblas_gemm_algo_standard,
-                           0,
-                           flag);
+                    d_stride,
+                    num_matrices,
+                    compute_type);
     }
-        else
+
+    /**
+     * Helper method to create that subset of a long rocBLAS argument list that is common
+     * to multiple "gemm_ex..." calls.
+     *
+     * The rocblas_gemm API handles inputs and output matrices as
+     *  column-major format. When doing a C = A * B, we actually do
+     *   C^T = (B^T) * (A^T). That is the reason we input args[1] as
+     *   A and args[0] as B in calling the rocblas_gemm.
+     *
+     * */
+    auto create_gemm_ex_args_common(context& ctx, const std::vector<argument>& args) const
     {
-            auto a_stride = get_batch_stride(args[0]);
-            auto b_stride = get_batch_stride(args[1]);
-            auto c_stride = get_batch_stride(args[2]);
-            auto d_stride = is_3inputs ? get_batch_stride(args[3]) : c_stride;
-            rocblas_invoke(&rocblas_gemm_strided_batched_ex,
-                           ctx.get_stream().get_rocblas(),
+        return pack(ctx.get_stream().get_rocblas(),
                     transb ? rocblas_operation_transpose : rocblas_operation_none,
                     transa ? rocblas_operation_transpose : rocblas_operation_none,
                     n,
                     m,
                     k,
-                           alpha_v,
-                           to_pointer(args.at(1)),
+                    get_alpha(),
+                    args[1].data(),
                     arg_type,
                     ldb,
-                           b_stride,
-                           to_pointer(args.at(0)),
+                    args[0].data(),
                     arg_type,
                     lda,
-                           a_stride,
-                           beta_v,
-                           to_pointer(args[2]),
+                    get_beta(),
+                    args[2].data(),
                     output_type,
                     ldc,
-                           c_stride,
-                           is_3inputs ? to_pointer(args[3]) : to_pointer(args[2]),
+                    is_3inputs ? args[3].data() : args[2].data(),
                     output_type,
                     ldd,
-                           d_stride,
-                           num_matrices,
-                           compute_type,
-                           rocblas_gemm_algo_standard,
-                           0,
-                           flag);
+                    compute_type);
     }
-    });
-}
+#ifdef ROCBLAS_BETA_FEATURES_API
+    /**
+     * Find best rocBLAS solution:  Get list of solutions and try them all, returning the index
+     * of the fastest one.
+     */
+    int tune(context& ctx, const std::vector<shape>& input_shapes) const
+    {
+        // tuning meta parameters
+        const int hot_calls  = 40;
+        const int cold_calls = 1;
+
+        std::vector<argument> input_args;
+        std::transform(input_shapes.begin(),
+                       input_shapes.end(),
+                       std::back_inserter(input_args),
+                       [](const shape& x) { return to_gpu(generate_argument(x)); });
+
+        // Get the solutions list in 2 rocBLAS steps:
+        // 1.  Find out how many solutions there are and allocate the array
+        // 2.  Get the solutions
+        //
+        rocblas_int list_size = 0;
+        std::vector<rocblas_int> solution_indices;
+        if(strided_batched)
+        {
+            auto common_args = create_strided_batched_args_common(ctx, input_args);
+            rocblas_invoke(&rocblas_gemm_strided_batched_ex_get_solutions,
+                           common_args,
+                           rocblas_gemm_algo_solution_index,
+                           int8_flag,
+                           nullptr,
+                           &list_size);
+            solution_indices.resize(list_size);
+
+            auto common_sol_args = create_strided_batched_args_common(ctx, input_args);
+            rocblas_invoke(&rocblas_gemm_strided_batched_ex_get_solutions,
+                           common_sol_args,
+                           rocblas_gemm_algo_solution_index,
+                           int8_flag,
+                           solution_indices.data(),
+                           &list_size);
+        }
+        else
+        {
+            auto common_args = create_gemm_ex_args_common(ctx, input_args);
+            rocblas_invoke(&rocblas_gemm_ex_get_solutions,
+                           common_args,
+                           rocblas_gemm_algo_solution_index,
+                           int8_flag,
+                           nullptr,
+                           &list_size);
+            solution_indices.resize(list_size);
+
+            auto common_sol_args = create_gemm_ex_args_common(ctx, input_args);
+            rocblas_invoke(&rocblas_gemm_ex_get_solutions,
+                           common_sol_args,
+                           rocblas_gemm_algo_solution_index,
+                           int8_flag,
+                           solution_indices.data(),
+                           &list_size);
+        }
+
+        double bestTime   = std::numeric_limits<double>::max();
+        double first_time = -1;
+        // Initialize to default solution index
+        rocblas_int bestSol = 0;
+        for(auto sol : solution_indices)
+        {
+            // Define the function to be timed
+            auto run_func = [&]() {
+                run(ctx, input_args, sol);
+                ctx.finish();
+            };
+
+            // Warmup: the first few calls to an op. may not be representative since there is
+            // more time taken initializing caches, etc. so we won't time them.
+            for(int cc = 0; cc < cold_calls; ++cc)
+            {
+                run_func();
+            }
+            double host_time = 0.0;
 
-void gemm(context& ctx,
+            for(int hc = 0; hc < hot_calls; ++hc)
+            {
+                ctx.finish();
+                host_time += time<microseconds>(run_func);
+            }
+            // todo:  Measured time dropped from 20 us to about 6.7 us when I raised hot_calls from
+            // 1 to 11. The higher the hot_calls value, the faster per-call time up to at least 25,
+            // and increasing cold_calls makes little or no difference.  Why?
+            host_time /= hot_calls;
+
+            // debug only: track time for first solution.
+            if(first_time < 0)
+                first_time = host_time;
+
+            // track current best
+            if(host_time < bestTime)
+            {
+                std::cout << " current best index " << sol << ", time " << host_time << std::endl;
+                bestSol  = sol;
+                bestTime = host_time;
+            }
+        }
+        std::cout << "Winner: " << bestSol << " in " << bestTime << " us, beats " << first_time
+                  << std::endl;
+        return bestSol;
+    }
+#endif
+    private:
+    size_t num_matrices;
+    rocblas_int m, n, k;
+    bool transa, transb;
+    T alpha, beta;
+
+    std::function<const void*()> get_alpha;
+    std::function<const void*()> get_beta;
+    flag_type int8_flag;
+    rocblas_int lda, ldb, ldc, ldd;
+    rocblas_int a_stride, b_stride, c_stride, d_stride;
+    rocblas_datatype compute_type, arg_type, output_type;
+    bool strided_batched = true, is_3inputs = true, compute_fp32 = true;
+}; // gemm_impl
+
+void gemm_compute(context& ctx,
                   const shape& output_shape,
                   const std::vector<argument>& args,
                   float alpha,
                   float beta,
                   bool int8_x4_format,
-          bool compute_fp32)
+                  bool compute_fp32,
+                  int32_t solution_idx)
 {
-    gemm_impl(ctx, output_shape, args, alpha, beta, int8_x4_format, compute_fp32);
+    std::vector<shape> input_shapes;
+    std::transform(args.begin(),
+                   args.end(),
+                   std::back_inserter(input_shapes),
+                   [](const argument& x) { return x.get_shape(); });
+    auto gemm_item =
+        gemm_impl<float>(output_shape, input_shapes, alpha, beta, int8_x4_format, compute_fp32);
+    gemm_item.run(ctx, args, solution_idx);
 }
 
-void gemm(context& ctx,
+void gemm_compute(context& ctx,
                   const shape& output_shape,
                   const std::vector<argument>& args,
                   int32_t alpha,
                   int32_t beta,
                   bool int8_x4_format,
-          bool compute_fp32)
+                  bool compute_fp32,
+                  int32_t solution_idx)
+{
+    std::vector<shape> input_shapes;
+    std::transform(args.begin(),
+                   args.end(),
+                   std::back_inserter(input_shapes),
+                   [](const argument& x) { return x.get_shape(); });
+    auto gemm_item =
+        gemm_impl<int32_t>(output_shape, input_shapes, alpha, beta, int8_x4_format, compute_fp32);
+    gemm_item.run(ctx, args, solution_idx);
+}
+
+int32_t gemm_finalize(context& ctx,
+                      const shape& output_shape,
+                      const std::vector<shape>& input_shapes,
+                      float alpha,
+                      float beta,
+                      bool int8_x4_format,
+                      bool compute_fp32,
+                      int32_t solution_idx)
 {
-    gemm_impl(ctx, output_shape, args, alpha, beta, int8_x4_format, compute_fp32);
+#ifdef ROCBLAS_BETA_FEATURES_API
+
+    if(ctx.get_exhaustive_tune_flag() && solution_idx == 0)
+    // if((true))
+    {
+        auto gemm_item =
+            gemm_impl<float>(output_shape, input_shapes, alpha, beta, int8_x4_format, compute_fp32);
+        solution_idx = gemm_item.tune(ctx, input_shapes);
+    }
+    else if(solution_idx != 0)
+    {
+        // If a tuned solution index is already given, don't tune again but validate
+        // in case the data was tuned with a different rocBLAS version
+        auto gemm_item =
+            gemm_impl<float>(output_shape, input_shapes, alpha, beta, int8_x4_format, compute_fp32);
+        solution_idx = gemm_item.validate(ctx, input_shapes, solution_idx);
+    }
+#else
+    // suppress compiler warnings
+    (void)ctx, (void)output_shape, (void)input_shapes;
+    (void)alpha, (void)beta, (void)int8_x4_format, (void)compute_fp32;
+#endif
+    return solution_idx;
+}
+
+int32_t gemm_finalize(context& ctx,
+                      const shape& output_shape,
+                      const std::vector<shape>& input_shapes,
+                      int32_t alpha,
+                      int32_t beta,
+                      bool int8_x4_format,
+                      bool compute_fp32,
+                      int32_t solution_idx)
+{
+#ifdef ROCBLAS_BETA_FEATURES_API
+
+    if(ctx.get_exhaustive_tune_flag() && solution_idx == 0)
+    {
+        auto gemm_item = gemm_impl<int32_t>(
+            output_shape, input_shapes, alpha, beta, int8_x4_format, compute_fp32);
+        solution_idx = gemm_item.tune(ctx, input_shapes);
+    }
+    else if(solution_idx != 0)
+    {
+        // If a tuned solution index is already given, don't tune again but validate
+        // in case the data was tuned with a different rocBLAS version
+        auto gemm_item = gemm_impl<int32_t>(
+            output_shape, input_shapes, alpha, beta, int8_x4_format, compute_fp32);
+        solution_idx = gemm_item.validate(ctx, input_shapes, solution_idx);
+    }
+#else
+    // suppress compiler warnings
+    (void)ctx, (void)output_shape, (void)input_shapes;
+    (void)alpha, (void)beta, (void)int8_x4_format, (void)compute_fp32;
+#endif
+    return solution_idx;
 }
 
 } // namespace gpu

src/targets/gpu/include/migraphx/gpu/convolution.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
@@ -27,6 +27,7 @@
 #include <migraphx/shape.hpp>
 #include <migraphx/generate.hpp>
 #include <migraphx/operation.hpp>
+#include <migraphx/register_op.hpp>
 #include <migraphx/gpu/miopen.hpp>
 #include <migraphx/op/identity.hpp>
 #include <migraphx/op/convolution.hpp>
@@ -35,6 +36,9 @@
 #include <unordered_map>
 #include <migraphx/reflect.hpp>
 #include <migraphx/gpu/context.hpp>
+
+MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_ENABLE_CONV_TUNING);
+
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
@@ -163,6 +167,20 @@ struct miopen_convolution
         auto w_desc                = make_tensor(reshape_if_1d(inputs[1]), int8_x4_format);
         auto y_desc                = make_tensor(reshape_if_1d(output_shape));
         std::size_t workspace_size = 0;
+
+        // TODO:  Using an environment variable to disable convolution tuning while still allowing
+        // other (i.e. GEMM) tuning with --exhaustive-tune, is a workaround for a problem in which
+        // convolution tuning takes multiple hours.
+        const bool convo_tune_enabled = enabled(MIGRAPHX_ENABLE_CONV_TUNING{});
+        if(ctx.get_exhaustive_tune_flag() and not convo_tune_enabled)
+        {
+            std::cerr
+                << "WARNING: MIGraphX convolution tuning not enabled. To run tuning, set both the "
+                   "env "
+                   "var MIGRAPHX_ENABLE_CONV_TUNING and the command argument --exhaustive-tune."
+                << std::endl;
+        }
+
 #ifdef MIGRAPHX_HAS_FIND_2_API
         {
             auto conv_problem = make_obj<miopen_problem>(
@@ -174,8 +192,9 @@ struct miopen_convolution
 
             auto* miopen_stream_handle = ctx.get_stream().get_miopen();
 
-            solution_ptr = find_solution(
-                miopen_stream_handle, conv_problem.get(), ctx.get_exhaustive_tune_flag());
+            solution_ptr = find_solution(miopen_stream_handle,
+                                         conv_problem.get(),
+                                         ctx.get_exhaustive_tune_flag() and convo_tune_enabled);
             auto status  = miopenGetSolutionWorkspaceSize(solution_ptr.get(), &workspace_size);
             if(status != miopenStatusSuccess)
                 MIGRAPHX_THROW("MIOpen" + op.name() + " : failed to get solution's workspace size");
@@ -233,7 +252,8 @@ struct miopen_convolution
                                                        &perf,
                                                        workspace.implicit(),
                                                        workspace_size,
-                                                       ctx.get_exhaustive_tune_flag());
+                                                       ctx.get_exhaustive_tune_flag() and
+                                                           convo_tune_enabled);
         if(status != miopenStatusSuccess)
             MIGRAPHX_THROW("MIOpen " + op.name() + " : find convolution failed");
         algo = perf.fwd_algo;

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
@@ -53,6 +52,7 @@ struct rocblas_gemm
     bool int8_x4_format  = true;
     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)
@@ -62,7 +62,8 @@ struct rocblas_gemm
                               f(self.beta, "beta"),
                               f(self.int8_x4_format, "int8_x4_format"),
                               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
@@ -113,17 +114,19 @@ struct rocblas_gemm
     {
         if(this->name() == "gpu::gemm")
         {
-            gemm(ctx, output_shape, args, alpha, beta, int8_x4_format, compute_fp32);
+            gemm_compute(
+                ctx, output_shape, args, alpha, beta, int8_x4_format, compute_fp32, solution_idx);
         }
         else
         {
-            gemm(ctx,
+            gemm_compute(ctx,
                          output_shape,
                          args,
                          int32_t(alpha),
                          int32_t(beta),
                          int8_x4_format,
-                 compute_fp32);
+                         compute_fp32,
+                         solution_idx);
         }
         return args.back();
     }
@@ -132,6 +135,40 @@ struct rocblas_gemm
     {
         return shapes.size() - 1;
     }
+
+    void finalize(context& ctx, const shape& output_shape, const std::vector<shape>& input_shapes)
+    {
+#ifdef ROCBLAS_BETA_FEATURES_API
+        if(ctx.get_exhaustive_tune_flag() && solution_idx == 0)
+        {
+            if(this->name() == "gpu::gemm")
+            {
+                solution_idx = gemm_finalize(ctx,
+                                             output_shape,
+                                             input_shapes,
+                                             alpha,
+                                             beta,
+                                             int8_x4_format,
+                                             compute_fp32,
+                                             solution_idx);
+            }
+            else
+            {
+                solution_idx = gemm_finalize(ctx,
+                                             output_shape,
+                                             input_shapes,
+                                             int32_t(alpha),
+                                             int32_t(beta),
+                                             int8_x4_format,
+                                             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,29 +24,76 @@
 #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/generate.hpp>
 #include <migraphx/gpu/context.hpp>
+#include <migraphx/reduce_dims.hpp>
+#include <migraphx/gpu/hip.hpp>
+#include <migraphx/time.hpp>
+
+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,
+#if ROCBLAS_VERSION_MAJOR >= 2 && ROCBLAS_VERSION_MINOR >= 38
+using flag_type = rocblas_gemm_flags;
+#else
+using flag_type = int;
+#endif
+
+/**
+ * @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 int8_x4_format .
+ * @param compute_fp32 .
+ */
+void gemm_compute(context& ctx,
                   const shape& output_shape,
                   const std::vector<argument>& args,
                   float alpha,
                   float beta,
                   bool int8_x4_format,
-          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 int8_x4_format,
+                  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 int8_x4_format,
                       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 int8_x4_format,
+                      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
@@ -25,6 +25,7 @@
 #define MIGRAPHX_GUARD_MIGRAPHLIB_ROCBLAS_HPP
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/config.hpp>
+// ROCBLAS_BETA_FEATURES_API is defined by CMake, if available.
 #include <rocblas/rocblas.h>
 
 namespace migraphx {

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/allocation_model.hpp>
+#include <migraphx/gpu/context.hpp>
+#include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/target.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 ROCBLAS_BETA_FEATURES_API
+    EXPECT(0 != solution_idx.to<std::size_t>());
+#else
+    EXPECT(0 == solution_idx.to<std::size_t>());
+#endif
+}
+
+// TODO:  make tests that run both strided and not-strided GEMMs.  Also, try tuning with an invalid
+// start value.
+
+int main(int argc, const char* argv[]) { test::run(argc, argv); }