Modify hipblaslt_invoke to also have the ability to return with statu…

…s in addition to throwing an exception
ROCm · Nov 7, 2024 · 5471ab9 · 5471ab9
1 parent 100da97
commit 5471ab9
Show file tree

Hide file tree

Showing 2 changed files with 7 additions and 4 deletions.
diff --git a/src/targets/gpu/hip_gemm_impl.cpp b/src/targets/gpu/hip_gemm_impl.cpp
@@ -463,7 +463,7 @@ struct hip_gemm_impl
     validate(context& ctx, const std::vector<argument>& input_args, int32_t solution_idx) // const
     {
         auto common_args = create_hipblaslt_args_common(ctx, input_args, solution_idx);
-        auto check_valid = hipblaslt_invoke(&hipblasLtMatmul, common_args);
+        auto check_valid = hipblaslt_invoke(&hipblasLtMatmul, common_args, false);
         if(check_valid != HIPBLAS_STATUS_SUCCESS)
         {
             std::cerr << "WARNING:  tuned solution is invalid; reverting to default" << std::endl;

diff --git a/src/targets/gpu/include/migraphx/gpu/hipblaslt.hpp b/src/targets/gpu/include/migraphx/gpu/hipblaslt.hpp
@@ -76,14 +76,17 @@ inline auto hipblaslt_invoke(F f, Ts... xs)
 }
 
 template <class F, class Pack, class... Ts>
-auto hipblaslt_invoke(F f, Pack p, Ts... xs)
+auto hipblaslt_invoke(F f, Pack p, Ts... xs, bool fatal_error = true)
 {
     return p([=](auto... ws) {
         auto status = f(ws..., xs...);
         if(status != HIPBLAS_STATUS_SUCCESS)
         {
-            MIGRAPHX_THROW("hipblaslt_invoke: hipBlasLt call failed with status " +
-                           std::to_string(status));
+            if(fatal_error)
+            {
+                MIGRAPHX_THROW("hipblaslt_invoke: hipBlasLt call failed with status " +
+                               std::to_string(status));
+            }
         }
         return status;
     });