mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-14 10:09:41 +00:00
Add custom type vector support (#1333)
* Add non_native_vector_type
* Add a test
* Add non-native vector type
* Fix CTOR
* Fix non-native vector type of 1
* Fix CTORs
* Use vector_type to cover non-native implementation as well
* Update the test
* Format
* Format
* Fix copyright years
* Remove BoolVecT so far
* Add AsType test cases
* Update assert error message
* Remove redundant type
* Update naming
* Add complex half type with tests
* Add tests for vector reshaping
* Add missing alignas
* Update test/data_type/test_custom_type.cpp
Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
* Compare custom types to built-in types
* Add default constructor test
* Add an alignment test
---------
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>
[ROCm/composable_kernel commit: 4cf70b36c1]
This commit is contained in:
Rostyslav Geyyer
@@ -1,5 +1,5 @@
|
||||
// SPDX-License-Identifier: MIT
|
||||
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
|
||||
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
#pragma once
|
||||
|
||||
@@ -13,8 +13,24 @@ using int4_t = _BitInt(4);
|
||||
using f8_t = _BitInt(8);
|
||||
using bf8_t = unsigned _BitInt(8);
|
||||
|
||||
inline constexpr auto next_pow2(uint32_t x)
|
||||
{
|
||||
// Precondition: x > 1.
|
||||
return x > 1u ? (1u << (32u - __builtin_clz(x - 1u))) : x;
|
||||
}
|
||||
|
||||
// native types: double, float, _Float16, ushort, int32_t, int8_t, uint8_t, f8_t, bf8_t, bool
|
||||
template <typename T>
|
||||
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;
|
||||
}
|
||||
|
||||
// vector_type
|
||||
template <typename T, index_t N>
|
||||
template <typename T, index_t N, typename Enable = void>
|
||||
struct vector_type;
|
||||
|
||||
// Caution: DO NOT REMOVE
|
||||
@@ -171,7 +187,7 @@ struct scalar_type<bool>
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct vector_type<T, 1>
|
||||
struct vector_type<T, 1, typename std::enable_if_t<is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
using type = d1_t;
|
||||
@@ -189,7 +205,8 @@ struct vector_type<T, 1>
|
||||
template <typename X>
|
||||
__host__ __device__ constexpr const auto& AsType() const
|
||||
{
|
||||
static_assert(is_same<X, d1_t>::value, "wrong!");
|
||||
static_assert(is_same<X, d1_t>::value,
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
return data_.d1x1_;
|
||||
}
|
||||
@@ -197,7 +214,8 @@ struct vector_type<T, 1>
|
||||
template <typename X>
|
||||
__host__ __device__ constexpr auto& AsType()
|
||||
{
|
||||
static_assert(is_same<X, d1_t>::value, "wrong!");
|
||||
static_assert(is_same<X, d1_t>::value,
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
return data_.d1x1_;
|
||||
}
|
||||
@@ -205,7 +223,7 @@ struct vector_type<T, 1>
|
||||
|
||||
__device__ int static err = 0;
|
||||
template <typename T>
|
||||
struct vector_type<T, 2>
|
||||
struct vector_type<T, 2, typename std::enable_if_t<is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
typedef T d2_t __attribute__((ext_vector_type(2)));
|
||||
@@ -226,7 +244,8 @@ struct vector_type<T, 2>
|
||||
template <typename X>
|
||||
__host__ __device__ constexpr const auto& AsType() const
|
||||
{
|
||||
static_assert(is_same<X, d1_t>::value || is_same<X, d2_t>::value, "wrong!");
|
||||
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)
|
||||
{
|
||||
@@ -245,7 +264,8 @@ struct vector_type<T, 2>
|
||||
template <typename X>
|
||||
__host__ __device__ constexpr auto& AsType()
|
||||
{
|
||||
static_assert(is_same<X, d1_t>::value || is_same<X, d2_t>::value, "wrong!");
|
||||
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)
|
||||
{
|
||||
@@ -263,7 +283,7 @@ struct vector_type<T, 2>
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct vector_type<T, 4>
|
||||
struct vector_type<T, 4, typename std::enable_if_t<is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
typedef T d2_t __attribute__((ext_vector_type(2)));
|
||||
@@ -287,7 +307,7 @@ struct vector_type<T, 4>
|
||||
__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,
|
||||
"wrong!");
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
@@ -311,7 +331,7 @@ struct vector_type<T, 4>
|
||||
__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,
|
||||
"wrong!");
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
@@ -333,7 +353,7 @@ struct vector_type<T, 4>
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct vector_type<T, 8>
|
||||
struct vector_type<T, 8, typename std::enable_if_t<is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
typedef T d2_t __attribute__((ext_vector_type(2)));
|
||||
@@ -360,7 +380,7 @@ struct vector_type<T, 8>
|
||||
{
|
||||
static_assert(is_same<X, d1_t>::value || is_same<X, d2_t>::value ||
|
||||
is_same<X, d4_t>::value || is_same<X, d8_t>::value,
|
||||
"wrong!");
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
@@ -389,7 +409,7 @@ struct vector_type<T, 8>
|
||||
{
|
||||
static_assert(is_same<X, d1_t>::value || is_same<X, d2_t>::value ||
|
||||
is_same<X, d4_t>::value || is_same<X, d8_t>::value,
|
||||
"wrong!");
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
@@ -415,7 +435,7 @@ struct vector_type<T, 8>
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct vector_type<T, 16>
|
||||
struct vector_type<T, 16, typename std::enable_if_t<is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
typedef T d2_t __attribute__((ext_vector_type(2)));
|
||||
@@ -445,7 +465,7 @@ struct vector_type<T, 16>
|
||||
static_assert(is_same<X, d1_t>::value || is_same<X, d2_t>::value ||
|
||||
is_same<X, d4_t>::value || is_same<X, d8_t>::value ||
|
||||
is_same<X, d16_t>::value,
|
||||
"wrong!");
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
@@ -479,7 +499,7 @@ struct vector_type<T, 16>
|
||||
static_assert(is_same<X, d1_t>::value || is_same<X, d2_t>::value ||
|
||||
is_same<X, d4_t>::value || is_same<X, d8_t>::value ||
|
||||
is_same<X, d16_t>::value,
|
||||
"wrong!");
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
@@ -509,7 +529,7 @@ struct vector_type<T, 16>
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct vector_type<T, 32>
|
||||
struct vector_type<T, 32, typename std::enable_if_t<is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
typedef T d2_t __attribute__((ext_vector_type(2)));
|
||||
@@ -541,7 +561,7 @@ struct vector_type<T, 32>
|
||||
static_assert(is_same<X, d1_t>::value || is_same<X, d2_t>::value ||
|
||||
is_same<X, d4_t>::value || is_same<X, d8_t>::value ||
|
||||
is_same<X, d16_t>::value || is_same<X, d32_t>::value,
|
||||
"wrong!");
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
@@ -579,7 +599,7 @@ struct vector_type<T, 32>
|
||||
static_assert(is_same<X, d1_t>::value || is_same<X, d2_t>::value ||
|
||||
is_same<X, d4_t>::value || is_same<X, d8_t>::value ||
|
||||
is_same<X, d16_t>::value || is_same<X, d32_t>::value,
|
||||
"wrong!");
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
@@ -613,7 +633,7 @@ struct vector_type<T, 32>
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct vector_type<T, 64>
|
||||
struct vector_type<T, 64, typename std::enable_if_t<is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
typedef T d2_t __attribute__((ext_vector_type(2)));
|
||||
@@ -648,7 +668,7 @@ struct vector_type<T, 64>
|
||||
is_same<X, d4_t>::value || is_same<X, d8_t>::value ||
|
||||
is_same<X, d16_t>::value || is_same<X, d32_t>::value ||
|
||||
is_same<X, d64_t>::value,
|
||||
"wrong!");
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
@@ -691,7 +711,7 @@ struct vector_type<T, 64>
|
||||
is_same<X, d4_t>::value || is_same<X, d8_t>::value ||
|
||||
is_same<X, d16_t>::value || is_same<X, d32_t>::value ||
|
||||
is_same<X, d64_t>::value,
|
||||
"wrong!");
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
@@ -729,7 +749,7 @@ struct vector_type<T, 64>
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct vector_type<T, 128>
|
||||
struct vector_type<T, 128, typename std::enable_if_t<is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
typedef T d2_t __attribute__((ext_vector_type(2)));
|
||||
@@ -766,7 +786,7 @@ struct vector_type<T, 128>
|
||||
is_same<X, d4_t>::value || is_same<X, d8_t>::value ||
|
||||
is_same<X, d16_t>::value || is_same<X, d32_t>::value ||
|
||||
is_same<X, d64_t>::value || is_same<X, d128_t>::value,
|
||||
"wrong!");
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
@@ -813,7 +833,7 @@ struct vector_type<T, 128>
|
||||
is_same<X, d4_t>::value || is_same<X, d8_t>::value ||
|
||||
is_same<X, d16_t>::value || is_same<X, d32_t>::value ||
|
||||
is_same<X, d64_t>::value || is_same<X, d128_t>::value,
|
||||
"wrong!");
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
@@ -855,7 +875,7 @@ struct vector_type<T, 128>
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct vector_type<T, 256>
|
||||
struct vector_type<T, 256, typename std::enable_if_t<is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
typedef T d2_t __attribute__((ext_vector_type(2)));
|
||||
@@ -894,7 +914,7 @@ struct vector_type<T, 256>
|
||||
is_same<X, d1_t>::value || is_same<X, d2_t>::value || is_same<X, d4_t>::value ||
|
||||
is_same<X, d8_t>::value || is_same<X, d16_t>::value || is_same<X, d32_t>::value ||
|
||||
is_same<X, d64_t>::value || is_same<X, d128_t>::value || is_same<X, d256_t>::value,
|
||||
"wrong!");
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
@@ -945,7 +965,7 @@ struct vector_type<T, 256>
|
||||
is_same<X, d1_t>::value || is_same<X, d2_t>::value || is_same<X, d4_t>::value ||
|
||||
is_same<X, d8_t>::value || is_same<X, d16_t>::value || is_same<X, d32_t>::value ||
|
||||
is_same<X, d64_t>::value || is_same<X, d128_t>::value || is_same<X, d256_t>::value,
|
||||
"wrong!");
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
@@ -990,6 +1010,581 @@ struct vector_type<T, 256>
|
||||
}
|
||||
};
|
||||
|
||||
template <typename T, index_t N>
|
||||
struct non_native_vector_base
|
||||
{
|
||||
using type = non_native_vector_base<T, N>;
|
||||
|
||||
__host__ __device__ non_native_vector_base() = default;
|
||||
__host__ __device__ non_native_vector_base(const type&) = default;
|
||||
__host__ __device__ non_native_vector_base(type&&) = default;
|
||||
__host__ __device__ ~non_native_vector_base() = default;
|
||||
|
||||
T d[N];
|
||||
};
|
||||
|
||||
// non-native vector_type implementation
|
||||
template <typename T>
|
||||
struct vector_type<T, 1, typename std::enable_if_t<!is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
using type = d1_t;
|
||||
|
||||
union alignas(next_pow2(1 * sizeof(T)))
|
||||
{
|
||||
d1_t d1_;
|
||||
StaticallyIndexedArray<d1_t, 1> d1x1_;
|
||||
} 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,
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
return data_.d1x1_;
|
||||
}
|
||||
|
||||
template <typename X>
|
||||
__host__ __device__ constexpr auto& AsType()
|
||||
{
|
||||
static_assert(is_same<X, d1_t>::value,
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
return data_.d1x1_;
|
||||
}
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct vector_type<T, 2, typename std::enable_if_t<!is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
using d2_t = non_native_vector_base<T, 2>;
|
||||
|
||||
using type = d2_t;
|
||||
|
||||
union alignas(next_pow2(2 * sizeof(T)))
|
||||
{
|
||||
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 <typename T>
|
||||
struct vector_type<T, 4, typename std::enable_if_t<!is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
using d2_t = non_native_vector_base<T, 2>;
|
||||
using d4_t = non_native_vector_base<T, 4>;
|
||||
|
||||
using type = d4_t;
|
||||
|
||||
union alignas(next_pow2(4 * sizeof(T)))
|
||||
{
|
||||
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;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct vector_type<T, 8, typename std::enable_if_t<!is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
using d2_t = non_native_vector_base<T, 2>;
|
||||
using d4_t = non_native_vector_base<T, 4>;
|
||||
using d8_t = non_native_vector_base<T, 8>;
|
||||
|
||||
using type = d8_t;
|
||||
|
||||
union alignas(next_pow2(8 * sizeof(T)))
|
||||
{
|
||||
d8_t d8_;
|
||||
StaticallyIndexedArray<d1_t, 8> d1x8_;
|
||||
StaticallyIndexedArray<d2_t, 4> d2x4_;
|
||||
StaticallyIndexedArray<d4_t, 2> d4x2_;
|
||||
StaticallyIndexedArray<d8_t, 1> d8x1_;
|
||||
} 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 || is_same<X, d8_t>::value,
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
return data_.d1x8_;
|
||||
}
|
||||
else if constexpr(is_same<X, d2_t>::value)
|
||||
{
|
||||
return data_.d2x4_;
|
||||
}
|
||||
else if constexpr(is_same<X, d4_t>::value)
|
||||
{
|
||||
return data_.d4x2_;
|
||||
}
|
||||
else if constexpr(is_same<X, d8_t>::value)
|
||||
{
|
||||
return data_.d8x1_;
|
||||
}
|
||||
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 || is_same<X, d8_t>::value,
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
return data_.d1x8_;
|
||||
}
|
||||
else if constexpr(is_same<X, d2_t>::value)
|
||||
{
|
||||
return data_.d2x4_;
|
||||
}
|
||||
else if constexpr(is_same<X, d4_t>::value)
|
||||
{
|
||||
return data_.d4x2_;
|
||||
}
|
||||
else if constexpr(is_same<X, d8_t>::value)
|
||||
{
|
||||
return data_.d8x1_;
|
||||
}
|
||||
else
|
||||
{
|
||||
return err;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct vector_type<T, 16, typename std::enable_if_t<!is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
using d2_t = non_native_vector_base<T, 2>;
|
||||
using d4_t = non_native_vector_base<T, 4>;
|
||||
using d8_t = non_native_vector_base<T, 8>;
|
||||
using d16_t = non_native_vector_base<T, 16>;
|
||||
|
||||
using type = d16_t;
|
||||
|
||||
union alignas(next_pow2(16 * sizeof(T)))
|
||||
{
|
||||
d16_t d16_;
|
||||
StaticallyIndexedArray<d1_t, 16> d1x16_;
|
||||
StaticallyIndexedArray<d2_t, 8> d2x8_;
|
||||
StaticallyIndexedArray<d4_t, 4> d4x4_;
|
||||
StaticallyIndexedArray<d8_t, 2> d8x2_;
|
||||
StaticallyIndexedArray<d16_t, 1> d16x1_;
|
||||
} 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 || is_same<X, d8_t>::value ||
|
||||
is_same<X, d16_t>::value,
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
return data_.d1x16_;
|
||||
}
|
||||
else if constexpr(is_same<X, d2_t>::value)
|
||||
{
|
||||
return data_.d2x8_;
|
||||
}
|
||||
else if constexpr(is_same<X, d4_t>::value)
|
||||
{
|
||||
return data_.d4x4_;
|
||||
}
|
||||
else if constexpr(is_same<X, d8_t>::value)
|
||||
{
|
||||
return data_.d8x2_;
|
||||
}
|
||||
else if constexpr(is_same<X, d16_t>::value)
|
||||
{
|
||||
return data_.d16x1_;
|
||||
}
|
||||
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 || is_same<X, d8_t>::value ||
|
||||
is_same<X, d16_t>::value,
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
return data_.d1x16_;
|
||||
}
|
||||
else if constexpr(is_same<X, d2_t>::value)
|
||||
{
|
||||
return data_.d2x8_;
|
||||
}
|
||||
else if constexpr(is_same<X, d4_t>::value)
|
||||
{
|
||||
return data_.d4x4_;
|
||||
}
|
||||
else if constexpr(is_same<X, d8_t>::value)
|
||||
{
|
||||
return data_.d8x2_;
|
||||
}
|
||||
else if constexpr(is_same<X, d16_t>::value)
|
||||
{
|
||||
return data_.d16x1_;
|
||||
}
|
||||
else
|
||||
{
|
||||
return err;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct vector_type<T, 32, typename std::enable_if_t<!is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
using d2_t = non_native_vector_base<T, 2>;
|
||||
using d4_t = non_native_vector_base<T, 4>;
|
||||
using d8_t = non_native_vector_base<T, 8>;
|
||||
using d16_t = non_native_vector_base<T, 16>;
|
||||
using d32_t = non_native_vector_base<T, 32>;
|
||||
|
||||
using type = d32_t;
|
||||
|
||||
union alignas(next_pow2(32 * sizeof(T)))
|
||||
{
|
||||
d32_t d32_;
|
||||
StaticallyIndexedArray<d1_t, 32> d1x32_;
|
||||
StaticallyIndexedArray<d2_t, 16> d2x16_;
|
||||
StaticallyIndexedArray<d4_t, 8> d4x8_;
|
||||
StaticallyIndexedArray<d8_t, 4> d8x4_;
|
||||
StaticallyIndexedArray<d16_t, 2> d16x2_;
|
||||
StaticallyIndexedArray<d32_t, 1> d32x1_;
|
||||
} 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 || is_same<X, d8_t>::value ||
|
||||
is_same<X, d16_t>::value || is_same<X, d32_t>::value,
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
return data_.d1x32_;
|
||||
}
|
||||
else if constexpr(is_same<X, d2_t>::value)
|
||||
{
|
||||
return data_.d2x16_;
|
||||
}
|
||||
else if constexpr(is_same<X, d4_t>::value)
|
||||
{
|
||||
return data_.d4x8_;
|
||||
}
|
||||
else if constexpr(is_same<X, d8_t>::value)
|
||||
{
|
||||
return data_.d8x4_;
|
||||
}
|
||||
else if constexpr(is_same<X, d16_t>::value)
|
||||
{
|
||||
return data_.d16x2_;
|
||||
}
|
||||
else if constexpr(is_same<X, d32_t>::value)
|
||||
{
|
||||
return data_.d32x1_;
|
||||
}
|
||||
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 || is_same<X, d8_t>::value ||
|
||||
is_same<X, d16_t>::value || is_same<X, d32_t>::value,
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
return data_.d1x32_;
|
||||
}
|
||||
else if constexpr(is_same<X, d2_t>::value)
|
||||
{
|
||||
return data_.d2x16_;
|
||||
}
|
||||
else if constexpr(is_same<X, d4_t>::value)
|
||||
{
|
||||
return data_.d4x8_;
|
||||
}
|
||||
else if constexpr(is_same<X, d8_t>::value)
|
||||
{
|
||||
return data_.d8x4_;
|
||||
}
|
||||
else if constexpr(is_same<X, d16_t>::value)
|
||||
{
|
||||
return data_.d16x2_;
|
||||
}
|
||||
else if constexpr(is_same<X, d32_t>::value)
|
||||
{
|
||||
return data_.d32x1_;
|
||||
}
|
||||
else
|
||||
{
|
||||
return err;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct vector_type<T, 64, typename std::enable_if_t<!is_native_type<T>()>>
|
||||
{
|
||||
using d1_t = T;
|
||||
using d2_t = non_native_vector_base<T, 2>;
|
||||
using d4_t = non_native_vector_base<T, 4>;
|
||||
using d8_t = non_native_vector_base<T, 8>;
|
||||
using d16_t = non_native_vector_base<T, 16>;
|
||||
using d32_t = non_native_vector_base<T, 32>;
|
||||
using d64_t = non_native_vector_base<T, 64>;
|
||||
|
||||
using type = d64_t;
|
||||
|
||||
union alignas(next_pow2(64 * sizeof(T)))
|
||||
{
|
||||
d64_t d64_;
|
||||
StaticallyIndexedArray<d1_t, 64> d1x64_;
|
||||
StaticallyIndexedArray<d2_t, 32> d2x32_;
|
||||
StaticallyIndexedArray<d4_t, 16> d4x16_;
|
||||
StaticallyIndexedArray<d8_t, 8> d8x8_;
|
||||
StaticallyIndexedArray<d16_t, 4> d16x4_;
|
||||
StaticallyIndexedArray<d32_t, 2> d32x2_;
|
||||
StaticallyIndexedArray<d64_t, 1> d64x1_;
|
||||
} 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 || is_same<X, d8_t>::value ||
|
||||
is_same<X, d16_t>::value || is_same<X, d32_t>::value ||
|
||||
is_same<X, d64_t>::value,
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
return data_.d1x64_;
|
||||
}
|
||||
else if constexpr(is_same<X, d2_t>::value)
|
||||
{
|
||||
return data_.d2x32_;
|
||||
}
|
||||
else if constexpr(is_same<X, d4_t>::value)
|
||||
{
|
||||
return data_.d4x16_;
|
||||
}
|
||||
else if constexpr(is_same<X, d8_t>::value)
|
||||
{
|
||||
return data_.d8x8_;
|
||||
}
|
||||
else if constexpr(is_same<X, d16_t>::value)
|
||||
{
|
||||
return data_.d16x4_;
|
||||
}
|
||||
else if constexpr(is_same<X, d32_t>::value)
|
||||
{
|
||||
return data_.d32x2_;
|
||||
}
|
||||
else if constexpr(is_same<X, d64_t>::value)
|
||||
{
|
||||
return data_.d64x1_;
|
||||
}
|
||||
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 || is_same<X, d8_t>::value ||
|
||||
is_same<X, d16_t>::value || is_same<X, d32_t>::value ||
|
||||
is_same<X, d64_t>::value,
|
||||
"Something went wrong, please check src and dst types.");
|
||||
|
||||
if constexpr(is_same<X, d1_t>::value)
|
||||
{
|
||||
return data_.d1x64_;
|
||||
}
|
||||
else if constexpr(is_same<X, d2_t>::value)
|
||||
{
|
||||
return data_.d2x32_;
|
||||
}
|
||||
else if constexpr(is_same<X, d4_t>::value)
|
||||
{
|
||||
return data_.d4x16_;
|
||||
}
|
||||
else if constexpr(is_same<X, d8_t>::value)
|
||||
{
|
||||
return data_.d8x8_;
|
||||
}
|
||||
else if constexpr(is_same<X, d16_t>::value)
|
||||
{
|
||||
return data_.d16x4_;
|
||||
}
|
||||
else if constexpr(is_same<X, d32_t>::value)
|
||||
{
|
||||
return data_.d32x2_;
|
||||
}
|
||||
else if constexpr(is_same<X, d64_t>::value)
|
||||
{
|
||||
return data_.d64x1_;
|
||||
}
|
||||
else
|
||||
{
|
||||
return err;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
using int64_t = long;
|
||||
|
||||
// fp64
|
||||
@@ -1051,8 +1646,8 @@ using bf8x8_t = typename vector_type<bf8_t, 8>::type;
|
||||
using bf8x16_t = typename vector_type<bf8_t, 16>::type;
|
||||
using bf8x32_t = typename vector_type<bf8_t, 32>::type;
|
||||
using bf8x64_t = typename vector_type<bf8_t, 64>::type;
|
||||
|
||||
// u8
|
||||
// i8
|
||||
using uint8x2_t = typename vector_type<uint8_t, 2>::type;
|
||||
using uint8x4_t = typename vector_type<uint8_t, 4>::type;
|
||||
using uint8x8_t = typename vector_type<uint8_t, 8>::type;
|
||||
|
||||
@@ -18,4 +18,9 @@ if(result EQUAL 0)
|
||||
target_link_libraries(test_bf8 PRIVATE utility)
|
||||
endif()
|
||||
|
||||
add_gtest_executable(test_custom_type test_custom_type.cpp)
|
||||
if(result EQUAL 0)
|
||||
target_link_libraries(test_custom_type PRIVATE utility)
|
||||
endif()
|
||||
|
||||
add_gtest_executable(test_type_convert_const type_convert_const.cpp)
|
||||
|
||||
874
test/data_type/test_custom_type.cpp
Normal file
874
test/data_type/test_custom_type.cpp
Normal file
@@ -0,0 +1,874 @@
|
||||
// SPDX-License-Identifier: MIT
|
||||
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
#include "gtest/gtest.h"
|
||||
#include "ck/utility/data_type.hpp"
|
||||
#include "ck/utility/type_convert.hpp"
|
||||
|
||||
using ck::bf8_t;
|
||||
using ck::bhalf_t;
|
||||
using ck::f8_t;
|
||||
using ck::half_t;
|
||||
using ck::Number;
|
||||
using ck::type_convert;
|
||||
using ck::vector_type;
|
||||
|
||||
TEST(Custom_bool, TestSize)
|
||||
{
|
||||
struct custom_bool_t
|
||||
{
|
||||
bool data;
|
||||
};
|
||||
ASSERT_EQ(sizeof(custom_bool_t), sizeof(bool));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bool_t, 2>), sizeof(vector_type<bool, 2>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bool_t, 4>), sizeof(vector_type<bool, 4>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bool_t, 8>), sizeof(vector_type<bool, 8>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bool_t, 16>), sizeof(vector_type<bool, 16>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bool_t, 32>), sizeof(vector_type<bool, 32>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bool_t, 64>), sizeof(vector_type<bool, 64>));
|
||||
}
|
||||
|
||||
TEST(Custom_bool, TestAsType)
|
||||
{
|
||||
struct custom_bool_t
|
||||
{
|
||||
using type = bool;
|
||||
type data;
|
||||
custom_bool_t() : data{type{}} {}
|
||||
custom_bool_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<bool> test_vec = {false, true, false, true};
|
||||
// reference vector
|
||||
vector_type<custom_bool_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<custom_bool_t>()(Number<i>{}).data, false);
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_bool_t>()(Number<i>{}) = custom_bool_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the vector
|
||||
vector_type<custom_bool_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<custom_bool_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_bool, TestAsTypeReshape)
|
||||
{
|
||||
struct custom_bool_t
|
||||
{
|
||||
using type = bool;
|
||||
type data;
|
||||
custom_bool_t() : data{type{}} {}
|
||||
custom_bool_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<bool> test_vec = {false, true, false, true};
|
||||
// reference vector
|
||||
vector_type<custom_bool_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<custom_bool_t>()(Number<i>{}).data, false);
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_bool_t>()(Number<i>{}) = custom_bool_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the first half of a vector
|
||||
vector_type<custom_bool_t, size / 2> left_vec{
|
||||
right_vec.template AsType<vector_type<custom_bool_t, size / 2>::type>()(Number<0>{})};
|
||||
// check if values were copied correctly
|
||||
ck::static_for<0, size / 2, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(left_vec.template AsType<custom_bool_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_int8, TestSize)
|
||||
{
|
||||
struct custom_int8_t
|
||||
{
|
||||
int8_t data;
|
||||
};
|
||||
ASSERT_EQ(sizeof(custom_int8_t), sizeof(int8_t));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_int8_t, 2>), sizeof(vector_type<int8_t, 2>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_int8_t, 4>), sizeof(vector_type<int8_t, 4>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_int8_t, 8>), sizeof(vector_type<int8_t, 8>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_int8_t, 16>), sizeof(vector_type<int8_t, 16>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_int8_t, 32>), sizeof(vector_type<int8_t, 32>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_int8_t, 64>), sizeof(vector_type<int8_t, 64>));
|
||||
}
|
||||
|
||||
TEST(Custom_int8, TestAsType)
|
||||
{
|
||||
struct custom_int8_t
|
||||
{
|
||||
using type = int8_t;
|
||||
type data;
|
||||
custom_int8_t() : data{type{}} {}
|
||||
custom_int8_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<int8_t> test_vec = {3, -6, 8, -2};
|
||||
// reference vector
|
||||
vector_type<custom_int8_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<custom_int8_t>()(Number<i>{}).data, 0);
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_int8_t>()(Number<i>{}) = custom_int8_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the vector
|
||||
vector_type<custom_int8_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<custom_int8_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_int8, TestAsTypeReshape)
|
||||
{
|
||||
struct custom_int8_t
|
||||
{
|
||||
using type = int8_t;
|
||||
type data;
|
||||
custom_int8_t() : data{type{}} {}
|
||||
custom_int8_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<int8_t> test_vec = {3, -6, 8, -2};
|
||||
// reference vector
|
||||
vector_type<custom_int8_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<custom_int8_t>()(Number<i>{}).data, 0);
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_int8_t>()(Number<i>{}) = custom_int8_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the first half of a vector
|
||||
vector_type<custom_int8_t, size / 2> left_vec{
|
||||
right_vec.template AsType<vector_type<custom_int8_t, size / 2>::type>()(Number<0>{})};
|
||||
// check if values were copied correctly
|
||||
ck::static_for<0, size / 2, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(left_vec.template AsType<custom_int8_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_uint8, TestSize)
|
||||
{
|
||||
struct custom_uint8_t
|
||||
{
|
||||
uint8_t data;
|
||||
};
|
||||
ASSERT_EQ(sizeof(custom_uint8_t), sizeof(uint8_t));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_uint8_t, 2>), sizeof(vector_type<uint8_t, 2>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_uint8_t, 4>), sizeof(vector_type<uint8_t, 4>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_uint8_t, 8>), sizeof(vector_type<uint8_t, 8>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_uint8_t, 16>), sizeof(vector_type<uint8_t, 16>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_uint8_t, 32>), sizeof(vector_type<uint8_t, 32>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_uint8_t, 64>), sizeof(vector_type<uint8_t, 64>));
|
||||
}
|
||||
|
||||
TEST(Custom_uint8, TestAsType)
|
||||
{
|
||||
struct custom_uint8_t
|
||||
{
|
||||
using type = uint8_t;
|
||||
type data;
|
||||
custom_uint8_t() : data{type{}} {}
|
||||
custom_uint8_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<uint8_t> test_vec = {3, 6, 8, 2};
|
||||
// reference vector
|
||||
vector_type<custom_uint8_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<custom_uint8_t>()(Number<i>{}).data, 0);
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_uint8_t>()(Number<i>{}) = custom_uint8_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the vector
|
||||
vector_type<custom_uint8_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<custom_uint8_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_uint8, TestAsTypeReshape)
|
||||
{
|
||||
struct custom_uint8_t
|
||||
{
|
||||
using type = uint8_t;
|
||||
type data;
|
||||
custom_uint8_t() : data{type{}} {}
|
||||
custom_uint8_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<uint8_t> test_vec = {3, 6, 8, 2};
|
||||
// reference vector
|
||||
vector_type<custom_uint8_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<custom_uint8_t>()(Number<i>{}).data, 0);
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_uint8_t>()(Number<i>{}) = custom_uint8_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the first half of a vector
|
||||
vector_type<custom_uint8_t, size / 2> left_vec{
|
||||
right_vec.template AsType<vector_type<custom_uint8_t, size / 2>::type>()(Number<0>{})};
|
||||
// check if values were copied correctly
|
||||
ck::static_for<0, size / 2, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(left_vec.template AsType<custom_uint8_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_f8, TestSize)
|
||||
{
|
||||
struct custom_f8_t
|
||||
{
|
||||
_BitInt(8) data;
|
||||
};
|
||||
ASSERT_EQ(sizeof(custom_f8_t), sizeof(_BitInt(8)));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_f8_t, 2>), sizeof(vector_type<_BitInt(8), 2>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_f8_t, 4>), sizeof(vector_type<_BitInt(8), 4>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_f8_t, 8>), sizeof(vector_type<_BitInt(8), 8>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_f8_t, 16>), sizeof(vector_type<_BitInt(8), 16>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_f8_t, 32>), sizeof(vector_type<_BitInt(8), 32>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_f8_t, 64>), sizeof(vector_type<_BitInt(8), 64>));
|
||||
}
|
||||
|
||||
TEST(Custom_f8, TestAsType)
|
||||
{
|
||||
struct custom_f8_t
|
||||
{
|
||||
using type = _BitInt(8);
|
||||
type data;
|
||||
custom_f8_t() : data{type{}} {}
|
||||
custom_f8_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<_BitInt(8)> test_vec = {type_convert<_BitInt(8)>(0.3f),
|
||||
type_convert<_BitInt(8)>(-0.6f),
|
||||
type_convert<_BitInt(8)>(0.8f),
|
||||
type_convert<_BitInt(8)>(-0.2f)};
|
||||
// reference vector
|
||||
vector_type<custom_f8_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}(
|
||||
[&](auto i) { ASSERT_EQ(right_vec.template AsType<custom_f8_t>()(Number<i>{}).data, 0); });
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_f8_t>()(Number<i>{}) = custom_f8_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the vector
|
||||
vector_type<custom_f8_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<custom_f8_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_f8, TestAsTypeReshape)
|
||||
{
|
||||
struct custom_f8_t
|
||||
{
|
||||
using type = _BitInt(8);
|
||||
type data;
|
||||
custom_f8_t() : data{type{}} {}
|
||||
custom_f8_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<_BitInt(8)> test_vec = {type_convert<_BitInt(8)>(0.3f),
|
||||
type_convert<_BitInt(8)>(-0.6f),
|
||||
type_convert<_BitInt(8)>(0.8f),
|
||||
type_convert<_BitInt(8)>(-0.2f)};
|
||||
// reference vector
|
||||
vector_type<custom_f8_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}(
|
||||
[&](auto i) { ASSERT_EQ(right_vec.template AsType<custom_f8_t>()(Number<i>{}).data, 0); });
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_f8_t>()(Number<i>{}) = custom_f8_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the first half of a vector
|
||||
vector_type<custom_f8_t, size / 2> left_vec{
|
||||
right_vec.template AsType<vector_type<custom_f8_t, size / 2>::type>()(Number<0>{})};
|
||||
// check if values were copied correctly
|
||||
ck::static_for<0, size / 2, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(left_vec.template AsType<custom_f8_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_bf8, TestSize)
|
||||
{
|
||||
struct custom_bf8_t
|
||||
{
|
||||
unsigned _BitInt(8) data;
|
||||
};
|
||||
ASSERT_EQ(sizeof(custom_bf8_t), sizeof(unsigned _BitInt(8)));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bf8_t, 2>), sizeof(vector_type<unsigned _BitInt(8), 2>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bf8_t, 4>), sizeof(vector_type<unsigned _BitInt(8), 4>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bf8_t, 8>), sizeof(vector_type<unsigned _BitInt(8), 8>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bf8_t, 16>), sizeof(vector_type<unsigned _BitInt(8), 16>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bf8_t, 32>), sizeof(vector_type<unsigned _BitInt(8), 32>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bf8_t, 64>), sizeof(vector_type<unsigned _BitInt(8), 64>));
|
||||
}
|
||||
|
||||
TEST(Custom_bf8, TestAsType)
|
||||
{
|
||||
struct custom_bf8_t
|
||||
{
|
||||
using type = unsigned _BitInt(8);
|
||||
type data;
|
||||
custom_bf8_t() : data{type{}} {}
|
||||
custom_bf8_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<unsigned _BitInt(8)> test_vec = {type_convert<unsigned _BitInt(8)>(0.3f),
|
||||
type_convert<unsigned _BitInt(8)>(-0.6f),
|
||||
type_convert<unsigned _BitInt(8)>(0.8f),
|
||||
type_convert<unsigned _BitInt(8)>(-0.2f)};
|
||||
// reference vector
|
||||
vector_type<custom_bf8_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}(
|
||||
[&](auto i) { ASSERT_EQ(right_vec.template AsType<custom_bf8_t>()(Number<i>{}).data, 0); });
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_bf8_t>()(Number<i>{}) = custom_bf8_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the vector
|
||||
vector_type<custom_bf8_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<custom_bf8_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_bf8, TestAsTypeReshape)
|
||||
{
|
||||
struct custom_bf8_t
|
||||
{
|
||||
using type = unsigned _BitInt(8);
|
||||
type data;
|
||||
custom_bf8_t() : data{type{}} {}
|
||||
custom_bf8_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<unsigned _BitInt(8)> test_vec = {type_convert<unsigned _BitInt(8)>(0.3f),
|
||||
type_convert<unsigned _BitInt(8)>(-0.6f),
|
||||
type_convert<unsigned _BitInt(8)>(0.8f),
|
||||
type_convert<unsigned _BitInt(8)>(-0.2f)};
|
||||
// reference vector
|
||||
vector_type<custom_bf8_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}(
|
||||
[&](auto i) { ASSERT_EQ(right_vec.template AsType<custom_bf8_t>()(Number<i>{}).data, 0); });
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_bf8_t>()(Number<i>{}) = custom_bf8_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the first half of a vector
|
||||
vector_type<custom_bf8_t, size / 2> left_vec{
|
||||
right_vec.template AsType<vector_type<custom_bf8_t, size / 2>::type>()(Number<0>{})};
|
||||
// check if values were copied correctly
|
||||
ck::static_for<0, size / 2, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(left_vec.template AsType<custom_bf8_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_half, TestSize)
|
||||
{
|
||||
struct custom_half_t
|
||||
{
|
||||
half_t data;
|
||||
};
|
||||
ASSERT_EQ(sizeof(custom_half_t), sizeof(half_t));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_half_t, 2>), sizeof(vector_type<half_t, 2>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_half_t, 4>), sizeof(vector_type<half_t, 4>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_half_t, 8>), sizeof(vector_type<half_t, 8>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_half_t, 16>), sizeof(vector_type<half_t, 16>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_half_t, 32>), sizeof(vector_type<half_t, 32>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_half_t, 64>), sizeof(vector_type<half_t, 64>));
|
||||
}
|
||||
|
||||
TEST(Custom_half, TestAsType)
|
||||
{
|
||||
struct custom_half_t
|
||||
{
|
||||
using type = half_t;
|
||||
type data;
|
||||
custom_half_t() : data{type{}} {}
|
||||
custom_half_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<half_t> test_vec = {half_t{0.3f}, half_t{-0.6f}, half_t{0.8f}, half_t{-0.2f}};
|
||||
// reference vector
|
||||
vector_type<custom_half_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<custom_half_t>()(Number<i>{}).data,
|
||||
type_convert<half_t>(0.0f));
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_half_t>()(Number<i>{}) = custom_half_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the vector
|
||||
vector_type<custom_half_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<custom_half_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_half, TestAsTypeReshape)
|
||||
{
|
||||
struct custom_half_t
|
||||
{
|
||||
using type = half_t;
|
||||
type data;
|
||||
custom_half_t() : data{type{}} {}
|
||||
custom_half_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<half_t> test_vec = {half_t{0.3f}, half_t{-0.6f}, half_t{0.8f}, half_t{-0.2f}};
|
||||
// reference vector
|
||||
vector_type<custom_half_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<custom_half_t>()(Number<i>{}).data,
|
||||
type_convert<half_t>(0.0f));
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_half_t>()(Number<i>{}) = custom_half_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the first half of a vector
|
||||
vector_type<custom_half_t, size / 2> left_vec{
|
||||
right_vec.template AsType<vector_type<custom_half_t, size / 2>::type>()(Number<0>{})};
|
||||
// check if values were copied correctly
|
||||
ck::static_for<0, size / 2, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(left_vec.template AsType<custom_half_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_bhalf, TestSize)
|
||||
{
|
||||
struct custom_bhalf_t
|
||||
{
|
||||
bhalf_t data;
|
||||
};
|
||||
ASSERT_EQ(sizeof(custom_bhalf_t), sizeof(bhalf_t));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bhalf_t, 2>), sizeof(vector_type<bhalf_t, 2>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bhalf_t, 4>), sizeof(vector_type<bhalf_t, 4>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bhalf_t, 8>), sizeof(vector_type<bhalf_t, 8>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bhalf_t, 16>), sizeof(vector_type<bhalf_t, 16>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bhalf_t, 32>), sizeof(vector_type<bhalf_t, 32>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_bhalf_t, 64>), sizeof(vector_type<bhalf_t, 64>));
|
||||
}
|
||||
|
||||
TEST(Custom_bhalf, TestAsType)
|
||||
{
|
||||
struct custom_bhalf_t
|
||||
{
|
||||
using type = bhalf_t;
|
||||
type data;
|
||||
custom_bhalf_t() : data{type{}} {}
|
||||
custom_bhalf_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<bhalf_t> test_vec = {type_convert<bhalf_t>(0.3f),
|
||||
type_convert<bhalf_t>(-0.6f),
|
||||
type_convert<bhalf_t>(0.8f),
|
||||
type_convert<bhalf_t>(-0.2f)};
|
||||
// reference vector
|
||||
vector_type<custom_bhalf_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<custom_bhalf_t>()(Number<i>{}).data,
|
||||
type_convert<bhalf_t>(0.0f));
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_bhalf_t>()(Number<i>{}) = custom_bhalf_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the vector
|
||||
vector_type<custom_bhalf_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<custom_bhalf_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_bhalf, TestAsTypeReshape)
|
||||
{
|
||||
struct custom_bhalf_t
|
||||
{
|
||||
using type = bhalf_t;
|
||||
type data;
|
||||
custom_bhalf_t() : data{type{}} {}
|
||||
custom_bhalf_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<bhalf_t> test_vec = {type_convert<bhalf_t>(0.3f),
|
||||
type_convert<bhalf_t>(-0.6f),
|
||||
type_convert<bhalf_t>(0.8f),
|
||||
type_convert<bhalf_t>(-0.2f)};
|
||||
// reference vector
|
||||
vector_type<custom_bhalf_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<custom_bhalf_t>()(Number<i>{}).data,
|
||||
type_convert<bhalf_t>(0.0f));
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_bhalf_t>()(Number<i>{}) = custom_bhalf_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the first half of a vector
|
||||
vector_type<custom_bhalf_t, size / 2> left_vec{
|
||||
right_vec.template AsType<vector_type<custom_bhalf_t, size / 2>::type>()(Number<0>{})};
|
||||
// check if values were copied correctly
|
||||
ck::static_for<0, size / 2, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(left_vec.template AsType<custom_bhalf_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_float, TestSize)
|
||||
{
|
||||
struct custom_float_t
|
||||
{
|
||||
float data;
|
||||
};
|
||||
ASSERT_EQ(sizeof(custom_float_t), sizeof(float));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_float_t, 2>), sizeof(vector_type<float, 2>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_float_t, 4>), sizeof(vector_type<float, 4>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_float_t, 8>), sizeof(vector_type<float, 8>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_float_t, 16>), sizeof(vector_type<float, 16>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_float_t, 32>), sizeof(vector_type<float, 32>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_float_t, 64>), sizeof(vector_type<float, 64>));
|
||||
}
|
||||
|
||||
TEST(Custom_float, TestAsType)
|
||||
{
|
||||
struct custom_float_t
|
||||
{
|
||||
using type = float;
|
||||
type data;
|
||||
custom_float_t() : data{type{}} {}
|
||||
custom_float_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<float> test_vec = {0.3f, -0.6f, 0.8f, -0.2f};
|
||||
// reference vector
|
||||
vector_type<custom_float_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<custom_float_t>()(Number<i>{}).data, 0.0f);
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_float_t>()(Number<i>{}) = custom_float_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the vector
|
||||
vector_type<custom_float_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<custom_float_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_float, TestAsTypeReshape)
|
||||
{
|
||||
struct custom_float_t
|
||||
{
|
||||
using type = float;
|
||||
type data;
|
||||
custom_float_t() : data{type{}} {}
|
||||
custom_float_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<float> test_vec = {0.3f, -0.6f, 0.8f, -0.2f};
|
||||
// reference vector
|
||||
vector_type<custom_float_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<custom_float_t>()(Number<i>{}).data, 0.0f);
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_float_t>()(Number<i>{}) = custom_float_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the first half of a vector
|
||||
vector_type<custom_float_t, size / 2> left_vec{
|
||||
right_vec.template AsType<vector_type<custom_float_t, size / 2>::type>()(Number<0>{})};
|
||||
// check if values were copied correctly
|
||||
ck::static_for<0, size / 2, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(left_vec.template AsType<custom_float_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_double, TestSize)
|
||||
{
|
||||
struct custom_double_t
|
||||
{
|
||||
double data;
|
||||
};
|
||||
ASSERT_EQ(sizeof(custom_double_t), sizeof(double));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_double_t, 2>), sizeof(vector_type<double, 2>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_double_t, 4>), sizeof(vector_type<double, 4>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_double_t, 8>), sizeof(vector_type<double, 8>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_double_t, 16>), sizeof(vector_type<double, 16>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_double_t, 32>), sizeof(vector_type<double, 32>));
|
||||
ASSERT_EQ(sizeof(vector_type<custom_double_t, 64>), sizeof(vector_type<double, 64>));
|
||||
}
|
||||
|
||||
TEST(Custom_double, TestAsType)
|
||||
{
|
||||
struct custom_double_t
|
||||
{
|
||||
using type = double;
|
||||
type data;
|
||||
custom_double_t() : data{type{}} {}
|
||||
custom_double_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<double> test_vec = {0.3, 0.6, 0.8, 0.2};
|
||||
// reference vector
|
||||
vector_type<custom_double_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<custom_double_t>()(Number<i>{}).data, 0.0);
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_double_t>()(Number<i>{}) = custom_double_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the vector
|
||||
vector_type<custom_double_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<custom_double_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Custom_double, TestAsTypeReshape)
|
||||
{
|
||||
struct custom_double_t
|
||||
{
|
||||
using type = double;
|
||||
type data;
|
||||
custom_double_t() : data{type{}} {}
|
||||
custom_double_t(type init) : data{init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
std::vector<double> test_vec = {0.3, 0.6, 0.8, 0.2};
|
||||
// reference vector
|
||||
vector_type<custom_double_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<custom_double_t>()(Number<i>{}).data, 0.0);
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<custom_double_t>()(Number<i>{}) = custom_double_t{test_vec.at(i)};
|
||||
});
|
||||
// copy the first half of a vector
|
||||
vector_type<custom_double_t, size / 2> left_vec{
|
||||
right_vec.template AsType<vector_type<custom_double_t, size / 2>::type>()(Number<0>{})};
|
||||
// check if values were copied correctly
|
||||
ck::static_for<0, size / 2, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(left_vec.template AsType<custom_double_t>()(Number<i>{}).data, test_vec.at(i));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Complex_half, TestSize)
|
||||
{
|
||||
struct complex_half_t
|
||||
{
|
||||
half_t real;
|
||||
half_t img;
|
||||
};
|
||||
ASSERT_EQ(sizeof(complex_half_t), sizeof(half_t) + sizeof(half_t));
|
||||
ASSERT_EQ(sizeof(vector_type<complex_half_t, 2>),
|
||||
sizeof(vector_type<half_t, 2>) + sizeof(vector_type<half_t, 2>));
|
||||
ASSERT_EQ(sizeof(vector_type<complex_half_t, 4>),
|
||||
sizeof(vector_type<half_t, 4>) + sizeof(vector_type<half_t, 4>));
|
||||
ASSERT_EQ(sizeof(vector_type<complex_half_t, 8>),
|
||||
sizeof(vector_type<half_t, 8>) + sizeof(vector_type<half_t, 8>));
|
||||
ASSERT_EQ(sizeof(vector_type<complex_half_t, 16>),
|
||||
sizeof(vector_type<half_t, 16>) + sizeof(vector_type<half_t, 16>));
|
||||
ASSERT_EQ(sizeof(vector_type<complex_half_t, 32>),
|
||||
sizeof(vector_type<half_t, 32>) + sizeof(vector_type<half_t, 32>));
|
||||
ASSERT_EQ(sizeof(vector_type<complex_half_t, 64>),
|
||||
sizeof(vector_type<half_t, 64>) + sizeof(vector_type<half_t, 64>));
|
||||
}
|
||||
|
||||
TEST(Complex_half, TestAlignment)
|
||||
{
|
||||
struct complex_half_t
|
||||
{
|
||||
half_t real;
|
||||
half_t img;
|
||||
};
|
||||
ASSERT_EQ(alignof(vector_type<complex_half_t, 2>),
|
||||
alignof(vector_type<half_t, 2>) + alignof(vector_type<half_t, 2>));
|
||||
ASSERT_EQ(alignof(vector_type<complex_half_t, 4>),
|
||||
alignof(vector_type<half_t, 4>) + alignof(vector_type<half_t, 4>));
|
||||
ASSERT_EQ(alignof(vector_type<complex_half_t, 8>),
|
||||
alignof(vector_type<half_t, 8>) + alignof(vector_type<half_t, 8>));
|
||||
ASSERT_EQ(alignof(vector_type<complex_half_t, 16>),
|
||||
alignof(vector_type<half_t, 16>) + alignof(vector_type<half_t, 16>));
|
||||
ASSERT_EQ(alignof(vector_type<complex_half_t, 32>),
|
||||
alignof(vector_type<half_t, 32>) + alignof(vector_type<half_t, 32>));
|
||||
ASSERT_EQ(alignof(vector_type<complex_half_t, 64>),
|
||||
alignof(vector_type<half_t, 64>) + alignof(vector_type<half_t, 64>));
|
||||
}
|
||||
|
||||
TEST(Complex_half, TestAsType)
|
||||
{
|
||||
struct complex_half_t
|
||||
{
|
||||
using type = half_t;
|
||||
type real;
|
||||
type img;
|
||||
complex_half_t() : real{type{}}, img{type{}} {}
|
||||
complex_half_t(type real_init, type img_init) : real{real_init}, img{img_init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
// custom type number of elements
|
||||
const int num_elem = sizeof(complex_half_t) / sizeof(complex_half_t::type);
|
||||
std::vector<half_t> test_vec = {half_t{0.3f},
|
||||
half_t{-0.6f},
|
||||
half_t{0.8f},
|
||||
half_t{-0.2f},
|
||||
half_t{0.5f},
|
||||
half_t{-0.7f},
|
||||
half_t{0.9f},
|
||||
half_t{-0.3f}};
|
||||
// reference vector
|
||||
vector_type<complex_half_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<complex_half_t>()(Number<i>{}).real,
|
||||
type_convert<half_t>(0.0f));
|
||||
ASSERT_EQ(right_vec.template AsType<complex_half_t>()(Number<i>{}).img,
|
||||
type_convert<half_t>(0.0f));
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<complex_half_t>()(Number<i>{}) =
|
||||
complex_half_t{test_vec.at(num_elem * i), test_vec.at(num_elem * i + 1)};
|
||||
});
|
||||
// copy the vector
|
||||
vector_type<complex_half_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<complex_half_t>()(Number<i>{}).real,
|
||||
test_vec.at(num_elem * i));
|
||||
ASSERT_EQ(left_vec.template AsType<complex_half_t>()(Number<i>{}).img,
|
||||
test_vec.at(num_elem * i + 1));
|
||||
});
|
||||
}
|
||||
|
||||
TEST(Complex_half, TestAsTypeReshape)
|
||||
{
|
||||
struct complex_half_t
|
||||
{
|
||||
using type = half_t;
|
||||
type real;
|
||||
type img;
|
||||
complex_half_t() : real{type{}}, img{type{}} {}
|
||||
complex_half_t(type real_init, type img_init) : real{real_init}, img{img_init} {}
|
||||
};
|
||||
|
||||
// test size
|
||||
const int size = 4;
|
||||
// custom type number of elements
|
||||
const int num_elem = sizeof(complex_half_t) / sizeof(complex_half_t::type);
|
||||
std::vector<half_t> test_vec = {half_t{0.3f},
|
||||
half_t{-0.6f},
|
||||
half_t{0.8f},
|
||||
half_t{-0.2f},
|
||||
half_t{0.5f},
|
||||
half_t{-0.7f},
|
||||
half_t{0.9f},
|
||||
half_t{-0.3f}};
|
||||
// reference vector
|
||||
vector_type<complex_half_t, size> right_vec;
|
||||
// check default CTOR
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(right_vec.template AsType<complex_half_t>()(Number<i>{}).real,
|
||||
type_convert<half_t>(0.0f));
|
||||
ASSERT_EQ(right_vec.template AsType<complex_half_t>()(Number<i>{}).img,
|
||||
type_convert<half_t>(0.0f));
|
||||
});
|
||||
// assign test values to the vector
|
||||
ck::static_for<0, size, 1>{}([&](auto i) {
|
||||
right_vec.template AsType<complex_half_t>()(Number<i>{}) =
|
||||
complex_half_t{test_vec.at(num_elem * i), test_vec.at(num_elem * i + 1)};
|
||||
});
|
||||
// copy the first half of a vector
|
||||
vector_type<complex_half_t, size / 2> left_vec{
|
||||
right_vec.template AsType<vector_type<complex_half_t, size / 2>::type>()(Number<0>{})};
|
||||
// check if values were copied correctly
|
||||
ck::static_for<0, size / 2, 1>{}([&](auto i) {
|
||||
ASSERT_EQ(left_vec.template AsType<complex_half_t>()(Number<i>{}).real,
|
||||
test_vec.at(num_elem * i));
|
||||
ASSERT_EQ(left_vec.template AsType<complex_half_t>()(Number<i>{}).img,
|
||||
test_vec.at(num_elem * i + 1));
|
||||
});
|
||||
}
|
||||
Reference in New Issue
Block a user