Provide single point of truth for FP8 INF and NAN checks

include/ck/utility/amd_ck_fp8.hpp

@@ -488,7 +488,30 @@ struct non_native_vector_base<bf8_ocp_t, 1>
     __host__ __device__ operator bf8_ocp_t() const { return bf8_ocp_t{d[0]}; }
 };
 
-namespace fp8_impl {
+template <typename T>
+__host__ __device__ static inline constexpr bool fp8_is_nan(T);
+
+template <>
+__host__ __device__ inline constexpr bool fp8_is_nan(f8_ocp_t a)
+{
+    return fp8_impl::ocp_f8_is_nan(a.data);
+}
+template <>
+__host__ __device__ inline constexpr bool fp8_is_nan(bf8_ocp_t a)
+{
+    return fp8_impl::ocp_bf8_is_nan(a.data);
+}
+template <>
+__host__ __device__ inline constexpr bool fp8_is_nan(f8_fnuz_t a)
+{
+    return fp8_impl::fnuz_f8_is_nan(a);
+}
+template <>
+__host__ __device__ inline constexpr bool fp8_is_nan(bf8_fnuz_t a)
+{
+    return fp8_impl::fnuz_bf8_is_nan(a);
+}
+
 template <typename T,
           std::enable_if_t<std::is_same_v<T, bf8_ocp_t> || std::is_same_v<T, f8_ocp_t> ||
                                std::is_same_v<T, bf8_fnuz_t> || std::is_same_v<T, f8_fnuz_t>,
@@ -503,6 +526,8 @@ __host__ __device__ inline constexpr bool fp8_is_inf(bf8_ocp_t a)
     return (a.data & 0x7f) == 0x7c;
 }
 
+namespace fp8_impl {
+
 // Assertions to check for supported conversion types
 #define __assert_ocp_support(interp)                                               \
     {                                                                              \

include/ck/utility/math_v2.hpp

@@ -80,7 +80,7 @@ static inline __host__ bool isnan(half_t x)
     return (xx & 0x7FFF) > 0x7C00;
 };
 
-static inline __host__ bool isnan(f8_t x) { return (x & 0x80); };
+static inline __host__ bool isnan(f8_t x) { return ck::fp8_is_nan(x); };
 
 #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
 static inline __host__ bool isnan(int4_t x)
@@ -531,7 +531,7 @@ static inline __device__ bool isnan(half_t x)
     return (xx & 0x7FFF) > 0x7C00;
 };
 
-static inline __device__ bool isnan(f8_t x) { return (x & 0x80); };
+static inline __device__ bool isnan(f8_t x) { return ck::fp8_is_nan(x); };
 
 static inline __device__ half_t sqrt(half_t x)
 {

test/data_type/test_bf8_ocp.cpp

@@ -23,8 +23,8 @@ TEST(BF8OCP, NumericLimits)
               type_convert<bf8_ocp_t>(0x7D).data); // 0b01111101
     EXPECT_FALSE(ck::NumericLimits<bf8_ocp_t>::QuietNaN() ==
                  ck::NumericLimits<bf8_ocp_t>::QuietNaN());
-    EXPECT_TRUE(ck::fp8_impl::fp8_is_inf(type_convert<bf8_ocp_t>(0xFC)) &&
-                ck::fp8_impl::fp8_is_inf(type_convert<bf8_ocp_t>(0x7C)));
+    EXPECT_TRUE(ck::fp8_is_inf(type_convert<bf8_ocp_t>(0xFC)) &&
+                ck::fp8_is_inf(type_convert<bf8_ocp_t>(0x7C)));
 }
 
 TEST(BF8OCP, ConvertFP32Nearest)