Skip to content

Commit 1bd83e0

Browse files
dm-vodopyanovdm-vodopyanov
committed
Enable __builtin_unique_stable_name in FE; Make -fsycl-use-footer enabled by default for any app
1 parent 1bd1e94 commit 1bd83e0

24 files changed

+149
-141
lines changed

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

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2610,9 +2610,9 @@ def fsycl_host_compiler_options_EQ : Joined<["-"], "fsycl-host-compiler-options=
26102610
Flags<[CoreOption]>, HelpText<"When performing the host compilation with "
26112611
"-fsycl-host-compiler specified, use the given options during that compile. "
26122612
"Options are expected to be a quoted list of space separated options.">;
2613-
def fsycl_use_footer : Flag<["-"], "fsycl-use-footer">, Flags<[CoreOption]>,
2614-
HelpText<"Enable usage of the integration footer during SYCL enabled "
2615-
"compilations.">;
2613+
def fno_sycl_use_footer : Flag<["-"], "fno-sycl-use-footer">, Flags<[CoreOption]>,
2614+
HelpText<"Disable usage of the integration footer during SYCL enabled "
2615+
"compilations.">;
26162616
def fsyntax_only : Flag<["-"], "fsyntax-only">,
26172617
Flags<[NoXarchOption,CoreOption,CC1Option,FC1Option]>, Group<Action_Group>;
26182618
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
@@ -5584,7 +5584,7 @@ Action *Driver::ConstructPhaseAction(
55845584
if (Args.hasArg(options::OPT_verify_pch))
55855585
return C.MakeAction<VerifyPCHJobAction>(Input, types::TY_Nothing);
55865586
if (Args.hasArg(options::OPT_fsycl) &&
5587-
Args.hasArg(options::OPT_fsycl_use_footer) &&
5587+
!Args.hasArg(options::OPT_fno_sycl_use_footer) &&
55885588
TargetDeviceOffloadKind == Action::OFK_None) {
55895589
// Performing a host compilation with -fsycl. Append the integrated
55905590
// footer to the preprocessed source file. We then add another

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4621,7 +4621,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
46214621
CmdArgs.push_back(Args.MakeArgString(HeaderOpt));
46224622
}
46234623

4624-
if (Args.hasArg(options::OPT_fsycl_use_footer)) {
4624+
if (!Args.hasArg(options::OPT_fno_sycl_use_footer)) {
46254625
// Add the integration footer option to generated the footer.
46264626
StringRef Footer(D.getIntegrationFooter(Input.getBaseInput()));
46274627
if (!Footer.empty()) {

‎clang/lib/Sema/SemaSYCL.cpp‎

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -4752,6 +4752,10 @@ bool SYCLIntegrationFooter::emit(StringRef IntHeaderName) {
47524752
}
47534753

47544754
void SYCLIntegrationFooter::emitSpecIDName(raw_ostream &O, const VarDecl *VD) {
4755+
PrintingPolicy Policy{S.getLangOpts()};
4756+
Policy.adjustForCPlusPlusFwdDecl();
4757+
Policy.SuppressTypedefs = true;
4758+
Policy.SuppressUnwrittenScope = true;
47554759
// FIXME: Figure out the spec-constant unique name here.
47564760
// Note that this changes based on the linkage of the variable.
47574761
// We typically want to use the __builtin_unique_stable_name for the variable
@@ -4762,7 +4766,7 @@ void SYCLIntegrationFooter::emitSpecIDName(raw_ostream &O, const VarDecl *VD) {
47624766
// ahead of it, so that we make sure it is unique across translation units.
47634767
// This name should come from the yet implemented__builtin_unique_stable_name
47644768
// feature that accepts variables and gives the mangling for that.
4765-
O << "";
4769+
VD->printQualifiedName(O, Policy);
47664770
}
47674771

47684772
template <typename BeforeFn, typename AfterFn>
@@ -4891,13 +4895,14 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
48914895
Policy.SuppressTypedefs = true;
48924896
Policy.SuppressUnwrittenScope = true;
48934897

4898+
OS << "#include <CL/sycl/detail/defines_elementary.hpp>\n";
4899+
48944900
// Used to uniquely name the 'shim's as we generate the names in each
48954901
// anonymous namespace.
48964902
unsigned ShimCounter = 0;
48974903
for (const VarDecl *VD : SpecConstants) {
48984904
VD = VD->getCanonicalDecl();
48994905
std::string TopShim = EmitSpecIdShims(OS, ShimCounter, VD);
4900-
OS << "#include <CL/sycl/detail/defines_elementary.hpp>\n";
49014906
OS << "__SYCL_INLINE_NAMESPACE(cl) {\n";
49024907
OS << "namespace sycl {\n";
49034908
OS << "namespace detail {\n";
@@ -4912,9 +4917,10 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
49124917
}
49134918

49144919
OS << ">() {\n";
4915-
OS << " return \"";
4920+
OS << " return "
4921+
"__builtin_unique_stable_name(specialization_id_name_generator<";
49164922
emitSpecIDName(OS, VD);
4917-
OS << "\";\n";
4923+
OS << ">);\n";
49184924
OS << "}\n";
49194925
OS << "} // namespace detail\n";
49204926
OS << "} // namespace sycl\n";

‎clang/test/CodeGenSYCL/anonymous_integration_footer.cpp‎

Lines changed: 38 additions & 37 deletions
Original file line numberDiff line numberDiff line change
@@ -14,12 +14,13 @@ using namespace cl;
1414
// variable.
1515
struct S1 {
1616
static constexpr sycl::specialization_id a{1};
17-
// CHECK: __SYCL_INLINE_NAMESPACE(cl) {
17+
// CHECK: #include <CL/sycl/detail/defines_elementary.hpp>
18+
// CHECK-NEXT: __SYCL_INLINE_NAMESPACE(cl) {
1819
// CHECK-NEXT: namespace sycl {
1920
// CHECK-NEXT: namespace detail {
2021
// CHECK-NEXT: template<>
2122
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S1::a>() {
22-
// CHECK-NEXT: return "";
23+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<S1::a>);
2324
// CHECK-NEXT: }
2425
// CHECK-NEXT: } // namespace detail
2526
// CHECK-NEXT: } // namespace sycl
@@ -31,7 +32,7 @@ constexpr sycl::specialization_id b{2};
3132
// CHECK-NEXT: namespace detail {
3233
// CHECK-NEXT: template<>
3334
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::b>() {
34-
// CHECK-NEXT: return "";
35+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<b>);
3536
// CHECK-NEXT: }
3637
// CHECK-NEXT: } // namespace detail
3738
// CHECK-NEXT: } // namespace sycl
@@ -42,7 +43,7 @@ inline constexpr sycl::specialization_id c{3};
4243
// CHECK-NEXT: namespace detail {
4344
// CHECK-NEXT: template<>
4445
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::c>() {
45-
// CHECK-NEXT: return "";
46+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<c>);
4647
// CHECK-NEXT: }
4748
// CHECK-NEXT: } // namespace detail
4849
// CHECK-NEXT: } // namespace sycl
@@ -53,7 +54,7 @@ static constexpr sycl::specialization_id d{4};
5354
// CHECK-NEXT: namespace detail {
5455
// CHECK-NEXT: template<>
5556
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::d>() {
56-
// CHECK-NEXT: return "";
57+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<d>);
5758
// CHECK-NEXT: }
5859
// CHECK-NEXT: } // namespace detail
5960
// CHECK-NEXT: } // namespace sycl
@@ -62,23 +63,23 @@ static constexpr sycl::specialization_id d{4};
6263
namespace {
6364
struct S2 {
6465
static constexpr sycl::specialization_id a{18};
65-
// CHECK-NEXT: namespace {
66-
// CHECK-NEXT: namespace __sycl_detail {
67-
// CHECK-NEXT: static constexpr decltype(S2::a) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() {
68-
// CHECK-NEXT: return S2::a;
69-
// CHECK-NEXT: }
70-
// CHECK-NEXT: } // namespace __sycl_detail
71-
// CHECK-NEXT: } // namespace
72-
// CHECK-NEXT: __SYCL_INLINE_NAMESPACE(cl) {
73-
// CHECK-NEXT: namespace sycl {
74-
// CHECK-NEXT: namespace detail {
75-
// CHECK-NEXT: template<>
76-
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
77-
// CHECK-NEXT: return "";
78-
// CHECK-NEXT: }
79-
// CHECK-NEXT: } // namespace detail
80-
// CHECK-NEXT: } // namespace sycl
81-
// CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
66+
// CHECK-NEXT: namespace {
67+
// CHECK-NEXT: namespace __sycl_detail {
68+
// CHECK-NEXT: static constexpr decltype(S2::a) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() {
69+
// CHECK-NEXT: return S2::a;
70+
// CHECK-NEXT: }
71+
// CHECK-NEXT: } // namespace __sycl_detail
72+
// CHECK-NEXT: } // namespace
73+
// CHECK-NEXT: __SYCL_INLINE_NAMESPACE(cl) {
74+
// CHECK-NEXT: namespace sycl {
75+
// CHECK-NEXT: namespace detail {
76+
// CHECK-NEXT: template<>
77+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
78+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<S2::a>);
79+
// CHECK-NEXT: }
80+
// CHECK-NEXT: } // namespace detail
81+
// CHECK-NEXT: } // namespace sycl
82+
// CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
8283
};
8384
} // namespace
8485

@@ -92,7 +93,7 @@ template class S3<1>;
9293
// CHECK-NEXT: namespace detail {
9394
// CHECK-NEXT: template<>
9495
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S3<1>::a>() {
95-
// CHECK-NEXT: return "";
96+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<S3<1>::a>);
9697
// CHECK-NEXT: }
9798
// CHECK-NEXT: } // namespace detail
9899
// CHECK-NEXT: } // namespace sycl
@@ -103,7 +104,7 @@ template class S3<2>;
103104
// CHECK-NEXT: namespace detail {
104105
// CHECK-NEXT: template<>
105106
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S3<2>::a>() {
106-
// CHECK-NEXT: return "";
107+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<S3<2>::a>);
107108
// CHECK-NEXT: }
108109
// CHECK-NEXT: } // namespace detail
109110
// CHECK-NEXT: } // namespace sycl
@@ -116,7 +117,7 @@ constexpr sycl::specialization_id same_name{5};
116117
// CHECK-NEXT: namespace detail {
117118
// CHECK-NEXT: template<>
118119
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::inner::same_name>() {
119-
// CHECK-NEXT: return "";
120+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<inner::same_name>);
120121
// CHECK-NEXT: }
121122
// CHECK-NEXT: } // namespace detail
122123
// CHECK-NEXT: } // namespace sycl
@@ -128,7 +129,7 @@ constexpr sycl::specialization_id same_name{6};
128129
// CHECK-NEXT: namespace detail {
129130
// CHECK-NEXT: template<>
130131
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::same_name>() {
131-
// CHECK-NEXT: return "";
132+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<same_name>);
132133
// CHECK-NEXT: }
133134
// CHECK-NEXT: } // namespace detail
134135
// CHECK-NEXT: } // namespace sycl
@@ -147,7 +148,7 @@ constexpr sycl::specialization_id same_name{7};
147148
// CHECK-NEXT: namespace detail {
148149
// CHECK-NEXT: template<>
149150
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
150-
// CHECK-NEXT: return "";
151+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<same_name>);
151152
// CHECK-NEXT: }
152153
// CHECK-NEXT: } // namespace detail
153154
// CHECK-NEXT: } // namespace sycl
@@ -168,7 +169,7 @@ constexpr sycl::specialization_id same_name{8};
168169
// CHECK-NEXT: namespace detail {
169170
// CHECK-NEXT: template<>
170171
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
171-
// CHECK-NEXT: return "";
172+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<inner::same_name>);
172173
// CHECK-NEXT: }
173174
// CHECK-NEXT: } // namespace detail
174175
// CHECK-NEXT: } // namespace sycl
@@ -192,7 +193,7 @@ constexpr sycl::specialization_id same_name{9};
192193
// CHECK-NEXT: namespace detail {
193194
// CHECK-NEXT: template<>
194195
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
195-
// CHECK-NEXT: return "";
196+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<inner::same_name>);
196197
// CHECK-NEXT: }
197198
// CHECK-NEXT: } // namespace detail
198199
// CHECK-NEXT: } // namespace sycl
@@ -207,7 +208,7 @@ constexpr sycl::specialization_id same_name{10};
207208
// CHECK-NEXT: namespace detail {
208209
// CHECK-NEXT: template<>
209210
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::same_name>() {
210-
// CHECK-NEXT: return "";
211+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<outer::same_name>);
211212
// CHECK-NEXT: }
212213
// CHECK-NEXT: } // namespace detail
213214
// CHECK-NEXT: } // namespace sycl
@@ -228,7 +229,7 @@ constexpr sycl::specialization_id same_name{11};
228229
// CHECK-NEXT: namespace detail {
229230
// CHECK-NEXT: template<>
230231
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
231-
// CHECK-NEXT: return "";
232+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<outer::same_name>);
232233
// CHECK-NEXT: }
233234
// CHECK-NEXT: } // namespace detail
234235
// CHECK-NEXT: } // namespace sycl
@@ -250,7 +251,7 @@ constexpr sycl::specialization_id same_name{12};
250251
// CHECK-NEXT: namespace detail {
251252
// CHECK-NEXT: template<>
252253
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
253-
// CHECK-NEXT: return "";
254+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<outer::inner::same_name>);
254255
// CHECK-NEXT: }
255256
// CHECK-NEXT: } // namespace detail
256257
// CHECK-NEXT: } // namespace sycl
@@ -288,7 +289,7 @@ constexpr sycl::specialization_id same_name{13};
288289
// CHECK-NEXT: namespace detail {
289290
// CHECK-NEXT: template<>
290291
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID_2]]()>() {
291-
// CHECK-NEXT: return "";
292+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<outer::inner::same_name>);
292293
// CHECK-NEXT: }
293294
// CHECK-NEXT: } // namespace detail
294295
// CHECK-NEXT: } // namespace sycl
@@ -313,7 +314,7 @@ constexpr sycl::specialization_id same_name{14};
313314
// CHECK-NEXT: namespace detail {
314315
// CHECK-NEXT: template<>
315316
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
316-
// CHECK-NEXT: return "";
317+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<outer::same_name>);
317318
// CHECK-NEXT: }
318319
// CHECK-NEXT: } // namespace detail
319320
// CHECK-NEXT: } // namespace sycl
@@ -343,7 +344,7 @@ constexpr sycl::specialization_id same_name{15};
343344
// CHECK-NEXT: namespace detail {
344345
// CHECK-NEXT: template<>
345346
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID2]]()>() {
346-
// CHECK-NEXT: return "";
347+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<outer::same_name>);
347348
// CHECK-NEXT: }
348349
// CHECK-NEXT: } // namespace detail
349350
// CHECK-NEXT: } // namespace sycl
@@ -373,7 +374,7 @@ constexpr sycl::specialization_id same_name{16};
373374
// CHECK-NEXT: namespace detail {
374375
// CHECK-NEXT: template<>
375376
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID2]]()>() {
376-
// CHECK-NEXT: return "";
377+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<outer::inner::same_name>);
377378
// CHECK-NEXT: }
378379
// CHECK-NEXT: } // namespace detail
379380
// CHECK-NEXT: } // namespace sycl
@@ -391,7 +392,7 @@ constexpr sycl::specialization_id same_name{17};
391392
// CHECK-NEXT: namespace detail {
392393
// CHECK-NEXT: template<>
393394
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::inner::same_name>() {
394-
// CHECK-NEXT: return "";
395+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<outer::inner::same_name>);
395396
// CHECK-NEXT: }
396397
// CHECK-NEXT: } // namespace detail
397398
// CHECK-NEXT: } // namespace sycl

‎clang/test/CodeGenSYCL/anonymous_integration_footer2.cpp‎

Lines changed: 11 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -13,12 +13,13 @@ using namespace cl;
1313
struct S1 {
1414
static constexpr sycl::specialization_id a{1};
1515
};
16-
// CHECK: __SYCL_INLINE_NAMESPACE(cl) {
16+
// CHECK: #include <CL/sycl/detail/defines_elementary.hpp>
17+
// CHECK-NEXT: __SYCL_INLINE_NAMESPACE(cl) {
1718
// CHECK-NEXT: namespace sycl {
1819
// CHECK-NEXT: namespace detail {
1920
// CHECK-NEXT: template<>
2021
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S1::a>() {
21-
// CHECK-NEXT: return "";
22+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<S1::a>);
2223
// CHECK-NEXT: }
2324
// CHECK-NEXT: } // namespace detail
2425
// CHECK-NEXT: } // namespace sycl
@@ -30,7 +31,7 @@ constexpr sycl::specialization_id b{202};
3031
// CHECK-NEXT: namespace detail {
3132
// CHECK-NEXT: template<>
3233
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::b>() {
33-
// CHECK-NEXT: return "";
34+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<b>);
3435
// CHECK-NEXT: }
3536
// CHECK-NEXT: } // namespace detail
3637
// CHECK-NEXT: } // namespace sycl
@@ -41,7 +42,7 @@ inline constexpr sycl::specialization_id c{3};
4142
// CHECK-NEXT: namespace detail {
4243
// CHECK-NEXT: template<>
4344
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::c>() {
44-
// CHECK-NEXT: return "";
45+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<c>);
4546
// CHECK-NEXT: }
4647
// CHECK-NEXT: } // namespace detail
4748
// CHECK-NEXT: } // namespace sycl
@@ -52,7 +53,7 @@ static constexpr sycl::specialization_id d{205};
5253
// CHECK-NEXT: namespace detail {
5354
// CHECK-NEXT: template<>
5455
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::d>() {
55-
// CHECK-NEXT: return "";
56+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<d>);
5657
// CHECK-NEXT: }
5758
// CHECK-NEXT: } // namespace detail
5859
// CHECK-NEXT: } // namespace sycl
@@ -65,7 +66,7 @@ constexpr sycl::specialization_id same_name{5};
6566
// CHECK-NEXT: namespace detail {
6667
// CHECK-NEXT: template<>
6768
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::inner::same_name>() {
68-
// CHECK-NEXT: return "";
69+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<inner::same_name>);
6970
// CHECK-NEXT: }
7071
// CHECK-NEXT: } // namespace detail
7172
// CHECK-NEXT: } // namespace sycl
@@ -77,7 +78,7 @@ constexpr sycl::specialization_id same_name{6};
7778
// CHECK-NEXT: namespace detail {
7879
// CHECK-NEXT: template<>
7980
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::same_name>() {
80-
// CHECK-NEXT: return "";
81+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<same_name>);
8182
// CHECK-NEXT: }
8283
// CHECK-NEXT: } // namespace detail
8384
// CHECK-NEXT: } // namespace sycl
@@ -96,7 +97,7 @@ constexpr sycl::specialization_id same_name{207};
9697
// CHECK-NEXT: namespace detail {
9798
// CHECK-NEXT: template<>
9899
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
99-
// CHECK-NEXT: return "";
100+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<same_name>);
100101
// CHECK-NEXT: }
101102
// CHECK-NEXT: } // namespace detail
102103
// CHECK-NEXT: } // namespace sycl
@@ -117,7 +118,7 @@ constexpr sycl::specialization_id same_name{208};
117118
// CHECK-NEXT: namespace detail {
118119
// CHECK-NEXT: template<>
119120
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
120-
// CHECK-NEXT: return "";
121+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<inner::same_name>);
121122
// CHECK-NEXT: }
122123
// CHECK-NEXT: } // namespace detail
123124
// CHECK-NEXT: } // namespace sycl
@@ -144,7 +145,7 @@ constexpr sycl::specialization_id same_name{209};
144145
// CHECK-NEXT: namespace detail {
145146
// CHECK-NEXT: template<>
146147
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
147-
// CHECK-NEXT: return "";
148+
// CHECK-NEXT: return __builtin_unique_stable_name(specialization_id_name_generator<outer::same_name>);
148149
// CHECK-NEXT: }
149150
// CHECK-NEXT: } // namespace detail
150151
// CHECK-NEXT: } // namespace sycl

0 commit comments

Comments
 (0)