[CUDA][HIP] Fix non-ODR-use of static device variable

[yaxunl]

Summary

Currently non-ODR-use of static device variable by host code
causes static device variable to be emitted and registered,
which should not.

This patch fixes that.

[tra]

LGTM with few nits.

clang/lib/Sema/SemaExpr.cpp

17145

This can be collapsed with the if (SemaRef.LangOpts.CUDA) above.

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

118–124

Positive checks should be -DAG and negative checks should be done in a separate RUN as we can't really guarantee the order in which registration calls are emitted.

[yaxunl]

clang/lib/Sema/SemaExpr.cpp

17145

This cannot be collapsed since there is else if.

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

118–124

will add separate run lines for negative tests

[tra]

Sam, this patch has apparently triggered some unwanted side effects. I'm still reducing the failures to something that could be used for debugging, but the rough symptoms are:

We now end up emitting the code for the host-only static member functions of instantiated class templates during device-side compilation.
Clang now complains about not allowed nontrivial static initializers for some variables in class templates.

[yaxunl]

In D102237#2767052, @tra wrote:

Sam, this patch has apparently triggered some unwanted side effects. I'm still reducing the failures to something that could be used for debugging, but the rough symptoms are:

We now end up emitting the code for the host-only static member functions of instantiated class templates during device-side compilation.
Clang now complains about not allowed nontrivial static initializers for some variables in class templates.

In D102237#2767052, @tra wrote:

Sam, this patch has apparently triggered some unwanted side effects. I'm still reducing the failures to something that could be used for debugging, but the rough symptoms are:

We now end up emitting the code for the host-only static member functions of instantiated class templates during device-side compilation.
Clang now complains about not allowed nontrivial static initializers for some variables in class templates.

We limit the variables to device or constant variables, e.g. https://godbolt.org/z/seh1n7KnY

xxx is emitted because it is implicitly constant, but yyy is not emitted.

It would be interesting to know what kind of variables are emitted.

Would you like the change reverted?

[tra]

In D102237#2767131, @yaxunl wrote:

It would be interesting to know what kind of variables are emitted.

I'm still reducing the failure. I'll send you a reproducer once I have it.

Would you like the change reverted?

We have a bit of time -- till the end of the week. Let's wait till we have a reproducer and a better idea what's going on and whether it's something easy to fix.

[tra]

Here's one example reproducer: https://godbolt.org/z/77M596W89
It's rather hairy, but should be usable for further debugging.

There are no CUDA attributes anywhere in sight, but we do end up emitting a host-only constructor for o_u which calls strlen.

[tra]

Here's a slightly simpler reproducer: https://godbolt.org/z/rW6P9e37s

[yaxunl]

In D102237#2767538, @tra wrote:

Here's a slightly simpler reproducer: https://godbolt.org/z/rW6P9e37s

I have a fix for this: https://reviews.llvm.org/D102801

It seems the issue was due to clang emits the implicit constant variable aw<ar>::c but it is invalid to be emitted on device side since it is initialized with a host function pointer. The fix is not to force emit implicit constant variables in device compilation. This basically limits our previous fix to explicit device variables.

[tra]

In D102237#2769475, @yaxunl wrote:

In D102237#2767538, @tra wrote:

Here's a slightly simpler reproducer: https://godbolt.org/z/rW6P9e37s

I have a fix for this: https://reviews.llvm.org/D102801

Here's the second regression introduced by this patch. This one triggers an assertion in clang:

class a {
public:
  a(char *);
};
void b() {
  [](char *c) {
    static a d(c);
    d;
  };
}

clang/include/clang/AST/ASTContext.h

#include "clang/Basic/RISCVVTypes.def"
 
   // Decl used to help define __builtin_va_list for some targets.
   // The decl is built when constructing 'BuiltinVaListDecl'.
   mutable Decl *VaListTagDecl = nullptr;
 
   // Implicitly-declared type 'struct _GUID'.
   mutable TagDecl *MSGuidTagDecl = nullptr;
 
-  /// Keep track of CUDA/HIP static device variables referenced by host code.
-  llvm::DenseSet<const VarDecl *> CUDAStaticDeviceVarReferencedByHost;
+  /// Keep track of CUDA/HIP device-side variables ODR-used by host code.
+  llvm::DenseSet<const VarDecl *> CUDADeviceVarODRUsedByHost;
 
   ASTContext(LangOptions &LOpts, SourceManager &SM, IdentifierTable &idents,
              SelectorTable &sels, Builtin::Context &builtins);
   ASTContext(const ASTContext &) = delete;
   ASTContext &operator=(const ASTContext &) = delete;
   ~ASTContext();
 
   /// Attach an external AST source to the AST context.

clang/lib/AST/ASTContext.cpp

bool ASTContext::mayExternalizeStaticVar(const Decl *D) const {
   // a declaration in IR, therefore cannot have internal linkage.
   return IsStaticVar &&
          (D->hasAttr<HIPManagedAttr>() || IsExplicitDeviceVar);
 }
 
 bool ASTContext::shouldExternalizeStaticVar(const Decl *D) const {
   return mayExternalizeStaticVar(D) &&
          (D->hasAttr<HIPManagedAttr>() ||
-          CUDAStaticDeviceVarReferencedByHost.count(cast<VarDecl>(D)));
+          CUDADeviceVarODRUsedByHost.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

void CGNVCUDARuntime::handleVarRegistration(const VarDecl *D,
   if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) {
     // Shadow variables and their properties must be registered with CUDA
     // runtime. Skip Extern global variables, which will be registered in
     // the TU where they are defined.
     //
     // Don't register a C++17 inline variable. The local symbol can be
     // discarded and referencing a discarded local symbol from outside the
     // comdat (__cuda_register_globals) is disallowed by the ELF spec.
-    // TODO: Reject __device__ constexpr and __device__ inline in Sema.
+    //
     // HIP managed variables need to be always recorded in device and host
     // compilations for transformation.
+    //
+    // HIP managed variables and variables in CUDADeviceVarODRUsedByHost are
+    // added to llvm.compiler-used, therefore they are safe to be registered.
     if ((!D->hasExternalStorage() && !D->isInline()) ||
+        CGM.getContext().CUDADeviceVarODRUsedByHost.contains(D) ||
         D->hasAttr<HIPManagedAttr>()) {
       registerDeviceVar(D, GV, !D->hasDefinition(),
                         D->hasAttr<CUDAConstantAttr>());
     }
   } else if (D->getType()->isCUDADeviceBuiltinSurfaceType() ||
              D->getType()->isCUDADeviceBuiltinTextureType()) {
     // Builtin surfaces and textures and their template arguments are
     // also registered with CUDA runtime.

clang/lib/CodeGen/CodeGenModule.cpp

if (!DeferredVTables.empty()) {
 
     // Emitting a vtable doesn't directly cause more vtables to
     // become deferred, although it can cause functions to be
     // emitted that then need those vtables.
     assert(DeferredVTables.empty());
   }
 
   // Emit CUDA/HIP static device variables referenced by host code only.
-  if (getLangOpts().CUDA)
-    for (auto V : getContext().CUDAStaticDeviceVarReferencedByHost)
+  if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice)
+    for (const auto *V : getContext().CUDADeviceVarODRUsedByHost)
       DeferredDeclsToEmit.push_back(V);
 
   // Stop if we're out of both deferred vtables and deferred declarations.
   if (DeferredDeclsToEmit.empty())
     return;
 
   // Grab the list of decls to emit. If EmitGlobalDefinition schedules more
   // work, it will not interfere with this.

clang/lib/Sema/SemaExpr.cpp

MarkVarDeclODRUsed(VarDecl *Var, SourceLocation Loc, Sema &SemaRef,
   if (SemaRef.LangOpts.OpenMP)
     SemaRef.tryCaptureOpenMPLambdas(Var);
   SemaRef.tryCaptureVariable(Var, Loc, Sema::TryCapture_Implicit,
     /*EllipsisLoc*/ SourceLocation(),
     /*BuildAndDiagnose*/ true,
     CaptureType, DeclRefType,
     FunctionScopeIndexToStopAt);
 
-  // Diagnose ODR-use of host global variables in device functions. Reference
-  // of device global variables in host functions is allowed through shadow
-  // variables therefore it is not diagnosed.
-  if (SemaRef.LangOpts.CUDA && SemaRef.LangOpts.CUDAIsDevice) {
+  if (SemaRef.LangOpts.CUDA) {
     auto *FD = dyn_cast_or_null<FunctionDecl>(SemaRef.CurContext);
     auto Target = SemaRef.IdentifyCUDATarget(FD);
     auto IsEmittedOnDeviceSide = [](VarDecl *Var) {
       if (Var->hasAttr<CUDADeviceAttr>() || Var->hasAttr<CUDAConstantAttr>() ||
           Var->hasAttr<CUDASharedAttr>() ||
           Var->getType()->isCUDADeviceBuiltinSurfaceType() ||
           Var->getType()->isCUDADeviceBuiltinTextureType())
         return true;
       // Function-scope static variable in device functions or kernels are
       // emitted on device side.
       if (auto *FD = dyn_cast<FunctionDecl>(Var->getDeclContext())) {
         return FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>();
       }
       return false;
     };
-    if (Var && Var->hasGlobalStorage() && !IsEmittedOnDeviceSide(Var)) {
+    if (Var && Var->hasGlobalStorage()) {
+      if (!IsEmittedOnDeviceSide(Var)) {
+        // Diagnose ODR-use of host global variables in device functions.
+        // Reference of device global variables in host functions is allowed
+        // through shadow variables therefore it is not diagnosed.
+        if (SemaRef.LangOpts.CUDAIsDevice)
           SemaRef.targetDiag(Loc, diag::err_ref_bad_target)
               << /*host*/ 2 << /*variable*/ 1 << Var << Target;
+      } else if ((Target == Sema::CFT_Host || Target == Sema::CFT_HostDevice) &&
+                 !Var->hasExternalStorage()) {
+        // Record a CUDA/HIP device side variable if it is ODR-used
+        // by host code. This is done conservatively, when the variable is
+        // referenced in any of the following contexts:
+        //   - a non-function context
+        //   - a host function
+        //   - a host device function
+        // This makes the ODR-use of the device side variable by host code to
+        // be visible in the device compilation for the compiler to be able to
+        // emit template variables instantiated by host code only and to
+        // externalize the static device side variable ODR-used by host code.
+        SemaRef.getASTContext().CUDADeviceVarODRUsedByHost.insert(Var);
+      }
     }
   }
 
   Var->markUsed(SemaRef.Context);
 }
 
 void Sema::MarkCaptureUsedInEnclosingContext(VarDecl *Capture,
                                              SourceLocation Loc,
static void DoMarkVarDeclReferenced(Sema &SemaRef, SourceLocation Loc,
   assert((!E || isa<DeclRefExpr>(E) || isa<MemberExpr>(E) ||
           isa<FunctionParmPackExpr>(E)) &&
          "Invalid Expr argument to DoMarkVarDeclReferenced");
   Var->setReferenced();
 
   if (Var->isInvalidDecl())
     return;
 
-  // Record a CUDA/HIP static device/constant variable if it is referenced
-  // by host code. This is done conservatively, when the variable is referenced
-  // in any of the following contexts:
-  //   - a non-function context
-  //   - a host function
-  //   - a host device function
-  // This also requires the reference of the static device/constant variable by
-  // host code to be visible in the device compilation for the compiler to be
-  // able to externalize the static device/constant variable.
-  if (SemaRef.getASTContext().mayExternalizeStaticVar(Var)) {
-    auto *CurContext = SemaRef.CurContext;
-    if (!CurContext || !isa<FunctionDecl>(CurContext) ||
-        cast<FunctionDecl>(CurContext)->hasAttr<CUDAHostAttr>() ||
-        (!cast<FunctionDecl>(CurContext)->hasAttr<CUDADeviceAttr>() &&
-         !cast<FunctionDecl>(CurContext)->hasAttr<CUDAGlobalAttr>()))
-      SemaRef.getASTContext().CUDAStaticDeviceVarReferencedByHost.insert(Var);
-  }
-
   auto *MSI = Var->getMemberSpecializationInfo();
   TemplateSpecializationKind TSK = MSI ? MSI->getTemplateSpecializationKind()
                                        : Var->getTemplateSpecializationKind();
 
   OdrUseContext OdrUse = isOdrUseContext(SemaRef);
   bool UsableInConstantExpr =
       Var->mightBeUsableInConstantExpressions(SemaRef.Context);
 

clang/test/CodeGenCUDA/device-stub.cu

 // NORDC-DAG: @ext_device_var_def = internal global i32 undef,
 // RDC-DAG: @ext_device_var_def = dso_local global i32 undef,
 // WIN-DAG: @"?ext_constant_var_def@@3HA" = internal global i32 undef
 __constant__ int ext_constant_var_def = 2;
 
 #if __cplusplus > 201402L
 // NORDC17: @inline_var = internal global i32 undef, comdat, align 4{{$}}
 // RDC17: @inline_var = linkonce_odr global i32 undef, comdat, align 4{{$}}
+// NORDC17-NOT: @inline_var2 =
+// RDC17-NOT: @inline_var2 =
 // NORDC17: @_ZN1C17member_inline_varE = internal constant i32 undef, comdat, align 4{{$}}
 // RDC17: @_ZN1C17member_inline_varE = linkonce_odr constant i32 undef, comdat, align 4{{$}}
+// Check inline variable ODR-used by host is emitted on host and registered.
 __device__ inline int inline_var = 3;
+// Check inline variable not ODR-used by host is not emitted on host or registered.
+__device__ inline int inline_var2 = 5;
 struct C {
   __device__ static constexpr int member_inline_var = 4;
 };
 #endif
 
 void use_pointers() {
   const int *p;
   p = &device_var;
   p = &constant_var;
   p = &shared_var;
   p = &host_var;
   p = &ext_device_var;
   p = &ext_constant_var;
   p = &ext_host_var;
 #if __cplusplus > 201402L
   p = &inline_var;
+  decltype(inline_var2) tmp;
   p = &C::member_inline_var;
 #endif
 }
 
+__device__ void device_use() {
+#if __cplusplus > 201402L
+  const int *p = &inline_var2;
+#endif
+}
+
 // Make sure that all parts of GPU code init/cleanup are there:
 // * constant unnamed string with the device-side kernel name to be passed to
 //   __hipRegisterFunction/__cudaRegisterFunction.
 // ALL: @0 = private unnamed_addr constant [18 x i8] c"_Z10kernelfunciii\00"
 // * constant unnamed string with the device-side kernel name to be passed to
 //   __hipRegisterVar/__cudaRegisterVar.
 // ALL: @1 = private unnamed_addr constant [11 x i8] c"device_var\00"
 // ALL: @2 = private unnamed_addr constant [13 x i8] c"constant_var\00"

 
 // Test that we've built a function to register kernels and global vars.
 // ALL: define internal void @__[[PREFIX]]_register_globals
 // ALL: call{{.*}}[[PREFIX]]RegisterFunction(i8** %0, {{.*}}kernelfunc{{[^,]*}}, {{[^@]*}}@0
 // ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}device_var{{[^,]*}}, {{[^@]*}}@1, {{.*}}i32 0, {{i32|i64}} 4, i32 0, i32 0
 // ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}constant_var{{[^,]*}}, {{[^@]*}}@2, {{.*}}i32 0, {{i32|i64}} 4, i32 1, i32 0
 // ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_device_var_def{{[^,]*}}, {{[^@]*}}@3, {{.*}}i32 0, {{i32|i64}} 4, i32 0, i32 0
 // ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_constant_var_def{{[^,]*}}, {{[^@]*}}@4, {{.*}}i32 0, {{i32|i64}} 4, i32 1, i32 0
-// LNX_17-NOT: [[PREFIX]]RegisterVar(i8** %0, {{.*}}inline_var
+// LNX_17-DAG: [[PREFIX]]RegisterVar(i8** %0, {{.*}}inline_var
+// LNX_17-NOT: [[PREFIX]]RegisterVar(i8** %0, {{.*}}inline_var2
 // ALL: ret void
 
 // Test that we've built a constructor.
 // LNX: define internal void @__[[PREFIX]]_module_ctor
 
 // In separate mode it calls __[[PREFIX]]RegisterFatBinary(&__[[PREFIX]]_fatbin_wrapper)
 // HIP only register fat binary once.
 // HIP: load i8**, i8*** @__hip_gpubin_handle

clang/test/CodeGenCUDA/host-used-device-var.cu

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
-// RUN:   -std=c++11 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \
-// RUN:   | FileCheck %s
+// RUN:   -std=c++17 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \
+// RUN:   | FileCheck -check-prefix=DEV %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \
+// RUN:   -std=c++17 -O3 -emit-llvm -o - | FileCheck -check-prefix=HOST %s
+
+// Negative tests.
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
+// RUN:   -std=c++17 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \
+// RUN:   | FileCheck -check-prefix=DEV-NEG %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \
+// RUN:   -std=c++17 -O3 -emit-llvm -o - | FileCheck -check-prefix=HOST-NEG %s
 
 #include "Inputs/cuda.h"
 
 // Check device variables used by neither host nor device functioins are not kept.
 
-// CHECK-NOT: @v1
+// DEV-NEG-NOT: @v1
 __device__ int v1;
 
-// CHECK-NOT: @v2
+// DEV-NEG-NOT: @v2
 __constant__ int v2;
 
-// CHECK-NOT: @_ZL2v3
+// DEV-NEG-NOT: @_ZL2v3
 static __device__ int v3;
 
 // Check device variables used by host functions are kept.
 
-// CHECK-DAG: @u1
+// DEV-DAG: @u1
 __device__ int u1;
 
-// CHECK-DAG: @u2
+// DEV-DAG: @u2
 __constant__ int u2;
 
 // Check host-used static device var is in llvm.compiler.used.
-// CHECK-DAG: @_ZL2u3
+// DEV-DAG: @_ZL2u3
 static __device__ int u3;
 
 // Check device-used static device var is emitted but is not in llvm.compiler.used.
-// CHECK-DAG: @_ZL2u4
+// DEV-DAG: @_ZL2u4
 static __device__ int u4;
 
 // Check device variables with used attribute are always kept.
-// CHECK-DAG: @u5
+// DEV-DAG: @u5
 __device__ __attribute__((used)) int u5;
 
-int fun1() {
-  return u1 + u2 + u3;
+// Test external device variable ODR-used by host code is not emitted or registered.
+// DEV-NEG-NOT: @ext_var
+extern __device__ int ext_var;
+
+// DEV-DAG: @inline_var = linkonce_odr addrspace(1) externally_initialized global i32 0
+__device__ inline int inline_var;
+
+template<typename T>
+using func_t = T (*) (T, T);
+
+template <typename T>
+__device__ T add_func (T x, T y)
+{
+  return x + y;
+}
+
+// DEV-DAG: @_Z10p_add_funcIiE = linkonce_odr addrspace(1) externally_initialized global i32 (i32, i32)* @_Z8add_funcIiET_S0_S0_
+template <typename T>
+__device__ func_t<T> p_add_func = add_func<T>;
+
+void use(func_t<int> p);
+void use(int *p);
+
+void fun1() {
+  use(&u1);
+  use(&u2);
+  use(&u3);
+  use(&ext_var);
+  use(&inline_var);
+  use(p_add_func<int>);
 }
 
 __global__ void kern1(int **x) {
   *x = &u4;
 }
+
 // Check the exact list of variables to ensure @_ZL2u4 is not among them.
-// CHECK: @llvm.compiler.used = {{[^@]*}} @_ZL2u3 {{[^@]*}} @u1 {{[^@]*}} @u2 {{[^@]*}} @u5
+// DEV: @llvm.compiler.used = {{[^@]*}} @_Z10p_add_funcIiE {{[^@]*}} @_ZL2u3 {{[^@]*}} @inline_var {{[^@]*}} @u1 {{[^@]*}} @u2 {{[^@]*}} @u5
+
+// HOST-DAG: hipRegisterVar{{.*}}@u1
+// HOST-DAG: hipRegisterVar{{.*}}@u2
+// HOST-DAG: hipRegisterVar{{.*}}@_ZL2u3
+// HOST-DAG: hipRegisterVar{{.*}}@u5
+// HOST-DAG: hipRegisterVar{{.*}}@inline_var
+// HOST-DAG: hipRegisterVar{{.*}}@_Z10p_add_funcIiE
+// HOST-NEG-NOT: hipRegisterVar{{.*}}@ext_var
+// HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZL2u4

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

 // REQUIRES: x86-registered-target
 // REQUIRES: 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=DEV %s
+// RUN:   -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV %s
 
 // RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
-// RUN:   -emit-llvm -o - -x hip %s | FileCheck \
-// RUN:   -check-prefixes=HOST %s
+// RUN:   -emit-llvm -o - -x hip %s | FileCheck -check-prefix=HOST %s
+
+// Negative tests.
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
+// RUN:   -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV-NEG %s
+
+// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
+// RUN:   -emit-llvm -o - -x hip %s | FileCheck -check-prefix=HOST-NEG %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
 // DEV-DAG: @_ZZ6kernelPiPPKiE21local_static_constant = internal addrspace(4) constant i32 42
 // DEV-DAG: @_ZZ6kernelPiPPKiE19local_static_device = internal addrspace(1) constant i32 43
 

 // DEV-DAG: @_ZL1y ={{.*}} addrspace(4) externally_initialized global i32 0
 // HOST-DAG: @_ZL1y = internal global i32 undef
 // HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y\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
+// DEV-NEG-NOT: @_ZL1z
 static int z;
 
 // Test implicit static constant variable, which should not be externalized.
 // HOST-DAG: @_ZL2z2 = internal constant i32 456
 // DEV-DAG: @_ZL2z2 = internal addrspace(4) constant i32 456
 
 static constexpr int z2 = 456;
 
 // 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
 
 // Check a static device variable referenced by host function only is externalized.
 // DEV-DAG: @_ZL1w ={{.*}} addrspace(1) externally_initialized global i32 0
 // HOST-DAG: @_ZL1w = internal global i32 undef
 // HOST-DAG: @[[DEVNAMEW:[0-9]+]] = {{.*}}c"_ZL1w\00"
 
 static __device__ int w;
 
+// Test non-ODR-use of static device var should not be emitted or registered.
+// DEV-NEG-NOT: @_ZL1u
+// HOST-NEG-NOT: @_ZL1u
+
+static __device__ int u;
+
 inline __device__ void devfun(const int ** b) {
   const static int p = 2;
   b[0] = &p;
   b[1] = &x2;
 }
 
 __global__ void kernel(int *a, const int **b) {
   const static int w = 1;
   const static __constant__ int local_static_constant = 42;
   const static __device__ int local_static_device = 43;
   a[0] = x;
   a[1] = y;
   a[2] = x2;
   a[3] = x3;
   a[4] = x4;
   a[5] = x5;
+  a[6] = sizeof(u);
   b[0] = &w;
   b[1] = &z2;
   b[2] = &local_static_constant;
   b[3] = &local_static_device;
   devfun(b);
 }
 
 __host__ __device__ void hdf(int *a) {
   a[0] = x3;
 }
 
 int* getDeviceSymbol(int *x);
 
 void foo(const int **a) {
   getDeviceSymbol(&x);
   getDeviceSymbol(&x5);
   getDeviceSymbol(&y);
   getDeviceSymbol(&w);
   z = 123;
   a[0] = &z2;
+  decltype(u) tmp;
 }
 
-// HOST: __hipRegisterVar({{.*}}@_ZL1x {{.*}}@[[DEVNAMEX]]
-// HOST: __hipRegisterVar({{.*}}@_ZL1y {{.*}}@[[DEVNAMEY]]
-// HOST: __hipRegisterVar({{.*}}@_ZL1w {{.*}}@[[DEVNAMEW]]
-// HOST-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w
-// HOST-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p
+// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1x {{.*}}@[[DEVNAMEX]]
+// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1y {{.*}}@[[DEVNAMEY]]
+// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1w {{.*}}@[[DEVNAMEW]]
+// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZL1u
+// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w
+// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p

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:   -std=c++11 -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:   -std=c++11 -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 > %t.dev
+// RUN:   -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.dev
 // RUN: cat %t.dev | FileCheck -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 > %t.host
+// RUN:   -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.host
 // RUN: cat %t.host | FileCheck -check-prefixes=HOST,EXT-HOST %s
 
 // Check host and device compilations use the same postfixes for static
 // variable names.
 
 // RUN: cat %t.dev %t.host | FileCheck -check-prefix=POSTFIX %s
 
+// Negative tests.
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN:   -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
+// RUN:   -check-prefix=DEV-NEG %s
+
+// RUN: %clang_cc1 -triple x86_64-gnu-linux \
+// RUN:   -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
+// RUN:   -check-prefix=HOST-NEG %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
+// RUN:   -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.dev
+// RUN: cat %t.dev | FileCheck -check-prefix=DEV-NEG %s
+
+// RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \
+// RUN:   -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.host
+// RUN: cat %t.host | FileCheck -check-prefix=HOST-NEG %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 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
+// DEV-NEG-NOT: @_ZL1z
 static int z;
 
+// Test non-ODR-use of static device variable is not emitted or registered.
+// DEV-NEG-NOT: @_ZL1u
+// HOST-NEG-NOT: @_ZL1u
+static __device__ int u;
+
 // 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;
+  a[2] = sizeof(u);
   b[0] = &w;
   b[1] = &x2;
   devfun(b);
 }
 
 int* getDeviceSymbol(int *x);
 
 void foo() {
   getDeviceSymbol(&x);
   getDeviceSymbol(&y);
   z = 123;
+  decltype(u) tmp;
 }
 
-// HOST: __hipRegisterVar({{.*}}@_ZL1x {{.*}}@[[DEVNAMEX]]
-// HOST: __hipRegisterVar({{.*}}@_ZL1y {{.*}}@[[DEVNAMEY]]
-// HOST-NOT: __hipRegisterVar({{.*}}@_ZL2x2
-// HOST-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w
-// HOST-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p
+// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1x {{.*}}@[[DEVNAMEX]]
+// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1y {{.*}}@[[DEVNAMEY]]
+// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZL2x2
+// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w
+// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p
+// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZL1u