[CUDA][HIP] Promote const variables to constant

[yaxunl]

Summary

Recently we added diagnosing ODR-use of host variables
in device functions, which includes ODR-use of const
host variables since they are not really emitted on
device side. This caused regressions since we used
to allow ODR-use of const host variables in device
functions.

This patch allows ODR-use of const variables in device
functions if the const variables can be statically initialized
and have an empty dtor. Such variables are marked with
implicit constant attrs and emitted on device side. This is
in line with what clang does for constexpr variables.

[yaxunl]

do not promote or check dependent variables since their ctor/dtor/initializers are not determined yet

[tra]

clang/lib/Sema/SemaCUDA.cpp

568

Nit: add an empty line to separate from the preceding function.

722

One of the issues with automatic promotion I've ran into in the past is that __constant__ is a very limited resource.

It would be nearly impossible for the end user to predict how much __constant__ space is available for their own use and if we use too much, the failure would only manifest at runtime when we fail to load the GPU code. That will be hard to troubleshoot.

We may need to add some sort of safeguard and report an error when we see too much constant data for a given GPU.
That would have to be done somewhere at the end of the compilation pipeline.

It's not likely to affect any/many users yet so it can be done later. It's probably worth a comment and a TODO here in case someone does run into this.

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

51

It would be great to have a test which verifies that we only emit const variables that we need.
It would be bad if we'd end up with *all* host-side consts in the GPU binary.

clang/test/SemaCUDA/device-use-host-var.cu

82

This diag is a bit misleading now. While it's technically true that we're not allowed to refer to a host variables from device in general, the real reason it's not allowed in *this* case is that the variable has a non-empty ctor.

I wonder if we could add a note pointing that out. Otherwise it would be rather confusing why I can access other host vars few lines above, but not here.

[yaxunl]

revised by Artem's comments

[yaxunl]

clang/lib/Sema/SemaCUDA.cpp

568

done

722

added a TODO

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

51

added test for that

clang/test/SemaCUDA/device-use-host-var.cu

82

done

[tra]

Overall looks good, though I've got one more question.

clang/test/SemaCUDA/device-use-host-var.cu

90

What happens if we have const int &ref_foo = something_we_cant_emit_on_device;?
Both if ref_foo is and isn't used from device code.

[yaxunl]

clang/test/SemaCUDA/device-use-host-var.cu

90

If we have const int &ref_foo = something_we_cant_emit_on_device, __constant__ attribute will not be added. If ref_foo is used in device code, it will be diagnosed. If ref_foo is not used in device code, no diagnostics is emitted.

[yaxunl]

fix test

[yaxunl]

ping

[tra]

LGTM. I would like to test the patch on our code first. Please wait a bit before landing the patch. I should be able to have the results tomorrow.

[tra]

I'm done with testing. The patch does not seem to break anything obvious. Tensorflow builds and works.

[sbc100]

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

68

This seems to break if the pathname where the llvm checkout lives contains the word "external"

https://logs.chromium.org/logs/emscripten-releases/buildbucket/cr-buildbucket.appspot.com/8845343598940893760/+/steps/LLVM_regression/0/stdout

--
/b/s/w/ir/cache/builder/emscripten-releases/llvm-project/clang/test/CodeGenCUDA/device-use-host-var.cu:68:13: error: NEG-NOT: excluded string found in input
// NEG-NOT: external
            ^
<stdin>:69:78: note: found here
!1 = !{!"clang version 13.0.0 (/b/s/w/ir/cache/git/chromium.googlesource.com-external-github.com-llvm-llvm--project db3e4faa4d2cadf204e67f42bccd98957496a87a)"}
                                                                             ^~~~~~~~

Unfortunate choice of directory names I know .. but probably worth fixing.

[yaxunl]

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

68

thanks. this has been fixed by https://reviews.llvm.org/D103658

[tra]

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

68

This should already be fixed by D103658

[sbc100]

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

68

Indeed! Our latest build is now passing. Sorry for the noise.

clang/include/clang/Basic/DiagnosticSemaKinds.td

def err_config_scalar_return : Error<
   "CUDA special function '%0' must have scalar return type">;
 def err_kern_call_not_global_function : Error<
   "kernel call to non-global function %0">;
 def err_global_call_not_config : Error<
   "call to global function %0 not configured">;
 def err_ref_bad_target : Error<
   "reference to %select{__device__|__global__|__host__|__host__ __device__}0 "
   "%select{function|variable}1 %2 in %select{__device__|__global__|__host__|__host__ __device__}3 function">;
+def note_cuda_const_var_unpromoted : Note<
+  "const variable cannot be emitted on device side due to dynamic initialization">;
+def note_cuda_host_var : Note<
+  "host variable declared here">;
 def err_ref_bad_target_global_initializer : Error<
   "reference to %select{__device__|__global__|__host__|__host__ __device__}0 "
   "function %1 in global initializer">;
 def err_capture_bad_target : Error<
   "capture host variable %0 by reference in device or host device lambda function">;
 def err_capture_bad_target_this_ptr : Error<
   "capture host side class data member by this pointer in device or host device lambda function">;
 def warn_kern_is_method : Extension<

clang/lib/Sema/SemaCUDA.cpp

 }
 
 /// IdentifyTarget - Determine the CUDA compilation target for this variable.
 Sema::CUDAVariableTarget Sema::IdentifyCUDATarget(const VarDecl *Var) {
   if (Var->hasAttr<HIPManagedAttr>())
     return CVT_Unified;
   if (Var->isConstexpr() && !hasExplicitAttr<CUDAConstantAttr>(Var))
     return CVT_Both;
+  if (Var->getType().isConstQualified() && Var->hasAttr<CUDAConstantAttr>() &&
+      !hasExplicitAttr<CUDAConstantAttr>(Var))
+    return CVT_Both;
   if (Var->hasAttr<CUDADeviceAttr>() || Var->hasAttr<CUDAConstantAttr>() ||
       Var->hasAttr<CUDASharedAttr>() ||
       Var->getType()->isCUDADeviceBuiltinSurfaceType() ||
       Var->getType()->isCUDADeviceBuiltinTextureType())
     return CVT_Device;
   // Function-scope static variable without explicit device or constant
   // attribute are emitted
   //  - on both sides in host device functions
if (!llvm::all_of(ClassDecl->fields(), [&](const FieldDecl *Field) {
           return isEmptyCudaDestructor(Loc, RD->getDestructor());
         return true;
       }))
     return false;
 
   return true;
 }
 
-void Sema::checkAllowedCUDAInitializer(VarDecl *VD) {
-  if (VD->isInvalidDecl() || !VD->hasInit() || !VD->hasGlobalStorage())
-    return;
+namespace {
+enum CUDAInitializerCheckKind {
+  CICK_DeviceOrConstant, // Check initializer for device/constant variable
+  CICK_Shared,           // Check initializer for shared variable
+};
+
+bool IsDependentVar(VarDecl *VD) {
+  if (VD->getType()->isDependentType())
+    return true;
+  if (const auto *Init = VD->getInit())
+    return Init->isValueDependent();
+  return false;
+}
+
+// Check whether a variable has an allowed initializer for a CUDA device side
+// variable with global storage. \p VD may be a host variable to be checked for
+// potential promotion to device side variable.
+//
+// CUDA/HIP allows only empty constructors as initializers for global
+// variables (see E.2.3.1, CUDA 7.5). The same restriction also applies to all
+// __shared__ variables whether they are local or not (they all are implicitly
+// static in CUDA). One exception is that CUDA allows constant initializers
+// for __constant__ and __device__ variables.
+bool HasAllowedCUDADeviceStaticInitializer(Sema &S, VarDecl *VD,
+                                           CUDAInitializerCheckKind CheckKind) {
+  assert(!VD->isInvalidDecl() && VD->hasGlobalStorage());
+  assert(!IsDependentVar(VD) && "do not check dependent var");
   const Expr *Init = VD->getInit();
-  if (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
-      VD->hasAttr<CUDASharedAttr>()) {
-    if (LangOpts.GPUAllowDeviceInit)
-      return;
-    bool AllowedInit = false;
-    if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
-      AllowedInit =
-          isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
-    // We'll allow constant initializers even if it's a non-empty
-    // constructor according to CUDA rules. This deviates from NVCC,
-    // but allows us to handle things like constexpr constructors.
-    if (!AllowedInit &&
-        (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>())) {
-      auto *Init = VD->getInit();
-      // isConstantInitializer cannot be called with dependent value, therefore
-      // we skip checking dependent value here. This is OK since
-      // checkAllowedCUDAInitializer is called again when the template is
-      // instantiated.
-      AllowedInit =
-          VD->getType()->isDependentType() || Init->isValueDependent() ||
-          Init->isConstantInitializer(Context,
+  auto IsEmptyInit = [&](const Expr *Init) {
+    if (!Init)
+      return true;
+    if (const auto *CE = dyn_cast<CXXConstructExpr>(Init)) {
+      return S.isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
+    }
+    return false;
+  };
+  auto IsConstantInit = [&](const Expr *Init) {
+    assert(Init);
+    return Init->isConstantInitializer(S.Context,
                                        VD->getType()->isReferenceType());
-    }
+  };
+  auto HasEmptyDtor = [&](VarDecl *VD) {
+    if (const auto *RD = VD->getType()->getAsCXXRecordDecl())
+      return S.isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
+    return true;
+  };
+  if (CheckKind == CICK_Shared)
+    return IsEmptyInit(Init) && HasEmptyDtor(VD);
+  return S.LangOpts.GPUAllowDeviceInit ||
+         ((IsEmptyInit(Init) || IsConstantInit(Init)) && HasEmptyDtor(VD));
+} // namespace
 
-    // Also make sure that destructor, if there is one, is empty.
-    if (AllowedInit)
-      if (CXXRecordDecl *RD = VD->getType()->getAsCXXRecordDecl())
-        AllowedInit =
-            isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
-
-    if (!AllowedInit) {
-      Diag(VD->getLocation(), VD->hasAttr<CUDASharedAttr>()
-                                  ? diag::err_shared_var_init
-                                  : diag::err_dynamic_var_init)
-          << Init->getSourceRange();
-      VD->setInvalidDecl();
-    }
+void Sema::checkAllowedCUDAInitializer(VarDecl *VD) {
+  // Do not check dependent variables since the ctor/dtor/initializer are not
+  // determined. Do it after instantiation.
+  if (VD->isInvalidDecl() || !VD->hasInit() || !VD->hasGlobalStorage() ||
+      IsDependentVar(VD))
+    return;
+  const Expr *Init = VD->getInit();
+  bool IsSharedVar = VD->hasAttr<CUDASharedAttr>();
+  bool IsDeviceOrConstantVar =
+      !IsSharedVar &&
+      (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>());
+  if (IsDeviceOrConstantVar || IsSharedVar) {
+    if (HasAllowedCUDADeviceStaticInitializer(
+            *this, VD, IsSharedVar ? CICK_Shared : CICK_DeviceOrConstant))
+      return;
+    Diag(VD->getLocation(),
+         IsSharedVar ? diag::err_shared_var_init : diag::err_dynamic_var_init)
   } else {
     // This is a host-side global variable.  Check that the initializer is
     // callable from the host side.
     const FunctionDecl *InitFn = nullptr;
     if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init)) {
       InitFn = CE->getConstructor();
     } else if (const CallExpr *CE = dyn_cast<CallExpr>(Init)) {
       InitFn = CE->getDirectCallee();
if (It != Previous.end()) {
     }
     return;
   }
 
   NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
   NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
 }
 
+// TODO: `__constant__` memory may be a limited resource for certain targets.
+// A safeguard may be needed at the end of compilation pipeline if
+// `__constant__` memory usage goes beyond limit.
 void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
-  if (getLangOpts().CUDAIsDevice && VD->isConstexpr() &&
+  // Do not promote dependent variables since the cotr/dtor/initializer are
+  // not determined. Do it after instantiation.
+  if (getLangOpts().CUDAIsDevice && !VD->hasAttr<CUDAConstantAttr>() &&
+      !VD->hasAttr<CUDAConstantAttr>() && !VD->hasAttr<CUDASharedAttr>() &&
       (VD->isFileVarDecl() || VD->isStaticDataMember()) &&
-      !VD->hasAttr<CUDAConstantAttr>()) {
+      !IsDependentVar(VD) &&
+      (VD->isConstexpr() || (VD->getType().isConstQualified() &&
+                             HasAllowedCUDADeviceStaticInitializer(
+                                 *this, VD, CICK_DeviceOrConstant)))) {
     VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
   }
 }
 
 Sema::SemaDiagnosticBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
                                                        unsigned DiagID) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   SemaDiagnosticBuilder::Kind DiagKind = [&] {
     if (!isa<FunctionDecl>(CurContext))

clang/lib/Sema/SemaDecl.cpp

NamedDecl *Sema::ActOnVariableDeclarator(
   case ConstexprSpecKind::Consteval:
     Diag(D.getDeclSpec().getConstexprSpecLoc(),
          diag::err_constexpr_wrong_decl_kind)
         << static_cast<int>(D.getDeclSpec().getConstexprSpecifier());
     LLVM_FALLTHROUGH;
 
   case ConstexprSpecKind::Constexpr:
     NewVD->setConstexpr(true);
-    MaybeAddCUDAConstantAttr(NewVD);
     // C++1z [dcl.spec.constexpr]p1:
     //   A static data member declared with the constexpr specifier is
     //   implicitly an inline variable.
     if (NewVD->isStaticDataMember() &&
         (getLangOpts().CPlusPlus17 ||
          Context.getTargetInfo().getCXXABI().isMicrosoft()))
       NewVD->setImplicitlyInline();
     break;
Sema::ActOnCXXForRangeIdentifier(Scope *S, SourceLocation IdentLoc,
   FinalizeDeclaration(Var);
   return ActOnDeclStmt(FinalizeDeclaratorGroup(S, DS, Var), IdentLoc,
                        AttrEnd.isValid() ? AttrEnd : IdentLoc);
 }
 
 void Sema::CheckCompleteVariableDeclaration(VarDecl *var) {
   if (var->isInvalidDecl()) return;
 
+  MaybeAddCUDAConstantAttr(var);
+
   if (getLangOpts().OpenCL) {
     // OpenCL v2.0 s6.12.5 - Every block variable declaration must have an
     // initialiser
     if (var->getTypeSourceInfo()->getType()->isBlockPointerType() &&
         !var->hasInit()) {
       Diag(var->getLocation(), diag::err_opencl_invalid_block_declaration)
           << 1 /*Init*/;
       var->setInvalidDecl();

clang/lib/Sema/SemaExpr.cpp

if (SemaRef.LangOpts.CUDA && Var && Var->hasGlobalStorage()) {
     auto VarTarget = SemaRef.IdentifyCUDATarget(Var);
     auto UserTarget = SemaRef.IdentifyCUDATarget(FD);
     if (VarTarget == Sema::CVT_Host &&
         (UserTarget == Sema::CFT_Device || UserTarget == Sema::CFT_HostDevice ||
          UserTarget == Sema::CFT_Global)) {
       // 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)
+      if (SemaRef.LangOpts.CUDAIsDevice) {
         SemaRef.targetDiag(Loc, diag::err_ref_bad_target)
             << /*host*/ 2 << /*variable*/ 1 << Var << UserTarget;
+        SemaRef.targetDiag(Var->getLocation(),
+                           Var->getType().isConstQualified()
+                               ? diag::note_cuda_const_var_unpromoted
+                               : diag::note_cuda_host_var);
+      }
     } else if (VarTarget == Sema::CVT_Device &&
                (UserTarget == Sema::CFT_Host ||
                 UserTarget == 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

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

if (OldVar->isLocalExternDecl()) {
     NewVar->setLexicalDeclContext(Owner);
   } else if (OldVar->isOutOfLine())
     NewVar->setLexicalDeclContext(OldVar->getLexicalDeclContext());
   NewVar->setTSCSpec(OldVar->getTSCSpec());
   NewVar->setInitStyle(OldVar->getInitStyle());
   NewVar->setCXXForRangeDecl(OldVar->isCXXForRangeDecl());
   NewVar->setObjCForDecl(OldVar->isObjCForDecl());
   NewVar->setConstexpr(OldVar->isConstexpr());
-  MaybeAddCUDAConstantAttr(NewVar);
   NewVar->setInitCapture(OldVar->isInitCapture());
   NewVar->setPreviousDeclInSameBlockScope(
       OldVar->isPreviousDeclInSameBlockScope());
   NewVar->setAccess(OldVar->getAccess());
 
   if (!OldVar->isStaticDataMember()) {
     if (OldVar->isUsed(false))
       NewVar->setIsUsed();

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

-// RUN: %clang_cc1 -std=c++11 -triple amdgcn-amd-amdhsa \
+// RUN: %clang_cc1 -std=c++14 -triple amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck %s
+// RUN: %clang_cc1 -std=c++14 -triple amdgcn-amd-amdhsa \
+// RUN:   -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck -check-prefix=NEG %s
 
 #include "Inputs/cuda.h"
 
 struct A {
   int x;
 };
 
+// Check the situation of B<T> has empty ctor but B<int> has non-empty ctor.
+// Make sure const B<int> variables are not promoted to constant variables.
+template<typename T>
+struct B {
+  T x;
+  B() {}
+  B(T _x) { x = _x; }
+  static const B<T> y;
+};
+
+template<>
+struct B<int> {
+  int x;
+  B() { x = 1; }
+  static const B<int> y;
+};
+
+template<typename T>
+const B<T> B<T>::y;
+
+const B<int> B<int>::y;
+
+template<typename T>
+T temp_fun(T x) {
+  return B<T>::y.x;
+}
+
+// Check template variable with empty default ctor but non-empty initializer
+// ctor is not promoted.
+template<typename T>
+const B<T> b = B<T>(-1);
+
 constexpr int constexpr_var = 1;
 constexpr A constexpr_struct{2};
 constexpr A constexpr_array[4] = {0, 0, 0, 3};
 constexpr char constexpr_str[] = "abcd";
 const int const_var = 4;
+const A const_struct{5};
+const A const_array[] = {0, 0, 0, 6};
+const char const_str[] = "xyz";
+
+// Check const variables used by host only are not emitted.
+const int var_host_only = 7;
 
 // CHECK-DAG: @_ZL13constexpr_str.const = private unnamed_addr addrspace(4) constant [5 x i8] c"abcd\00"
 // CHECK-DAG: @_ZL13constexpr_var = internal addrspace(4) constant i32 1
 // CHECK-DAG: @_ZL16constexpr_struct = internal addrspace(4) constant %struct.A { i32 2 }
 // CHECK-DAG: @_ZL15constexpr_array = internal addrspace(4) constant [4 x %struct.A] [%struct.A zeroinitializer, %struct.A zeroinitializer, %struct.A zeroinitializer, %struct.A { i32 3 }]
-// CHECK-NOT: external
+// CHECK-DAG: @_ZL9const_var = internal addrspace(4) constant i32 4
+// CHECK-DAG: @_ZL12const_struct = internal addrspace(4) constant %struct.A { i32 5 }
+// CHECK-DAG: @_ZL11const_array = internal addrspace(4) constant [4 x %struct.A] [%struct.A zeroinitializer, %struct.A zeroinitializer, %struct.A zeroinitializer, %struct.A { i32 6 }]
+// CHECK-DAG: @_ZL9const_str = internal addrspace(4) constant [4 x i8] c"xyz\00"
+
+// NEG-NOT: @_ZN1BIiE1yE
+// NEG-NOT: @_Z1bIdE
+// NEG-NOT: @_ZL13var_host_only
+// NEG-NOT: external
 
 // CHECK-LABEL: define{{.*}}@_Z7dev_funPiPPKi
 // CHECK: store i32 1
 // CHECK: store i32 2
 // CHECK: store i32 3
-// CHECK: store i32 4
 // CHECK: load i8, i8* getelementptr {{.*}} @_ZL13constexpr_str.const
+// CHECK: store i32 4
+// CHECK: store i32 5
+// CHECK: store i32 6
+// CHECK: load i8, i8* getelementptr {{.*}} @_ZL9const_str
 // CHECK: store i32* {{.*}}@_ZL13constexpr_var
 // CHECK: store i32* getelementptr {{.*}} @_ZL16constexpr_struct
 // CHECK: store i32* getelementptr {{.*}} @_ZL15constexpr_array
+// CHECK: store i32* {{.*}}@_ZL9const_var
+// CHECK: store i32* getelementptr {{.*}} @_ZL12const_struct
+// CHECK: store i32* getelementptr {{.*}} @_ZL11const_array
 __device__ void dev_fun(int *out, const int **out2) {
   *out = constexpr_var;
   *out = constexpr_struct.x;
   *out = constexpr_array[3].x;
-  *out = const_var;
   *out = constexpr_str[3];
+  *out = const_var;
+  *out = const_struct.x;
+  *out = const_array[3].x;
+  *out = const_str[3];
   *out2 = &constexpr_var;
   *out2 = &constexpr_struct.x;
   *out2 = &constexpr_array[3].x;
+  *out2 = &const_var;
+  *out2 = &const_struct.x;
+  *out2 = &const_array[3].x;
+}
+
+void fun() {
+  temp_fun(1);
+  (void) b<double>;
+  (void) var_host_only;
 }

clang/test/SemaCUDA/device-use-host-var.cu

 // RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify=dev %s
 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify=host %s
 
 // host-no-diagnostics
 
 #include "Inputs/cuda.h"
 
+int func();
+
 struct A {
   int x;
   static int host_var;
 };
 
-int A::host_var;
+int A::host_var; // dev-note {{host variable declared here}}
 
 namespace X {
-  int host_var;
+  int host_var; // dev-note {{host variable declared here}}
 }
 
-static int static_host_var;
+// struct with non-empty ctor.
+struct B1 {
+  int x;
+  B1() { x = 1; }
+};
+
+// struct with non-empty dtor.
+struct B2 {
+  int x;
+  B2() {}
+  ~B2() { x = 0; }
+};
+
+static int static_host_var; // dev-note {{host variable declared here}}
 
 __device__ int global_dev_var;
 __constant__ int global_constant_var;
 __shared__ int global_shared_var;
 
-int global_host_var;
+int global_host_var; // dev-note 8{{host variable declared here}}
 const int global_const_var = 1;
 constexpr int global_constexpr_var = 1;
 
-int global_host_array[2] = {1, 2};
+int global_host_array[2] = {1, 2}; // dev-note {{host variable declared here}}
 const int global_const_array[2] = {1, 2};
 constexpr int global_constexpr_array[2] = {1, 2};
 
-A global_host_struct_var{1};
+A global_host_struct_var{1}; // dev-note 2{{host variable declared here}}
 const A global_const_struct_var{1};
 constexpr A global_constexpr_struct_var{1};
 
+// Check const host var initialized with non-empty ctor is not allowed in
+// device function.
+const B1 b1; // dev-note {{const variable cannot be emitted on device side due to dynamic initialization}}
+
+// Check const host var having non-empty dtor is not allowed in device function.
+const B2 b2; // dev-note {{const variable cannot be emitted on device side due to dynamic initialization}}
+
+// Check const host var initialized by non-constant initializer is not allowed
+// in device function.
+const int b3 = func(); // dev-note {{const variable cannot be emitted on device side due to dynamic initialization}}
+
 template<typename F>
 __global__ void kernel(F f) { f(); } // dev-note2 {{called by 'kernel<(lambda}}
 
 __device__ void dev_fun(int *out) {
   // Check access device variables are allowed.
   int &ref_dev_var = global_dev_var;
   int &ref_constant_var = global_constant_var;
   int &ref_shared_var = global_shared_var;
   *out = ref_dev_var;
   *out = ref_constant_var;
   *out = ref_shared_var;
   *out = global_dev_var;
   *out = global_constant_var;
   *out = global_shared_var;
 
   // Check access of non-const host variables are not allowed.
   *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}
   *out = global_const_var;
   *out = global_constexpr_var;
+  *out = b1.x; // dev-error {{reference to __host__ variable 'b1' in __device__ function}}
+  *out = b2.x; // dev-error {{reference to __host__ variable 'b2' in __device__ function}}
+  *out = b3; // dev-error {{reference to __host__ variable 'b3' in __device__ function}}
   global_host_var = 1; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}
 
   // Check reference of non-constexpr host variables are not allowed.
   int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}
-  const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __device__ function}}
+  const int &ref_const_var = global_const_var;
   const int &ref_constexpr_var = global_constexpr_var;
   *out = ref_host_var;
   *out = ref_constexpr_var;
   *out = ref_const_var;
 
   // Check access member of non-constexpr struct type host variable is not allowed.
   *out = global_host_struct_var.x; // dev-error {{reference to __host__ variable 'global_host_struct_var' in __device__ function}}
-  *out = global_const_struct_var.x; // dev-error {{reference to __host__ variable 'global_const_struct_var' in __device__ function}}
+  *out = global_const_struct_var.x;
   *out = global_constexpr_struct_var.x;
   global_host_struct_var.x = 1; // dev-error {{reference to __host__ variable 'global_host_struct_var' in __device__ function}}
 
   // Check address taking of non-constexpr host variables is not allowed.
   int *p = &global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}
-  const int *cp = &global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __device__ function}}
+  const int *cp = &global_const_var;
   const int *cp2 = &global_constexpr_var;
 
   // Check access elements of non-constexpr host array is not allowed.
   *out = global_host_array[1]; // dev-error {{reference to __host__ variable 'global_host_array' in __device__ function}}
-  *out = global_const_array[1]; // dev-error {{reference to __host__ variable 'global_const_array' in __device__ function}}
+  *out = global_const_array[1];
   *out = global_constexpr_array[1];
 
   // Check ODR-use of host variables in namespace is not allowed.
   *out = X::host_var; // dev-error {{reference to __host__ variable 'host_var' in __device__ function}}
 
   // Check ODR-use of static host varables in class or file scope is not allowed.
   *out = A::host_var; // dev-error {{reference to __host__ variable 'host_var' in __device__ function}}
   *out = static_host_var; // dev-error {{reference to __host__ variable 'static_host_var' in __device__ function}}

 }
 
 __global__ void global_fun(int *out) {
   int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __global__ function}}
   int &ref_dev_var = global_dev_var;
   int &ref_constant_var = global_constant_var;
   int &ref_shared_var = global_shared_var;
   const int &ref_constexpr_var = global_constexpr_var;
-  const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __global__ function}}
+  const int &ref_const_var = global_const_var;
 
   *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __global__ function}}
   *out = global_dev_var;
   *out = global_constant_var;
   *out = global_shared_var;
   *out = global_constexpr_var;
   *out = global_const_var;
 
   *out = ref_host_var;
   *out = ref_dev_var;
   *out = ref_constant_var;
   *out = ref_shared_var;
   *out = ref_constexpr_var;
   *out = ref_const_var;
 }
 
 __host__ __device__ void host_dev_fun(int *out) {
   int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
   int &ref_dev_var = global_dev_var;
   int &ref_constant_var = global_constant_var;
   int &ref_shared_var = global_shared_var;
   const int &ref_constexpr_var = global_constexpr_var;
-  const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __host__ __device__ function}}
+  const int &ref_const_var = global_const_var;
 
   *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
   *out = global_dev_var;
   *out = global_constant_var;
   *out = global_shared_var;
   *out = global_constexpr_var;
   *out = global_const_var;
 

 
 void dev_lambda_capture_by_ref(int *out) {
   int &ref_host_var = global_host_var;
   kernel<<<1,1>>>([&]() {
   int &ref_dev_var = global_dev_var;
   int &ref_constant_var = global_constant_var;
   int &ref_shared_var = global_shared_var;
   const int &ref_constexpr_var = global_constexpr_var;
-  const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __host__ __device__ function}}
+  const int &ref_const_var = global_const_var;
 
   *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
                           // dev-error@-1 {{capture host variable 'out' by reference in device or host device lambda function}}
   *out = global_dev_var;
   *out = global_constant_var;
   *out = global_shared_var;
   *out = global_constexpr_var;
   *out = global_const_var;

 
 void dev_lambda_capture_by_copy(int *out) {
   int &ref_host_var = global_host_var;
   kernel<<<1,1>>>([=]() {
   int &ref_dev_var = global_dev_var;
   int &ref_constant_var = global_constant_var;
   int &ref_shared_var = global_shared_var;
   const int &ref_constexpr_var = global_constexpr_var;
-  const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __host__ __device__ function}}
+  const int &ref_const_var = global_const_var;
 
   *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
   *out = global_dev_var;
   *out = global_constant_var;
   *out = global_shared_var;
   *out = global_constexpr_var;
   *out = global_const_var;
 

 struct  not_a_texture {
   static not_a_texture<int> ref;
   __device__ void c() {
     auto &x = ref; // dev-error {{reference to __host__ variable 'ref' in __device__ function}}
   }
 };
 
 template<>
-not_a_texture<int> not_a_texture<int>::ref;
+not_a_texture<int> not_a_texture<int>::ref; // dev-note {{host variable declared here}}
 
 __device__ void test_not_a_texture() {
   not_a_texture<int> inst;
   inst.c(); // dev-note {{in instantiation of member function 'not_a_texture<int, 1, 1>::c' requested here}}
 }
 
 // Test static variable in host function used by device function.
 void test_static_var_host() {
   for (int i = 0; i < 10; i++) {
-    static int x;
+    static int x; // dev-note {{host variable declared here}}
     struct A {
       __device__ int f() {
         return x; // dev-error{{reference to __host__ variable 'x' in __device__ function}}
       }
     };
   }
 }
 

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

 }
 
 // Checks static device and constant variables are allowed in device and
 // host functions, and static host variables are not allowed in device
 // functions.
 
 static __device__ int x;
 static __constant__ int y;
-static int z;
+static int z; // dev-note {{host variable declared here}}
 
 __global__ void kernel(int *a) {
   a[0] = x;
   a[1] = y;
   a[2] = z;
   // dev-error@-1 {{reference to __host__ variable 'z' in __global__ function}}
 }