[HIP] Fix managed variable linkage

[yaxunl]

Summary

Currently managed variables are emitted as undefined symbols, which
causes difficulty for diagnosing undefined symbols for non-managed
variables.

This patch transforms managed variables in device compilation so that
they can be emitted as normal variables.

[yaxunl]

ping

[yaxunl]

keep managed var in llvm.compiler.used since they need runtime handling even if they are not used in device code.

[tra]

clang/lib/CodeGen/CGCUDANV.cpp

1019

A comment about how exactly we're transforming the vars would be helpful.

clang/lib/CodeGen/CGCUDARuntime.h

103

Can we fold it into makeModuleCtorFunction ? It does not look like something we would need to customize independently of other glue generating machinery.

clang/lib/CodeGen/CodeGenModule.cpp

463–466

This does not look right.
Was transformManagedVars() supposed to *replace* AddGlobalCtor(CudaCtorFunction)?

I would expect it to be in addition to the ctor we create now.

if (Context.getLangOpts().CUDAIsDevice)
  CUDARuntime->transformManagedVars();
if (llvm::Function *CudaCtorFunction =
             CUDARuntime->makeModuleCtorFunction())
     AddGlobalCtor(CudaCtorFunction);

[yaxunl]

clang/lib/CodeGen/CGCUDANV.cpp

1019

done

clang/lib/CodeGen/CGCUDARuntime.h

103

makeModuleCtorFunction is creating module ctor function, but transformManagedVars has nothing to do with module ctor function.

Also, makeModuleCtorFunction is called in host compilation only.

Do we want to rename the API here finalizeModule ? And let it call makeModuleCtorFunction and transformManagedVars as static functions in CGCUDANV.cpp ?

clang/lib/CodeGen/CodeGenModule.cpp

463–466

originally, makeModuleCtorFunction is called for host compilation only.

[yaxunl]

refactor CGCUDARuntime interface not to expose transformManagedVar.

[tra]

clang/lib/CodeGen/CGCUDARuntime.h

103

SGTM.

clang/lib/CodeGen/CGCUDANV.cpp

private:
   void registerDeviceTex(const VarDecl *VD, llvm::GlobalVariable &Var,
                          bool Extern, int Type, bool Normalized) {
     DeviceVars.push_back({&Var,
                           VD,
                           {DeviceVarFlags::Texture, Extern, /*Constant*/ false,
                            /*Managed*/ false, Normalized, Type}});
   }
 
+  /// Creates module constructor function
+  llvm::Function *makeModuleCtorFunction();
+  /// Creates module destructor function
+  llvm::Function *makeModuleDtorFunction();
+  /// Transform managed variables for device compilation.
+  void transformManagedVars();
+
 public:
   CGNVCUDARuntime(CodeGenModule &CGM);
 
   void emitDeviceStub(CodeGenFunction &CGF, FunctionArgList &Args) override;
   void handleVarRegistration(const VarDecl *VD,
                              llvm::GlobalVariable &Var) override;
-
-  /// Creates module constructor function
-  llvm::Function *makeModuleCtorFunction() override;
-  /// Creates module destructor function
-  llvm::Function *makeModuleDtorFunction() override;
   void
   internalizeDeviceSideVar(const VarDecl *D,
                            llvm::GlobalValue::LinkageTypes &Linkage) override;
+
+  llvm::Function *finalizeModule() override;
 };
 
 }
 
 std::string CGNVCUDARuntime::addPrefixToName(StringRef FuncName) const {
   if (CGM.getLangOpts().HIP)
     return ((Twine("hip") + Twine(FuncName)).str());
   return ((Twine("cuda") + Twine(FuncName)).str());
llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
   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;
+    assert((!Var->isDeclaration() || Info.Flags.isManaged()) &&
+           "External variables should not show up here, except HIP managed "
+           "variables");
     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,
+            /*Init=*/Var->isDeclaration()
+                ? nullptr
+                : llvm::ConstantPointerNull::get(Var->getType()),
+            /*Name=*/"", /*InsertBefore=*/nullptr,
             llvm::GlobalVariable::NotThreadLocal);
         ManagedVar->setDSOLocal(Var->isDSOLocal());
         ManagedVar->setVisibility(Var->getVisibility());
+        ManagedVar->setExternallyInitialized(true);
+        ManagedVar->takeName(Var);
+        Var->setName(Twine(ManagedVar->getName() + ".managed"));
         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())};
+        if (!Var->isDeclaration())
           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),
if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) {
     // Shadow variables and their properties must be registered with CUDA
     // runtime. Skip Extern global variables, which will be registered in
     // the TU where they are defined.
     //
     // Don't register a C++17 inline variable. The local symbol can be
     // discarded and referencing a discarded local symbol from outside the
     // comdat (__cuda_register_globals) is disallowed by the ELF spec.
     // TODO: Reject __device__ constexpr and __device__ inline in Sema.
-    if (!D->hasExternalStorage() && !D->isInline())
+    // HIP managed variables need to be always recorded in device and host
+    // compilations for transformation.
+    if ((!D->hasExternalStorage() && !D->isInline()) ||
+        D->hasAttr<HIPManagedAttr>()) {
       registerDeviceVar(D, GV, !D->hasDefinition(),
                         D->hasAttr<CUDAConstantAttr>());
+    }
   } else if (D->getType()->isCUDADeviceBuiltinSurfaceType() ||
              D->getType()->isCUDADeviceBuiltinTextureType()) {
     // Builtin surfaces and textures and their template arguments are
     // also registered with CUDA runtime.
     const ClassTemplateSpecializationDecl *TD =
         cast<ClassTemplateSpecializationDecl>(
             D->getType()->getAs<RecordType>()->getDecl());
     const TemplateArgumentList &Args = TD->getTemplateArgs();
if (TD->hasAttr<CUDADeviceBuiltinSurfaceTypeAttr>()) {
       auto TexType = Args[1].getAsIntegral();
       auto Normalized = Args[2].getAsIntegral();
       if (!D->hasExternalStorage())
         registerDeviceTex(D, GV, !D->hasDefinition(), TexType.getSExtValue(),
                           Normalized.getZExtValue());
     }
   }
 }
+
+// Transform managed variables to pointers to managed variables in device code.
+// Each use of the original managed variable is replaced by a load from the
+// transformed managed variable. The transformed managed variable contains
+// the address of managed memory which will be allocated by the runtime.
+void CGNVCUDARuntime::transformManagedVars() {
+  for (auto &&Info : DeviceVars) {
+    llvm::GlobalVariable *Var = Info.Var;
+    if (Info.Flags.getKind() == DeviceVarFlags::Variable &&
+        Info.Flags.isManaged()) {
+      auto ManagedVar = new llvm::GlobalVariable(
+          CGM.getModule(), Var->getType(),
+          /*isConstant=*/false, Var->getLinkage(),
+          /*Init=*/Var->isDeclaration()
+              ? nullptr
+              : llvm::ConstantPointerNull::get(Var->getType()),
+          /*Name=*/"", /*InsertBefore=*/nullptr,
+          llvm::GlobalVariable::NotThreadLocal,
+          CGM.getContext().getTargetAddressSpace(LangAS::cuda_device));
+      ManagedVar->setDSOLocal(Var->isDSOLocal());
+      ManagedVar->setVisibility(Var->getVisibility());
+      ManagedVar->setExternallyInitialized(true);
+      replaceManagedVar(Var, ManagedVar);
+      ManagedVar->takeName(Var);
+      Var->setName(Twine(ManagedVar->getName()) + ".managed");
+      // Keep managed variables even if they are not used in device code since
+      // they need to be allocated by the runtime.
+      if (!Var->isDeclaration()) {
+        assert(!ManagedVar->isDeclaration());
+        CGM.addCompilerUsedGlobal(Var);
+        CGM.addCompilerUsedGlobal(ManagedVar);
+      }
+    }
+  }
+}
+
+// Returns module constructor to be added.
+llvm::Function *CGNVCUDARuntime::finalizeModule() {
+  if (CGM.getLangOpts().CUDAIsDevice) {
+    transformManagedVars();
+    return nullptr;
+  }
+  return makeModuleCtorFunction();
+}

clang/lib/CodeGen/CGCUDARuntime.h

public:
 
   /// Emits a kernel launch stub.
   virtual void emitDeviceStub(CodeGenFunction &CGF, FunctionArgList &Args) = 0;
 
   /// Check whether a variable is a device variable and register it if true.
   virtual void handleVarRegistration(const VarDecl *VD,
                                      llvm::GlobalVariable &Var) = 0;
 
-  /// Constructs and returns a module initialization function or nullptr if it's
-  /// not needed. Must be called after all kernels have been emitted.
-  virtual llvm::Function *makeModuleCtorFunction() = 0;
-
-  /// Returns a module cleanup function or nullptr if it's not needed.
-  /// Must be called after ModuleCtorFunction
-  virtual llvm::Function *makeModuleDtorFunction() = 0;
+  /// Finalize generated LLVM module. Returns a module constructor function
+  /// to be added or a null pointer.
+  virtual llvm::Function *finalizeModule() = 0;
 
   /// Returns function or variable name on device side even if the current
   /// compilation is for host.
   virtual std::string getDeviceSideName(const NamedDecl *ND) = 0;
 
   /// Adjust linkage of shadow variables in host compilation.
   virtual void
   internalizeDeviceSideVar(const VarDecl *D,
                            llvm::GlobalValue::LinkageTypes &Linkage) = 0;
 };
 
 /// Creates an instance of a CUDA runtime class.
 CGCUDARuntime *CreateNVCUDARuntime(CodeGenModule &CGM);
 
 }
 }
 
 #endif

clang/lib/CodeGen/CodeGenModule.cpp

void CodeGenModule::Release() {
   emitMultiVersionFunctions();
   EmitCXXGlobalInitFunc();
   EmitCXXGlobalCleanUpFunc();
   registerGlobalDtorsWithAtExit();
   EmitCXXThreadLocalInitFunc();
   if (ObjCRuntime)
     if (llvm::Function *ObjCInitFunction = ObjCRuntime->ModuleInitFunction())
       AddGlobalCtor(ObjCInitFunction);
-  if (Context.getLangOpts().CUDA && !Context.getLangOpts().CUDAIsDevice &&
-      CUDARuntime) {
-    if (llvm::Function *CudaCtorFunction =
-            CUDARuntime->makeModuleCtorFunction())
+  if (Context.getLangOpts().CUDA && CUDARuntime) {
+    if (llvm::Function *CudaCtorFunction = CUDARuntime->finalizeModule())
       AddGlobalCtor(CudaCtorFunction);
   }
   if (OpenMPRuntime) {
     if (llvm::Function *OpenMPRequiresDirectiveRegFun =
             OpenMPRuntime->emitRequiresDirectiveRegFun()) {
       AddGlobalCtor(OpenMPRequiresDirectiveRegFun, 0);
     }
     OpenMPRuntime->createOffloadEntriesAndInfoMetadata();
     OpenMPRuntime->clear();
   }
   if (PGOReader) {
if (Context.getLangOpts().CPlusPlus && GV->hasExternalLinkage() &&
             }
             emitter.finalize(GV);
           }
         }
       }
     }
   }
 
-  if (GV->isDeclaration())
+  if (GV->isDeclaration()) {
     getTargetCodeGenInfo().setTargetAttributes(D, GV, *this);
+    // External HIP managed variables needed to be recorded for transformation
+    // in both device and host compilations.
+    if (getLangOpts().CUDA && D && D->hasAttr<HIPManagedAttr>() &&
+        D->hasExternalStorage())
+      getCUDARuntime().handleVarRegistration(D, *GV);
+  }
 
   LangAS ExpectedAS =
       D ? D->getType().getAddressSpace()
         : (LangOpts.OpenCL ? LangAS::opencl_global : LangAS::Default);
   assert(getContext().getTargetAddressSpace(ExpectedAS) ==
          Ty->getPointerAddressSpace());
   if (AddrSpace != ExpectedAS)
     return getTargetCodeGenInfo().performAddrSpaceCast(*this, GV, AddrSpace,
if (LangOpts.OpenMPIsDevice && OpenMPRuntime &&
       OpenMPRuntime->emitTargetGlobalVariable(D))
     return;
 
   llvm::Constant *Init = nullptr;
   bool NeedsGlobalCtor = false;
   bool NeedsGlobalDtor =
       D->needsDestruction(getContext()) == QualType::DK_cxx_destructor;
 
-  bool IsHIPManagedVarOnDevice =
-      getLangOpts().CUDAIsDevice && D->hasAttr<HIPManagedAttr>();
-
   const VarDecl *InitDecl;
-  const Expr *InitExpr =
-      IsHIPManagedVarOnDevice ? nullptr : D->getAnyInitializer(InitDecl);
+  const Expr *InitExpr = D->getAnyInitializer(InitDecl);
 
   Optional<ConstantEmitter> emitter;
 
   // 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.
+  // Managed Variables should be initialized on both host side and device side.
   bool IsCUDAShadowVar =
       !getLangOpts().CUDAIsDevice && !D->hasAttr<HIPManagedAttr>() &&
       (D->hasAttr<CUDAConstantAttr>() || D->hasAttr<CUDADeviceAttr>() ||
        D->hasAttr<CUDASharedAttr>());
   bool IsCUDADeviceShadowVar =
-      getLangOpts().CUDAIsDevice &&
+      getLangOpts().CUDAIsDevice && !D->hasAttr<HIPManagedAttr>() &&
       (D->getType()->isCUDADeviceBuiltinSurfaceType() ||
-       D->getType()->isCUDADeviceBuiltinTextureType() ||
-       D->hasAttr<HIPManagedAttr>());
+       D->getType()->isCUDADeviceBuiltinTextureType());
   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) {
     // This is a tentative definition; tentative definitions are
     // implicitly initialized with { 0 }.
void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
   // / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol())."
   if (GV && LangOpts.CUDA) {
     if (LangOpts.CUDAIsDevice) {
       if (Linkage != llvm::GlobalValue::InternalLinkage &&
           (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()))
         GV->setExternallyInitialized(true);
     } else {
       getCUDARuntime().internalizeDeviceSideVar(D, Linkage);
-      getCUDARuntime().handleVarRegistration(D, *GV);
     }
+    getCUDARuntime().handleVarRegistration(D, *GV);
   }
 
-  // HIP managed variables need to be emitted as declarations in device
-  // compilation.
-  if (!IsHIPManagedVarOnDevice)
-    GV->setInitializer(Init);
   if (emitter)
     emitter->finalize(GV);
 
   // If it is safe to mark the global 'constant', do so now.
   GV->setConstant(!NeedsGlobalCtor && !NeedsGlobalDtor &&
                   isTypeConstant(D->getType(), true));
 
   // If it is in a read-only section, mark it 'constant'.

clang/test/CodeGenCUDA/device-var-linkage.cu

 // DEV-DAG: @v1 = dso_local addrspace(1) externally_initialized global i32 0
 // NORDC-H-DAG: @v1 = internal global i32 undef
 // RDC-H-DAG: @v1 = dso_local global i32 undef
 __device__ int v1;
 // DEV-DAG: @v2 = dso_local addrspace(4) externally_initialized global i32 0
 // NORDC-H-DAG: @v2 = internal global i32 undef
 // RDC-H-DAG: @v2 = dso_local global i32 undef
 __constant__ int v2;
-// DEV-DAG: @v3 = external addrspace(1) externally_initialized global i32
-// NORDC-H-DAG: @v3 = internal global i32 0
-// RDC-H-DAG: @v3 = dso_local global i32 0
+// DEV-DAG: @v3 = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
+// NORDC-H-DAG: @v3 = internal externally_initialized global i32* null
+// RDC-H-DAG: @v3 = dso_local externally_initialized global i32* null
 __managed__ int v3;
 
 // DEV-DAG: @ev1 = external addrspace(1) global i32
 // HOST-DAG: @ev1 = external global i32
 extern __device__ int ev1;
 // DEV-DAG: @ev2 = external addrspace(4) global i32
 // HOST-DAG: @ev2 = external global i32
 extern __constant__ int ev2;
-// DEV-DAG: @ev3 = external addrspace(1) global i32
-// HOST-DAG: @ev3 = external global i32
+// DEV-DAG: @ev3 = external addrspace(1) externally_initialized global i32 addrspace(1)*
+// HOST-DAG: @ev3 = external externally_initialized global i32*
 extern __managed__ int ev3;
 
 // NORDC-DAG: @_ZL3sv1 = dso_local addrspace(1) externally_initialized global i32 0
 // RDC-DAG: @_ZL3sv1 = internal addrspace(1) global i32 0
 // HOST-DAG: @_ZL3sv1 = internal global i32 undef
 static __device__ int sv1;
 // NORDC-DAG: @_ZL3sv2 = dso_local addrspace(4) externally_initialized global i32 0
 // RDC-DAG: @_ZL3sv2 = internal addrspace(4) global i32 0
 // HOST-DAG: @_ZL3sv2 = internal global i32 undef
 static __constant__ int sv2;
-// DEV-DAG: @_ZL3sv3 = external addrspace(1) externally_initialized global i32
-// HOST-DAG: @_ZL3sv3 = internal global i32 0
+// DEV-DAG: @_ZL3sv3 = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
+// HOST-DAG: @_ZL3sv3 = internal externally_initialized global i32* null
 static __managed__ int sv3;
 
 __device__ __host__ int work(int *x);
 
 __device__ __host__ int fun1() {
   return work(&ev1) + work(&ev2) + work(&ev3) + work(&sv1) + work(&sv2) + work(&sv3);
 }
 

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:   -check-prefixes=COMMON,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:   -check-prefixes=COMMON,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,NORDC %s
+// RUN:   -check-prefixes=COMMON,HOST,NORDC %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,RDC %s
+// RUN:   -check-prefixes=COMMON,HOST,RDC %s
 
 #include "Inputs/cuda.h"
 
-// DEV-DAG: @x = external addrspace(1) externally_initialized global i32
-// NORDC-DAG: @x = internal global i32 1
-// RDC-DAG: @x = dso_local global i32 1
-// NORDC-DAG: @x.managed = internal global i32* null
-// RDC-DAG: @x.managed = dso_local global i32* null
-// HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"x\00"
-
 struct vec {
   float x,y,z;
 };
 
+// DEV-DAG: @x.managed = dso_local addrspace(1) externally_initialized global i32 1, align 4
+// DEV-DAG: @x = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
+// NORDC-DAG: @x.managed = internal global i32 1
+// RDC-DAG: @x.managed = dso_local global i32 1
+// NORDC-DAG: @x = internal externally_initialized global i32* null
+// RDC-DAG: @x = dso_local externally_initialized global i32* null
+// HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"x\00"
 __managed__ int x = 1;
+
+// DEV-DAG: @v.managed = dso_local addrspace(1) externally_initialized global [100 x %struct.vec] zeroinitializer, align 4
+// DEV-DAG: @v = dso_local addrspace(1) externally_initialized global [100 x %struct.vec] addrspace(1)* null
 __managed__ vec v[100];
+
+// DEV-DAG: @v2.managed = dso_local addrspace(1) externally_initialized global <{ %struct.vec, [99 x %struct.vec] }> <{ %struct.vec { float 1.000000e+00, float 1.000000e+00, float 1.000000e+00 }, [99 x %struct.vec] zeroinitializer }>, align 4
+// DEV-DAG: @v2 = dso_local addrspace(1) externally_initialized global <{ %struct.vec, [99 x %struct.vec] }> addrspace(1)* null
 __managed__ vec v2[100] = {{1, 1, 1}};
 
-// DEV-DAG: @ex = external addrspace(1) global i32
-// HOST-DAG: @ex = external global i32
+// DEV-DAG: @ex.managed = external addrspace(1) global i32, align 4
+// DEV-DAG: @ex = external addrspace(1) externally_initialized global i32 addrspace(1)*
+// HOST-DAG: @ex.managed = external global i32
+// HOST-DAG: @ex = external externally_initialized global i32*
 extern __managed__ int ex;
 
-// DEV-DAG: @_ZL2sx = external addrspace(1) externally_initialized global i32
-// HOST-DAG: @_ZL2sx = internal global i32 1
-// HOST-DAG: @_ZL2sx.managed = internal global i32* null
+// DEV-DAG: @_ZL2sx.managed = dso_local addrspace(1) externally_initialized global i32 1, align 4
+// DEV-DAG: @_ZL2sx = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
+// HOST-DAG: @_ZL2sx.managed = internal global i32 1
+// HOST-DAG: @_ZL2sx = internal externally_initialized global i32* null
 static __managed__ int sx = 1;
 
-// HOST-NOT: @ex.managed
+// DEV-DAG: @llvm.compiler.used
+// DEV-SAME-DAG: @x.managed
+// DEV-SAME-DAG: @x
+// DEV-SAME-DAG: @v.managed
+// DEV-SAME-DAG: @v
+// DEV-SAME-DAG: @_ZL2sx.managed
+// DEV-SAME-DAG: @_ZL2sx
 
 // Force ex and sx mitted in device compilation.
 __global__ void foo(int *z) {
   *z = x + ex + sx;
   v[1].x = 2;
 }
 
 // Force ex and sx emitted in host compilatioin.
 int foo2() {
   return ex + sx;
 }
 
-// HOST-LABEL: define {{.*}}@_Z4loadv()
-// HOST:  %ld.managed = load i32*, i32** @x.managed, align 4
+// COMMON-LABEL: define {{.*}}@_Z4loadv()
+// DEV:  %ld.managed = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(1)* @x, align 4
+// DEV:  %0 = addrspacecast i32 addrspace(1)* %ld.managed to i32*
+// DEV:  %1 = load i32, i32* %0, align 4
+// DEV:  ret i32 %1
+// HOST:  %ld.managed = load i32*, i32** @x, align 4
 // HOST:  %0 = load i32, i32* %ld.managed, align 4
 // HOST:  ret i32 %0
-int load() {
+__device__ __host__ int load() {
   return x;
 }
 
-// HOST-LABEL: define {{.*}}@_Z5storev()
-// HOST:  %ld.managed = load i32*, i32** @x.managed, align 4
+// COMMON-LABEL: define {{.*}}@_Z5storev()
+// DEV:  %ld.managed = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(1)* @x, align 4
+// DEV:  %0 = addrspacecast i32 addrspace(1)* %ld.managed to i32*
+// DEV:  store i32 2, i32* %0, align 4
+// HOST:  %ld.managed = load i32*, i32** @x, align 4
 // HOST:  store i32 2, i32* %ld.managed, align 4
-void store() {
+__device__ __host__ void store() {
   x = 2;
 }
 
-// HOST-LABEL: define {{.*}}@_Z10addr_takenv()
-// HOST:  %ld.managed = load i32*, i32** @x.managed, align 4
+// COMMON-LABEL: define {{.*}}@_Z10addr_takenv()
+// DEV:  %0 = addrspacecast i32 addrspace(1)* %ld.managed to i32*
+// DEV:  store i32* %0, i32** %p.ascast, align 8
+// DEV:  %1 = load i32*, i32** %p.ascast, align 8
+// DEV:  store i32 3, i32* %1, align 4
+// HOST:  %ld.managed = load i32*, i32** @x, 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() {
+__device__ __host__ 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: %ld.managed = load [100 x %struct.vec]*, [100 x %struct.vec]** @v, 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() {
+__device__ __host__ 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:  %ld.managed = load <{ %struct.vec, [99 x %struct.vec] }>*, <{ %struct.vec, [99 x %struct.vec] }>** @v2, 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:  %ld.managed = load [100 x %struct.vec]*, [100 x %struct.vec]** @v, 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:  %ld.managed1 = load <{ %struct.vec, [99 x %struct.vec] }>*, <{ %struct.vec, [99 x %struct.vec] }>** @v2, 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: __hipRegisterManagedVar({{.*}}@_ZL2sx.managed {{.*}}@_ZL2sx
-// HOST-NOT: __hipRegisterManagedVar({{.*}}@ex.managed {{.*}}@ex
+// COMMON-LABEL: define {{.*}}@_Z5load4v()
+// DEV:  %ld.managed = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(1)* @ex, align 4
+// DEV:  %0 = addrspacecast i32 addrspace(1)* %ld.managed to i32*
+// DEV:  %1 = load i32, i32* %0, align 4
+// DEV:  ret i32 %1
+// HOST:  %ld.managed = load i32*, i32** @ex, align 4
+// HOST:  %0 = load i32, i32* %ld.managed, align 4
+// HOST:  ret i32 %0
+__device__ __host__ int load4() {
+  return ex;
+}
+
+// HOST-DAG: __hipRegisterManagedVar({{.*}}@x {{.*}}@x.managed {{.*}}@[[DEVNAMEX]]{{.*}}, i64 4, i32 4)
+// HOST-DAG: __hipRegisterManagedVar({{.*}}@_ZL2sx {{.*}}@_ZL2sx.managed
+// HOST-NOT: __hipRegisterManagedVar({{.*}}@ex {{.*}}@ex.managed
 // HOST-DAG: declare void @__hipRegisterManagedVar(i8**, i8*, i8*, i8*, i64, i32)

llvm/lib/IR/ReplaceConstant.cpp

Instruction *createReplacementInstr(ConstantExpr *CE, Instruction *Instr) {
   case Instruction::FPToSI:
   case Instruction::UIToFP:
   case Instruction::SIToFP:
   case Instruction::FPTrunc:
   case Instruction::FPExt:
   case Instruction::PtrToInt:
   case Instruction::IntToPtr:
   case Instruction::BitCast:
+  case Instruction::AddrSpaceCast:
     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