[HIP] Support `__managed__` attribute

[yaxunl]

Summary

This patch implements codegen for __managed__ variable attribute for HIP.

Diagnostics will be added later.

[tra]

Presumably, __managed__ variables would have to be memory-mapped into the host address space.

clang/lib/CodeGen/CGCUDANV.cpp

391

VarUse ?

395

WorkItem

552–553

This will always be false, because all cases where Info.Flags.isManaged() is true are handled in the other branch of this if.

clang/lib/Sema/SemaDeclAttr.cpp

7823

The code changes in the patch appear to treat __managed__ variable as a superset of a __device__ var. If that's indeed the case, adding an implicit __device__ attribute here would help to simplify the code. This way the existing code can handle generic __device__ var functionality without additional changes, and would use __managed__ checks for the cases specific for managed vars only.

llvm/lib/IR/ReplaceConstant.cpp

21

The function could use a descriptive comment.

Despite the comment above, it does seem to do any replacing itself. Presumably it creates a an instruction that would perform an equivalent calculation at runtime.

[aaron.ballman]

clang/include/clang/Basic/AttrDocs.td

5426

5427–5429

clang/lib/Sema/SemaDeclAttr.cpp

7823

I think you're missing changes to the CUDA global attr to check for mutual exclusions with __managed__ as well. Also, I think this won't do the right thing for redeclarations, like:

__device__ extern int i;
__managed__ extern int i;

7823

This way the existing code can handle generic device var functionality without additional changes, and would use managed checks for the cases specific for managed vars only.

Another alternative to consider is to not create a new semantic attribute named HIPManagedAttr but to instead add a new spelling to CUDADevice in Attr.td, giving the class an Accessor to distinguish which spelling the user wrote in code, and use that accessor for the specific cases for managed vars.

[yaxunl]

clang/lib/Sema/SemaDeclAttr.cpp

7823

I will add an implicit device attr for managed attr. I think this will have a more readable AST.

7823

Actually __managed__ does not need exclude with __global__ since __global__ cannot be used with variable. However, __managed__ need to exclude with __constant__ and __shared__. I will add handling for that.

__managed__ should behave the same as __device__ regarding compatibility with extern. It is allowed to be with extern. Currently for extern __device__ var, on device side, it is treated as a normal device variable with definition; on host side, the shadow variable is emitted with internal linkage. This is fine for the default -fno-gpu-rdc mode. However, this does not work for -fgpu-rdc mode. For -fgpu-rdc mode, both device variable and shadow variable should be declarations. Since this issue is orthogonal to __managed__, I will create a separate patch to fix it.

[yaxunl]

revised by Artem's and Aaron's comments.

[tra]

LGTM.

[aaron.ballman]

The attribute parts LGTM aside from some small nits.

clang/include/clang/Basic/DiagnosticSemaKinds.td

8250–8251

Since we're modifying this line anyway, can you wrap it for the 80-col limit?

clang/test/SemaCUDA/managed-var.cu

46

I think you're missing tests that check that the new managed attribute accepts no arguments and that it doesn't apply to things other than variables (like a function declaration).

[yaxunl]

clang/include/clang/Basic/DiagnosticSemaKinds.td

8250–8251

will fix when committing

clang/test/SemaCUDA/managed-var.cu

46

will add tests when committing

clang/include/clang/Basic/Attr.td

class LangOpt<string name, code customCode = [{}]> {
 
   // A custom predicate, written as an expression evaluated in a context with
   // "LangOpts" bound.
   code CustomCode = 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">;
 def OpenCL : LangOpt<"OpenCL">;
 def RenderScript : LangOpt<"RenderScript">;
 def ObjC : LangOpt<"ObjC">;
 def BlocksSupported : LangOpt<"Blocks">;
 def ObjCAutoRefCount : LangOpt<"ObjCAutoRefCount">;

 
 def CUDAHost : InheritableAttr {
   let Spellings = [GNU<"host">, Declspec<"__host__">];
   let Subjects = SubjectList<[Function]>;
   let LangOpts = [CUDA];
   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]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
 }
 
 def CUDALaunchBounds : InheritableAttr {

clang/include/clang/Basic/AttrDocs.td

 The ``device_builtin_texture_type`` attribute can be applied to a class
 template when declaring the texture reference. A texture reference variable
 could be accessed on the host side and, on the device side, might be translated
 into an internal texture object, which is established through texture bind and
 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 = [{
 .. Note:: This attribute is experimental and its effect on analysis is subject to change in
   a future version of clang.
 
 The attribute ``[[gsl::Owner(T)]]`` applies to structs and classes that own an
 object of type ``T``:

clang/include/clang/Basic/DiagnosticSemaKinds.td

def err_cuda_unattributed_constexpr_cannot_overload_device : Error<
   "attribute, or build with -fno-cuda-host-device-constexpr.">;
 def note_cuda_conflicting_device_function_declared_here : Note<
   "conflicting __device__ function declared here">;
 def err_cuda_device_exceptions : Error<
   "cannot use '%0' in "
   "%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<
     "cannot use variable-length arrays in "
     "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
 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">;
 def note_cuda_ovl_candidate_target_mismatch : Note<
     "candidate template ignored: target attributes do not match">;
 
 def err_cuda_device_builtin_surftex_cls_template : Error<
     "illegal device builtin %select{surface|texture}0 reference "

clang/lib/CodeGen/CGCUDANV.cpp

 #include "CodeGenModule.h"
 #include "clang/AST/Decl.h"
 #include "clang/Basic/Cuda.h"
 #include "clang/CodeGen/CodeGenABITypes.h"
 #include "clang/CodeGen/ConstantInitBuilder.h"
 #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;
 using namespace CodeGen;
 
 namespace {
 constexpr unsigned CudaFatMagic = 0x466243b1;
 constexpr unsigned HIPFatMagic = 0x48495046; // "HIPF"
public:
   CGNVCUDARuntime(CodeGenModule &CGM);
 
   void emitDeviceStub(CodeGenFunction &CGF, FunctionArgList &Args) override;
   void registerDeviceVar(const VarDecl *VD, llvm::GlobalVariable &Var,
                          bool Extern, bool Constant) override {
     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,
                          bool Extern, int Type, bool Normalized) override {
     DeviceVars.push_back({&Var,
                           VD,
                           {DeviceVarFlags::Texture, Extern, /*Constant*/ false,
-                           Normalized, Type}});
+                           /*Managed*/ false, Normalized, Type}});
   }
 
   /// Creates module constructor function
   llvm::Function *makeModuleCtorFunction() override;
   /// Creates module destructor function
   llvm::Function *makeModuleDtorFunction() override;
 };
 
void CGNVCUDARuntime::emitDeviceStubBodyLegacy(CodeGenFunction &CGF,
   llvm::FunctionCallee cudaLaunchFn = getLaunchFn();
   llvm::Value *Arg = CGF.Builder.CreatePointerCast(CGF.CurFn, CharPtrTy);
   CGF.EmitRuntimeCallOrInvoke(cudaLaunchFn, Arg);
   CGF.EmitBranch(EndBlock);
 
   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
 /// runtime.
 /// \code
 /// void __cuda_register_globals(void** GpuBinaryHandle) {
 ///    __cudaRegisterFunction(GpuBinaryHandle,Kernel0,...);
 ///    ...
llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
   // void __cudaRegisterVar(void **, char *, char *, const char *,
   //                        int, int, int, int)
   llvm::Type *RegisterVarParams[] = {VoidPtrPtrTy, CharPtrTy, CharPtrTy,
                                      CharPtrTy,    IntTy,     VarSizeTy,
                                      IntTy,        IntTy};
   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(
       llvm::FunctionType::get(
           VoidTy, {VoidPtrPtrTy, VoidPtrTy, CharPtrTy, CharPtrTy, IntTy, IntTy},
           false),
       addUnderscoredPrefixToName("RegisterSurface"));
   // void __cudaRegisterTexture(void **, const struct textureReference *,
   //                            const void **, const char *, int, int, int)
   llvm::FunctionCallee RegisterTex = CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(
           VoidTy,
           {VoidPtrPtrTy, VoidPtrTy, CharPtrTy, CharPtrTy, IntTy, IntTy, IntTy},
           false),
       addUnderscoredPrefixToName("RegisterTexture"));
   for (auto &&Info : DeviceVars) {
     llvm::GlobalVariable *Var = Info.Var;
     llvm::Constant *VarName = makeConstantString(getDeviceSideName(Info.D));
     switch (Info.Flags.getKind()) {
     case DeviceVarFlags::Variable: {
       uint64_t VarSize =
           CGM.getDataLayout().getTypeAllocSize(Var->getValueType());
+      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:
       Builder.CreateCall(
           RegisterSurf,
           {&GpuBinaryHandlePtr, Builder.CreateBitCast(Var, VoidPtrTy), VarName,
            VarName, llvm::ConstantInt::get(IntTy, Info.Flags.getSurfTexType()),
            llvm::ConstantInt::get(IntTy, Info.Flags.isExtern())});

clang/lib/CodeGen/CGCUDARuntime.h

enum DeviceVarKind {
       Surface,  // Builtin surface
       Texture,  // Builtin texture
     };
 
   private:
     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; }
   };
 
   CGCUDARuntime(CodeGenModule &CGM) : CGM(CGM) {}
   virtual ~CGCUDARuntime();
 
   virtual RValue EmitCUDAKernelCallExpr(CodeGenFunction &CGF,

clang/lib/CodeGen/CodeGenModule.cpp

void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
   // CUDA E.2.4.1 "__shared__ variables cannot have an initialization
   // as part of their declaration."  Sema has already checked for
   // error cases, so we just need to set Init to UndefValue.
   bool IsCUDASharedVar =
       getLangOpts().CUDAIsDevice && D->hasAttr<CUDASharedAttr>();
   // 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 &&
       (IsCUDASharedVar || IsCUDAShadowVar || IsCUDADeviceShadowVar))
     Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
   else if (D->hasAttr<LoaderUninitializedAttr>())
     Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
   else if (!InitExpr) {

clang/lib/Sema/SemaDeclAttr.cpp

 }
 
 static void handleOptimizeNoneAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   if (OptimizeNoneAttr *Optnone = S.mergeOptimizeNoneAttr(D, AL))
     D->addAttr(Optnone);
 }
 
 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()) {
     S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
     return;
   }
   D->addAttr(::new (S.Context) CUDAConstantAttr(S.Context, 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.
   // "int x[]").
   if (!S.getLangOpts().GPURelocatableDeviceCode && VD->hasExternalStorage() &&
       !isa<IncompleteArrayType>(VD->getType())) {
     S.Diag(AL.getLoc(), diag::err_cuda_extern_shared) << VD;
     return;
static void handleDeviceAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   }
 
   if (const auto *VD = dyn_cast<VarDecl>(D)) {
     if (VD->hasLocalStorage()) {
       S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
       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()) {
     S.Diag(AL.getLoc(), diag::warn_gnu_inline_attribute_requires_inline);
     return;
   }
 
   if (S.LangOpts.CPlusPlus && Fn->getStorageClass() != SC_Extern)
case ParsedAttr::AT_CUDAGlobal:
     handleGlobalAttr(S, D, AL);
     break;
   case ParsedAttr::AT_CUDADevice:
     handleDeviceAttr(S, D, AL);
     break;
   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,
                                                                           AL);
     break;
   case ParsedAttr::AT_CUDADeviceBuiltinTextureType:
     handleSimpleAttributeWithExclusions<CUDADeviceBuiltinTextureTypeAttr,
                                         CUDADeviceBuiltinSurfaceTypeAttr>(S, D,

clang/test/AST/Inputs/cuda.h

 /* 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__

clang/test/AST/ast-dump-managed-var.cu

+// 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;

clang/test/CodeGenCUDA/Inputs/cuda.h

 /* Minimal declarations for CUDA support.  Testing purposes only. */
 
 #include <stddef.h>
 
 #define __constant__ __attribute__((constant))
 #define __device__ __attribute__((device))
 #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 {
   unsigned x, y, z;
   __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}
 };
 
 #ifdef __HIP__

clang/test/CodeGenCUDA/managed-var.cu

+// 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)

clang/test/Misc/pragma-attribute-supported-attributes-list.test

 // CHECK-NEXT: EnforceTCB (SubjectMatchRule_function)
 // CHECK-NEXT: EnforceTCBLeaf (SubjectMatchRule_function)
 // CHECK-NEXT: EnumExtensibility (SubjectMatchRule_enum)
 // CHECK-NEXT: ExcludeFromExplicitInstantiation (SubjectMatchRule_variable, SubjectMatchRule_function, SubjectMatchRule_record)
 // CHECK-NEXT: ExternalSourceSymbol ((SubjectMatchRule_record, SubjectMatchRule_enum, SubjectMatchRule_enum_constant, SubjectMatchRule_field, SubjectMatchRule_function, SubjectMatchRule_namespace, SubjectMatchRule_objc_category, SubjectMatchRule_objc_implementation, SubjectMatchRule_objc_interface, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property, SubjectMatchRule_objc_protocol, SubjectMatchRule_record, SubjectMatchRule_type_alias, SubjectMatchRule_variable))
 // 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)
 // CHECK-NEXT: InitPriority (SubjectMatchRule_variable)
 // CHECK-NEXT: InternalLinkage (SubjectMatchRule_variable, SubjectMatchRule_function, SubjectMatchRule_record)
 // CHECK-NEXT: LTOVisibilityPublic (SubjectMatchRule_record)
 // CHECK-NEXT: Leaf (SubjectMatchRule_function)
 // CHECK-NEXT: LoaderUninitialized (SubjectMatchRule_variable_is_global)

clang/test/SemaCUDA/Inputs/cuda.h

 /* 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__

clang/test/SemaCUDA/bad-attributes.cu

 
 __global__ static inline void foobar() {};
 #ifdef EXPECT_INLINE_WARNING
 // expected-warning@-2 {{ignored 'inline' attribute on kernel function 'foobar'}}
 #endif
 
 __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}}
 typedef __attribute__((device_builtin_texture_type)) unsigned long long t0_ty; // expected-warning {{'device_builtin_texture_type' attribute only applies to classes}}
 
 struct __attribute__((device_builtin_surface_type)) s1_ref {}; // expected-error {{illegal device builtin surface reference type 's1_ref' declared here}}
 // expected-note@-1 {{'s1_ref' needs to be instantiated from a class template with proper template arguments}}
 struct __attribute__((device_builtin_texture_type)) t1_ref {}; // expected-error {{illegal device builtin texture reference type 't1_ref' declared here}}

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

 
 // Use the types we share with CodeGen tests.
 #include "Inputs/cuda-initializers.h"
 
 __shared__ int s_v_i = 1;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 
 __device__ int d_v_f = f();
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ int s_v_f = f();
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ int c_v_f = f();
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __shared__ T s_t_i = {2};
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __device__ T d_t_i = {2};
 __constant__ T c_t_i = {2};
 
 __device__ ECD d_ecd_i{};
 __shared__ ECD s_ecd_i{};
 __constant__ ECD c_ecd_i{};
 
 __device__ EC d_ec_i(3);
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ EC s_ec_i(3);
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ EC c_ec_i(3);
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ EC d_ec_i2 = {3};
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ EC s_ec_i2 = {3};
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ EC c_ec_i2 = {3};
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ ETC d_etc_i(3);
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ ETC s_etc_i(3);
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ ETC c_etc_i(3);
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ ETC d_etc_i2 = {3};
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ ETC s_etc_i2 = {3};
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ ETC c_etc_i2 = {3};
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ UC d_uc;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ UC s_uc;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ UC c_uc;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ UD d_ud;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ UD s_ud;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ UD c_ud;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ ECI d_eci;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ ECI s_eci;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ ECI c_eci;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ NEC d_nec;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ NEC s_nec;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NEC c_nec;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ NED d_ned;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ NED s_ned;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NED c_ned;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ NCV d_ncv;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ NCV s_ncv;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NCV c_ncv;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ VD d_vd;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ VD s_vd;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ VD c_vd;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ NCF d_ncf;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ NCF s_ncf;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NCF c_ncf;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __shared__ NCFS s_ncfs;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 
 __device__ UTC d_utc;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ UTC s_utc;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ UTC c_utc;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ UTC d_utc_i(3);
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ UTC s_utc_i(3);
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ UTC c_utc_i(3);
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ NETC d_netc;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ NETC s_netc;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NETC c_netc;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ NETC d_netc_i(3);
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ NETC s_netc_i(3);
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NETC c_netc_i(3);
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ EC_I_EC1 d_ec_i_ec1;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ EC_I_EC1 s_ec_i_ec1;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ EC_I_EC1 c_ec_i_ec1;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ T_V_T d_t_v_t;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ T_V_T s_t_v_t;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_V_T c_t_v_t;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ T_B_NEC d_t_b_nec;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ T_B_NEC s_t_b_nec;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_B_NEC c_t_b_nec;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ T_F_NEC d_t_f_nec;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ T_F_NEC s_t_f_nec;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_F_NEC c_t_f_nec;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ T_FA_NEC d_t_fa_nec;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ T_FA_NEC s_t_fa_nec;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_FA_NEC c_t_fa_nec;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ T_B_NED d_t_b_ned;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ T_B_NED s_t_b_ned;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_B_NED c_t_b_ned;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ T_F_NED d_t_f_ned;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ T_F_NED s_t_f_ned;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_F_NED c_t_f_ned;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ T_FA_NED d_t_fa_ned;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __shared__ T_FA_NED s_t_fa_ned;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_FA_NED c_t_fa_ned;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-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.
 // __shared__ can't be initialized at all and others don't support dynamic initialization.
 __device__ void df_sema() {
   static __device__ int ds;
   static __constant__ int dc;
   static int v;
   static const int cv = 1;
   static const __device__ int cds = 1;
   static const __constant__ int cdc = 1;
 
 
   // __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@-1 {{__constant__ and __device__ are not allowed on non-static local variables}}
+  // expected-error@-1 {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}}
   __device__ int ldi;
-  // expected-error@-1 {{__constant__ and __device__ are not allowed on non-static local variables}}
+  // expected-error@-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@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ int s_v_f = f();
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ int c_v_f = f();
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __shared__ T s_t_i = {2};
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __device__ T d_t_i = {2};
   static __constant__ T c_t_i = {2};
 
   static __device__ ECD d_ecd_i;
   static __shared__ ECD s_ecd_i;
   static __constant__ ECD c_ecd_i;
 
   static __device__ EC d_ec_i(3);
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ EC s_ec_i(3);
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ EC c_ec_i(3);
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ EC d_ec_i2 = {3};
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ EC s_ec_i2 = {3};
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ EC c_ec_i2 = {3};
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ ETC d_etc_i(3);
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ ETC s_etc_i(3);
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ ETC c_etc_i(3);
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ ETC d_etc_i2 = {3};
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ ETC s_etc_i2 = {3};
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ ETC c_etc_i2 = {3};
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ UC d_uc;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ UC s_uc;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ UC c_uc;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ UD d_ud;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ UD s_ud;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ UD c_ud;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ ECI d_eci;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ ECI s_eci;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ ECI c_eci;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ NEC d_nec;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ NEC s_nec;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ NEC c_nec;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ NED d_ned;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ NED s_ned;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ NED c_ned;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ NCV d_ncv;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ NCV s_ncv;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ NCV c_ncv;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ VD d_vd;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ VD s_vd;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ VD c_vd;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ NCF d_ncf;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ NCF s_ncf;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ NCF c_ncf;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __shared__ NCFS s_ncfs;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 
   static __device__ UTC d_utc;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ UTC s_utc;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ UTC c_utc;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ UTC d_utc_i(3);
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ UTC s_utc_i(3);
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ UTC c_utc_i(3);
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ NETC d_netc;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ NETC s_netc;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ NETC c_netc;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ NETC d_netc_i(3);
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ NETC s_netc_i(3);
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ NETC c_netc_i(3);
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ EC_I_EC1 d_ec_i_ec1;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ EC_I_EC1 s_ec_i_ec1;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ EC_I_EC1 c_ec_i_ec1;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ T_V_T d_t_v_t;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ T_V_T s_t_v_t;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ T_V_T c_t_v_t;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ T_B_NEC d_t_b_nec;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ T_B_NEC s_t_b_nec;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ T_B_NEC c_t_b_nec;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ T_F_NEC d_t_f_nec;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ T_F_NEC s_t_f_nec;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ T_F_NEC c_t_f_nec;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ T_FA_NEC d_t_fa_nec;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ T_FA_NEC s_t_fa_nec;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ T_FA_NEC c_t_fa_nec;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ T_B_NED d_t_b_ned;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ T_B_NED s_t_b_ned;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ T_B_NED c_t_b_ned;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ T_F_NED d_t_f_ned;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ T_F_NED s_t_f_ned;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ T_F_NED c_t_f_ned;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
   static __device__ T_FA_NED d_t_fa_ned;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
   static __shared__ T_FA_NED s_t_fa_ned;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __constant__ T_FA_NED c_t_fa_ned;
-  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 }
 
 __host__ __device__ void hd_sema() {
   static int x = 42;
 }
 
 inline __host__ __device__ void hd_emitted_host_only() {
   static int x = 42; // no error on device because this is never codegen'ed there.

clang/test/SemaCUDA/function-overload.cu

 // Verify that function overloading doesn't prune candidate wrongly.
 int test_constexpr_overload(C2 &x, C2 &y) {
   return constexpr_overload(x, y);
 }
 
 // Verify no ambiguity for new operator.
 void *a = new int;
 __device__ void *b = new int;
-// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 // Verify no ambiguity for new operator.
 template<typename _Tp> _Tp&& f();
 template<typename _Tp, typename = decltype(new _Tp(f<_Tp>()))>
 void __test();
 
 void foo() {
   __test<int>();

clang/test/SemaCUDA/managed-var.cu

+// 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@-1 {{'constant' and 'managed' attributes are not compatible}}
+// expected-note@-2 {{conflicting attribute is here}}
+
+__constant__ __managed__ int m4a;
+// expected-error@-1 {{'managed' and 'constant' attributes are not compatible}}
+// expected-note@-2 {{conflicting attribute is here}}
+
+__managed__ __shared__ int m5;
+// expected-error@-1 {{'shared' and 'managed' attributes are not compatible}}
+// expected-note@-2 {{conflicting attribute is here}}
+
+__shared__ __managed__ int m5a;
+// expected-error@-1 {{'managed' and 'shared' attributes are not compatible}}
+// expected-note@-2 {{conflicting attribute is here}}
+
+__managed__ __global__ int m6;
+// expected-warning@-1 {{'global' attribute only applies to functions}}
+
+void func() {
+  __managed__ int m7;
+  // expected-error@-1 {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}}
+}
+
+__attribute__((managed(1))) int m8;
+// expected-error@-1 {{'managed' attribute takes no arguments}}
+
+__managed__ void func2() {}
+// expected-warning@-1 {{'managed' attribute only applies to variables}}
+
+typedef __managed__ int managed_int;
+// expected-warning@-1 {{'managed' attribute only applies to variables}}
+
+__managed__ A a;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}

clang/test/SemaCUDA/union-init.cu

 union D {
   A a;
   __device__ D() { }
   __device__ ~D() { a.a = 4; }
 };
 
 __device__ B b;
 __device__ C c;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 __device__ D d;
-// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}}
 
 __device__ void foo() {
   __shared__ B b;
   __shared__ C c;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   __shared__ D d;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 }

llvm/include/llvm/IR/ReplaceConstant.h

+//===- 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

llvm/lib/IR/CMakeLists.txt

add_llvm_component_library(LLVMCore
   PassInstrumentation.cpp
   PassManager.cpp
   PassRegistry.cpp
   PassTimingInfo.cpp
   PrintPasses.cpp
   SafepointIRVerifier.cpp
   ProfileSummary.cpp
   PseudoProbe.cpp
+  ReplaceConstant.cpp
   Statepoint.cpp
   StructuralHash.cpp
   Type.cpp
   TypeFinder.cpp
   Use.cpp
   User.cpp
   Value.cpp
   ValueSymbolTable.cpp

llvm/lib/IR/ReplaceConstant.cpp

+//===- 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

llvm/lib/Target/XCore/XCoreLowerThreadLocal.cpp

 #include "llvm/IR/Constants.h"
 #include "llvm/IR/DerivedTypes.h"
 #include "llvm/IR/GlobalVariable.h"
 #include "llvm/IR/IRBuilder.h"
 #include "llvm/IR/Intrinsics.h"
 #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"
 #include "llvm/Transforms/Utils/BasicBlockUtils.h"
 
 #define DEBUG_TYPE "xcore-lower-thread-local"
 
 using namespace llvm;

 createLoweredInitializer(ArrayType *NewType, Constant *OriginalInitializer) {
   SmallVector<Constant *, 8> Elements(MaxThreads);
   for (unsigned i = 0; i != MaxThreads; ++i) {
     Elements[i] = 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 {
     SmallVector<WeakTrackingVH, 8> WUsers(CE->users());
     llvm::sort(WUsers);
     WUsers.erase(std::unique(WUsers.begin(), WUsers.end()), WUsers.end());
     while (!WUsers.empty())
       if (WeakTrackingVH WU = WUsers.pop_back_val()) {