Add size checks in pack function

include/ck/utility/data_type.hpp

@@ -62,8 +62,12 @@ struct f6x16_pk_t
         return data_union.f6_array[I];
     }
 
+    // V is vector_size and E is number of elements
+    template <size_t V, size_t E>
     __host__ __device__ inline type pack(f6_t* x)
     {
+        static_assert(V == 1 || V == 2, "Vector size must be 1 or 2.");
+        static_assert(E == 16, "Number of elements must be 16.");
         type* retval = reinterpret_cast<type*>(x);
         return *retval;
     }
@@ -92,8 +96,12 @@ struct f6x32_pk_t
         return data_union.f6_array[I];
     }
 
+    // V is vector_size and E is number of elements
+    template <size_t V, size_t E>
     __host__ __device__ inline type pack(f6_t* x)
     {
+        static_assert(V == 1, "Vector size must be 1.");
+        static_assert(E == 32, "Number of elements must be 16.");
         type* retval = reinterpret_cast<type*>(x);
         return *retval;
     }
@@ -122,8 +130,12 @@ struct bf6x16_pk_t
         return data_union.bf6_array[I];
     }
 
+    // V is vector_size and E is number of elements
+    template <size_t V, size_t E>
     __host__ __device__ inline type pack(bf6_t* x)
     {
+        static_assert(V == 1 || V == 2, "Vector size must be 1 or 2.");
+        static_assert(E == 16, "Number of elements must be 16.");
         type* retval = reinterpret_cast<type*>(x);
         return *retval;
     }
@@ -152,8 +164,12 @@ struct bf6x32_pk_t
         return data_union.bf6_array[I];
     }
 
+    // V is vector_size and E is number of elements
+    template <size_t V, size_t E>
     __host__ __device__ inline type pack(bf6_t* x)
     {
+        static_assert(V == 1, "Vector size must be 1.");
+        static_assert(E == 32, "Number of elements must be 16.");
         type* retval = reinterpret_cast<type*>(x);
         return *retval;
     }

test/data_type/test_bf6.cpp

@@ -271,7 +271,8 @@ TEST(BF6, TestAsType16x1)
     });
     // assign test values to the vector
     ck::static_for<0, vector_size, 1>{}([&](auto i) {
-        right_vec.template AsType<bf6x16_pk_t>()(Number<i>{}) = bf6x16_pk_t{}.pack(test_vec.data());
+        right_vec.template AsType<bf6x16_pk_t>()(Number<i>{}) =
+            bf6x16_pk_t{}.pack<vector_size, packed_size>(test_vec.data());
     });
     // copy the vector
     vector_type<bf6x16_pk_t, vector_size> left_vec{right_vec};
@@ -310,7 +311,7 @@ TEST(BF6, TestAsType16x2)
     // assign test values to the vector
     ck::static_for<0, vector_size, 1>{}([&](auto i) {
         right_vec.template AsType<bf6x16_pk_t>()(Number<i>{}) =
-            bf6x16_pk_t{}.pack(test_vec.data() + i * packed_size);
+            bf6x16_pk_t{}.pack<vector_size, packed_size>(test_vec.data() + i * packed_size);
     });
     // copy the vector
     vector_type<bf6x16_pk_t, vector_size> left_vec{right_vec};
@@ -348,7 +349,8 @@ TEST(BF6, TestAsType32x1)
     });
     // assign test values to the vector
     ck::static_for<0, vector_size, 1>{}([&](auto i) {
-        right_vec.template AsType<bf6x32_pk_t>()(Number<i>{}) = bf6x32_pk_t{}.pack(test_vec.data());
+        right_vec.template AsType<bf6x32_pk_t>()(Number<i>{}) =
+            bf6x32_pk_t{}.pack<vector_size, packed_size>(test_vec.data());
     });
     // copy the vector
     vector_type<bf6x32_pk_t, vector_size> left_vec{right_vec};

test/data_type/test_fp6.cpp

@@ -269,7 +269,8 @@ TEST(FP6, TestAsType16x1)
     });
     // assign test values to the vector
     ck::static_for<0, vector_size, 1>{}([&](auto i) {
-        right_vec.template AsType<f6x16_pk_t>()(Number<i>{}) = f6x16_pk_t{}.pack(test_vec.data());
+        right_vec.template AsType<f6x16_pk_t>()(Number<i>{}) =
+            f6x16_pk_t{}.pack<vector_size, packed_size>(test_vec.data());
     });
     // copy the vector
     vector_type<f6x16_pk_t, vector_size> left_vec{right_vec};
@@ -308,7 +309,7 @@ TEST(FP6, TestAsType16x2)
     // assign test values to the vector
     ck::static_for<0, vector_size, 1>{}([&](auto i) {
         right_vec.template AsType<f6x16_pk_t>()(Number<i>{}) =
-            f6x16_pk_t{}.pack(test_vec.data() + i * packed_size);
+            f6x16_pk_t{}.pack<vector_size, packed_size>(test_vec.data() + i * packed_size);
     });
     // copy the vector
     vector_type<f6x16_pk_t, vector_size> left_vec{right_vec};
@@ -345,7 +346,8 @@ TEST(FP6, TestAsType32x1)
     });
     // assign test values to the vector
     ck::static_for<0, vector_size, 1>{}([&](auto i) {
-        right_vec.template AsType<f6x32_pk_t>()(Number<i>{}) = f6x32_pk_t{}.pack(test_vec.data());
+        right_vec.template AsType<f6x32_pk_t>()(Number<i>{}) =
+            f6x32_pk_t{}.pack<vector_size, packed_size>(test_vec.data());
     });
     // copy the vector
     vector_type<f6x32_pk_t, vector_size> left_vec{right_vec};