Fix tidy and other errors

src/targets/cpu/dnnl.cpp

@@ -68,6 +68,7 @@ dnnl::memory::data_type to_dnnl_memory_data_type(shape::type_t t)
     case st::int32_type: return dt::s32;
     case st::int8_type: return dt::s8;
     case st::uint8_type: return dt::u8;
+    case st::fp8e4m3fnuz_type: return dt : u8;
     default: MIGRAPHX_THROW("Unsupported data type");
     }
 }

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

@@ -29,7 +29,7 @@
 #define MIGRAPHX_HIP_DEVICE __device__
 
 // We are clipping in down conversion by default
-#define MIGRAPHX_F8_DOWNCAST_CLIPPING 1
+#define MIGRAPHX_F8_DOWNCAST_CLIPPING 1 // NOLINT
 
 #include <migraphx/kernels/types.hpp>
 #include <migraphx/kernels/float8_impl.hpp>
@@ -178,7 +178,7 @@ struct float8
 
     // convert to float
 // #if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
-#if 0 // need constexpr operator(). This version can't be constexpr
+#if 0 // need constexpr operator(). This version can't be constexpr // NOLINT
     // upcast using device specific intrinsic
     inline MIGRAPHX_HIP_DEVICE operator float() const
     {
@@ -264,6 +264,7 @@ struct float8
         }
     }
 
+// NOLINTNEXTLINE
 #define MIGRAPHX_FP8_SHORT_UNARY_OP(unary_op, binary_op)                              \
     constexpr float8& MIGRAPHX_HIP_DEVICE operator unary_op(const float8& rhs)        \
     {                                                                                 \
@@ -324,13 +325,14 @@ using fp8e5m2fnuz = float8<migraphx::fp8::f8_type::bf8, true>;
     }
 
 // NOLINTNEXTLINE
-#define MIGRAPHX_FP8_UNARY_OP(unary_op, T)               \
-    inline constexpr MIGRAPHX_HIP_DEVICE T unary_op(T v) \
+#define MIGRAPHX_FP8_FABS(T)                         \
+    inline constexpr MIGRAPHX_HIP_DEVICE T fabs(T v) \
     {                                                \
         v.data = v.data & 0x7f;                      \
         return v;                                    \
     }
 
+// NOLINTNEXTLINE
 #define MIGRAPHX_FP8_GEN_OP_OVERLOADS(T) \
     MIGRAPHX_FP8_BINARY_OP(*, T, T)      \
     MIGRAPHX_FP8_BINARY_OP(-, T, T)      \
@@ -342,7 +344,7 @@ using fp8e5m2fnuz = float8<migraphx::fp8::f8_type::bf8, true>;
     MIGRAPHX_FP8_BINARY_OP(>, T, bool)   \
     MIGRAPHX_FP8_BINARY_OP(<, T, bool)   \
     MIGRAPHX_FP8_BINARY_OP(!=, T, bool)  \
-    MIGRAPHX_FP8_UNARY_OP(fabs, T)
+    MIGRAPHX_FP8_FABS(T)
 
 MIGRAPHX_FP8_GEN_OP_OVERLOADS(fp8e5m2)
 MIGRAPHX_FP8_GEN_OP_OVERLOADS(fp8e5m2fnuz)
@@ -453,10 +455,10 @@ class numeric_limits<fp8e5m2>
         return fp8e5m2(0x34, fp8e5m2::from_bits());
     }
     // 7D, 7E, 7F are positive NaNs and FD, FE, FF are negative NaNs
-    static constexpr MIGRAPHX_HIP_DEVICE fp8e5m2 quiet_NaN()
+    static constexpr MIGRAPHX_HIP_DEVICE fp8e5m2 quiet_NaN() // NOLINT
     {
         return fp8e5m2(0xFF, fp8e5m2::from_bits());
-    } // NOLINT
+    }
 
     static constexpr MIGRAPHX_HIP_DEVICE fp8e5m2 max()
     {

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

@@ -47,6 +47,7 @@ struct conditional<false, T, F>
 namespace fp8 {
 namespace impl {
 
+// NOLINTBEGIN
 template <int Wm, int We, typename T, bool NegativeZeroNan, bool Clip>
 __device__ constexpr uint8_t cast_to_f8(T f_x, bool stoch = false, uint32_t rng = 0)
 {
@@ -256,6 +257,7 @@ __device__ constexpr uint8_t cast_to_f8(T f_x, bool stoch = false, uint32_t rng
     mantissa &= (1 << Wm) - 1;
     return (sign << 7) | (f8_exponent << Wm) | mantissa;
 }
+// NOLINTEND
 
 template <int Wm, int We, typename T, bool NegativeZeroNan>
 __device__ constexpr T cast_from_f8(uint8_t x)

test/verify/test_literal_limits.cpp

@@ -26,6 +26,7 @@
 #include <migraphx/program.hpp>
 #include <migraphx/make_op.hpp>
 #include <limits>
+#include <type_traits>
 
 template <migraphx::shape::type_t Q, typename T>
 struct test_literal_limits : verify_program<test_literal_limits<Q, T>>
@@ -36,10 +37,14 @@ struct test_literal_limits : verify_program<test_literal_limits<Q, T>>
         auto* mm          = p.get_main_module();
         auto input_s      = migraphx::shape(Q, {3, 1});
         auto infinity_val = std::numeric_limits<T>::max();
-        if constexpr(std::numeric_limits<T>::has_infinity)
+        if constexpr(std::numeric_limits<T>::has_infinity and std::is_floating_point<T>{})
         {
             infinity_val = std::numeric_limits<T>::infinity();
         }
+        else
+        { // for the interger vals, infinity doesn't exist
+            infinity_val = 0;
+        }
         std::vector<T> s_data{
             infinity_val, static_cast<T>(-infinity_val), std::numeric_limits<T>::quiet_NaN()};