Author: Yaxun (Sam) Liu Date: 2021-01-22T11:43:58-05:00 New Revision: 622eaa4a4cea17c2cec6942d9702b010deae392b
URL: https://github.com/llvm/llvm-project/commit/622eaa4a4cea17c2cec6942d9702b010deae392b DIFF: https://github.com/llvm/llvm-project/commit/622eaa4a4cea17c2cec6942d9702b010deae392b.diff LOG: [HIP] Support __managed__ attribute This patch implements codegen for __managed__ variable attribute for HIP. Diagnostics will be added later. Differential Revision: https://reviews.llvm.org/D94814 Added: clang/test/AST/Inputs/cuda.h clang/test/AST/ast-dump-managed-var.cu clang/test/CodeGenCUDA/managed-var.cu clang/test/SemaCUDA/managed-var.cu llvm/include/llvm/IR/ReplaceConstant.h llvm/lib/IR/ReplaceConstant.cpp Modified: clang/include/clang/Basic/Attr.td clang/include/clang/Basic/AttrDocs.td clang/include/clang/Basic/DiagnosticSemaKinds.td clang/lib/CodeGen/CGCUDANV.cpp clang/lib/CodeGen/CGCUDARuntime.h clang/lib/CodeGen/CodeGenModule.cpp clang/lib/Sema/SemaDeclAttr.cpp clang/test/CodeGenCUDA/Inputs/cuda.h clang/test/Misc/pragma-attribute-supported-attributes-list.test clang/test/SemaCUDA/Inputs/cuda.h clang/test/SemaCUDA/bad-attributes.cu clang/test/SemaCUDA/device-var-init.cu clang/test/SemaCUDA/function-overload.cu clang/test/SemaCUDA/union-init.cu llvm/lib/IR/CMakeLists.txt llvm/lib/Target/XCore/XCoreLowerThreadLocal.cpp Removed: ################################################################################ diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index b30b91d3d4a6..bfd50f6a6779 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -324,6 +324,7 @@ class LangOpt<string name, code customCode = [{}]> { def MicrosoftExt : LangOpt<"MicrosoftExt">; def Borland : LangOpt<"Borland">; def CUDA : LangOpt<"CUDA">; +def HIP : LangOpt<"HIP">; def SYCL : LangOpt<"SYCLIsDevice">; def COnly : LangOpt<"", "!LangOpts.CPlusPlus">; def CPlusPlus : LangOpt<"CPlusPlus">; @@ -1115,6 +1116,13 @@ def CUDAHost : InheritableAttr { let Documentation = [Undocumented]; } +def HIPManaged : InheritableAttr { + let Spellings = [GNU<"managed">, Declspec<"__managed__">]; + let Subjects = SubjectList<[Var]>; + let LangOpts = [HIP]; + let Documentation = [HIPManagedAttrDocs]; +} + def CUDAInvalidTarget : InheritableAttr { let Spellings = []; let Subjects = SubjectList<[Function]>; diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index fffede41db1e..170a0fe3d4c4 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -5419,6 +5419,17 @@ unbind runtime APIs. }]; } +def HIPManagedAttrDocs : Documentation { + let Category = DocCatDecl; + let Content = [{ +The ``__managed__`` attribute can be applied to a global variable declaration in HIP. +A managed variable is emitted as an undefined global symbol in the device binary and is +registered by ``__hipRegisterManagedVariable`` in init functions. The HIP runtime allocates +managed memory and uses it to define the symbol when loading the device binary. +A managed variable can be accessed in both device and host code. + }]; +} + def LifetimeOwnerDocs : Documentation { let Category = DocCatDecl; let Content = [{ diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 758b2ed3e90b..67c59f3ca09a 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -8237,7 +8237,7 @@ def err_cuda_device_exceptions : Error< "%select{__device__|__global__|__host__|__host__ __device__}1 function">; def err_dynamic_var_init : Error< "dynamic initialization is not supported for " - "__device__, __constant__, and __shared__ variables.">; + "__device__, __constant__, __shared__, and __managed__ variables.">; def err_shared_var_init : Error< "initialization is not supported for __shared__ variables.">; def err_cuda_vla : Error< @@ -8247,7 +8247,8 @@ def err_cuda_extern_shared : Error<"__shared__ variable %0 cannot be 'extern'">; def err_cuda_host_shared : Error< "__shared__ local variables not allowed in " "%select{__device__|__global__|__host__|__host__ __device__}0 functions">; -def err_cuda_nonstatic_constdev: Error<"__constant__ and __device__ are not allowed on non-static local variables">; +def err_cuda_nonstatic_constdev: Error<"__constant__, __device__, and " + "__managed__ are not allowed on non-static local variables">; def err_cuda_ovl_target : Error< "%select{__device__|__global__|__host__|__host__ __device__}0 function %1 " "cannot overload %select{__device__|__global__|__host__|__host__ __device__}2 function %3">; diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp index 7c5ab39a85ec..33a2d6f4483e 100644 --- a/clang/lib/CodeGen/CGCUDANV.cpp +++ b/clang/lib/CodeGen/CGCUDANV.cpp @@ -21,6 +21,7 @@ #include "llvm/IR/BasicBlock.h" #include "llvm/IR/Constants.h" #include "llvm/IR/DerivedTypes.h" +#include "llvm/IR/ReplaceConstant.h" #include "llvm/Support/Format.h" using namespace clang; @@ -128,13 +129,15 @@ class CGNVCUDARuntime : public CGCUDARuntime { DeviceVars.push_back({&Var, VD, {DeviceVarFlags::Variable, Extern, Constant, - /*Normalized*/ false, /*Type*/ 0}}); + VD->hasAttr<HIPManagedAttr>(), + /*Normalized*/ false, 0}}); } void registerDeviceSurf(const VarDecl *VD, llvm::GlobalVariable &Var, bool Extern, int Type) override { DeviceVars.push_back({&Var, VD, {DeviceVarFlags::Surface, Extern, /*Constant*/ false, + /*Managed*/ false, /*Normalized*/ false, Type}}); } void registerDeviceTex(const VarDecl *VD, llvm::GlobalVariable &Var, @@ -142,7 +145,7 @@ class CGNVCUDARuntime : public CGCUDARuntime { DeviceVars.push_back({&Var, VD, {DeviceVarFlags::Texture, Extern, /*Constant*/ false, - Normalized, Type}}); + /*Managed*/ false, Normalized, Type}}); } /// Creates module constructor function @@ -380,6 +383,47 @@ void CGNVCUDARuntime::emitDeviceStubBodyLegacy(CodeGenFunction &CGF, CGF.EmitBlock(EndBlock); } +// Replace the original variable Var with the address loaded from variable +// ManagedVar populated by HIP runtime. +static void replaceManagedVar(llvm::GlobalVariable *Var, + llvm::GlobalVariable *ManagedVar) { + SmallVector<SmallVector<llvm::User *, 8>, 8> WorkList; + for (auto &&VarUse : Var->uses()) { + WorkList.push_back({VarUse.getUser()}); + } + while (!WorkList.empty()) { + auto &&WorkItem = WorkList.pop_back_val(); + auto *U = WorkItem.back(); + if (isa<llvm::ConstantExpr>(U)) { + for (auto &&UU : U->uses()) { + WorkItem.push_back(UU.getUser()); + WorkList.push_back(WorkItem); + WorkItem.pop_back(); + } + continue; + } + if (auto *I = dyn_cast<llvm::Instruction>(U)) { + llvm::Value *OldV = Var; + llvm::Instruction *NewV = + new llvm::LoadInst(Var->getType(), ManagedVar, "ld.managed", false, + llvm::Align(Var->getAlignment()), I); + WorkItem.pop_back(); + // Replace constant expressions directly or indirectly using the managed + // variable with instructions. + for (auto &&Op : WorkItem) { + auto *CE = cast<llvm::ConstantExpr>(Op); + auto *NewInst = llvm::createReplacementInstr(CE, I); + NewInst->replaceUsesOfWith(OldV, NewV); + OldV = CE; + NewV = NewInst; + } + I->replaceUsesOfWith(OldV, NewV); + } else { + llvm_unreachable("Invalid use of managed variable"); + } + } +} + /// Creates a function that sets up state on the host side for CUDA objects that /// have a presence on both the host and device sides. Specifically, registers /// the host side of kernel functions and device global variables with the CUDA @@ -452,6 +496,13 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() { llvm::FunctionCallee RegisterVar = CGM.CreateRuntimeFunction( llvm::FunctionType::get(VoidTy, RegisterVarParams, false), addUnderscoredPrefixToName("RegisterVar")); + // void __hipRegisterManagedVar(void **, char *, char *, const char *, + // size_t, unsigned) + llvm::Type *RegisterManagedVarParams[] = {VoidPtrPtrTy, CharPtrTy, CharPtrTy, + CharPtrTy, VarSizeTy, IntTy}; + llvm::FunctionCallee RegisterManagedVar = CGM.CreateRuntimeFunction( + llvm::FunctionType::get(VoidTy, RegisterManagedVarParams, false), + addUnderscoredPrefixToName("RegisterManagedVar")); // void __cudaRegisterSurface(void **, const struct surfaceReference *, // const void **, const char *, int, int); llvm::FunctionCallee RegisterSurf = CGM.CreateRuntimeFunction( @@ -474,16 +525,34 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() { case DeviceVarFlags::Variable: { uint64_t VarSize = CGM.getDataLayout().getTypeAllocSize(Var->getValueType()); - llvm::Value *Args[] = { - &GpuBinaryHandlePtr, - Builder.CreateBitCast(Var, VoidPtrTy), - VarName, - VarName, - llvm::ConstantInt::get(IntTy, Info.Flags.isExtern()), - llvm::ConstantInt::get(VarSizeTy, VarSize), - llvm::ConstantInt::get(IntTy, Info.Flags.isConstant()), - llvm::ConstantInt::get(IntTy, 0)}; - Builder.CreateCall(RegisterVar, Args); + if (Info.Flags.isManaged()) { + auto ManagedVar = new llvm::GlobalVariable( + CGM.getModule(), Var->getType(), + /*isConstant=*/false, Var->getLinkage(), + /*Init=*/llvm::ConstantPointerNull::get(Var->getType()), + Twine(Var->getName() + ".managed"), /*InsertBefore=*/nullptr, + llvm::GlobalVariable::NotThreadLocal); + replaceManagedVar(Var, ManagedVar); + llvm::Value *Args[] = { + &GpuBinaryHandlePtr, + Builder.CreateBitCast(ManagedVar, VoidPtrTy), + Builder.CreateBitCast(Var, VoidPtrTy), + VarName, + llvm::ConstantInt::get(VarSizeTy, VarSize), + llvm::ConstantInt::get(IntTy, Var->getAlignment())}; + Builder.CreateCall(RegisterManagedVar, Args); + } else { + llvm::Value *Args[] = { + &GpuBinaryHandlePtr, + Builder.CreateBitCast(Var, VoidPtrTy), + VarName, + VarName, + llvm::ConstantInt::get(IntTy, Info.Flags.isExtern()), + llvm::ConstantInt::get(VarSizeTy, VarSize), + llvm::ConstantInt::get(IntTy, Info.Flags.isConstant()), + llvm::ConstantInt::get(IntTy, 0)}; + Builder.CreateCall(RegisterVar, Args); + } break; } case DeviceVarFlags::Surface: diff --git a/clang/lib/CodeGen/CGCUDARuntime.h b/clang/lib/CodeGen/CGCUDARuntime.h index 19e70a2022a5..ba3404ead368 100644 --- a/clang/lib/CodeGen/CGCUDARuntime.h +++ b/clang/lib/CodeGen/CGCUDARuntime.h @@ -54,16 +54,19 @@ class CGCUDARuntime { unsigned Kind : 2; unsigned Extern : 1; unsigned Constant : 1; // Constant variable. + unsigned Managed : 1; // Managed variable. unsigned Normalized : 1; // Normalized texture. int SurfTexType; // Type of surface/texutre. public: - DeviceVarFlags(DeviceVarKind K, bool E, bool C, bool N, int T) - : Kind(K), Extern(E), Constant(C), Normalized(N), SurfTexType(T) {} + DeviceVarFlags(DeviceVarKind K, bool E, bool C, bool M, bool N, int T) + : Kind(K), Extern(E), Constant(C), Managed(M), Normalized(N), + SurfTexType(T) {} DeviceVarKind getKind() const { return static_cast<DeviceVarKind>(Kind); } bool isExtern() const { return Extern; } bool isConstant() const { return Constant; } + bool isManaged() const { return Managed; } bool isNormalized() const { return Normalized; } int getSurfTexType() const { return SurfTexType; } }; diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index bee51715bdc6..18d633911f55 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4152,13 +4152,14 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, // Shadows of initialized device-side global variables are also left // undefined. bool IsCUDAShadowVar = - !getLangOpts().CUDAIsDevice && + !getLangOpts().CUDAIsDevice && !D->hasAttr<HIPManagedAttr>() && (D->hasAttr<CUDAConstantAttr>() || D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDASharedAttr>()); bool IsCUDADeviceShadowVar = getLangOpts().CUDAIsDevice && (D->getType()->isCUDADeviceBuiltinSurfaceType() || - D->getType()->isCUDADeviceBuiltinTextureType()); + D->getType()->isCUDADeviceBuiltinTextureType() || + D->hasAttr<HIPManagedAttr>()); // HIP pinned shadow of initialized host-side global variables are also // left undefined. if (getLangOpts().CUDA && diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index bd8ec2bdef76..30d08b3d4ac0 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -4493,7 +4493,8 @@ static void handleOptimizeNoneAttr(Sema &S, Decl *D, const ParsedAttr &AL) { } static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) { - if (checkAttrMutualExclusion<CUDASharedAttr>(S, D, AL)) + if (checkAttrMutualExclusion<CUDASharedAttr>(S, D, AL) || + checkAttrMutualExclusion<HIPManagedAttr>(S, D, AL)) return; const auto *VD = cast<VarDecl>(D); if (VD->hasLocalStorage()) { @@ -4504,7 +4505,8 @@ static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) { } static void handleSharedAttr(Sema &S, Decl *D, const ParsedAttr &AL) { - if (checkAttrMutualExclusion<CUDAConstantAttr>(S, D, AL)) + if (checkAttrMutualExclusion<CUDAConstantAttr>(S, D, AL) || + checkAttrMutualExclusion<HIPManagedAttr>(S, D, AL)) return; const auto *VD = cast<VarDecl>(D); // extern __shared__ is only allowed on arrays with no length (e.g. @@ -4569,9 +4571,33 @@ static void handleDeviceAttr(Sema &S, Decl *D, const ParsedAttr &AL) { return; } } + + if (auto *A = D->getAttr<CUDADeviceAttr>()) { + if (!A->isImplicit()) + return; + D->dropAttr<CUDADeviceAttr>(); + } D->addAttr(::new (S.Context) CUDADeviceAttr(S.Context, AL)); } +static void handleManagedAttr(Sema &S, Decl *D, const ParsedAttr &AL) { + if (checkAttrMutualExclusion<CUDAConstantAttr>(S, D, AL) || + checkAttrMutualExclusion<CUDASharedAttr>(S, D, AL)) { + return; + } + + if (const auto *VD = dyn_cast<VarDecl>(D)) { + if (VD->hasLocalStorage()) { + S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev); + return; + } + } + if (!D->hasAttr<HIPManagedAttr>()) + D->addAttr(::new (S.Context) HIPManagedAttr(S.Context, AL)); + if (!D->hasAttr<CUDADeviceAttr>()) + D->addAttr(CUDADeviceAttr::CreateImplicit(S.Context)); +} + static void handleGNUInlineAttr(Sema &S, Decl *D, const ParsedAttr &AL) { const auto *Fn = cast<FunctionDecl>(D); if (!Fn->isInlineSpecified()) { @@ -7793,6 +7819,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case ParsedAttr::AT_CUDAHost: handleSimpleAttributeWithExclusions<CUDAHostAttr, CUDAGlobalAttr>(S, D, AL); break; + case ParsedAttr::AT_HIPManaged: + handleManagedAttr(S, D, AL); + break; case ParsedAttr::AT_CUDADeviceBuiltinSurfaceType: handleSimpleAttributeWithExclusions<CUDADeviceBuiltinSurfaceTypeAttr, CUDADeviceBuiltinTextureTypeAttr>(S, D, diff --git a/clang/test/AST/Inputs/cuda.h b/clang/test/AST/Inputs/cuda.h new file mode 100644 index 000000000000..405ef8bb807d --- /dev/null +++ b/clang/test/AST/Inputs/cuda.h @@ -0,0 +1,54 @@ +/* Minimal declarations for CUDA support. Testing purposes only. */ + +#include <stddef.h> + +// Make this file work with nvcc, for testing compatibility. + +#ifndef __NVCC__ +#define __constant__ __attribute__((constant)) +#define __device__ __attribute__((device)) +#define __global__ __attribute__((global)) +#define __host__ __attribute__((host)) +#define __shared__ __attribute__((shared)) +#define __managed__ __attribute__((managed)) +#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__))) + +struct dim3 { + unsigned x, y, z; + __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {} +}; + +#ifdef __HIP__ +typedef struct hipStream *hipStream_t; +typedef enum hipError {} hipError_t; +int hipConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0, + hipStream_t stream = 0); +extern "C" hipError_t __hipPushCallConfiguration(dim3 gridSize, dim3 blockSize, + size_t sharedSize = 0, + hipStream_t stream = 0); +extern "C" hipError_t hipLaunchKernel(const void *func, dim3 gridDim, + dim3 blockDim, void **args, + size_t sharedMem, + hipStream_t stream); +#else +typedef struct cudaStream *cudaStream_t; +typedef enum cudaError {} cudaError_t; + +extern "C" int cudaConfigureCall(dim3 gridSize, dim3 blockSize, + size_t sharedSize = 0, + cudaStream_t stream = 0); +extern "C" int __cudaPushCallConfiguration(dim3 gridSize, dim3 blockSize, + size_t sharedSize = 0, + cudaStream_t stream = 0); +extern "C" cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim, + dim3 blockDim, void **args, + size_t sharedMem, cudaStream_t stream); +#endif + +// Host- and device-side placement new overloads. +void *operator new(__SIZE_TYPE__, void *p) { return p; } +void *operator new[](__SIZE_TYPE__, void *p) { return p; } +__device__ void *operator new(__SIZE_TYPE__, void *p) { return p; } +__device__ void *operator new[](__SIZE_TYPE__, void *p) { return p; } + +#endif // !__NVCC__ diff --git a/clang/test/AST/ast-dump-managed-var.cu b/clang/test/AST/ast-dump-managed-var.cu new file mode 100644 index 000000000000..862a70c81f9e --- /dev/null +++ b/clang/test/AST/ast-dump-managed-var.cu @@ -0,0 +1,28 @@ +// RUN: %clang_cc1 -ast-dump -x hip %s | FileCheck %s +// RUN: %clang_cc1 -ast-dump -fcuda-is-device -x hip %s | FileCheck %s + +#include "Inputs/cuda.h" + +// CHECK-LABEL: VarDecl {{.*}} m1 'int' +// CHECK-NEXT: HIPManagedAttr +// CHECK-NEXT: CUDADeviceAttr {{.*}}Implicit +__managed__ int m1; + +// CHECK-LABEL: VarDecl {{.*}} m2 'int' +// CHECK-NEXT: HIPManagedAttr +// CHECK-NEXT: CUDADeviceAttr {{.*}}Implicit +// CHECK-NOT: HIPManagedAttr +// CHECK-NOT: CUDADeviceAttr +__managed__ __managed__ int m2; + +// CHECK-LABEL: VarDecl {{.*}} m3 'int' +// CHECK-NEXT: HIPManagedAttr +// CHECK-NEXT: CUDADeviceAttr {{.*}}line +// CHECK-NOT: CUDADeviceAttr {{.*}}Implicit +__managed__ __device__ int m3; + +// CHECK-LABEL: VarDecl {{.*}} m3a 'int' +// CHECK-NEXT: CUDADeviceAttr {{.*}}cuda.h +// CHECK-NEXT: HIPManagedAttr +// CHECK-NOT: CUDADeviceAttr {{.*}}Implicit +__device__ __managed__ int m3a; diff --git a/clang/test/CodeGenCUDA/Inputs/cuda.h b/clang/test/CodeGenCUDA/Inputs/cuda.h index 5d73b81041ab..daa6328c9499 100644 --- a/clang/test/CodeGenCUDA/Inputs/cuda.h +++ b/clang/test/CodeGenCUDA/Inputs/cuda.h @@ -7,6 +7,9 @@ #define __global__ __attribute__((global)) #define __host__ __attribute__((host)) #define __shared__ __attribute__((shared)) +#if __HIP__ +#define __managed__ __attribute__((managed)) +#endif #define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__))) struct dim3 { diff --git a/clang/test/CodeGenCUDA/managed-var.cu b/clang/test/CodeGenCUDA/managed-var.cu new file mode 100644 index 000000000000..c95139784085 --- /dev/null +++ b/clang/test/CodeGenCUDA/managed-var.cu @@ -0,0 +1,100 @@ +// REQUIRES: x86-registered-target, amdgpu-registered-target + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck \ +// RUN: -check-prefixes=DEV %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \ +// RUN: -emit-llvm -fgpu-rdc -o - -x hip %s | FileCheck \ +// RUN: -check-prefixes=DEV %s + +// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck \ +// RUN: -check-prefixes=HOST %s + +// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \ +// RUN: -emit-llvm -fgpu-rdc -o - -x hip %s | FileCheck \ +// RUN: -check-prefixes=HOST %s + +#include "Inputs/cuda.h" + +// DEV-DAG: @x = {{.*}}addrspace(1) externally_initialized global i32 undef +// HOST-DAG: @x = internal global i32 1 +// HOST-DAG: @x.managed = internal global i32* null +// HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"x\00" + +struct vec { + float x,y,z; +}; + +__managed__ int x = 1; +__managed__ vec v[100]; +__managed__ vec v2[100] = {{1, 1, 1}}; + +__global__ void foo(int *z) { + *z = x; + v[1].x = 2; +} + +// HOST-LABEL: define {{.*}}@_Z4loadv() +// HOST: %ld.managed = load i32*, i32** @x.managed, align 4 +// HOST: %0 = load i32, i32* %ld.managed, align 4 +// HOST: ret i32 %0 +int load() { + return x; +} + +// HOST-LABEL: define {{.*}}@_Z5storev() +// HOST: %ld.managed = load i32*, i32** @x.managed, align 4 +// HOST: store i32 2, i32* %ld.managed, align 4 +void store() { + x = 2; +} + +// HOST-LABEL: define {{.*}}@_Z10addr_takenv() +// HOST: %ld.managed = load i32*, i32** @x.managed, align 4 +// HOST: store i32* %ld.managed, i32** %p, align 8 +// HOST: %0 = load i32*, i32** %p, align 8 +// HOST: store i32 3, i32* %0, align 4 +void addr_taken() { + int *p = &x; + *p = 3; +} + +// HOST-LABEL: define {{.*}}@_Z5load2v() +// HOST: %ld.managed = load [100 x %struct.vec]*, [100 x %struct.vec]** @v.managed, align 16 +// HOST: %0 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %ld.managed, i64 0, i64 1, i32 0 +// HOST: %1 = load float, float* %0, align 4 +// HOST: ret float %1 +float load2() { + return v[1].x; +} + +// HOST-LABEL: define {{.*}}@_Z5load3v() +// HOST: %ld.managed = load <{ %struct.vec, [99 x %struct.vec] }>*, <{ %struct.vec, [99 x %struct.vec] }>** @v2.managed, align 16 +// HOST: %0 = bitcast <{ %struct.vec, [99 x %struct.vec] }>* %ld.managed to [100 x %struct.vec]* +// HOST: %1 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %0, i64 0, i64 1, i32 1 +// HOST: %2 = load float, float* %1, align 4 +// HOST: ret float %2 +float load3() { + return v2[1].y; +} + +// HOST-LABEL: define {{.*}}@_Z11addr_taken2v() +// HOST: %ld.managed = load [100 x %struct.vec]*, [100 x %struct.vec]** @v.managed, align 16 +// HOST: %0 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %ld.managed, i64 0, i64 1, i32 0 +// HOST: %1 = ptrtoint float* %0 to i64 +// HOST: %ld.managed1 = load <{ %struct.vec, [99 x %struct.vec] }>*, <{ %struct.vec, [99 x %struct.vec] }>** @v2.managed, align 16 +// HOST: %2 = bitcast <{ %struct.vec, [99 x %struct.vec] }>* %ld.managed1 to [100 x %struct.vec]* +// HOST: %3 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %2, i64 0, i64 1, i32 1 +// HOST: %4 = ptrtoint float* %3 to i64 +// HOST: %5 = sub i64 %4, %1 +// HOST: %6 = sdiv i64 %5, 4 +// HOST: %7 = sitofp i64 %6 to float +// HOST: ret float %7 +float addr_taken2() { + return (float)reinterpret_cast<long>(&(v2[1].y)-&(v[1].x)); +} + +// HOST-DAG: __hipRegisterManagedVar({{.*}}@x.managed {{.*}}@x {{.*}}@[[DEVNAMEX]]{{.*}}, i64 4, i32 4) +// HOST-DAG: declare void @__hipRegisterManagedVar(i8**, i8*, i8*, i8*, i64, i32) diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index 270b377c2064..cb62f56912aa 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -66,6 +66,7 @@ // CHECK-NEXT: FlagEnum (SubjectMatchRule_enum) // CHECK-NEXT: Flatten (SubjectMatchRule_function) // CHECK-NEXT: GNUInline (SubjectMatchRule_function) +// CHECK-NEXT: HIPManaged (SubjectMatchRule_variable) // CHECK-NEXT: Hot (SubjectMatchRule_function) // CHECK-NEXT: IBAction (SubjectMatchRule_objc_method_is_instance) // CHECK-NEXT: IFunc (SubjectMatchRule_function) diff --git a/clang/test/SemaCUDA/Inputs/cuda.h b/clang/test/SemaCUDA/Inputs/cuda.h index 901f8e0c17cf..405ef8bb807d 100644 --- a/clang/test/SemaCUDA/Inputs/cuda.h +++ b/clang/test/SemaCUDA/Inputs/cuda.h @@ -10,6 +10,7 @@ #define __global__ __attribute__((global)) #define __host__ __attribute__((host)) #define __shared__ __attribute__((shared)) +#define __managed__ __attribute__((managed)) #define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__))) struct dim3 { diff --git a/clang/test/SemaCUDA/bad-attributes.cu b/clang/test/SemaCUDA/bad-attributes.cu index a990598e7b7e..8ac6fa1ab5c2 100644 --- a/clang/test/SemaCUDA/bad-attributes.cu +++ b/clang/test/SemaCUDA/bad-attributes.cu @@ -64,11 +64,11 @@ __global__ static inline void foobar() {}; __constant__ int global_constant; void host_fn() { - __constant__ int c; // expected-error {{__constant__ and __device__ are not allowed on non-static local variables}} + __constant__ int c; // expected-error {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}} __shared__ int s; // expected-error {{__shared__ local variables not allowed in __host__ functions}} } __device__ void device_fn() { - __constant__ int c; // expected-error {{__constant__ and __device__ are not allowed on non-static local variables}} + __constant__ int c; // expected-error {{__constant__, __device__, and __managed__ are not allowed on non-static local variables}} } typedef __attribute__((device_builtin_surface_type)) unsigned long long s0_ty; // expected-warning {{'device_builtin_surface_type' attribute only applies to classes}} diff --git a/clang/test/SemaCUDA/device-var-init.cu b/clang/test/SemaCUDA/device-var-init.cu index 88350f56651c..9d499bddbe1b 100644 --- a/clang/test/SemaCUDA/device-var-init.cu +++ b/clang/test/SemaCUDA/device-var-init.cu @@ -16,11 +16,11 @@ __shared__ int s_v_i = 1; // expected-error@-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.}} @@ -32,175 +32,175 @@ __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. @@ -216,20 +216,20 @@ __device__ void df_sema() { // __shared__ does not need to be explicitly static. __shared__ int lsi; - // __constant__ and __device__ can not be non-static local + // __constant__, __device__, and __managed__ can not be non-static local __constant__ int lci; - // expected-error@-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.}} @@ -241,175 +241,175 @@ __device__ void df_sema() { 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() { diff --git a/clang/test/SemaCUDA/function-overload.cu b/clang/test/SemaCUDA/function-overload.cu index 574b65ee7fd8..822e25996820 100644 --- a/clang/test/SemaCUDA/function-overload.cu +++ b/clang/test/SemaCUDA/function-overload.cu @@ -463,7 +463,7 @@ int test_constexpr_overload(C2 &x, C2 &y) { // Verify no ambiguity for new operator. void *a = new int; __device__ void *b = new int; -// expected-error@-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(); diff --git a/clang/test/SemaCUDA/managed-var.cu b/clang/test/SemaCUDA/managed-var.cu new file mode 100644 index 000000000000..3f699b79a043 --- /dev/null +++ b/clang/test/SemaCUDA/managed-var.cu @@ -0,0 +1,54 @@ +// RUN: %clang_cc1 -fsyntax-only -verify -x hip %s +// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify -x hip %s +// RUN: %clang_cc1 -fsyntax-only -fgpu-rdc -verify -x hip %s +// RUN: %clang_cc1 -fsyntax-only -fgpu-rdc -fcuda-is-device -verify -x hip %s + +#include "Inputs/cuda.h" + +struct A { + int a; + A() { a = 1; } +}; + +__managed__ int m1; + +__managed__ __managed__ int m2; + +__managed__ __device__ int m3; +__device__ __managed__ int m3a; + +__managed__ __constant__ int m4; +// expected-error@-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}} diff --git a/clang/test/SemaCUDA/union-init.cu b/clang/test/SemaCUDA/union-init.cu index a633975e3776..9e4d14a71069 100644 --- a/clang/test/SemaCUDA/union-init.cu +++ b/clang/test/SemaCUDA/union-init.cu @@ -31,9 +31,9 @@ union D { __device__ B b; __device__ C c; -// expected-error@-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; diff --git a/llvm/include/llvm/IR/ReplaceConstant.h b/llvm/include/llvm/IR/ReplaceConstant.h new file mode 100644 index 000000000000..753f6d558ef8 --- /dev/null +++ b/llvm/include/llvm/IR/ReplaceConstant.h @@ -0,0 +1,28 @@ +//===- ReplaceConstant.h - Replacing LLVM constant expressions --*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file declares the utility function for replacing LLVM constant +// expressions by instructions. +// +//===----------------------------------------------------------------------===// + +#ifndef LLVM_IR_REPLACECONSTANT_H +#define LLVM_IR_REPLACECONSTANT_H + +#include "llvm/IR/Constants.h" +#include "llvm/IR/Instruction.h" + +namespace llvm { + +/// Create a replacement instruction for constant expression \p CE and insert +/// it before \p Instr. +Instruction *createReplacementInstr(ConstantExpr *CE, Instruction *Instr); + +} // end namespace llvm + +#endif // LLVM_IR_REPLACECONSTANT_H diff --git a/llvm/lib/IR/CMakeLists.txt b/llvm/lib/IR/CMakeLists.txt index ca570121460b..fb4993742e85 100644 --- a/llvm/lib/IR/CMakeLists.txt +++ b/llvm/lib/IR/CMakeLists.txt @@ -49,6 +49,7 @@ add_llvm_component_library(LLVMCore SafepointIRVerifier.cpp ProfileSummary.cpp PseudoProbe.cpp + ReplaceConstant.cpp Statepoint.cpp StructuralHash.cpp Type.cpp diff --git a/llvm/lib/IR/ReplaceConstant.cpp b/llvm/lib/IR/ReplaceConstant.cpp new file mode 100644 index 000000000000..7efa525d427e --- /dev/null +++ b/llvm/lib/IR/ReplaceConstant.cpp @@ -0,0 +1,70 @@ +//===- ReplaceConstant.cpp - Replace LLVM constant expression--------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file implements a utility function for replacing LLVM constant +// expressions by instructions. +// +//===----------------------------------------------------------------------===// + +#include "llvm/IR/ReplaceConstant.h" +#include "llvm/IR/IRBuilder.h" +#include "llvm/IR/Instructions.h" +#include "llvm/IR/NoFolder.h" + +namespace llvm { +// Replace a constant expression by instructions with equivalent operations at +// a specified location. +Instruction *createReplacementInstr(ConstantExpr *CE, Instruction *Instr) { + IRBuilder<NoFolder> Builder(Instr); + unsigned OpCode = CE->getOpcode(); + switch (OpCode) { + case Instruction::GetElementPtr: { + SmallVector<Value *, 4> CEOpVec(CE->operands()); + ArrayRef<Value *> CEOps(CEOpVec); + return dyn_cast<Instruction>( + Builder.CreateInBoundsGEP(cast<GEPOperator>(CE)->getSourceElementType(), + CEOps[0], CEOps.slice(1))); + } + case Instruction::Add: + case Instruction::Sub: + case Instruction::Mul: + case Instruction::UDiv: + case Instruction::SDiv: + case Instruction::FDiv: + case Instruction::URem: + case Instruction::SRem: + case Instruction::FRem: + case Instruction::Shl: + case Instruction::LShr: + case Instruction::AShr: + case Instruction::And: + case Instruction::Or: + case Instruction::Xor: + return dyn_cast<Instruction>( + Builder.CreateBinOp((Instruction::BinaryOps)OpCode, CE->getOperand(0), + CE->getOperand(1), CE->getName())); + case Instruction::Trunc: + case Instruction::ZExt: + case Instruction::SExt: + case Instruction::FPToUI: + case Instruction::FPToSI: + case Instruction::UIToFP: + case Instruction::SIToFP: + case Instruction::FPTrunc: + case Instruction::FPExt: + case Instruction::PtrToInt: + case Instruction::IntToPtr: + case Instruction::BitCast: + return dyn_cast<Instruction>( + Builder.CreateCast((Instruction::CastOps)OpCode, CE->getOperand(0), + CE->getType(), CE->getName())); + default: + llvm_unreachable("Unhandled constant expression!\n"); + } +} +} // namespace llvm diff --git a/llvm/lib/Target/XCore/XCoreLowerThreadLocal.cpp b/llvm/lib/Target/XCore/XCoreLowerThreadLocal.cpp index bd269f7f4c43..6528154ab0e2 100644 --- a/llvm/lib/Target/XCore/XCoreLowerThreadLocal.cpp +++ b/llvm/lib/Target/XCore/XCoreLowerThreadLocal.cpp @@ -21,6 +21,7 @@ #include "llvm/IR/IntrinsicsXCore.h" #include "llvm/IR/Module.h" #include "llvm/IR/NoFolder.h" +#include "llvm/IR/ReplaceConstant.h" #include "llvm/IR/ValueHandle.h" #include "llvm/Pass.h" #include "llvm/Support/CommandLine.h" @@ -74,57 +75,6 @@ createLoweredInitializer(ArrayType *NewType, Constant *OriginalInitializer) { return ConstantArray::get(NewType, Elements); } -static Instruction * -createReplacementInstr(ConstantExpr *CE, Instruction *Instr) { - IRBuilder<NoFolder> Builder(Instr); - unsigned OpCode = CE->getOpcode(); - switch (OpCode) { - case Instruction::GetElementPtr: { - SmallVector<Value *, 4> CEOpVec(CE->operands()); - ArrayRef<Value *> CEOps(CEOpVec); - return dyn_cast<Instruction>(Builder.CreateInBoundsGEP( - cast<GEPOperator>(CE)->getSourceElementType(), CEOps[0], - CEOps.slice(1))); - } - case Instruction::Add: - case Instruction::Sub: - case Instruction::Mul: - case Instruction::UDiv: - case Instruction::SDiv: - case Instruction::FDiv: - case Instruction::URem: - case Instruction::SRem: - case Instruction::FRem: - case Instruction::Shl: - case Instruction::LShr: - case Instruction::AShr: - case Instruction::And: - case Instruction::Or: - case Instruction::Xor: - return dyn_cast<Instruction>( - Builder.CreateBinOp((Instruction::BinaryOps)OpCode, - CE->getOperand(0), CE->getOperand(1), - CE->getName())); - case Instruction::Trunc: - case Instruction::ZExt: - case Instruction::SExt: - case Instruction::FPToUI: - case Instruction::FPToSI: - case Instruction::UIToFP: - case Instruction::SIToFP: - case Instruction::FPTrunc: - case Instruction::FPExt: - case Instruction::PtrToInt: - case Instruction::IntToPtr: - case Instruction::BitCast: - return dyn_cast<Instruction>( - Builder.CreateCast((Instruction::CastOps)OpCode, - CE->getOperand(0), CE->getType(), - CE->getName())); - default: - llvm_unreachable("Unhandled constant expression!\n"); - } -} static bool replaceConstantExprOp(ConstantExpr *CE, Pass *P) { do { _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
