Fixes

src/include/migraphx/migraphx_float8.hpp

@@ -474,10 +474,12 @@ template <>
 class NumericLimits<migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::fp8>>
 {
     public:
+    // TODO :figure out epsilon in Hex to make it constexpr
     static constexpr MIGRAPHX_HIP_HOST_DEVICE migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::fp8>
     epsilon()
     {
-        return static_cast<migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::fp8>>(float(0.0625));
+        return migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::fp8>(
+            0x28, migraphx_fp8::hip_f8<>::from_bits());
     }
 
     static constexpr MIGRAPHX_HIP_HOST_DEVICE migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::fp8>
@@ -493,13 +495,6 @@ class NumericLimits<migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::fp8>>
         return migraphx_fp8::F8_Max<migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::fp8>>();
     }
 
-    static constexpr MIGRAPHX_HIP_HOST_DEVICE migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::fp8>
-    min()
-    {
-        return static_cast<migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::fp8>>(-1.0f) *
-               migraphx_fp8::F8_Max<migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::fp8>>();
-    }
-
     static constexpr MIGRAPHX_HIP_HOST_DEVICE migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::fp8>
     lowest()
     {
@@ -521,7 +516,8 @@ class NumericLimits<migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::bf8>>
     static constexpr MIGRAPHX_HIP_HOST_DEVICE migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::bf8>
     epsilon()
     {
-        return static_cast<migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::bf8>>(float(0.125));
+        return migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::bf8>(
+            0x34, migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::bf8>::from_bits());
     }
 
     static constexpr MIGRAPHX_HIP_HOST_DEVICE migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::bf8>
@@ -538,12 +534,6 @@ class NumericLimits<migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::bf8>>
         return static_cast<migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::bf8>>(
             migraphx_fp8::F8_Max<migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::bf8>>());
     }
-    static constexpr MIGRAPHX_HIP_HOST_DEVICE migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::bf8>
-    min()
-    {
-        return static_cast<migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::bf8>>(float(-1.0f)) *
-               migraphx_fp8::F8_Max<migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::bf8>>();
-    }
 
     static constexpr MIGRAPHX_HIP_HOST_DEVICE migraphx_fp8::hip_f8<migraphx_fp8::hip_f8_type::bf8>
     lowest()

src/include/migraphx/requires.hpp

@@ -38,9 +38,6 @@ struct and_ : std::is_same<and_<Bs...>, and_<(Bs or true)...>> // NOLINT
 template <bool B>
 using bool_c = std::integral_constant<bool, B>;
 
-template <class From, class To>
-using is_convertible = std::is_convertible<From, To>;
-
 #define MIGRAPHX_REQUIRES_PRIMITIVE_CAT(x, y) x##y
 #define MIGRAPHX_REQUIRES_CAT(x, y) MIGRAPHX_REQUIRES_PRIMITIVE_CAT(x, y)
 

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

@@ -36,7 +36,8 @@ namespace migraphx {
 template <class F, class T, class... Ts>
 __device__ void pointwise_tensor(index idx, F f, T out, Ts... xs)
 {
-    idx.global_stride(out.get_shape().elements(), [&](auto i) { out[i] = f(xs[i]...); });
+    idx.global_stride(out.get_shape().elements(),
+                      [&](auto i) { out[i] = implicit_conversion(f(xs[i]...)); });
 }
 
 template <class... Transforms>

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

@@ -244,8 +244,9 @@ 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];
+            });
         }
     }
 

test/py/test_shape.py

@@ -81,7 +81,7 @@ def test_create_dyn_shape():
 
 def test_type_enum():
     mgx_types = [
-        'bool_type', 'double_type', 'float_type', 'half_type', 'float_type', 'int16_type',
+        'bool_type', 'double_type', 'float_type', 'half_type', 'float8_type', 'int16_type',
         'int32_type', 'int64_type', 'int8_type', 'uint16_type', 'uint32_type',
         'uint64_type', 'uint8_type'
     ]

test/verify/test_literal_limits.cpp

@@ -22,10 +22,10 @@
  * THE SOFTWARE.
  */
 
-#include "migraphx/migraphx_float8.hpp"
 #include "verify_program.hpp"
 #include <migraphx/program.hpp>
 #include <migraphx/make_op.hpp>
+#include <migraphx/migraphx_float8.hpp>
 #include <limits>
 
 template <migraphx::shape::type_t Q, typename T>