Add vector types and tests

include/ck/utility/data_type.hpp

@@ -13,7 +13,37 @@ using int4_t       = _BitInt(4);
 using f4_t    = unsigned _BitInt(4);
 using f8_t    = _BitInt(8);
 using bf8_t   = unsigned _BitInt(8);
-using e8m0_scale_t = uint8_t;
+
+struct e8m0_scale_t
+{
+    using type = uint8_t;
+    type data;
+    constexpr e8m0_scale_t() : data{type{}} {}
+    constexpr e8m0_scale_t(type init) : data{init} {}
+
+    bool operator==(const e8m0_scale_t& other) const { return (data == other.data); }
+};
+
+struct f4x2_pk_t
+{
+    using type = uint8_t;
+    type data;
+    f4x2_pk_t() : data{type{}} {}
+    f4x2_pk_t(type init) : data{init} {}
+
+    __host__ __device__ inline type unpack(const size_t index)
+    {
+        // if (index == 0) return (data << 4) >> 4;
+        if(index == 0)
+            return data & 0b00001111;
+        return (data >> 4);
+    }
+
+    __host__ __device__ inline type pack(const type x0, const type x1)
+    {
+        return (x1 << 4) | (x0 & 0b00001111);
+    }
+};
 
 inline constexpr auto next_pow2(uint32_t x)
 {
@@ -28,7 +58,7 @@ inline constexpr bool is_native_type()
     return is_same<T, double>::value || is_same<T, float>::value || is_same<T, half_t>::value ||
            is_same<T, bhalf_t>::value || is_same<T, int32_t>::value || is_same<T, int8_t>::value ||
            is_same<T, uint8_t>::value || is_same<T, f8_t>::value || is_same<T, bf8_t>::value ||
-           is_same<T, bool>::value;
+           is_same<T, bool>::value || is_same<T, f4_t>::value;
 }
 
 // vector_type
@@ -1587,136 +1617,6 @@ struct vector_type<T, 64, typename std::enable_if_t<!is_native_type<T>()>>
     }
 };
 
-template <>
-struct vector_type<f4_t, 2, typename std::enable_if_t<!is_native_type<f4_t>()>>
-{
-    using d1_t = f4_t;
-    using d2_t = uint8_t;
-
-    using type = d2_t;
-
-    union alignas(next_pow2(sizeof(type)))
-    {
-        d2_t d2_;
-        StaticallyIndexedArray<d1_t, 2> d1x2_;
-        StaticallyIndexedArray<d2_t, 1> d2x1_;
-    } data_;
-
-    __host__ __device__ constexpr vector_type() : data_{type{}} {}
-
-    __host__ __device__ constexpr vector_type(type v) : data_{v} {}
-
-    template <typename X>
-    __host__ __device__ constexpr const auto& AsType() const
-    {
-        static_assert(is_same<X, d1_t>::value || is_same<X, d2_t>::value,
-                      "Something went wrong, please check src and dst types.");
-
-        if constexpr(is_same<X, d1_t>::value)
-        {
-            return data_.d1x2_;
-        }
-        else if constexpr(is_same<X, d2_t>::value)
-        {
-            return data_.d2x1_;
-        }
-        else
-        {
-            return err;
-        }
-    }
-
-    template <typename X>
-    __host__ __device__ constexpr auto& AsType()
-    {
-        static_assert(is_same<X, d1_t>::value || is_same<X, d2_t>::value,
-                      "Something went wrong, please check src and dst types.");
-
-        if constexpr(is_same<X, d1_t>::value)
-        {
-            return data_.d1x2_;
-        }
-        else if constexpr(is_same<X, d2_t>::value)
-        {
-            return data_.d2x1_;
-        }
-        else
-        {
-            return err;
-        }
-    }
-};
-
-template <>
-struct vector_type<f4_t, 4, typename std::enable_if_t<!is_native_type<f4_t>()>>
-{
-    using d1_t = f4_t;
-    using d2_t = uint8_t;
-    using d4_t = uint16_t;
-
-    using type = d4_t;
-
-    union alignas(next_pow2(sizeof(type)))
-    {
-        d4_t d4_;
-        StaticallyIndexedArray<d1_t, 4> d1x4_;
-        StaticallyIndexedArray<d2_t, 2> d2x2_;
-        StaticallyIndexedArray<d4_t, 1> d4x1_;
-    } data_;
-
-    __host__ __device__ constexpr vector_type() : data_{type{}} {}
-
-    __host__ __device__ constexpr vector_type(type v) : data_{v} {}
-
-    template <typename X>
-    __host__ __device__ constexpr const auto& AsType() const
-    {
-        static_assert(is_same<X, d1_t>::value || is_same<X, d2_t>::value || is_same<X, d4_t>::value,
-                      "Something went wrong, please check src and dst types.");
-
-        if constexpr(is_same<X, d1_t>::value)
-        {
-            return data_.d1x4_;
-        }
-        else if constexpr(is_same<X, d2_t>::value)
-        {
-            return data_.d2x2_;
-        }
-        else if constexpr(is_same<X, d4_t>::value)
-        {
-            return data_.d4x1_;
-        }
-        else
-        {
-            return err;
-        }
-    }
-
-    template <typename X>
-    __host__ __device__ constexpr auto& AsType()
-    {
-        static_assert(is_same<X, d1_t>::value || is_same<X, d2_t>::value || is_same<X, d4_t>::value,
-                      "Something went wrong, please check src and dst types.");
-
-        if constexpr(is_same<X, d1_t>::value)
-        {
-            return data_.d1x4_;
-        }
-        else if constexpr(is_same<X, d2_t>::value)
-        {
-            return data_.d2x2_;
-        }
-        else if constexpr(is_same<X, d4_t>::value)
-        {
-            return data_.d4x1_;
-        }
-        else
-        {
-            return err;
-        }
-    }
-};
-
 using int64_t = long;
 
 // fp64
@@ -1787,6 +1687,14 @@ using uint8x16_t = typename vector_type<uint8_t, 16>::type;
 using uint8x32_t = typename vector_type<uint8_t, 32>::type;
 using uint8x64_t = typename vector_type<uint8_t, 64>::type;
 
+// f4
+using f4x2_t  = typename vector_type<f4x2_pk_t, 1>::type;
+using f4x4_t  = typename vector_type<f4x2_pk_t, 2>::type;
+using f4x8_t  = typename vector_type<f4x2_pk_t, 4>::type;
+using f4x16_t = typename vector_type<f4x2_pk_t, 8>::type;
+using f4x32_t = typename vector_type<f4x2_pk_t, 16>::type;
+using f4x64_t = typename vector_type<f4x2_pk_t, 32>::type;
+
 template <typename T>
 struct NumericLimits
 {

include/ck/utility/e8m0_utils.hpp

@@ -18,4 +18,14 @@ __host__ __device__ inline e8m0_scale_t cast_from_float(float const scale)
     return static_cast<uint8_t>(std::log2(scale) + NumericUtils<e8m0_scale_t>::bias);
 }
 
+template <>
+__host__ __device__ inline int get_exponent_value<e8m0_scale_t>(e8m0_scale_t x)
+{
+    x.data >>= NumericUtils<e8m0_scale_t>::mant;
+
+    x.data &= ((1 << NumericUtils<e8m0_scale_t>::exp) - 1);
+
+    return static_cast<int>(x.data);
+}
+
 } // namespace ck::utils

include/ck/utility/type_convert.hpp

@@ -566,6 +566,18 @@ inline __host__ __device__ float type_convert<float, f4_t>(f4_t data)
 #endif
 }
 
+template <>
+inline __host__ __device__ float type_convert<float, e8m0_scale_t>(e8m0_scale_t scale)
+{
+    return utils::cast_to_float(scale);
+}
+
+template <>
+inline __host__ __device__ e8m0_scale_t type_convert<e8m0_scale_t, float>(float scale)
+{
+    return utils::cast_from_float(scale);
+}
+
 // 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);

test/data_type/test_fp4.cpp

@@ -9,6 +9,7 @@ using ck::e8m0_scale_t;
 using ck::f4_convert_rne;
 using ck::f4_convert_sr;
 using ck::f4_t;
+using ck::f4x2_pk_t;
 using ck::Number;
 using ck::scaled_type_convert;
 using ck::type_convert;
@@ -17,7 +18,7 @@ using ck::vector_type;
 using ck::utils::cast_from_float;
 using ck::utils::cast_to_float;
 
-TEST(FP8, NumericLimits)
+TEST(FP4, NumericLimits)
 {
     // constants given for negative zero nan mode
     EXPECT_EQ(ck::NumericLimits<f4_t>::Min(), f4_t{0x2});
@@ -89,7 +90,7 @@ TEST(FP4, ScaledConvertFP32Nearest)
     float max_fp4 = 6.0f;
     // set maximum scale
     float max_scale = std::pow(2,
-                               ck::NumericLimits<e8m0_scale_t>::Max() -
+                               ck::NumericLimits<e8m0_scale_t>::Max().data -
                                    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
@@ -161,7 +162,7 @@ TEST(FP4, ScaledConvertFP32Stochastic)
     float max_fp4 = 6.0f;
     // set maximum scale
     float max_scale = std::pow(2,
-                               ck::NumericLimits<e8m0_scale_t>::Max() -
+                               ck::NumericLimits<e8m0_scale_t>::Max().data -
                                    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
@@ -225,54 +226,242 @@ TEST(FP4, ScaledConvertFP32Stochastic)
 
 TEST(FP4, TestSize)
 {
-    ASSERT_EQ(1, sizeof(f4_t));
-    ASSERT_EQ(2, sizeof(vector_type<f4_t, 2>));
-    ASSERT_EQ(4, sizeof(vector_type<f4_t, 4>));
+    ASSERT_EQ(1, sizeof(f4x2_pk_t));
+    ASSERT_EQ(1, sizeof(vector_type<f4x2_pk_t, 1>));
+    ASSERT_EQ(2, sizeof(vector_type<f4x2_pk_t, 2>));
+    ASSERT_EQ(4, sizeof(vector_type<f4x2_pk_t, 4>));
+    ASSERT_EQ(8, sizeof(vector_type<f4x2_pk_t, 8>));
+    ASSERT_EQ(16, sizeof(vector_type<f4x2_pk_t, 16>));
+    ASSERT_EQ(32, sizeof(vector_type<f4x2_pk_t, 32>));
 }
 
 TEST(FP4, TestAlignment)
 {
-    ASSERT_EQ(1, alignof(f4_t));
-    ASSERT_EQ(1, alignof(vector_type<f4_t, 2>));
-    ASSERT_EQ(2, alignof(vector_type<f4_t, 4>));
+    ASSERT_EQ(1, alignof(f4x2_pk_t));
+    ASSERT_EQ(1, alignof(vector_type<f4x2_pk_t, 1>));
+    ASSERT_EQ(2, alignof(vector_type<f4x2_pk_t, 2>));
+    ASSERT_EQ(4, alignof(vector_type<f4x2_pk_t, 4>));
+    ASSERT_EQ(8, alignof(vector_type<f4x2_pk_t, 8>));
+    ASSERT_EQ(16, alignof(vector_type<f4x2_pk_t, 16>));
+    ASSERT_EQ(32, alignof(vector_type<f4x2_pk_t, 32>));
 }
 
+// test vector of 1 f4x2_pk_t, contains 2 f4_t
+TEST(FP4, TestAsType1)
+{
+    // test size
+    const int size                        = 1;
+    std::vector<f4x2_pk_t::type> test_vec = {f4x2_pk_t::type{0b0010}, f4x2_pk_t::type{0b1001}};
+    // reference vector
+    vector_type<f4x2_pk_t, size> right_vec;
+    // check default CTOR
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(0), 0);
+        ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(1), 0);
+    });
+    // assign test values to the vector
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        right_vec.template AsType<f4x2_pk_t>()(Number<i>{}) =
+            f4x2_pk_t{}.pack(test_vec.at(i), test_vec.at(i + 1));
+    });
+    // copy the vector
+    vector_type<f4x2_pk_t, size> left_vec{right_vec};
+    // check if values were copied correctly
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(0), test_vec.at(i));
+        ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(1), test_vec.at(i + 1));
+    });
+}
+
+// test vector of 2 f4x2_pk_t, contains 4 f4_t
 TEST(FP4, TestAsType2)
 {
     // test size
     const int size                        = 2;
-    std::vector<f4_t> test_vec = {f4_t{0b0010}, f4_t{0b1001}};
+    std::vector<f4x2_pk_t::type> test_vec = {f4x2_pk_t::type{0b0010},
+                                             f4x2_pk_t::type{0b1001},
+                                             f4x2_pk_t::type{0b0001},
+                                             f4x2_pk_t::type{0b0111}};
     // reference vector
-    vector_type<f4_t, size> right_vec;
+    vector_type<f4x2_pk_t, size> right_vec;
     // check default CTOR
-    ck::static_for<0, size, 1>{}(
-        [&](auto i) { ASSERT_EQ(right_vec.template AsType<f4_t>()(Number<i>{}), 0); });
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(0), 0);
+        ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(1), 0);
+    });
     // assign test values to the vector
-    ck::static_for<0, size, 1>{}(
-        [&](auto i) { right_vec.template AsType<f4_t>()(Number<i>{}) = test_vec.at(i); });
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        right_vec.template AsType<f4x2_pk_t>()(Number<i>{}) =
+            f4x2_pk_t{}.pack(test_vec.at(i), test_vec.at(i + 1));
+    });
     // copy the vector
-    vector_type<f4_t, size> left_vec{right_vec};
+    vector_type<f4x2_pk_t, size> left_vec{right_vec};
     // check if values were copied correctly
-    ck::static_for<0, size, 1>{}(
-        [&](auto i) { ASSERT_EQ(left_vec.template AsType<f4_t>()(Number<i>{}), test_vec.at(i)); });
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(0), test_vec.at(i));
+        ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(1), test_vec.at(i + 1));
+    });
 }
 
+// test vector of 4 f4x2_pk_t, contains 8 f4_t
 TEST(FP4, TestAsType4)
 {
     // test size
     const int size                        = 4;
-    std::vector<f4_t> test_vec = {f4_t{0b0010}, f4_t{0b1001}, f4_t{0b0001}, f4_t{0b0111}};
+    std::vector<f4x2_pk_t::type> test_vec = {f4x2_pk_t::type{0b0010},
+                                             f4x2_pk_t::type{0b1001},
+                                             f4x2_pk_t::type{0b0001},
+                                             f4x2_pk_t::type{0b0111},
+                                             f4x2_pk_t::type{0b1010},
+                                             f4x2_pk_t::type{0b0001},
+                                             f4x2_pk_t::type{0b1001},
+                                             f4x2_pk_t::type{0b1111}};
+    // reference vector
+    vector_type<f4x2_pk_t, size> right_vec;
+    // check default CTOR
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(0), 0);
+        ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(1), 0);
+    });
+    // assign test values to the vector
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        right_vec.template AsType<f4x2_pk_t>()(Number<i>{}) =
+            f4x2_pk_t{}.pack(test_vec.at(i), test_vec.at(i + 1));
+    });
+    // copy the vector
+    vector_type<f4x2_pk_t, size> left_vec{right_vec};
+    // check if values were copied correctly
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(0), test_vec.at(i));
+        ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(1), test_vec.at(i + 1));
+    });
+}
+
+// test vector of 8 f4x2_pk_t, contains 16 f4_t
+TEST(FP4, TestAsType8)
+{
+    // test size
+    const int size                        = 8;
+    std::vector<f4x2_pk_t::type> test_vec = {f4x2_pk_t::type{0b0010},
+                                             f4x2_pk_t::type{0b1001},
+                                             f4x2_pk_t::type{0b0001},
+                                             f4x2_pk_t::type{0b0111},
+                                             f4x2_pk_t::type{0b1010},
+                                             f4x2_pk_t::type{0b0001},
+                                             f4x2_pk_t::type{0b1001},
+                                             f4x2_pk_t::type{0b1111},
+                                             f4x2_pk_t::type{0b0001},
+                                             f4x2_pk_t::type{0b0111},
+                                             f4x2_pk_t::type{0b1010},
+                                             f4x2_pk_t::type{0b0001},
+                                             f4x2_pk_t::type{0b0010},
+                                             f4x2_pk_t::type{0b1001},
+                                             f4x2_pk_t::type{0b1001},
+                                             f4x2_pk_t::type{0b1111}};
+    // reference vector
+    vector_type<f4x2_pk_t, size> right_vec;
+    // check default CTOR
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(0), 0);
+        ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(1), 0);
+    });
+    // assign test values to the vector
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        right_vec.template AsType<f4x2_pk_t>()(Number<i>{}) =
+            f4x2_pk_t{}.pack(test_vec.at(i), test_vec.at(i + 1));
+    });
+    // copy the vector
+    vector_type<f4x2_pk_t, size> left_vec{right_vec};
+    // check if values were copied correctly
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(0), test_vec.at(i));
+        ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(1), test_vec.at(i + 1));
+    });
+}
+
+// test vector of 16 f4x2_pk_t, contains 32 f4_t
+TEST(FP4, TestAsType16)
+{
+    // test size
+    const int size                        = 16;
+    std::vector<f4x2_pk_t::type> test_vec = {
+        f4x2_pk_t::type{0b0010}, f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b0001},
+        f4x2_pk_t::type{0b0111}, f4x2_pk_t::type{0b1010}, f4x2_pk_t::type{0b0001},
+        f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b1111}, f4x2_pk_t::type{0b0001},
+        f4x2_pk_t::type{0b0111}, f4x2_pk_t::type{0b1010}, f4x2_pk_t::type{0b0001},
+        f4x2_pk_t::type{0b0010}, f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b1001},
+        f4x2_pk_t::type{0b1111}, f4x2_pk_t::type{0b0010}, f4x2_pk_t::type{0b1001},
+        f4x2_pk_t::type{0b0001}, f4x2_pk_t::type{0b0111}, f4x2_pk_t::type{0b1010},
+        f4x2_pk_t::type{0b0001}, f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b1111},
+        f4x2_pk_t::type{0b0001}, f4x2_pk_t::type{0b0111}, f4x2_pk_t::type{0b1010},
+        f4x2_pk_t::type{0b0001}, f4x2_pk_t::type{0b0010}, f4x2_pk_t::type{0b1001},
+        f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b1111}};
+    // reference vector
+    vector_type<f4x2_pk_t, size> right_vec;
+    // check default CTOR
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(0), 0);
+        ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(1), 0);
+    });
+    // assign test values to the vector
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        right_vec.template AsType<f4x2_pk_t>()(Number<i>{}) =
+            f4x2_pk_t{}.pack(test_vec.at(i), test_vec.at(i + 1));
+    });
+    // copy the vector
+    vector_type<f4x2_pk_t, size> left_vec{right_vec};
+    // check if values were copied correctly
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(0), test_vec.at(i));
+        ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(1), test_vec.at(i + 1));
+    });
+}
+
+// test vector of 32 f4x2_pk_t, contains 64 f4_t
+TEST(FP4, TestAsType32)
+{
+    // test size
+    const int size                        = 32;
+    std::vector<f4x2_pk_t::type> test_vec = {
+        f4x2_pk_t::type{0b0010}, f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b0001},
+        f4x2_pk_t::type{0b0111}, f4x2_pk_t::type{0b1010}, f4x2_pk_t::type{0b0001},
+        f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b1111}, f4x2_pk_t::type{0b0001},
+        f4x2_pk_t::type{0b0111}, f4x2_pk_t::type{0b1010}, f4x2_pk_t::type{0b0001},
+        f4x2_pk_t::type{0b0010}, f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b1001},
+        f4x2_pk_t::type{0b1111}, f4x2_pk_t::type{0b0010}, f4x2_pk_t::type{0b1001},
+        f4x2_pk_t::type{0b0001}, f4x2_pk_t::type{0b0111}, f4x2_pk_t::type{0b1010},
+        f4x2_pk_t::type{0b0001}, f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b1111},
+        f4x2_pk_t::type{0b0001}, f4x2_pk_t::type{0b0111}, f4x2_pk_t::type{0b1010},
+        f4x2_pk_t::type{0b0001}, f4x2_pk_t::type{0b0010}, f4x2_pk_t::type{0b1001},
+        f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b1111}, f4x2_pk_t::type{0b0010},
+        f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b0001}, f4x2_pk_t::type{0b0111},
+        f4x2_pk_t::type{0b1010}, f4x2_pk_t::type{0b0001}, f4x2_pk_t::type{0b1001},
+        f4x2_pk_t::type{0b1111}, f4x2_pk_t::type{0b0001}, f4x2_pk_t::type{0b0111},
+        f4x2_pk_t::type{0b1010}, f4x2_pk_t::type{0b0001}, f4x2_pk_t::type{0b0010},
+        f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b1111},
+        f4x2_pk_t::type{0b0010}, f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b0001},
+        f4x2_pk_t::type{0b0111}, f4x2_pk_t::type{0b1010}, f4x2_pk_t::type{0b0001},
+        f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b1111}, f4x2_pk_t::type{0b0001},
+        f4x2_pk_t::type{0b0111}, f4x2_pk_t::type{0b1010}, f4x2_pk_t::type{0b0001},
+        f4x2_pk_t::type{0b0010}, f4x2_pk_t::type{0b1001}, f4x2_pk_t::type{0b1001},
+        f4x2_pk_t::type{0b1111}};
     // reference vector
-    vector_type<f4_t, size> right_vec;
+    vector_type<f4x2_pk_t, size> right_vec;
     // check default CTOR
-    ck::static_for<0, size, 1>{}(
-        [&](auto i) { ASSERT_EQ(right_vec.template AsType<f4_t>()(Number<i>{}), 0); });
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(0), 0);
+        ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(1), 0);
+    });
     // assign test values to the vector
-    ck::static_for<0, size, 1>{}(
-        [&](auto i) { right_vec.template AsType<f4_t>()(Number<i>{}) = test_vec.at(i); });
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        right_vec.template AsType<f4x2_pk_t>()(Number<i>{}) =
+            f4x2_pk_t{}.pack(test_vec.at(i), test_vec.at(i + 1));
+    });
     // copy the vector
-    vector_type<f4_t, size> left_vec{right_vec};
+    vector_type<f4x2_pk_t, size> left_vec{right_vec};
     // check if values were copied correctly
-    ck::static_for<0, size, 1>{}(
-        [&](auto i) { ASSERT_EQ(left_vec.template AsType<f4_t>()(Number<i>{}), test_vec.at(i)); });
+    ck::static_for<0, size, 1>{}([&](auto i) {
+        ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(0), test_vec.at(i));
+        ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).unpack(1), test_vec.at(i + 1));
+    });
 }