Skip to content

Commit cfdb8e9

Browse files
HPS-1HPS-1
committed
[SYCL] Restore
Signed-off-by: Hu, Peisen <[email protected]>
1 parent b7a52b6 commit cfdb8e9

File tree

1 file changed

+35
-56
lines changed

1 file changed

+35
-56
lines changed

‎sycl/test-e2e/GroupAlgorithm/root_group.cpp‎

Lines changed: 35 additions & 56 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
// XFAIL: (opencl && !cpu && !accelerator)
33
// XFAIL-TRACKER: https://github.com/intel/llvm/issues/14641
44

5-
// RUN: %{build} -Wno-deprecated-declarations -I . -o %t.out %if any-device-is-cuda %{ -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70 %}
5+
// RUN: %{build} -I . -o %t.out %if any-device-is-cuda %{ -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70 %}
66
// RUN: %{run} %t.out
77

88
// Disabled temporarily while investigation into the failure is ongoing.
@@ -17,7 +17,7 @@
1717
#include <sycl/ext/oneapi/experimental/root_group.hpp>
1818
#include <sycl/group_barrier.hpp>
1919
#include <sycl/kernel_bundle.hpp>
20-
20+
// TODO: AAAAA
2121
static constexpr int WorkGroupSize = 32;
2222

2323
void testFeatureMacro() {
@@ -42,14 +42,9 @@ void testQueriesAndProperties() {
4242
.ext_oneapi_get_info<sycl::ext::oneapi::experimental::info::
4343
kernel_queue_specific::max_num_work_groups>(
4444
q, wgRange, wgRange.size() * sizeof(int));
45-
struct TestKernel0 {
46-
void operator()() const {}
47-
auto get(sycl::ext::oneapi::experimental::properties_tag) {
48-
return sycl::ext::oneapi::experimental::properties{
49-
sycl::ext::oneapi::experimental::use_root_sync};
50-
}
51-
};
52-
q.single_task<class QueryKernel>(TestKernel0{});
45+
const auto props = sycl::ext::oneapi::experimental::properties{
46+
sycl::ext::oneapi::experimental::use_root_sync};
47+
q.single_task<class QueryKernel>(props, []() {});
5348

5449
static auto check_max_num_work_group_sync = [](auto Result) {
5550
static_assert(std::is_same_v<std::remove_cv_t<decltype(Result)>, size_t>,
@@ -70,15 +65,14 @@ void testRootGroup() {
7065
.ext_oneapi_get_info<sycl::ext::oneapi::experimental::info::
7166
kernel_queue_specific::max_num_work_groups>(
7267
q, WorkGroupSize, 0);
68+
const auto props = sycl::ext::oneapi::experimental::properties{
69+
sycl::ext::oneapi::experimental::use_root_sync};
7370
sycl::buffer<int> dataBuf{sycl::range{maxWGs * WorkGroupSize}};
7471
const auto range = sycl::nd_range<1>{maxWGs * WorkGroupSize, WorkGroupSize};
75-
struct TestKernel1 {
76-
sycl::buffer<int> *m_dataBuf;
77-
sycl::handler *m_h;
78-
TestKernel1(sycl::buffer<int> *dataBuf, sycl::handler *h)
79-
: m_dataBuf(dataBuf), m_h(h) {}
80-
void operator()(sycl::nd_item<1> it) const {
81-
sycl::accessor data{*m_dataBuf, *m_h};
72+
q.submit([&](sycl::handler &h) {
73+
sycl::accessor data{dataBuf, h};
74+
h.parallel_for<
75+
class RootGroupKernel>(range, props, [=](sycl::nd_item<1> it) {
8276
volatile float X = 1.0f;
8377
volatile float Y = 1.0f;
8478
auto root = it.ext_oneapi_get_root_group();
@@ -96,15 +90,7 @@ void testRootGroup() {
9690
data[root.get_local_range() - root.get_local_id() - 1];
9791
sycl::group_barrier(root);
9892
data[root.get_local_id()] = sum;
99-
}
100-
auto get(sycl::ext::oneapi::experimental::properties_tag) {
101-
return sycl::ext::oneapi::experimental::properties{
102-
sycl::ext::oneapi::experimental::use_root_sync};
103-
;
104-
}
105-
};
106-
q.submit([&](sycl::handler &h) {
107-
h.parallel_for<class RootGroupKernel>(range, TestKernel1(&dataBuf, &h));
93+
});
10894
});
10995
sycl::host_accessor data{dataBuf};
11096
const int workItemCount = static_cast<int>(range.get_global_range().size());
@@ -123,41 +109,34 @@ void testRootGroupFunctions() {
123109
.ext_oneapi_get_info<sycl::ext::oneapi::experimental::info::
124110
kernel_queue_specific::max_num_work_groups>(
125111
q, WorkGroupSize, 0);
112+
const auto props = sycl::ext::oneapi::experimental::properties{
113+
sycl::ext::oneapi::experimental::use_root_sync};
114+
126115
constexpr int testCount = 9;
127116
sycl::buffer<bool> testResultsBuf{sycl::range{testCount}};
128117
const auto range = sycl::nd_range<1>{maxWGs * WorkGroupSize, WorkGroupSize};
129-
struct TestKernel2 {
130-
sycl::buffer<bool> *m_testResultsBuf;
131-
sycl::handler *m_h;
132-
TestKernel2(sycl::buffer<bool> *testResultsBuf, sycl::handler *h)
133-
: m_testResultsBuf(testResultsBuf), m_h(h) {}
134-
void operator()(sycl::nd_item<1> it) const {
135-
sycl::accessor testResults{*m_testResultsBuf, *m_h};
136-
const auto root = it.ext_oneapi_get_root_group();
137-
if (root.leader() || root.get_local_id() == 3) {
138-
testResults[0] = root.get_group_id() == sycl::id<1>(0);
139-
testResults[1] = root.leader() ? root.get_local_id() == sycl::id<1>(0)
140-
: root.get_local_id() == sycl::id<1>(3);
141-
testResults[2] = root.get_group_range() == sycl::range<1>(1);
142-
testResults[3] = root.get_local_range() == it.get_global_range();
143-
testResults[4] = root.get_max_local_range() == root.get_local_range();
144-
testResults[5] = root.get_group_linear_id() == 0;
145-
testResults[6] =
146-
root.get_local_linear_id() == root.get_local_id().get(0);
147-
testResults[7] = root.get_group_linear_range() == 1;
148-
testResults[8] =
149-
root.get_local_linear_range() == root.get_local_range().size();
150-
}
151-
}
152-
auto get(sycl::ext::oneapi::experimental::properties_tag) {
153-
return sycl::ext::oneapi::experimental::properties{
154-
sycl::ext::oneapi::experimental::use_root_sync};
155-
}
156-
};
157-
158118
q.submit([&](sycl::handler &h) {
119+
sycl::accessor testResults{testResultsBuf, h};
159120
h.parallel_for<class RootGroupFunctionsKernel>(
160-
range, TestKernel2(&testResultsBuf, &h));
121+
range, props, [=](sycl::nd_item<1> it) {
122+
const auto root = it.ext_oneapi_get_root_group();
123+
if (root.leader() || root.get_local_id() == 3) {
124+
testResults[0] = root.get_group_id() == sycl::id<1>(0);
125+
testResults[1] = root.leader()
126+
? root.get_local_id() == sycl::id<1>(0)
127+
: root.get_local_id() == sycl::id<1>(3);
128+
testResults[2] = root.get_group_range() == sycl::range<1>(1);
129+
testResults[3] = root.get_local_range() == it.get_global_range();
130+
testResults[4] =
131+
root.get_max_local_range() == root.get_local_range();
132+
testResults[5] = root.get_group_linear_id() == 0;
133+
testResults[6] =
134+
root.get_local_linear_id() == root.get_local_id().get(0);
135+
testResults[7] = root.get_group_linear_range() == 1;
136+
testResults[8] =
137+
root.get_local_linear_range() == root.get_local_range().size();
138+
}
139+
});
161140
});
162141
sycl::host_accessor testResults{testResultsBuf};
163142
for (int i = 0; i < testCount; i++) {

0 commit comments

Comments
 (0)