Skip to content

GitLab

Commit 3e520bbd authored by Rostyslav Geyyer's avatar Rostyslav Geyyer
Browse files

Update unpack method

parent 48d58131
Hide whitespace changes
Inline Side-by-side
Showing with 124 additions and 98 deletions
include/ck/utility/data_type.hpp
View file @ 3e520bbd
......@@ -32,11 +32,13 @@ struct f4x2_pk_t
f4x2_pk_t() : data{type{}} {}
f4x2_pk_t(type init) : data{init} {}
__host__ __device__ inline type unpack(const size_t index)
template <index_t I>
__host__ __device__ inline type unpack()
{
if(index == 0)
if constexpr(I == 0)
return data & 0b00001111;
return (data >> 4);
else
return (data >> 4);
}
__host__ __device__ inline type pack(const type x0, const type x1)
......
include/ck/utility/type_convert.hpp
View file @ 3e520bbd
......@@ -1018,8 +1018,8 @@ inline __host__ __device__ float2_t type_convert<float2_t, f4x2_t>(f4x2_t x)
float scale = 1.0f;
return __builtin_amdgcn_cvt_scalef32_pk_f32_fp4(value.bitwise, scale, 0);
#else
float2_t ret{utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(), x.unpack(1)),
utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(), x.unpack(0))};
float2_t ret{utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(), x.unpack<1>()),
utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(), x.unpack<0>())};
return ret;
#endif
}
......@@ -1154,72 +1154,72 @@ inline __host__ __device__ float32_t type_convert<float32_t, f4x32_t>(f4x32_t x)
} f4_values{bit_cast<__uint128_t>(x)};
// TODO: pack in a loop
float_values.float_array[0] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[0].unpack(0));
f4_values.f4x2_array[0].unpack<0>());
float_values.float_array[1] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[0].unpack(1));
f4_values.f4x2_array[0].unpack<1>());
float_values.float_array[2] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[1].unpack(0));
f4_values.f4x2_array[1].unpack<0>());
float_values.float_array[3] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[1].unpack(1));
f4_values.f4x2_array[1].unpack<1>());
float_values.float_array[4] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[2].unpack(0));
f4_values.f4x2_array[2].unpack<0>());
float_values.float_array[5] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[2].unpack(1));
f4_values.f4x2_array[2].unpack<1>());
float_values.float_array[6] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[3].unpack(0));
f4_values.f4x2_array[3].unpack<0>());
float_values.float_array[7] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[3].unpack(1));
f4_values.f4x2_array[3].unpack<1>());
float_values.float_array[0] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[4].unpack(0));
f4_values.f4x2_array[4].unpack<0>());
float_values.float_array[1] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[4].unpack(1));
f4_values.f4x2_array[4].unpack<1>());
float_values.float_array[2] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[5].unpack(0));
f4_values.f4x2_array[5].unpack<0>());
float_values.float_array[3] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[5].unpack(1));
f4_values.f4x2_array[5].unpack<1>());
float_values.float_array[4] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[6].unpack(0));
f4_values.f4x2_array[6].unpack<0>());
float_values.float_array[5] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[6].unpack(1));
f4_values.f4x2_array[6].unpack<1>());
float_values.float_array[6] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[7].unpack(0));
f4_values.f4x2_array[7].unpack<0>());
float_values.float_array[7] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[7].unpack(1));
f4_values.f4x2_array[7].unpack<1>());
float_values.float_array[0] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[8].unpack(0));
f4_values.f4x2_array[8].unpack<0>());
float_values.float_array[1] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[8].unpack(1));
f4_values.f4x2_array[8].unpack<1>());
float_values.float_array[2] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[9].unpack(0));
f4_values.f4x2_array[9].unpack<0>());
float_values.float_array[3] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[9].unpack(1));
f4_values.f4x2_array[9].unpack<1>());
float_values.float_array[4] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[10].unpack(0));
f4_values.f4x2_array[10].unpack<0>());
float_values.float_array[5] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[10].unpack(1));
f4_values.f4x2_array[10].unpack<1>());
float_values.float_array[6] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[11].unpack(0));
f4_values.f4x2_array[11].unpack<0>());
float_values.float_array[7] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[11].unpack(1));
f4_values.f4x2_array[11].unpack<1>());
float_values.float_array[0] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[12].unpack(0));
f4_values.f4x2_array[12].unpack<0>());
float_values.float_array[1] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[12].unpack(1));
f4_values.f4x2_array[12].unpack<1>());
float_values.float_array[2] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[13].unpack(0));
f4_values.f4x2_array[13].unpack<0>());
float_values.float_array[3] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[13].unpack(1));
f4_values.f4x2_array[13].unpack<1>());
float_values.float_array[4] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[14].unpack(0));
f4_values.f4x2_array[14].unpack<0>());
float_values.float_array[5] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[14].unpack(1));
f4_values.f4x2_array[14].unpack<1>());
float_values.float_array[6] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[15].unpack(0));
f4_values.f4x2_array[15].unpack<0>());
float_values.float_array[7] = utils::to_float<f4_t>(NumericLimits<e8m0_bexp_t>::Binary_1(),
f4_values.f4x2_array[15].unpack(1));
f4_values.f4x2_array[15].unpack<1>());
return float_values.float32_array;
#endif
......@@ -1273,8 +1273,8 @@ inline __host__ __device__ float2_t scaled_type_convert<float2_t, f4x2_t>(e8m0_b
value.f4x2_array[0] = x;
return __builtin_amdgcn_cvt_scalef32_pk_f32_fp4(value.bitwise, type_convert<float>(scale), 0);
#else
float2_t ret{utils::to_float<f4_t>(scale, x.unpack(1)),
utils::to_float<f4_t>(scale, x.unpack(0))};
float2_t ret{utils::to_float<f4_t>(scale, x.unpack<1>()),
utils::to_float<f4_t>(scale, x.unpack<0>())};
return ret;
#endif
}
......@@ -1408,41 +1408,53 @@ inline __host__ __device__ float32_t scaled_type_convert<float32_t, f4x32_t>(e8m
f4x32_t f4x32_array;
} f4_values{bit_cast<__uint128_t>(x)};
// TODO: pack in a loop
float_values.float_array[0] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[0].unpack(0));
float_values.float_array[1] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[0].unpack(1));
float_values.float_array[2] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[1].unpack(0));
float_values.float_array[3] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[1].unpack(1));
float_values.float_array[4] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[2].unpack(0));
float_values.float_array[5] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[2].unpack(1));
float_values.float_array[6] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[3].unpack(0));
float_values.float_array[7] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[3].unpack(1));
float_values.float_array[0] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[4].unpack(0));
float_values.float_array[1] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[4].unpack(1));
float_values.float_array[2] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[5].unpack(0));
float_values.float_array[3] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[5].unpack(1));
float_values.float_array[4] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[6].unpack(0));
float_values.float_array[5] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[6].unpack(1));
float_values.float_array[6] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[7].unpack(0));
float_values.float_array[7] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[7].unpack(1));
float_values.float_array[0] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[8].unpack(0));
float_values.float_array[1] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[8].unpack(1));
float_values.float_array[2] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[9].unpack(0));
float_values.float_array[3] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[9].unpack(1));
float_values.float_array[4] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[10].unpack(0));
float_values.float_array[5] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[10].unpack(1));
float_values.float_array[6] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[11].unpack(0));
float_values.float_array[7] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[11].unpack(1));
float_values.float_array[0] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[12].unpack(0));
float_values.float_array[1] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[12].unpack(1));
float_values.float_array[2] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[13].unpack(0));
float_values.float_array[3] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[13].unpack(1));
float_values.float_array[4] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[14].unpack(0));
float_values.float_array[5] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[14].unpack(1));
float_values.float_array[6] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[15].unpack(0));
float_values.float_array[7] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[15].unpack(1));
float_values.float_array[0] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[0].unpack<0>());
float_values.float_array[1] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[0].unpack<1>());
float_values.float_array[2] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[1].unpack<0>());
float_values.float_array[3] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[1].unpack<1>());
float_values.float_array[4] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[2].unpack<0>());
float_values.float_array[5] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[2].unpack<1>());
float_values.float_array[6] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[3].unpack<0>());
float_values.float_array[7] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[3].unpack<1>());
float_values.float_array[0] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[4].unpack<0>());
float_values.float_array[1] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[4].unpack<1>());
float_values.float_array[2] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[5].unpack<0>());
float_values.float_array[3] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[5].unpack<1>());
float_values.float_array[4] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[6].unpack<0>());
float_values.float_array[5] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[6].unpack<1>());
float_values.float_array[6] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[7].unpack<0>());
float_values.float_array[7] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[7].unpack<1>());
float_values.float_array[0] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[8].unpack<0>());
float_values.float_array[1] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[8].unpack<1>());
float_values.float_array[2] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[9].unpack<0>());
float_values.float_array[3] = utils::to_float<f4_t>(scale, f4_values.f4x2_array[9].unpack<1>());
float_values.float_array[4] =
utils::to_float<f4_t>(scale, f4_values.f4x2_array[10].unpack<0>());
float_values.float_array[5] =
utils::to_float<f4_t>(scale, f4_values.f4x2_array[10].unpack<1>());
float_values.float_array[6] =
utils::to_float<f4_t>(scale, f4_values.f4x2_array[11].unpack<0>());
float_values.float_array[7] =
utils::to_float<f4_t>(scale, f4_values.f4x2_array[11].unpack<1>());
float_values.float_array[0] =
utils::to_float<f4_t>(scale, f4_values.f4x2_array[12].unpack<0>());
float_values.float_array[1] =
utils::to_float<f4_t>(scale, f4_values.f4x2_array[12].unpack<1>());
float_values.float_array[2] =
utils::to_float<f4_t>(scale, f4_values.f4x2_array[13].unpack<0>());
float_values.float_array[3] =
utils::to_float<f4_t>(scale, f4_values.f4x2_array[13].unpack<1>());
float_values.float_array[4] =
utils::to_float<f4_t>(scale, f4_values.f4x2_array[14].unpack<0>());
float_values.float_array[5] =
utils::to_float<f4_t>(scale, f4_values.f4x2_array[14].unpack<1>());
float_values.float_array[6] =
utils::to_float<f4_t>(scale, f4_values.f4x2_array[15].unpack<0>());
float_values.float_array[7] =
utils::to_float<f4_t>(scale, f4_values.f4x2_array[15].unpack<1>());
return float_values.float32_array;
#endif
......
test/data_type/test_fp4.cpp
View file @ 3e520bbd
......@@ -256,8 +256,8 @@ TEST(FP4, TestAsType1)
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);
ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<0>(), 0);
ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<1>(), 0);
});
// assign test values to the vector
ck::static_for<0, size, 1>{}([&](auto i) {
......@@ -268,8 +268,10 @@ TEST(FP4, TestAsType1)
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));
ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<0>(),
test_vec.at(i));
ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<1>(),
test_vec.at(i + 1));
});
}
......@@ -286,8 +288,8 @@ TEST(FP4, TestAsType2)
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);
ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<0>(), 0);
ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<1>(), 0);
});
// assign test values to the vector
ck::static_for<0, size, 1>{}([&](auto i) {
......@@ -298,8 +300,10 @@ TEST(FP4, TestAsType2)
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));
ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<0>(),
test_vec.at(i));
ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<1>(),
test_vec.at(i + 1));
});
}
......@@ -320,8 +324,8 @@ TEST(FP4, TestAsType4)
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);
ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<0>(), 0);
ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<1>(), 0);
});
// assign test values to the vector
ck::static_for<0, size, 1>{}([&](auto i) {
......@@ -332,8 +336,10 @@ TEST(FP4, TestAsType4)
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));
ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<0>(),
test_vec.at(i));
ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<1>(),
test_vec.at(i + 1));
});
}
......@@ -362,8 +368,8 @@ TEST(FP4, TestAsType8)
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);
ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<0>(), 0);
ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<1>(), 0);
});
// assign test values to the vector
ck::static_for<0, size, 1>{}([&](auto i) {
......@@ -374,8 +380,10 @@ TEST(FP4, TestAsType8)
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));
ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<0>(),
test_vec.at(i));
ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<1>(),
test_vec.at(i + 1));
});
}
......@@ -400,8 +408,8 @@ TEST(FP4, TestAsType16)
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);
ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<0>(), 0);
ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<1>(), 0);
});
// assign test values to the vector
ck::static_for<0, size, 1>{}([&](auto i) {
......@@ -412,8 +420,10 @@ TEST(FP4, TestAsType16)
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));
ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<0>(),
test_vec.at(i));
ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<1>(),
test_vec.at(i + 1));
});
}
......@@ -449,8 +459,8 @@ TEST(FP4, TestAsType32)
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);
ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<0>(), 0);
ASSERT_EQ(right_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<1>(), 0);
});
// assign test values to the vector
ck::static_for<0, size, 1>{}([&](auto i) {
......@@ -461,7 +471,9 @@ TEST(FP4, TestAsType32)
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));
ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<0>(),
test_vec.at(i));
ASSERT_EQ(left_vec.template AsType<f4x2_pk_t>()(Number<i>{}).template unpack<1>(),
test_vec.at(i + 1));
});
}
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