Skip to content

Commit 1d118bf

Browse files
nstesternstester
authored
Auto pulldown and update tc files for xmain branch
Auto pulldown and update tc files for xmain branch on 20220712
2 parents cfc373d + 04ded06 commit 1d118bf

39 files changed

+514
-29
lines changed

‎SYCL/Assert/assert_in_kernels_ndebug.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
1-
// FIXME unsupported on CUDA and HIP until fallback libdevice becomes available
2-
// UNSUPPORTED: cuda || hip
1+
// FIXME unsupported on HIP until fallback libdevice becomes available
2+
// UNSUPPORTED: hip
33
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DNDEBUG %S/assert_in_kernels.cpp -o %t.out
44
// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
55
// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER

‎SYCL/Basic/accessor/accessor.cpp

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -556,6 +556,40 @@ int main() {
556556
return 1;
557557
}
558558
}
559+
560+
// placeholder accessor exception // SYCL2020 4.7.6.9
561+
{
562+
sycl::queue q;
563+
// host device executes kernels via a different method and there
564+
// is no good way to throw an exception at this time.
565+
if (!q.is_host()) {
566+
sycl::range<1> r(4);
567+
sycl::buffer<int, 1> b(r);
568+
try {
569+
sycl::accessor<int, 1, sycl::access::mode::read_write,
570+
sycl::access::target::device,
571+
sycl::access::placeholder::true_t>
572+
acc(b);
573+
574+
q.submit([&](sycl::handler &cgh) {
575+
// we do NOT call .require(acc) without which we should throw a
576+
// synchronous exception with errc::kernel_argument
577+
cgh.parallel_for<class ph>(
578+
r, [=](sycl::id<1> index) { acc[index] = 0; });
579+
});
580+
q.wait_and_throw();
581+
assert(false && "we should not be here, missing exception");
582+
} catch (sycl::exception &e) {
583+
std::cout << "exception received: " << e.what() << std::endl;
584+
assert(e.code() == sycl::errc::kernel_argument &&
585+
"incorrect error code");
586+
} catch (...) {
587+
std::cout << "some other exception" << std::endl;
588+
return 1;
589+
}
590+
}
591+
}
592+
559593
{
560594
try {
561595
int data = -1;

‎SYCL/DeviceLib/built-ins/ext_native_math.cpp

Lines changed: 155 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,155 @@
1+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel %s -o %t.out
2+
// RUN: %HOST_RUN_PLACEHOLDER %t.out
3+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
4+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
5+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
6+
7+
// OpenCL CPU driver does not support cl_khr_fp16 extension for this reason this
8+
// test is compiled with the -fsycl-device-code-split flag
9+
10+
#include <CL/sycl.hpp>
11+
#include <cassert>
12+
13+
template <typename T> void assert_out_of_bound(T val, T lower, T upper) {
14+
assert(sycl::all(lower < val && val < upper));
15+
}
16+
17+
template <>
18+
void assert_out_of_bound<float>(float val, float lower, float upper) {
19+
assert(lower < val && val < upper);
20+
}
21+
22+
template <>
23+
void assert_out_of_bound<sycl::half>(sycl::half val, sycl::half lower,
24+
sycl::half upper) {
25+
assert(lower < val && val < upper);
26+
}
27+
28+
template <typename T>
29+
void native_tanh_tester(sycl::queue q, T val, T up, T lo) {
30+
T r = val;
31+
32+
#ifdef SYCL_EXT_ONEAPI_NATIVE_MATH
33+
{
34+
sycl::buffer<T, 1> BufR(&r, sycl::range<1>(1));
35+
q.submit([&](sycl::handler &cgh) {
36+
auto AccR = BufR.template get_access<sycl::access::mode::read_write>(cgh);
37+
cgh.single_task([=]() {
38+
AccR[0] = sycl::ext::oneapi::experimental::native::tanh(AccR[0]);
39+
});
40+
});
41+
}
42+
43+
assert_out_of_bound(r, up, lo);
44+
#else
45+
assert(!"SYCL_EXT_ONEAPI_NATIVE_MATH not supported");
46+
#endif
47+
}
48+
49+
template <typename T>
50+
void native_exp2_tester(sycl::queue q, T val, T up, T lo) {
51+
T r = val;
52+
53+
#ifdef SYCL_EXT_ONEAPI_NATIVE_MATH
54+
{
55+
sycl::buffer<T, 1> BufR(&r, sycl::range<1>(1));
56+
q.submit([&](sycl::handler &cgh) {
57+
auto AccR = BufR.template get_access<sycl::access::mode::read_write>(cgh);
58+
cgh.single_task([=]() {
59+
AccR[0] = sycl::ext::oneapi::experimental::native::exp2(AccR[0]);
60+
});
61+
});
62+
}
63+
64+
assert_out_of_bound(r, up, lo);
65+
#else
66+
assert(!"SYCL_EXT_ONEAPI_NATIVE_MATH not supported");
67+
#endif
68+
}
69+
70+
int main() {
71+
72+
sycl::queue q;
73+
74+
const double tv[16] = {-2.0, -1.5, -1.0, 0.0, 2.0, 1.5, 1.0, 0.0,
75+
-1.7, 1.7, -1.2, 1.2, -3.0, 3.0, -10.0, 10.0};
76+
const double tl[16] = {-0.97, -0.91, -0.77, -0.1, 0.95, 0.89, 0.75, -0.1,
77+
-0.94, 0.92, -0.84, 0.82, -1.0, 0.98, -1.10, 0.98};
78+
const double tu[16] = {-0.95, -0.89, -0.75, 0.1, 0.97, 0.91, 0.77, 0.1,
79+
-0.92, 0.94, -0.82, 0.84, -0.98, 1.00, -0.98, 1.10};
80+
81+
native_tanh_tester<float>(q, tv[0], tl[0], tu[0]);
82+
native_tanh_tester<sycl::float2>(q, {tv[0], tv[1]}, {tl[0], tl[1]},
83+
{tu[0], tu[1]});
84+
native_tanh_tester<sycl::float3>(
85+
q, {tv[0], tv[1], tv[2]}, {tl[0], tl[1], tl[2]}, {tu[0], tu[1], tu[2]});
86+
native_tanh_tester<sycl::float4>(q, {tv[0], tv[1], tv[2], tv[3]},
87+
{tl[0], tl[1], tl[2], tl[3]},
88+
{tu[0], tu[1], tu[2], tu[3]});
89+
native_tanh_tester<sycl::float8>(
90+
q, {tv[0], tv[1], tv[2], tv[3], tv[4], tv[5], tv[6], tv[7]},
91+
{tl[0], tl[1], tl[2], tl[3], tl[4], tl[5], tl[6], tl[7]},
92+
{tu[0], tu[1], tu[2], tu[3], tu[4], tu[5], tu[6], tu[7]});
93+
native_tanh_tester<sycl::float16>(
94+
q,
95+
{tv[0], tv[1], tv[2], tv[3], tv[4], tv[5], tv[6], tv[7], tv[8], tv[9],
96+
tv[10], tv[11], tv[12], tv[13], tv[14], tv[15]},
97+
{tl[0], tl[1], tl[2], tl[3], tl[4], tl[5], tl[6], tl[7], tl[8], tl[9],
98+
tl[10], tl[11], tl[12], tl[13], tl[14], tl[15]},
99+
{tu[0], tu[1], tu[2], tu[3], tu[4], tu[5], tu[6], tu[7], tu[8], tu[9],
100+
tu[10], tu[11], tu[12], tu[13], tu[14], tu[15]});
101+
102+
if (q.get_device().has(sycl::aspect::fp16)) {
103+
104+
native_tanh_tester<sycl::half>(q, tv[0], tl[0], tu[0]);
105+
native_tanh_tester<sycl::half2>(q, {tv[0], tv[1]}, {tl[0], tl[1]},
106+
{tu[0], tu[1]});
107+
native_tanh_tester<sycl::half3>(
108+
q, {tv[0], tv[1], tv[2]}, {tl[0], tl[1], tl[2]}, {tu[0], tu[1], tu[2]});
109+
native_tanh_tester<sycl::half4>(q, {tv[0], tv[1], tv[2], tv[3]},
110+
{tl[0], tl[1], tl[2], tl[3]},
111+
{tu[0], tu[1], tu[2], tu[3]});
112+
native_tanh_tester<sycl::half8>(
113+
q, {tv[0], tv[1], tv[2], tv[3], tv[4], tv[5], tv[6], tv[7]},
114+
{tl[0], tl[1], tl[2], tl[3], tl[4], tl[5], tl[6], tl[7]},
115+
{tu[0], tu[1], tu[2], tu[3], tu[4], tu[5], tu[6], tu[7]});
116+
native_tanh_tester<sycl::half16>(
117+
q,
118+
{tv[0], tv[1], tv[2], tv[3], tv[4], tv[5], tv[6], tv[7], tv[8], tv[9],
119+
tv[10], tv[11], tv[12], tv[13], tv[14], tv[15]},
120+
{tl[0], tl[1], tl[2], tl[3], tl[4], tl[5], tl[6], tl[7], tl[8], tl[9],
121+
tl[10], tl[11], tl[12], tl[13], tl[14], tl[15]},
122+
{tu[0], tu[1], tu[2], tu[3], tu[4], tu[5], tu[6], tu[7], tu[8], tu[9],
123+
tu[10], tu[11], tu[12], tu[13], tu[14], tu[15]});
124+
125+
const double ev[16] = {-2.0, -1.5, -1.0, 0.0, 2.0, 1.5, 1.0, 0.0,
126+
-2.0, -1.5, -1.0, 0.0, 2.0, 1.5, 1.0, 0.0};
127+
const double el[16] = {0.1, 0.34, 0.4, -0.9, 3.9, 2.7, 1.9, -0.9,
128+
0.1, 0.34, 0.4, -0.9, 3.9, 2.7, 1.9, -0.9};
129+
const double eu[16] = {0.3, 0.36, 0.6, 1.1, 4.1, 2.9, 2.1, 1.1,
130+
0.3, 0.36, 0.6, 1.1, 4.1, 2.9, 2.1, 1.1};
131+
132+
native_exp2_tester<sycl::half>(q, ev[0], el[0], eu[0]);
133+
native_exp2_tester<sycl::half2>(q, {ev[0], ev[1]}, {el[0], el[1]},
134+
{eu[0], eu[1]});
135+
native_exp2_tester<sycl::half3>(
136+
q, {ev[0], ev[1], ev[2]}, {el[0], el[1], el[2]}, {eu[0], eu[1], eu[2]});
137+
native_exp2_tester<sycl::half4>(q, {ev[0], ev[1], ev[2], ev[3]},
138+
{el[0], el[1], el[2], el[3]},
139+
{eu[0], eu[1], eu[2], eu[3]});
140+
native_exp2_tester<sycl::half8>(
141+
q, {ev[0], ev[1], ev[2], ev[3], ev[4], ev[5], ev[6], ev[7]},
142+
{el[0], el[1], el[2], el[3], el[4], el[5], el[6], el[7]},
143+
{eu[0], eu[1], eu[2], eu[3], eu[4], eu[5], eu[6], eu[7]});
144+
native_exp2_tester<sycl::half16>(
145+
q,
146+
{ev[0], ev[1], ev[2], ev[3], ev[4], ev[5], ev[6], ev[7], ev[8], ev[9],
147+
ev[10], ev[11], ev[12], ev[13], ev[14], ev[15]},
148+
{el[0], el[1], el[2], el[3], el[4], el[5], el[6], el[7], el[8], el[9],
149+
el[10], el[11], el[12], el[13], el[14], el[15]},
150+
{eu[0], eu[1], eu[2], eu[3], eu[4], eu[5], eu[6], eu[7], eu[8], eu[9],
151+
eu[10], eu[11], eu[12], eu[13], eu[14], eu[15]});
152+
}
153+
154+
return 0;
155+
}

‎SYCL/DeviceLib/built-ins/fast-math-flag.cpp

Lines changed: 82 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,82 @@
1+
// RUN: %clangxx -fsycl -ffast-math -fsycl-targets=%sycl_triple %s -o %t.out
2+
// RUN: %HOST_RUN_PLACEHOLDER %t.out
3+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
4+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
5+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
6+
7+
#include <CL/sycl.hpp>
8+
#include <cassert>
9+
10+
#define __TEST_FFMATH_BINARY(func) \
11+
int test_ffmath_##func() { \
12+
sycl::float4 r[2]; \
13+
sycl::float4 val[2] = {{1.0004f, 1e-4f, 1.4f, 14.0f}, \
14+
{1.0004f, 1e-4f, 1.4f, 14.0f}}; \
15+
{ \
16+
sycl::buffer<sycl::float4, 1> output(&r[0], sycl::range<1>(2)); \
17+
sycl::buffer<sycl::float4, 1> input(&val[0], sycl::range<1>(2)); \
18+
sycl::queue q; \
19+
q.submit([&](sycl::handler &cgh) { \
20+
auto AccO = \
21+
output.template get_access<sycl::access::mode::write>(cgh); \
22+
auto AccI = input.template get_access<sycl::access::mode::read>(cgh); \
23+
cgh.single_task([=]() { \
24+
AccO[0] = sycl::func(AccI[0], AccI[1]); \
25+
AccO[1] = sycl::native::func(AccI[0], AccI[1]); \
26+
}); \
27+
}); \
28+
} \
29+
return sycl::all(r[0] == r[1]); \
30+
}
31+
32+
#define __TEST_FFMATH_UNARY(func) \
33+
int test_ffmath_##func() { \
34+
sycl::float4 val = {1.0004f, 1e-4f, 1.4f, 14.0f}; \
35+
sycl::float4 r[2]; \
36+
{ \
37+
sycl::buffer<sycl::float4, 1> output(&r[0], sycl::range<1>(2)); \
38+
sycl::buffer<sycl::float4, 1> input(&val, sycl::range<1>(1)); \
39+
sycl::queue q; \
40+
q.submit([&](sycl::handler &cgh) { \
41+
auto AccO = \
42+
output.template get_access<sycl::access::mode::write>(cgh); \
43+
auto AccI = input.template get_access<sycl::access::mode::read>(cgh); \
44+
cgh.single_task([=]() { \
45+
AccO[0] = sycl::func(AccI[0]); \
46+
AccO[1] = sycl::native::func(AccI[0]); \
47+
}); \
48+
}); \
49+
} \
50+
return sycl::all(r[0] == r[1]); \
51+
}
52+
53+
__TEST_FFMATH_UNARY(cos)
54+
__TEST_FFMATH_UNARY(exp)
55+
__TEST_FFMATH_UNARY(exp2)
56+
__TEST_FFMATH_UNARY(exp10)
57+
__TEST_FFMATH_UNARY(log)
58+
__TEST_FFMATH_UNARY(log2)
59+
__TEST_FFMATH_UNARY(log10)
60+
__TEST_FFMATH_BINARY(powr)
61+
__TEST_FFMATH_UNARY(rsqrt)
62+
__TEST_FFMATH_UNARY(sin)
63+
__TEST_FFMATH_UNARY(sqrt)
64+
__TEST_FFMATH_UNARY(tan)
65+
66+
int main() {
67+
68+
assert(test_ffmath_cos());
69+
assert(test_ffmath_exp());
70+
assert(test_ffmath_exp2());
71+
assert(test_ffmath_exp10());
72+
assert(test_ffmath_log());
73+
assert(test_ffmath_log2());
74+
assert(test_ffmath_log10());
75+
assert(test_ffmath_powr());
76+
assert(test_ffmath_rsqrt());
77+
assert(test_ffmath_sin());
78+
assert(test_ffmath_sqrt());
79+
assert(test_ffmath_tan());
80+
81+
return 0;
82+
}

‎SYCL/DeviceLib/cmath_fp64_test.cpp

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

6+
// RUN: %clangxx -fsycl -fsycl-device-lib-jit-link %s -o %t.out
7+
// RUN: %HOST_RUN_PLACEHOLDER %t.out
8+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
9+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
10+
611
#include "math_utils.hpp"
712
#include <CL/sycl.hpp>
813
#include <cmath>

‎SYCL/DeviceLib/cmath_test.cpp

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

6+
// RUN: %clangxx -fsycl -fno-builtin -fsycl-device-lib-jit-link %s -o %t.out
7+
// RUN: %HOST_RUN_PLACEHOLDER %t.out
8+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
9+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
10+
611
#include "math_utils.hpp"
712
#include <CL/sycl.hpp>
813
#include <cmath>

‎SYCL/DeviceLib/math_fp64_test.cpp

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

6+
// RUN: %clangxx -fsycl -fsycl-device-lib-jit-link %s -o %t.out
7+
// RUN: %HOST_RUN_PLACEHOLDER %t.out
8+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
9+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
10+
611
#include "math_utils.hpp"
712
#include <CL/sycl.hpp>
813
#include <cstdint>

‎SYCL/DeviceLib/math_test.cpp

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

6+
// RUN: %clangxx -fsycl -fsycl-device-lib-jit-link %s -o %t.out
7+
// RUN: %HOST_RUN_PLACEHOLDER %t.out
8+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
9+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
10+
611
#include "math_utils.hpp"
712
#include <CL/sycl.hpp>
813
#include <cstdint>

‎SYCL/DeviceLib/std_complex_math_fp64_test.cpp

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

6+
// RUN: %clangxx -fsycl -fsycl-device-lib-jit-link %s -o %t.out
7+
// RUN: %HOST_RUN_PLACEHOLDER %t.out
8+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
9+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
10+
611
#include <CL/sycl.hpp>
712
#include <array>
813
#include <cassert>

‎SYCL/DeviceLib/std_complex_math_test.cpp

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

6+
// RUN: %clangxx -fsycl -fsycl-device-lib-jit-link %s -o %t.out
7+
// RUN: %HOST_RUN_PLACEHOLDER %t.out
8+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
9+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
10+
611
#include <CL/sycl.hpp>
712
#include <array>
813
#include <cassert>

‎SYCL/DeviceLib/string_test.cpp

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

6+
// RUN: %clangxx -fsycl -fno-builtin -fsycl-device-lib-jit-link %s -o %t.out
7+
// RUN: %HOST_RUN_PLACEHOLDER %t.out
8+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
9+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
10+
611
#include <CL/sycl.hpp>
712
#include <cassert>
813
#include <cstdint>

0 commit comments

Comments
 (0)