Add math and reduce tests

src/targets/gpu/kernels/include/migraphx/kernels/reduce.hpp

@@ -106,7 +106,7 @@ __device__ auto block_reduce(index idx, Op op, T init, Index n, F f)
 #endif
     using type = decltype(index::invoke_loop(f, 0, _c<0>));
     __shared__ type buffer[idx.max_nlocal() / lanes_per_thread];
-    type x = init;
+    type x = type(init);
     idx.local_stride(n, [&](auto i, auto d) { x = op(x, index::invoke_loop(f, i, d)); });
     dpp_reduce(x, op);
 
@@ -117,7 +117,7 @@ __device__ auto block_reduce(index idx, Op op, T init, Index n, F f)
     }
     __syncthreads();
 
-    type y = init;
+    type y = type(init);
     for(index_int i = 0; i < idx.nlocal() / lanes_per_thread; i++)
     {
         y = op(y, buffer[i]);
@@ -244,9 +244,8 @@ struct reducer_base
         {
             auto&& derived = static_cast<const Derived&>(*this);
             auto t         = derived.slice(x);
-            return make_storage_access<typename decltype(t)::type>([=](auto i, auto...) -> auto& {
-                return t[i];
-            });
+            return make_storage_access<typename decltype(t)::type>(
+                [=](auto i, auto...) -> auto& { return t[i]; });
         }
     }
 
@@ -482,7 +481,7 @@ struct lane
         __device__ auto reduce_impl(Op op, T init, Read read, N n, U&& x, Us&&... xs) const
         {
             using type = remove_reference_t<decltype(x(0, _c<0>))>;
-            type r     = init;
+            type r     = type(init);
             for(index_int j = 0; j < n; j++)
             {
                 r = op(r, read(x(j, _c<0>), xs(j, _c<0>)...));

test/verify/test_acosh.cpp

@@ -27,13 +27,14 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_acosh : verify_program<test_acosh>
+template <migraphx::shape::type_t DType>
+struct test_acosh : verify_program<test_acosh<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape s{migraphx::shape::float_type, {16}};
+        migraphx::shape s{DType, {16}};
         auto x       = mm->add_parameter("x", s);
         auto min_val = mm->add_literal(1.1f);
         auto max_val = mm->add_literal(100.0f);
@@ -46,3 +47,7 @@ struct test_acosh : verify_program<test_acosh>
         return p;
     }
 };
+
+template struct test_acosh<migraphx::shape::float_type>;
+// template struct test_acosh<migraphx::shape::half_type>;
+// template struct test_acosh<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/test_asin.cpp

@@ -27,15 +27,20 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_asin : verify_program<test_asin>
+template <migraphx::shape::type_t DType>
+struct test_asin : verify_program<test_asin<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape s{migraphx::shape::float_type, {16}};
+        migraphx::shape s{DType, {16}};
         auto x = mm->add_parameter("x", s);
         mm->add_instruction(migraphx::make_op("asin"), x);
         return p;
     }
 };
+
+template struct test_asin<migraphx::shape::float_type>;
+template struct test_asin<migraphx::shape::half_type>;
+template struct test_asin<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/test_asinh.cpp

@@ -27,15 +27,20 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_asinh : verify_program<test_asinh>
+template <migraphx::shape::type_t DType>
+struct test_asinh : verify_program<test_asinh<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape s{migraphx::shape::float_type, {16}};
+        migraphx::shape s{DType, {16}};
         auto x = mm->add_parameter("x", s);
         mm->add_instruction(migraphx::make_op("asinh"), x);
         return p;
     }
 };
+
+template struct test_asinh<migraphx::shape::float_type>;
+template struct test_asinh<migraphx::shape::half_type>;
+template struct test_asinh<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/test_atan.cpp

@@ -27,15 +27,20 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_atan : verify_program<test_atan>
+template <migraphx::shape::type_t DType>
+struct test_atan : verify_program<test_atan<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape s{migraphx::shape::float_type, {16}};
+        migraphx::shape s{DType, {16}};
         auto x = mm->add_parameter("x", s);
         mm->add_instruction(migraphx::make_op("atan"), x);
         return p;
     }
 };
+
+template struct test_atan<migraphx::shape::float_type>;
+template struct test_atan<migraphx::shape::half_type>;
+template struct test_atan<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/test_atanh.cpp

@@ -27,13 +27,14 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_atanh : verify_program<test_atanh>
+template <migraphx::shape::type_t DType>
+struct test_atanh : verify_program<test_atanh<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape s{migraphx::shape::float_type, {16}};
+        migraphx::shape s{DType, {16}};
         auto x       = mm->add_parameter("x", s);
         auto min_val = mm->add_literal(-0.95f);
         auto max_val = mm->add_literal(0.95f);
@@ -46,3 +47,7 @@ struct test_atanh : verify_program<test_atanh>
         return p;
     }
 };
+
+template struct test_atanh<migraphx::shape::float_type>;
+// template struct test_atanh<migraphx::shape::half_type>;
+// template struct test_atanh<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/test_ceil.cpp

@@ -27,16 +27,21 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_ceil : verify_program<test_ceil>
+template <migraphx::shape::type_t DType>
+struct test_ceil : verify_program<test_ceil<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
 
-        migraphx::shape s{migraphx::shape::double_type, {2, 3, 4, 6}};
+        migraphx::shape s{DType, {2, 3, 4, 6}};
         auto param = mm->add_parameter("x", s);
         mm->add_instruction(migraphx::make_op("ceil"), param);
         return p;
     };
 };
+
+template struct test_ceil<migraphx::shape::float_type>;
+template struct test_ceil<migraphx::shape::half_type>;
+template struct test_ceil<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/test_cos.cpp

@@ -27,15 +27,20 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_cos : verify_program<test_cos>
+template <migraphx::shape::type_t DType>
+struct test_cos : verify_program<test_cos<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape s{migraphx::shape::float_type, {8}};
+        migraphx::shape s{DType, {8}};
         auto x = mm->add_parameter("x", s);
         mm->add_instruction(migraphx::make_op("cos"), x);
         return p;
     }
 };
+
+template struct test_cos<migraphx::shape::float_type>;
+template struct test_cos<migraphx::shape::half_type>;
+template struct test_cos<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/test_cosh.cpp

@@ -27,15 +27,20 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_cosh : verify_program<test_cosh>
+template <migraphx::shape::type_t DType>
+struct test_cosh : verify_program<test_cosh<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape s{migraphx::shape::float_type, {16}};
+        migraphx::shape s{DType, {16}};
         auto x = mm->add_parameter("x", s);
         mm->add_instruction(migraphx::make_op("cosh"), x);
         return p;
     }
 };
+
+template struct test_cosh<migraphx::shape::float_type>;
+template struct test_cosh<migraphx::shape::half_type>;
+template struct test_cosh<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/test_erf.cpp

@@ -27,15 +27,20 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_erf : verify_program<test_erf>
+template <migraphx::shape::type_t DType>
+struct test_erf : verify_program<test_erf<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape s{migraphx::shape::float_type, {2, 3, 4, 6}};
+        migraphx::shape s{DType, {2, 3, 4, 6}};
         auto param = mm->add_parameter("x", s);
         mm->add_instruction(migraphx::make_op("erf"), param);
         return p;
     }
 };
+
+template struct test_erf<migraphx::shape::float_type>;
+template struct test_erf<migraphx::shape::half_type>;
+template struct test_erf<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/test_exp.cpp

@@ -27,15 +27,20 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_exp : verify_program<test_exp>
+template <migraphx::shape::type_t DType>
+struct test_exp : verify_program<test_exp<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape s{migraphx::shape::float_type, {6}};
+        migraphx::shape s{DType, {6}};
         auto x = mm->add_instruction(migraphx::make_op("abs"), mm->add_parameter("x", s));
         mm->add_instruction(migraphx::make_op("exp"), x);
         return p;
     }
 };
+
+template struct test_exp<migraphx::shape::float_type>;
+template struct test_exp<migraphx::shape::half_type>;
+template struct test_exp<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/test_floor.cpp

@@ -27,16 +27,21 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_floor : verify_program<test_floor>
+template <migraphx::shape::type_t DType>
+struct test_floor : verify_program<test_floor<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
 
-        migraphx::shape s{migraphx::shape::float_type, {2, 3, 4, 6}};
+        migraphx::shape s{DType, {2, 3, 4, 6}};
         auto param = mm->add_parameter("x", s);
         mm->add_instruction(migraphx::make_op("floor"), param);
         return p;
     };
 };
+
+template struct test_floor<migraphx::shape::float_type>;
+template struct test_floor<migraphx::shape::half_type>;
+template struct test_floor<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/test_fmod_mod.cpp

@@ -71,3 +71,4 @@ struct test_mod : verify_program<test_mod>
         return p;
     }
 };
+// TODO: check if requires FP8 test
\ No newline at end of file

test/verify/test_gathernd_default.cpp

@@ -27,7 +27,7 @@
 #include <migraphx/make_op.hpp>
 
 template <migraphx::shape::type_t DType>
-struct test_gathernd_default : verify_program<test_gathernd_default, DType>
+struct test_gathernd_default : verify_program<test_gathernd_default<DType>>
 {
     migraphx::program create_program() const
     {

test/verify/test_layernorm.cpp

@@ -117,6 +117,18 @@ struct test_layernorm_fp16 : verify_program<test_layernorm_fp16>
     }
 };
 
+// struct test_layernorm_fp8 : verify_program<test_layernorm_fp8>
+// {
+//     migraphx::program create_program() const
+//     {
+//         migraphx::program p;
+//         auto* mm                 = p.get_main_module();
+//         std::vector<size_t> dims = {1, 24, 64};
+//         auto x = mm->add_parameter("x", migraphx::shape{migraphx::shape::fp8e4m3fnuz_type,
+//         dims}); add_layernorm(*mm, x, dims); return p;
+//     }
+// };
+
 struct test_layernorm_eps : verify_program<test_layernorm_eps>
 {
     migraphx::program create_program() const

test/verify/test_log.cpp

@@ -27,15 +27,20 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_log : verify_program<test_log>
+template <migraphx::shape::type_t DType>
+struct test_log : verify_program<test_log<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape s{migraphx::shape::float_type, {6}};
+        migraphx::shape s{DType, {6}};
         auto x = mm->add_instruction(migraphx::make_op("abs"), mm->add_parameter("x", s));
         mm->add_instruction(migraphx::make_op("log"), x);
         return p;
     }
 };
+
+template struct test_log<migraphx::shape::float_type>;
+template struct test_log<migraphx::shape::half_type>;
+template struct test_log<migraphx::shape::fp8e4m3fnuz_type>;

test/verify/test_min_max.cpp

@@ -46,7 +46,9 @@ struct test_min_max : verify_program<test_min_max<Op, T>>
 template struct test_min_max<migraphx::op::max, migraphx::shape::float_type>;
 template struct test_min_max<migraphx::op::max, migraphx::shape::half_type>;
 template struct test_min_max<migraphx::op::max, migraphx::shape::double_type>;
+template struct test_min_max<migraphx::op::max, migraphx::shape::fp8e4m3fnuz_type>;
 
 template struct test_min_max<migraphx::op::min, migraphx::shape::float_type>;
 template struct test_min_max<migraphx::op::min, migraphx::shape::half_type>;
 template struct test_min_max<migraphx::op::min, migraphx::shape::double_type>;
+template struct test_min_max<migraphx::op::min, migraphx::shape::fp8e4m3fnuz_type>;

test/verify/test_nearbyint.cpp

@@ -22,6 +22,7 @@
  * THE SOFTWARE.
  */
 
+#include "migraphx/float8.hpp"
 #include "verify_program.hpp"
 #include <migraphx/program.hpp>
 #include <migraphx/generate.hpp>
@@ -45,3 +46,4 @@ struct test_nearbyint : verify_program<test_nearbyint<T>>
 
 template struct test_nearbyint<migraphx::half>;
 template struct test_nearbyint<float>;
+template struct test_nearbyint<migraphx::fp8::fp8e4m3fnuz>;

test/verify/test_pad.cpp

@@ -27,13 +27,14 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_pad : verify_program<test_pad>
+template <migraphx::shape::type_t DType>
+struct test_pad : verify_program<test_pad<DType>>
 {
     migraphx::program create_program() const
     {
         migraphx::program p;
         auto* mm = p.get_main_module();
-        migraphx::shape s0{migraphx::shape::int32_type, {1, 96, 165, 165}};
+        migraphx::shape s0{DType, {1, 96, 165, 165}};
         std::vector<int64_t> pads0 = {0, 0, 0, 0, 0, 0, 1, 1};
         std::vector<int64_t> pads1 = {0, 0, 0, 0, 1, 1, 1, 1};
         std::vector<int64_t> pads2 = {1, 1, 1, 1, 0, 0, 0, 0};
@@ -46,3 +47,8 @@ struct test_pad : verify_program<test_pad>
         return p;
     }
 };
+
+template struct test_pad<migraphx::shape::int32_type>;
+template struct test_pad<migraphx::shape::float_type>;
+template struct test_pad<migraphx::shape::half_type>;
+// template struct test_pad<migraphx::shape::fp8e4m3fnuz_type>;