Fix UB related to get_spec_constant_symbolic_ID

Added a wrapper, which is defined in integration footer only after we
define all specializations for `get_spec_constant_symbolic_ID`. By using
it we ensure that we do not try to instantiate
`get_spec_constant_symbolic_ID` until after we made all required
specializations.

Author: AlexeySachkov

clang/lib/Sema/SemaSYCL.cpp

@@ -4986,7 +4986,6 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
   Policy.SuppressUnwrittenScope = true;
 
   OS << "#include <CL/sycl/detail/defines_elementary.hpp>\n";
-  OS << "#include <CL/sycl/detail/spec_const_integration.hpp>\n";
 
   llvm::SmallSet<const VarDecl *, 8> VisitedSpecConstants;
 
@@ -5029,6 +5028,9 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
     OS << "} // namespace sycl\n";
     OS << "} // __SYCL_INLINE_NAMESPACE(cl)\n";
   }
+
+  OS << "#include <CL/sycl/detail/spec_const_integration.hpp>\n";
+
   return true;
 }
 

clang/test/CodeGenSYCL/anonymous_integration_footer.cpp

@@ -10,7 +10,6 @@ int main() {
 }
 
 // CHECK: #include <CL/sycl/detail/defines_elementary.hpp>
-// CHECK-NEXT: #include <CL/sycl/detail/spec_const_integration.hpp>
 
 using namespace cl;
 
@@ -402,3 +401,5 @@ constexpr sycl::specialization_id same_name{17};
 // CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
 }
 } // namespace outer
+
+// CHECK: #include <CL/sycl/detail/spec_const_integration.hpp>

clang/test/CodeGenSYCL/anonymous_integration_footer2.cpp

@@ -10,7 +10,6 @@ int main() {
 }
 
 // CHECK: #include <CL/sycl/detail/defines_elementary.hpp>
-// CHECK-NEXT: #include <CL/sycl/detail/spec_const_integration.hpp>
 
 using namespace cl;
 
@@ -155,3 +154,5 @@ constexpr sycl::specialization_id same_name{209};
 // CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
 }
 }
+
+// CHECK: #include <CL/sycl/detail/spec_const_integration.hpp>

clang/test/CodeGenSYCL/integration_footer.cpp

@@ -8,7 +8,6 @@ int main() {
 }
 
 // CHECK: #include <CL/sycl/detail/defines_elementary.hpp>
-// CHECK-NEXT: #include <CL/sycl/detail/spec_const_integration.hpp>
 
 using namespace cl::sycl;
 
@@ -190,3 +189,5 @@ auto x = HasVarTemplate::VarTempl<int, 2>.getDefaultValue();
 // CHECK-NEXT: } // namespace detail
 // CHECK-NEXT: } // namespace sycl
 // CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
+
+// CHECK: #include <CL/sycl/detail/spec_const_integration.hpp>

sycl/include/CL/sycl/detail/kernel_desc.hpp

@@ -50,10 +50,17 @@ template <class Name> struct SpecConstantInfo {
   static constexpr const char *getName() { return ""; }
 };
 
+// This guard is needed because the libsycl.so can be compiled with C++ <=14
+// while the code requires C++17. This code is not supposed to be used by the
+// libsycl.so so it should not be a problem.
 #if __cplusplus >= 201703L
 // Translates SYCL 2020 specialization constant type to its name.
 // Definition is in spec_const_integration.hpp
 template <auto &SpecName> const char *get_spec_constant_symbolic_ID();
+// Wrapper is needed to delay instantiation of 'get_spec_constant_symbolic_ID'
+// until after we have encountered all specializations for it generated by the
+// compiler in integration footer.
+template <auto &SpecName> const char *get_spec_constant_symbolic_ID_wrapper();
 #endif
 
 #ifndef __SYCL_UNNAMED_LAMBDA__

sycl/include/CL/sycl/detail/spec_const_integration.hpp

@@ -17,12 +17,8 @@ namespace detail {
 
 #if __cplusplus >= 201703L
 // Translates SYCL 2020 specialization constant type to its name.
-template <auto &SpecName> const char *get_spec_constant_symbolic_ID() {
-#ifdef SYCL_LANGUAGE_VERSION
-  return __builtin_sycl_unique_stable_id(SpecName);
-#else
-  return "";
-#endif
+template <auto &SpecName> const char *get_spec_constant_symbolic_ID_wrapper() {
+  return get_spec_constant_symbolic_ID<SpecName>();
 }
 #endif
 

sycl/include/CL/sycl/kernel_bundle.hpp

@@ -263,7 +263,8 @@ class kernel_bundle : public detail::kernel_bundle_plain {
   /// \returns true if any device image in the kernel_bundle uses specialization
   /// constant whose address is SpecName
   template <auto &SpecName> bool has_specialization_constant() const noexcept {
-    const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
+    const char *SpecSymName =
+        detail::get_spec_constant_symbolic_ID_wrapper<SpecName>();
     return has_specialization_constant_impl(SpecSymName);
   }
 
@@ -274,7 +275,8 @@ class kernel_bundle : public detail::kernel_bundle_plain {
             typename = detail::enable_if_t<_State == bundle_state::input>>
   void set_specialization_constant(
       typename std::remove_reference_t<decltype(SpecName)>::value_type Value) {
-    const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
+    const char *SpecSymName =
+        detail::get_spec_constant_symbolic_ID_wrapper<SpecName>();
     set_specialization_constant_impl(SpecSymName, &Value,
                                      sizeof(decltype(Value)));
   }
@@ -284,7 +286,8 @@ class kernel_bundle : public detail::kernel_bundle_plain {
   template <auto &SpecName>
   typename std::remove_reference_t<decltype(SpecName)>::value_type
   get_specialization_constant() const {
-    const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
+    const char *SpecSymName =
+        detail::get_spec_constant_symbolic_ID_wrapper<SpecName>();
     if (!is_specialization_constant_set(SpecSymName))
       return SpecName.getDefaultValue();
 

sycl/unittests/spec_constants/DefaultValues.cpp

@@ -33,7 +33,7 @@ template <> struct KernelInfo<TestKernel> {
   static constexpr bool callsAnyThisFreeFunction() { return false; }
 };
 
-template <> const char *get_spec_constant_symbolic_ID<SpecConst1>() {
+template <> const char *get_spec_constant_symbolic_ID_wrapper<SpecConst1>() {
   return "SC1";
 }
 } // namespace detail