diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index b9aa2d2cfadb55..dc2a439702d0a4 100644
--- a/clang/include/clang/AST/ASTContext.h
+++ b/clang/include/clang/AST/ASTContext.h
@@ -656,6 +656,20 @@ class ASTContext : public RefCountedBase<ASTContext> {
   /// Returns the clang bytecode interpreter context.
   interp::Context &getInterpContext();
 
+  struct CUDAConstantEvalContext {
+    /// Do not allow wrong-sided variables in constant expressions.
+    bool NoWrongSidedVars = false;
+  } CUDAConstantEvalCtx;
+  struct CUDAConstantEvalContextRAII {
+    ASTContext &Ctx;
+    CUDAConstantEvalContext SavedCtx;
+    CUDAConstantEvalContextRAII(ASTContext &Ctx_, bool NoWrongSidedVars)
+        : Ctx(Ctx_), SavedCtx(Ctx_.CUDAConstantEvalCtx) {
+      Ctx_.CUDAConstantEvalCtx.NoWrongSidedVars = NoWrongSidedVars;
+    }
+    ~CUDAConstantEvalContextRAII() { Ctx.CUDAConstantEvalCtx = SavedCtx; }
+  };
+
   /// Returns the dynamic AST node parent map context.
   ParentMapContext &getParentMapContext();
 
diff --git a/clang/lib/AST/ExprConstant.cpp b/clang/lib/AST/ExprConstant.cpp
index fa0d22064f0eb4..dc9627d9d1d01f 100644
--- a/clang/lib/AST/ExprConstant.cpp
+++ b/clang/lib/AST/ExprConstant.cpp
@@ -983,6 +983,8 @@ namespace {
       discardCleanups();
     }
 
+    ASTContext &getCtx() const override { return Ctx; }
+
     void setEvaluatingDecl(APValue::LValueBase Base, APValue &Value,
                            EvaluatingDeclKind EDK = EvaluatingDeclKind::Ctor) {
       EvaluatingDecl = Base;
@@ -1116,8 +1118,6 @@ namespace {
 
     Expr::EvalStatus &getEvalStatus() const override { return EvalStatus; }
 
-    ASTContext &getCtx() const override { return Ctx; }
-
     // If we have a prior diagnostic, it will be noting that the expression
     // isn't a constant expression. This diagnostic is more important,
     // unless we require this evaluation to produce a constant expression.
@@ -2215,6 +2215,19 @@ static bool CheckLValueConstantExpression(EvalInfo &Info, SourceLocation Loc,
       if (!isForManglingOnly(Kind) && Var->hasAttr<DLLImportAttr>())
         // FIXME: Diagnostic!
         return false;
+
+      // In CUDA/HIP device compilation, only device side variables have
+      // constant addresses.
+      if (Info.getCtx().getLangOpts().CUDA &&
+          Info.getCtx().getLangOpts().CUDAIsDevice &&
+          Info.getCtx().CUDAConstantEvalCtx.NoWrongSidedVars) {
+        if ((!Var->hasAttr<CUDADeviceAttr>() &&
+             !Var->hasAttr<CUDAConstantAttr>() &&
+             !Var->getType()->isCUDADeviceBuiltinSurfaceType() &&
+             !Var->getType()->isCUDADeviceBuiltinTextureType()) ||
+            Var->hasAttr<HIPManagedAttr>())
+          return false;
+      }
     }
     if (const auto *FD = dyn_cast<const FunctionDecl>(BaseVD)) {
       // __declspec(dllimport) must be handled very carefully:
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 840b3daae63cfc..7387dae2d6749f 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -590,6 +590,8 @@ bool HasAllowedCUDADeviceStaticInitializer(Sema &S, VarDecl *VD,
   };
   auto IsConstantInit = [&](const Expr *Init) {
     assert(Init);
+    ASTContext::CUDAConstantEvalContextRAII EvalCtx(S.Context,
+                                                    /*NoWronSidedVars=*/true);
     return Init->isConstantInitializer(S.Context,
                                        VD->getType()->isReferenceType());
   };
diff --git a/clang/test/CodeGenCUDA/const-var.cu b/clang/test/CodeGenCUDA/const-var.cu
new file mode 100644
index 00000000000000..5ea5696ba8b8b6
--- /dev/null
+++ b/clang/test/CodeGenCUDA/const-var.cu
@@ -0,0 +1,54 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
+// RUN:   -emit-llvm -o - | FileCheck -check-prefix=DEV %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \
+// RUN:   -emit-llvm -o - | FileCheck -check-prefix=HOST %s
+
+// Negative tests.
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
+// RUN:   -emit-llvm -o - | FileCheck -check-prefix=DEV-NEG %s
+
+#include "Inputs/cuda.h"
+
+// Test const var initialized with address of a const var.
+// Both are promoted to device side.
+
+// DEV-DAG: @_ZN5Test1L1aE = internal addrspace(4) constant i32 1
+// DEV-DAG: @_ZN5Test11B2p1E = addrspace(4) externally_initialized constant i32* addrspacecast (i32 addrspace(4)* @_ZN5Test1L1aE to i32*)
+// DEV-DAG: @_ZN5Test11B2p2E = addrspace(4) externally_initialized constant i32* addrspacecast (i32 addrspace(4)* @_ZN5Test1L1aE to i32*)
+// DEV-DAG: @_ZN5Test12b2E = addrspace(1) externally_initialized global i32 1
+// HOST-DAG: @_ZN5Test1L1aE = internal constant i32 1
+// HOST-DAG: @_ZN5Test11B2p1E = constant i32* @_ZN5Test1L1aE
+// HOST-DAG: @_ZN5Test11B2p2E = internal constant i32* undef
+// HOST-DAG: @_ZN5Test12b1E = global i32 1
+// HOST-DAG: @_ZN5Test12b2E = internal global i32 undef
+namespace Test1 {
+const int a = 1;
+
+struct B {
+    static const int *const p1;
+    static __device__ const int *const p2;
+};
+const int *const B::p1 = &a;
+__device__ const int *const B::p2 = &a;
+int b1 = B::p1 == B::p2;
+__device__ int b2 = B::p1 == B::p2;
+}
+
+// Test const var initialized with address of a non-cost var.
+// Neither is promoted to device side.
+
+// DEV-NEG-NOT: @_ZN5Test2L1aE
+// DEV-NEG-NOT: @_ZN5Test21B1pE
+// HOST-DAG: @_ZN5Test21aE = global i32 1
+// HOST-DAG: @_ZN5Test21B1pE = constant i32* @_ZN5Test21aE
+
+namespace Test2 {
+int a = 1;
+
+struct B {
+    static int *const p;
+};
+int *const B::p = &a;
+}
diff --git a/clang/test/SemaCUDA/const-var.cu b/clang/test/SemaCUDA/const-var.cu
new file mode 100644
index 00000000000000..a4803aeae11f42
--- /dev/null
+++ b/clang/test/SemaCUDA/const-var.cu
@@ -0,0 +1,111 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
+// RUN:   -fsyntax-only -verify
+// RUN: %clang_cc1 -triple x86_64 -x hip %s \
+// RUN:   -fsyntax-only -verify=host
+
+// host-no-diagnostics
+
+#include "Inputs/cuda.h"
+
+// Test const var initialized with address of a const var.
+// Both are promoted to device side.
+
+namespace Test1 {
+const int a = 1;
+
+struct B {
+    static const int *const p;
+    __device__ static const int *const p2;
+};
+const int *const B::p = &a;
+// Const variable 'a' is treated as __constant__ on device side,
+// therefore its address can be used as initializer for another
+// device variable.
+__device__ const int *const B::p2 = &a;
+
+__device__ void f() {
+  int y = a;
+  const int *x = B::p;
+  const int *z = B::p2;
+}
+}
+
+// Test const var initialized with address of a non-cost var.
+// Neither is promoted to device side.
+
+namespace Test2 {
+int a = 1;
+// expected-note@-1{{host variable declared here}}
+
+struct B {
+    static int *const p;
+};
+int *const B::p = &a;
+// expected-note@-1{{const variable cannot be emitted on device side due to dynamic initialization}}
+
+__device__ void f() {
+  int y = a;
+  // expected-error@-1{{reference to __host__ variable 'a' in __device__ function}}
+  const int *x = B::p;
+  // expected-error@-1{{reference to __host__ variable 'p' in __device__ function}}
+}
+}
+
+// Test device var initialized with address of a non-const host var, __shared var,
+// __managed__ var, __device__ var, __constant__ var, texture var, surface var.
+
+namespace Test3 {
+struct textureReference {
+  int desc;
+};
+
+enum ReadMode {
+  ElementType = 0,
+  NormalizedFloat = 1
+};
+
+template <typename T, int dim = 1, enum ReadMode mode = ElementType>
+struct __attribute__((device_builtin_texture_type)) texture : public textureReference {
+};
+
+struct surfaceReference {
+  int desc;
+};
+
+template <typename T, int dim = 1>
+struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference {
+};
+
+// Partial specialization over `void`.
+template<int dim>
+struct __attribute__((device_builtin_surface_type)) surface<void, dim> : public surfaceReference {
+};
+
+texture<float, 2, ElementType> tex;
+surface<void, 2> surf;
+
+int a = 1;
+__shared__ int b;
+__managed__ int c = 1;
+__device__ int d = 1;
+__constant__ int e = 1;
+struct B {
+    __device__ static int *const p1;
+    __device__ static int *const p2;
+    __device__ static int *const p3;
+    __device__ static int *const p4;
+    __device__ static int *const p5;
+    __device__ static texture<float, 2, ElementType> *const p6;
+    __device__ static surface<void, 2> *const p7;
+};
+__device__ int *const B::p1 = &a;
+// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+__device__ int *const B::p2 = &b;
+// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+__device__ int *const B::p3 = &c;
+// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+__device__ int *const B::p4 = &d;
+__device__ int *const B::p5 = &e;
+__device__ texture<float, 2, ElementType> *const B::p6 = &tex;
+__device__ surface<void, 2> *const B::p7 = &surf;
+}