Index: clang/include/clang/Basic/Attr.td =================================================================== --- clang/include/clang/Basic/Attr.td +++ clang/include/clang/Basic/Attr.td @@ -324,6 +324,7 @@ 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">; @@ -1108,6 +1109,13 @@ 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]>; Index: clang/include/clang/Basic/AttrDocs.td =================================================================== --- clang/include/clang/Basic/AttrDocs.td +++ clang/include/clang/Basic/AttrDocs.td @@ -5419,6 +5419,17 @@ }]; } +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 a undefined global symbol in device binary and +registered by ``__hipRegisterManagedVariable`` in init functions. HIP runtime allocates +managed memory and use 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 = [{ Index: clang/include/clang/Sema/SemaInternal.h =================================================================== --- clang/include/clang/Sema/SemaInternal.h +++ clang/include/clang/Sema/SemaInternal.h @@ -44,9 +44,9 @@ inline bool DeclAttrsMatchCUDAMode(const LangOptions &LangOpts, Decl *D) { if (!LangOpts.CUDA || !D) return true; - bool isDeviceSideDecl = D->hasAttr() || - D->hasAttr() || - D->hasAttr(); + bool isDeviceSideDecl = + D->hasAttr() || D->hasAttr() || + D->hasAttr() || D->hasAttr(); return isDeviceSideDecl == LangOpts.CUDAIsDevice; } Index: clang/lib/CodeGen/CGCUDANV.cpp =================================================================== --- clang/lib/CodeGen/CGCUDANV.cpp +++ 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 @@ DeviceVars.push_back({&Var, VD, {DeviceVarFlags::Variable, Extern, Constant, - /*Normalized*/ false, /*Type*/ 0}}); + VD->hasAttr(), + /*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 @@ DeviceVars.push_back({&Var, VD, {DeviceVarFlags::Texture, Extern, /*Constant*/ false, - Normalized, Type}}); + /*Managed*/ false, Normalized, Type}}); } /// Creates module constructor function @@ -380,6 +383,47 @@ 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, 8> WorkList; + for (auto &&UU : Var->uses()) { + WorkList.push_back({UU.getUser()}); + } + while (!WorkList.empty()) { + auto &&US = WorkList.pop_back_val(); + auto *U = US.back(); + if (isa(U)) { + for (auto &&UU : U->uses()) { + US.push_back(UU.getUser()); + WorkList.push_back(US); + US.pop_back(); + } + continue; + } + if (auto *I = dyn_cast(U)) { + llvm::Value *OldV = Var; + llvm::Instruction *NewV = + new llvm::LoadInst(Var->getType(), ManagedVar, "ld.managed", false, + llvm::Align(Var->getAlignment()), I); + US.pop_back(); + // Replace constant expressions directly or indirectly using the managed + // variable with instructions. + for (auto &&Op : US) { + auto *CE = cast(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::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 @@ 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, llvm::GlobalValue::ExternalLinkage, nullptr, + Twine(Var->getName() + ".managed"), 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, CGM.getLangOpts().HIP && + Info.Flags.isManaged())}; + Builder.CreateCall(RegisterVar, Args); + } break; } case DeviceVarFlags::Surface: Index: clang/lib/CodeGen/CGCUDARuntime.h =================================================================== --- clang/lib/CodeGen/CGCUDARuntime.h +++ clang/lib/CodeGen/CGCUDARuntime.h @@ -54,16 +54,19 @@ 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(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; } }; Index: clang/lib/CodeGen/CGDeclCXX.cpp =================================================================== --- clang/lib/CodeGen/CGDeclCXX.cpp +++ clang/lib/CodeGen/CGDeclCXX.cpp @@ -460,7 +460,7 @@ // are allowed are empty and we just need to ignore them here. if (getLangOpts().CUDAIsDevice && !getLangOpts().GPUAllowDeviceInit && (D->hasAttr() || D->hasAttr() || - D->hasAttr())) + D->hasAttr() || D->hasAttr())) return; if (getLangOpts().OpenMP && Index: clang/lib/CodeGen/CodeGenModule.cpp =================================================================== --- clang/lib/CodeGen/CodeGenModule.cpp +++ clang/lib/CodeGen/CodeGenModule.cpp @@ -2740,6 +2740,7 @@ !Global->hasAttr() && !Global->hasAttr() && !Global->hasAttr() && + !Global->hasAttr() && !Global->getType()->isCUDADeviceBuiltinSurfaceType() && !Global->getType()->isCUDADeviceBuiltinTextureType()) return; @@ -3993,7 +3994,8 @@ return LangAS::cuda_constant; else if (D && D->hasAttr()) return LangAS::cuda_shared; - else if (D && D->hasAttr()) + else if (D && + (D->hasAttr() || D->hasAttr())) return LangAS::cuda_device; else if (D && D->getType().isConstQualified()) return LangAS::cuda_constant; @@ -4151,7 +4153,8 @@ bool IsCUDADeviceShadowVar = getLangOpts().CUDAIsDevice && (D->getType()->isCUDADeviceBuiltinSurfaceType() || - D->getType()->isCUDADeviceBuiltinTextureType()); + D->getType()->isCUDADeviceBuiltinTextureType() || + D->hasAttr()); // HIP pinned shadow of initialized host-side global variables are also // left undefined. if (getLangOpts().CUDA && @@ -4256,14 +4259,16 @@ if (GV && LangOpts.CUDA) { if (LangOpts.CUDAIsDevice) { if (Linkage != llvm::GlobalValue::InternalLinkage && - (D->hasAttr() || D->hasAttr())) + (D->hasAttr() || D->hasAttr() || + D->hasAttr())) GV->setExternallyInitialized(true); } else { // Host-side shadows of external declarations of device-side // global variables become internal definitions. These have to // be internal in order to prevent name conflicts with global // host variables with the same name in a different TUs. - if (D->hasAttr() || D->hasAttr()) { + if (D->hasAttr() || D->hasAttr() || + D->hasAttr()) { Linkage = llvm::GlobalValue::InternalLinkage; // Shadow variables and their properties must be registered with CUDA // runtime. Skip Extern global variables, which will be registered in Index: clang/lib/CodeGen/TargetInfo.cpp =================================================================== --- clang/lib/CodeGen/TargetInfo.cpp +++ clang/lib/CodeGen/TargetInfo.cpp @@ -8958,6 +8958,7 @@ (isa(D) && D->hasAttr()) || (isa(D) && (D->hasAttr() || D->hasAttr() || + D->hasAttr() || cast(D)->getType()->isCUDADeviceBuiltinSurfaceType() || cast(D)->getType()->isCUDADeviceBuiltinTextureType())); } Index: clang/lib/Sema/SemaCUDA.cpp =================================================================== --- clang/lib/Sema/SemaCUDA.cpp +++ clang/lib/Sema/SemaCUDA.cpp @@ -516,7 +516,7 @@ return; const Expr *Init = VD->getInit(); if (VD->hasAttr() || VD->hasAttr() || - VD->hasAttr()) { + VD->hasAttr() || VD->hasAttr()) { if (LangOpts.GPUAllowDeviceInit) return; bool AllowedInit = false; @@ -527,7 +527,8 @@ // constructor according to CUDA rules. This deviates from NVCC, // but allows us to handle things like constexpr constructors. if (!AllowedInit && - (VD->hasAttr() || VD->hasAttr())) { + (VD->hasAttr() || VD->hasAttr() || + VD->hasAttr())) { auto *Init = VD->getInit(); AllowedInit = ((VD->getType()->isDependentType() || Init->isValueDependent()) && Index: clang/lib/Sema/SemaDeclAttr.cpp =================================================================== --- clang/lib/Sema/SemaDeclAttr.cpp +++ clang/lib/Sema/SemaDeclAttr.cpp @@ -7724,6 +7724,10 @@ case ParsedAttr::AT_CUDAHost: handleSimpleAttributeWithExclusions(S, D, AL); break; + case ParsedAttr::AT_HIPManaged: + handleSimpleAttributeWithExclusions(S, D, + AL); + break; case ParsedAttr::AT_CUDADeviceBuiltinSurfaceType: handleSimpleAttributeWithExclusions(S, D, Index: clang/lib/Sema/SemaExpr.cpp =================================================================== --- clang/lib/Sema/SemaExpr.cpp +++ clang/lib/Sema/SemaExpr.cpp @@ -364,6 +364,7 @@ const auto *VD = dyn_cast(D); if (VD && VD->hasGlobalStorage() && !VD->hasAttr() && !VD->hasAttr() && !VD->hasAttr() && + !VD->hasAttr() && !VD->getType()->isCUDADeviceBuiltinSurfaceType() && !VD->getType()->isCUDADeviceBuiltinTextureType() && !VD->isConstexpr() && !VD->getType().isConstQualified()) Index: clang/test/CodeGenCUDA/Inputs/cuda.h =================================================================== --- clang/test/CodeGenCUDA/Inputs/cuda.h +++ 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 { Index: clang/test/CodeGenCUDA/managed-var.cu =================================================================== --- /dev/null +++ clang/test/CodeGenCUDA/managed-var.cu @@ -0,0 +1,94 @@ +// 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 x86_64-gnu-linux -std=c++11 \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck \ +// RUN: -check-prefixes=HOST %s + +#include "Inputs/cuda.h" + + +// Check a static device variable referenced by host function is externalized. +// DEV-DAG: @x = dso_local addrspace(1) externally_initialized global i32 undef +// HOST-DAG: @x = internal global i32 1 +// HOST-DAG: @x.managed = external global i32* +// 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(&(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) Index: clang/test/Misc/pragma-attribute-supported-attributes-list.test =================================================================== --- clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -64,6 +64,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) Index: llvm/include/llvm/IR/ReplaceConstant.h =================================================================== --- /dev/null +++ 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 Index: llvm/lib/IR/CMakeLists.txt =================================================================== --- llvm/lib/IR/CMakeLists.txt +++ llvm/lib/IR/CMakeLists.txt @@ -49,6 +49,7 @@ SafepointIRVerifier.cpp ProfileSummary.cpp PseudoProbe.cpp + ReplaceConstant.cpp Statepoint.cpp StructuralHash.cpp Type.cpp Index: llvm/lib/IR/ReplaceConstant.cpp =================================================================== --- /dev/null +++ llvm/lib/IR/ReplaceConstant.cpp @@ -0,0 +1,68 @@ +//===- 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 { +Instruction *createReplacementInstr(ConstantExpr *CE, Instruction *Instr) { + IRBuilder Builder(Instr); + unsigned OpCode = CE->getOpcode(); + switch (OpCode) { + case Instruction::GetElementPtr: { + SmallVector CEOpVec(CE->operands()); + ArrayRef CEOps(CEOpVec); + return dyn_cast( + Builder.CreateInBoundsGEP(cast(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( + 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( + Builder.CreateCast((Instruction::CastOps)OpCode, CE->getOperand(0), + CE->getType(), CE->getName())); + default: + llvm_unreachable("Unhandled constant expression!\n"); + } +} +} // namespace llvm Index: llvm/lib/Target/XCore/XCoreLowerThreadLocal.cpp =================================================================== --- llvm/lib/Target/XCore/XCoreLowerThreadLocal.cpp +++ 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 @@ return ConstantArray::get(NewType, Elements); } -static Instruction * -createReplacementInstr(ConstantExpr *CE, Instruction *Instr) { - IRBuilder Builder(Instr); - unsigned OpCode = CE->getOpcode(); - switch (OpCode) { - case Instruction::GetElementPtr: { - SmallVector CEOpVec(CE->operands()); - ArrayRef CEOps(CEOpVec); - return dyn_cast(Builder.CreateInBoundsGEP( - cast(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( - 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( - 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 {