Skip to content

Commit f3a34af

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Add CTAD for launch_config to avoid GCC warning (#15004)
This commit adds a CTAD work-around for the launch_config class added in the [sycl_ext_oneapi_enqueue_functions](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc) extension to avoid a warning from GCC about potentially unintended CTAD. This CTAD should never be hit, but will make the compiler accept that even the implicit CTAD is intended. --------- Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
1 parent 0070c93 commit f3a34af

File tree

11 files changed

+137
-107
lines changed

11 files changed

+137
-107
lines changed

‎sycl/include/sycl/detail/common.hpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -374,6 +374,11 @@ static constexpr std::array<T, N> RepeatValue(const T &Arg) {
374374
#define __SYCL_REPORT_EXCEPTION_TO_STREAM(str, e)
375375
#endif
376376

377+
// Tag to help create CTAD definition to avoid ctad-maybe-unsupported warning
378+
// in GCC when relying on default deductions on non-template ctors in template
379+
// classes.
380+
struct AllowCTADTag;
381+
377382
} // namespace detail
378383
} // namespace _V1
379384
} // namespace sycl

‎sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -74,6 +74,13 @@ class launch_config {
7474
friend struct detail::LaunchConfigAccess;
7575
};
7676

77+
#ifdef __cpp_deduction_guides
78+
// CTAD work-around to avoid warning from GCC when using default deduction
79+
// guidance.
80+
launch_config(detail::AllowCTADTag)
81+
-> launch_config<void, empty_properties_t, void>;
82+
#endif // __cpp_deduction_guides
83+
7784
namespace detail {
7885
// Helper for accessing the members of launch_config.
7986
template <typename LCRangeT, typename LCPropertiesT> struct LaunchConfigAccess {

‎sycl/test-e2e/EnqueueFunctions/kernel_shortcut_with_kb.cpp

Lines changed: 18 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -94,58 +94,53 @@ int main() {
9494
Failed += Check(Memory, 48, I, "3D nd_launch shortcut");
9595

9696
// 1D parallel_for shortcut with launch config
97-
oneapiext::parallel_for(
98-
Q, oneapiext::launch_config<sycl::range<1>>{sycl::range<1>{N}}, Kernel1D,
99-
49, Memory);
97+
oneapiext::parallel_for(Q, oneapiext::launch_config{sycl::range<1>{N}},
98+
Kernel1D, 49, Memory);
10099
Q.wait();
101100
for (size_t I = 0; I < N; ++I)
102101
Failed +=
103102
Check(Memory, 49, I, "1D parallel_for shortcut with launch config");
104103

105104
// 2D parallel_for shortcut with launch config
106-
oneapiext::parallel_for(
107-
Q, oneapiext::launch_config<sycl::range<2>>{sycl::range<2>{8, N / 8}},
108-
Kernel2D, 50, Memory);
105+
oneapiext::parallel_for(Q, oneapiext::launch_config{sycl::range<2>{8, N / 8}},
106+
Kernel2D, 50, Memory);
109107
Q.wait();
110108
for (size_t I = 0; I < N; ++I)
111109
Failed +=
112110
Check(Memory, 50, I, "2D parallel_for shortcut with launch config");
113111

114112
// 3D parallel_for shortcut with launch config
115113
oneapiext::parallel_for(
116-
Q, oneapiext::launch_config<sycl::range<3>>{sycl::range<3>{8, 8, N / 64}},
117-
Kernel3D, 51, Memory);
114+
Q, oneapiext::launch_config{sycl::range<3>{8, 8, N / 64}}, Kernel3D, 51,
115+
Memory);
118116
Q.wait();
119117
for (size_t I = 0; I < N; ++I)
120118
Failed +=
121119
Check(Memory, 51, I, "3D parallel_for shortcut with launch config");
122120

123121
// 1D nd_launch shortcut with launch config
124-
oneapiext::nd_launch(
125-
Q,
126-
oneapiext::launch_config<sycl::nd_range<1>>{
127-
sycl::nd_range<1>{sycl::range<1>{N}, sycl::range{8}}},
128-
Kernel1D, 52, Memory);
122+
oneapiext::nd_launch(Q,
123+
oneapiext::launch_config{sycl::nd_range<1>{
124+
sycl::range<1>{N}, sycl::range{8}}},
125+
Kernel1D, 52, Memory);
129126
Q.wait();
130127
for (size_t I = 0; I < N; ++I)
131128
Failed += Check(Memory, 52, I, "1D nd_launch shortcut with launch config");
132129

133130
// 2D nd_launch shortcut with launch config
134-
oneapiext::nd_launch(
135-
Q,
136-
oneapiext::launch_config<sycl::nd_range<2>>{
137-
sycl::nd_range<2>{sycl::range<2>{8, N / 8}, sycl::range{8, 8}}},
138-
Kernel2D, 53, Memory);
131+
oneapiext::nd_launch(Q,
132+
oneapiext::launch_config{sycl::nd_range<2>{
133+
sycl::range<2>{8, N / 8}, sycl::range{8, 8}}},
134+
Kernel2D, 53, Memory);
139135
Q.wait();
140136
for (size_t I = 0; I < N; ++I)
141137
Failed += Check(Memory, 53, I, "2D nd_launch shortcut with launch config");
142138

143139
// 3D nd_launch shortcut with launch config
144-
oneapiext::nd_launch(
145-
Q,
146-
oneapiext::launch_config<sycl::nd_range<3>>{sycl::nd_range<3>{
147-
sycl::range<3>{8, 8, N / 64}, sycl::range{8, 8, 8}}},
148-
Kernel3D, 54, Memory);
140+
oneapiext::nd_launch(Q,
141+
oneapiext::launch_config{sycl::nd_range<3>{
142+
sycl::range<3>{8, 8, N / 64}, sycl::range{8, 8, 8}}},
143+
Kernel3D, 54, Memory);
149144
Q.wait();
150145
for (size_t I = 0; I < N; ++I)
151146
Failed += Check(Memory, 54, I, "3D nd_launch shortcut with launch config");

‎sycl/test-e2e/EnqueueFunctions/kernel_shortcuts.cpp

Lines changed: 8 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -77,17 +77,16 @@ int main() {
7777
Failed += Check(Memory, 48, I, "3D nd_launch shortcut");
7878

7979
// 1D parallel_for shortcut with launch config
80-
oneapiext::parallel_for(
81-
Q, oneapiext::launch_config<sycl::range<1>>{sycl::range<1>{N}},
82-
[=](sycl::item<1> Item) { Memory[Item] = 49; });
80+
oneapiext::parallel_for(Q, oneapiext::launch_config{sycl::range<1>{N}},
81+
[=](sycl::item<1> Item) { Memory[Item] = 49; });
8382
Q.wait();
8483
for (size_t I = 0; I < N; ++I)
8584
Failed +=
8685
Check(Memory, 49, I, "1D parallel_for shortcut with launch config");
8786

8887
// 2D parallel_for shortcut with launch config
8988
oneapiext::parallel_for(
90-
Q, oneapiext::launch_config<sycl::range<2>>{sycl::range<2>{8, N / 8}},
89+
Q, oneapiext::launch_config{sycl::range<2>{8, N / 8}},
9190
[=](sycl::item<2> Item) { Memory[Item.get_linear_id()] = 50; });
9291
Q.wait();
9392
for (size_t I = 0; I < N; ++I)
@@ -96,7 +95,7 @@ int main() {
9695

9796
// 3D parallel_for shortcut with launch config
9897
oneapiext::parallel_for(
99-
Q, oneapiext::launch_config<sycl::range<3>>{sycl::range<3>{8, 8, N / 64}},
98+
Q, oneapiext::launch_config{sycl::range<3>{8, 8, N / 64}},
10099
[=](sycl::item<3> Item) { Memory[Item.get_linear_id()] = 51; });
101100
Q.wait();
102101
for (size_t I = 0; I < N; ++I)
@@ -106,7 +105,7 @@ int main() {
106105
// 1D nd_launch shortcut with launch config
107106
oneapiext::nd_launch(
108107
Q,
109-
oneapiext::launch_config<sycl::nd_range<1>>{
108+
oneapiext::launch_config{
110109
sycl::nd_range<1>{sycl::range<1>{N}, sycl::range{8}}},
111110
[=](sycl::nd_item<1> Item) { Memory[Item.get_global_linear_id()] = 52; });
112111
Q.wait();
@@ -116,7 +115,7 @@ int main() {
116115
// 2D nd_launch shortcut with launch config
117116
oneapiext::nd_launch(
118117
Q,
119-
oneapiext::launch_config<sycl::nd_range<2>>{
118+
oneapiext::launch_config{
120119
sycl::nd_range<2>{sycl::range<2>{8, N / 8}, sycl::range{8, 8}}},
121120
[=](sycl::nd_item<2> Item) { Memory[Item.get_global_linear_id()] = 53; });
122121
Q.wait();
@@ -126,8 +125,8 @@ int main() {
126125
// 3D nd_launch shortcut with launch config
127126
oneapiext::nd_launch(
128127
Q,
129-
oneapiext::launch_config<sycl::nd_range<3>>{sycl::nd_range<3>{
130-
sycl::range<3>{8, 8, N / 64}, sycl::range{8, 8, 8}}},
128+
oneapiext::launch_config{sycl::nd_range<3>{sycl::range<3>{8, 8, N / 64},
129+
sycl::range{8, 8, 8}}},
131130
[=](sycl::nd_item<3> Item) { Memory[Item.get_global_linear_id()] = 54; });
132131
Q.wait();
133132
for (size_t I = 0; I < N; ++I)

‎sycl/test-e2e/EnqueueFunctions/kernel_submit.cpp

Lines changed: 8 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -134,9 +134,9 @@ int main() {
134134
{
135135
oneapiext::submit(Q, [&](sycl::handler &CGH) {
136136
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
137-
oneapiext::parallel_for(
138-
CGH, oneapiext::launch_config<sycl::range<1>>{sycl::range<1>{N}},
139-
[=](sycl::item<1> Item) { MemAcc[Item] = 49; });
137+
oneapiext::parallel_for(CGH,
138+
oneapiext::launch_config{sycl::range<1>{N}},
139+
[=](sycl::item<1> Item) { MemAcc[Item] = 49; });
140140
});
141141
}
142142
}
@@ -150,8 +150,7 @@ int main() {
150150
{
151151
oneapiext::submit(Q, [&](sycl::handler &CGH) {
152152
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
153-
oneapiext::parallel_for(CGH,
154-
oneapiext::launch_config<sycl::range<2>>{Range},
153+
oneapiext::parallel_for(CGH, oneapiext::launch_config{Range},
155154
[=](sycl::item<2> Item) { MemAcc[Item] = 50; });
156155
});
157156
}
@@ -166,8 +165,7 @@ int main() {
166165
{
167166
oneapiext::submit(Q, [&](sycl::handler &CGH) {
168167
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
169-
oneapiext::parallel_for(CGH,
170-
oneapiext::launch_config<sycl::range<3>>{Range},
168+
oneapiext::parallel_for(CGH, oneapiext::launch_config{Range},
171169
[=](sycl::item<3> Item) { MemAcc[Item] = 51; });
172170
});
173171
}
@@ -183,7 +181,7 @@ int main() {
183181
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
184182
oneapiext::nd_launch(
185183
CGH,
186-
oneapiext::launch_config<sycl::nd_range<1>>{
184+
oneapiext::launch_config{
187185
sycl::nd_range<1>{sycl::range<1>{N}, sycl::range{8}}},
188186
[=](sycl::nd_item<1> Item) { MemAcc[Item.get_global_id()] = 52; });
189187
});
@@ -201,7 +199,7 @@ int main() {
201199
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
202200
oneapiext::nd_launch(
203201
CGH,
204-
oneapiext::launch_config<sycl::nd_range<2>>{
202+
oneapiext::launch_config{
205203
sycl::nd_range<2>{Range, sycl::range{8, 8}}},
206204
[=](sycl::nd_item<2> Item) { MemAcc[Item.get_global_id()] = 53; });
207205
});
@@ -219,7 +217,7 @@ int main() {
219217
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
220218
oneapiext::nd_launch(
221219
CGH,
222-
oneapiext::launch_config<sycl::nd_range<3>>{
220+
oneapiext::launch_config{
223221
sycl::nd_range<3>{Range, sycl::range{8, 8, 8}}},
224222
[=](sycl::nd_item<3> Item) { MemAcc[Item.get_global_id()] = 54; });
225223
});

‎sycl/test-e2e/EnqueueFunctions/kernel_submit_with_event.cpp

Lines changed: 8 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -134,9 +134,9 @@ int main() {
134134
{
135135
oneapiext::submit_with_event(Q, [&](sycl::handler &CGH) {
136136
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
137-
oneapiext::parallel_for(
138-
CGH, oneapiext::launch_config<sycl::range<1>>{sycl::range<1>{N}},
139-
[=](sycl::item<1> Item) { MemAcc[Item] = 49; });
137+
oneapiext::parallel_for(CGH,
138+
oneapiext::launch_config{sycl::range<1>{N}},
139+
[=](sycl::item<1> Item) { MemAcc[Item] = 49; });
140140
}).wait();
141141
}
142142
}
@@ -150,8 +150,7 @@ int main() {
150150
{
151151
oneapiext::submit_with_event(Q, [&](sycl::handler &CGH) {
152152
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
153-
oneapiext::parallel_for(CGH,
154-
oneapiext::launch_config<sycl::range<2>>{Range},
153+
oneapiext::parallel_for(CGH, oneapiext::launch_config{Range},
155154
[=](sycl::item<2> Item) { MemAcc[Item] = 50; });
156155
}).wait();
157156
}
@@ -166,8 +165,7 @@ int main() {
166165
{
167166
oneapiext::submit_with_event(Q, [&](sycl::handler &CGH) {
168167
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
169-
oneapiext::parallel_for(CGH,
170-
oneapiext::launch_config<sycl::range<3>>{Range},
168+
oneapiext::parallel_for(CGH, oneapiext::launch_config{Range},
171169
[=](sycl::item<3> Item) { MemAcc[Item] = 51; });
172170
}).wait();
173171
}
@@ -183,7 +181,7 @@ int main() {
183181
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
184182
oneapiext::nd_launch(
185183
CGH,
186-
oneapiext::launch_config<sycl::nd_range<1>>{
184+
oneapiext::launch_config{
187185
sycl::nd_range<1>{sycl::range<1>{N}, sycl::range{8}}},
188186
[=](sycl::nd_item<1> Item) { MemAcc[Item.get_global_id()] = 52; });
189187
}).wait();
@@ -201,7 +199,7 @@ int main() {
201199
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
202200
oneapiext::nd_launch(
203201
CGH,
204-
oneapiext::launch_config<sycl::nd_range<2>>{
202+
oneapiext::launch_config{
205203
sycl::nd_range<2>{Range, sycl::range{8, 8}}},
206204
[=](sycl::nd_item<2> Item) { MemAcc[Item.get_global_id()] = 53; });
207205
}).wait();
@@ -219,7 +217,7 @@ int main() {
219217
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
220218
oneapiext::nd_launch(
221219
CGH,
222-
oneapiext::launch_config<sycl::nd_range<3>>{
220+
oneapiext::launch_config{
223221
sycl::nd_range<3>{Range, sycl::range{8, 8, 8}}},
224222
[=](sycl::nd_item<3> Item) { MemAcc[Item.get_global_id()] = 54; });
225223
}).wait();

‎sycl/test-e2e/EnqueueFunctions/kernel_submit_with_event_and_kb.cpp

Lines changed: 16 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -149,9 +149,9 @@ int main() {
149149
{
150150
oneapiext::submit_with_event(Q, [&](sycl::handler &CGH) {
151151
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
152-
oneapiext::parallel_for(
153-
CGH, oneapiext::launch_config<sycl::range<1>>{sycl::range<1>{N}},
154-
Kernel1D, 49, MemAcc);
152+
oneapiext::parallel_for(CGH,
153+
oneapiext::launch_config{sycl::range<1>{N}},
154+
Kernel1D, 49, MemAcc);
155155
}).wait();
156156
}
157157
}
@@ -165,9 +165,8 @@ int main() {
165165
{
166166
oneapiext::submit_with_event(Q, [&](sycl::handler &CGH) {
167167
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
168-
oneapiext::parallel_for(CGH,
169-
oneapiext::launch_config<sycl::range<2>>{Range},
170-
Kernel2D, 50, MemAcc);
168+
oneapiext::parallel_for(CGH, oneapiext::launch_config{Range}, Kernel2D,
169+
50, MemAcc);
171170
}).wait();
172171
}
173172
}
@@ -181,9 +180,8 @@ int main() {
181180
{
182181
oneapiext::submit_with_event(Q, [&](sycl::handler &CGH) {
183182
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
184-
oneapiext::parallel_for(CGH,
185-
oneapiext::launch_config<sycl::range<3>>{Range},
186-
Kernel3D, 51, MemAcc);
183+
oneapiext::parallel_for(CGH, oneapiext::launch_config{Range}, Kernel3D,
184+
51, MemAcc);
187185
}).wait();
188186
}
189187
}
@@ -196,11 +194,10 @@ int main() {
196194
{
197195
oneapiext::submit_with_event(Q, [&](sycl::handler &CGH) {
198196
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
199-
oneapiext::nd_launch(
200-
CGH,
201-
oneapiext::launch_config<sycl::nd_range<1>>{
202-
sycl::nd_range<1>{sycl::range<1>{N}, sycl::range{8}}},
203-
Kernel1D, 52, MemAcc);
197+
oneapiext::nd_launch(CGH,
198+
oneapiext::launch_config{sycl::nd_range<1>{
199+
sycl::range<1>{N}, sycl::range{8}}},
200+
Kernel1D, 52, MemAcc);
204201
}).wait();
205202
}
206203
}
@@ -215,7 +212,7 @@ int main() {
215212
oneapiext::submit_with_event(Q, [&](sycl::handler &CGH) {
216213
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
217214
oneapiext::nd_launch(CGH,
218-
oneapiext::launch_config<sycl::nd_range<2>>{
215+
oneapiext::launch_config{
219216
sycl::nd_range<2>{Range, sycl::range{8, 8}}},
220217
Kernel2D, 53, MemAcc);
221218
}).wait();
@@ -231,11 +228,10 @@ int main() {
231228
{
232229
oneapiext::submit_with_event(Q, [&](sycl::handler &CGH) {
233230
sycl::accessor MemAcc{MemBuf, CGH, sycl::write_only};
234-
oneapiext::nd_launch(
235-
CGH,
236-
oneapiext::launch_config<sycl::nd_range<3>>{
237-
sycl::nd_range<3>{Range, sycl::range{8, 8, 8}}},
238-
Kernel3D, 54, MemAcc);
231+
oneapiext::nd_launch(CGH,
232+
oneapiext::launch_config{sycl::nd_range<3>{
233+
Range, sycl::range{8, 8, 8}}},
234+
Kernel3D, 54, MemAcc);
239235
}).wait();
240236
}
241237
}

0 commit comments

Comments
 (0)