Skip to content

Commit b7a52b6

Browse files
HPS-1HPS-1
committed
aaa
Signed-off-by: Hu, Peisen <[email protected]>
1 parent 67d669c commit b7a52b6

File tree

1 file changed

+56
-35
lines changed

1 file changed

+56
-35
lines changed

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

Lines changed: 56 additions & 35 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} -I . -o %t.out %if any-device-is-cuda %{ -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70 %}
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 %}
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-
// TODO: AAAAA
20+
2121
static constexpr int WorkGroupSize = 32;
2222

2323
void testFeatureMacro() {
@@ -42,9 +42,14 @@ 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-
const auto props = sycl::ext::oneapi::experimental::properties{
46-
sycl::ext::oneapi::experimental::use_root_sync};
47-
q.single_task<class QueryKernel>(props, []() {});
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{});
4853

4954
static auto check_max_num_work_group_sync = [](auto Result) {
5055
static_assert(std::is_same_v<std::remove_cv_t<decltype(Result)>, size_t>,
@@ -65,14 +70,15 @@ void testRootGroup() {
6570
.ext_oneapi_get_info<sycl::ext::oneapi::experimental::info::
6671
kernel_queue_specific::max_num_work_groups>(
6772
q, WorkGroupSize, 0);
68-
const auto props = sycl::ext::oneapi::experimental::properties{
69-
sycl::ext::oneapi::experimental::use_root_sync};
7073
sycl::buffer<int> dataBuf{sycl::range{maxWGs * WorkGroupSize}};
7174
const auto range = sycl::nd_range<1>{maxWGs * WorkGroupSize, WorkGroupSize};
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) {
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};
7682
volatile float X = 1.0f;
7783
volatile float Y = 1.0f;
7884
auto root = it.ext_oneapi_get_root_group();
@@ -90,7 +96,15 @@ void testRootGroup() {
9096
data[root.get_local_range() - root.get_local_id() - 1];
9197
sycl::group_barrier(root);
9298
data[root.get_local_id()] = sum;
93-
});
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));
94108
});
95109
sycl::host_accessor data{dataBuf};
96110
const int workItemCount = static_cast<int>(range.get_global_range().size());
@@ -109,34 +123,41 @@ void testRootGroupFunctions() {
109123
.ext_oneapi_get_info<sycl::ext::oneapi::experimental::info::
110124
kernel_queue_specific::max_num_work_groups>(
111125
q, WorkGroupSize, 0);
112-
const auto props = sycl::ext::oneapi::experimental::properties{
113-
sycl::ext::oneapi::experimental::use_root_sync};
114-
115126
constexpr int testCount = 9;
116127
sycl::buffer<bool> testResultsBuf{sycl::range{testCount}};
117128
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+
118158
q.submit([&](sycl::handler &h) {
119-
sycl::accessor testResults{testResultsBuf, h};
120159
h.parallel_for<class RootGroupFunctionsKernel>(
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-
});
160+
range, TestKernel2(&testResultsBuf, &h));
140161
});
141162
sycl::host_accessor testResults{testResultsBuf};
142163
for (int i = 0; i < testCount; i++) {

0 commit comments

Comments
 (0)