Skip to content

Commit 227f987

Browse files
stijnhstijnh
committed
Add more tests
1 parent 6a7bd2e commit 227f987

File tree

6 files changed

+129
-193
lines changed

6 files changed

+129
-193
lines changed

‎single_include/kernel_float.h‎

Lines changed: 62 additions & 44 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
//================================================================================
22
// this file has been auto-generated, do not modify its contents!
3-
// date: 2023-09-19 20:45:16.880746
4-
// git hash: da0a46b533ef9d25638748eb951284f14e7c48bb
3+
// date: 2023-09-21 09:37:28.638971
4+
// git hash: 07af0ad9ff5c16595790d579577244bc482f0999
55
//================================================================================
66

77
#ifndef KERNEL_FLOAT_MACROS_H
@@ -497,7 +497,7 @@ struct extent<N> {
497497
};
498498

499499
template<typename T>
500-
struct into_vector_traits {
500+
struct into_vector_impl {
501501
using value_type = T;
502502
using extent_type = extent<1>;
503503

@@ -508,7 +508,7 @@ struct into_vector_traits {
508508
};
509509

510510
template<typename T, size_t N>
511-
struct into_vector_traits<T[N]> {
511+
struct into_vector_impl<T[N]> {
512512
using value_type = T;
513513
using extent_type = extent<N>;
514514

@@ -526,19 +526,19 @@ struct into_vector_traits<T[N]> {
526526
};
527527

528528
template<typename V>
529-
struct into_vector_traits<const V>: into_vector_traits<V> {};
529+
struct into_vector_impl<const V>: into_vector_impl<V> {};
530530

531531
template<typename V>
532-
struct into_vector_traits<V&>: into_vector_traits<V> {};
532+
struct into_vector_impl<V&>: into_vector_impl<V> {};
533533

534534
template<typename V>
535-
struct into_vector_traits<const V&>: into_vector_traits<V> {};
535+
struct into_vector_impl<const V&>: into_vector_impl<V> {};
536536

537537
template<typename V>
538-
struct into_vector_traits<V&&>: into_vector_traits<V> {};
538+
struct into_vector_impl<V&&>: into_vector_impl<V> {};
539539

540540
template<typename T, size_t N, size_t A>
541-
struct into_vector_traits<aligned_array<T, N, A>> {
541+
struct into_vector_impl<aligned_array<T, N, A>> {
542542
using value_type = T;
543543
using extent_type = extent<N>;
544544

@@ -550,7 +550,7 @@ struct into_vector_traits<aligned_array<T, N, A>> {
550550

551551
#define KERNEL_FLOAT_DEFINE_VECTOR_TYPE(T, T1, T2, T3, T4) \
552552
template<> \
553-
struct into_vector_traits<::T1> { \
553+
struct into_vector_impl<::T1> { \
554554
using value_type = T; \
555555
using extent_type = extent<1>; \
556556
\
@@ -561,7 +561,7 @@ struct into_vector_traits<aligned_array<T, N, A>> {
561561
}; \
562562
\
563563
template<> \
564-
struct into_vector_traits<::T2> { \
564+
struct into_vector_impl<::T2> { \
565565
using value_type = T; \
566566
using extent_type = extent<2>; \
567567
\
@@ -572,7 +572,7 @@ struct into_vector_traits<aligned_array<T, N, A>> {
572572
}; \
573573
\
574574
template<> \
575-
struct into_vector_traits<::T3> { \
575+
struct into_vector_impl<::T3> { \
576576
using value_type = T; \
577577
using extent_type = extent<3>; \
578578
\
@@ -583,7 +583,7 @@ struct into_vector_traits<aligned_array<T, N, A>> {
583583
}; \
584584
\
585585
template<> \
586-
struct into_vector_traits<::T4> { \
586+
struct into_vector_impl<::T4> { \
587587
using value_type = T; \
588588
using extent_type = extent<4>; \
589589
\
@@ -612,7 +612,7 @@ template<typename T, typename E, typename S = vector_storage<T, E::size>>
612612
struct vector;
613613

614614
template<typename T, typename E, typename S>
615-
struct into_vector_traits<vector<T, E, S>> {
615+
struct into_vector_impl<vector<T, E, S>> {
616616
using value_type = T;
617617
using extent_type = E;
618618

@@ -634,10 +634,10 @@ struct vector_traits<vector<T, E, S>> {
634634
};
635635

636636
template<typename V>
637-
using vector_value_type = typename into_vector_traits<V>::value_type;
637+
using vector_value_type = typename into_vector_impl<V>::value_type;
638638

639639
template<typename V>
640-
using vector_extent_type = typename into_vector_traits<V>::extent_type;
640+
using vector_extent_type = typename into_vector_impl<V>::extent_type;
641641

642642
template<typename V>
643643
static constexpr size_t vector_extent = vector_extent_type<V>::value;
@@ -653,7 +653,7 @@ using promoted_vector_value_type = promote_t<vector_value_type<Vs>...>;
653653

654654
template<typename V>
655655
KERNEL_FLOAT_INLINE vector_storage_type<V> into_vector_storage(V&& input) {
656-
return into_vector_traits<V>::call(std::forward<V>(input));
656+
return into_vector_impl<V>::call(std::forward<V>(input));
657657
}
658658

659659
} // namespace kernel_float
@@ -1732,7 +1732,10 @@ namespace kernel_float {
17321732
template<typename T = double>
17331733
struct constant {
17341734
template<typename R>
1735-
KERNEL_FLOAT_INLINE explicit constexpr constant(const constant<R>& that) : value_(that.get()) {}
1735+
KERNEL_FLOAT_INLINE explicit constexpr constant(const constant<R>& that) {
1736+
auto f = ops::cast<R, T>();
1737+
value_ = f(that.get());
1738+
}
17361739

17371740
KERNEL_FLOAT_INLINE
17381741
constexpr constant(T value = {}) : value_(value) {}
@@ -1793,28 +1796,43 @@ struct cast<constant<T>, R, m> {
17931796
};
17941797
} // namespace ops
17951798

1796-
#define KERNEL_FLOAT_CONSTANT_DEFINE_OP(OP) \
1797-
template<typename L, typename R> \
1798-
R operator OP(const constant<L>& left, const R& right) { \
1799-
using T = vector_value_type<R>; \
1800-
return operator OP(T(left.get()), right); \
1801-
} \
1802-
\
1803-
template<typename L, typename R> \
1804-
L operator OP(const L& left, const constant<R>& right) { \
1805-
using T = vector_value_type<L>; \
1806-
return operator OP(left, T(right.get())); \
1807-
} \
1808-
\
1809-
template<typename L, typename R, typename T = promote_t<L, R>> \
1810-
constant<T> operator OP(const constant<L>& left, const constant<R>& right) { \
1811-
return constant<T>(operator OP(T(left.get()), T(right.get()))); \
1812-
}
1813-
1814-
//KERNEL_FLOAT_CONSTANT_DEFINE_OP(+)
1815-
//KERNEL_FLOAT_CONSTANT_DEFINE_OP(-)
1816-
//KERNEL_FLOAT_CONSTANT_DEFINE_OP(*)
1817-
//KERNEL_FLOAT_CONSTANT_DEFINE_OP(/)
1799+
#define KERNEL_FLOAT_CONSTANT_DEFINE_OP(OP) \
1800+
template<typename L, typename R> \
1801+
KERNEL_FLOAT_INLINE auto operator OP(const constant<L>& left, const R& right) { \
1802+
auto f = ops::cast<L, vector_value_type<R>>(); \
1803+
return f(left.get()) OP right; \
1804+
} \
1805+
\
1806+
template<typename L, typename R> \
1807+
KERNEL_FLOAT_INLINE auto operator OP(const L& left, const constant<R>& right) { \
1808+
auto f = ops::cast<R, vector_value_type<L>>(); \
1809+
return left OP f(right.get()); \
1810+
} \
1811+
\
1812+
template<typename L, typename R, typename E> \
1813+
KERNEL_FLOAT_INLINE auto operator OP(const constant<L>& left, const vector<R, E>& right) { \
1814+
auto f = ops::cast<L, R>(); \
1815+
return f(left.get()) OP right; \
1816+
} \
1817+
\
1818+
template<typename L, typename R, typename E> \
1819+
KERNEL_FLOAT_INLINE auto operator OP(const vector<L, E>& left, const constant<R>& right) { \
1820+
auto f = ops::cast<R, L>(); \
1821+
return left OP f(right.get()); \
1822+
} \
1823+
\
1824+
template<typename L, typename R, typename T = promote_t<L, R>> \
1825+
KERNEL_FLOAT_INLINE constant<T> operator OP( \
1826+
const constant<L>& left, \
1827+
const constant<R>& right) { \
1828+
return constant<T>(left.get()) OP constant<T>(right.get()); \
1829+
}
1830+
1831+
KERNEL_FLOAT_CONSTANT_DEFINE_OP(+)
1832+
KERNEL_FLOAT_CONSTANT_DEFINE_OP(-)
1833+
KERNEL_FLOAT_CONSTANT_DEFINE_OP(*)
1834+
KERNEL_FLOAT_CONSTANT_DEFINE_OP(/)
1835+
KERNEL_FLOAT_CONSTANT_DEFINE_OP(%)
18181836

18191837
} // namespace kernel_float
18201838

@@ -2731,7 +2749,7 @@ namespace ops {
27312749
template<typename T>
27322750
struct fma {
27332751
KERNEL_FLOAT_INLINE T operator()(T a, T b, T c) {
2734-
return a + b * c;
2752+
return a * b + c;
27352753
}
27362754
};
27372755

@@ -3066,7 +3084,7 @@ struct vector: public S {
30663084
*/
30673085
template<typename V>
30683086
KERNEL_FLOAT_INLINE into_vector_type<V> into_vector(V&& input) {
3069-
return into_vector_traits<V>::call(std::forward<V>(input));
3087+
return into_vector_impl<V>::call(std::forward<V>(input));
30703088
}
30713089

30723090
template<typename T>
@@ -3136,7 +3154,7 @@ KERNEL_FLOAT_DEFINE_PROMOTED_TYPE(float, __half)
31363154
KERNEL_FLOAT_DEFINE_PROMOTED_TYPE(double, __half)
31373155

31383156
template<>
3139-
struct into_vector_traits<__half2> {
3157+
struct into_vector_impl<__half2> {
31403158
using value_type = __half;
31413159
using extent_type = extent<2>;
31423160

@@ -3440,7 +3458,7 @@ KERNEL_FLOAT_DEFINE_PROMOTED_TYPE(float, __nv_bfloat16)
34403458
KERNEL_FLOAT_DEFINE_PROMOTED_TYPE(double, __nv_bfloat16)
34413459

34423460
template<>
3443-
struct into_vector_traits<__nv_bfloat162> {
3461+
struct into_vector_impl<__nv_bfloat162> {
34443462
using value_type = __nv_bfloat16;
34453463
using extent_type = extent<2>;
34463464

‎tests/basics.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -106,7 +106,7 @@ struct creation_tests {
106106

107107
// into_vector on scalar
108108
{
109-
kf::vec<float, 1> a = into_vector(int(5.0f));
109+
kf::vec<float, 1> a = into_vector(float(5.0f));
110110
ASSERT(a[0] == 5.0f);
111111
}
112112

‎tests/common.h‎

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -76,6 +76,19 @@ struct equals_helper<__nv_bfloat16> {
7676
}
7777
};
7878

79+
template<typename T, size_t N>
80+
struct equals_helper<kf::vec<T, N>> {
81+
static __host__ __device__ bool call(const kf::vec<T, N>& left, const kf::vec<T, N>& right) {
82+
for (int i = 0; i < N; i++) {
83+
if (!equals_helper<T>::call(left[i], right[i])) {
84+
return false;
85+
}
86+
}
87+
88+
return true;
89+
}
90+
};
91+
7992
} // namespace detail
8093

8194
template<typename T>
@@ -346,11 +359,13 @@ void run_tests_device(F fun, type_sequence<T>, size_sequence<Ns...>) {
346359
#define REGISTER_TEST_CASE_CPU(NAME, F, ...) \
347360
TEMPLATE_TEST_CASE(NAME " - CPU", "", __VA_ARGS__) { \
348361
run_tests_host(F {}, type_sequence<TestType> {}, default_size_sequence {}); \
362+
CHECK("done"); \
349363
}
350364

351365
#define REGISTER_TEST_CASE_GPU(NAME, F, ...) \
352366
TEMPLATE_TEST_CASE(NAME " - GPU", "[GPU]", __VA_ARGS__) { \
353367
run_tests_device(F {}, type_sequence<TestType> {}, default_size_sequence {}); \
368+
CHECK("done"); \
354369
}
355370

356371
#undef REGISTER_TEST_CASE

‎tests/constant.cu‎

Lines changed: 35 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -1,30 +1,40 @@
11
#include "common.h"
22

3-
struct triops_tests {
4-
template<typename T, size_t... I, size_t N = sizeof...(I)>
5-
__host__ __device__ void operator()(generator<T> gen, std::index_sequence<I...>) {
6-
T x[N] = {gen.next(I)...};
7-
T y[N] = {gen.next(I)...};
8-
T z[N] = {gen.next(I)...};
9-
10-
kf::vec<T, N> a = {x[I]...};
11-
kf::vec<T, N> b = {y[I]...};
12-
kf::vec<T, N> c = {z[I]...};
13-
14-
kf::vec<T, N> answer = kf::where(a, b, c);
15-
ASSERT_EQ_ALL(answer[I], bool(x[I]) ? y[I] : z[I]);
16-
17-
answer = kf::where(a, b);
18-
ASSERT_EQ_ALL(answer[I], bool(x[I]) ? y[I] : T());
19-
20-
answer = kf::where(a);
21-
ASSERT_EQ_ALL(answer[I], T(bool(x[I])));
22-
23-
answer = kf::fma(a, b, c);
24-
ASSERT_EQ_ALL(answer[I], x[I] * y[I] + z[I]);
25-
3+
#define ASSERT_TYPE(A, B) ASSERT(std::is_same<decltype(A), B>::value);
4+
5+
struct constant_tests {
6+
template<typename T>
7+
__host__ __device__ void operator()(generator<T> gen) {
8+
T value = gen.next();
9+
kf::vec<T, 2> vector = {gen.next(), gen.next()};
10+
11+
ASSERT_EQ(kf::make_constant(5.0) + value, T(5) + value);
12+
ASSERT_EQ(value + kf::make_constant(5.0), value + T(5));
13+
ASSERT_EQ(kf::make_constant(5.0) + vector, T(5) + vector);
14+
ASSERT_EQ(vector + kf::make_constant(5.0), vector + T(5));
15+
16+
ASSERT_EQ(kf::make_constant(5.0) - value, T(5) - value);
17+
ASSERT_EQ(value - kf::make_constant(5.0), value - T(5));
18+
ASSERT_EQ(kf::make_constant(5.0) - vector, T(5) - vector);
19+
ASSERT_EQ(vector - kf::make_constant(5.0), vector - T(5));
20+
21+
ASSERT_EQ(kf::make_constant(5.0) * value, T(5) * value);
22+
ASSERT_EQ(value * kf::make_constant(5.0), value * T(5));
23+
ASSERT_EQ(kf::make_constant(5.0) * vector, T(5) * vector);
24+
ASSERT_EQ(vector * kf::make_constant(5.0), vector * T(5));
25+
26+
// These results in division by zero for integers
27+
// ASSERT_EQ(kf::make_constant(5.0) / value, T(5) / value);
28+
// ASSERT_EQ(value / kf::make_constant(5.0), value / T(5));
29+
// ASSERT_EQ(kf::make_constant(5.0) / vector, T(5) / vector);
30+
// ASSERT_EQ(vector / kf::make_constant(5.0), vector / T(5));
31+
//
32+
// ASSERT_EQ(kf::make_constant(5.0) % value, T(5) % value);
33+
// ASSERT_EQ(value % kf::make_constant(5.0), value % T(5));
34+
// ASSERT_EQ(kf::make_constant(5.0) % vector, T(5) % vector);
35+
// ASSERT_EQ(vector % kf::make_constant(5.0), vector % T(5));
2636
}
2737
};
2838

29-
REGISTER_TEST_CASE("ternary operators", triops_tests, int, float, double)
30-
REGISTER_TEST_CASE_GPU("ternary operators", triops_tests, __half, __nv_bfloat16)
39+
REGISTER_TEST_CASE("constant tests", constant_tests, int, float, double)
40+
REGISTER_TEST_CASE_GPU("constant tests", constant_tests, __half, __nv_bfloat16)

‎tests/promotion.cu‎

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,8 @@
11
#include "common.h"
22

3-
// Check if combining type `A` and `B` results in `C`
4-
#define CHECK_PROMOTION(A, B, C) CHECK(std::is_same<kernel_float::promote_t<A, B>, C>::value);
3+
// Check if combining type `vec<A, N>` and `vec<B, N>` results in `vec<C, N>`
4+
#define CHECK_PROMOTION(A, B, C) \
5+
CHECK(std::is_same<decltype(kf::vec<A, 2>() + kf::vec<B, 2>()), kf::vec<C, 2>>::value);
56

67
TEST_CASE("type promotion") {
78
CHECK_PROMOTION(int, int, int);

0 commit comments

Comments
 (0)