Add scaled conversions with tests

include/ck/utility/type_convert.hpp

@@ -4,6 +4,7 @@
 #pragma once
 
 #include "ck/utility/data_type.hpp"
+#include "ck/utility/e8m0_utils.hpp"
 #include "ck/utility/f8_utils.hpp"
 #include "ck/utility/mxf4_utils.hpp"
 #include "ck/utility/random_gen.hpp"
@@ -578,6 +579,25 @@ inline __host__ __device__ float type_convert<float, f4_t>(f4_t data)
 #endif
 }
 
+// Declare a template function for scaled conversion
+template <typename Y, typename X>
+__host__ __device__ constexpr Y scaled_type_convert(e8m0_scale_t scale, X x);
+
+// convert fp4 to fp32
+template <>
+inline __host__ __device__ float scaled_type_convert<float, f4_t>(e8m0_scale_t scale, f4_t data)
+{
+#if defined(__gfx94__)
+    // float fval;
+    // uint32_t i32val = static_cast<uint32_t>(x);
+    // fval            = __builtin_amdgcn_cvt_f32_fp8(i32val, 0);
+    // // asm volatile("v_cvt_f32_fp8 %0, %1 src0_sel:BYTE_0" : "=v"(fval) : "v"(i32val));
+    // return fval;
+#else
+    return utils::to_float<f4_t>(scale, data);
+#endif
+}
+
 template <typename Y, typename X, std::size_t NumElems>
 inline __host__ __device__ void array_convert(std::array<Y, NumElems>& y,
                                               const std::array<X, NumElems>& x)

test/data_type/test_fp4.cpp

@@ -5,11 +5,16 @@
 #include "ck/utility/data_type.hpp"
 #include "ck/utility/type_convert.hpp"
 
+using ck::e8m0_scale_t;
 using ck::f4_convert_rne;
 using ck::f4_convert_sr;
 using ck::f4_t;
+using ck::scaled_type_convert;
 using ck::type_convert;
 
+using ck::utils::cast_from_float;
+using ck::utils::cast_to_float;
+
 TEST(FP8, NumericLimits)
 {
     // constants given for negative zero nan mode
@@ -73,3 +78,145 @@ TEST(FP4, ConvertFP32Stochastic)
     neg_float = -0.5f;
     ASSERT_NEAR(neg_float, type_convert<float>(f4_convert_sr(neg_float)), abs_tol);
 }
+
+TEST(FP4, ScaledConvertFP32Nearest)
+{
+    // fix the tolerance value
+    float abs_tol = 1e-6;
+    // set maximum fp4 value
+    float max_fp4 = 6.0f;
+    // set maximum scale
+    float max_scale = std::pow(2,
+                               ck::NumericLimits<e8m0_scale_t>::Max() -
+                                   ck::NumericUtils<e8m0_scale_t>::bias); // 0xFE -> float
+    // set minimum scale
+    float min_scale = std::pow(2, -ck::NumericUtils<e8m0_scale_t>::bias); // 0x00 -> float
+    // set arbitrary scale to 256.0
+    float test_scale = 256.0f; // 0b10000111
+    // convert 0 float to fp4 and back with maximal scale, check if holds
+    ASSERT_NEAR(0.0f,
+                scaled_type_convert<float>(cast_from_float(max_scale), f4_convert_rne(0.0f)),
+                abs_tol);
+    // convert 0 float to fp4 and back with minimal scale, check if holds
+    ASSERT_NEAR(0.0f,
+                scaled_type_convert<float>(cast_from_float(min_scale), f4_convert_rne(0.0f)),
+                abs_tol);
+    // convert maximal f4_t with minimal scale to float and check if equal to minimal float
+    ASSERT_NEAR(ck::NumericLimits<float>::Min(),
+                scaled_type_convert<float>(cast_from_float(min_scale), f4_convert_rne(max_fp4)),
+                abs_tol);
+    // positive norm float value to fp4 and back with various scales, check if holds
+    float pos_float = 1.0f;
+    ASSERT_NEAR(pos_float * test_scale,
+                scaled_type_convert<float>(cast_from_float(test_scale), f4_convert_rne(pos_float)),
+                abs_tol);
+    ASSERT_NEAR(pos_float * max_scale,
+                scaled_type_convert<float>(cast_from_float(max_scale), f4_convert_rne(pos_float)),
+                abs_tol);
+    ASSERT_NEAR(pos_float * min_scale,
+                scaled_type_convert<float>(cast_from_float(min_scale), f4_convert_rne(pos_float)),
+                abs_tol);
+    // negative norm float value to fp4 and back with various scales, check if holds
+    float neg_float = -1.5f;
+    ASSERT_NEAR(neg_float * test_scale,
+                scaled_type_convert<float>(cast_from_float(test_scale), f4_convert_rne(neg_float)),
+                abs_tol);
+    ASSERT_NEAR(neg_float * max_scale,
+                scaled_type_convert<float>(cast_from_float(max_scale), f4_convert_rne(neg_float)),
+                abs_tol);
+    ASSERT_NEAR(neg_float * min_scale,
+                scaled_type_convert<float>(cast_from_float(min_scale), f4_convert_rne(neg_float)),
+                abs_tol);
+    // positive subnorm float value to fp4 and back with various scales, check if holds
+    pos_float = 0.5f;
+    ASSERT_NEAR(pos_float * test_scale,
+                scaled_type_convert<float>(cast_from_float(test_scale), f4_convert_rne(pos_float)),
+                abs_tol);
+    ASSERT_NEAR(pos_float * max_scale,
+                scaled_type_convert<float>(cast_from_float(max_scale), f4_convert_rne(pos_float)),
+                abs_tol);
+    ASSERT_NEAR(pos_float * min_scale,
+                scaled_type_convert<float>(cast_from_float(min_scale), f4_convert_rne(pos_float)),
+                abs_tol);
+    // negative subnorm float value to fp4 and back with various scales, check if holds
+    neg_float = -0.5f;
+    ASSERT_NEAR(neg_float * test_scale,
+                scaled_type_convert<float>(cast_from_float(test_scale), f4_convert_rne(neg_float)),
+                abs_tol);
+    ASSERT_NEAR(neg_float * max_scale,
+                scaled_type_convert<float>(cast_from_float(max_scale), f4_convert_rne(neg_float)),
+                abs_tol);
+    ASSERT_NEAR(neg_float * min_scale,
+                scaled_type_convert<float>(cast_from_float(min_scale), f4_convert_rne(neg_float)),
+                abs_tol);
+}
+
+TEST(FP4, ScaledConvertFP32Stochastic)
+{
+    // fix the tolerance value
+    float abs_tol = 1e-6;
+    // set maximum fp4 value
+    float max_fp4 = 6.0f;
+    // set maximum scale
+    float max_scale = std::pow(2,
+                               ck::NumericLimits<e8m0_scale_t>::Max() -
+                                   ck::NumericUtils<e8m0_scale_t>::bias); // 0xFE -> float
+    // set minimum scale
+    float min_scale = std::pow(2, -ck::NumericUtils<e8m0_scale_t>::bias); // 0x00 -> float
+    // set arbitrary scale to 256.0
+    float test_scale = 256.0f; // 0b10000111
+    // convert 0 float to fp4 and back with maximal scale, check if holds
+    ASSERT_NEAR(
+        0.0f, scaled_type_convert<float>(cast_from_float(max_scale), f4_convert_sr(0.0f)), abs_tol);
+    // convert 0 float to fp4 and back with minimal scale, check if holds
+    ASSERT_NEAR(
+        0.0f, scaled_type_convert<float>(cast_from_float(min_scale), f4_convert_sr(0.0f)), abs_tol);
+    // convert maximal f4_t with minimal scale to float and check if equal to minimal float
+    ASSERT_NEAR(ck::NumericLimits<float>::Min(),
+                scaled_type_convert<float>(cast_from_float(min_scale), f4_convert_sr(max_fp4)),
+                abs_tol);
+    // positive norm float value to fp4 and back with various scales, check if holds
+    float pos_float = 1.0f;
+    ASSERT_NEAR(pos_float * test_scale,
+                scaled_type_convert<float>(cast_from_float(test_scale), f4_convert_sr(pos_float)),
+                abs_tol);
+    ASSERT_NEAR(pos_float * max_scale,
+                scaled_type_convert<float>(cast_from_float(max_scale), f4_convert_sr(pos_float)),
+                abs_tol);
+    ASSERT_NEAR(pos_float * min_scale,
+                scaled_type_convert<float>(cast_from_float(min_scale), f4_convert_sr(pos_float)),
+                abs_tol);
+    // negative norm float value to fp4 and back with various scales, check if holds
+    float neg_float = -1.5f;
+    ASSERT_NEAR(neg_float * test_scale,
+                scaled_type_convert<float>(cast_from_float(test_scale), f4_convert_sr(neg_float)),
+                abs_tol);
+    ASSERT_NEAR(neg_float * max_scale,
+                scaled_type_convert<float>(cast_from_float(max_scale), f4_convert_sr(neg_float)),
+                abs_tol);
+    ASSERT_NEAR(neg_float * min_scale,
+                scaled_type_convert<float>(cast_from_float(min_scale), f4_convert_sr(neg_float)),
+                abs_tol);
+    // positive subnorm float value to fp4 and back with various scales, check if holds
+    pos_float = 0.5f;
+    ASSERT_NEAR(pos_float * test_scale,
+                scaled_type_convert<float>(cast_from_float(test_scale), f4_convert_sr(pos_float)),
+                abs_tol);
+    ASSERT_NEAR(pos_float * max_scale,
+                scaled_type_convert<float>(cast_from_float(max_scale), f4_convert_sr(pos_float)),
+                abs_tol);
+    ASSERT_NEAR(pos_float * min_scale,
+                scaled_type_convert<float>(cast_from_float(min_scale), f4_convert_sr(pos_float)),
+                abs_tol);
+    // negative subnorm float value to fp4 and back with various scales, check if holds
+    neg_float = -0.5f;
+    ASSERT_NEAR(neg_float * test_scale,
+                scaled_type_convert<float>(cast_from_float(test_scale), f4_convert_sr(neg_float)),
+                abs_tol);
+    ASSERT_NEAR(neg_float * max_scale,
+                scaled_type_convert<float>(cast_from_float(max_scale), f4_convert_sr(neg_float)),
+                abs_tol);
+    ASSERT_NEAR(neg_float * min_scale,
+                scaled_type_convert<float>(cast_from_float(min_scale), f4_convert_sr(neg_float)),
+                abs_tol);
+}