[clang] Modify linkage and register initialization of device_global

clang/lib/CodeGen/CodeGenModule.cpp

@@ -907,6 +907,7 @@ void CodeGenModule::Release() {
   DeferredDecls.insert(EmittedDeferredDecls.begin(),
                        EmittedDeferredDecls.end());
   EmittedDeferredDecls.clear();
+  RenameSYCLStaticDeviceGlobalVariablesDeferred();
   EmitVTablesOpportunistically();
   applyGlobalValReplacements();
   applyReplacements();
@@ -2816,19 +2817,6 @@ void CodeGenModule::SetCommonAttributes(GlobalDecl GD, llvm::GlobalValue *GV) {
        (CodeGenOpts.KeepStaticConsts && VD->getStorageDuration() == SD_Static &&
         VD->getType().isConstQualified())))
     addUsedOrCompilerUsedGlobal(GV);
-
-  if (getLangOpts().SYCLIsDevice) {
-    // Add internal device_global variables to llvm.compiler.used array to
-    // prevent early optimizations from removing these variables from the
-    // module.
-    if (D && isa<VarDecl>(D)) {
-      const auto *VD = cast<VarDecl>(D);
-      const RecordDecl *RD = VD->getType()->getAsRecordDecl();
-      if (RD && RD->hasAttr<SYCLDeviceGlobalAttr>() &&
-          VD->getFormalLinkage() == Linkage::Internal)
-        addUsedOrCompilerUsedGlobal(GV);
-    }
-  }
 }
 
 bool CodeGenModule::GetCPUAndFeaturesAttributes(GlobalDecl GD,
@@ -3566,6 +3554,11 @@ void CodeGenModule::EmitVTablesOpportunistically() {
   OpportunisticVTables.clear();
 }
 
+void CodeGenModule::RenameSYCLStaticDeviceGlobalVariablesDeferred() {
+  for (auto &Entry : SYCLStaticDeviceGlobalsToRename)
+    Entry.second->setName(Entry.first());
+}
+
 void CodeGenModule::EmitGlobalAnnotations() {
   for (const auto& [MangledName, VD] : DeferredAnnotations) {
     llvm::GlobalValue *GV = GetGlobalValue(MangledName);
@@ -6186,6 +6179,9 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
   if (getLangOpts().SYCLIsDevice)
     addGlobalIntelFPGAAnnotation(D, GV);
 
+  // Set the llvm linkage type as appropriate.
+  llvm::GlobalValue::LinkageTypes Linkage = getLLVMLinkageVarDefinition(D);
+
   if (getLangOpts().SYCLIsDevice) {
     const RecordDecl *RD = D->getType()->getAsRecordDecl();
 
@@ -6196,8 +6192,30 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
         AddGlobalSYCLIRAttributes(GV, RD);
       // If VarDecl has a type decorated with SYCL device_global attribute 
       // emit IR attribute 'sycl-unique-id'.
-      if (RD->hasAttr<SYCLDeviceGlobalAttr>())
+      if (RD->hasAttr<SYCLDeviceGlobalAttr>()) {
         addSYCLUniqueID(GV, D, Context);
+        if (Linkage == llvm::GlobalValue::InternalLinkage) {
+          // Despite being `static`, static device globals need to be linked
+          // externally as symbols must persist across the host-device boundary
+          // and not removed or tampered with in later optimizations. Since we
+          // are linking these symbols externally we need a way to 'hide' the
+          // static device global symbols from other TUs that may have the same
+          // static symbols. We can 'hide' the symbols from other TUs despite
+          // being linked externally by prefixing the symbol with a UID string.
+          //
+          // Also prefix the symbol with __ so the symbol is not in user space.
+          //
+          // Global vars will be renamed at the end of module codegen in
+          // RenameSYCLStaticDeviceGlobalVariablesDeferred()
+          auto builtinString =
+              "__" + SYCLUniqueStableIdExpr::ComputeName(Context, D);
+          SYCLStaticDeviceGlobalsToRename[builtinString] = GV;
+          Linkage = llvm::GlobalValue::ExternalLinkage;
+        }
+        // SYCL device globals are initialized externally
+        GV->setExternallyInitialized(true);
+      }
+
       // If VarDecl type is SYCLTypeAttr::host_pipe, emit the IR attribute 
       // 'sycl-unique-id'.
       if (const auto *Attr = RD->getAttr<SYCLTypeAttr>())
@@ -6217,9 +6235,6 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
     GlobalsRestrict->addOperand(Node);
   }
 
-  // Set the llvm linkage type as appropriate.
-  llvm::GlobalValue::LinkageTypes Linkage = getLLVMLinkageVarDefinition(D);
-
   // CUDA B.2.1 "The __device__ qualifier declares a variable that resides on
   // the device. [...]"
   // CUDA B.2.2 "The __constant__ qualifier, optionally used together with

clang/lib/CodeGen/CodeGenModule.h

@@ -416,6 +416,10 @@ class CodeGenModule : public CodeGenTypeCache {
   /// A queue of (optional) vtables to consider emitting.
   std::vector<const CXXRecordDecl*> DeferredVTables;
 
+  /// A queue of static SYCL device global vars to rename once module has been
+  /// parsed
+  llvm::StringMap<llvm::GlobalVariable *> SYCLStaticDeviceGlobalsToRename;
+
   /// A queue of (optional) vtables that may be emitted opportunistically.
   std::vector<const CXXRecordDecl *> OpportunisticVTables;
 
@@ -1809,6 +1813,10 @@ class CodeGenModule : public CodeGenTypeCache {
   /// lazily.
   void EmitVTablesOpportunistically();
 
+  // SYCL static device global vars need to be renamed at the end of module
+  // codegen
+  void RenameSYCLStaticDeviceGlobalVariablesDeferred();
+
   /// Call replaceAllUsesWith on all pairs in Replacements.
   void applyReplacements();