[CUDA][HIP] Support accessing static device variable in host code for -fgpu-rdc

[yaxunl]

Summary

This is separated from https://reviews.llvm.org/D80858

For -fgpu-rdc mode, static device vars in different TU's may have the same name.
To support accessing file-scope static device variables in host code, we need to give them
a distinct name and external linkage. This can be done by postfixing each static device variable with
a distinct CUID (Compilation Unit ID) hash.

Since the static device variables have different name across compilation units, now we let
them have external linkage so that they can be looked up by the runtime.

[JonChesterfield]

I concede that making the variables external, and trying to give them unique names, does work around static variables not working. I believe static variables are subjected to more aggressive optimisation than external ones but the effect might not be significant.

This "works" in cuda today because the loader ignores the local annotation when accessing the variable. There is some probably unintended behaviour when multiple static variables have the same name in that the first one wins.

The corresponding change to the hsa loader is trivial. Why is making the symbols external, with the associated complexity in picking non-conflicting names, considered better than changing the loader?

[yaxunl]

In D85223#2452363, @JonChesterfield wrote:

I concede that making the variables external, and trying to give them unique names, does work around static variables not working. I believe static variables are subjected to more aggressive optimisation than external ones but the effect might not be significant.

This "works" in cuda today because the loader ignores the local annotation when accessing the variable. There is some probably unintended behaviour when multiple static variables have the same name in that the first one wins.

The corresponding change to the hsa loader is trivial. Why is making the symbols external, with the associated complexity in picking non-conflicting names, considered better than changing the loader?

Three reasons:

1. The loader would like to look up dynsym only, which conforms better to the standard dynamic linker behavior and is more efficient than looking up all symbols.

2. There could be symbols with the same name from different compilation units and they end up as local symbols with the same name in the binary. How does the loader know which is which.

3. If a device symbol is static but actually accessed by the host code in the same compilation unit, the device symbol has de facto external linkage since it is truly accessed by some one out side of the device object (this is due to the unfortunate fact that a single source file ends up with a host object and a device object even though they are supposed to be the same compilation unit). Keeping the device symbol with internal linkage will cause the compiler over optimize the device code.

[tra]

I'd propose splitting the patch into two. One with the addition of CUID and the other that changes the way we havdle static vars.
CUID is useful on its own and is relatively uncontroversial.

Externalizing static vars is a more interesting issue and I'm not sure what's the best way to handle it yet. On one hand it is necessary for visibility across host/device, on the other, externalizing all static vars will almost always have negative effect as very few of the static vars actually need this. As already pointed out in the #if 0 section of the patch, ideally we should externalize only the vars that need it. Generally speaking, I do not think we will be able to do that, because with -fgpu-rdc it may be used from the host code in some other TU.

We may need to explicitly annotate such the static variables that need to be visible on both sides and only apply externalization to the variables annotated this way. E.g. require them to be __host__ __device__.

WDYT?

[yaxunl]

In D85223#2507518, @tra wrote:

I'd propose splitting the patch into two. One with the addition of CUID and the other that changes the way we havdle static vars.
CUID is useful on its own and is relatively uncontroversial.

Externalizing static vars is a more interesting issue and I'm not sure what's the best way to handle it yet. On one hand it is necessary for visibility across host/device, on the other, externalizing all static vars will almost always have negative effect as very few of the static vars actually need this. As already pointed out in the #if 0 section of the patch, ideally we should externalize only the vars that need it. Generally speaking, I do not think we will be able to do that, because with -fgpu-rdc it may be used from the host code in some other TU.

We may need to explicitly annotate such the static variables that need to be visible on both sides and only apply externalization to the variables annotated this way. E.g. require them to be __host__ __device__.

WDYT?

Agree that CUID may be useful for other situations. Will separate it to another review.

[yaxunl]

separate CUID patch.

[tra]

clang/lib/AST/ASTContext.cpp

11446–11447

!(getLangOpts().GPURelocatableDeviceCode && getLangOpts().CUID.empty()).

Maybe this should be broken down into something easier to read.

// Applies only to -fgpu-rdc or when we were given a CUID
if (!getLangOpts().GPURelocatableDeviceCode || !getLangOpts().CUID.empty()))
    return false;
// .. only file-scope static vars...
auto *VD = dyn_cast<VarDecl>(D);
if (!(VD && VD->isFileVarDecl() && VD->getStorageClass() == SC_Static))
    return false;
// .. with explicit __device__ or __constant__ attributes.
return ((D->hasAttr<CUDADeviceAttr>() && !D->getAttr<CUDADeviceAttr>()->isImplicit()) ||
            (D->hasAttr<CUDAConstantAttr>() &&!D->getAttr<CUDAConstantAttr>()->isImplicit()));

11446–11447

BTW, does this mean that we'll externalize & uniquify the vars even w/o -fgpu-rdc if CUID is given?

IMO -fgpu-rdc should remain the flag to control whether externalization is needed.
CUID controls the value of a unique suffix, if we need it, but should not automatically enable externalization.

clang/lib/CodeGen/CodeGenModule.cpp

2895–2896

Is this code needed?

[yaxunl]

clang/lib/AST/ASTContext.cpp

11446–11447

done

11446–11447

mayExternalizeStaticVar returns true does not mean the static var must be externalized. mayExternalizeStaticVar only indicates the static var may be externalized. It is used to enable checking whether this var is used by host code.

For -fno-gpu-rdc, we only externalize a static variable if it is referenced by host code. If a static var is referenced by host code, -fno-gpu-rdc will change its linkage to external, but does not need to make the symbol unique because each TU ends up as a different device binary.

clang/lib/CodeGen/CodeGenModule.cpp

2895–2896

this code is not needed. removed.

[yaxunl]

Revised by Artem's comments. Use CUID hash as postfix for static variable name.

[tra]

LGTM with new test nits.

@JonChesterfield -- are you OK with the patch?

clang/test/CodeGenCUDA/device-var-linkage.cu

40

It should probably be a regex after HASH:, not the hash value itself.

clang/test/CodeGenCUDA/managed-var.cu

42

Same here.

clang/test/CodeGenCUDA/static-device-var-rdc.cu

34

ditto.

clang/test/SemaCUDA/static-device-var.cu

10

A comment explaining what we're testing would be helpful. no-diagnostics gives no clues about what is it we're looking for here.

14–22

So, this verifies that we're allowed to use static local vars in device code. A comment would be useful.

23–37

And this verifies that global static vars can be referenced from both host and device.
I'd also add a negative test with static int host_only; and would verify that we still don't allow accessing it from the device.

[JonChesterfield]

This works around the limitations of the binary format nvptx and amdgpu are using in the compiler. It's the wrong place in the stack to fix it - we could introduce another symbol table in the binary to capture the per-tu-between-arch scoping.

However, if we later reach consensus on what to do in the elf instead, we can still do that. In particular, embedding an elf for one arch in a named section of an elf for a host arch is crude. This workaround seems acceptable in the meantime.

[tra]

What breaks existing abstractions is that we produce N ELF objects from a single TU and the meaning of static becomes fuzzy. On one hand, we don't want that static symbol to be visible across objects on the same target, at the same time we do want it to be visible across host/device objects compiled from the same TU. ELF does not have a way to express it. Making such symbols visible with an unique suffix seems to be a reasonable tradeoff. We probably have more options available for AMDGPU. E.g. as you've suggested, give runtime extra clues about referencing these variables across host/device boundary without resorting to making them externally visible. However, we don't have that flexibility for NVPTX.

[yaxunl]

In D85223#2551894, @JonChesterfield wrote:

This works around the limitations of the binary format nvptx and amdgpu are using in the compiler. It's the wrong place in the stack to fix it - we could introduce another symbol table in the binary to capture the per-tu-between-arch scoping.

However, if we later reach consensus on what to do in the elf instead, we can still do that. In particular, embedding an elf for one arch in a named section of an elf for a host arch is crude. This workaround seems acceptable in the meantime.

Yes we should revisit this if there is a better solution.

clang/test/CodeGenCUDA/device-var-linkage.cu

40

will do

clang/test/CodeGenCUDA/managed-var.cu

42

will do

clang/test/CodeGenCUDA/static-device-var-rdc.cu

34

will do

clang/test/SemaCUDA/static-device-var.cu

10

will do

14–22

will do

23–37

will do

[Hahnfeld]

clang/lib/CodeGen/CodeGenModule.cpp

6292–6295

I've tried to use this with CUDA, but it errors out because . is not allowed in identifiers. Could you check if https://reviews.llvm.org/D108456 also works for HIP?

[yaxunl]

clang/lib/CodeGen/CodeGenModule.cpp

6292–6295

I will try it with our CI and get back to you.

clang/include/clang/AST/ASTContext.h

class ASTContext : public RefCountedBase<ASTContext> {
   mutable llvm::FoldingSet<TemplateParamObjectDecl> TemplateParamObjectDecls;
 
   /// A cache mapping a string value to a StringLiteral object with the same
   /// value.
   ///
   /// This is lazily created.  This is intentionally not serialized.
   mutable llvm::StringMap<StringLiteral *> StringLiteralCache;
 
+  /// MD5 hash of CUID. It is calculated when first used and cached by this
+  /// data member.
+  mutable std::string CUIDHash;
+
   /// Representation of a "canonical" template template parameter that
   /// is used in canonical template names.
   class CanonicalTemplateTemplateParm : public llvm::FoldingSetNode {
     TemplateTemplateParmDecl *Parm;
 
   public:
     CanonicalTemplateTemplateParm(TemplateTemplateParmDecl *Parm)
         : Parm(Parm) {}
public:
   OMPTraitInfo &getNewOMPTraitInfo();
 
   /// Whether a C++ static variable may be externalized.
   bool mayExternalizeStaticVar(const Decl *D) const;
 
   /// Whether a C++ static variable should be externalized.
   bool shouldExternalizeStaticVar(const Decl *D) const;
 
+  StringRef getCUIDHash() const;
+
 private:
   /// All OMPTraitInfo objects live in this collection, one per
   /// `pragma omp [begin] declare variant` directive.
   SmallVector<std::unique_ptr<OMPTraitInfo>, 4> OMPTraitInfoVector;
 };
 
 /// Insertion operator for diagnostics.
 const StreamingDiagnostic &operator<<(const StreamingDiagnostic &DB,

clang/lib/AST/ASTContext.cpp

 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/StringExtras.h"
 #include "llvm/ADT/StringRef.h"
 #include "llvm/ADT/Triple.h"
 #include "llvm/Support/Capacity.h"
 #include "llvm/Support/Casting.h"
 #include "llvm/Support/Compiler.h"
 #include "llvm/Support/ErrorHandling.h"
+#include "llvm/Support/MD5.h"
 #include "llvm/Support/MathExtras.h"
 #include "llvm/Support/raw_ostream.h"
 #include <algorithm>
 #include <cassert>
 #include <cstddef>
 #include <cstdint>
 #include <cstdlib>
 #include <map>
static GVALinkage adjustGVALinkageForAttributes(const ASTContext &Context,
   } else if (Context.getLangOpts().CUDA && Context.getLangOpts().CUDAIsDevice) {
     // Device-side functions with __global__ attribute must always be
     // visible externally so they can be launched from host.
     if (D->hasAttr<CUDAGlobalAttr>() &&
         (L == GVA_DiscardableODR || L == GVA_Internal))
       return GVA_StrongODR;
     // Single source offloading languages like CUDA/HIP need to be able to
     // access static device variables from host code of the same compilation
-    // unit. This is done by externalizing the static variable.
+    // unit. This is done by externalizing the static variable with a shared
+    // name between the host and device compilation which is the same for the
+    // same compilation unit whereas different among different compilation
+    // units.
     if (Context.shouldExternalizeStaticVar(D))
       return GVA_StrongExternal;
   }
   return L;
 }
 
 /// Adjust the GVALinkage for a declaration based on what an external AST source
 /// knows about whether there can be other definitions of this declaration.
operator<<(const StreamingDiagnostic &DB,
   return DB << "a prior #pragma section";
 }
 
 bool ASTContext::mayExternalizeStaticVar(const Decl *D) const {
   bool IsStaticVar =
       isa<VarDecl>(D) && cast<VarDecl>(D)->getStorageClass() == SC_Static;
   bool IsExplicitDeviceVar = (D->hasAttr<CUDADeviceAttr>() &&
                               !D->getAttr<CUDADeviceAttr>()->isImplicit()) ||
                              (D->hasAttr<CUDAConstantAttr>() &&
                               !D->getAttr<CUDAConstantAttr>()->isImplicit());
   // CUDA/HIP: static managed variables need to be externalized since it is
   // a declaration in IR, therefore cannot have internal linkage.
-  // ToDo: externalize static variables for -fgpu-rdc.
   return IsStaticVar &&
-         (D->hasAttr<HIPManagedAttr>() ||
-          (!getLangOpts().GPURelocatableDeviceCode && IsExplicitDeviceVar));
+         (D->hasAttr<HIPManagedAttr>() || IsExplicitDeviceVar);
 }
 
 bool ASTContext::shouldExternalizeStaticVar(const Decl *D) const {
   return mayExternalizeStaticVar(D) &&
          (D->hasAttr<HIPManagedAttr>() ||
           CUDAStaticDeviceVarReferencedByHost.count(cast<VarDecl>(D)));
 }
+
+StringRef ASTContext::getCUIDHash() const {
+  if (!CUIDHash.empty())
+    return CUIDHash;
+  if (LangOpts.CUID.empty())
+    return StringRef();
+  CUIDHash = llvm::utohexstr(llvm::MD5Hash(LangOpts.CUID), /*LowerCase=*/true);
+  return CUIDHash;
+}

clang/lib/CodeGen/CGCUDANV.cpp

std::string CGNVCUDARuntime::getDeviceSideName(const NamedDecl *ND) {
   std::string DeviceSideName;
   if (DeviceMC->shouldMangleDeclName(ND)) {
     SmallString<256> Buffer;
     llvm::raw_svector_ostream Out(Buffer);
     DeviceMC->mangleName(GD, Out);
     DeviceSideName = std::string(Out.str());
   } else
     DeviceSideName = std::string(ND->getIdentifier()->getName());
+
+  // Make unique name for device side static file-scope variable for HIP.
+  if (CGM.getContext().shouldExternalizeStaticVar(ND) &&
+      CGM.getLangOpts().GPURelocatableDeviceCode &&
+      !CGM.getLangOpts().CUID.empty()) {
+    SmallString<256> Buffer;
+    llvm::raw_svector_ostream Out(Buffer);
+    Out << DeviceSideName;
+    CGM.printPostfixForExternalizedStaticVar(Out);
+    DeviceSideName = std::string(Out.str());
+  }
   return DeviceSideName;
 }
 
 void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
                                      FunctionArgList &Args) {
   EmittedKernels.push_back({CGF.CurFn, CGF.CurFuncDecl});
   if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(),
                          CudaFeature::CUDA_USES_NEW_LAUNCH) ||

clang/lib/CodeGen/CodeGenModule.h

CharUnits getNaturalTypeAlignment(QualType T,
                                     LValueBaseInfo *BaseInfo = nullptr,
                                     TBAAAccessInfo *TBAAInfo = nullptr,
                                     bool forPointeeType = false);
   CharUnits getNaturalPointeeTypeAlignment(QualType T,
                                            LValueBaseInfo *BaseInfo = nullptr,
                                            TBAAAccessInfo *TBAAInfo = nullptr);
   bool stopAutoInit();
 
+  /// Print the postfix for externalized static variable for single source
+  /// offloading languages CUDA and HIP.
+  void printPostfixForExternalizedStaticVar(llvm::raw_ostream &OS) const;
+
 private:
   llvm::Constant *GetOrCreateLLVMFunction(
       StringRef MangledName, llvm::Type *Ty, GlobalDecl D, bool ForVTable,
       bool DontDefer = false, bool IsThunk = false,
       llvm::AttributeList ExtraAttrs = llvm::AttributeList(),
       ForDefinition_t IsForDefinition = NotForDefinition);
 
   llvm::Constant *GetOrCreateMultiVersionResolver(GlobalDecl GD,

clang/lib/CodeGen/CodeGenModule.cpp

if (FD->isMultiVersion() && !OmitMultiVersionMangling) {
       case MultiVersionKind::Target:
         AppendTargetMangling(CGM, FD->getAttr<TargetAttr>(), Out);
         break;
       case MultiVersionKind::None:
         llvm_unreachable("None multiversion type isn't valid here");
       }
     }
 
+  // Make unique name for device side static file-scope variable for HIP.
+  if (CGM.getContext().shouldExternalizeStaticVar(ND) &&
+      CGM.getLangOpts().GPURelocatableDeviceCode &&
+      CGM.getLangOpts().CUDAIsDevice && !CGM.getLangOpts().CUID.empty())
+    CGM.printPostfixForExternalizedStaticVar(Out);
   return std::string(Out.str());
 }
 
 void CodeGenModule::UpdateMultiVersionNames(GlobalDecl GD,
                                             const FunctionDecl *FD) {
   if (!FD->isMultiVersion())
     return;
 
if (const auto *CD = dyn_cast<CXXConstructorDecl>(CanonicalGD.getDecl())) {
     if (!getTarget().getCXXABI().hasConstructorVariants()) {
       CXXCtorType OrigCtorType = GD.getCtorType();
       assert(OrigCtorType == Ctor_Base || OrigCtorType == Ctor_Complete);
       if (OrigCtorType == Ctor_Base)
         CanonicalGD = GlobalDecl(CD, Ctor_Complete);
     }
   }
 
+  // In CUDA/HIP device compilation with -fgpu-rdc, the mangled name of a
+  // static device variable depends on whether the variable is referenced by
+  // a host or device host function. Therefore the mangled name cannot be
+  // cached.
+  if (!LangOpts.CUDAIsDevice ||
+      !getContext().mayExternalizeStaticVar(GD.getDecl())) {
     auto FoundName = MangledDeclNames.find(CanonicalGD);
     if (FoundName != MangledDeclNames.end())
       return FoundName->second;
+  }
 
   // Keep the first result in the case of a mangling collision.
   const auto *ND = cast<NamedDecl>(GD.getDecl());
   std::string MangledName = getMangledNameImpl(*this, GD, ND);
 
   // Ensure either we have different ABIs between host and device compilations,
   // says host compilation following MSVC ABI but device compilation follows
   // Itanium C++ ABI or, if they follow the same ABI, kernel names after
void CodeGenModule::EmitGlobal(GlobalDecl GD) {
   // If we're deferring emission of a C++ variable with an
   // initializer, remember the order in which it appeared in the file.
   if (getLangOpts().CPlusPlus && isa<VarDecl>(Global) &&
       cast<VarDecl>(Global)->hasInit()) {
     DelayedCXXInitPosition[Global] = CXXGlobalInits.size();
     CXXGlobalInits.push_back(nullptr);
   }
 
   StringRef MangledName = getMangledName(GD);
   if (GetGlobalValue(MangledName) != nullptr) {
     // The value has already been used and should therefore be emitted.
     addDeferredDeclToEmit(GD);
   } else if (MustBeEmitted(Global)) {
     // The value must be emitted, but cannot be emitted eagerly.
     assert(!MayBeEmittedEagerly(Global));
     addDeferredDeclToEmit(GD);
   } else {
     // Otherwise, remember that we saw a deferred decl with this name.  The
if (!NumAutoVarInit) {
                       LangOptions::TrivialAutoVarInitKind::Zero
                   ? "zero"
                   : "pattern");
     }
     ++NumAutoVarInit;
   }
   return false;
 }
+
+void CodeGenModule::printPostfixForExternalizedStaticVar(
+    llvm::raw_ostream &OS) const {
+  OS << ".static." << getContext().getCUIDHash();
+}

clang/test/CodeGenCUDA/device-var-linkage.cu

 // RUN: %clang_cc1 -triple nvptx -fcuda-is-device \
 // RUN:   -emit-llvm -o - -x hip %s \
 // RUN:   | FileCheck -check-prefixes=DEV,NORDC %s
 // RUN: %clang_cc1 -triple nvptx -fcuda-is-device \
-// RUN:   -fgpu-rdc -emit-llvm -o - -x hip %s \
+// RUN:   -fgpu-rdc -cuid=abc -emit-llvm -o - -x hip %s \
 // RUN:   | FileCheck -check-prefixes=DEV,RDC %s
 // RUN: %clang_cc1 -triple nvptx \
 // RUN:   -emit-llvm -o - -x hip %s \
 // RUN:   | FileCheck -check-prefixes=HOST,NORDC-H %s
 // RUN: %clang_cc1 -triple nvptx \
-// RUN:   -fgpu-rdc -emit-llvm -o - -x hip %s \
+// RUN:   -fgpu-rdc -cuid=abc -emit-llvm -o - -x hip %s \
 // RUN:   | FileCheck -check-prefixes=HOST,RDC-H %s
 
 #include "Inputs/cuda.h"
 
 // DEV-DAG: @v1 = dso_local addrspace(1) externally_initialized global i32 0
 // NORDC-H-DAG: @v1 = internal global i32 undef
 // RDC-H-DAG: @v1 = dso_local global i32 undef
 __device__ int v1;

 // DEV-DAG: @ev2 = external addrspace(4) global i32
 // HOST-DAG: @ev2 = external global i32
 extern __constant__ int ev2;
 // DEV-DAG: @ev3 = external addrspace(1) externally_initialized global i32 addrspace(1)*
 // HOST-DAG: @ev3 = external externally_initialized global i32*
 extern __managed__ int ev3;
 
 // NORDC-DAG: @_ZL3sv1 = dso_local addrspace(1) externally_initialized global i32 0
-// RDC-DAG: @_ZL3sv1 = internal addrspace(1) global i32 0
+// RDC-DAG: @_ZL3sv1.static.[[HASH:b04fd23c98500190]] = dso_local addrspace(1) externally_initialized global i32 0
 // HOST-DAG: @_ZL3sv1 = internal global i32 undef
 static __device__ int sv1;
 // NORDC-DAG: @_ZL3sv2 = dso_local addrspace(4) externally_initialized global i32 0
-// RDC-DAG: @_ZL3sv2 = internal addrspace(4) global i32 0
+// RDC-DAG: @_ZL3sv2.static.[[HASH]] = dso_local addrspace(4) externally_initialized global i32 0
 // HOST-DAG: @_ZL3sv2 = internal global i32 undef
 static __constant__ int sv2;
-// DEV-DAG: @_ZL3sv3 = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
+// NORDC-DAG: @_ZL3sv3 = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
+// RDC-DAG: @_ZL3sv3.static.[[HASH]] = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
 // HOST-DAG: @_ZL3sv3 = internal externally_initialized global i32* null
 static __managed__ int sv3;
 
 __device__ __host__ int work(int *x);
 
 __device__ __host__ int fun1() {
   return work(&ev1) + work(&ev2) + work(&ev3) + work(&sv1) + work(&sv2) + work(&sv3);
 }

clang/test/CodeGenCUDA/managed-var.cu

 // REQUIRES: x86-registered-target, amdgpu-registered-target
 
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck \
-// RUN:   -check-prefixes=COMMON,DEV %s
+// RUN:   -check-prefixes=COMMON,DEV,NORDC-D %s
 
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
-// RUN:   -emit-llvm -fgpu-rdc -o - -x hip %s | FileCheck \
-// RUN:   -check-prefixes=COMMON,DEV %s
+// RUN:   -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s | FileCheck \
+// RUN:   -check-prefixes=COMMON,DEV,RDC-D %s
 
 // RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=COMMON,HOST,NORDC %s
 
 // RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
-// RUN:   -emit-llvm -fgpu-rdc -o - -x hip %s | FileCheck \
+// RUN:   -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=COMMON,HOST,RDC %s
 
 #include "Inputs/cuda.h"
 
 // DEV-DAG: @x = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
 // NORDC-DAG: @x.init = internal global i32 1
 // RDC-DAG: @x.init = dso_local global i32 1
 // NORDC-DAG: @x = internal externally_initialized global i32* null
 // RDC-DAG: @x = dso_local externally_initialized global i32* null
 // HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"x\00"
 
 struct vec {
   float x,y,z;
 };
 
 __managed__ int x = 1;
 __managed__ vec v[100];
 __managed__ vec v2[100] = {{1, 1, 1}};
 
 // DEV-DAG: @ex = external addrspace(1) externally_initialized global i32 addrspace(1)*
 // HOST-DAG: @ex.init = external global i32
 // HOST-DAG: @ex = external externally_initialized global i32*
 extern __managed__ int ex;
 
-// DEV-DAG: @_ZL2sx = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
+// NORDC-D-DAG: @_ZL2sx = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
+// RDC-D-DAG: @_ZL2sx.static.[[HASH:b04fd23c98500190]] = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
 // HOST-DAG: @_ZL2sx.init = internal global i32 1
 // HOST-DAG: @_ZL2sx = internal externally_initialized global i32* null
+// NORDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx\00"
+// RDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx.static.[[HASH:b04fd23c98500190]]\00"
 static __managed__ int sx = 1;
 
 // Force ex and sx mitted in device compilation.
 __global__ void foo(int *z) {
   *z = x + ex + sx;
   v[1].x = 2;
 }
 

 // HOST:  %ld.managed = load i32*, i32** @ex, align 4
 // HOST:  %0 = load i32, i32* %ld.managed, align 4
 // HOST:  ret i32 %0
 __device__ __host__ int load4() {
   return ex;
 }
 
 // HOST-DAG: __hipRegisterManagedVar({{.*}}@x {{.*}}@x.init {{.*}}@[[DEVNAMEX]]{{.*}}, i64 4, i32 4)
-// HOST-DAG: __hipRegisterManagedVar({{.*}}@_ZL2sx {{.*}}@_ZL2sx.init
+// HOST-DAG: __hipRegisterManagedVar({{.*}}@_ZL2sx {{.*}}@_ZL2sx.init {{.*}}@[[DEVNAMESX]]
 // HOST-NOT: __hipRegisterManagedVar({{.*}}@ex {{.*}}@ex.init
 // HOST-DAG: declare void @__hipRegisterManagedVar(i8**, i8*, i8*, i8*, i64, i32)

clang/test/CodeGenCUDA/static-device-var-rdc.cu

+// REQUIRES: x86-registered-target
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN:   -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
+// RUN:   -check-prefixes=DEV,INT-DEV %s
+
+// RUN: %clang_cc1 -triple x86_64-gnu-linux \
+// RUN:   -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
+// RUN:   -check-prefixes=HOST,INT-HOST %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
+// RUN:   -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
+// RUN:   -check-prefixes=DEV,EXT-DEV %s
+
+// RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \
+// RUN:   -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
+// RUN:   -check-prefixes=HOST,EXT-HOST %s
+
+#include "Inputs/cuda.h"
+
+// Test function scope static device variable, which should not be externalized.
+// DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 1
+
+
+// HOST-DAG: @_ZL1x = internal global i32 undef
+// HOST-DAG: @_ZL1y = internal global i32 undef
+
+// Test normal static device variables
+// INT-DEV-DAG: @_ZL1x = dso_local addrspace(1) externally_initialized global i32 0
+// INT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x\00"
+
+// Test externalized static device variables
+// EXT-DEV-DAG: @_ZL1x.static.[[HASH:b04fd23c98500190]] = dso_local addrspace(1) externally_initialized global i32 0
+// EXT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[HASH:b04fd23c98500190]]\00"
+
+static __device__ int x;
+
+// Test static device variables not used by host code should not be externalized
+// DEV-DAG: @_ZL2x2 = internal addrspace(1) global i32 0
+
+static __device__ int x2;
+
+// Test normal static device variables
+// INT-DEV-DAG: @_ZL1y = dso_local addrspace(4) externally_initialized global i32 0
+// INT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y\00"
+
+// Test externalized static device variables
+// EXT-DEV-DAG: @_ZL1y.static.[[HASH]] = dso_local addrspace(4) externally_initialized global i32 0
+// EXT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y.static.[[HASH]]\00"
+
+static __constant__ int y;
+
+// Test static host variable, which should not be externalized nor registered.
+// HOST-DAG: @_ZL1z = internal global i32 0
+// DEV-NOT: @_ZL1z
+static int z;
+
+// Test static device variable in inline function, which should not be
+// externalized nor registered.
+// DEV-DAG: @_ZZ6devfunPPKiE1p = linkonce_odr addrspace(4) constant i32 2, comdat
+
+inline __device__ void devfun(const int ** b) {
+  const static int p = 2;
+  b[0] = &p;
+}
+
+__global__ void kernel(int *a, const int **b) {
+  const static int w = 1;
+  a[0] = x;
+  a[1] = y;
+  b[0] = &w;
+  b[1] = &x2;
+  devfun(b);
+}
+
+int* getDeviceSymbol(int *x);
+
+void foo() {
+  getDeviceSymbol(&x);
+  getDeviceSymbol(&y);
+  z = 123;
+}
+
+// HOST: __hipRegisterVar({{.*}}@_ZL1x {{.*}}@[[DEVNAMEX]]
+// HOST: __hipRegisterVar({{.*}}@_ZL1y {{.*}}@[[DEVNAMEY]]
+// HOST-NOT: __hipRegisterVar({{.*}}@_ZL2x2
+// HOST-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w
+// HOST-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p

clang/test/SemaCUDA/static-device-var.cu

+// REQUIRES: x86-registered-target
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -triple nvptx -fcuda-is-device \
+// RUN:    -emit-llvm -o - %s -fsyntax-only -verify
+
+// RUN: %clang_cc1 -triple x86_64-gnu-linux \
+// RUN:    -emit-llvm -o - %s -fsyntax-only -verify
+
+// expected-no-diagnostics
+
+#include "Inputs/cuda.h"
+
+__device__ void f1() {
+  const static int b = 123;
+  static int a;
+}
+
+__global__ void k1() {
+  const static int b = 123;
+  static int a;
+}
+
+static __device__ int x;
+static __constant__ int y;
+
+__global__ void kernel(int *a) {
+  a[0] = x;
+  a[1] = y;
+}
+
+int* getDeviceSymbol(int *x);
+
+void foo() {
+  getDeviceSymbol(&x);
+  getDeviceSymbol(&y);
+}