Skip to content

Commit 50d7a9c

Browse files
author
Sergey Kanaev
committed
Merge remote-tracking branch 'public/sycl' into fix-tests
2 parents b497bbb + 807fdae commit 50d7a9c

File tree

63 files changed

+928
-599
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

63 files changed

+928
-599
lines changed

‎buildbot/dependency.conf

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -4,8 +4,8 @@ ocl_cpu_rt_ver=2021.12.6.0.19
44
# https://github.com/intel/llvm/releases/download/2021-WW26/win-oclcpuexp-2021.12.6.0.19_rel.zip
55
ocl_cpu_rt_ver_win=2021.12.6.0.19
66
# Same GPU driver supports Level Zero and OpenCL
7-
# https://github.com/intel/compute-runtime/releases/tag/21.23.20043
8-
ocl_gpu_rt_ver=21.23.20043
7+
# https://github.com/intel/compute-runtime/releases/tag/21.24.20098
8+
ocl_gpu_rt_ver=21.24.20098
99
# Same GPU driver supports Level Zero and OpenCL
1010
# https://downloadmirror.intel.com/30381/a08/igfx_win10_100.9466.zip
1111
ocl_gpu_rt_ver_win=27.20.100.9466
@@ -30,7 +30,7 @@ ocloc_ver_win=27.20.100.9168
3030
[DRIVER VERSIONS]
3131
cpu_driver_lin=2021.12.6.0.19
3232
cpu_driver_win=2021.12.6.0.19
33-
gpu_driver_lin=21.23.20043
33+
gpu_driver_lin=21.24.20098
3434
gpu_driver_win=27.20.100.9466
3535
fpga_driver_lin=2021.12.6.0.19
3636
fpga_driver_win=2021.12.6.0.19

‎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/include/clang/Driver/Types.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -114,4 +114,5 @@ TYPE("fpga_aocr_emu", FPGA_AOCR_EMU, INVALID, "aocr", phases
114114
TYPE("fpga_aoco", FPGA_AOCO, INVALID, "aoco", phases::Compile, phases::Backend, phases::Assemble, phases::Link)
115115
TYPE("fpga_dep", FPGA_Dependencies, INVALID, "d", phases::Compile, phases::Backend, phases::Assemble, phases::Link)
116116
TYPE("fpga_dep_list", FPGA_Dependencies_List, INVALID, "txt", phases::Compile, phases::Backend, phases::Assemble, phases::Link)
117+
TYPE("host_dep_image", Host_Dependencies_Image, INVALID, "out", phases::Compile, phases::Backend, phases::Assemble, phases::Link)
117118
TYPE("none", Nothing, INVALID, nullptr, phases::Compile, phases::Backend, phases::Assemble, phases::Link)

‎clang/lib/Driver/Driver.cpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -5015,7 +5015,8 @@ class OffloadingActionBuilder final {
50155015
/// they can add it to the device linker inputs.
50165016
void addDeviceLinkDependenciesFromHost(ActionList &LinkerInputs) {
50175017
// Link image for reading dependencies from it.
5018-
auto *LA = C.MakeAction<LinkJobAction>(LinkerInputs, types::TY_Image);
5018+
auto *LA = C.MakeAction<LinkJobAction>(LinkerInputs,
5019+
types::TY_Host_Dependencies_Image);
50195020

50205021
// Calculate all the offload kinds used in the current compilation.
50215022
unsigned ActiveOffloadKinds = 0u;
@@ -5612,7 +5613,7 @@ Action *Driver::ConstructPhaseAction(
56125613
}
56135614
types::ID HostPPType = types::getPreprocessedType(Input->getType());
56145615
if (Args.hasArg(options::OPT_fsycl) && HostPPType != types::TY_INVALID &&
5615-
Args.hasArg(options::OPT_fsycl_use_footer) &&
5616+
!Args.hasArg(options::OPT_fno_sycl_use_footer) &&
56165617
TargetDeviceOffloadKind == Action::OFK_None) {
56175618
// Performing a host compilation with -fsycl. Append the integration
56185619
// footer to the preprocessed source file. We then add another

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

Lines changed: 15 additions & 20 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()) {
@@ -8608,7 +8608,11 @@ void SPIRVTranslator::ConstructJob(Compilation &C, const JobAction &JA,
86088608
TranslatorArgs.push_back(Output.getFilename());
86098609
if (getToolChain().getTriple().isSYCLDeviceEnvironment()) {
86108610
TranslatorArgs.push_back("-spirv-max-version=1.3");
8611-
TranslatorArgs.push_back("-spirv-debug-info-version=ocl-100");
8611+
// TODO: align debug info for FPGA H/W when its SPIR-V consumer is ready
8612+
if (C.getDriver().isFPGAEmulationMode())
8613+
TranslatorArgs.push_back("-spirv-debug-info-version=ocl-100");
8614+
else
8615+
TranslatorArgs.push_back("-spirv-debug-info-version=legacy");
86128616
// Prevent crash in the translator if input IR contains DIExpression
86138617
// operations which don't have mapping to OpenCL.DebugInfo.100 spec.
86148618
TranslatorArgs.push_back("-spirv-allow-extra-diexpressions");
@@ -8639,23 +8643,14 @@ void SPIRVTranslator::ConstructJob(Compilation &C, const JobAction &JA,
86398643
",+SPV_INTEL_long_constant_composite"
86408644
",+SPV_INTEL_fpga_invocation_pipelining_attributes";
86418645
ExtArg = ExtArg + DefaultExtArg + INTELExtArg;
8642-
if (getToolChain().getTriple().getSubArch() ==
8643-
llvm::Triple::SPIRSubArch_fpga) {
8644-
for (auto *A : TCArgs) {
8645-
if (A->getOption().matches(options::OPT_Xs_separate) ||
8646-
A->getOption().matches(options::OPT_Xs)) {
8647-
StringRef ArgString(A->getValue());
8648-
// Enable SPV_INTEL_usm_storage_classes only for FPGA hardware,
8649-
// since it adds new storage classes that represent global_device and
8650-
// global_host address spaces, which are not supported for all
8651-
// targets. With the extension disabled the storage classes will be
8652-
// lowered to CrossWorkgroup storage class that is mapped to just
8653-
// global address space.
8654-
if (ArgString == "hardware" || ArgString == "simulation")
8655-
ExtArg += ",+SPV_INTEL_usm_storage_classes";
8656-
}
8657-
}
8658-
}
8646+
if (!C.getDriver().isFPGAEmulationMode())
8647+
// Enable SPV_INTEL_usm_storage_classes only for FPGA hardware,
8648+
// since it adds new storage classes that represent global_device and
8649+
// global_host address spaces, which are not supported for all
8650+
// targets. With the extension disabled the storage classes will be
8651+
// lowered to CrossWorkgroup storage class that is mapped to just
8652+
// global address space.
8653+
ExtArg += ",+SPV_INTEL_usm_storage_classes";
86598654
TranslatorArgs.push_back(TCArgs.MakeArgString(ExtArg));
86608655
}
86618656
for (auto I : Inputs) {

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

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -571,6 +571,14 @@ void visualstudio::Linker::ConstructJob(Compilation &C, const JobAction &JA,
571571

572572
Args.AddAllArgValues(CmdArgs, options::OPT__SLASH_link);
573573

574+
// A user can add the -out: option to the /link sequence on the command line
575+
// which we do not want to use when we are performing the host link when
576+
// gathering dependencies used for device compilation. Add an additional
577+
// -out: to override in case it was seen.
578+
if (JA.getType() == types::TY_Host_Dependencies_Image && Output.isFilename())
579+
CmdArgs.push_back(
580+
Args.MakeArgString(std::string("-out:") + Output.getFilename()));
581+
574582
// Control Flow Guard checks
575583
if (Arg *A = Args.getLastArg(options::OPT__SLASH_guard)) {
576584
StringRef GuardArgs = A->getValue();

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

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -66,7 +66,10 @@ const char *SYCL::Linker::constructLLVMSpirvCommand(
6666
} else {
6767
CmdArgs.push_back("-spirv-max-version=1.3");
6868
CmdArgs.push_back("-spirv-ext=+all");
69-
CmdArgs.push_back("-spirv-debug-info-version=ocl-100");
69+
if (!C.getDriver().isFPGAEmulationMode())
70+
CmdArgs.push_back("-spirv-debug-info-version=legacy");
71+
else
72+
CmdArgs.push_back("-spirv-debug-info-version=ocl-100");
7073
CmdArgs.push_back("-spirv-allow-extra-diexpressions");
7174
CmdArgs.push_back("-spirv-allow-unknown-intrinsics=llvm.genx.");
7275
CmdArgs.push_back("-o");

‎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

0 commit comments

Comments
 (0)