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