Index: clang/lib/CodeGen/CGCUDANV.cpp =================================================================== --- clang/lib/CodeGen/CGCUDANV.cpp +++ clang/lib/CodeGen/CGCUDANV.cpp @@ -158,6 +158,8 @@ void internalizeDeviceSideVar(const VarDecl *D, llvm::GlobalValue::LinkageTypes &Linkage) override; + + void transformManagedVars() override; }; } @@ -530,6 +532,9 @@ 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: { @@ -539,11 +544,16 @@ 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() + ".init")); replaceManagedVar(Var, ManagedVar); llvm::Value *Args[] = { &GpuBinaryHandlePtr, @@ -552,7 +562,8 @@ VarName, llvm::ConstantInt::get(VarSizeTy, VarSize), llvm::ConstantInt::get(IntTy, Var->getAlignment())}; - Builder.CreateCall(RegisterManagedVar, Args); + if (!Var->isDeclaration()) + Builder.CreateCall(RegisterManagedVar, Args); } else { llvm::Value *Args[] = { &GpuBinaryHandlePtr, @@ -964,9 +975,13 @@ // 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()) { registerDeviceVar(D, GV, !D->hasDefinition(), D->hasAttr()); + } } else if (D->getType()->isCUDADeviceBuiltinSurfaceType() || D->getType()->isCUDADeviceBuiltinTextureType()) { // Builtin surfaces and textures and their template arguments are @@ -994,3 +1009,26 @@ } } } + +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); + } + } +} Index: clang/lib/CodeGen/CGCUDARuntime.h =================================================================== --- clang/lib/CodeGen/CGCUDARuntime.h +++ clang/lib/CodeGen/CGCUDARuntime.h @@ -102,6 +102,9 @@ virtual void internalizeDeviceSideVar(const VarDecl *D, llvm::GlobalValue::LinkageTypes &Linkage) = 0; + + /// Transform managed variables in device compilation. + virtual void transformManagedVars() = 0; }; /// Creates an instance of a CUDA runtime class. Index: clang/lib/CodeGen/CodeGenModule.cpp =================================================================== --- clang/lib/CodeGen/CodeGenModule.cpp +++ clang/lib/CodeGen/CodeGenModule.cpp @@ -459,10 +459,11 @@ 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 (Context.getLangOpts().CUDAIsDevice) + CUDARuntime->transformManagedVars(); + else if (llvm::Function *CudaCtorFunction = + CUDARuntime->makeModuleCtorFunction()) AddGlobalCtor(CudaCtorFunction); } if (OpenMPRuntime) { @@ -3860,8 +3861,14 @@ } } - 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() && + D->hasExternalStorage()) + getCUDARuntime().handleVarRegistration(D, *GV); + } LangAS ExpectedAS = D ? D->getType().getAddressSpace() @@ -4169,12 +4176,8 @@ bool NeedsGlobalDtor = D->needsDestruction(getContext()) == QualType::DK_cxx_destructor; - bool IsHIPManagedVarOnDevice = - getLangOpts().CUDAIsDevice && D->hasAttr(); - const VarDecl *InitDecl; - const Expr *InitExpr = - IsHIPManagedVarOnDevice ? nullptr : D->getAnyInitializer(InitDecl); + const Expr *InitExpr = D->getAnyInitializer(InitDecl); Optional emitter; @@ -4300,14 +4303,11 @@ 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); + GV->setInitializer(Init); if (emitter) emitter->finalize(GV); Index: clang/test/CodeGenCUDA/device-var-linkage.cu =================================================================== --- clang/test/CodeGenCUDA/device-var-linkage.cu +++ clang/test/CodeGenCUDA/device-var-linkage.cu @@ -21,9 +21,9 @@ // 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 @@ -32,8 +32,8 @@ // 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 @@ -44,8 +44,8 @@ // 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); Index: clang/test/CodeGenCUDA/managed-var.cu =================================================================== --- clang/test/CodeGenCUDA/managed-var.cu +++ clang/test/CodeGenCUDA/managed-var.cu @@ -2,27 +2,27 @@ // 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 +// DEV-DAG: @x = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null +// NORDC-DAG: @x.init = internal global i32 1 +// RDC-DAG: @x.init = 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" struct vec { @@ -33,17 +33,16 @@ __managed__ vec v[100]; __managed__ vec v2[100] = {{1, 1, 1}}; -// DEV-DAG: @ex = external addrspace(1) global i32 -// HOST-DAG: @ex = external global i32 +// DEV-DAG: @ex = external addrspace(1) externally_initialized global i32 addrspace(1)* +// HOST-DAG: @ex.init = 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 = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null +// HOST-DAG: @_ZL2sx.init = internal global i32 1 +// HOST-DAG: @_ZL2sx = internal externally_initialized global i32* null static __managed__ int sx = 1; -// HOST-NOT: @ex.managed - // Force ex and sx mitted in device compilation. __global__ void foo(int *z) { *z = x + ex + sx; @@ -55,42 +54,54 @@ 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 @@ -100,10 +111,10 @@ } // 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 @@ -115,7 +126,19 @@ return (float)reinterpret_cast(&(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.init {{.*}}@[[DEVNAMEX]]{{.*}}, i64 4, i32 4) +// HOST-DAG: __hipRegisterManagedVar({{.*}}@_ZL2sx {{.*}}@_ZL2sx.init +// HOST-NOT: __hipRegisterManagedVar({{.*}}@ex {{.*}}@ex.init // HOST-DAG: declare void @__hipRegisterManagedVar(i8**, i8*, i8*, i8*, i64, i32) Index: llvm/lib/IR/ReplaceConstant.cpp =================================================================== --- llvm/lib/IR/ReplaceConstant.cpp +++ llvm/lib/IR/ReplaceConstant.cpp @@ -60,6 +60,7 @@ case Instruction::PtrToInt: case Instruction::IntToPtr: case Instruction::BitCast: + case Instruction::AddrSpaceCast: return dyn_cast( Builder.CreateCast((Instruction::CastOps)OpCode, CE->getOperand(0), CE->getType(), CE->getName()));