Implement ConvertFP32Nearest test.

include/ck/utility/data_type.hpp

@@ -728,11 +728,11 @@ struct f8_ocp_t
     using data_type = fp8_storage_t;
     data_type data;
 
-    constexpr static internal::ck_saturation_t default_saturation        = internal::CK_SATFINITE;
-    constexpr static internal::ck_fp8_interpretation_t default_interpret = internal::CK_E4M3_OCP;
+    static constexpr internal::ck_saturation_t default_saturation        = internal::CK_SATFINITE;
+    static constexpr internal::ck_fp8_interpretation_t default_interpret = internal::CK_E4M3_OCP;
 
-    constexpr static unsigned int we = 4; // exponent width
-    constexpr static unsigned int wm = 3; // mantissa width
+    static constexpr unsigned int we = 4; // exponent width
+    static constexpr unsigned int wm = 3; // mantissa width
 
     __host__ __device__ constexpr bool operator==(const f8_ocp_t& other) const
     {
@@ -773,7 +773,33 @@ struct bf8_ocp_t
 {
     using data_type = fp8_storage_t;
     data_type data;
-};
+
+    static constexpr internal::ck_saturation_t default_saturation        = internal::CK_SATFINITE;
+    static constexpr internal::ck_fp8_interpretation_t default_interpret = internal::CK_E5M2_OCP;
+
+    static constexpr unsigned int we = 5; // exponent width
+    static constexpr unsigned int wm = 2; // mantissa width
+
+    __host__ __device__ constexpr bool operator==(const bf8_ocp_t& other) const
+    {
+        return (data == other.data) && (internal::ocp_bf8_is_nan(data) == false); // NaN != NaN
+    }
+
+#if CK_USE_OCP_FP8
+    __host__ __device__ explicit operator float() const {
+#else
+    __host__ explicit operator float() const
+    {
+#endif
+#if CK_FP8_CVT_FAST_PATH
+        return internal::cast_to_f32_from_f8<default_interpret>(this->data);
+#else
+        return internal::cast_from_f8<float, wm, we, false>(
+            this->data); // XXX: clip==false must be consistent with operator half_t
+#endif
+}
+}
+;
 
 namespace internal {
 template <typename T,
@@ -802,6 +828,15 @@ inline __host__ __device__ f8_ocp_t f8_convert_rne<f8_ocp_t, float>(float x)
     return f8_ocp_t{
         internal::cvt_float_to_fp8<f8_ocp_t::default_interpret, f8_ocp_t::default_saturation>(x)};
 }
+
+// convert fp32 to bf8 with rounding to nearest even
+template <>
+inline __host__ __device__ bf8_ocp_t f8_convert_rne<bf8_ocp_t, float>(float x)
+{
+    return bf8_ocp_t{
+        internal::cvt_float_to_fp8<bf8_ocp_t::default_interpret, bf8_ocp_t::default_saturation>(x)};
+}
+
 // convert half_t to fp8 with rounding to nearest even
 template <>
 inline __host__ __device__ f8_ocp_t f8_convert_rne<f8_ocp_t, half_t>(half_t x)

include/ck/utility/type_convert.hpp

@@ -106,6 +106,12 @@ inline __host__ __device__ constexpr f8_ocp_t type_convert<f8_ocp_t, int>(int x)
     return f8_ocp_t{type_convert<f8_ocp_t::data_type>(x)};
 }
 
+template <>
+inline __host__ __device__ constexpr bf8_ocp_t type_convert<bf8_ocp_t, int>(int x)
+{
+    return bf8_ocp_t{type_convert<bf8_ocp_t::data_type>(x)};
+}
+
 // Convert X to Y
 template <typename Y, typename X>
 __host__ __device__ constexpr Y type_convert_sp(X x)

test/data_type/test_bf8_ocp.cpp

@@ -11,12 +11,79 @@ using ck::f8_convert_sr;
 using ck::half_t;
 using ck::type_convert;
 
-TEST(BF8OCP, NumericLimits) {}
+TEST(BF8OCP, NumericLimits)
+{ // constants given for OCP FP8
+    EXPECT_EQ(ck::NumericLimits<bf8_ocp_t>::Min(),
+              type_convert<bf8_ocp_t>(0x04)); // 0b00000100 = 2^-14
+    EXPECT_EQ(ck::NumericLimits<bf8_ocp_t>::Max(),
+              type_convert<bf8_ocp_t>(0x7B)); // 0b01111011 = 57344
+    EXPECT_EQ(ck::NumericLimits<bf8_ocp_t>::Lowest(),
+              type_convert<bf8_ocp_t>(0xFB)); // 0b11111011 = -57344
+    EXPECT_EQ(ck::NumericLimits<bf8_ocp_t>::QuietNaN().data,
+              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::internal::fp8_is_inf(type_convert<bf8_ocp_t>(0xFC)) &&
+                ck::internal::fp8_is_inf(type_convert<bf8_ocp_t>(0x7C)));
+}
 
-TEST(BF8OCP, ConvertFP32Nearest) {}
+TEST(BF8OCP, ConvertFP32Nearest)
+{
+    // fix the tolerance value
+    float abs_tol = 1e-6;
 
-TEST(BF8OCP, ConvertFP32Stochastic) {}
+    // convert 0 float to bfp8 and back, check if holds
+    ASSERT_NEAR(0.0f, type_convert<float>(f8_convert_rne<bf8_ocp_t>(0.0f)), 0.0f);
 
-TEST(BF8OCP, ConvertFP16Nearest) {}
+    // convert minimal float to bf8 and back, check if holds
+    ASSERT_NEAR(std::numeric_limits<float>::min(),
+                type_convert<float>(f8_convert_rne<bf8_ocp_t>(std::numeric_limits<float>::min())),
+                abs_tol);
 
-TEST(BF8OCP, ConvertFP16Stochastic) {}
+    const auto max_bf8_t_float = type_convert<float>(ck::NumericLimits<bf8_ocp_t>::Max());
+
+    // convert maximal bf8_ocp_t to float and check if equal to bf8 max
+    ASSERT_NEAR(
+        max_bf8_t_float, type_convert<float>(f8_convert_rne<bf8_ocp_t>(max_bf8_t_float)), 0.0f);
+
+    // convert maximal float to bf8 and back, check if clipped to bf8 max (saturation to finite)
+    ASSERT_NEAR(max_bf8_t_float,
+                type_convert<float>(f8_convert_rne<bf8_ocp_t>(std::numeric_limits<float>::max())),
+                0.0f);
+
+    // convert float infinity to bf8_ocp_t and check if it is max value (saturation to finite)
+    ASSERT_EQ(ck::NumericLimits<bf8_ocp_t>::Max(),
+              f8_convert_rne<bf8_ocp_t>(std::numeric_limits<float>::infinity()));
+
+    // positive normal float value to bf8 and back, check if holds
+    float pos_float = 0.0000762939f; // 10*2^-17
+    ASSERT_NEAR(pos_float, type_convert<float>(f8_convert_rne<bf8_ocp_t>(pos_float)), abs_tol);
+
+    // negative smallest normal bf8 value to bf8 and back, check if holds
+    constexpr auto neg_min_bf8 = -0.00006103515625f; //-2^-14
+    ASSERT_NEAR(neg_min_bf8, type_convert<float>(f8_convert_rne<bf8_ocp_t>(neg_min_bf8)), 0.0f);
+
+    // positive subnorm float value to bf8 and back, check if holds
+    constexpr auto pos_subnorm_bf8 = 0.000030517578125f; // 2^-15
+    ASSERT_NEAR(
+        pos_subnorm_bf8, type_convert<float>(f8_convert_rne<bf8_ocp_t>(pos_subnorm_bf8)), 0.0f);
+
+    // min subnorm bf8 value to bf8 and back, check if holds
+    constexpr auto min_subnorm_bf8 = -0.0000152587890625f; //-2^-16
+    ASSERT_NEAR(
+        min_subnorm_bf8, type_convert<float>(f8_convert_rne<bf8_ocp_t>(min_subnorm_bf8)), 0.0f);
+
+    // smaller than min subnorm bf8 value to bf8 must be zero
+    constexpr auto less_than_min_subnorm = 0.00000762939453125f; // 2^-17
+    ASSERT_EQ(0.0f, type_convert<float>(f8_convert_rne<bf8_ocp_t>(less_than_min_subnorm)));
+
+    // convert quiet NaN to bf8_ocp_t and check if it is quiet NaN
+    const auto bf8_nan = f8_convert_rne<bf8_ocp_t>(std::numeric_limits<float>::quiet_NaN());
+    ASSERT_TRUE(ck::internal::ocp_bf8_is_nan(bf8_nan.data));
+}
+
+TEST(BF8OCP, ConvertFP32Stochastic) { ASSERT_TRUE(false) << "Not implemented"; }
+
+TEST(BF8OCP, ConvertFP16Nearest) { ASSERT_TRUE(false) << "Not implemented"; }
+
+TEST(BF8OCP, ConvertFP16Stochastic) { ASSERT_TRUE(false) << "Not implemented"; }