Skip to content

GitLab

Commit df1bad99 authored by Rostyslav Geyyer's avatar Rostyslav Geyyer
Browse files

Update tests and pack functions

parent 544aad11
Show whitespace changes
Inline Side-by-side
Showing with 156 additions and 122 deletions
include/ck/utility/data_type.hpp
View file @ df1bad99
......@@ -45,6 +45,7 @@ struct f6x16_pk_t
using element_type = uint32_t;
using type = StaticallyIndexedArray_v2<element_type, 3>;
type data;
typedef int8_t test_vec_t __attribute__((ext_vector_type(16)));
f6x16_pk_t() : data{type{}} {}
f6x16_pk_t(type init) : data{init} {}
......@@ -71,18 +72,13 @@ struct f6x16_pk_t
return static_cast<f6_t>(bits & 0x3F);
}
// 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)
__host__ __device__ inline type pack(const test_vec_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 packed{};
// for each of the 16 f6_t values, place its 6 bits in the correct position
ck::static_for<0, 16, 1>{}([&](auto i) {
uint32_t bits = static_cast<uint32_t>(x[i]) & 0x3F;
uint32_t bits = static_cast<uint32_t>(x[static_cast<int>(i)]) & 0x3F;
constexpr int num_bits_elem = 6;
constexpr int num_bits_vec_elem = 32;
constexpr int vector_size = 3;
......@@ -115,6 +111,7 @@ struct f6x32_pk_t
using element_type = uint32_t;
using type = StaticallyIndexedArray_v2<element_type, 6>;
type data;
typedef int8_t test_vec_t __attribute__((ext_vector_type(32)));
f6x32_pk_t() : data{type{}} {}
f6x32_pk_t(type init) : data{init} {}
......@@ -141,18 +138,13 @@ struct f6x32_pk_t
return static_cast<f6_t>(bits & 0x3F);
}
// 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)
__host__ __device__ inline type pack(const test_vec_t& x)
{
static_assert(V == 1, "Vector size must be 1.");
static_assert(E == 32, "Number of elements must be 32.");
type packed{};
// for each of the 32 f6_t values, place its 6 bits in the correct position
ck::static_for<0, 32, 1>{}([&](auto i) {
uint32_t bits = static_cast<uint32_t>(x[i]) & 0x3F;
uint32_t bits = static_cast<uint32_t>(x[static_cast<int>(i)]) & 0x3F;
constexpr int num_bits_elem = 6;
constexpr int num_bits_vec_elem = 32;
constexpr int vector_size = 6;
......@@ -185,6 +177,7 @@ struct bf6x16_pk_t
using element_type = uint32_t;
using type = StaticallyIndexedArray_v2<element_type, 3>;
type data;
typedef int8_t test_vec_t __attribute__((ext_vector_type(16)));
bf6x16_pk_t() : data{type{}} {}
bf6x16_pk_t(type init) : data{init} {}
......@@ -211,18 +204,13 @@ struct bf6x16_pk_t
return static_cast<bf6_t>(bits & 0x3F);
}
// 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)
__host__ __device__ inline type pack(const test_vec_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 packed{};
// for each of the 16 bf6_t values, place its 6 bits in the correct position
ck::static_for<0, 16, 1>{}([&](auto i) {
uint32_t bits = static_cast<uint32_t>(x[i]) & 0x3F;
uint32_t bits = static_cast<uint32_t>(x[static_cast<int>(i)]) & 0x3F;
constexpr int num_bits_elem = 6;
constexpr int num_bits_vec_elem = 32;
constexpr int vector_size = 3;
......@@ -255,6 +243,7 @@ struct bf6x32_pk_t
using element_type = uint32_t;
using type = StaticallyIndexedArray_v2<element_type, 6>;
type data;
typedef int8_t test_vec_t __attribute__((ext_vector_type(32)));
bf6x32_pk_t() : data{type{}} {}
bf6x32_pk_t(type init) : data{init} {}
......@@ -281,18 +270,13 @@ struct bf6x32_pk_t
return static_cast<bf6_t>(bits & 0x3F);
}
// 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)
__host__ __device__ inline type pack(const test_vec_t& x)
{
static_assert(V == 1, "Vector size must be 1.");
static_assert(E == 32, "Number of elements must be 32.");
type packed{};
// for each of the 32 bf6_t values, place its 6 bits in the correct position
ck::static_for<0, 32, 1>{}([&](auto i) {
uint32_t bits = static_cast<uint32_t>(x[i]) & 0x3F;
uint32_t bits = static_cast<uint32_t>(x[static_cast<int>(i)]) & 0x3F;
constexpr int num_bits_elem = 6;
constexpr int num_bits_vec_elem = 32;
constexpr int vector_size = 6;
......
test/data_type/test_bf6.cpp
View file @ df1bad99
......@@ -245,7 +245,8 @@ TEST(BF6, TestAsType16x1)
// test size
const int vector_size = 1;
const int packed_size = 16;
std::vector<bf6_t> test_vec = {bf6_t(0b000000),
typedef int8_t test_vec_t __attribute__((ext_vector_type(16)));
test_vec_t test_vec = {bf6_t(0b000000),
bf6_t(0b100000),
bf6_t(0b000001),
bf6_t(0b100001),
......@@ -271,8 +272,7 @@ 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<vector_size, packed_size>(test_vec.data());
right_vec.template AsType<bf6x16_pk_t>()(Number<i>{}) = bf6x16_pk_t{}.pack(test_vec);
});
// copy the vector
vector_type<bf6x16_pk_t, vector_size> left_vec{right_vec};
......@@ -280,7 +280,7 @@ TEST(BF6, TestAsType16x1)
ck::static_for<0, packed_size, 1>{}([&](auto i) {
ASSERT_EQ(
left_vec.template AsType<bf6x16_pk_t>()(Number<0>{}).template unpack<>(Number<i>{}),
test_vec.at(i));
static_cast<bf6_t>(test_vec[static_cast<int>(i)]));
});
}
......@@ -290,14 +290,40 @@ TEST(BF6, TestAsType16x2)
// test size
const int vector_size = 2;
const int packed_size = 16;
std::vector<bf6_t> test_vec = {
bf6_t(0b000000), bf6_t(0b100000), bf6_t(0b000001), bf6_t(0b100001), bf6_t(0b000010),
bf6_t(0b100010), bf6_t(0b000011), bf6_t(0b100011), bf6_t(0b000100), bf6_t(0b100100),
bf6_t(0b000101), bf6_t(0b100101), bf6_t(0b000110), bf6_t(0b100110), bf6_t(0b001011),
bf6_t(0b101011), bf6_t(0b010000), bf6_t(0b110000), bf6_t(0b010001), bf6_t(0b110001),
bf6_t(0b010010), bf6_t(0b110010), bf6_t(0b010011), bf6_t(0b110011), bf6_t(0b010100),
bf6_t(0b110100), bf6_t(0b010101), bf6_t(0b110101), bf6_t(0b010110), bf6_t(0b110110),
bf6_t(0b011011), bf6_t(0b111011)};
typedef int8_t test_vec_t __attribute__((ext_vector_type(16)));
test_vec_t test_vec[2];
test_vec[0] = {bf6_t(0b000000),
bf6_t(0b100000),
bf6_t(0b000001),
bf6_t(0b100001),
bf6_t(0b000010),
bf6_t(0b100010),
bf6_t(0b000011),
bf6_t(0b100011),
bf6_t(0b000100),
bf6_t(0b100100),
bf6_t(0b000101),
bf6_t(0b100101),
bf6_t(0b000110),
bf6_t(0b100110),
bf6_t(0b001011),
bf6_t(0b101011)};
test_vec[1] = {bf6_t(0b010000),
bf6_t(0b110000),
bf6_t(0b010001),
bf6_t(0b110001),
bf6_t(0b010010),
bf6_t(0b110010),
bf6_t(0b010011),
bf6_t(0b110011),
bf6_t(0b010100),
bf6_t(0b110100),
bf6_t(0b010101),
bf6_t(0b110101),
bf6_t(0b010110),
bf6_t(0b110110),
bf6_t(0b011011),
bf6_t(0b111011)};
// reference vector
vector_type<bf6x16_pk_t, vector_size> right_vec;
// check default CTOR
......@@ -310,8 +336,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<vector_size, packed_size>(test_vec.data() + i * packed_size);
right_vec.template AsType<bf6x16_pk_t>()(Number<i>{}) = bf6x16_pk_t{}.pack(test_vec[i]);
});
// copy the vector
vector_type<bf6x16_pk_t, vector_size> left_vec{right_vec};
......@@ -320,7 +345,7 @@ TEST(BF6, TestAsType16x2)
ck::static_for<0, packed_size, 1>{}([&](auto idx_element) {
ASSERT_EQ(left_vec.template AsType<bf6x16_pk_t>()(Number<idx_vector>{})
.template unpack<>(Number<idx_element>{}),
test_vec.at(idx_vector * packed_size + idx_element));
static_cast<bf6_t>(test_vec[idx_vector][static_cast<int>(idx_element)]));
});
});
}
......@@ -331,14 +356,15 @@ TEST(BF6, TestAsType32x1)
// test size
const int vector_size = 1;
const int packed_size = 32;
std::vector<bf6_t> test_vec = {
bf6_t(0b000000), bf6_t(0b100000), bf6_t(0b000001), bf6_t(0b100001), bf6_t(0b000010),
bf6_t(0b100010), bf6_t(0b000011), bf6_t(0b100011), bf6_t(0b000100), bf6_t(0b100100),
bf6_t(0b000101), bf6_t(0b100101), bf6_t(0b000110), bf6_t(0b100110), bf6_t(0b001011),
bf6_t(0b101011), bf6_t(0b010000), bf6_t(0b110000), bf6_t(0b010001), bf6_t(0b110001),
bf6_t(0b010010), bf6_t(0b110010), bf6_t(0b010011), bf6_t(0b110011), bf6_t(0b010100),
bf6_t(0b110100), bf6_t(0b010101), bf6_t(0b110101), bf6_t(0b010110), bf6_t(0b110110),
bf6_t(0b011011), bf6_t(0b111011)};
typedef int8_t test_vec_t __attribute__((ext_vector_type(32)));
test_vec_t test_vec = {bf6_t(0b000000), bf6_t(0b100000), bf6_t(0b000001), bf6_t(0b100001),
bf6_t(0b000010), bf6_t(0b100010), bf6_t(0b000011), bf6_t(0b100011),
bf6_t(0b000100), bf6_t(0b100100), bf6_t(0b000101), bf6_t(0b100101),
bf6_t(0b000110), bf6_t(0b100110), bf6_t(0b001011), bf6_t(0b101011),
bf6_t(0b010000), bf6_t(0b110000), bf6_t(0b010001), bf6_t(0b110001),
bf6_t(0b010010), bf6_t(0b110010), bf6_t(0b010011), bf6_t(0b110011),
bf6_t(0b010100), bf6_t(0b110100), bf6_t(0b010101), bf6_t(0b110101),
bf6_t(0b010110), bf6_t(0b110110), bf6_t(0b011011), bf6_t(0b111011)};
// reference vector
vector_type<bf6x32_pk_t, vector_size> right_vec;
// check default CTOR
......@@ -349,8 +375,7 @@ 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<vector_size, packed_size>(test_vec.data());
right_vec.template AsType<bf6x32_pk_t>()(Number<i>{}) = bf6x32_pk_t{}.pack(test_vec);
});
// copy the vector
vector_type<bf6x32_pk_t, vector_size> left_vec{right_vec};
......@@ -358,6 +383,6 @@ TEST(BF6, TestAsType32x1)
ck::static_for<0, packed_size, 1>{}([&](auto i) {
ASSERT_EQ(
left_vec.template AsType<bf6x32_pk_t>()(Number<0>{}).template unpack<>(Number<i>{}),
test_vec.at(i));
static_cast<bf6_t>(test_vec[static_cast<int>(i)]));
});
}
test/data_type/test_fp6.cpp
View file @ df1bad99
......@@ -244,7 +244,8 @@ TEST(FP6, TestAsType16x1)
// test size
const int vector_size = 1;
const int packed_size = 16;
std::vector<f6_t> test_vec = {f6_t(0b000000),
typedef int8_t test_vec_t __attribute__((ext_vector_type(16)));
test_vec_t test_vec = {f6_t(0b000000),
f6_t(0b100000),
f6_t(0b000001),
f6_t(0b100001),
......@@ -269,8 +270,7 @@ 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<vector_size, packed_size>(test_vec.data());
right_vec.template AsType<f6x16_pk_t>()(Number<i>{}) = f6x16_pk_t{}.pack(test_vec);
});
// copy the vector
vector_type<f6x16_pk_t, vector_size> left_vec{right_vec};
......@@ -278,7 +278,7 @@ TEST(FP6, TestAsType16x1)
ck::static_for<0, packed_size, 1>{}([&](auto i) {
ASSERT_EQ(
left_vec.template AsType<f6x16_pk_t>()(Number<0>{}).template unpack<>(Number<i>{}),
test_vec.at(i));
static_cast<f6_t>(test_vec[static_cast<int>(i)]));
});
}
......@@ -288,14 +288,40 @@ TEST(FP6, TestAsType16x2)
// test size
const int vector_size = 2;
const int packed_size = 16;
std::vector<f6_t> test_vec = {f6_t(0b000000), f6_t(0b100000), f6_t(0b000001), f6_t(0b100001),
f6_t(0b000010), f6_t(0b100010), f6_t(0b000011), f6_t(0b100011),
f6_t(0b000100), f6_t(0b100100), f6_t(0b000101), f6_t(0b100101),
f6_t(0b000110), f6_t(0b100110), f6_t(0b001011), f6_t(0b101011),
f6_t(0b010000), f6_t(0b110000), f6_t(0b010001), f6_t(0b110001),
f6_t(0b010010), f6_t(0b110010), f6_t(0b010011), f6_t(0b110011),
f6_t(0b010100), f6_t(0b110100), f6_t(0b010101), f6_t(0b110101),
f6_t(0b010110), f6_t(0b110110), f6_t(0b011011), f6_t(0b111011)};
typedef int8_t test_vec_t __attribute__((ext_vector_type(16)));
test_vec_t test_vec[2];
test_vec[0] = {f6_t(0b000000),
f6_t(0b100000),
f6_t(0b000001),
f6_t(0b100001),
f6_t(0b000010),
f6_t(0b100010),
f6_t(0b000011),
f6_t(0b100011),
f6_t(0b000100),
f6_t(0b100100),
f6_t(0b000101),
f6_t(0b100101),
f6_t(0b000110),
f6_t(0b100110),
f6_t(0b001011),
f6_t(0b101011)};
test_vec[1] = {f6_t(0b010000),
f6_t(0b110000),
f6_t(0b010001),
f6_t(0b110001),
f6_t(0b010010),
f6_t(0b110010),
f6_t(0b010011),
f6_t(0b110011),
f6_t(0b010100),
f6_t(0b110100),
f6_t(0b010101),
f6_t(0b110101),
f6_t(0b010110),
f6_t(0b110110),
f6_t(0b011011),
f6_t(0b111011)};
// reference vector
vector_type<f6x16_pk_t, vector_size> right_vec;
// check default CTOR
......@@ -308,8 +334,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<vector_size, packed_size>(test_vec.data() + i * packed_size);
right_vec.template AsType<f6x16_pk_t>()(Number<i>{}) = f6x16_pk_t{}.pack(test_vec[i]);
});
// copy the vector
vector_type<f6x16_pk_t, vector_size> left_vec{right_vec};
......@@ -318,7 +343,7 @@ TEST(FP6, TestAsType16x2)
ck::static_for<0, packed_size, 1>{}([&](auto idx_element) {
ASSERT_EQ(left_vec.template AsType<f6x16_pk_t>()(Number<idx_vector>{})
.template unpack<>(Number<idx_element>{}),
test_vec.at(idx_vector * packed_size + idx_element));
static_cast<f6_t>(test_vec[idx_vector][static_cast<int>(idx_element)]));
});
});
}
......@@ -329,7 +354,8 @@ TEST(FP6, TestAsType32x1)
// test size
const int vector_size = 1;
const int packed_size = 32;
std::vector<f6_t> test_vec = {f6_t(0b000000), f6_t(0b100000), f6_t(0b000001), f6_t(0b100001),
typedef int8_t test_vec_t __attribute__((ext_vector_type(32)));
test_vec_t test_vec = {f6_t(0b000000), f6_t(0b100000), f6_t(0b000001), f6_t(0b100001),
f6_t(0b000010), f6_t(0b100010), f6_t(0b000011), f6_t(0b100011),
f6_t(0b000100), f6_t(0b100100), f6_t(0b000101), f6_t(0b100101),
f6_t(0b000110), f6_t(0b100110), f6_t(0b001011), f6_t(0b101011),
......@@ -346,8 +372,7 @@ 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<vector_size, packed_size>(test_vec.data());
right_vec.template AsType<f6x32_pk_t>()(Number<i>{}) = f6x32_pk_t{}.pack(test_vec);
});
// copy the vector
vector_type<f6x32_pk_t, vector_size> left_vec{right_vec};
......@@ -355,6 +380,6 @@ TEST(FP6, TestAsType32x1)
ck::static_for<0, packed_size, 1>{}([&](auto i) {
ASSERT_EQ(
left_vec.template AsType<f6x32_pk_t>()(Number<0>{}).template unpack<>(Number<i>{}),
test_vec.at(i));
static_cast<f6_t>(test_vec[static_cast<int>(i)]));
});
}
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment