[clang] 622eaa4 - [HIP] Support __managed__ attribute

Yaxun Liu via cfe-commits cfe-commits at lists.llvm.org
Fri Jan 22 08:46:57 PST 2021


Author: Yaxun (Sam) Liu
Date: 2021-01-22T11:43:58-05:00
New Revision: 622eaa4a4cea17c2cec6942d9702b010deae392b

URL: https://github.com/llvm/llvm-project/commit/622eaa4a4cea17c2cec6942d9702b010deae392b
DIFF: https://github.com/llvm/llvm-project/commit/622eaa4a4cea17c2cec6942d9702b010deae392b.diff

LOG: [HIP] Support __managed__ attribute

This patch implements codegen for __managed__ variable attribute for HIP.

Diagnostics will be added later.

Differential Revision: https://reviews.llvm.org/D94814

Added: 
    clang/test/AST/Inputs/cuda.h
    clang/test/AST/ast-dump-managed-var.cu
    clang/test/CodeGenCUDA/managed-var.cu
    clang/test/SemaCUDA/managed-var.cu
    llvm/include/llvm/IR/ReplaceConstant.h
    llvm/lib/IR/ReplaceConstant.cpp

Modified: 
    clang/include/clang/Basic/Attr.td
    clang/include/clang/Basic/AttrDocs.td
    clang/include/clang/Basic/DiagnosticSemaKinds.td
    clang/lib/CodeGen/CGCUDANV.cpp
    clang/lib/CodeGen/CGCUDARuntime.h
    clang/lib/CodeGen/CodeGenModule.cpp
    clang/lib/Sema/SemaDeclAttr.cpp
    clang/test/CodeGenCUDA/Inputs/cuda.h
    clang/test/Misc/pragma-attribute-supported-attributes-list.test
    clang/test/SemaCUDA/Inputs/cuda.h
    clang/test/SemaCUDA/bad-attributes.cu
    clang/test/SemaCUDA/device-var-init.cu
    clang/test/SemaCUDA/function-overload.cu
    clang/test/SemaCUDA/union-init.cu
    llvm/lib/IR/CMakeLists.txt
    llvm/lib/Target/XCore/XCoreLowerThreadLocal.cpp

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index b30b91d3d4a6..bfd50f6a6779 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -324,6 +324,7 @@ class LangOpt<string name, code customCode = [{}]> {
 def MicrosoftExt : LangOpt<"MicrosoftExt">;
 def Borland : LangOpt<"Borland">;
 def CUDA : LangOpt<"CUDA">;
+def HIP : LangOpt<"HIP">;
 def SYCL : LangOpt<"SYCLIsDevice">;
 def COnly : LangOpt<"", "!LangOpts.CPlusPlus">;
 def CPlusPlus : LangOpt<"CPlusPlus">;
@@ -1115,6 +1116,13 @@ def CUDAHost : InheritableAttr {
   let Documentation = [Undocumented];
 }
 
+def HIPManaged : InheritableAttr {
+  let Spellings = [GNU<"managed">, Declspec<"__managed__">];
+  let Subjects = SubjectList<[Var]>;
+  let LangOpts = [HIP];
+  let Documentation = [HIPManagedAttrDocs];
+}
+
 def CUDAInvalidTarget : InheritableAttr {
   let Spellings = [];
   let Subjects = SubjectList<[Function]>;

diff  --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index fffede41db1e..170a0fe3d4c4 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -5419,6 +5419,17 @@ unbind runtime APIs.
   }];
 }
 
+def HIPManagedAttrDocs : Documentation {
+  let Category = DocCatDecl;
+  let Content = [{
+The ``__managed__`` attribute can be applied to a global variable declaration in HIP.
+A managed variable is emitted as an undefined global symbol in the device binary and is
+registered by ``__hipRegisterManagedVariable`` in init functions. The HIP runtime allocates
+managed memory and uses it to define the symbol when loading the device binary.
+A managed variable can be accessed in both device and host code.
+  }];
+}
+
 def LifetimeOwnerDocs : Documentation {
   let Category = DocCatDecl;
   let Content = [{

diff  --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 758b2ed3e90b..67c59f3ca09a 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8237,7 +8237,7 @@ def err_cuda_device_exceptions : Error<
   "%select{__device__|__global__|__host__|__host__ __device__}1 function">;
 def err_dynamic_var_init : Error<
     "dynamic initialization is not supported for "
-    "__device__, __constant__, and __shared__ variables.">;
+    "__device__, __constant__, __shared__, and __managed__ variables.">;
 def err_shared_var_init : Error<
     "initialization is not supported for __shared__ variables.">;
 def err_cuda_vla : Error<
@@ -8247,7 +8247,8 @@ def err_cuda_extern_shared : Error<"__shared__ variable %0 cannot be 'extern'">;
 def err_cuda_host_shared : Error<
     "__shared__ local variables not allowed in "
     "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
-def err_cuda_nonstatic_constdev: Error<"__constant__ and __device__ are not allowed on non-static local variables">;
+def err_cuda_nonstatic_constdev: Error<"__constant__, __device__, and "
+    "__managed__ are not allowed on non-static local variables">;
 def err_cuda_ovl_target : Error<
   "%select{__device__|__global__|__host__|__host__ __device__}0 function %1 "
   "cannot overload %select{__device__|__global__|__host__|__host__ __device__}2 function %3">;

diff  --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index 7c5ab39a85ec..33a2d6f4483e 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -21,6 +21,7 @@
 #include "llvm/IR/BasicBlock.h"
 #include "llvm/IR/Constants.h"
 #include "llvm/IR/DerivedTypes.h"
+#include "llvm/IR/ReplaceConstant.h"
 #include "llvm/Support/Format.h"
 
 using namespace clang;
@@ -128,13 +129,15 @@ class CGNVCUDARuntime : public CGCUDARuntime {
     DeviceVars.push_back({&Var,
                           VD,
                           {DeviceVarFlags::Variable, Extern, Constant,
-                           /*Normalized*/ false, /*Type*/ 0}});
+                           VD->hasAttr<HIPManagedAttr>(),
+                           /*Normalized*/ false, 0}});
   }
   void registerDeviceSurf(const VarDecl *VD, llvm::GlobalVariable &Var,
                           bool Extern, int Type) override {
     DeviceVars.push_back({&Var,
                           VD,
                           {DeviceVarFlags::Surface, Extern, /*Constant*/ false,
+                           /*Managed*/ false,
                            /*Normalized*/ false, Type}});
   }
   void registerDeviceTex(const VarDecl *VD, llvm::GlobalVariable &Var,
@@ -142,7 +145,7 @@ class CGNVCUDARuntime : public CGCUDARuntime {
     DeviceVars.push_back({&Var,
                           VD,
                           {DeviceVarFlags::Texture, Extern, /*Constant*/ false,
-                           Normalized, Type}});
+                           /*Managed*/ false, Normalized, Type}});
   }
 
   /// Creates module constructor function
@@ -380,6 +383,47 @@ void CGNVCUDARuntime::emitDeviceStubBodyLegacy(CodeGenFunction &CGF,
   CGF.EmitBlock(EndBlock);
 }
 
+// Replace the original variable Var with the address loaded from variable
+// ManagedVar populated by HIP runtime.
+static void replaceManagedVar(llvm::GlobalVariable *Var,
+                              llvm::GlobalVariable *ManagedVar) {
+  SmallVector<SmallVector<llvm::User *, 8>, 8> WorkList;
+  for (auto &&VarUse : Var->uses()) {
+    WorkList.push_back({VarUse.getUser()});
+  }
+  while (!WorkList.empty()) {
+    auto &&WorkItem = WorkList.pop_back_val();
+    auto *U = WorkItem.back();
+    if (isa<llvm::ConstantExpr>(U)) {
+      for (auto &&UU : U->uses()) {
+        WorkItem.push_back(UU.getUser());
+        WorkList.push_back(WorkItem);
+        WorkItem.pop_back();
+      }
+      continue;
+    }
+    if (auto *I = dyn_cast<llvm::Instruction>(U)) {
+      llvm::Value *OldV = Var;
+      llvm::Instruction *NewV =
+          new llvm::LoadInst(Var->getType(), ManagedVar, "ld.managed", false,
+                             llvm::Align(Var->getAlignment()), I);
+      WorkItem.pop_back();
+      // Replace constant expressions directly or indirectly using the managed
+      // variable with instructions.
+      for (auto &&Op : WorkItem) {
+        auto *CE = cast<llvm::ConstantExpr>(Op);
+        auto *NewInst = llvm::createReplacementInstr(CE, I);
+        NewInst->replaceUsesOfWith(OldV, NewV);
+        OldV = CE;
+        NewV = NewInst;
+      }
+      I->replaceUsesOfWith(OldV, NewV);
+    } else {
+      llvm_unreachable("Invalid use of managed variable");
+    }
+  }
+}
+
 /// Creates a function that sets up state on the host side for CUDA objects that
 /// have a presence on both the host and device sides. Specifically, registers
 /// the host side of kernel functions and device global variables with the CUDA
@@ -452,6 +496,13 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
   llvm::FunctionCallee RegisterVar = CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(VoidTy, RegisterVarParams, false),
       addUnderscoredPrefixToName("RegisterVar"));
+  // void __hipRegisterManagedVar(void **, char *, char *, const char *,
+  //                              size_t, unsigned)
+  llvm::Type *RegisterManagedVarParams[] = {VoidPtrPtrTy, CharPtrTy, CharPtrTy,
+                                            CharPtrTy,    VarSizeTy, IntTy};
+  llvm::FunctionCallee RegisterManagedVar = CGM.CreateRuntimeFunction(
+      llvm::FunctionType::get(VoidTy, RegisterManagedVarParams, false),
+      addUnderscoredPrefixToName("RegisterManagedVar"));
   // void __cudaRegisterSurface(void **, const struct surfaceReference *,
   //                            const void **, const char *, int, int);
   llvm::FunctionCallee RegisterSurf = CGM.CreateRuntimeFunction(
@@ -474,16 +525,34 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
     case DeviceVarFlags::Variable: {
       uint64_t VarSize =
           CGM.getDataLayout().getTypeAllocSize(Var->getValueType());
-      llvm::Value *Args[] = {
-          &GpuBinaryHandlePtr,
-          Builder.CreateBitCast(Var, VoidPtrTy),
-          VarName,
-          VarName,
-          llvm::ConstantInt::get(IntTy, Info.Flags.isExtern()),
-          llvm::ConstantInt::get(VarSizeTy, VarSize),
-          llvm::ConstantInt::get(IntTy, Info.Flags.isConstant()),
-          llvm::ConstantInt::get(IntTy, 0)};
-      Builder.CreateCall(RegisterVar, Args);
+      if (Info.Flags.isManaged()) {
+        auto ManagedVar = new llvm::GlobalVariable(
+            CGM.getModule(), Var->getType(),
+            /*isConstant=*/false, Var->getLinkage(),
+            /*Init=*/llvm::ConstantPointerNull::get(Var->getType()),
+            Twine(Var->getName() + ".managed"), /*InsertBefore=*/nullptr,
+            llvm::GlobalVariable::NotThreadLocal);
+        replaceManagedVar(Var, ManagedVar);
+        llvm::Value *Args[] = {
+            &GpuBinaryHandlePtr,
+            Builder.CreateBitCast(ManagedVar, VoidPtrTy),
+            Builder.CreateBitCast(Var, VoidPtrTy),
+            VarName,
+            llvm::ConstantInt::get(VarSizeTy, VarSize),
+            llvm::ConstantInt::get(IntTy, Var->getAlignment())};
+        Builder.CreateCall(RegisterManagedVar, Args);
+      } else {
+        llvm::Value *Args[] = {
+            &GpuBinaryHandlePtr,
+            Builder.CreateBitCast(Var, VoidPtrTy),
+            VarName,
+            VarName,
+            llvm::ConstantInt::get(IntTy, Info.Flags.isExtern()),
+            llvm::ConstantInt::get(VarSizeTy, VarSize),
+            llvm::ConstantInt::get(IntTy, Info.Flags.isConstant()),
+            llvm::ConstantInt::get(IntTy, 0)};
+        Builder.CreateCall(RegisterVar, Args);
+      }
       break;
     }
     case DeviceVarFlags::Surface:

diff  --git a/clang/lib/CodeGen/CGCUDARuntime.h b/clang/lib/CodeGen/CGCUDARuntime.h
index 19e70a2022a5..ba3404ead368 100644
--- a/clang/lib/CodeGen/CGCUDARuntime.h
+++ b/clang/lib/CodeGen/CGCUDARuntime.h
@@ -54,16 +54,19 @@ class CGCUDARuntime {
     unsigned Kind : 2;
     unsigned Extern : 1;
     unsigned Constant : 1;   // Constant variable.
+    unsigned Managed : 1;    // Managed variable.
     unsigned Normalized : 1; // Normalized texture.
     int SurfTexType;         // Type of surface/texutre.
 
   public:
-    DeviceVarFlags(DeviceVarKind K, bool E, bool C, bool N, int T)
-        : Kind(K), Extern(E), Constant(C), Normalized(N), SurfTexType(T) {}
+    DeviceVarFlags(DeviceVarKind K, bool E, bool C, bool M, bool N, int T)
+        : Kind(K), Extern(E), Constant(C), Managed(M), Normalized(N),
+          SurfTexType(T) {}
 
     DeviceVarKind getKind() const { return static_cast<DeviceVarKind>(Kind); }
     bool isExtern() const { return Extern; }
     bool isConstant() const { return Constant; }
+    bool isManaged() const { return Managed; }
     bool isNormalized() const { return Normalized; }
     int getSurfTexType() const { return SurfTexType; }
   };

diff  --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index bee51715bdc6..18d633911f55 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -4152,13 +4152,14 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
   // Shadows of initialized device-side global variables are also left
   // undefined.
   bool IsCUDAShadowVar =
-      !getLangOpts().CUDAIsDevice &&
+      !getLangOpts().CUDAIsDevice && !D->hasAttr<HIPManagedAttr>() &&
       (D->hasAttr<CUDAConstantAttr>() || D->hasAttr<CUDADeviceAttr>() ||
        D->hasAttr<CUDASharedAttr>());
   bool IsCUDADeviceShadowVar =
       getLangOpts().CUDAIsDevice &&
       (D->getType()->isCUDADeviceBuiltinSurfaceType() ||
-       D->getType()->isCUDADeviceBuiltinTextureType());
+       D->getType()->isCUDADeviceBuiltinTextureType() ||
+       D->hasAttr<HIPManagedAttr>());
   // HIP pinned shadow of initialized host-side global variables are also
   // left undefined.
   if (getLangOpts().CUDA &&

diff  --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index bd8ec2bdef76..30d08b3d4ac0 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -4493,7 +4493,8 @@ static void handleOptimizeNoneAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
 }
 
 static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
-  if (checkAttrMutualExclusion<CUDASharedAttr>(S, D, AL))
+  if (checkAttrMutualExclusion<CUDASharedAttr>(S, D, AL) ||
+      checkAttrMutualExclusion<HIPManagedAttr>(S, D, AL))
     return;
   const auto *VD = cast<VarDecl>(D);
   if (VD->hasLocalStorage()) {
@@ -4504,7 +4505,8 @@ static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
 }
 
 static void handleSharedAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
-  if (checkAttrMutualExclusion<CUDAConstantAttr>(S, D, AL))
+  if (checkAttrMutualExclusion<CUDAConstantAttr>(S, D, AL) ||
+      checkAttrMutualExclusion<HIPManagedAttr>(S, D, AL))
     return;
   const auto *VD = cast<VarDecl>(D);
   // extern __shared__ is only allowed on arrays with no length (e.g.
@@ -4569,9 +4571,33 @@ static void handleDeviceAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
       return;
     }
   }
+
+  if (auto *A = D->getAttr<CUDADeviceAttr>()) {
+    if (!A->isImplicit())
+      return;
+    D->dropAttr<CUDADeviceAttr>();
+  }
   D->addAttr(::new (S.Context) CUDADeviceAttr(S.Context, AL));
 }
 
+static void handleManagedAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
+  if (checkAttrMutualExclusion<CUDAConstantAttr>(S, D, AL) ||
+      checkAttrMutualExclusion<CUDASharedAttr>(S, D, AL)) {
+    return;
+  }
+
+  if (const auto *VD = dyn_cast<VarDecl>(D)) {
+    if (VD->hasLocalStorage()) {
+      S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
+      return;
+    }
+  }
+  if (!D->hasAttr<HIPManagedAttr>())
+    D->addAttr(::new (S.Context) HIPManagedAttr(S.Context, AL));
+  if (!D->hasAttr<CUDADeviceAttr>())
+    D->addAttr(CUDADeviceAttr::CreateImplicit(S.Context));
+}
+
 static void handleGNUInlineAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   const auto *Fn = cast<FunctionDecl>(D);
   if (!Fn->isInlineSpecified()) {
@@ -7793,6 +7819,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case ParsedAttr::AT_CUDAHost:
     handleSimpleAttributeWithExclusions<CUDAHostAttr, CUDAGlobalAttr>(S, D, AL);
     break;
+  case ParsedAttr::AT_HIPManaged:
+    handleManagedAttr(S, D, AL);
+    break;
   case ParsedAttr::AT_CUDADeviceBuiltinSurfaceType:
     handleSimpleAttributeWithExclusions<CUDADeviceBuiltinSurfaceTypeAttr,
                                         CUDADeviceBuiltinTextureTypeAttr>(S, D,

diff  --git a/clang/test/AST/Inputs/cuda.h b/clang/test/AST/Inputs/cuda.h
new file mode 100644
index 000000000000..405ef8bb807d
--- /dev/null
+++ b/clang/test/AST/Inputs/cuda.h
@@ -0,0 +1,54 @@
+/* Minimal declarations for CUDA support.  Testing purposes only. */
+
+#include <stddef.h>
+
+// Make this file work with nvcc, for testing compatibility.
+
+#ifndef __NVCC__
+#define __constant__ __attribute__((constant))
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+#define __host__ __attribute__((host))
+#define __shared__ __attribute__((shared))
+#define __managed__ __attribute__((managed))
+#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
+
+struct dim3 {
+  unsigned x, y, z;
+  __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}
+};
+
+#ifdef __HIP__
+typedef struct hipStream *hipStream_t;
+typedef enum hipError {} hipError_t;
+int hipConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
+                     hipStream_t stream = 0);
+extern "C" hipError_t __hipPushCallConfiguration(dim3 gridSize, dim3 blockSize,
+                                                 size_t sharedSize = 0,
+                                                 hipStream_t stream = 0);
+extern "C" hipError_t hipLaunchKernel(const void *func, dim3 gridDim,
+                                      dim3 blockDim, void **args,
+                                      size_t sharedMem,
+                                      hipStream_t stream);
+#else
+typedef struct cudaStream *cudaStream_t;
+typedef enum cudaError {} cudaError_t;
+
+extern "C" int cudaConfigureCall(dim3 gridSize, dim3 blockSize,
+                                 size_t sharedSize = 0,
+                                 cudaStream_t stream = 0);
+extern "C" int __cudaPushCallConfiguration(dim3 gridSize, dim3 blockSize,
+                                           size_t sharedSize = 0,
+                                           cudaStream_t stream = 0);
+extern "C" cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim,
+                                        dim3 blockDim, void **args,
+                                        size_t sharedMem, cudaStream_t stream);
+#endif
+
+// Host- and device-side placement new overloads.
+void *operator new(__SIZE_TYPE__, void *p) { return p; }
+void *operator new[](__SIZE_TYPE__, void *p) { return p; }
+__device__ void *operator new(__SIZE_TYPE__, void *p) { return p; }
+__device__ void *operator new[](__SIZE_TYPE__, void *p) { return p; }
+
+#endif // !__NVCC__

diff  --git a/clang/test/AST/ast-dump-managed-var.cu b/clang/test/AST/ast-dump-managed-var.cu
new file mode 100644
index 000000000000..862a70c81f9e
--- /dev/null
+++ b/clang/test/AST/ast-dump-managed-var.cu
@@ -0,0 +1,28 @@
+// RUN: %clang_cc1 -ast-dump -x hip %s | FileCheck %s
+// RUN: %clang_cc1 -ast-dump -fcuda-is-device -x hip %s | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// CHECK-LABEL: VarDecl {{.*}} m1 'int'
+// CHECK-NEXT: HIPManagedAttr
+// CHECK-NEXT: CUDADeviceAttr {{.*}}Implicit
+__managed__ int m1;
+
+// CHECK-LABEL: VarDecl {{.*}} m2 'int'
+// CHECK-NEXT: HIPManagedAttr
+// CHECK-NEXT: CUDADeviceAttr {{.*}}Implicit
+// CHECK-NOT: HIPManagedAttr
+// CHECK-NOT: CUDADeviceAttr
+__managed__ __managed__ int m2;
+
+// CHECK-LABEL: VarDecl {{.*}} m3 'int'
+// CHECK-NEXT: HIPManagedAttr
+// CHECK-NEXT: CUDADeviceAttr {{.*}}line
+// CHECK-NOT: CUDADeviceAttr {{.*}}Implicit
+__managed__ __device__ int m3;
+
+// CHECK-LABEL: VarDecl {{.*}} m3a 'int'
+// CHECK-NEXT: CUDADeviceAttr {{.*}}cuda.h
+// CHECK-NEXT: HIPManagedAttr
+// CHECK-NOT: CUDADeviceAttr {{.*}}Implicit
+__device__ __managed__ int m3a;

diff  --git a/clang/test/CodeGenCUDA/Inputs/cuda.h b/clang/test/CodeGenCUDA/Inputs/cuda.h
index 5d73b81041ab..daa6328c9499 100644
--- a/clang/test/CodeGenCUDA/Inputs/cuda.h
+++ b/clang/test/CodeGenCUDA/Inputs/cuda.h
@@ -7,6 +7,9 @@
 #define __global__ __attribute__((global))
 #define __host__ __attribute__((host))
 #define __shared__ __attribute__((shared))
+#if __HIP__
+#define __managed__ __attribute__((managed))
+#endif
 #define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
 
 struct dim3 {

diff  --git a/clang/test/CodeGenCUDA/managed-var.cu b/clang/test/CodeGenCUDA/managed-var.cu
new file mode 100644
index 000000000000..c95139784085
--- /dev/null
+++ b/clang/test/CodeGenCUDA/managed-var.cu
@@ -0,0 +1,100 @@
+// 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=DEV %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=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: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
+// RUN:   -emit-llvm -fgpu-rdc -o - -x hip %s | FileCheck \
+// RUN:   -check-prefixes=HOST %s
+
+#include "Inputs/cuda.h"
+
+// DEV-DAG: @x = {{.*}}addrspace(1) externally_initialized global i32 undef
+// HOST-DAG: @x = internal global i32 1
+// HOST-DAG: @x.managed = internal 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}};
+
+__global__ void foo(int *z) {
+  *z = x;
+  v[1].x = 2;
+}
+
+// HOST-LABEL: define {{.*}}@_Z4loadv()
+// HOST:  %ld.managed = load i32*, i32** @x.managed, align 4
+// HOST:  %0 = load i32, i32* %ld.managed, align 4
+// HOST:  ret i32 %0
+int load() {
+  return x;
+}
+
+// HOST-LABEL: define {{.*}}@_Z5storev()
+// HOST:  %ld.managed = load i32*, i32** @x.managed, align 4
+// HOST:  store i32 2, i32* %ld.managed, align 4
+void store() {
+  x = 2;
+}
+
+// HOST-LABEL: define {{.*}}@_Z10addr_takenv()
+// HOST:  %ld.managed = load i32*, i32** @x.managed, align 4
+// HOST:  store i32* %ld.managed, i32** %p, align 8
+// HOST:  %0 = load i32*, i32** %p, align 8
+// HOST:  store i32 3, i32* %0, align 4
+void addr_taken() {
+  int *p = &x;
+  *p = 3;
+}
+
+// HOST-LABEL: define {{.*}}@_Z5load2v()
+// HOST: %ld.managed = load [100 x %struct.vec]*, [100 x %struct.vec]** @v.managed, align 16
+// HOST:  %0 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %ld.managed, i64 0, i64 1, i32 0
+// HOST:  %1 = load float, float* %0, align 4
+// HOST:  ret float %1
+float load2() {
+  return v[1].x;
+}
+
+// HOST-LABEL: define {{.*}}@_Z5load3v()
+// HOST:  %ld.managed = load <{ %struct.vec, [99 x %struct.vec] }>*, <{ %struct.vec, [99 x %struct.vec] }>** @v2.managed, align 16
+// HOST:  %0 = bitcast <{ %struct.vec, [99 x %struct.vec] }>* %ld.managed to [100 x %struct.vec]*
+// HOST:  %1 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %0, i64 0, i64 1, i32 1
+// HOST:  %2 = load float, float* %1, align 4
+// HOST:  ret float %2
+float load3() {
+  return v2[1].y;
+}
+
+// HOST-LABEL: define {{.*}}@_Z11addr_taken2v()
+// HOST:  %ld.managed = load [100 x %struct.vec]*, [100 x %struct.vec]** @v.managed, align 16
+// HOST:  %0 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %ld.managed, i64 0, i64 1, i32 0
+// HOST:  %1 = ptrtoint float* %0 to i64
+// HOST:  %ld.managed1 = load <{ %struct.vec, [99 x %struct.vec] }>*, <{ %struct.vec, [99 x %struct.vec] }>** @v2.managed, align 16
+// HOST:  %2 = bitcast <{ %struct.vec, [99 x %struct.vec] }>* %ld.managed1 to [100 x %struct.vec]*
+// HOST:  %3 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %2, i64 0, i64 1, i32 1
+// HOST:  %4 = ptrtoint float* %3 to i64
+// HOST:  %5 = sub i64 %4, %1
+// HOST:  %6 = sdiv i64 %5, 4
+// HOST:  %7 = sitofp i64 %6 to float
+// HOST:  ret float %7
+float addr_taken2() {
+  return (float)reinterpret_cast<long>(&(v2[1].y)-&(v[1].x));
+}
+
+// HOST-DAG: __hipRegisterManagedVar({{.*}}@x.managed {{.*}}@x {{.*}}@[[DEVNAMEX]]{{.*}}, i64 4, i32 4)
+// HOST-DAG: declare void @__hipRegisterManagedVar(i8**, i8*, i8*, i8*, i64, i32)

diff  --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index 270b377c2064..cb62f56912aa 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -66,6 +66,7 @@
 // CHECK-NEXT: FlagEnum (SubjectMatchRule_enum)
 // CHECK-NEXT: Flatten (SubjectMatchRule_function)
 // CHECK-NEXT: GNUInline (SubjectMatchRule_function)
+// CHECK-NEXT: HIPManaged (SubjectMatchRule_variable)
 // CHECK-NEXT: Hot (SubjectMatchRule_function)
 // CHECK-NEXT: IBAction (SubjectMatchRule_objc_method_is_instance)
 // CHECK-NEXT: IFunc (SubjectMatchRule_function)

diff  --git a/clang/test/SemaCUDA/Inputs/cuda.h b/clang/test/SemaCUDA/Inputs/cuda.h
index 901f8e0c17cf..405ef8bb807d 100644
--- a/clang/test/SemaCUDA/Inputs/cuda.h
+++ b/clang/test/SemaCUDA/Inputs/cuda.h
@@ -10,6 +10,7 @@
 #define __global__ __attribute__((global))
 #define __host__ __attribute__((host))
 #define __shared__ __attribute__((shared))
+#define __managed__ __attribute__((managed))
 #define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
 
 struct dim3 {

diff  --git a/clang/test/SemaCUDA/bad-attributes.cu b/clang/test/SemaCUDA/bad-attributes.cu
index a990598e7b7e..8ac6fa1ab5c2 100644
--- a/clang/test/SemaCUDA/bad-attributes.cu
+++ b/clang/test/SemaCUDA/bad-attributes.cu
@@ -64,11 +64,11 @@ __global__ static inline void foobar() {};
 
 __constant__ int global_constant;
 void host_fn() {
-  __constant__ int c; // expected-error {{__constant__ and __device__ are not allowed on non-static local variables}}
+  __constant__ int c; // expected-error {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}}
   __shared__ int s; // expected-error {{__shared__ local variables not allowed in __host__ functions}}
 }
 __device__ void device_fn() {
-  __constant__ int c; // expected-error {{__constant__ and __device__ are not allowed on non-static local variables}}
+  __constant__ int c; // expected-error {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}}
 }
 
 typedef __attribute__((device_builtin_surface_type)) unsigned long long s0_ty; // expected-warning {{'device_builtin_surface_type' attribute only applies to classes}}

diff  --git a/clang/test/SemaCUDA/device-var-init.cu b/clang/test/SemaCUDA/device-var-init.cu
index 88350f56651c..9d499bddbe1b 100644
--- a/clang/test/SemaCUDA/device-var-init.cu
+++ b/clang/test/SemaCUDA/device-var-init.cu
@@ -16,11 +16,11 @@ __shared__ int s_v_i = 1;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 
 __device__ int d_v_f = f();
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ int s_v_f = f();
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ int c_v_f = f();
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __shared__ T s_t_i = {2};
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
@@ -32,175 +32,175 @@ __shared__ ECD s_ecd_i{};
 __constant__ ECD c_ecd_i{};
 
 __device__ EC d_ec_i(3);
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ EC s_ec_i(3);
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ EC c_ec_i(3);
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ EC d_ec_i2 = {3};
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ EC s_ec_i2 = {3};
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ EC c_ec_i2 = {3};
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ ETC d_etc_i(3);
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ ETC s_etc_i(3);
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ ETC c_etc_i(3);
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ ETC d_etc_i2 = {3};
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ ETC s_etc_i2 = {3};
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ ETC c_etc_i2 = {3};
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ UC d_uc;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ UC s_uc;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ UC c_uc;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ UD d_ud;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ UD s_ud;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ UD c_ud;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ ECI d_eci;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ ECI s_eci;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ ECI c_eci;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ NEC d_nec;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ NEC s_nec;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NEC c_nec;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ NED d_ned;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ NED s_ned;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NED c_ned;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ NCV d_ncv;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ NCV s_ncv;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NCV c_ncv;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ VD d_vd;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ VD s_vd;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ VD c_vd;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ NCF d_ncf;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ NCF s_ncf;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NCF c_ncf;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __shared__ NCFS s_ncfs;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 
 __device__ UTC d_utc;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ UTC s_utc;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ UTC c_utc;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ UTC d_utc_i(3);
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ UTC s_utc_i(3);
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ UTC c_utc_i(3);
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ NETC d_netc;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ NETC s_netc;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NETC c_netc;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ NETC d_netc_i(3);
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ NETC s_netc_i(3);
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NETC c_netc_i(3);
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ EC_I_EC1 d_ec_i_ec1;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ EC_I_EC1 s_ec_i_ec1;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ EC_I_EC1 c_ec_i_ec1;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ T_V_T d_t_v_t;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ T_V_T s_t_v_t;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_V_T c_t_v_t;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ T_B_NEC d_t_b_nec;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ T_B_NEC s_t_b_nec;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_B_NEC c_t_b_nec;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ T_F_NEC d_t_f_nec;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ T_F_NEC s_t_f_nec;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_F_NEC c_t_f_nec;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ T_FA_NEC d_t_fa_nec;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ T_FA_NEC s_t_fa_nec;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_FA_NEC c_t_fa_nec;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ T_B_NED d_t_b_ned;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ T_B_NED s_t_b_ned;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_B_NED c_t_b_ned;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ T_F_NED d_t_f_ned;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ T_F_NED s_t_f_ned;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_F_NED c_t_f_ned;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ T_FA_NED d_t_fa_ned;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ T_FA_NED s_t_fa_ned;
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_FA_NED c_t_fa_ned;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 // Verify that local variables may be static on device
 // side and that they conform to the initialization constraints.
@@ -216,20 +216,20 @@ __device__ void df_sema() {
 
   // __shared__ does not need to be explicitly static.
   __shared__ int lsi;
-  // __constant__ and __device__ can not be non-static local
+  // __constant__, __device__, and __managed__ can not be non-static local
   __constant__ int lci;
-  // expected-error at -1 {{__constant__ and __device__ are not allowed on non-static local variables}}
+  // expected-error at -1 {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}}
   __device__ int ldi;
-  // expected-error at -1 {{__constant__ and __device__ are not allowed on non-static local variables}}
+  // expected-error at -1 {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}}
 
   // Same test cases as for the globals above.
 
   static __device__ int d_v_f = f();
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ int s_v_f = f();
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ int c_v_f = f();
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __shared__ T s_t_i = {2};
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
@@ -241,175 +241,175 @@ __device__ void df_sema() {
   static __constant__ ECD c_ecd_i;
 
   static __device__ EC d_ec_i(3);
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ EC s_ec_i(3);
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ EC c_ec_i(3);
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ EC d_ec_i2 = {3};
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ EC s_ec_i2 = {3};
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ EC c_ec_i2 = {3};
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ ETC d_etc_i(3);
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ ETC s_etc_i(3);
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ ETC c_etc_i(3);
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ ETC d_etc_i2 = {3};
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ ETC s_etc_i2 = {3};
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ ETC c_etc_i2 = {3};
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ UC d_uc;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ UC s_uc;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ UC c_uc;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ UD d_ud;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ UD s_ud;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ UD c_ud;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ ECI d_eci;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ ECI s_eci;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ ECI c_eci;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ NEC d_nec;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ NEC s_nec;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ NEC c_nec;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ NED d_ned;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ NED s_ned;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ NED c_ned;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ NCV d_ncv;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ NCV s_ncv;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ NCV c_ncv;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ VD d_vd;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ VD s_vd;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ VD c_vd;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ NCF d_ncf;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ NCF s_ncf;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ NCF c_ncf;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __shared__ NCFS s_ncfs;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 
   static __device__ UTC d_utc;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ UTC s_utc;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ UTC c_utc;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ UTC d_utc_i(3);
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ UTC s_utc_i(3);
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ UTC c_utc_i(3);
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ NETC d_netc;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ NETC s_netc;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ NETC c_netc;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ NETC d_netc_i(3);
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ NETC s_netc_i(3);
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ NETC c_netc_i(3);
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ EC_I_EC1 d_ec_i_ec1;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ EC_I_EC1 s_ec_i_ec1;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ EC_I_EC1 c_ec_i_ec1;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ T_V_T d_t_v_t;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ T_V_T s_t_v_t;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ T_V_T c_t_v_t;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ T_B_NEC d_t_b_nec;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ T_B_NEC s_t_b_nec;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ T_B_NEC c_t_b_nec;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ T_F_NEC d_t_f_nec;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ T_F_NEC s_t_f_nec;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ T_F_NEC c_t_f_nec;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ T_FA_NEC d_t_fa_nec;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ T_FA_NEC s_t_fa_nec;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ T_FA_NEC c_t_fa_nec;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ T_B_NED d_t_b_ned;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ T_B_NED s_t_b_ned;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ T_B_NED c_t_b_ned;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ T_F_NED d_t_f_ned;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ T_F_NED s_t_f_ned;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ T_F_NED c_t_f_ned;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ T_FA_NED d_t_fa_ned;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ T_FA_NED s_t_fa_ned;
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ T_FA_NED c_t_fa_ned;
-  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 }
 
 __host__ __device__ void hd_sema() {

diff  --git a/clang/test/SemaCUDA/function-overload.cu b/clang/test/SemaCUDA/function-overload.cu
index 574b65ee7fd8..822e25996820 100644
--- a/clang/test/SemaCUDA/function-overload.cu
+++ b/clang/test/SemaCUDA/function-overload.cu
@@ -463,7 +463,7 @@ int test_constexpr_overload(C2 &x, C2 &y) {
 // Verify no ambiguity for new operator.
 void *a = new int;
 __device__ void *b = new int;
-// expected-error at -1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 // Verify no ambiguity for new operator.
 template<typename _Tp> _Tp&& f();

diff  --git a/clang/test/SemaCUDA/managed-var.cu b/clang/test/SemaCUDA/managed-var.cu
new file mode 100644
index 000000000000..3f699b79a043
--- /dev/null
+++ b/clang/test/SemaCUDA/managed-var.cu
@@ -0,0 +1,54 @@
+// RUN: %clang_cc1 -fsyntax-only -verify -x hip %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify -x hip %s
+// RUN: %clang_cc1 -fsyntax-only -fgpu-rdc -verify -x hip %s
+// RUN: %clang_cc1 -fsyntax-only -fgpu-rdc -fcuda-is-device -verify -x hip %s
+
+#include "Inputs/cuda.h"
+
+struct A {
+  int a;
+  A() { a = 1; }
+};
+
+__managed__ int m1;
+
+__managed__ __managed__ int m2;
+
+__managed__ __device__ int m3;
+__device__ __managed__ int m3a;
+
+__managed__ __constant__ int m4;
+// expected-error at -1 {{'constant' and 'managed' attributes are not compatible}}
+// expected-note at -2 {{conflicting attribute is here}}
+
+__constant__ __managed__ int m4a;
+// expected-error at -1 {{'managed' and 'constant' attributes are not compatible}}
+// expected-note at -2 {{conflicting attribute is here}}
+
+__managed__ __shared__ int m5;
+// expected-error at -1 {{'shared' and 'managed' attributes are not compatible}}
+// expected-note at -2 {{conflicting attribute is here}}
+
+__shared__ __managed__ int m5a;
+// expected-error at -1 {{'managed' and 'shared' attributes are not compatible}}
+// expected-note at -2 {{conflicting attribute is here}}
+
+__managed__ __global__ int m6;
+// expected-warning at -1 {{'global' attribute only applies to functions}}
+
+void func() {
+  __managed__ int m7;
+  // expected-error at -1 {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}}
+}
+
+__attribute__((managed(1))) int m8;
+// expected-error at -1 {{'managed' attribute takes no arguments}}
+
+__managed__ void func2() {}
+// expected-warning at -1 {{'managed' attribute only applies to variables}}
+
+typedef __managed__ int managed_int;
+// expected-warning at -1 {{'managed' attribute only applies to variables}}
+
+__managed__ A a;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}

diff  --git a/clang/test/SemaCUDA/union-init.cu b/clang/test/SemaCUDA/union-init.cu
index a633975e3776..9e4d14a71069 100644
--- a/clang/test/SemaCUDA/union-init.cu
+++ b/clang/test/SemaCUDA/union-init.cu
@@ -31,9 +31,9 @@ union D {
 
 __device__ B b;
 __device__ C c;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __device__ D d;
-// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ void foo() {
   __shared__ B b;

diff  --git a/llvm/include/llvm/IR/ReplaceConstant.h b/llvm/include/llvm/IR/ReplaceConstant.h
new file mode 100644
index 000000000000..753f6d558ef8
--- /dev/null
+++ b/llvm/include/llvm/IR/ReplaceConstant.h
@@ -0,0 +1,28 @@
+//===- ReplaceConstant.h - Replacing LLVM constant expressions --*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file declares the utility function for replacing LLVM constant
+// expressions by instructions.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_IR_REPLACECONSTANT_H
+#define LLVM_IR_REPLACECONSTANT_H
+
+#include "llvm/IR/Constants.h"
+#include "llvm/IR/Instruction.h"
+
+namespace llvm {
+
+/// Create a replacement instruction for constant expression \p CE and insert
+/// it before \p Instr.
+Instruction *createReplacementInstr(ConstantExpr *CE, Instruction *Instr);
+
+} // end namespace llvm
+
+#endif // LLVM_IR_REPLACECONSTANT_H

diff  --git a/llvm/lib/IR/CMakeLists.txt b/llvm/lib/IR/CMakeLists.txt
index ca570121460b..fb4993742e85 100644
--- a/llvm/lib/IR/CMakeLists.txt
+++ b/llvm/lib/IR/CMakeLists.txt
@@ -49,6 +49,7 @@ add_llvm_component_library(LLVMCore
   SafepointIRVerifier.cpp
   ProfileSummary.cpp
   PseudoProbe.cpp
+  ReplaceConstant.cpp
   Statepoint.cpp
   StructuralHash.cpp
   Type.cpp

diff  --git a/llvm/lib/IR/ReplaceConstant.cpp b/llvm/lib/IR/ReplaceConstant.cpp
new file mode 100644
index 000000000000..7efa525d427e
--- /dev/null
+++ b/llvm/lib/IR/ReplaceConstant.cpp
@@ -0,0 +1,70 @@
+//===- ReplaceConstant.cpp - Replace LLVM constant expression--------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements a utility function for replacing LLVM constant
+// expressions by instructions.
+//
+//===----------------------------------------------------------------------===//
+
+#include "llvm/IR/ReplaceConstant.h"
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/NoFolder.h"
+
+namespace llvm {
+// Replace a constant expression by instructions with equivalent operations at
+// a specified location.
+Instruction *createReplacementInstr(ConstantExpr *CE, Instruction *Instr) {
+  IRBuilder<NoFolder> Builder(Instr);
+  unsigned OpCode = CE->getOpcode();
+  switch (OpCode) {
+  case Instruction::GetElementPtr: {
+    SmallVector<Value *, 4> CEOpVec(CE->operands());
+    ArrayRef<Value *> CEOps(CEOpVec);
+    return dyn_cast<Instruction>(
+        Builder.CreateInBoundsGEP(cast<GEPOperator>(CE)->getSourceElementType(),
+                                  CEOps[0], CEOps.slice(1)));
+  }
+  case Instruction::Add:
+  case Instruction::Sub:
+  case Instruction::Mul:
+  case Instruction::UDiv:
+  case Instruction::SDiv:
+  case Instruction::FDiv:
+  case Instruction::URem:
+  case Instruction::SRem:
+  case Instruction::FRem:
+  case Instruction::Shl:
+  case Instruction::LShr:
+  case Instruction::AShr:
+  case Instruction::And:
+  case Instruction::Or:
+  case Instruction::Xor:
+    return dyn_cast<Instruction>(
+        Builder.CreateBinOp((Instruction::BinaryOps)OpCode, CE->getOperand(0),
+                            CE->getOperand(1), CE->getName()));
+  case Instruction::Trunc:
+  case Instruction::ZExt:
+  case Instruction::SExt:
+  case Instruction::FPToUI:
+  case Instruction::FPToSI:
+  case Instruction::UIToFP:
+  case Instruction::SIToFP:
+  case Instruction::FPTrunc:
+  case Instruction::FPExt:
+  case Instruction::PtrToInt:
+  case Instruction::IntToPtr:
+  case Instruction::BitCast:
+    return dyn_cast<Instruction>(
+        Builder.CreateCast((Instruction::CastOps)OpCode, CE->getOperand(0),
+                           CE->getType(), CE->getName()));
+  default:
+    llvm_unreachable("Unhandled constant expression!\n");
+  }
+}
+} // namespace llvm

diff  --git a/llvm/lib/Target/XCore/XCoreLowerThreadLocal.cpp b/llvm/lib/Target/XCore/XCoreLowerThreadLocal.cpp
index bd269f7f4c43..6528154ab0e2 100644
--- a/llvm/lib/Target/XCore/XCoreLowerThreadLocal.cpp
+++ b/llvm/lib/Target/XCore/XCoreLowerThreadLocal.cpp
@@ -21,6 +21,7 @@
 #include "llvm/IR/IntrinsicsXCore.h"
 #include "llvm/IR/Module.h"
 #include "llvm/IR/NoFolder.h"
+#include "llvm/IR/ReplaceConstant.h"
 #include "llvm/IR/ValueHandle.h"
 #include "llvm/Pass.h"
 #include "llvm/Support/CommandLine.h"
@@ -74,57 +75,6 @@ createLoweredInitializer(ArrayType *NewType, Constant *OriginalInitializer) {
   return ConstantArray::get(NewType, Elements);
 }
 
-static Instruction *
-createReplacementInstr(ConstantExpr *CE, Instruction *Instr) {
-  IRBuilder<NoFolder> Builder(Instr);
-  unsigned OpCode = CE->getOpcode();
-  switch (OpCode) {
-    case Instruction::GetElementPtr: {
-      SmallVector<Value *, 4> CEOpVec(CE->operands());
-      ArrayRef<Value *> CEOps(CEOpVec);
-      return dyn_cast<Instruction>(Builder.CreateInBoundsGEP(
-          cast<GEPOperator>(CE)->getSourceElementType(), CEOps[0],
-          CEOps.slice(1)));
-    }
-    case Instruction::Add:
-    case Instruction::Sub:
-    case Instruction::Mul:
-    case Instruction::UDiv:
-    case Instruction::SDiv:
-    case Instruction::FDiv:
-    case Instruction::URem:
-    case Instruction::SRem:
-    case Instruction::FRem:
-    case Instruction::Shl:
-    case Instruction::LShr:
-    case Instruction::AShr:
-    case Instruction::And:
-    case Instruction::Or:
-    case Instruction::Xor:
-      return dyn_cast<Instruction>(
-                  Builder.CreateBinOp((Instruction::BinaryOps)OpCode,
-                                      CE->getOperand(0), CE->getOperand(1),
-                                      CE->getName()));
-    case Instruction::Trunc:
-    case Instruction::ZExt:
-    case Instruction::SExt:
-    case Instruction::FPToUI:
-    case Instruction::FPToSI:
-    case Instruction::UIToFP:
-    case Instruction::SIToFP:
-    case Instruction::FPTrunc:
-    case Instruction::FPExt:
-    case Instruction::PtrToInt:
-    case Instruction::IntToPtr:
-    case Instruction::BitCast:
-      return dyn_cast<Instruction>(
-                  Builder.CreateCast((Instruction::CastOps)OpCode,
-                                     CE->getOperand(0), CE->getType(),
-                                     CE->getName()));
-    default:
-      llvm_unreachable("Unhandled constant expression!\n");
-  }
-}
 
 static bool replaceConstantExprOp(ConstantExpr *CE, Pass *P) {
   do {


        


More information about the cfe-commits mailing list