Skip to content

Commit b158e52

Browse files
HPS-1HPS-1
committed
[SYCL] Rewrite e2e tests affected by deprecation
Signed-off-by: Hu, Peisen <[email protected]>
1 parent 046c807 commit b158e52

File tree

4 files changed

+70
-83
lines changed

4 files changed

+70
-83
lines changed

‎sycl/test-e2e/Basic/max_linear_work_group_size_props.cpp‎

Lines changed: 9 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -57,17 +57,15 @@ template <size_t I> struct KernelFunctorWithMaxWGSizeProp {
5757
}
5858
};
5959

60-
template <Variant KernelVariant, size_t I, typename PropertiesT,
61-
typename KernelType>
62-
int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
60+
template <Variant KernelVariant, size_t I, typename KernelType>
61+
int test(queue &Q, KernelType KernelFunc) {
6362
constexpr size_t Dims = 1;
6463

6564
// Positive test case: Specify local size that matches required size.
6665
try {
6766
Q.submit([&](handler &CGH) {
6867
CGH.parallel_for<MaxLinearWGSizePositive<KernelVariant, false, I>>(
69-
nd_range<Dims>(repeatRange<Dims>(8), range<Dims>(I)), Props,
70-
KernelFunc);
68+
nd_range<Dims>(repeatRange<Dims>(8), range<Dims>(I)), KernelFunc);
7169
});
7270
Q.wait_and_throw();
7371
} catch (exception &E) {
@@ -80,8 +78,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
8078
// Same as above but using the queue shortcuts.
8179
try {
8280
Q.parallel_for<MaxLinearWGSizePositive<KernelVariant, true, I>>(
83-
nd_range<Dims>(repeatRange<Dims>(8), range<Dims>(I)), Props,
84-
KernelFunc);
81+
nd_range<Dims>(repeatRange<Dims>(8), range<Dims>(I)), KernelFunc);
8582
Q.wait_and_throw();
8683
} catch (exception &E) {
8784
std::cerr
@@ -96,7 +93,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
9693
try {
9794
Q.submit([&](handler &CGH) {
9895
CGH.parallel_for<MaxLinearWGSizeNoLocalPositive<KernelVariant, false, I>>(
99-
repeatRange<Dims>(16), Props, KernelFunc);
96+
repeatRange<Dims>(16), KernelFunc);
10097
});
10198
Q.wait_and_throw();
10299
} catch (exception &E) {
@@ -108,7 +105,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
108105

109106
try {
110107
Q.parallel_for<MaxLinearWGSizeNoLocalPositive<KernelVariant, true, I>>(
111-
repeatRange<Dims>(16), Props, KernelFunc);
108+
repeatRange<Dims>(16), KernelFunc);
112109
Q.wait_and_throw();
113110
} catch (exception &E) {
114111
std::cerr << "Test case MaxLinearWGSizeNoLocalPositive shortcut failed: "
@@ -121,7 +118,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
121118
try {
122119
Q.submit([&](handler &CGH) {
123120
CGH.parallel_for<MaxLinearWGSizeNegative<KernelVariant, false, I>>(
124-
nd_range<Dims>(repeatRange<Dims>(16), repeatRange<Dims>(8)), Props,
121+
nd_range<Dims>(repeatRange<Dims>(16), repeatRange<Dims>(8)),
125122
KernelFunc);
126123
});
127124
Q.wait_and_throw();
@@ -146,7 +143,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
146143
// Same as above but using the queue shortcuts.
147144
try {
148145
Q.parallel_for<MaxLinearWGSizeNegative<KernelVariant, true, I>>(
149-
nd_range<Dims>(repeatRange<Dims>(16), repeatRange<Dims>(8)), Props,
146+
nd_range<Dims>(repeatRange<Dims>(16), repeatRange<Dims>(8)),
150147
KernelFunc);
151148
Q.wait_and_throw();
152149
std::cerr
@@ -173,17 +170,10 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
173170
}
174171

175172
template <size_t I> int test_max(queue &Q) {
176-
auto Props = ext::oneapi::experimental::properties{
177-
ext::oneapi::experimental::max_linear_work_group_size<I>};
178-
auto KernelFunction = [](auto) {};
179-
180-
auto EmptyProps = ext::oneapi::experimental::properties{};
181173
KernelFunctorWithMaxWGSizeProp<I> KernelFunctor;
182174

183175
int Res = 0;
184-
Res += test<Variant::Function, I>(Q, Props, KernelFunction);
185-
Res += test<Variant::Functor, I>(Q, EmptyProps, KernelFunctor);
186-
Res += test<Variant::FunctorAndProperty, I>(Q, Props, KernelFunctor);
176+
Res += test<Variant::Functor, I>(Q, KernelFunctor);
187177
return Res;
188178
}
189179

‎sycl/test-e2e/Basic/max_work_group_size_props.cpp‎

Lines changed: 9 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -48,17 +48,15 @@ template <size_t... Is> struct KernelFunctorWithMaxWGSizeProp {
4848
}
4949
};
5050

51-
template <Variant KernelVariant, size_t... Is, typename PropertiesT,
52-
typename KernelType>
53-
int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
51+
template <Variant KernelVariant, size_t... Is, typename KernelType>
52+
int test(queue &Q, KernelType KernelFunc) {
5453
constexpr size_t Dims = sizeof...(Is);
5554

5655
// Positive test case: Specify local size that matches required size.
5756
try {
5857
Q.submit([&](handler &CGH) {
5958
CGH.parallel_for<MaxWGSizePositive<KernelVariant, false, Is...>>(
60-
nd_range<Dims>(repeatRange<Dims>(8), range<Dims>(Is...)), Props,
61-
KernelFunc);
59+
nd_range<Dims>(repeatRange<Dims>(8), range<Dims>(Is...)), KernelFunc);
6260
});
6361
Q.wait_and_throw();
6462
} catch (exception &E) {
@@ -70,8 +68,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
7068
// Same as above but using the queue shortcuts.
7169
try {
7270
Q.parallel_for<MaxWGSizePositive<KernelVariant, true, Is...>>(
73-
nd_range<Dims>(repeatRange<Dims>(8), range<Dims>(Is...)), Props,
74-
KernelFunc);
71+
nd_range<Dims>(repeatRange<Dims>(8), range<Dims>(Is...)), KernelFunc);
7572
Q.wait_and_throw();
7673
} catch (exception &E) {
7774
std::cerr << "Test case MaxWGSizePositive shortcut failed: unexpected "
@@ -85,7 +82,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
8582
try {
8683
Q.submit([&](handler &CGH) {
8784
CGH.parallel_for<MaxWGSizeNoLocalPositive<KernelVariant, false, Is...>>(
88-
repeatRange<Dims>(16), Props, KernelFunc);
85+
repeatRange<Dims>(16), KernelFunc);
8986
});
9087
Q.wait_and_throw();
9188
} catch (exception &E) {
@@ -97,7 +94,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
9794

9895
try {
9996
Q.parallel_for<MaxWGSizeNoLocalPositive<KernelVariant, true, Is...>>(
100-
repeatRange<Dims>(16), Props, KernelFunc);
97+
repeatRange<Dims>(16), KernelFunc);
10198
Q.wait_and_throw();
10299
} catch (exception &E) {
103100
std::cerr << "Test case MaxWGSizeNoLocalPositive shortcut failed: "
@@ -110,7 +107,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
110107
try {
111108
Q.submit([&](handler &CGH) {
112109
CGH.parallel_for<MaxWGSizeNegative<KernelVariant, false, Is...>>(
113-
nd_range<Dims>(repeatRange<Dims>(16), repeatRange<Dims>(8)), Props,
110+
nd_range<Dims>(repeatRange<Dims>(16), repeatRange<Dims>(8)),
114111
KernelFunc);
115112
});
116113
Q.wait_and_throw();
@@ -133,7 +130,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
133130
// Same as above but using the queue shortcuts.
134131
try {
135132
Q.parallel_for<MaxWGSizeNegative<KernelVariant, true, Is...>>(
136-
nd_range<Dims>(repeatRange<Dims>(16), repeatRange<Dims>(8)), Props,
133+
nd_range<Dims>(repeatRange<Dims>(16), repeatRange<Dims>(8)),
137134
KernelFunc);
138135
Q.wait_and_throw();
139136
std::cerr << "Test case MaxWGSizeNegative shortcut failed: no exception "
@@ -158,17 +155,10 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
158155
}
159156

160157
template <size_t... Is> int test_max(queue &Q) {
161-
auto Props = ext::oneapi::experimental::properties{
162-
ext::oneapi::experimental::max_work_group_size<Is...>};
163-
auto KernelFunction = [](auto) {};
164-
165-
auto EmptyProps = ext::oneapi::experimental::properties{};
166158
KernelFunctorWithMaxWGSizeProp<Is...> KernelFunctor;
167159

168160
int Res = 0;
169-
Res += test<Variant::Function, Is...>(Q, Props, KernelFunction);
170-
Res += test<Variant::Functor, Is...>(Q, EmptyProps, KernelFunctor);
171-
Res += test<Variant::FunctorAndProperty, Is...>(Q, Props, KernelFunctor);
161+
Res += test<Variant::Functor, Is...>(Q, KernelFunctor);
172162
return Res;
173163
}
174164

‎sycl/test-e2e/Basic/work_group_size_prop.cpp‎

Lines changed: 9 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -45,9 +45,8 @@ template <size_t... Is> struct KernelFunctorWithWGSizeProp {
4545
}
4646
};
4747

48-
template <Variant KernelVariant, size_t... Is, typename PropertiesT,
49-
typename KernelType>
50-
int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
48+
template <Variant KernelVariant, size_t... Is, typename KernelType>
49+
int test(queue &Q, KernelType KernelFunc) {
5150
constexpr size_t Dims = sizeof...(Is);
5251

5352
bool IsOpenCL = (Q.get_backend() == backend::opencl);
@@ -56,8 +55,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
5655
try {
5756
Q.submit([&](handler &CGH) {
5857
CGH.parallel_for<ReqdWGSizePositiveA<KernelVariant, false, Is...>>(
59-
nd_range<Dims>(repeatRange<Dims>(8), range<Dims>(Is...)), Props,
60-
KernelFunc);
58+
nd_range<Dims>(repeatRange<Dims>(8), range<Dims>(Is...)), KernelFunc);
6159
});
6260
Q.wait_and_throw();
6361
} catch (exception &E) {
@@ -69,8 +67,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
6967
// Same as above but using the queue shortcuts.
7068
try {
7169
Q.parallel_for<ReqdWGSizePositiveA<KernelVariant, true, Is...>>(
72-
nd_range<Dims>(repeatRange<Dims>(8), range<Dims>(Is...)), Props,
73-
KernelFunc);
70+
nd_range<Dims>(repeatRange<Dims>(8), range<Dims>(Is...)), KernelFunc);
7471
Q.wait_and_throw();
7572
} catch (exception &E) {
7673
std::cerr << "Test case ReqdWGSizePositiveA shortcut failed: unexpected "
@@ -87,7 +84,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
8784
Q.submit([&](handler &CGH) {
8885
CGH.parallel_for<
8986
ReqdWGSizeNoLocalPositive<KernelVariant, false, Is...>>(
90-
repeatRange<Dims>(16), Props, KernelFunc);
87+
repeatRange<Dims>(16), KernelFunc);
9188
});
9289
Q.wait_and_throw();
9390
} catch (exception &E) {
@@ -99,7 +96,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
9996

10097
try {
10198
Q.parallel_for<ReqdWGSizeNoLocalPositive<KernelVariant, true, Is...>>(
102-
repeatRange<Dims>(16), Props, KernelFunc);
99+
repeatRange<Dims>(16), KernelFunc);
103100
Q.wait_and_throw();
104101
} catch (exception &E) {
105102
std::cerr << "Test case ReqdWGSizeNoLocalPositive shortcut failed: "
@@ -113,7 +110,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
113110
try {
114111
Q.submit([&](handler &CGH) {
115112
CGH.parallel_for<ReqdWGSizeNegativeA<KernelVariant, false, Is...>>(
116-
nd_range<Dims>(repeatRange<Dims>(16), repeatRange<Dims>(8)), Props,
113+
nd_range<Dims>(repeatRange<Dims>(16), repeatRange<Dims>(8)),
117114
KernelFunc);
118115
});
119116
Q.wait_and_throw();
@@ -137,7 +134,7 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
137134
// Same as above but using the queue shortcuts.
138135
try {
139136
Q.parallel_for<ReqdWGSizeNegativeA<KernelVariant, true, Is...>>(
140-
nd_range<Dims>(repeatRange<Dims>(16), repeatRange<Dims>(8)), Props,
137+
nd_range<Dims>(repeatRange<Dims>(16), repeatRange<Dims>(8)),
141138
KernelFunc);
142139
Q.wait_and_throw();
143140
std::cerr << "Test case ReqdWGSizeNegativeA shortcut failed: no exception "
@@ -162,17 +159,10 @@ int test(queue &Q, PropertiesT Props, KernelType KernelFunc) {
162159
}
163160

164161
template <size_t... Is> int test(queue &Q) {
165-
auto Props = ext::oneapi::experimental::properties{
166-
ext::oneapi::experimental::work_group_size<Is...>};
167-
auto KernelFunction = [](auto) {};
168-
169-
auto EmptyProps = ext::oneapi::experimental::properties{};
170162
KernelFunctorWithWGSizeProp<Is...> KernelFunctor;
171163

172164
int Res = 0;
173-
Res += test<Variant::Function, Is...>(Q, Props, KernelFunction);
174-
Res += test<Variant::Functor, Is...>(Q, EmptyProps, KernelFunctor);
175-
Res += test<Variant::FunctorAndProperty, Is...>(Q, Props, KernelFunctor);
165+
Res += test<Variant::Functor, Is...>(Q, KernelFunctor);
176166
return Res;
177167
}
178168

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

Lines changed: 43 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -44,7 +44,11 @@ void testQueriesAndProperties() {
4444
q, wgRange, wgRange.size() * sizeof(int));
4545
const auto props = sycl::ext::oneapi::experimental::properties{
4646
sycl::ext::oneapi::experimental::use_root_sync};
47-
q.single_task<class QueryKernel>(props, []() {});
47+
struct TestKernel0 {
48+
void operator()() const {}
49+
auto get(sycl::ext::oneapi::experimental::properties_tag) { return props; }
50+
};
51+
q.single_task<class QueryKernel>(TestKernel0{});
4852

4953
static auto check_max_num_work_group_sync = [](auto Result) {
5054
static_assert(std::is_same_v<std::remove_cv_t<decltype(Result)>, size_t>,
@@ -69,10 +73,10 @@ void testRootGroup() {
6973
sycl::ext::oneapi::experimental::use_root_sync};
7074
sycl::buffer<int> dataBuf{sycl::range{maxWGs * WorkGroupSize}};
7175
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) {
76+
struct TestKernel1 {
77+
sycl::accessor data;
78+
TestKernel1(sycl::accessor data_param) { data = data_param; }
79+
void operator()(sycl::nd_item<1> it) const {
7680
volatile float X = 1.0f;
7781
volatile float Y = 1.0f;
7882
auto root = it.ext_oneapi_get_root_group();
@@ -90,7 +94,12 @@ void testRootGroup() {
9094
data[root.get_local_range() - root.get_local_id() - 1];
9195
sycl::group_barrier(root);
9296
data[root.get_local_id()] = sum;
93-
});
97+
}
98+
auto get(sycl::ext::oneapi::experimental::properties_tag) { return props; }
99+
};
100+
q.submit([&](sycl::handler &h) {
101+
sycl::accessor data{dataBuf, h};
102+
h.parallel_for<class RootGroupKernel>(range, TestKernel1(data));
94103
});
95104
sycl::host_accessor data{dataBuf};
96105
const int workItemCount = static_cast<int>(range.get_global_range().size());
@@ -115,28 +124,36 @@ void testRootGroupFunctions() {
115124
constexpr int testCount = 9;
116125
sycl::buffer<bool> testResultsBuf{sycl::range{testCount}};
117126
const auto range = sycl::nd_range<1>{maxWGs * WorkGroupSize, WorkGroupSize};
127+
128+
struct TestKernel2 {
129+
sycl::accessor testResults;
130+
TestKernel2(sycl::accessor testResults_param) {
131+
testResults = testResults_param;
132+
}
133+
void operator()(sycl::nd_item<1> it) const {
134+
const auto root = it.ext_oneapi_get_root_group();
135+
if (root.leader() || root.get_local_id() == 3) {
136+
testResults[0] = root.get_group_id() == sycl::id<1>(0);
137+
testResults[1] = root.leader() ? root.get_local_id() == sycl::id<1>(0)
138+
: root.get_local_id() == sycl::id<1>(3);
139+
testResults[2] = root.get_group_range() == sycl::range<1>(1);
140+
testResults[3] = root.get_local_range() == it.get_global_range();
141+
testResults[4] = root.get_max_local_range() == root.get_local_range();
142+
testResults[5] = root.get_group_linear_id() == 0;
143+
testResults[6] =
144+
root.get_local_linear_id() == root.get_local_id().get(0);
145+
testResults[7] = root.get_group_linear_range() == 1;
146+
testResults[8] =
147+
root.get_local_linear_range() == root.get_local_range().size();
148+
}
149+
}
150+
auto get(sycl::ext::oneapi::experimental::properties_tag) { return props; }
151+
};
152+
118153
q.submit([&](sycl::handler &h) {
119154
sycl::accessor testResults{testResultsBuf, h};
120-
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-
});
155+
h.parallel_for<class RootGroupFunctionsKernel>(range,
156+
TestKernel2(testResults));
140157
});
141158
sycl::host_accessor testResults{testResultsBuf};
142159
for (int i = 0; i < testCount; i++) {

0 commit comments

Comments
 (0)