diff --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h
index eceb3c91543bd..e05e763d11a00 100644
--- a/clang/include/clang/Basic/LangOptions.h
+++ b/clang/include/clang/Basic/LangOptions.h
@@ -354,6 +354,12 @@ class LangOptions : public LangOptionsBase {
   /// SYCL integration footer to be generated by the device compiler
   std::string SYCLIntFooter;
 
+  /// A driver-provided unique string for this translation unit that is used to
+  /// generate unique names for SYCL names.  This is provided by the driver so
+  /// that the case of multiple-offload can have each device compilation share a
+  /// name.
+  std::string SYCLUniquePrefix;
+
   LangOptions();
 
   // Define accessors/mutators for language options of enumeration type.
diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td
index dfea7fd128f8f..9b96d6b6564cb 100644
--- a/clang/include/clang/Driver/Options.td
+++ b/clang/include/clang/Driver/Options.td
@@ -5698,6 +5698,11 @@ def fsycl_int_footer : Separate<["-"], "fsycl-int-footer">,
   MarshallingInfoString<LangOpts<"SYCLIntFooter">>;
 def fsycl_int_footer_EQ : Joined<["-"], "fsycl-int-footer=">,
   Alias<fsycl_int_footer>;
+def fsycl_unique_prefix_EQ
+    : Joined<["-"], "fsycl-unique-prefix=">,
+      HelpText<"A unique prefix for this translation unit across devices, used "
+               "to generate a unique name for local variables.">,
+      MarshallingInfoString<LangOpts<"SYCLUniquePrefix">>;
 def fsycl_std_layout_kernel_params: Flag<["-"], "fsycl-std-layout-kernel-params">,
   HelpText<"Enable standard layout requirement for SYCL kernel parameters.">,
   MarshallingInfoFlag<LangOpts<"SYCLStdLayoutKernelParams">>;
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 070a1174d0e98..8303aaddd0942 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -434,10 +434,13 @@ class SYCLIntegrationFooter {
 public:
   SYCLIntegrationFooter(Sema &S) : S(S) {}
   bool emit(StringRef MainSrc);
+  void addVarDecl(const VarDecl *VD);
 
 private:
   bool emit(raw_ostream &O);
   Sema &S;
+  llvm::SmallVector<const VarDecl *> SpecConstants;
+  void emitSpecIDName(raw_ostream &O, const VarDecl *VD);
 };
 
 /// Tracks expected type during expression parsing, for use in code completion.
@@ -13162,6 +13165,11 @@ class Sema final {
     return *SyclIntFooter.get();
   }
 
+  void addSyclVarDecl(VarDecl *VD) {
+    if (LangOpts.SYCLIsDevice && !LangOpts.SYCLIntFooter.empty())
+      getSyclIntegrationFooter().addVarDecl(VD);
+  }
+
   enum SYCLRestrictKind {
     KernelGlobalVariable,
     KernelRTTI,
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 71ceb8aefefc5..a6d494d644ec8 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -7532,6 +7532,7 @@ NamedDecl *Sema::ActOnVariableDeclarator(
   if (IsMemberSpecialization && !NewVD->isInvalidDecl())
     CompleteMemberSpecialization(NewVD, Previous);
 
+  addSyclVarDecl(NewVD);
   return NewVD;
 }
 
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 774045145f699..b3afa12cd9368 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -122,6 +122,10 @@ class Util {
   /// specialization constant class.
   static bool isSyclSpecConstantType(QualType Ty);
 
+  /// Checks whether given clang type is a full specialization of the SYCL
+  /// specialization id class.
+  static bool isSyclSpecIdType(QualType Ty);
+
   /// Checks whether given clang type is a full specialization of the SYCL
   /// kernel_handler class.
   static bool isSyclKernelHandlerType(QualType Ty);
@@ -4291,6 +4295,22 @@ SYCLIntegrationHeader::SYCLIntegrationHeader(bool _UnnamedLambdaSupport,
                                              Sema &_S)
     : UnnamedLambdaSupport(_UnnamedLambdaSupport), S(_S) {}
 
+void SYCLIntegrationFooter::addVarDecl(const VarDecl *VD) {
+  // Step 1: ensure that this is of the correct type-spec-constant template
+  // specialization).
+  if (!Util::isSyclSpecIdType(VD->getType()))
+    return;
+  // Step 2: ensure that this is a static member, or a namespace-scope.
+  // Note that isLocalVarDeclorParm excludes thread-local and static-local
+  // intentionally, as there is no way to 'spell' one of those in the
+  // specialization. We just don't generate the specialization for those, and
+  // let an error happen during host compilation.
+  if (!VD->hasGlobalStorage() || VD->isLocalVarDeclOrParm())
+    return;
+  // Step 3: Add to SpecConstants collection.
+  SpecConstants.push_back(VD);
+}
+
 // Post-compile integration header support.
 bool SYCLIntegrationFooter::emit(StringRef IntHeaderName) {
   if (IntHeaderName.empty())
@@ -4307,8 +4327,40 @@ bool SYCLIntegrationFooter::emit(StringRef IntHeaderName) {
   return emit(Out);
 }
 
+void SYCLIntegrationFooter::emitSpecIDName(raw_ostream &O, const VarDecl *VD) {
+  // FIXME: Figure out the spec-constant unique name here.
+  // Note that this changes based on the linkage of the variable.
+  // We typically want to use the __builtin_unique_stable_name for the variable
+  // (or the newer-equivilent for values, see the JIRA), but we also have to
+  // figure out if this has internal or external linkage.  In external-case this
+  // should be the same as the the unique-name.  However, this isn't the case
+  // with local-linkage, where we want to put the driver-provided random-value
+  // ahead of it, so that we make sure it is unique across translation units.
+  // This name should come from the yet implemented__builtin_unique_stable_name
+  // feature that accepts variables and gives the mangling for that.
+  O << "";
+}
+
 bool SYCLIntegrationFooter::emit(raw_ostream &O) {
-  O << "// Integration Footer contents to go here.\n";
+  PrintingPolicy Policy{S.getLangOpts()};
+  Policy.adjustForCPlusPlusFwdDecl();
+  Policy.SuppressTypedefs = true;
+  Policy.SuppressUnwrittenScope = true;
+
+  for (const VarDecl *D : SpecConstants) {
+    O << "template<>\n";
+    O << "inline const char *get_spec_constant_symbolic_ID<";
+    // Emit the FQN for this, but we probably need to do some funny-business for
+    // anonymous namespaces.
+    D->printQualifiedName(O, Policy);
+    O << ">() {\n";
+    O << "  return \"";
+    emitSpecIDName(O, D);
+    O << "\";\n";
+    O << "}\n";
+  }
+
+  O << "#include <CL/sycl/detail/spec_const_integration.hpp>\n";
   return true;
 }
 
@@ -4345,6 +4397,15 @@ bool Util::isSyclSpecConstantType(QualType Ty) {
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
+bool Util::isSyclSpecIdType(QualType Ty) {
+  std::array<DeclContextDesc, 3> Scopes = {
+      Util::MakeDeclContextDesc(clang::Decl::Kind::Namespace, "cl"),
+      Util::MakeDeclContextDesc(clang::Decl::Kind::Namespace, "sycl"),
+      Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization,
+                                "specialization_id")};
+  return matchQualifiedTypeName(Ty, Scopes);
+}
+
 bool Util::isSyclKernelHandlerType(QualType Ty) {
   std::array<DeclContextDesc, 3> Scopes = {
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 1600c9c5e1f0e..5effaaf5a3d58 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -1341,6 +1341,7 @@ Decl *TemplateDeclInstantiator::VisitVarDecl(VarDecl *D,
   if (Var->isStaticLocal())
     SemaRef.CheckStaticLocalForDllExport(Var);
 
+  SemaRef.addSyclVarDecl(Var);
   return Var;
 }
 
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index fc3ca2c146ad6..f041bd230a9ff 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -295,6 +295,23 @@ class kernel_handler {
   void __init_specialization_constants_buffer(char *specialization_constants_buffer) {}
 };
 
+template <typename T> class specialization_id {
+public:
+  using value_type = T;
+
+  template <class... Args>
+  explicit constexpr specialization_id(Args &&...args)
+      : MDefaultValue(args...) {}
+
+  specialization_id(const specialization_id &rhs) = delete;
+  specialization_id(specialization_id &&rhs) = delete;
+  specialization_id &operator=(const specialization_id &rhs) = delete;
+  specialization_id &operator=(specialization_id &&rhs) = delete;
+
+private:
+  T MDefaultValue;
+};
+
 #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
 template <typename KernelName = auto_name, typename KernelType>
 ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { // #KernelSingleTask
diff --git a/clang/test/CodeGenSYCL/integration_footer.cpp b/clang/test/CodeGenSYCL/integration_footer.cpp
index 9d95ae0f20274..41b37f7b5fd68 100644
--- a/clang/test/CodeGenSYCL/integration_footer.cpp
+++ b/clang/test/CodeGenSYCL/integration_footer.cpp
@@ -1,10 +1,92 @@
 // RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-footer=%t.h %s -emit-llvm -o %t.ll
 // RUN: FileCheck -input-file=%t.h %s
 
-// CHECK: // Integration Footer contents to go here.
-
 #include "Inputs/sycl.hpp"
 
 int main() {
   cl::sycl::kernel_single_task<class first_kernel>([]() {});
 }
+
+using namespace cl::sycl;
+
+cl::sycl::specialization_id<int> GlobalSpecID;
+// CHECK: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<GlobalSpecID>() {
+// CHECK-NEXT: return "";
+// CHECK-NEXT: }
+
+struct Wrapper {
+  static specialization_id<int> WrapperSpecID;
+  // CHECK: template<>
+  // CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<Wrapper::WrapperSpecID>() {
+  // CHECK-NEXT: return "";
+  // CHECK-NEXT: }
+};
+
+template <typename T>
+struct WrapperTemplate {
+  static specialization_id<T> WrapperSpecID;
+};
+template class WrapperTemplate<int>;
+// CHECK: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<WrapperTemplate<int>::WrapperSpecID>() {
+// CHECK-NEXT: return "";
+// CHECK-NEXT: }
+template class WrapperTemplate<double>;
+// CHECK: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<WrapperTemplate<double>::WrapperSpecID>() {
+// CHECK-NEXT: return "";
+// CHECK-NEXT: }
+
+namespace Foo {
+specialization_id<int> NSSpecID;
+// CHECK: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<Foo::NSSpecID>() {
+// CHECK-NEXT: return "";
+// CHECK-NEXT: }
+inline namespace Bar {
+specialization_id<int> InlineNSSpecID;
+// CHECK: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<Foo::InlineNSSpecID>() {
+// CHECK-NEXT: return "";
+// CHECK-NEXT: }
+specialization_id<int> NSSpecID;
+// CHECK: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<Foo::Bar::NSSpecID>() {
+// CHECK-NEXT: return "";
+// CHECK-NEXT: }
+
+struct Wrapper {
+  static specialization_id<int> WrapperSpecID;
+  // CHECK: template<>
+  // CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<Foo::Wrapper::WrapperSpecID>() {
+  // CHECK-NEXT: return "";
+  // CHECK-NEXT: }
+};
+
+template <typename T>
+struct WrapperTemplate {
+  static specialization_id<T> WrapperSpecID;
+};
+template class WrapperTemplate<int>;
+// CHECK: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<Foo::WrapperTemplate<int>::WrapperSpecID>() {
+// CHECK-NEXT: return "";
+// CHECK-NEXT: }
+template class WrapperTemplate<double>;
+// CHECK: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<Foo::WrapperTemplate<double>::WrapperSpecID>() {
+// CHECK-NEXT: return "";
+// CHECK-NEXT: }
+} // namespace Bar
+namespace {
+specialization_id<int> AnonNSSpecID;
+// CHECK: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<Foo::AnonNSSpecID>() {
+// CHECK-NEXT: return "";
+// CHECK-NEXT: }
+} // namespace
+
+} // namespace Foo
+
+// CHECK: #include <CL/sycl/detail/spec_const_integration.hpp>