Skip to content

[SYCL] Switch to using integration footer by default #3777

New issue

Have a question about this project? # for a free GitHub account to open an issue and contact its maintainers and the community.

#

By clicking “#”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? # to your account

Merged
merged 34 commits into from
Jun 30, 2021
Merged

[SYCL] Switch to using integration footer by default #3777

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
34 commits
Select commit Hold shift + click to select a range
a6f5b4e
[SYCL] Move __buildin_unique_stable_name dependent funcs to a separat…
dm-vodopyanov May 18, 2021
dd66b44
Add missing EOL
dm-vodopyanov May 18, 2021
22f9009
Remove extra endif
dm-vodopyanov May 18, 2021
3db201f
Merge branch 'intel:sycl' into private/dvodopya/move-__buildin_unique…
dm-vodopyanov May 20, 2021
68187a5
Merge branch 'intel:sycl' into private/dvodopya/move-__buildin_unique…
dm-vodopyanov May 27, 2021
1bd1e94
Add -fsycl-use-footer to tests and include missing header to int footer
dm-vodopyanov May 31, 2021
1bd83e0
Enable __builtin_unique_stable_name in FE; Make -fsycl-use-footer ena…
dm-vodopyanov Jun 1, 2021
0c9e8e2
Merge branch 'sycl' into private/dvodopya/move-__buildin_unique_stabl…
dm-vodopyanov Jun 9, 2021
0d3e000
Fix some test failures
dm-vodopyanov Jun 9, 2021
224eb1a
Merge branch 'sycl' into private/dvodopya/move-__buildin_unique_stabl…
dm-vodopyanov Jun 15, 2021
ed1686d
Update PR to use new builtin
dm-vodopyanov Jun 15, 2021
feb54c4
Fix SYCL :: basic_tests/SYCL-2020*** tests
dm-vodopyanov Jun 15, 2021
e56ee41
Fix check-clang
AlexeySachkov Jun 16, 2021
30ce3d0
Merge remote-tracking branch 'sycl' into private/dvodopya/move-__buil…
AlexeySachkov Jun 16, 2021
7669de4
Add "-fsycl-device-only" to ESIMD diagnostics tests
AlexeySachkov Jun 16, 2021
d86c37a
temp
AlexeySachkov Jun 16, 2021
33f1ed9
Fix host-only tests for diagnostics
AlexeySachkov Jun 18, 2021
dcf1839
Merge remote-tracking branch 'sycl' into private/dvodopya/move-__buil…
AlexeySachkov Jun 18, 2021
b3a7076
Fix one more ESIMD test after merge from upstream
AlexeySachkov Jun 18, 2021
a513769
temp
AlexeySachkov Jun 21, 2021
e94914d
Merge branch 'intel:sycl' into private/dvodopya/move-__buildin_unique…
dm-vodopyanov Jun 23, 2021
7050d43
Merge remote-tracking branch 'sycl' into private/dvodopya/move-__buil…
AlexeySachkov Jun 25, 2021
b69f889
Fix UB related to get_spec_constant_symbolic_ID
AlexeySachkov Jun 25, 2021
3ef8887
Merge branch 'private/dvodopya/move-__buildin_unique_stable_name-to-s…
AlexeySachkov Jun 25, 2021
77e7988
Merge remote-tracking branch 'sycl' into private/dvodopya/move-__buil…
AlexeySachkov Jun 25, 2021
68c1f1a
Adjust sycl-offload-amdgcn.cpp test
AlexeySachkov Jun 28, 2021
30de609
Merge remote-tracking branch 'sycl' into private/dvodopya/move-__buil…
AlexeySachkov Jun 28, 2021
37be540
Add a test for -fno-sycl-use-footer
AlexeySachkov Jun 28, 2021
48e59c9
Rename get_spec_constant_symbolic_ID
AlexeySachkov Jun 28, 2021
b73a756
Fix comment and clang-format
dm-vodopyanov Jun 28, 2021
d5765f4
Merge remote-tracking branch 'sycl' into private/dvodopya/move-__buil…
AlexeySachkov Jun 29, 2021
8e58b19
Add test for compiler phases w/ and w/o integration footer
AlexeySachkov Jun 29, 2021
d5e10a2
Merge branch 'private/dvodopya/move-__buildin_unique_stable_name-to-s…
AlexeySachkov Jun 29, 2021
9476403
An attempt to make sycl-int-footer.cpp OS-independent
AlexeySachkov Jun 29, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions clang/include/clang/Driver/Options.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2619,9 +2619,9 @@ def fsycl_host_compiler_options_EQ : Joined<["-"], "fsycl-host-compiler-options=
Flags<[CoreOption]>, HelpText<"When performing the host compilation with "
"-fsycl-host-compiler specified, use the given options during that compile. "
"Options are expected to be a quoted list of space separated options.">;
def fsycl_use_footer : Flag<["-"], "fsycl-use-footer">, Flags<[CoreOption]>,
HelpText<"Enable usage of the integration footer during SYCL enabled "
"compilations.">;
def fno_sycl_use_footer : Flag<["-"], "fno-sycl-use-footer">, Flags<[CoreOption]>,
HelpText<"Disable usage of the integration footer during SYCL enabled "
"compilations.">;
def fsyntax_only : Flag<["-"], "fsyntax-only">,
Flags<[NoXarchOption,CoreOption,CC1Option,FC1Option]>, Group<Action_Group>;
def ftabstop_EQ : Joined<["-"], "ftabstop=">, Group<f_Group>;
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/Driver/Driver.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5612,7 +5612,7 @@ Action *Driver::ConstructPhaseAction(
}
types::ID HostPPType = types::getPreprocessedType(Input->getType());
if (Args.hasArg(options::OPT_fsycl) && HostPPType != types::TY_INVALID &&
Args.hasArg(options::OPT_fsycl_use_footer) &&
!Args.hasArg(options::OPT_fno_sycl_use_footer) &&
TargetDeviceOffloadKind == Action::OFK_None) {
// Performing a host compilation with -fsycl. Append the integration
// footer to the preprocessed source file. We then add another
Expand Down
4 changes: 2 additions & 2 deletions clang/lib/Driver/ToolChains/Clang.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1189,7 +1189,7 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
// When preprocessing using the integration footer, add the comments
// to the first preprocessing step.
if (JA.isOffloading(Action::OFK_SYCL) && !ContainsAppendFooterAction(&JA) &&
Args.hasArg(options::OPT_fsycl_use_footer) &&
!Args.hasArg(options::OPT_fno_sycl_use_footer) &&
JA.isDeviceOffloading(Action::OFK_None))
CmdArgs.push_back("-C");

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

if (Args.hasArg(options::OPT_fsycl_use_footer)) {
if (!Args.hasArg(options::OPT_fno_sycl_use_footer)) {
// Add the integration footer option to generated the footer.
StringRef Footer(D.getIntegrationFooter(Input.getBaseInput()));
if (!Footer.empty()) {
Expand Down
5 changes: 4 additions & 1 deletion clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4965,6 +4965,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
Policy.SuppressTypedefs = true;
Policy.SuppressUnwrittenScope = true;

OS << "#include <CL/sycl/detail/defines_elementary.hpp>\n";

llvm::SmallSet<const VarDecl *, 8> VisitedSpecConstants;

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

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

OS << "#include <CL/sycl/detail/spec_const_integration.hpp>\n";

return true;
}

Expand Down
43 changes: 23 additions & 20 deletions clang/test/CodeGenSYCL/anonymous_integration_footer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,10 @@ int main() {
cl::sycl::kernel_single_task<class first_kernel>([]() {});
}

// CHECK: #include <CL/sycl/detail/defines_elementary.hpp>

using namespace cl;

// Example ways in which the application can declare a "specialization_id"
// variable.
struct S1 {
Expand All @@ -18,7 +21,7 @@ struct S1 {
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S1::a>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::S1::a>() {
// CHECK-NEXT: return "_ZN2S11aE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -30,7 +33,7 @@ constexpr sycl::specialization_id b{2};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::b>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::b>() {
// CHECK-NEXT: return "____ZL1b";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -41,7 +44,7 @@ inline constexpr sycl::specialization_id c{3};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::c>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::c>() {
// CHECK-NEXT: return "_Z1c";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -52,7 +55,7 @@ static constexpr sycl::specialization_id d{4};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::d>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::d>() {
// CHECK-NEXT: return "____ZL1d";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -73,7 +76,7 @@ struct S2 {
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: return "____ZN12_GLOBAL__N_12S21aE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -91,7 +94,7 @@ template class S3<1>;
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S3<1>::a>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::S3<1>::a>() {
// CHECK-NEXT: return "_ZN2S3ILi1EE1aE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -102,7 +105,7 @@ template class S3<2>;
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S3<2>::a>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::S3<2>::a>() {
// CHECK-NEXT: return "_ZN2S3ILi2EE1aE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -115,7 +118,7 @@ constexpr sycl::specialization_id same_name{5};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::inner::same_name>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::inner::same_name>() {
// CHECK-NEXT: return "____ZN5innerL9same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -127,7 +130,7 @@ constexpr sycl::specialization_id same_name{6};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::same_name>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::same_name>() {
// CHECK-NEXT: return "____ZL9same_name";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -146,7 +149,7 @@ constexpr sycl::specialization_id same_name{7};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: return "____ZN12_GLOBAL__N_19same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -167,7 +170,7 @@ constexpr sycl::specialization_id same_name{8};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: return "____ZN12_GLOBAL__N_15inner9same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -191,7 +194,7 @@ constexpr sycl::specialization_id same_name{9};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: return "____ZN5inner12_GLOBAL__N_19same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -206,7 +209,7 @@ constexpr sycl::specialization_id same_name{10};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::same_name>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::same_name>() {
// CHECK-NEXT: return "____ZN5outerL9same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -227,7 +230,7 @@ constexpr sycl::specialization_id same_name{11};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: return "____ZN5outer12_GLOBAL__N_19same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -249,7 +252,7 @@ constexpr sycl::specialization_id same_name{12};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: return "____ZN5outer12_GLOBAL__N_15inner9same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand Down Expand Up @@ -287,7 +290,7 @@ constexpr sycl::specialization_id same_name{13};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID_2]]()>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID_2]]()>() {
// CHECK-NEXT: return "____ZN5outer12_GLOBAL__N_15inner12_GLOBAL__N_19same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -312,7 +315,7 @@ constexpr sycl::specialization_id same_name{14};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: return "____ZN12_GLOBAL__N_15outer9same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand Down Expand Up @@ -342,7 +345,7 @@ constexpr sycl::specialization_id same_name{15};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID2]]()>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID2]]()>() {
// CHECK-NEXT: return "____ZN12_GLOBAL__N_15outer12_GLOBAL__N_19same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand Down Expand Up @@ -372,7 +375,7 @@ constexpr sycl::specialization_id same_name{16};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID2]]()>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID2]]()>() {
// CHECK-NEXT: return "____ZN12_GLOBAL__N_15outer12_GLOBAL__N_15inner9same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -390,7 +393,7 @@ constexpr sycl::specialization_id same_name{17};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::inner::same_name>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::inner::same_name>() {
// CHECK-NEXT: return "____ZN5outer5innerL9same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand Down
21 changes: 12 additions & 9 deletions clang/test/CodeGenSYCL/anonymous_integration_footer2.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,9 @@
int main() {
cl::sycl::kernel_single_task<class first_kernel>([]() {});
}

// CHECK: #include <CL/sycl/detail/defines_elementary.hpp>

using namespace cl;

struct S1 {
Expand All @@ -17,7 +20,7 @@ struct S1 {
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S1::a>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::S1::a>() {
// CHECK-NEXT: return "_ZN2S11aE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -29,7 +32,7 @@ constexpr sycl::specialization_id b{202};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::b>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::b>() {
// CHECK-NEXT: return "____ZL1b";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -40,7 +43,7 @@ inline constexpr sycl::specialization_id c{3};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::c>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::c>() {
// CHECK-NEXT: return "_Z1c";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -51,7 +54,7 @@ static constexpr sycl::specialization_id d{205};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::d>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::d>() {
// CHECK-NEXT: return "____ZL1d";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -64,7 +67,7 @@ constexpr sycl::specialization_id same_name{5};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::inner::same_name>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::inner::same_name>() {
// CHECK-NEXT: return "____ZN5innerL9same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -76,7 +79,7 @@ constexpr sycl::specialization_id same_name{6};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::same_name>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::same_name>() {
// CHECK-NEXT: return "____ZL9same_name";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -95,7 +98,7 @@ constexpr sycl::specialization_id same_name{207};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: return "____ZN12_GLOBAL__N_19same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -116,7 +119,7 @@ constexpr sycl::specialization_id same_name{208};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: return "____ZN12_GLOBAL__N_15inner9same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand All @@ -143,7 +146,7 @@ constexpr sycl::specialization_id same_name{209};
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: return "____ZN5outer5inner12_GLOBAL__N_19same_nameE";
// CHECK-NEXT: }
// CHECK-NEXT: } // namespace detail
Expand Down
Loading