Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit f6dffae

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Prevent use of fp64 and fp16 when unsupported in more tests (#1353)
A number of tests currently use double and sycl::half despite them not being supported by the used device. Although it does not necessarily fail to run on those devices, they are not considered supported by the SYCL specification and should not be used in testing on those devices. This commit prevent parts of a selection of tests from running on devices that do not support the types used. Signed-off-by: Larsen, Steffen <[email protected]>
1 parent 26e147d commit f6dffae

33 files changed

+550
-251
lines changed

‎SYCL/Basic/bit_cast/bit_cast.cpp

Lines changed: 22 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
1+
// RUN: %clangxx -fsycl-device-code-split=per_kernel -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
22
// RUN: %CPU_RUN_PLACEHOLDER %t.out
33
// RUN: %GPU_RUN_PLACEHOLDER %t.out
44
// RUN: %ACC_RUN_PLACEHOLDER %t.out
@@ -13,11 +13,11 @@ constexpr sycl::access::mode sycl_write = sycl::access::mode::write;
1313

1414
template <typename To, typename From> class BitCastKernel;
1515

16-
template <typename To, typename From> To doBitCast(const From &ValueToConvert) {
16+
template <typename To, typename From>
17+
To doBitCast(sycl::queue Queue, const From &ValueToConvert) {
1718
std::vector<To> Vec(1);
1819
{
1920
sycl::buffer<To, 1> Buf(Vec.data(), 1);
20-
sycl::queue Queue;
2121
Queue.submit([&](sycl::handler &cgh) {
2222
auto acc = Buf.template get_access<sycl_write>(cgh);
2323
cgh.single_task<class BitCastKernel<To, From>>([=]() {
@@ -28,8 +28,10 @@ template <typename To, typename From> To doBitCast(const From &ValueToConvert) {
2828
return Vec[0];
2929
}
3030

31-
template <typename To, typename From> int test(const From &Value) {
32-
auto ValueConvertedTwoTimes = doBitCast<From>(doBitCast<To>(Value));
31+
template <typename To, typename From>
32+
int test(sycl::queue Queue, const From &Value) {
33+
auto ValueConvertedTwoTimes =
34+
doBitCast<From>(Queue, doBitCast<To>(Queue, Value));
3335
bool isOriginalValueEqualsToConvertedTwoTimes = false;
3436
if (std::is_integral<From>::value) {
3537
isOriginalValueEqualsToConvertedTwoTimes = Value == ValueConvertedTwoTimes;
@@ -54,31 +56,34 @@ template <typename To, typename From> int test(const From &Value) {
5456
}
5557

5658
int main() {
59+
sycl::queue Queue;
5760
int ReturnCode = 0;
5861

59-
std::cout << "sycl::half to unsigned short ...\n";
60-
ReturnCode += test<unsigned short>(sycl::half(1.0f));
62+
if (Queue.get_device().has(sycl::aspect::fp16)) {
63+
std::cout << "sycl::half to unsigned short ...\n";
64+
ReturnCode += test<unsigned short>(Queue, sycl::half(1.0f));
6165

62-
std::cout << "unsigned short to sycl::half ...\n";
63-
ReturnCode += test<sycl::half>(static_cast<unsigned short>(16384));
66+
std::cout << "unsigned short to sycl::half ...\n";
67+
ReturnCode += test<sycl::half>(Queue, static_cast<unsigned short>(16384));
6468

65-
std::cout << "sycl::half to short ...\n";
66-
ReturnCode += test<short>(sycl::half(1.0f));
69+
std::cout << "sycl::half to short ...\n";
70+
ReturnCode += test<short>(Queue, sycl::half(1.0f));
6771

68-
std::cout << "short to sycl::half ...\n";
69-
ReturnCode += test<sycl::half>(static_cast<short>(16384));
72+
std::cout << "short to sycl::half ...\n";
73+
ReturnCode += test<sycl::half>(Queue, static_cast<short>(16384));
74+
}
7075

7176
std::cout << "int to float ...\n";
72-
ReturnCode += test<float>(static_cast<int>(2));
77+
ReturnCode += test<float>(Queue, static_cast<int>(2));
7378

7479
std::cout << "float to int ...\n";
75-
ReturnCode += test<int>(static_cast<float>(-2.4f));
80+
ReturnCode += test<int>(Queue, static_cast<float>(-2.4f));
7681

7782
std::cout << "unsigned int to float ...\n";
78-
ReturnCode += test<float>(static_cast<unsigned int>(6));
83+
ReturnCode += test<float>(Queue, static_cast<unsigned int>(6));
7984

8085
std::cout << "float to unsigned int ...\n";
81-
ReturnCode += test<unsigned int>(static_cast<float>(-2.4f));
86+
ReturnCode += test<unsigned int>(Queue, static_cast<float>(-2.4f));
8287

8388
return ReturnCode;
8489
}

‎SYCL/Basic/built-ins.cpp

Lines changed: 18 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,11 @@
33
// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
44
// RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
55

6+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -D__SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ %s -o %t_nonvar.out
7+
// RUN: %CPU_RUN_PLACEHOLDER %t_nonvar.out %CPU_CHECK_PLACEHOLDER
8+
// RUN: %GPU_RUN_PLACEHOLDER %t_nonvar.out %GPU_CHECK_PLACEHOLDER
9+
// RUN: %ACC_RUN_PLACEHOLDER %t_nonvar.out %ACC_CHECK_PLACEHOLDER
10+
611
// CUDA does not support printf.
712
// UNSUPPORTED: cuda
813
//
@@ -28,15 +33,25 @@ static const CONSTANT char format[] = "Hello, World! %d %f\n";
2833
int main() {
2934
s::queue q{};
3035

36+
#ifndef __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
37+
if (!q.get_device().has(sycl::aspect::fp64)) {
38+
std::cout
39+
<< "Test without __SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ defined is "
40+
"skipped because the device did not have fp64."
41+
<< std::endl;
42+
return 0;
43+
}
44+
#endif
45+
3146
// Test printf
3247
q.submit([&](s::handler &CGH) {
3348
CGH.single_task<class printf>([=]() {
34-
s::ext::oneapi::experimental::printf(format, 123, 1.23);
49+
s::ext::oneapi::experimental::printf(format, 123, 1.23f);
3550
// CHECK: {{(Hello, World! 123 1.23)?}}
3651
});
3752
}).wait();
3853

39-
s::ext::oneapi::experimental::printf(format, 321, 3.21);
54+
s::ext::oneapi::experimental::printf(format, 321, 3.21f);
4055
// CHECK: {{(Hello, World! 123 1.23)?}}
4156

4257
// Test common
@@ -47,7 +62,7 @@ int main() {
4762
auto AccMin = BufMin.get_access<s::access::mode::write>(cgh);
4863
auto AccMax = BufMax.get_access<s::access::mode::write>(cgh);
4964
cgh.single_task<class common>([=]() {
50-
AccMax[0] = s::max(s::cl_float2{0.5f, 2.5}, s::cl_float2{2.3f, 2.3});
65+
AccMax[0] = s::max(s::cl_float2{0.5f, 2.5f}, s::cl_float2{2.3f, 2.3f});
5166
AccMin[0] = s::min(s::cl_float{0.5f}, s::cl_float{2.3f});
5267
});
5368
});

‎SYCL/Basic/half_builtins.cpp

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -165,6 +165,14 @@ template <int N> bool check(vec<float, N> a, vec<float, N> b) {
165165

166166
int main() {
167167
queue q;
168+
169+
if (!q.get_device().has(sycl::aspect::fp16)) {
170+
std::cout
171+
<< "Test was skipped because the selected device does not support fp16"
172+
<< std::endl;
173+
return 0;
174+
}
175+
168176
float16 a, b, c, d;
169177
for (int i = 0; i < SZ_max; i++) {
170178
a[i] = i / (float)SZ_max;

‎SYCL/Basic/scalar_vec_access.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,7 @@ int main() {
2929
// Test that it is possible to get a reference to single element of the
3030
// vector type. This behavior could possibly change in the future, this
3131
// test is necessary to track that.
32-
float4_t my_float4 = {0.0, 1.0, 2.0, 3.0};
32+
float4_t my_float4 = {0.0f, 1.0f, 2.0f, 3.0f};
3333
float f[4];
3434
for (int i = 0; i < 4; ++i) {
3535
f[i] = reinterpret_cast<float *>(&my_float4)[i];
@@ -40,14 +40,14 @@ int main() {
4040
}
4141

4242
// Test that there is no template resolution error
43-
sycl::float4 a = {1.0, 2.0, 3.0, 4.0};
43+
sycl::float4 a = {1.0f, 2.0f, 3.0f, 4.0f};
4444
out << sycl::native::recip(a.x()) << sycl::endl;
4545
});
4646
});
4747
Q.wait();
4848

4949
// Test that there is no ambiguity in overload resolution.
50-
sycl::float4 a = {1.0, 2.0, 3.0, 4.0};
50+
sycl::float4 a = {1.0f, 2.0f, 3.0f, 4.0f};
5151
std::cout << a.x() << std::endl;
5252

5353
return 0;

‎SYCL/Basic/stream/stream.cpp

Lines changed: 29 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
1+
// RUN: %clangxx -fsycl -fsycl-device-code-split=per_kernel -fsycl-targets=%sycl_triple %s -o %t.out
22
// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
33
// RUN: %GPU_RUN_ON_LINUX_PLACEHOLDER %t.out %GPU_CHECK_ON_LINUX_PLACEHOLDER
44
// RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
@@ -97,26 +97,14 @@ int main() {
9797
// CHECK-NEXT: -12345678901245
9898
// CHECK-NEXT: 12345678901245
9999

100-
// Floating point types
100+
// Floats
101101
Out << 33.4f << endl;
102-
Out << 5.2 << endl;
103102
Out << -33.4f << endl;
104-
Out << -5.2 << endl;
105-
Out << 0.0003 << endl;
106-
Out << -1.0 / 0.0 << endl;
107-
Out << 1.0 / 0.0 << endl;
108-
Out << sycl::sqrt(-1.0) << endl;
109103
Out << -1.0f / 0.0f << endl;
110104
Out << 1.0f / 0.0f << endl;
111105
Out << sycl::sqrt(-1.0f) << endl;
112106
// CHECK-NEXT: 33.4
113-
// CHECK-NEXT: 5.2
114107
// CHECK-NEXT: -33.4
115-
// CHECK-NEXT: -5.2
116-
// CHECK-NEXT: 0.0003
117-
// CHECK-NEXT: -inf
118-
// CHECK-NEXT: inf
119-
// CHECK-NEXT: nan
120108
// CHECK-NEXT: -inf
121109
// CHECK-NEXT: inf
122110
// CHECK-NEXT: nan
@@ -205,6 +193,33 @@ int main() {
205193
});
206194
Queue.wait();
207195

196+
if (Queue.get_device().has(sycl::aspect::fp64)) {
197+
Queue.submit([&](handler &CGH) {
198+
stream Out(1024, 80, CGH);
199+
CGH.single_task<class doubles>([=]() {
200+
// Double
201+
Out << 5.2 << endl;
202+
Out << -5.2 << endl;
203+
Out << 0.0003 << endl;
204+
Out << -1.0 / 0.0 << endl;
205+
Out << 1.0 / 0.0 << endl;
206+
Out << sycl::sqrt(-1.0) << endl;
207+
});
208+
});
209+
Queue.wait();
210+
} else {
211+
// Repeat skipped message same number of times as the number of skipped
212+
// output lines.
213+
for (size_t I = 0; I < 6; ++I)
214+
std::cout << "Skipped double test." << std::endl;
215+
}
216+
// CHECK-NEXT: {{(5.2|Skipped double test.)}}
217+
// CHECK-NEXT: {{(-5.2|Skipped double test.)}}
218+
// CHECK-NEXT: {{(0.0003|Skipped double test.)}}
219+
// CHECK-NEXT: {{(-inf|Skipped double test.)}}
220+
// CHECK-NEXT: {{(inf|Skipped double test.)}}
221+
// CHECK-NEXT: {{(nan|Skipped double test.)}}
222+
208223
// Stream in parallel_for
209224
Queue.submit([&](handler &CGH) {
210225
stream Out(1024, 80, CGH);

‎SYCL/Basic/vector_operators.cpp

Lines changed: 4 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
1+
// RUN: %clangxx -fsycl-device-code-split=per_kernel -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
22
// RUN: %CPU_RUN_PLACEHOLDER %t.out
33
// RUN: %GPU_RUN_PLACEHOLDER %t.out
44
// RUN: %ACC_RUN_PLACEHOLDER %t.out
@@ -38,6 +38,7 @@ template <typename T, int N> void check_vector_size() {
3838
}
3939

4040
int main() {
41+
s::queue Queue;
4142

4243
/* Separate checks for NumElements=1 edge case */
4344

@@ -46,7 +47,6 @@ int main() {
4647
vec_type res;
4748
{
4849
s::buffer<vec_type, 1> Buf(&res, s::range<1>(1));
49-
s::queue Queue;
5050
Queue.submit([&](s::handler &cgh) {
5151
auto Acc = Buf.get_access<s::access::mode::write>(cgh);
5252
cgh.single_task<class isequal_vec_op_1_elem>([=]() {
@@ -67,7 +67,6 @@ int main() {
6767
vec_type res;
6868
{
6969
s::buffer<vec_type, 1> Buf(&res, s::range<1>(1));
70-
s::queue Queue;
7170
Queue.submit([&](s::handler &cgh) {
7271
auto Acc = Buf.get_access<s::access::mode::write>(cgh);
7372
cgh.single_task<class isequal_vec_op_1_elem_scalar>([=]() {
@@ -94,7 +93,6 @@ int main() {
9493
res_vec_type res;
9594
{
9695
s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
97-
s::queue Queue;
9896
Queue.submit([&](s::handler &cgh) {
9997
auto Acc = Buf.get_access<s::access::mode::write>(cgh);
10098
cgh.single_task<class isequal_vec_op>([=]() {
@@ -109,12 +107,11 @@ int main() {
109107
}
110108

111109
// Operator <, cl_double
112-
{
110+
if (Queue.get_device().has(sycl::aspect::fp64)) {
113111
using res_vec_type = s::vec<s::cl_long, 4>;
114112
res_vec_type res;
115113
{
116114
s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
117-
s::queue Queue;
118115
Queue.submit([&](s::handler &cgh) {
119116
auto Acc = Buf.get_access<s::access::mode::write>(cgh);
120117
cgh.single_task<class isless_vec_op>([=]() {
@@ -134,7 +131,6 @@ int main() {
134131
res_vec_type res;
135132
{
136133
s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
137-
s::queue Queue;
138134
Queue.submit([&](s::handler &cgh) {
139135
auto Acc = Buf.get_access<s::access::mode::write>(cgh);
140136
cgh.single_task<class isgreater_vec_op>([=]() {
@@ -149,12 +145,11 @@ int main() {
149145
}
150146

151147
// Operator <=, cl_half
152-
{
148+
if (Queue.get_device().has(sycl::aspect::fp16)) {
153149
using res_vec_type = s::vec<s::cl_short, 4>;
154150
res_vec_type res;
155151
{
156152
s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
157-
s::queue Queue;
158153
Queue.submit([&](s::handler &cgh) {
159154
auto Acc = Buf.get_access<s::access::mode::write>(cgh);
160155
cgh.single_task<class isnotgreater_vec_op>([=]() {
@@ -176,7 +171,6 @@ int main() {
176171
res_vec_type res;
177172
{
178173
s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
179-
s::queue Queue;
180174
Queue.submit([&](s::handler &cgh) {
181175
auto Acc = Buf.get_access<s::access::mode::write>(cgh);
182176
cgh.single_task<class isnotless_vec_op>([=]() {
@@ -196,7 +190,6 @@ int main() {
196190
res_vec_type res;
197191
{
198192
s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
199-
s::queue Queue;
200193
Queue.submit([&](s::handler &cgh) {
201194
auto Acc = Buf.get_access<s::access::mode::write>(cgh);
202195
cgh.single_task<class isnotequal_vec_op>([=]() {
@@ -216,7 +209,6 @@ int main() {
216209
res_vec_type res;
217210
{
218211
s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
219-
s::queue Queue;
220212
Queue.submit([&](s::handler &cgh) {
221213
auto Acc = Buf.get_access<s::access::mode::write>(cgh);
222214
cgh.single_task<class logical_and_vec_op>([=]() {
@@ -236,7 +228,6 @@ int main() {
236228
res_vec_type res;
237229
{
238230
s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
239-
s::queue Queue;
240231
Queue.submit([&](s::handler &cgh) {
241232
auto Acc = Buf.get_access<s::access::mode::write>(cgh);
242233
cgh.single_task<class logical_or_vec_op>([=]() {
@@ -257,7 +248,6 @@ int main() {
257248
res_vec_type res;
258249
{
259250
s::buffer<res_vec_type, 1> Buf(&res, s::range<1>(1));
260-
s::queue Queue;
261251
Queue.submit([&](s::handler &cgh) {
262252
auto Acc = Buf.get_access<s::access::mode::write>(cgh);
263253
cgh.single_task<class as_op>([=]() {

0 commit comments

Comments
 (0)