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

Commit 9c31c38

Browse files
cperkinsintelcperkinsintel
committed
Merge branch 'intel' into cperkins-expanded-group_broadcast-testing
2 parents 7340c89 + 7e1bb59 commit 9c31c38

File tree

70 files changed

+549
-393
lines changed

Some content is hidden

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

70 files changed

+549
-393
lines changed

‎SYCL/AtomicRef/accessor.cpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -38,10 +38,10 @@ template <typename T> void accessor_test(queue q, size_t N) {
3838
cgh.parallel_for(range<1>(N), [=](item<1> it) {
3939
int gid = it.get_id(0);
4040
static_assert(
41-
std::is_same<
42-
decltype(sum[0]),
43-
atomic_ref<T, memory_order::relaxed, memory_scope::device,
44-
access::address_space::global_space>>::value,
41+
std::is_same<decltype(sum[0]),
42+
::sycl::ext::oneapi::atomic_ref<
43+
T, memory_order::relaxed, memory_scope::device,
44+
access::address_space::global_space>>::value,
4545
"atomic_accessor returns incorrect atomic_ref");
4646
out[gid] = sum[0].fetch_add(T(1));
4747
});
@@ -78,10 +78,10 @@ void local_accessor_test(queue q, size_t N, size_t L = 8) {
7878
sum[0].store(0);
7979
it.barrier();
8080
static_assert(
81-
std::is_same<
82-
decltype(sum[0]),
83-
atomic_ref<T, memory_order::relaxed, memory_scope::device,
84-
access::address_space::local_space>>::value,
81+
std::is_same<decltype(sum[0]),
82+
::sycl::ext::oneapi::atomic_ref<
83+
T, memory_order::relaxed, memory_scope::device,
84+
access::address_space::local_space>>::value,
8585
"local atomic_accessor returns incorrect atomic_ref");
8686
T result = sum[0].fetch_add(T(1));
8787
if (result == it.get_local_range(0) - 1) {

‎SYCL/AtomicRef/add.h

Lines changed: 12 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -24,8 +24,9 @@ void add_fetch_test(queue q, size_t N) {
2424
output_buf.template get_access<access::mode::discard_write>(cgh);
2525
cgh.parallel_for(range<1>(N), [=](item<1> it) {
2626
int gid = it.get_id(0);
27-
auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
28-
access::address_space::global_space>(sum[0]);
27+
auto atm = ::sycl::ext::oneapi::atomic_ref<
28+
T, memory_order::relaxed, memory_scope::device,
29+
access::address_space::global_space>(sum[0]);
2930
out[gid] = atm.fetch_add(Difference(1));
3031
});
3132
});
@@ -59,8 +60,9 @@ void add_plus_equal_test(queue q, size_t N) {
5960
output_buf.template get_access<access::mode::discard_write>(cgh);
6061
cgh.parallel_for(range<1>(N), [=](item<1> it) {
6162
int gid = it.get_id(0);
62-
auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
63-
access::address_space::global_space>(sum[0]);
63+
auto atm = ::sycl::ext::oneapi::atomic_ref<
64+
T, memory_order::relaxed, memory_scope::device,
65+
access::address_space::global_space>(sum[0]);
6466
out[gid] = atm += Difference(1);
6567
});
6668
});
@@ -94,8 +96,9 @@ void add_pre_inc_test(queue q, size_t N) {
9496
output_buf.template get_access<access::mode::discard_write>(cgh);
9597
cgh.parallel_for(range<1>(N), [=](item<1> it) {
9698
int gid = it.get_id(0);
97-
auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
98-
access::address_space::global_space>(sum[0]);
99+
auto atm = ::sycl::ext::oneapi::atomic_ref<
100+
T, memory_order::relaxed, memory_scope::device,
101+
access::address_space::global_space>(sum[0]);
99102
out[gid] = ++atm;
100103
});
101104
});
@@ -129,8 +132,9 @@ void add_post_inc_test(queue q, size_t N) {
129132
output_buf.template get_access<access::mode::discard_write>(cgh);
130133
cgh.parallel_for(range<1>(N), [=](item<1> it) {
131134
int gid = it.get_id(0);
132-
auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
133-
access::address_space::global_space>(sum[0]);
135+
auto atm = ::sycl::ext::oneapi::atomic_ref<
136+
T, memory_order::relaxed, memory_scope::device,
137+
access::address_space::global_space>(sum[0]);
134138
out[gid] = atm++;
135139
});
136140
});

‎SYCL/AtomicRef/assignment.h

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -21,9 +21,10 @@ template <typename T> void assignment_test(queue q, size_t N) {
2121
assignment_buf.template get_access<access::mode::read_write>(cgh);
2222
cgh.parallel_for<assignment_kernel<T>>(range<1>(N), [=](item<1> it) {
2323
size_t gid = it.get_id(0);
24-
auto atm = atomic_ref<T, ext::oneapi::memory_order::relaxed,
25-
ext::oneapi::memory_scope::device,
26-
access::address_space::global_space>(st[0]);
24+
auto atm = ::sycl::ext::oneapi::atomic_ref<
25+
T, ext::oneapi::memory_order::relaxed,
26+
ext::oneapi::memory_scope::device,
27+
access::address_space::global_space>(st[0]);
2728
atm = T(gid);
2829
});
2930
});

‎SYCL/AtomicRef/atomic_memory_order_acq_rel.h

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -19,9 +19,9 @@ template <typename T> void acq_rel_test(queue q, size_t N) {
1919
auto a_acc = a_buf.template get_access<access::mode::read_write>(cgh);
2020
cgh.parallel_for<atomic_memory_order_acq_rel_kernel<T>>(
2121
range<1>(N), [=](item<1> it) {
22-
auto aar =
23-
atomic_ref<T, memory_order::acq_rel, memory_scope::device,
24-
access::address_space::global_space>(a_acc[0]);
22+
auto aar = ::sycl::ext::oneapi::atomic_ref<
23+
T, memory_order::acq_rel, memory_scope::device,
24+
access::address_space::global_space>(a_acc[0]);
2525
auto ld = aar.load();
2626
ld += 1;
2727
aar.store(ld);

‎SYCL/AtomicRef/atomic_memory_order_seq_cst.h

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -22,12 +22,12 @@ template <typename T> void seq_cst_test(queue q, size_t N) {
2222
auto b_acc = b_buf.template get_access<access::mode::read_write>(cgh);
2323
cgh.parallel_for<atomic_memory_order_seq_cst_kernel<T>>(
2424
range<1>(N), [=](item<1> it) {
25-
auto aar =
26-
atomic_ref<T, memory_order::seq_cst, memory_scope::device,
27-
access::address_space::global_space>(a_acc[0]);
28-
auto bar =
29-
atomic_ref<T, memory_order::seq_cst, memory_scope::device,
30-
access::address_space::global_space>(b_acc[0]);
25+
auto aar = ::sycl::ext::oneapi::atomic_ref<
26+
T, memory_order::seq_cst, memory_scope::device,
27+
access::address_space::global_space>(a_acc[0]);
28+
auto bar = ::sycl::ext::oneapi::atomic_ref<
29+
T, memory_order::seq_cst, memory_scope::device,
30+
access::address_space::global_space>(b_acc[0]);
3131
auto ald = aar.load();
3232
auto bld = bar.load();
3333
ald += 1;

‎SYCL/AtomicRef/compare_exchange.h

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -29,9 +29,9 @@ template <typename T> void compare_exchange_test(queue q, size_t N) {
2929
cgh.parallel_for<compare_exchange_kernel<T>>(
3030
range<1>(N), [=](item<1> it) {
3131
size_t gid = it.get_id(0);
32-
auto atm =
33-
atomic_ref<T, memory_order::relaxed, memory_scope::device,
34-
access::address_space::global_space>(exc[0]);
32+
auto atm = ::sycl::ext::oneapi::atomic_ref<
33+
T, memory_order::relaxed, memory_scope::device,
34+
access::address_space::global_space>(exc[0]);
3535
T result = T(N); // Avoid copying pointer
3636
bool success = atm.compare_exchange_strong(result, (T)gid);
3737
if (success) {

‎SYCL/AtomicRef/exchange.h

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -27,8 +27,9 @@ template <typename T> void exchange_test(queue q, size_t N) {
2727
output_buf.template get_access<access::mode::discard_write>(cgh);
2828
cgh.parallel_for<exchange_kernel<T>>(range<1>(N), [=](item<1> it) {
2929
size_t gid = it.get_id(0);
30-
auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
31-
access::address_space::global_space>(exc[0]);
30+
auto atm = ::sycl::ext::oneapi::atomic_ref<
31+
T, memory_order::relaxed, memory_scope::device,
32+
access::address_space::global_space>(exc[0]);
3233
out[gid] = atm.exchange(T(gid));
3334
});
3435
});

‎SYCL/AtomicRef/load.h

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -26,8 +26,9 @@ template <typename T> void load_test(queue q, size_t N) {
2626
output_buf.template get_access<access::mode::discard_write>(cgh);
2727
cgh.parallel_for<load_kernel<T>>(range<1>(N), [=](item<1> it) {
2828
size_t gid = it.get_id(0);
29-
auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
30-
access::address_space::global_space>(ld[0]);
29+
auto atm = ::sycl::ext::oneapi::atomic_ref<
30+
T, memory_order::relaxed, memory_scope::device,
31+
access::address_space::global_space>(ld[0]);
3132
out[gid] = atm.load();
3233
});
3334
});

‎SYCL/AtomicRef/max.h

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -24,9 +24,9 @@ template <typename T> void max_test(queue q, size_t N) {
2424
output_buf.template get_access<access::mode::discard_write>(cgh);
2525
cgh.parallel_for(range<1>(N), [=](item<1> it) {
2626
int gid = it.get_id(0);
27-
auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
28-
access::address_space::global_space>(val[0]);
29-
27+
auto atm = ::sycl::ext::oneapi::atomic_ref<
28+
T, memory_order::relaxed, memory_scope::device,
29+
access::address_space::global_space>(val[0]);
3030
// +1 accounts for lowest() returning 0 for unsigned types
3131
out[gid] = atm.fetch_max(T(gid) + 1);
3232
});

‎SYCL/AtomicRef/min.h

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -24,8 +24,9 @@ template <typename T> void min_test(queue q, size_t N) {
2424
output_buf.template get_access<access::mode::discard_write>(cgh);
2525
cgh.parallel_for(range<1>(N), [=](item<1> it) {
2626
int gid = it.get_id(0);
27-
auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
28-
access::address_space::global_space>(val[0]);
27+
auto atm = ::sycl::ext::oneapi::atomic_ref<
28+
T, memory_order::relaxed, memory_scope::device,
29+
access::address_space::global_space>(val[0]);
2930
out[gid] = atm.fetch_min(T(gid));
3031
});
3132
});

‎SYCL/AtomicRef/store.h

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -20,8 +20,9 @@ template <typename T> void store_test(queue q, size_t N) {
2020
auto st = store_buf.template get_access<access::mode::read_write>(cgh);
2121
cgh.parallel_for<store_kernel<T>>(range<1>(N), [=](item<1> it) {
2222
size_t gid = it.get_id(0);
23-
auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
24-
access::address_space::global_space>(st[0]);
23+
auto atm = ::sycl::ext::oneapi::atomic_ref<
24+
T, memory_order::relaxed, memory_scope::device,
25+
access::address_space::global_space>(st[0]);
2526
atm.store(T(gid));
2627
});
2728
});

‎SYCL/AtomicRef/sub.h

Lines changed: 12 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -24,8 +24,9 @@ void sub_fetch_test(queue q, size_t N) {
2424
output_buf.template get_access<access::mode::discard_write>(cgh);
2525
cgh.parallel_for(range<1>(N), [=](item<1> it) {
2626
int gid = it.get_id(0);
27-
auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
28-
access::address_space::global_space>(val[0]);
27+
auto atm = ::sycl::ext::oneapi::atomic_ref<
28+
T, memory_order::relaxed, memory_scope::device,
29+
access::address_space::global_space>(val[0]);
2930
out[gid] = atm.fetch_sub(Difference(1));
3031
});
3132
});
@@ -59,8 +60,9 @@ void sub_plus_equal_test(queue q, size_t N) {
5960
output_buf.template get_access<access::mode::discard_write>(cgh);
6061
cgh.parallel_for(range<1>(N), [=](item<1> it) {
6162
int gid = it.get_id(0);
62-
auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
63-
access::address_space::global_space>(val[0]);
63+
auto atm = ::sycl::ext::oneapi::atomic_ref<
64+
T, memory_order::relaxed, memory_scope::device,
65+
access::address_space::global_space>(val[0]);
6466
out[gid] = atm -= Difference(1);
6567
});
6668
});
@@ -94,8 +96,9 @@ void sub_pre_dec_test(queue q, size_t N) {
9496
output_buf.template get_access<access::mode::discard_write>(cgh);
9597
cgh.parallel_for(range<1>(N), [=](item<1> it) {
9698
int gid = it.get_id(0);
97-
auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
98-
access::address_space::global_space>(val[0]);
99+
auto atm = ::sycl::ext::oneapi::atomic_ref<
100+
T, memory_order::relaxed, memory_scope::device,
101+
access::address_space::global_space>(val[0]);
99102
out[gid] = --atm;
100103
});
101104
});
@@ -129,8 +132,9 @@ void sub_post_dec_test(queue q, size_t N) {
129132
output_buf.template get_access<access::mode::discard_write>(cgh);
130133
cgh.parallel_for(range<1>(N), [=](item<1> it) {
131134
int gid = it.get_id(0);
132-
auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
133-
access::address_space::global_space>(val[0]);
135+
auto atm = ::sycl::ext::oneapi::atomic_ref<
136+
T, memory_order::relaxed, memory_scope::device,
137+
access::address_space::global_space>(val[0]);
134138
out[gid] = atm--;
135139
});
136140
});

‎SYCL/Basic/kernel_bundle/kernel_bundle_api.cpp

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -369,5 +369,28 @@ int main() {
369369
"is empty");
370370
}
371371

372+
{
373+
// no duplicate devices
374+
sycl::kernel_bundle KernelBundleDupTest =
375+
sycl::get_kernel_bundle<sycl::bundle_state::input>(Ctx, {Dev, Dev},
376+
{Kernel1ID});
377+
assert(KernelBundleDupTest.get_devices().size() == 1);
378+
379+
sycl::kernel_bundle<sycl::bundle_state::object>
380+
KernelBundleDupeTestCompiled =
381+
sycl::compile(KernelBundleDupTest, {Dev, Dev});
382+
assert(KernelBundleDupeTestCompiled.get_devices().size() == 1);
383+
384+
sycl::kernel_bundle<sycl::bundle_state::executable>
385+
KernelBundleDupeTestLinked =
386+
sycl::link({KernelBundleDupeTestCompiled}, {Dev, Dev});
387+
assert(KernelBundleDupeTestLinked.get_devices().size() == 1);
388+
389+
sycl::kernel_bundle<sycl::bundle_state::executable>
390+
KernelBundleDupeTestBuilt =
391+
sycl::build(KernelBundleDupTest, {Dev, Dev});
392+
assert(KernelBundleDupeTestBuilt.get_devices().size() == 1);
393+
}
394+
372395
return 0;
373396
}

‎SYCL/Basic/kernel_info.cpp

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -5,9 +5,6 @@
55
//
66
// Fail is flaky for level_zero, enable when fixed.
77
// UNSUPPORTED: level_zero
8-
//
9-
// Failing on HIP AMD and HIP Nvidia
10-
// XFAIL: hip_amd || hip_nvidia
118

129
//==--- kernel_info.cpp - SYCL kernel info test ----------------------------==//
1310
//

‎SYCL/Basic/linear-sub_group.cpp

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4,8 +4,6 @@
44
// RUN: %GPU_RUN_PLACEHOLDER %t.out
55
// RUN: %ACC_RUN_PLACEHOLDER %t.out
66
//
7-
// Missing built-ins on AMD
8-
// XFAIL: hip_amd
97
//==--------------- linear-sub_group.cpp - SYCL linear id test -------------==//
108
//
119
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.

‎SYCL/Basic/multi_ptr.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,7 @@ template <typename T> class testMultPtrKernel;
3030
template <typename T> class testMultPtrArrowOperatorKernel;
3131

3232
template <typename T> struct point {
33-
point(const point &rhs) : x(rhs.x), y(rhs.y) {}
33+
point(const point &rhs) = default;
3434
point(T x, T y) : x(x), y(y) {}
3535
point(T v) : x(v), y(v) {}
3636
point() : x(0), y(0) {}

‎SYCL/DeprecatedFeatures/get_backend.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19,6 +19,7 @@ bool check(backend be) {
1919
case backend::level_zero:
2020
case backend::cuda:
2121
case backend::host:
22+
case backend::hip:
2223
return true;
2324
default:
2425
return false;

‎SYCL/DeprecatedFeatures/kernel_interop.cpp

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -52,10 +52,9 @@ int main() {
5252

5353
// Try to create kernel with another context
5454
bool Pass = false;
55-
queue Queue1;
56-
context Context1 = Queue1.get_context();
55+
context OtherContext{Context.get_devices()[0]};
5756
try {
58-
kernel Kernel(ClKernel, Context1);
57+
kernel Kernel(ClKernel, OtherContext);
5958
} catch (cl::sycl::invalid_parameter_error e) {
6059
Pass = true;
6160
}

‎SYCL/ESIMD/accessor_gather_scatter.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -70,10 +70,10 @@ template <typename T, unsigned VL, unsigned STRIDE> bool test(queue q) {
7070
cgh.parallel_for(glob_range, kernel);
7171
});
7272
e.wait();
73-
} catch (cl::sycl::exception const &e) {
73+
} catch (sycl::exception const &e) {
7474
std::cout << "SYCL exception caught: " << e.what() << '\n';
7575
delete[] A;
76-
return e.get_cl_code();
76+
return false; // not success
7777
}
7878

7979
int err_cnt = 0;

‎SYCL/ESIMD/accessor_load_store.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -62,10 +62,10 @@ template <typename T> bool test(queue q, size_t size) {
6262
Kernel<T> kernel(acc);
6363
cgh.parallel_for(glob_range, kernel);
6464
});
65-
} catch (cl::sycl::exception const &e) {
65+
} catch (sycl::exception const &e) {
6666
std::cout << "SYCL exception caught: " << e.what() << '\n';
6767
delete[] A;
68-
return e.get_cl_code();
68+
return false; // not success
6969
}
7070

7171
int err_cnt = 0;

‎SYCL/ESIMD/api/simd_binop_integer_promotion.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -74,9 +74,9 @@ template <typename T> bool test(queue q) {
7474
});
7575
});
7676
q.wait_and_throw();
77-
} catch (cl::sycl::exception const &e) {
77+
} catch (sycl::exception const &e) {
7878
std::cout << "SYCL exception caught: " << e.what() << '\n';
79-
return e.get_cl_code();
79+
return false; // not success
8080
}
8181

8282
int err_cnt = 0;

‎SYCL/ESIMD/api/simd_memory_access.cpp

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -79,11 +79,10 @@ template <typename T, int N, bool IsAcc> bool test(queue q, size_t size) {
7979
std::cout << "Testing T=" << typeid(T).name() << ", N=" << N
8080
<< " using accessor=" << IsAcc << "...\n";
8181
T *A;
82-
if constexpr (IsAcc) {
82+
if constexpr (IsAcc)
8383
A = new T[size];
84-
} else {
85-
A = reinterpret_cast<T *>(sycl::malloc_shared(size, q));
86-
}
84+
else
85+
A = sycl::malloc_shared<T>(size, q);
8786

8887
for (unsigned i = 0; i < size; ++i) {
8988
A[i] = i; // should not be zero to test `copy_from` really works
@@ -108,10 +107,10 @@ template <typename T, int N, bool IsAcc> bool test(queue q, size_t size) {
108107
});
109108
}
110109
q.wait_and_throw();
111-
} catch (cl::sycl::exception const &e) {
110+
} catch (sycl::exception const &e) {
112111
std::cout << "SYCL exception caught: " << e.what() << '\n';
113112
free_mem<IsAcc>(A, q);
114-
return e.get_cl_code();
113+
return false; // not success
115114
}
116115

117116
int err_cnt = 0;

0 commit comments

Comments
 (0)