Skip to content

Commit 4a6d21f

Browse files
dm-vodopyanovdm-vodopyanov
authored
[SYCL] Switch to using integration footer by default (#3777)
Added spec_constant_integration.hpp header file, which is included from integration footer. Added get_spec_constant_symbolic_ID_wrapper which in combination with spec_constant_integration.hpp allows us to avoid exploiting a UB related to instantiating a template before we have seen all its specializations. Replaced -fsycl-use-footer compiler flag with -fno-sycl-use-footer flag with the opposite semantics.
1 parent 0833052 commit 4a6d21f

26 files changed

+250
-168
lines changed

‎clang/include/clang/Driver/Options.td

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2619,9 +2619,9 @@ def fsycl_host_compiler_options_EQ : Joined<["-"], "fsycl-host-compiler-options=
26192619
Flags<[CoreOption]>, HelpText<"When performing the host compilation with "
26202620
"-fsycl-host-compiler specified, use the given options during that compile. "
26212621
"Options are expected to be a quoted list of space separated options.">;
2622-
def fsycl_use_footer : Flag<["-"], "fsycl-use-footer">, Flags<[CoreOption]>,
2623-
HelpText<"Enable usage of the integration footer during SYCL enabled "
2624-
"compilations.">;
2622+
def fno_sycl_use_footer : Flag<["-"], "fno-sycl-use-footer">, Flags<[CoreOption]>,
2623+
HelpText<"Disable usage of the integration footer during SYCL enabled "
2624+
"compilations.">;
26252625
def fsyntax_only : Flag<["-"], "fsyntax-only">,
26262626
Flags<[NoXarchOption,CoreOption,CC1Option,FC1Option]>, Group<Action_Group>;
26272627
def ftabstop_EQ : Joined<["-"], "ftabstop=">, Group<f_Group>;

‎clang/lib/Driver/Driver.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -5613,7 +5613,7 @@ Action *Driver::ConstructPhaseAction(
56135613
}
56145614
types::ID HostPPType = types::getPreprocessedType(Input->getType());
56155615
if (Args.hasArg(options::OPT_fsycl) && HostPPType != types::TY_INVALID &&
5616-
Args.hasArg(options::OPT_fsycl_use_footer) &&
5616+
!Args.hasArg(options::OPT_fno_sycl_use_footer) &&
56175617
TargetDeviceOffloadKind == Action::OFK_None) {
56185618
// Performing a host compilation with -fsycl. Append the integration
56195619
// footer to the preprocessed source file. We then add another

‎clang/lib/Driver/ToolChains/Clang.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1189,7 +1189,7 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
11891189
// When preprocessing using the integration footer, add the comments
11901190
// to the first preprocessing step.
11911191
if (JA.isOffloading(Action::OFK_SYCL) && !ContainsAppendFooterAction(&JA) &&
1192-
Args.hasArg(options::OPT_fsycl_use_footer) &&
1192+
!Args.hasArg(options::OPT_fno_sycl_use_footer) &&
11931193
JA.isDeviceOffloading(Action::OFK_None))
11941194
CmdArgs.push_back("-C");
11951195

@@ -4643,7 +4643,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
46434643
CmdArgs.push_back(Args.MakeArgString(HeaderOpt));
46444644
}
46454645

4646-
if (Args.hasArg(options::OPT_fsycl_use_footer)) {
4646+
if (!Args.hasArg(options::OPT_fno_sycl_use_footer)) {
46474647
// Add the integration footer option to generated the footer.
46484648
StringRef Footer(D.getIntegrationFooter(Input.getBaseInput()));
46494649
if (!Footer.empty()) {

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4965,6 +4965,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
49654965
Policy.SuppressTypedefs = true;
49664966
Policy.SuppressUnwrittenScope = true;
49674967

4968+
OS << "#include <CL/sycl/detail/defines_elementary.hpp>\n";
4969+
49684970
llvm::SmallSet<const VarDecl *, 8> VisitedSpecConstants;
49694971

49704972
// Used to uniquely name the 'shim's as we generate the names in each
@@ -4988,7 +4990,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
49884990
OS << "namespace sycl {\n";
49894991
OS << "namespace detail {\n";
49904992
OS << "template<>\n";
4991-
OS << "inline const char *get_spec_constant_symbolic_ID<";
4993+
OS << "inline const char *get_spec_constant_symbolic_ID_impl<";
49924994

49934995
if (VD->isInAnonymousNamespace()) {
49944996
OS << TopShim;
@@ -5008,6 +5010,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
50085010
}
50095011

50105012
OS << "#include <CL/sycl/detail/spec_const_integration.hpp>\n";
5013+
50115014
return true;
50125015
}
50135016

‎clang/test/CodeGenSYCL/anonymous_integration_footer.cpp

Lines changed: 23 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,10 @@ int main() {
99
cl::sycl::kernel_single_task<class first_kernel>([]() {});
1010
}
1111

12+
// CHECK: #include <CL/sycl/detail/defines_elementary.hpp>
13+
1214
using namespace cl;
15+
1316
// Example ways in which the application can declare a "specialization_id"
1417
// variable.
1518
struct S1 {
@@ -18,7 +21,7 @@ struct S1 {
1821
// CHECK-NEXT: namespace sycl {
1922
// CHECK-NEXT: namespace detail {
2023
// CHECK-NEXT: template<>
21-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S1::a>() {
24+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::S1::a>() {
2225
// CHECK-NEXT: return "_ZN2S11aE";
2326
// CHECK-NEXT: }
2427
// CHECK-NEXT: } // namespace detail
@@ -30,7 +33,7 @@ constexpr sycl::specialization_id b{2};
3033
// CHECK-NEXT: namespace sycl {
3134
// CHECK-NEXT: namespace detail {
3235
// CHECK-NEXT: template<>
33-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::b>() {
36+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::b>() {
3437
// CHECK-NEXT: return "____ZL1b";
3538
// CHECK-NEXT: }
3639
// CHECK-NEXT: } // namespace detail
@@ -41,7 +44,7 @@ inline constexpr sycl::specialization_id c{3};
4144
// CHECK-NEXT: namespace sycl {
4245
// CHECK-NEXT: namespace detail {
4346
// CHECK-NEXT: template<>
44-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::c>() {
47+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::c>() {
4548
// CHECK-NEXT: return "_Z1c";
4649
// CHECK-NEXT: }
4750
// CHECK-NEXT: } // namespace detail
@@ -52,7 +55,7 @@ static constexpr sycl::specialization_id d{4};
5255
// CHECK-NEXT: namespace sycl {
5356
// CHECK-NEXT: namespace detail {
5457
// CHECK-NEXT: template<>
55-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::d>() {
58+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::d>() {
5659
// CHECK-NEXT: return "____ZL1d";
5760
// CHECK-NEXT: }
5861
// CHECK-NEXT: } // namespace detail
@@ -73,7 +76,7 @@ struct S2 {
7376
// CHECK-NEXT: namespace sycl {
7477
// CHECK-NEXT: namespace detail {
7578
// CHECK-NEXT: template<>
76-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
79+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
7780
// CHECK-NEXT: return "____ZN12_GLOBAL__N_12S21aE";
7881
// CHECK-NEXT: }
7982
// CHECK-NEXT: } // namespace detail
@@ -91,7 +94,7 @@ template class S3<1>;
9194
// CHECK-NEXT: namespace sycl {
9295
// CHECK-NEXT: namespace detail {
9396
// CHECK-NEXT: template<>
94-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S3<1>::a>() {
97+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::S3<1>::a>() {
9598
// CHECK-NEXT: return "_ZN2S3ILi1EE1aE";
9699
// CHECK-NEXT: }
97100
// CHECK-NEXT: } // namespace detail
@@ -102,7 +105,7 @@ template class S3<2>;
102105
// CHECK-NEXT: namespace sycl {
103106
// CHECK-NEXT: namespace detail {
104107
// CHECK-NEXT: template<>
105-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S3<2>::a>() {
108+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::S3<2>::a>() {
106109
// CHECK-NEXT: return "_ZN2S3ILi2EE1aE";
107110
// CHECK-NEXT: }
108111
// CHECK-NEXT: } // namespace detail
@@ -115,7 +118,7 @@ constexpr sycl::specialization_id same_name{5};
115118
// CHECK-NEXT: namespace sycl {
116119
// CHECK-NEXT: namespace detail {
117120
// CHECK-NEXT: template<>
118-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::inner::same_name>() {
121+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::inner::same_name>() {
119122
// CHECK-NEXT: return "____ZN5innerL9same_nameE";
120123
// CHECK-NEXT: }
121124
// CHECK-NEXT: } // namespace detail
@@ -127,7 +130,7 @@ constexpr sycl::specialization_id same_name{6};
127130
// CHECK-NEXT: namespace sycl {
128131
// CHECK-NEXT: namespace detail {
129132
// CHECK-NEXT: template<>
130-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::same_name>() {
133+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::same_name>() {
131134
// CHECK-NEXT: return "____ZL9same_name";
132135
// CHECK-NEXT: }
133136
// CHECK-NEXT: } // namespace detail
@@ -146,7 +149,7 @@ constexpr sycl::specialization_id same_name{7};
146149
// CHECK-NEXT: namespace sycl {
147150
// CHECK-NEXT: namespace detail {
148151
// CHECK-NEXT: template<>
149-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
152+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
150153
// CHECK-NEXT: return "____ZN12_GLOBAL__N_19same_nameE";
151154
// CHECK-NEXT: }
152155
// CHECK-NEXT: } // namespace detail
@@ -167,7 +170,7 @@ constexpr sycl::specialization_id same_name{8};
167170
// CHECK-NEXT: namespace sycl {
168171
// CHECK-NEXT: namespace detail {
169172
// CHECK-NEXT: template<>
170-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
173+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
171174
// CHECK-NEXT: return "____ZN12_GLOBAL__N_15inner9same_nameE";
172175
// CHECK-NEXT: }
173176
// CHECK-NEXT: } // namespace detail
@@ -191,7 +194,7 @@ constexpr sycl::specialization_id same_name{9};
191194
// CHECK-NEXT: namespace sycl {
192195
// CHECK-NEXT: namespace detail {
193196
// CHECK-NEXT: template<>
194-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
197+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
195198
// CHECK-NEXT: return "____ZN5inner12_GLOBAL__N_19same_nameE";
196199
// CHECK-NEXT: }
197200
// CHECK-NEXT: } // namespace detail
@@ -206,7 +209,7 @@ constexpr sycl::specialization_id same_name{10};
206209
// CHECK-NEXT: namespace sycl {
207210
// CHECK-NEXT: namespace detail {
208211
// CHECK-NEXT: template<>
209-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::same_name>() {
212+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::same_name>() {
210213
// CHECK-NEXT: return "____ZN5outerL9same_nameE";
211214
// CHECK-NEXT: }
212215
// CHECK-NEXT: } // namespace detail
@@ -227,7 +230,7 @@ constexpr sycl::specialization_id same_name{11};
227230
// CHECK-NEXT: namespace sycl {
228231
// CHECK-NEXT: namespace detail {
229232
// CHECK-NEXT: template<>
230-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
233+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
231234
// CHECK-NEXT: return "____ZN5outer12_GLOBAL__N_19same_nameE";
232235
// CHECK-NEXT: }
233236
// CHECK-NEXT: } // namespace detail
@@ -249,7 +252,7 @@ constexpr sycl::specialization_id same_name{12};
249252
// CHECK-NEXT: namespace sycl {
250253
// CHECK-NEXT: namespace detail {
251254
// CHECK-NEXT: template<>
252-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
255+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
253256
// CHECK-NEXT: return "____ZN5outer12_GLOBAL__N_15inner9same_nameE";
254257
// CHECK-NEXT: }
255258
// CHECK-NEXT: } // namespace detail
@@ -287,7 +290,7 @@ constexpr sycl::specialization_id same_name{13};
287290
// CHECK-NEXT: namespace sycl {
288291
// CHECK-NEXT: namespace detail {
289292
// CHECK-NEXT: template<>
290-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID_2]]()>() {
293+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID_2]]()>() {
291294
// CHECK-NEXT: return "____ZN5outer12_GLOBAL__N_15inner12_GLOBAL__N_19same_nameE";
292295
// CHECK-NEXT: }
293296
// CHECK-NEXT: } // namespace detail
@@ -312,7 +315,7 @@ constexpr sycl::specialization_id same_name{14};
312315
// CHECK-NEXT: namespace sycl {
313316
// CHECK-NEXT: namespace detail {
314317
// CHECK-NEXT: template<>
315-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
318+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
316319
// CHECK-NEXT: return "____ZN12_GLOBAL__N_15outer9same_nameE";
317320
// CHECK-NEXT: }
318321
// CHECK-NEXT: } // namespace detail
@@ -342,7 +345,7 @@ constexpr sycl::specialization_id same_name{15};
342345
// CHECK-NEXT: namespace sycl {
343346
// CHECK-NEXT: namespace detail {
344347
// CHECK-NEXT: template<>
345-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID2]]()>() {
348+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID2]]()>() {
346349
// CHECK-NEXT: return "____ZN12_GLOBAL__N_15outer12_GLOBAL__N_19same_nameE";
347350
// CHECK-NEXT: }
348351
// CHECK-NEXT: } // namespace detail
@@ -372,7 +375,7 @@ constexpr sycl::specialization_id same_name{16};
372375
// CHECK-NEXT: namespace sycl {
373376
// CHECK-NEXT: namespace detail {
374377
// CHECK-NEXT: template<>
375-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID2]]()>() {
378+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID2]]()>() {
376379
// CHECK-NEXT: return "____ZN12_GLOBAL__N_15outer12_GLOBAL__N_15inner9same_nameE";
377380
// CHECK-NEXT: }
378381
// CHECK-NEXT: } // namespace detail
@@ -390,7 +393,7 @@ constexpr sycl::specialization_id same_name{17};
390393
// CHECK-NEXT: namespace sycl {
391394
// CHECK-NEXT: namespace detail {
392395
// CHECK-NEXT: template<>
393-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::inner::same_name>() {
396+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::inner::same_name>() {
394397
// CHECK-NEXT: return "____ZN5outer5innerL9same_nameE";
395398
// CHECK-NEXT: }
396399
// CHECK-NEXT: } // namespace detail

‎clang/test/CodeGenSYCL/anonymous_integration_footer2.cpp

Lines changed: 12 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,9 @@
88
int main() {
99
cl::sycl::kernel_single_task<class first_kernel>([]() {});
1010
}
11+
12+
// CHECK: #include <CL/sycl/detail/defines_elementary.hpp>
13+
1114
using namespace cl;
1215

1316
struct S1 {
@@ -17,7 +20,7 @@ struct S1 {
1720
// CHECK-NEXT: namespace sycl {
1821
// CHECK-NEXT: namespace detail {
1922
// CHECK-NEXT: template<>
20-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S1::a>() {
23+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::S1::a>() {
2124
// CHECK-NEXT: return "_ZN2S11aE";
2225
// CHECK-NEXT: }
2326
// CHECK-NEXT: } // namespace detail
@@ -29,7 +32,7 @@ constexpr sycl::specialization_id b{202};
2932
// CHECK-NEXT: namespace sycl {
3033
// CHECK-NEXT: namespace detail {
3134
// CHECK-NEXT: template<>
32-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::b>() {
35+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::b>() {
3336
// CHECK-NEXT: return "____ZL1b";
3437
// CHECK-NEXT: }
3538
// CHECK-NEXT: } // namespace detail
@@ -40,7 +43,7 @@ inline constexpr sycl::specialization_id c{3};
4043
// CHECK-NEXT: namespace sycl {
4144
// CHECK-NEXT: namespace detail {
4245
// CHECK-NEXT: template<>
43-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::c>() {
46+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::c>() {
4447
// CHECK-NEXT: return "_Z1c";
4548
// CHECK-NEXT: }
4649
// CHECK-NEXT: } // namespace detail
@@ -51,7 +54,7 @@ static constexpr sycl::specialization_id d{205};
5154
// CHECK-NEXT: namespace sycl {
5255
// CHECK-NEXT: namespace detail {
5356
// CHECK-NEXT: template<>
54-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::d>() {
57+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::d>() {
5558
// CHECK-NEXT: return "____ZL1d";
5659
// CHECK-NEXT: }
5760
// CHECK-NEXT: } // namespace detail
@@ -64,7 +67,7 @@ constexpr sycl::specialization_id same_name{5};
6467
// CHECK-NEXT: namespace sycl {
6568
// CHECK-NEXT: namespace detail {
6669
// CHECK-NEXT: template<>
67-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::inner::same_name>() {
70+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::inner::same_name>() {
6871
// CHECK-NEXT: return "____ZN5innerL9same_nameE";
6972
// CHECK-NEXT: }
7073
// CHECK-NEXT: } // namespace detail
@@ -76,7 +79,7 @@ constexpr sycl::specialization_id same_name{6};
7679
// CHECK-NEXT: namespace sycl {
7780
// CHECK-NEXT: namespace detail {
7881
// CHECK-NEXT: template<>
79-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::same_name>() {
82+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::same_name>() {
8083
// CHECK-NEXT: return "____ZL9same_name";
8184
// CHECK-NEXT: }
8285
// CHECK-NEXT: } // namespace detail
@@ -95,7 +98,7 @@ constexpr sycl::specialization_id same_name{207};
9598
// CHECK-NEXT: namespace sycl {
9699
// CHECK-NEXT: namespace detail {
97100
// CHECK-NEXT: template<>
98-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
101+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
99102
// CHECK-NEXT: return "____ZN12_GLOBAL__N_19same_nameE";
100103
// CHECK-NEXT: }
101104
// CHECK-NEXT: } // namespace detail
@@ -116,7 +119,7 @@ constexpr sycl::specialization_id same_name{208};
116119
// CHECK-NEXT: namespace sycl {
117120
// CHECK-NEXT: namespace detail {
118121
// CHECK-NEXT: template<>
119-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
122+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
120123
// CHECK-NEXT: return "____ZN12_GLOBAL__N_15inner9same_nameE";
121124
// CHECK-NEXT: }
122125
// CHECK-NEXT: } // namespace detail
@@ -143,7 +146,7 @@ constexpr sycl::specialization_id same_name{209};
143146
// CHECK-NEXT: namespace sycl {
144147
// CHECK-NEXT: namespace detail {
145148
// CHECK-NEXT: template<>
146-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
149+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
147150
// CHECK-NEXT: return "____ZN5outer5inner12_GLOBAL__N_19same_nameE";
148151
// CHECK-NEXT: }
149152
// CHECK-NEXT: } // namespace detail

0 commit comments

Comments
 (0)