Skip to content

Commit 4824f97

Browse files
stijnhstijnh
committed
Compile tests with sm_70 and sm_80
1 parent 48beccb commit 4824f97

File tree

7 files changed

+163
-155
lines changed

7 files changed

+163
-155
lines changed

‎tests/CMakeLists.txt‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@ file(GLOB FILES *.cu)
77
add_executable(kernel_float_tests ${FILES})
88
target_link_libraries(kernel_float_tests PRIVATE kernel_float)
99
target_compile_options(kernel_float_tests PRIVATE "--extended-lambda")
10-
set_target_properties(kernel_float_tests PROPERTIES CUDA_ARCHITECTURES "80")
10+
set_target_properties(kernel_float_tests PROPERTIES CUDA_ARCHITECTURES "70;80")
1111

1212
target_compile_options(kernel_float_tests PRIVATE "-ftime-report -ftime-report-details")
1313

‎tests/binops.cu‎

Lines changed: 16 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -29,23 +29,25 @@ struct binops_tests {
2929
// ASSERT(equals(T(x[I] % y[I]), c[I]) && ...);
3030

3131
// Comparison
32+
auto from_bool = [](bool x) { return x ? T(1.0) : T(0.0); };
33+
3234
c = a < b;
33-
ASSERT(equals(T(x[I] < y[I]), c[I]) && ...);
35+
ASSERT(equals(from_bool(x[I] < y[I]), c[I]) && ...);
3436

3537
c = a > b;
36-
ASSERT(equals(T(x[I] > y[I]), c[I]) && ...);
38+
ASSERT(equals(from_bool(x[I] > y[I]), c[I]) && ...);
3739

3840
c = a <= b;
39-
ASSERT(equals(T(x[I] <= y[I]), c[I]) && ...);
41+
ASSERT(equals(from_bool(x[I] <= y[I]), c[I]) && ...);
4042

4143
c = a >= b;
42-
ASSERT(equals(T(x[I] >= y[I]), c[I]) && ...);
44+
ASSERT(equals(from_bool(x[I] >= y[I]), c[I]) && ...);
4345

4446
c = a == b;
45-
ASSERT(equals(T(x[I] == y[I]), c[I]) && ...);
47+
ASSERT(equals(from_bool(x[I] == y[I]), c[I]) && ...);
4648

4749
c = a != b;
48-
ASSERT(equals(T(x[I] != y[I]), c[I]) && ...);
50+
ASSERT(equals(from_bool(x[I] != y[I]), c[I]) && ...);
4951

5052
// Assignment
5153
c = a;
@@ -108,8 +110,11 @@ struct minmax_tests {
108110
ASSERT(equals(fminf(a[I], b[I]), lo[I]) && ...);
109111
ASSERT(equals(fmaxf(a[I], b[I]), hi[I]) && ...);
110112
} else if constexpr (is_one_of<T, __half, __nv_bfloat16>) {
113+
// __hmin/__hmax are only supported in CC >= 8
114+
#if KERNEL_FLOAT_CUDA_ARCH >= 800
111115
ASSERT(equals(__hmin(a[I], b[I]), lo[I]) && ...);
112116
ASSERT(equals(__hmax(a[I], b[I]), hi[I]) && ...);
117+
#endif
113118
} else {
114119
ASSERT(equals(x[I] < y[I] ? x[I] : y[I], lo[I]) && ...);
115120
ASSERT(equals(x[I] < y[I] ? y[I] : x[I], hi[I]) && ...);
@@ -123,13 +128,13 @@ REGISTER_TEST_CASE_GPU("min/max functions", minmax_tests, __half, __nv_bfloat16)
123128
struct cross_test {
124129
template<typename T>
125130
__host__ __device__ void operator()(generator<T> gen) {
126-
kf::vec<T, 3> a = {1, 2, 3};
127-
kf::vec<T, 3> b = {4, 5, 6};
131+
kf::vec<T, 3> a = {T(1.0), T(2.0), T(3.0)};
132+
kf::vec<T, 3> b = {T(4.0), T(5.0), T(6.0)};
128133
kf::vec<T, 3> c = cross(a, b);
129134

130-
ASSERT(c[0] == T(-3));
131-
ASSERT(c[1] == T(6));
132-
ASSERT(c[2] == T(-3));
135+
ASSERT(c[0] == T(-3.0));
136+
ASSERT(c[1] == T(6.0));
137+
ASSERT(c[2] == T(-3.0));
133138
}
134139
};
135140

‎tests/constant.cu‎

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -8,20 +8,20 @@ struct constant_tests {
88
T value = gen.next();
99
kf::vec<T, 2> vector = {gen.next(), gen.next()};
1010

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));
11+
ASSERT_EQ(kf::make_constant(5.0) + value, T(5.0) + value);
12+
ASSERT_EQ(value + kf::make_constant(5.0), value + T(5.0));
13+
ASSERT_EQ(kf::make_constant(5.0) + vector, T(5.0) + vector);
14+
ASSERT_EQ(vector + kf::make_constant(5.0), vector + T(5.0));
1515

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));
16+
ASSERT_EQ(kf::make_constant(5.0) - value, T(5.0) - value);
17+
ASSERT_EQ(value - kf::make_constant(5.0), value - T(5.0));
18+
ASSERT_EQ(kf::make_constant(5.0) - vector, T(5.0) - vector);
19+
ASSERT_EQ(vector - kf::make_constant(5.0), vector - T(5.0));
2020

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));
21+
ASSERT_EQ(kf::make_constant(5.0) * value, T(5.0) * value);
22+
ASSERT_EQ(value * kf::make_constant(5.0), value * T(5.0));
23+
ASSERT_EQ(kf::make_constant(5.0) * vector, T(5.0) * vector);
24+
ASSERT_EQ(vector * kf::make_constant(5.0), vector * T(5.0));
2525

2626
// These results in division by zero for integers
2727
// ASSERT_EQ(kf::make_constant(5.0) / value, T(5) / value);

‎tests/memory.cu‎

Lines changed: 49 additions & 49 deletions
Original file line numberDiff line numberDiff line change
@@ -3,40 +3,40 @@
33
struct load_test {
44
template<typename T>
55
__host__ __device__ void operator()(generator<T> gen) {
6-
T data[8] = {T(0), T(1), T(2), T(3), T(4), T(5), T(6), T(7)};
6+
T data[8] = {T(0.0), T(1.0), T(2.0), T(3.0), T(4.0), T(5.0), T(6.0), T(7.0)};
77

88
{
9-
auto expected = kf::make_vec(T(3), T(2), T(7));
9+
auto expected = kf::make_vec(T(3.0), T(2.0), T(7.0));
1010
auto output = kf::load(data, kf::make_vec(3, 2, 7));
1111
ASSERT_EQ(expected, output);
1212
}
1313

1414
{
15-
auto expected = kf::make_vec(T(3), T(2), T(7));
15+
auto expected = kf::make_vec(T(3.0), T(2.0), T(7.0));
1616
auto output = kf::load(data, kf::make_vec(3, 2, 7), kf::make_vec(true, true, true));
1717
ASSERT_EQ(expected, output);
1818
}
1919

2020
{
21-
auto expected = kf::make_vec(T(3), T(), T(7));
21+
auto expected = kf::make_vec(T(3.0), T(), T(7.0));
2222
auto output = kf::load(data, kf::make_vec(3, 100, 7), kf::make_vec(true, false, true));
2323
ASSERT_EQ(expected, output);
2424
}
2525

2626
{
27-
auto expected = kf::make_vec(T(0), T(1), T(2));
27+
auto expected = kf::make_vec(T(0.0), T(1.0), T(2.0));
2828
auto output = kf::loadn<3>(data);
2929
ASSERT_EQ(expected, output);
3030
}
3131

3232
{
33-
auto expected = kf::make_vec(T(2), T(3), T(4));
33+
auto expected = kf::make_vec(T(2.0), T(3.0), T(4.0));
3434
auto output = kf::loadn<3>(data, 2);
3535
ASSERT_EQ(expected, output);
3636
}
3737

3838
{
39-
auto expected = kf::make_vec(T(6), T(7), T());
39+
auto expected = kf::make_vec(T(6.0), T(7.0), T());
4040
auto output = kf::loadn<3>(data, 6, 8);
4141
ASSERT_EQ(expected, output);
4242
}
@@ -50,78 +50,78 @@ struct store_test {
5050
template<typename T>
5151
__host__ __device__ void operator()(generator<T> gen) {
5252
{
53-
T data[4] = {T(0), T(1), T(2), T(3)};
54-
auto values = kf::make_vec(T(100), T(200));
53+
T data[4] = {T(0.0), T(1.0), T(2.0), T(3.0)};
54+
auto values = kf::make_vec(T(100.0), T(200.0));
5555
auto offsets = kf::make_vec(1, 3);
5656
kf::store(values, data, offsets);
57-
ASSERT_EQ(data[0], T(0));
58-
ASSERT_EQ(data[1], T(100));
59-
ASSERT_EQ(data[2], T(2));
60-
ASSERT_EQ(data[3], T(200));
57+
ASSERT_EQ(data[0], T(0.0));
58+
ASSERT_EQ(data[1], T(100.0));
59+
ASSERT_EQ(data[2], T(2.0));
60+
ASSERT_EQ(data[3], T(200.0));
6161
}
6262

6363
{
64-
T data[4] = {T(0), T(1), T(2), T(3)};
65-
auto values = kf::make_vec(T(100), T(200));
64+
T data[4] = {T(0.0), T(1.0), T(2.0), T(3.0)};
65+
auto values = kf::make_vec(T(100.0), T(200.0));
6666
auto offsets = kf::make_vec(1, 3);
6767
auto mask = kf::make_vec(true, true);
6868
kf::store(values, data, offsets, mask);
69-
ASSERT_EQ(data[0], T(0));
70-
ASSERT_EQ(data[1], T(100));
71-
ASSERT_EQ(data[2], T(2));
72-
ASSERT_EQ(data[3], T(200));
69+
ASSERT_EQ(data[0], T(0.0));
70+
ASSERT_EQ(data[1], T(100.0));
71+
ASSERT_EQ(data[2], T(2.0));
72+
ASSERT_EQ(data[3], T(200.0));
7373
}
7474

7575
{
76-
T data[4] = {T(0), T(1), T(2), T(3)};
77-
auto values = kf::make_vec(T(100), T(200));
76+
T data[4] = {T(0.0), T(1.0), T(2.0), T(3.0)};
77+
auto values = kf::make_vec(T(100.0), T(200.0));
7878
auto offsets = kf::make_vec(1, 3);
7979
auto mask = kf::make_vec(true, false);
8080
kf::store(values, data, offsets, mask);
81-
ASSERT_EQ(data[0], T(0));
82-
ASSERT_EQ(data[1], T(100));
83-
ASSERT_EQ(data[2], T(2));
84-
ASSERT_EQ(data[3], T(3));
81+
ASSERT_EQ(data[0], T(0.0));
82+
ASSERT_EQ(data[1], T(100.0));
83+
ASSERT_EQ(data[2], T(2.0));
84+
ASSERT_EQ(data[3], T(3.0));
8585
}
8686

8787
{
88-
T data[4] = {T(0), T(1), T(2), T(3)};
89-
auto values = kf::make_vec(T(100), T(200));
88+
T data[4] = {T(0.0), T(1.0), T(2.0), T(3.0)};
89+
auto values = kf::make_vec(T(100.0), T(200.0));
9090
kf::storen(values, data);
91-
ASSERT_EQ(data[0], T(100));
92-
ASSERT_EQ(data[1], T(200));
93-
ASSERT_EQ(data[2], T(2));
94-
ASSERT_EQ(data[3], T(3));
91+
ASSERT_EQ(data[0], T(100.0));
92+
ASSERT_EQ(data[1], T(200.0));
93+
ASSERT_EQ(data[2], T(2.0));
94+
ASSERT_EQ(data[3], T(3.0));
9595
}
9696

9797
{
98-
T data[4] = {T(0), T(1), T(2), T(3)};
99-
auto values = kf::make_vec(T(100), T(200));
98+
T data[4] = {T(0.0), T(1.0), T(2.0), T(3.0)};
99+
auto values = kf::make_vec(T(100.0), T(200.0));
100100
kf::storen(values, data, 1);
101-
ASSERT_EQ(data[0], T(0));
102-
ASSERT_EQ(data[1], T(100));
103-
ASSERT_EQ(data[2], T(200));
104-
ASSERT_EQ(data[3], T(3));
101+
ASSERT_EQ(data[0], T(0.0));
102+
ASSERT_EQ(data[1], T(100.0));
103+
ASSERT_EQ(data[2], T(200.0));
104+
ASSERT_EQ(data[3], T(3.0));
105105
}
106106

107107
{
108-
T data[4] = {T(0), T(1), T(2), T(3)};
109-
auto values = kf::make_vec(T(100), T(200));
108+
T data[4] = {T(0.0), T(1.0), T(2.0), T(3.0)};
109+
auto values = kf::make_vec(T(100.0), T(200.0));
110110
kf::storen(values, data, 1, 4);
111-
ASSERT_EQ(data[0], T(0));
112-
ASSERT_EQ(data[1], T(100));
113-
ASSERT_EQ(data[2], T(200));
114-
ASSERT_EQ(data[3], T(3));
111+
ASSERT_EQ(data[0], T(0.0));
112+
ASSERT_EQ(data[1], T(100.0));
113+
ASSERT_EQ(data[2], T(200.0));
114+
ASSERT_EQ(data[3], T(3.0));
115115
}
116116

117117
{
118-
T data[4] = {T(0), T(1), T(2), T(3)};
119-
auto values = kf::make_vec(T(100), T(200));
118+
T data[4] = {T(0.0), T(1.0), T(2.0), T(3.0)};
119+
auto values = kf::make_vec(T(100.0), T(200.0));
120120
kf::storen(values, data, 3, 4);
121-
ASSERT_EQ(data[0], T(0));
122-
ASSERT_EQ(data[1], T(1));
123-
ASSERT_EQ(data[2], T(2));
124-
ASSERT_EQ(data[3], T(100));
121+
ASSERT_EQ(data[0], T(0.0));
122+
ASSERT_EQ(data[1], T(1.0));
123+
ASSERT_EQ(data[2], T(2.0));
124+
ASSERT_EQ(data[3], T(100.0));
125125
}
126126
}
127127
};

0 commit comments

Comments
 (0)