Skip to content

Commit a4444bb

Browse files
mylermyler
authored
Merge pull request intel#1650 from bb-sycl/xmain
Auto pulldown and update tc files for xmain branch on 20230316
2 parents 3073468 + 6c0ade2 commit a4444bb

File tree

111 files changed

+1431
-457
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

111 files changed

+1431
-457
lines changed

‎SYCL/BFloat16/bfloat16_builtins.cpp‎

Lines changed: 15 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -28,6 +28,8 @@ bool check(float a, float b) {
2828
return fabs(2 * (a - b) / (a + b)) > bf16_eps * 2;
2929
}
3030

31+
bool check(bool a, bool b) { return (a != b); }
32+
3133
#define TEST_BUILTIN_1_SCAL_IMPL(NAME) \
3234
{ \
3335
buffer<float> a_buf(&a[0], N); \
@@ -45,7 +47,7 @@ bool check(float a, float b) {
4547
} \
4648
assert(err == 0);
4749

48-
#define TEST_BUILTIN_1_ARR_IMPL(NAME, SZ) \
50+
#define TEST_BUILTIN_1_ARR_IMPL(NAME, SZ, RETTY) \
4951
{ \
5052
buffer<float, 2> a_buf{range<2>{N / SZ, SZ}}; \
5153
buffer<int> err_buf(&err, 1); \
@@ -58,7 +60,7 @@ bool check(float a, float b) {
5860
for (int i = 0; i < SZ; i++) { \
5961
arg[i] = A[index][i]; \
6062
} \
61-
marray<bfloat16, SZ> res = NAME(arg); \
63+
marray<RETTY, SZ> res = NAME(arg); \
6264
for (int i = 0; i < SZ; i++) { \
6365
if (check(res[i], NAME(A[index][i]))) { \
6466
ERR[0] = 1; \
@@ -69,13 +71,13 @@ bool check(float a, float b) {
6971
} \
7072
assert(err == 0);
7173

72-
#define TEST_BUILTIN_1(NAME) \
74+
#define TEST_BUILTIN_1(NAME, RETTY) \
7375
TEST_BUILTIN_1_SCAL_IMPL(NAME) \
74-
TEST_BUILTIN_1_ARR_IMPL(NAME, 1) \
75-
TEST_BUILTIN_1_ARR_IMPL(NAME, 2) \
76-
TEST_BUILTIN_1_ARR_IMPL(NAME, 3) \
77-
TEST_BUILTIN_1_ARR_IMPL(NAME, 4) \
78-
TEST_BUILTIN_1_ARR_IMPL(NAME, 5)
76+
TEST_BUILTIN_1_ARR_IMPL(NAME, 1, RETTY) \
77+
TEST_BUILTIN_1_ARR_IMPL(NAME, 2, RETTY) \
78+
TEST_BUILTIN_1_ARR_IMPL(NAME, 3, RETTY) \
79+
TEST_BUILTIN_1_ARR_IMPL(NAME, 4, RETTY) \
80+
TEST_BUILTIN_1_ARR_IMPL(NAME, 5, RETTY)
7981

8082
#define TEST_BUILTIN_2_SCAL_IMPL(NAME) \
8183
{ \
@@ -232,14 +234,18 @@ int main() {
232234
c[i] = (float)(3 * i);
233235
}
234236

235-
TEST_BUILTIN_1(fabs);
237+
TEST_BUILTIN_1(fabs, bfloat16);
236238
TEST_BUILTIN_2(fmin);
237239
TEST_BUILTIN_2(fmax);
238240
TEST_BUILTIN_3(fma);
239241

240242
float check_nan = 0;
241243
TEST_BUILTIN_2_NAN(fmin);
242244
TEST_BUILTIN_2_NAN(fmax);
245+
246+
// Insert NAN value in a to test isnan
247+
a[0] = a[N - 1] = NAN;
248+
TEST_BUILTIN_1(isnan, bool);
243249
}
244250
return 0;
245251
}

‎SYCL/Basic/accessor/accessor.cpp‎

Lines changed: 76 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -653,7 +653,7 @@ int main() {
653653
}
654654
}
655655

656-
// placeholder accessor exception // SYCL2020 4.7.6.9
656+
// placeholder accessor exception (1) // SYCL2020 4.7.6.9
657657
{
658658
sycl::queue q;
659659
// host device executes kernels via a different method and there
@@ -669,16 +669,88 @@ int main() {
669669
q.submit([&](sycl::handler &cgh) {
670670
// we do NOT call .require(acc) without which we should throw a
671671
// synchronous exception with errc::kernel_argument
672-
cgh.parallel_for<class ph>(r,
673-
[=](sycl::id<1> index) { acc[index] = 0; });
672+
cgh.parallel_for<class ph1>(r,
673+
[=](sycl::id<1> index) { acc[index] = 0; });
674674
});
675675
q.wait_and_throw();
676676
assert(false && "we should not be here, missing exception");
677677
} catch (sycl::exception &e) {
678678
std::cout << "exception received: " << e.what() << std::endl;
679679
assert(e.code() == sycl::errc::kernel_argument && "incorrect error code");
680680
} catch (...) {
681-
std::cout << "some other exception" << std::endl;
681+
std::cout << "Some other exception (line " << __LINE__ << ")"
682+
<< std::endl;
683+
return 1;
684+
}
685+
}
686+
687+
// placeholder accessor exception (2) // SYCL2020 4.7.6.9
688+
{
689+
sycl::queue q;
690+
// host device executes kernels via a different method and there
691+
// is no good way to throw an exception at this time.
692+
sycl::range<1> r(4);
693+
sycl::buffer<int, 1> b(r);
694+
try {
695+
using AccT = sycl::accessor<int, 1, sycl::access::mode::read_write,
696+
sycl::access::target::device,
697+
sycl::access::placeholder::true_t>;
698+
AccT acc(b);
699+
700+
q.submit([&](sycl::handler &cgh) {
701+
// we do NOT call .require(acc) without which we should throw a
702+
// synchronous exception with errc::kernel_argument
703+
// The difference with the previous test is that the use of acc
704+
// is usually optimized away for this particular scenario, but the
705+
// exception should be thrown because of passing it, not because of
706+
// using it
707+
cgh.single_task<class ph2>([=] { int x = acc[0]; });
708+
});
709+
q.wait_and_throw();
710+
assert(false && "we should not be here, missing exception");
711+
} catch (sycl::exception &e) {
712+
std::cout << "exception received: " << e.what() << std::endl;
713+
assert(e.code() == sycl::errc::kernel_argument && "incorrect error code");
714+
} catch (...) {
715+
std::cout << "Some other exception (line " << __LINE__ << ")"
716+
<< std::endl;
717+
return 1;
718+
}
719+
}
720+
721+
// placeholder accessor exception (3) // SYCL2020 4.7.6.9
722+
{
723+
sycl::queue q;
724+
// host device executes kernels via a different method and there
725+
// is no good way to throw an exception at this time.
726+
sycl::range<1> r(4);
727+
sycl::buffer<int, 1> b(r);
728+
try {
729+
using AccT = sycl::accessor<int, 1, sycl::access::mode::read_write,
730+
sycl::access::target::device,
731+
sycl::access::placeholder::true_t>;
732+
AccT acc(b);
733+
734+
q.submit([&](sycl::handler &cgh) {
735+
AccT acc2(b, cgh);
736+
// we do NOT call .require(acc) without which we should throw a
737+
// synchronous exception with errc::kernel_argument
738+
// The particularity of this test is that it passes to a command
739+
// one bound accessor and one unbound accessor. In the past, this
740+
// has led to throw the wrong exception.
741+
cgh.single_task<class ph3>([=] {
742+
volatile int x = acc[0];
743+
volatile int y = acc2[0];
744+
});
745+
});
746+
q.wait_and_throw();
747+
assert(false && "we should not be here, missing exception");
748+
} catch (sycl::exception &e) {
749+
std::cout << "exception received: " << e.what() << std::endl;
750+
assert(e.code() == sycl::errc::kernel_argument && "incorrect error code");
751+
} catch (...) {
752+
std::cout << "Some other exception (line " << __LINE__ << ")"
753+
<< std::endl;
682754
return 1;
683755
}
684756
}

‎SYCL/Basic/buffer/buffer_release.cpp‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,6 @@
1+
// UNSUPPORTED: windows
2+
// DeferredMemory Destruction not presently supported on Windows.
3+
14
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
25
// RUN: %CPU_RUN_PLACEHOLDER %t.out
36

‎SYCL/Basic/built-ins.cpp‎

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -8,9 +8,6 @@
88
// RUN: %GPU_RUN_PLACEHOLDER %t_nonvar.out %GPU_CHECK_PLACEHOLDER
99
// RUN: %ACC_RUN_PLACEHOLDER %t_nonvar.out %ACC_CHECK_PLACEHOLDER
1010

11-
// CUDA does not support printf.
12-
// UNSUPPORTED: cuda
13-
//
1411
// Hits an assertion with AMD:
1512
// XFAIL: hip_amd
1613

0 commit comments

Comments
 (0)