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

Commit cbccaf7

Browse files
mylermyler
authored
Merge branch 'intel' into split_double_tests
2 parents 2fbfbbe + d89605a commit cbccaf7

Some content is hidden

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

51 files changed

+460
-145
lines changed

‎SYCL/Assert/assert_in_simultaneously_multiple_tus_one_ndebug.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,5 @@
1+
// FIXME flaky fail on CUDA
2+
// UNSUPPORTED: cuda
13
// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -DDEFINE_NDEBUG_INFILE2 -I %S/Inputs %S/assert_in_simultaneously_multiple_tus.cpp %S/Inputs/kernels_in_file2.cpp -o %t.out %threads_lib
24
// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
35
// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt

‎SYCL/AtomicRef/add.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void add_fetch_local_test(queue q, size_t N) {
3030
auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/and.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void and_local_test(queue q) {
3030
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/atomic_memory_order_acq_rel.cpp

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -76,8 +76,7 @@ template <memory_order order> void test_acquire_local() {
7676
q.submit([&](handler &cgh) {
7777
auto error =
7878
error_buf.template get_access<access::mode::read_write>(cgh);
79-
accessor<int, 1, access::mode::read_write, access::target::local> val(
80-
2, cgh);
79+
local_accessor<int, 1> val(2, cgh);
8180
cgh.parallel_for(
8281
nd_range<1>(global_size, local_size), [=](nd_item<1> it) {
8382
size_t lid = it.get_local_id(0);
@@ -168,8 +167,7 @@ template <memory_order order> void test_release_local() {
168167
q.submit([&](handler &cgh) {
169168
auto error =
170169
error_buf.template get_access<access::mode::read_write>(cgh);
171-
accessor<int, 1, access::mode::read_write, access::target::local> val(
172-
2, cgh);
170+
local_accessor<int, 1> val(2, cgh);
173171
cgh.parallel_for(
174172
nd_range<1>(global_size, local_size), [=](nd_item<1> it) {
175173
size_t lid = it.get_local_id(0);

‎SYCL/AtomicRef/atomic_memory_order_seq_cst.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -120,8 +120,7 @@ template <memory_order order> void test_local() {
120120

121121
q.submit([&](handler &cgh) {
122122
auto res = res_buf.template get_access<access::mode::discard_write>(cgh);
123-
accessor<int, 1, access::mode::read_write, access::target::local> val(2,
124-
cgh);
123+
local_accessor<int, 1> val(2, cgh);
125124
cgh.parallel_for(nd_range<1>(N_items, N_items), [=](nd_item<1> it) {
126125
val[0] = 0;
127126
it.barrier(access::fence_space::local_space);

‎SYCL/AtomicRef/compare_exchange.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -32,8 +32,7 @@ void compare_exchange_local_test(queue q, size_t N) {
3232
cgh);
3333
auto out =
3434
output_buf.template get_access<access::mode::discard_write>(cgh);
35-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
36-
cgh);
35+
local_accessor<T, 1> loc(1, cgh);
3736

3837
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3938
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/exchange.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void exchange_local_test(queue q, size_t N) {
3030
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/load.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -31,8 +31,7 @@ void load_local_test(queue q, size_t N) {
3131
auto ld = load_buf.template get_access<access::mode::read_write>(cgh);
3232
auto out =
3333
output_buf.template get_access<access::mode::discard_write>(cgh);
34-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
35-
cgh);
34+
local_accessor<T, 1> loc(1, cgh);
3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);
3837
if (gid == 0)

‎SYCL/AtomicRef/max.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void max_local_test(queue q, size_t N) {
3030
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/min.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void min_local_test(queue q, size_t N) {
3030
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/or.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void or_local_test(queue q) {
3030
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/store.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -51,8 +51,7 @@ void store_local_test(queue q, size_t N) {
5151
buffer<T> store_buf(&store, 1);
5252
q.submit([&](handler &cgh) {
5353
auto st = store_buf.template get_access<access::mode::read_write>(cgh);
54-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
55-
cgh);
54+
local_accessor<T, 1> loc(1, cgh);
5655
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
5756
size_t gid = it.get_global_id(0);
5857
auto atm = AtomicRef<T, memory_order::relaxed, scope, space>(loc[0]);

‎SYCL/AtomicRef/sub.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void sub_fetch_local_test(queue q, size_t N) {
3030
auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/xor.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void xor_local_test(queue q) {
3030
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/Basic/device_event.cpp

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,13 @@
55
// TODO: nd_item::barrier() is not implemented on HOST
66
// RUNx: %HOST_RUN_PLACEHOLDER %t.run
77
//
8+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DUSE_DEPRECATED_LOCAL_ACC %s -o %t.run
9+
// RUN: %GPU_RUN_PLACEHOLDER %t.run
10+
// RUN: %CPU_RUN_PLACEHOLDER %t.run
11+
// RUN: %ACC_RUN_PLACEHOLDER %t.run
12+
// TODO: nd_item::barrier() is not implemented on HOST
13+
// RUNx: %HOST_RUN_PLACEHOLDER %t.run
14+
//
815
// Returns error "Barrier is not supported on the host device
916
// yet." with Nvidia.
1017
// XFAIL: hip_nvidia
@@ -76,8 +83,12 @@ int test_strideN(size_t stride) {
7683

7784
myQueue.submit([&](handler &cgh) {
7885
auto out_ptr = out_buf.get_access<access::mode::write>(cgh);
86+
#ifdef USE_DEPRECATED_LOCAL_ACC
7987
accessor<sycl::cl_int, 1, access::mode::read_write, access::target::local>
8088
local_acc(range<1>(16), cgh);
89+
#else
90+
local_accessor<sycl::cl_int, 1> local_acc(range<1>(16), cgh);
91+
#endif
8192

8293
// Create work-groups with 16 work items in each group.
8394
auto myRange = nd_range<1>(range<1>(nElems), range<1>(workGroupSize));

‎SYCL/Basic/group_async_copy.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -112,8 +112,7 @@ template <typename T> int test(size_t Stride) {
112112
Q.submit([&](handler &CGH) {
113113
auto In = InBuf.template get_access<access::mode::read>(CGH);
114114
auto Out = OutBuf.template get_access<access::mode::write>(CGH);
115-
accessor<T, 1, access::mode::read_write, access::target::local> Local(
116-
range<1>{WorkGroupSize}, CGH);
115+
local_accessor<T, 1> Local(range<1>{WorkGroupSize}, CGH);
117116

118117
nd_range<1> NDR{range<1>(NElems), range<1>(WorkGroupSize)};
119118
CGH.parallel_for<KernelName<T>>(NDR, [=](nd_item<1> NDId) {

‎SYCL/Basic/multi_ptr.cpp

Lines changed: 2 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -71,8 +71,7 @@ template <typename T> void testMultPtr() {
7171
accessor<T, 1, access::mode::read_write, access::target::device,
7272
access::placeholder::false_t>
7373
accessorData_2(bufferData_2, cgh);
74-
accessor<T, 1, access::mode::read_write, access::target::local>
75-
localAccessor(numOfItems, cgh);
74+
local_accessor<T, 1> localAccessor(numOfItems, cgh);
7675

7776
cgh.parallel_for<class testMultPtrKernel<T>>(range<1>{10}, [=](id<1> wiID) {
7877
auto ptr_1 = make_ptr<T, access::address_space::global_space>(
@@ -136,9 +135,7 @@ template <typename T> void testMultPtrArrowOperator() {
136135
accessor<point<T>, 1, access::mode::read, access::target::constant_buffer,
137136
access::placeholder::false_t>
138137
accessorData_2(bufferData_2, cgh);
139-
accessor<point<T>, 1, access::mode::read_write, access::target::local,
140-
access::placeholder::false_t>
141-
accessorData_3(1, cgh);
138+
local_accessor<point<T>, 1> accessorData_3(1, cgh);
142139
accessor<point<T>, 1, access::mode::read, access::target::device,
143140
access::placeholder::false_t>
144141
accessorData_4(bufferData_4, cgh);

‎SYCL/DeviceLib/ITTAnnotations/barrier.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -24,8 +24,7 @@ int main() {
2424
// ITT start/finish annotations and ITT wg_barrier/wi_resume annotations.
2525
q.submit([&](handler &cgh) {
2626
auto acc = buf.get_access<access::mode::read_write>(cgh);
27-
accessor<int, 1, access::mode::read_write, access::target::local>
28-
local_acc(local_range, cgh);
27+
local_accessor<int, 1> local_acc(local_range, cgh);
2928
cgh.parallel_for<class simple_barrier_kernel>(
3029
nd_range<1>(num_items, local_range), [=](nd_item<1> item) {
3130
size_t idx = item.get_global_linear_id();

‎SYCL/DeviceLib/assert.cpp

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
1-
// REQUIRES: cpu,linux
2-
// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl %s -o %t.out
1+
// REQUIRES: (cpu || cuda ) && linux
2+
// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
33
// (see the other RUN lines below; it is a bit complicated)
44
//
55
// assert() call in device code guarantees nothing: on some devices it behaves
@@ -72,6 +72,7 @@
7272
// Overall this sounds stable enough. What could possibly go wrong?
7373
//
7474
// RUN: %CPU_RUN_PLACEHOLDER env SYCL_PI_TRACE=2 SHOULD_CRASH=1 EXPECTED_SIGNAL=SIGABRT %t.out 2>%t.stderr.native
75+
// RUN: %GPU_RUN_PLACEHOLDER env SHOULD_CRASH=1 EXPECTED_SIGNAL=SIGIOT %t.out 2>%t.stderr.native
7576
// RUN: FileCheck %s --input-file %t.stderr.native --check-prefixes=CHECK-MESSAGE || FileCheck %s --input-file %t.stderr.native --check-prefix CHECK-NOTSUPPORTED
7677
//
7778
// Skip the test if the CPU RT doesn't support the extension yet:
@@ -181,6 +182,8 @@ int main() {
181182
expected = SIGABRT;
182183
} else if (0 == strcmp(env, "SIGSEGV")) {
183184
expected = SIGSEGV;
185+
} else if (0 == strcmp(env, "SIGIOT")) {
186+
expected = SIGIOT;
184187
}
185188
if (!expected) {
186189
fprintf(stderr, "EXPECTED_SIGNAL should be set to either \"SIGABRT\", "

‎SYCL/DeviceLib/cmath_test.cpp

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,8 @@
1-
// RUN: %clangxx -fsycl -fno-builtin %s -o %t.out
1+
// UNSUPPORTED: hip
2+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fno-builtin %s -o %t.out
23
// RUN: %HOST_RUN_PLACEHOLDER %t.out
34
// RUN: %CPU_RUN_PLACEHOLDER %t.out
5+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
46
// RUN: %ACC_RUN_PLACEHOLDER %t.out
57

68
// RUN: %clangxx -fsycl -fno-builtin -fsycl-device-lib-jit-link %s -o %t.out

‎SYCL/DeviceLib/math_fp64_test.cpp

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,8 @@
1-
// RUN: %clangxx -fsycl %s -o %t.out
1+
// UNSUPPORTED: hip
2+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
23
// RUN: %HOST_RUN_PLACEHOLDER %t.out
34
// RUN: %CPU_RUN_PLACEHOLDER %t.out
5+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
46
// RUN: %ACC_RUN_PLACEHOLDER %t.out
57

68
// RUN: %clangxx -fsycl -fsycl-device-lib-jit-link %s -o %t.out

‎SYCL/DeviceLib/math_test.cpp

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,8 @@
1-
// RUN: %clangxx -fsycl %s -o %t.out
1+
// UNSUPPORTED: hip
2+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
23
// RUN: %HOST_RUN_PLACEHOLDER %t.out
34
// RUN: %CPU_RUN_PLACEHOLDER %t.out
5+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
46
// RUN: %ACC_RUN_PLACEHOLDER %t.out
57

68
// RUN: %clangxx -fsycl -fsycl-device-lib-jit-link %s -o %t.out

‎SYCL/DeviceLib/string_test.cpp

Lines changed: 4 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,8 @@
1-
// RUN: %clangxx -fsycl -fno-builtin %s -o %t.out
1+
// UNSUPPORTED: hip
2+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fno-builtin %s -o %t.out
23
// RUN: %HOST_RUN_PLACEHOLDER %t.out
34
// RUN: %CPU_RUN_PLACEHOLDER %t.out
5+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
46
// RUN: %ACC_RUN_PLACEHOLDER %t.out
57

68
// RUN: %clangxx -fsycl -fno-builtin -fsycl-device-lib-jit-link %s -o %t.out
@@ -395,10 +397,7 @@ bool kernel_test_memcpy_addr_space(sycl::queue &deviceQueue) {
395397
sycl::access::placeholder::false_t>
396398
src_acc(buffer1, cgh);
397399

398-
sycl::accessor<char, 1, sycl::access::mode::read_write,
399-
sycl::access::target::local,
400-
sycl::access::placeholder::false_t>
401-
local_acc(sycl::range<1>(16), cgh);
400+
sycl::local_accessor<char, 1> local_acc(sycl::range<1>(16), cgh);
402401

403402
sycl::accessor<char, 1, sycl::access::mode::write,
404403
sycl::access::target::device,

‎SYCL/DiscardEvents/discard_events_accessors.cpp

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -60,10 +60,7 @@ int main(int Argc, const char *Argv[]) {
6060
RunKernelHelper(Q, [&](int *Harray) {
6161
Q.submit([&](sycl::handler &CGH) {
6262
const size_t LocalMemSize = BUFFER_SIZE;
63-
using LocalAccessor =
64-
sycl::accessor<int, 1, sycl::access::mode::read_write,
65-
sycl::access::target::local>;
66-
LocalAccessor LocalAcc(LocalMemSize, CGH);
63+
sycl::local_accessor<int, 1> LocalAcc(LocalMemSize, CGH);
6764

6865
CGH.parallel_for<class kernel_using_local_memory>(
6966
Range, [=](sycl::item<1> itemID) {

0 commit comments

Comments
 (0)