Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -11656,6 +11656,10 @@
   void maybeAddCUDAHostDeviceAttrs(FunctionDecl *FD,
                                    const LookupResult &Previous);
 
+  /// May add implicit CUDAConstantAttr attribute to VD, depending on VD
+  /// and current compilation settings.
+  void MaybeAddCUDAConstantAttr(VarDecl *VD);
+
 public:
   /// Check whether we're allowed to call Callee from the current context.
   ///
Index: clang/lib/Sema/SemaCUDA.cpp
===================================================================
--- clang/lib/Sema/SemaCUDA.cpp
+++ clang/lib/Sema/SemaCUDA.cpp
@@ -528,9 +528,12 @@
     // 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>()))
-      AllowedInit = VD->getInit()->isConstantInitializer(
-          Context, VD->getType()->isReferenceType());
+        (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>())) {
+      auto *Init = VD->getInit();
+      AllowedInit = (Init->isValueDependent() && VD->isConstexpr()) ||
+                    Init->isConstantInitializer(
+                        Context, VD->getType()->isReferenceType());
+    }
 
     // Also make sure that destructor, if there is one, is empty.
     if (AllowedInit)
@@ -627,6 +630,13 @@
   NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
 }
 
+void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
+  if (getLangOpts().CUDAIsDevice && VD->isConstexpr() &&
+      (VD->isFileVarDecl() || VD->isStaticDataMember())) {
+    VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
+  }
+}
+
 Sema::DeviceDiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
                                                    unsigned DiagID) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
Index: clang/lib/Sema/SemaDecl.cpp
===================================================================
--- clang/lib/Sema/SemaDecl.cpp
+++ clang/lib/Sema/SemaDecl.cpp
@@ -7081,6 +7081,7 @@
 
   case CSK_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.
Index: clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
===================================================================
--- clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -4836,6 +4836,7 @@
   NewVar->setCXXForRangeDecl(OldVar->isCXXForRangeDecl());
   NewVar->setObjCForDecl(OldVar->isObjCForDecl());
   NewVar->setConstexpr(OldVar->isConstexpr());
+  MaybeAddCUDAConstantAttr(NewVar);
   NewVar->setInitCapture(OldVar->isInitCapture());
   NewVar->setPreviousDeclInSameBlockScope(
       OldVar->isPreviousDeclInSameBlockScope());
Index: clang/test/CodeGenCUDA/constexpr-variables.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/constexpr-variables.cu
@@ -0,0 +1,43 @@
+// RUN: %clang_cc1 -std=c++14 %s -emit-llvm -o - -triple nvptx \
+// RUN:   -fcuda-is-device | FileCheck --check-prefixes=CXX14 %s
+// RUN: %clang_cc1 -std=c++17 %s -emit-llvm -o - -triple nvptx \
+// RUN:   -fcuda-is-device | FileCheck --check-prefixes=CXX17 %s
+
+#include "Inputs/cuda.h"
+
+// COM: @_ZL1a = internal {{.*}}constant i32 7
+constexpr int a = 7;
+__constant__ const int &use_a = a;
+
+namespace B {
+ // COM: @_ZN1BL1bE = internal {{.*}}constant i32 9
+  constexpr int b = 9;
+}
+__constant__ const int &use_B_b = B::b;
+
+struct Q {
+  // CXX14: @_ZN1Q2k2E = {{.*}}externally_initialized constant i32 6
+  // CXX17: @_ZN1Q2k2E = internal {{.*}}constant i32 6
+  // CXX14: @_ZN1Q2k1E = available_externally {{.*}}constant i32 5
+  // CXX17: @_ZN1Q2k1E = linkonce_odr {{.*}}constant i32 5
+  static constexpr int k1 = 5;
+  static constexpr int k2 = 6;
+};
+constexpr int Q::k2;
+
+__constant__ const int &use_Q_k1 = Q::k1;
+__constant__ const int &use_Q_k2 = Q::k2;
+
+template<typename T> struct X {
+  // CXX14: @_ZN1XIiE1aE = available_externally {{.*}}constant i32 123
+  // CXX17: @_ZN1XIiE1aE = linkonce_odr {{.*}}constant i32 123
+  static constexpr int a = 123;
+};
+__constant__ const int &use_X_a = X<int>::a;
+
+template <typename T, T a, T b> struct A {
+  // CXX14: @_ZN1AIiLi1ELi2EE1xE = available_externally {{.*}}constant i32 2
+  // CXX17: @_ZN1AIiLi1ELi2EE1xE = linkonce_odr {{.*}}constant i32 2
+  constexpr static T x = a * b;
+};
+__constant__ const int &y = A<int, 1, 2>::x;
Index: clang/test/SemaCUDA/constexpr-variables.cu
===================================================================
--- /dev/null
+++ clang/test/SemaCUDA/constexpr-variables.cu
@@ -0,0 +1,50 @@
+// RUN: %clang_cc1 -std=c++14 %s -emit-llvm -o - -triple nvptx64-nvidia-cuda \
+// RUN:   -fcuda-is-device -verify -fsyntax-only
+// RUN: %clang_cc1 -std=c++17 %s -emit-llvm -o - -triple nvptx64-nvidia-cuda \
+// RUN:   -fcuda-is-device -verify -fsyntax-only
+// RUN: %clang_cc1 -std=c++14 %s -emit-llvm -o - \
+// RUN:   -triple x86_64-unknown-linux-gnu -verify -fsyntax-only
+// RUN: %clang_cc1 -std=c++17 %s -emit-llvm -o - \
+// RUN:   -triple x86_64-unknown-linux-gnu -verify -fsyntax-only
+
+// expected-no-diagnostics
+
+#include "Inputs/cuda.h"
+
+// Check constexpr local variable is not made static variable.
+template<typename T>
+__host__ __device__ void foo(const T **a) {
+  static const T b = sizeof(a);
+  static constexpr T c = sizeof(a);
+  const T d = sizeof(a);
+  constexpr T e = sizeof(a);
+  a[0] = &b;
+  a[1] = &c;
+  a[2] = &d;
+  a[3] = &e;
+}
+
+__device__ void device_fun(const int **a) {
+  constexpr int b = sizeof(a);
+  static constexpr int c = sizeof(a);
+  a[0] = &b;
+  a[1] = &c;
+  foo(a);
+}
+
+void host_fun(const int **a) {
+  constexpr int b = sizeof(a);
+  static constexpr int c = sizeof(a);
+  a[0] = &b;
+  a[1] = &c;
+  foo(a);
+}
+
+__host__ __device__ void host_device_fun(const int **a) {
+  constexpr int b = sizeof(a);
+  static constexpr int c = sizeof(a);
+  a[0] = &b;
+  a[1] = &c;
+  foo(a);
+}
+