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

Commit 2211aac

Browse files
JackAKirkJackAKirk
committed
Merge branch 'intel' into bfloat16-class-tests
2 parents e817dba + 258a40d commit 2211aac

33 files changed

+661
-34
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/BFloat16/bfloat16_type_cuda.cpp

Lines changed: 11 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,6 @@
11
// REQUIRES: gpu, cuda
22
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Xsycl-target-backend --cuda-gpu-arch=sm_80 %s -o %t.out
3-
// TODO: Currently the CI does not have a sm_80 capable machine. Enable the test
4-
// execution once it does.
5-
// RUNx: %t.out
3+
// RUN: %t.out
64

75
//==--------- bfloat16_type_cuda.cpp - SYCL bfloat16 type test -------------==//
86
//
@@ -14,4 +12,13 @@
1412

1513
#include "bfloat16_type.hpp"
1614

17-
int main() { return run_tests(); }
15+
int main() {
16+
bool has_bfloat16_aspect = false;
17+
for (const auto &plt : sycl::platform::get_platforms()) {
18+
if (plt.has(aspect::ext_oneapi_bfloat16))
19+
has_bfloat16_aspect = true;
20+
}
21+
22+
if (has_bfloat16_aspect)
23+
return run_tests();
24+
}

‎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/Basic/aspects.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -57,6 +57,9 @@ int main() {
5757
if (plt.has(aspect::fp64)) {
5858
std::cout << " fp64" << std::endl;
5959
}
60+
if (plt.has(aspect::ext_oneapi_bfloat16)) {
61+
std::cout << " ext_oneapi_bfloat16" << std::endl;
62+
}
6063
if (plt.has(aspect::int64_base_atomics)) {
6164
std::cout << " base atomic operations" << std::endl;
6265
}

‎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/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>

‎SYCL/ESIMD/histogram_raw_send.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,7 +6,7 @@
66
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
77
//
88
//===----------------------------------------------------------------------===//
9-
// REQUIRES: gpu
9+
// REQUIRES: gpu-intel-gen9
1010
// UNSUPPORTED: gpu-intel-dg1,cuda,hip
1111
// UNSUPPORTED: ze_debug-1,ze_debug4
1212
// RUN: %clangxx -fsycl %s -o %t.out

0 commit comments

Comments
 (0)