Add fp4 type with constants

include/ck/utility/data_type.hpp

@@ -10,6 +10,7 @@ namespace ck {
 using bhalf_t = ushort;
 using half_t  = _Float16;
 using int4_t  = _BitInt(4);
+using f4_t    = unsigned _BitInt(4);
 using f8_t    = _BitInt(8);
 using bf8_t   = unsigned _BitInt(8);
 
@@ -1152,6 +1153,22 @@ struct NumericLimits<bf8_t>
     __host__ __device__ static constexpr bf8_t QuietNaN() { return bf8_t(binary_qnan); }
 };
 
+template <>
+struct NumericLimits<f4_t>
+{
+    static constexpr uint8_t binary_min_normal    = 0x2; // 0b0010
+    static constexpr uint8_t binary_max_normal    = 0x7; // 0b0111
+    static constexpr uint8_t binary_lowest_normal = 0xF; // 0b1111
+    static constexpr uint8_t binary_min_subnorm   = 0x1; // 0b0001
+    static constexpr uint8_t binary_max_subnorm   = 0x1; // 0b0001
+
+    __host__ __device__ static constexpr f4_t Min() { return f4_t(binary_min_normal); }
+    __host__ __device__ static constexpr f4_t Max() { return f4_t(binary_max_normal); }
+    __host__ __device__ static constexpr f4_t Lowest() { return f4_t(binary_lowest_normal); }
+    __host__ __device__ static constexpr f4_t MinSubnorm() { return f4_t(binary_min_subnorm); }
+    __host__ __device__ static constexpr f4_t MaxSubnorm() { return f4_t(binary_max_subnorm); }
+};
+
 template <typename T>
 struct NumericUtils
 {
@@ -1208,4 +1225,12 @@ struct NumericUtils<bf8_t>
     static constexpr int bias = 16; // negative zero nan mode
     // static constexpr int bias = 15; // ieee mode
 };
+
+template <>
+struct NumericUtils<f4_t>
+{
+    static constexpr int exp  = 2;
+    static constexpr int mant = 1;
+    static constexpr int bias = 1;
+};
 } // namespace ck

test/data_type/CMakeLists.txt

@@ -17,5 +17,9 @@ add_gtest_executable(test_bf8 test_bf8.cpp)
 if(result EQUAL 0)
   target_link_libraries(test_bf8 PRIVATE utility)
 endif()
+add_gtest_executable(test_fp4 test_fp4.cpp)
+if(result EQUAL 0)
+  target_link_libraries(test_fp4 PRIVATE utility)
+endif()
 
 add_gtest_executable(test_type_convert_const type_convert_const.cpp)

test/data_type/test_fp4.cpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "gtest/gtest.h"
+#include "ck/utility/data_type.hpp"
+#include "ck/utility/type_convert.hpp"
+
+using ck::f4_t;
+using ck::type_convert;
+
+TEST(FP8, NumericLimits)
+{
+    // constants given for negative zero nan mode
+    EXPECT_EQ(ck::NumericLimits<f4_t>::Min(), f4_t{0x2});
+    EXPECT_EQ(ck::NumericLimits<f4_t>::Max(), f4_t{0x7});
+    EXPECT_EQ(ck::NumericLimits<f4_t>::Lowest(), f4_t{0xF});
+    EXPECT_EQ(ck::NumericLimits<f4_t>::MinSubnorm(), f4_t{0x1});
+    EXPECT_EQ(ck::NumericLimits<f4_t>::MaxSubnorm(), f4_t{0x1});
+}