add tests and fix cmake

src/targets/gpu/CMakeLists.txt

@@ -254,7 +254,7 @@ check_library_exists(MIOpen "miopenFindSolutions" "${MIOPEN_LOCATION}" HAS_FIND_
 # Beta API for automated GEMM tuning
 check_library_exists(roc::rocblas "rocblas_gemm_ex_get_solutions" "${ROCBLAS_LOCATION}" HAS_ROCBLAS_TUNING_BETA_FEATURE_API)
 # rocblas FP8 API
-check_library_exists(roc::rocblas "rocblas_gemm_ex3" "${ROCBLAS_LOCATION}" HAS_ROCBLAS_FP8_BETA_API)
+check_library_exists(roc::rocblas "rocblas_gemm_strided_batched_ex3" "${ROCBLAS_LOCATION}" HAS_ROCBLAS_FP8_BETA_API)
 
 set(MIGRAPHX_USE_FIND_2_API "${HAS_FIND_2_API}" CACHE BOOL "")
 

test/verify/gemm_2args_mm_2.cpp

@@ -22,19 +22,21 @@
  * THE SOFTWARE.
  */
 
+#include "migraphx/shape.hpp"
 #include "verify_program.hpp"
 #include <migraphx/program.hpp>
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct gemm_2args_mm_2 : verify_program<gemm_2args_mm_2>
+template <migraphx::shape::type_t DType>
+struct gemm_2args_mm_2 : verify_program<gemm_2args_mm_2<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape m1_shape{migraphx::shape::float_type, {2, 2, 3}};
-        migraphx::shape m2_shape{migraphx::shape::float_type, {3, 4}};
+        migraphx::shape m1_shape{DType, {2, 2, 3}};
+        migraphx::shape m2_shape{DType, {3, 4}};
         auto l1 = mm->add_parameter("1", m1_shape);
         auto l2 = mm->add_parameter("2", m2_shape);
         auto bl2 =
@@ -45,3 +47,6 @@ struct gemm_2args_mm_2 : verify_program<gemm_2args_mm_2>
         return p;
     }
 };
+
+template struct gemm_2args_mm_2<migraphx::shape::float_type>;
+template struct gemm_2args_mm_2<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/gemm_2args_mm_3.cpp

@@ -22,19 +22,21 @@
  * THE SOFTWARE.
  */
 
+#include "migraphx/shape.hpp"
 #include "verify_program.hpp"
 #include <migraphx/program.hpp>
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct gemm_2args_mm_3 : verify_program<gemm_2args_mm_3>
+template <migraphx::shape::type_t DType>
+struct gemm_2args_mm_3 : verify_program<gemm_2args_mm_3<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape m1_shape{migraphx::shape::float_type, {1, 2, 3}};
-        migraphx::shape m2_shape{migraphx::shape::float_type, {3, 3, 4}};
+        migraphx::shape m1_shape{DType, {1, 2, 3}};
+        migraphx::shape m2_shape{DType, {3, 3, 4}};
         auto l1 = mm->add_parameter("1", m1_shape);
         auto bl1 =
             mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", {3, 2, 3}}}), l1);
@@ -45,3 +47,6 @@ struct gemm_2args_mm_3 : verify_program<gemm_2args_mm_3>
         return p;
     }
 };
+
+template struct gemm_2args_mm_3<migraphx::shape::float_type>;
+template struct gemm_2args_mm_3<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/gemm_2args_mm_4.cpp

@@ -22,19 +22,21 @@
  * THE SOFTWARE.
  */
 
+#include "migraphx/shape.hpp"
 #include "verify_program.hpp"
 #include <migraphx/program.hpp>
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct gemm_2args_mm_4 : verify_program<gemm_2args_mm_4>
+template <migraphx::shape::type_t DType>
+struct gemm_2args_mm_4 : verify_program<gemm_2args_mm_4<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape m1_shape{migraphx::shape::float_type, {2, 3}};
-        migraphx::shape m2_shape{migraphx::shape::float_type, {3, 3, 4}};
+        migraphx::shape m1_shape{DType, {2, 3}};
+        migraphx::shape m2_shape{DType, {3, 3, 4}};
         auto l1 = mm->add_parameter("1", m1_shape);
         auto bl1 =
             mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", {3, 2, 3}}}), l1);
@@ -45,3 +47,6 @@ struct gemm_2args_mm_4 : verify_program<gemm_2args_mm_4>
         return p;
     }
 };
+
+template struct gemm_2args_mm_4<migraphx::shape::float_type>;
+template struct gemm_2args_mm_4<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/gemm_2args_mm_5.cpp

@@ -27,14 +27,15 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct gemm_2args_mm_5 : verify_program<gemm_2args_mm_5>
+template <migraphx::shape::type_t DType>
+struct gemm_2args_mm_5 : verify_program<gemm_2args_mm_5<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape m1_shape{migraphx::shape::float_type, {2, 1, 2, 3}};
-        migraphx::shape m2_shape{migraphx::shape::float_type, {2, 3, 3, 4}};
+        migraphx::shape m1_shape{DType, {2, 1, 2, 3}};
+        migraphx::shape m2_shape{DType, {2, 3, 3, 4}};
         auto l1  = mm->add_parameter("1", m1_shape);
         auto bl1 = mm->add_instruction(
             migraphx::make_op("multibroadcast", {{"out_lens", {2, 3, 2, 3}}}), l1);
@@ -45,3 +46,6 @@ struct gemm_2args_mm_5 : verify_program<gemm_2args_mm_5>
         return p;
     }
 };
+
+template struct gemm_2args_mm_5<migraphx::shape::float_type>;
+template struct gemm_2args_mm_5<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/gemm_2args_mm_6.cpp

@@ -27,14 +27,16 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct gemm_2args_mm_6 : verify_program<gemm_2args_mm_6>
+template <migraphx::shape::type_t DType>
+
+struct gemm_2args_mm_6 : verify_program<gemm_2args_mm_6<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape m1_shape{migraphx::shape::float_type, {2, 1, 2, 3}};
-        migraphx::shape m2_shape{migraphx::shape::float_type, {1, 3, 3, 4}};
+        migraphx::shape m1_shape{DType, {2, 1, 2, 3}};
+        migraphx::shape m2_shape{DType, {1, 3, 3, 4}};
         auto l1  = mm->add_parameter("1", m1_shape);
         auto bl1 = mm->add_instruction(
             migraphx::make_op("multibroadcast", {{"out_lens", {2, 3, 2, 3}}}), l1);
@@ -47,3 +49,6 @@ struct gemm_2args_mm_6 : verify_program<gemm_2args_mm_6>
         return p;
     }
 };
+
+template struct gemm_2args_mm_6<migraphx::shape::float_type>;
+template struct gemm_2args_mm_6<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/gemm_2args_mm_7.cpp

@@ -27,14 +27,15 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct gemm_2args_mm_7 : verify_program<gemm_2args_mm_7>
+template <migraphx::shape::type_t DType>
+struct gemm_2args_mm_7 : verify_program<gemm_2args_mm_7<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape m1_shape{migraphx::shape::float_type, {2, 3}};
-        migraphx::shape m2_shape{migraphx::shape::float_type, {2, 3, 3, 4}};
+        migraphx::shape m1_shape{DType, {2, 3}};
+        migraphx::shape m2_shape{DType, {2, 3, 3, 4}};
         auto l1  = mm->add_parameter("1", m1_shape);
         auto bl1 = mm->add_instruction(
             migraphx::make_op("multibroadcast", {{"out_lens", {2, 3, 2, 3}}}), l1);
@@ -45,3 +46,6 @@ struct gemm_2args_mm_7 : verify_program<gemm_2args_mm_7>
         return p;
     }
 };
+
+template struct gemm_2args_mm_7<migraphx::shape::float_type>;
+template struct gemm_2args_mm_7<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/gemm_2args_mm_8.cpp

@@ -27,14 +27,15 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct gemm_2args_mm_8 : verify_program<gemm_2args_mm_8>
+template <migraphx::shape::type_t DType>
+struct gemm_2args_mm_8 : verify_program<gemm_2args_mm_8<DType>>
 {
     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}};
+        migraphx::shape a_shape{DType, {2, 128, 32}, {4096, 1, 128}};
+        migraphx::shape b_shape{DType, {32, 32}};
         auto a  = mm->add_parameter("a", a_shape);
         auto b  = mm->add_parameter("b", b_shape);
         auto bb = mm->add_instruction(
@@ -45,3 +46,6 @@ struct gemm_2args_mm_8 : verify_program<gemm_2args_mm_8>
         return p;
     }
 };
+
+template struct gemm_2args_mm_8<migraphx::shape::float_type>;
+template struct gemm_2args_mm_8<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/gemm_2args_mv.cpp

@@ -27,14 +27,15 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct gemm_2args_mv : verify_program<gemm_2args_mv>
+template <migraphx::shape::type_t DType>
+struct gemm_2args_mv : verify_program<gemm_2args_mv<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape m1_shape{migraphx::shape::float_type, {3, 5}};
-        migraphx::shape m2_shape{migraphx::shape::float_type, {5}};
+        migraphx::shape m1_shape{DType, {3, 5}};
+        migraphx::shape m2_shape{DType, {5}};
         auto l1  = mm->add_parameter("1", m1_shape);
         auto l2  = mm->add_parameter("2", m2_shape);
         auto ul2 = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1}}}), l2);
@@ -44,3 +45,6 @@ struct gemm_2args_mv : verify_program<gemm_2args_mv>
         return p;
     }
 };
+
+template struct gemm_2args_mv<migraphx::shape::float_type>;
+template struct gemm_2args_mv<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/gemm_2args_vbm.cpp

@@ -27,14 +27,15 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct gemm_2args_vbm : verify_program<gemm_2args_vbm>
+template <migraphx::shape::type_t DType>
+struct gemm_2args_vbm : verify_program<gemm_2args_vbm<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape m1_shape{migraphx::shape::float_type, {5}};
-        migraphx::shape m2_shape{migraphx::shape::float_type, {2, 2, 5, 4}};
+        migraphx::shape m1_shape{DType, {5}};
+        migraphx::shape m2_shape{DType, {2, 2, 5, 4}};
         auto l1   = mm->add_parameter("1", m1_shape);
         auto ul1  = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {0}}}), l1);
         auto bul1 = mm->add_instruction(
@@ -48,3 +49,6 @@ struct gemm_2args_vbm : verify_program<gemm_2args_vbm>
         return p;
     }
 };
+
+template struct gemm_2args_vbm<migraphx::shape::float_type>;
+template struct gemm_2args_vbm<migraphx::shape::fp8e4m3fnuz_type>;