bool AMDGPUCodeGenPrepare::visitLoadInst(LoadInst &I) {
if (!WidenLoads)
return false;
if ((I.getPointerAddressSpace() == AMDGPUASI.CONSTANT_ADDRESS ||
I.getPointerAddressSpace() == AMDGPUASI.CONSTANT_ADDRESS_32BIT) &&
canWidenScalarExtLoad(I)) {
IRBuilder<> Builder(&I);
Builder.SetCurrentDebugLocation(I.getDebugLoc());
Type *I32Ty = Builder.getInt32Ty();
Type *PT = PointerType::get(I32Ty, I.getPointerAddressSpace());
Value *BitCast= Builder.CreateBitCast(I.getPointerOperand(), PT);
LoadInst *WidenLoad = Builder.CreateLoad(BitCast);
WidenLoad->copyMetadata(I);
// If we have range metadata, we need to convert the type, and not make
// assumptions about the high bits.
if (auto *Range = WidenLoad->getMetadata(LLVMContext::MD_range)) {
ConstantInt *Lower =
mdconst::extract<ConstantInt>(Range->getOperand(0));
if (Lower->getValue().isNullValue()) {
WidenLoad->setMetadata(LLVMContext::MD_range, nullptr);
} else {
Metadata *LowAndHigh[] = {
ConstantAsMetadata::get(ConstantInt::get(I32Ty, Lower->getValue().zext(32))),
// Don't make assumptions about the high bits.
ConstantAsMetadata::get(ConstantInt::get(I32Ty, 0))
};
WidenLoad->setMetadata(LLVMContext::MD_range,
MDNode::get(Mod->getContext(), LowAndHigh));
}
}
int TySize = Mod->getDataLayout().getTypeSizeInBits(I.getType());
Type *IntNTy = Builder.getIntNTy(TySize);
Value *ValTrunc = Builder.CreateTrunc(WidenLoad, IntNTy);
Value *ValOrig = Builder.CreateBitCast(ValTrunc, I.getType());
I.replaceAllUsesWith(ValOrig);
I.eraseFromParent();
return true;
}
return false;
}
void SanitizerCoverageModule::InjectCoverageAtBlock(Function &F,
BasicBlock &BB) {
BasicBlock::iterator IP = BB.getFirstInsertionPt(), BE = BB.end();
// Skip static allocas at the top of the entry block so they don't become
// dynamic when we split the block. If we used our optimized stack layout,
// then there will only be one alloca and it will come first.
for (; IP != BE; ++IP) {
AllocaInst *AI = dyn_cast<AllocaInst>(IP);
if (!AI || !AI->isStaticAlloca())
break;
}
bool IsEntryBB = &BB == &F.getEntryBlock();
DebugLoc EntryLoc =
IsEntryBB ? IP->getDebugLoc().getFnDebugLoc(*C) : IP->getDebugLoc();
IRBuilder<> IRB(IP);
IRB.SetCurrentDebugLocation(EntryLoc);
SmallVector<Value *, 1> Indices;
Value *GuardP = IRB.CreateAdd(
IRB.CreatePointerCast(GuardArray, IntptrTy),
ConstantInt::get(IntptrTy, (1 + SanCovFunction->getNumUses()) * 4));
Type *Int32PtrTy = PointerType::getUnqual(IRB.getInt32Ty());
GuardP = IRB.CreateIntToPtr(GuardP, Int32PtrTy);
LoadInst *Load = IRB.CreateLoad(GuardP);
Load->setAtomic(Monotonic);
Load->setAlignment(4);
Load->setMetadata(F.getParent()->getMDKindID("nosanitize"),
MDNode::get(*C, None));
Value *Cmp = IRB.CreateICmpSGE(Constant::getNullValue(Load->getType()), Load);
Instruction *Ins = SplitBlockAndInsertIfThen(
Cmp, IP, false, MDBuilder(*C).createBranchWeights(1, 100000));
IRB.SetInsertPoint(Ins);
IRB.SetCurrentDebugLocation(EntryLoc);
// __sanitizer_cov gets the PC of the instruction using GET_CALLER_PC.
IRB.CreateCall(SanCovFunction, GuardP);
IRB.CreateCall(EmptyAsm); // Avoids callback merge.
if (ClExperimentalTracing) {
// Experimental support for tracing.
// Insert a callback with the same guard variable as used for coverage.
IRB.SetInsertPoint(IP);
IRB.CreateCall(IsEntryBB ? SanCovTraceEnter : SanCovTraceBB, GuardP);
}
}
/// \brief Helper to combine a load to a new type.
///
/// This just does the work of combining a load to a new type. It handles
/// metadata, etc., and returns the new instruction. The \c NewTy should be the
/// loaded *value* type. This will convert it to a pointer, cast the operand to
/// that pointer type, load it, etc.
///
/// Note that this will create all of the instructions with whatever insert
/// point the \c InstCombiner currently is using.
static LoadInst *combineLoadToNewType(InstCombiner &IC, LoadInst &LI, Type *NewTy) {
Value *Ptr = LI.getPointerOperand();
unsigned AS = LI.getPointerAddressSpace();
SmallVector<std::pair<unsigned, MDNode *>, 8> MD;
LI.getAllMetadata(MD);
LoadInst *NewLoad = IC.Builder->CreateAlignedLoad(
IC.Builder->CreateBitCast(Ptr, NewTy->getPointerTo(AS)),
LI.getAlignment(), LI.getName());
for (const auto &MDPair : MD) {
unsigned ID = MDPair.first;
MDNode *N = MDPair.second;
// Note, essentially every kind of metadata should be preserved here! This
// routine is supposed to clone a load instruction changing *only its type*.
// The only metadata it makes sense to drop is metadata which is invalidated
// when the pointer type changes. This should essentially never be the case
// in LLVM, but we explicitly switch over only known metadata to be
// conservatively correct. If you are adding metadata to LLVM which pertains
// to loads, you almost certainly want to add it here.
switch (ID) {
case LLVMContext::MD_dbg:
case LLVMContext::MD_tbaa:
case LLVMContext::MD_prof:
case LLVMContext::MD_fpmath:
case LLVMContext::MD_tbaa_struct:
case LLVMContext::MD_invariant_load:
case LLVMContext::MD_alias_scope:
case LLVMContext::MD_noalias:
case LLVMContext::MD_nontemporal:
case LLVMContext::MD_mem_parallel_loop_access:
case LLVMContext::MD_nonnull:
// All of these directly apply.
NewLoad->setMetadata(ID, N);
break;
case LLVMContext::MD_range:
// FIXME: It would be nice to propagate this in some way, but the type
// conversions make it hard.
break;
}
}
return NewLoad;
}
void Closure::unpack_struct(Scope<Value *> &dst,
Value *src,
IRBuilder<> *builder) {
// src should be a pointer to a struct of the type returned by build_type
int idx = 0;
LLVMContext &context = builder->getContext();
vector<string> nm = names();
for (size_t i = 0; i < nm.size(); i++) {
Value *ptr = builder->CreateConstInBoundsGEP2_32(src, 0, idx++);
LoadInst *load = builder->CreateLoad(ptr);
if (load->getType()->isPointerTy()) {
// Give it a unique type so that tbaa tells llvm that this can't alias anything
load->setMetadata("tbaa", MDNode::get(context,
vec<Value *>(MDString::get(context, nm[i]))));
}
dst.push(nm[i], load);
load->setName(nm[i]);
}
}
std::pair<Value *, Value *>
AMDGPUPromoteAlloca::getLocalSizeYZ(IRBuilder<> &Builder) {
if (!IsAMDHSA) {
Function *LocalSizeYFn
= Intrinsic::getDeclaration(Mod, Intrinsic::r600_read_local_size_y);
Function *LocalSizeZFn
= Intrinsic::getDeclaration(Mod, Intrinsic::r600_read_local_size_z);
CallInst *LocalSizeY = Builder.CreateCall(LocalSizeYFn, {});
CallInst *LocalSizeZ = Builder.CreateCall(LocalSizeZFn, {});
LocalSizeY->setMetadata(LLVMContext::MD_range, MaxWorkGroupSizeRange);
LocalSizeZ->setMetadata(LLVMContext::MD_range, MaxWorkGroupSizeRange);
return std::make_pair(LocalSizeY, LocalSizeZ);
}
// We must read the size out of the dispatch pointer.
assert(IsAMDGCN);
// We are indexing into this struct, and want to extract the workgroup_size_*
// fields.
//
// typedef struct hsa_kernel_dispatch_packet_s {
// uint16_t header;
// uint16_t setup;
// uint16_t workgroup_size_x ;
// uint16_t workgroup_size_y;
// uint16_t workgroup_size_z;
// uint16_t reserved0;
// uint32_t grid_size_x ;
// uint32_t grid_size_y ;
// uint32_t grid_size_z;
//
// uint32_t private_segment_size;
// uint32_t group_segment_size;
// uint64_t kernel_object;
//
// #ifdef HSA_LARGE_MODEL
// void *kernarg_address;
// #elif defined HSA_LITTLE_ENDIAN
// void *kernarg_address;
// uint32_t reserved1;
// #else
// uint32_t reserved1;
// void *kernarg_address;
// #endif
// uint64_t reserved2;
// hsa_signal_t completion_signal; // uint64_t wrapper
// } hsa_kernel_dispatch_packet_t
//
Function *DispatchPtrFn
= Intrinsic::getDeclaration(Mod, Intrinsic::amdgcn_dispatch_ptr);
CallInst *DispatchPtr = Builder.CreateCall(DispatchPtrFn, {});
DispatchPtr->addAttribute(AttributeSet::ReturnIndex, Attribute::NoAlias);
DispatchPtr->addAttribute(AttributeSet::ReturnIndex, Attribute::NonNull);
// Size of the dispatch packet struct.
DispatchPtr->addDereferenceableAttr(AttributeSet::ReturnIndex, 64);
Type *I32Ty = Type::getInt32Ty(Mod->getContext());
Value *CastDispatchPtr = Builder.CreateBitCast(
DispatchPtr, PointerType::get(I32Ty, AMDGPUAS::CONSTANT_ADDRESS));
// We could do a single 64-bit load here, but it's likely that the basic
// 32-bit and extract sequence is already present, and it is probably easier
// to CSE this. The loads should be mergable later anyway.
Value *GEPXY = Builder.CreateConstInBoundsGEP1_64(CastDispatchPtr, 1);
LoadInst *LoadXY = Builder.CreateAlignedLoad(GEPXY, 4);
Value *GEPZU = Builder.CreateConstInBoundsGEP1_64(CastDispatchPtr, 2);
LoadInst *LoadZU = Builder.CreateAlignedLoad(GEPZU, 4);
MDNode *MD = llvm::MDNode::get(Mod->getContext(), None);
LoadXY->setMetadata(LLVMContext::MD_invariant_load, MD);
LoadZU->setMetadata(LLVMContext::MD_invariant_load, MD);
LoadZU->setMetadata(LLVMContext::MD_range, MaxWorkGroupSizeRange);
// Extract y component. Upper half of LoadZU should be zero already.
Value *Y = Builder.CreateLShr(LoadXY, 16);
return std::make_pair(Y, LoadZU);
}
/// \brief Helper to combine a load to a new type.
///
/// This just does the work of combining a load to a new type. It handles
/// metadata, etc., and returns the new instruction. The \c NewTy should be the
/// loaded *value* type. This will convert it to a pointer, cast the operand to
/// that pointer type, load it, etc.
///
/// Note that this will create all of the instructions with whatever insert
/// point the \c InstCombiner currently is using.
static LoadInst *combineLoadToNewType(InstCombiner &IC, LoadInst &LI, Type *NewTy,
const Twine &Suffix = "") {
Value *Ptr = LI.getPointerOperand();
unsigned AS = LI.getPointerAddressSpace();
SmallVector<std::pair<unsigned, MDNode *>, 8> MD;
LI.getAllMetadata(MD);
LoadInst *NewLoad = IC.Builder->CreateAlignedLoad(
IC.Builder->CreateBitCast(Ptr, NewTy->getPointerTo(AS)),
LI.getAlignment(), LI.getName() + Suffix);
MDBuilder MDB(NewLoad->getContext());
for (const auto &MDPair : MD) {
unsigned ID = MDPair.first;
MDNode *N = MDPair.second;
// Note, essentially every kind of metadata should be preserved here! This
// routine is supposed to clone a load instruction changing *only its type*.
// The only metadata it makes sense to drop is metadata which is invalidated
// when the pointer type changes. This should essentially never be the case
// in LLVM, but we explicitly switch over only known metadata to be
// conservatively correct. If you are adding metadata to LLVM which pertains
// to loads, you almost certainly want to add it here.
switch (ID) {
case LLVMContext::MD_dbg:
case LLVMContext::MD_tbaa:
case LLVMContext::MD_prof:
case LLVMContext::MD_fpmath:
case LLVMContext::MD_tbaa_struct:
case LLVMContext::MD_invariant_load:
case LLVMContext::MD_alias_scope:
case LLVMContext::MD_noalias:
case LLVMContext::MD_nontemporal:
case LLVMContext::MD_mem_parallel_loop_access:
// All of these directly apply.
NewLoad->setMetadata(ID, N);
break;
case LLVMContext::MD_nonnull:
// This only directly applies if the new type is also a pointer.
if (NewTy->isPointerTy()) {
NewLoad->setMetadata(ID, N);
break;
}
// If it's integral now, translate it to !range metadata.
if (NewTy->isIntegerTy()) {
auto *ITy = cast<IntegerType>(NewTy);
auto *NullInt = ConstantExpr::getPtrToInt(
ConstantPointerNull::get(cast<PointerType>(Ptr->getType())), ITy);
auto *NonNullInt =
ConstantExpr::getAdd(NullInt, ConstantInt::get(ITy, 1));
NewLoad->setMetadata(LLVMContext::MD_range,
MDB.createRange(NonNullInt, NullInt));
}
break;
case LLVMContext::MD_align:
case LLVMContext::MD_dereferenceable:
case LLVMContext::MD_dereferenceable_or_null:
// These only directly apply if the new type is also a pointer.
if (NewTy->isPointerTy())
NewLoad->setMetadata(ID, N);
break;
case LLVMContext::MD_range:
// FIXME: It would be nice to propagate this in some way, but the type
// conversions make it hard. If the new type is a pointer, we could
// translate it to !nonnull metadata.
break;
}
}
return NewLoad;
}
/// DoPromotion - This method actually performs the promotion of the specified
/// arguments, and returns the new function. At this point, we know that it's
/// safe to do so.
CallGraphNode *ArgPromotion::DoPromotion(Function *F,
SmallPtrSet<Argument*, 8> &ArgsToPromote,
SmallPtrSet<Argument*, 8> &ByValArgsToTransform) {
// Start by computing a new prototype for the function, which is the same as
// the old function, but has modified arguments.
const FunctionType *FTy = F->getFunctionType();
std::vector<const Type*> Params;
typedef std::set<IndicesVector> ScalarizeTable;
// ScalarizedElements - If we are promoting a pointer that has elements
// accessed out of it, keep track of which elements are accessed so that we
// can add one argument for each.
//
// Arguments that are directly loaded will have a zero element value here, to
// handle cases where there are both a direct load and GEP accesses.
//
std::map<Argument*, ScalarizeTable> ScalarizedElements;
// OriginalLoads - Keep track of a representative load instruction from the
// original function so that we can tell the alias analysis implementation
// what the new GEP/Load instructions we are inserting look like.
std::map<IndicesVector, LoadInst*> OriginalLoads;
// Attributes - Keep track of the parameter attributes for the arguments
// that we are *not* promoting. For the ones that we do promote, the parameter
// attributes are lost
SmallVector<AttributeWithIndex, 8> AttributesVec;
const AttrListPtr &PAL = F->getAttributes();
// Add any return attributes.
if (Attributes attrs = PAL.getRetAttributes())
AttributesVec.push_back(AttributeWithIndex::get(0, attrs));
// First, determine the new argument list
unsigned ArgIndex = 1;
for (Function::arg_iterator I = F->arg_begin(), E = F->arg_end(); I != E;
++I, ++ArgIndex) {
if (ByValArgsToTransform.count(I)) {
// Simple byval argument? Just add all the struct element types.
const Type *AgTy = cast<PointerType>(I->getType())->getElementType();
const StructType *STy = cast<StructType>(AgTy);
for (unsigned i = 0, e = STy->getNumElements(); i != e; ++i)
Params.push_back(STy->getElementType(i));
++NumByValArgsPromoted;
} else if (!ArgsToPromote.count(I)) {
// Unchanged argument
Params.push_back(I->getType());
if (Attributes attrs = PAL.getParamAttributes(ArgIndex))
AttributesVec.push_back(AttributeWithIndex::get(Params.size(), attrs));
} else if (I->use_empty()) {
// Dead argument (which are always marked as promotable)
++NumArgumentsDead;
} else {
// Okay, this is being promoted. This means that the only uses are loads
// or GEPs which are only used by loads
// In this table, we will track which indices are loaded from the argument
// (where direct loads are tracked as no indices).
ScalarizeTable &ArgIndices = ScalarizedElements[I];
for (Value::use_iterator UI = I->use_begin(), E = I->use_end(); UI != E;
++UI) {
Instruction *User = cast<Instruction>(*UI);
assert(isa<LoadInst>(User) || isa<GetElementPtrInst>(User));
IndicesVector Indices;
Indices.reserve(User->getNumOperands() - 1);
// Since loads will only have a single operand, and GEPs only a single
// non-index operand, this will record direct loads without any indices,
// and gep+loads with the GEP indices.
for (User::op_iterator II = User->op_begin() + 1, IE = User->op_end();
II != IE; ++II)
Indices.push_back(cast<ConstantInt>(*II)->getSExtValue());
// GEPs with a single 0 index can be merged with direct loads
if (Indices.size() == 1 && Indices.front() == 0)
Indices.clear();
ArgIndices.insert(Indices);
LoadInst *OrigLoad;
if (LoadInst *L = dyn_cast<LoadInst>(User))
OrigLoad = L;
else
// Take any load, we will use it only to update Alias Analysis
OrigLoad = cast<LoadInst>(User->use_back());
OriginalLoads[Indices] = OrigLoad;
}
// Add a parameter to the function for each element passed in.
for (ScalarizeTable::iterator SI = ArgIndices.begin(),
E = ArgIndices.end(); SI != E; ++SI) {
// not allowed to dereference ->begin() if size() is 0
Params.push_back(GetElementPtrInst::getIndexedType(I->getType(),
SI->begin(),
SI->end()));
assert(Params.back());
}
if (ArgIndices.size() == 1 && ArgIndices.begin()->empty())
++NumArgumentsPromoted;
else
++NumAggregatesPromoted;
}
}
// Add any function attributes.
if (Attributes attrs = PAL.getFnAttributes())
AttributesVec.push_back(AttributeWithIndex::get(~0, attrs));
const Type *RetTy = FTy->getReturnType();
// Work around LLVM bug PR56: the CWriter cannot emit varargs functions which
// have zero fixed arguments.
bool ExtraArgHack = false;
if (Params.empty() && FTy->isVarArg()) {
ExtraArgHack = true;
Params.push_back(Type::getInt32Ty(F->getContext()));
}
// Construct the new function type using the new arguments.
FunctionType *NFTy = FunctionType::get(RetTy, Params, FTy->isVarArg());
// Create the new function body and insert it into the module.
Function *NF = Function::Create(NFTy, F->getLinkage(), F->getName());
NF->copyAttributesFrom(F);
DEBUG(dbgs() << "ARG PROMOTION: Promoting to:" << *NF << "\n"
<< "From: " << *F);
// Recompute the parameter attributes list based on the new arguments for
// the function.
NF->setAttributes(AttrListPtr::get(AttributesVec.begin(),
AttributesVec.end()));
AttributesVec.clear();
F->getParent()->getFunctionList().insert(F, NF);
NF->takeName(F);
// Get the alias analysis information that we need to update to reflect our
// changes.
AliasAnalysis &AA = getAnalysis<AliasAnalysis>();
// Get the callgraph information that we need to update to reflect our
// changes.
CallGraph &CG = getAnalysis<CallGraph>();
// Get a new callgraph node for NF.
CallGraphNode *NF_CGN = CG.getOrInsertFunction(NF);
// Loop over all of the callers of the function, transforming the call sites
// to pass in the loaded pointers.
//
SmallVector<Value*, 16> Args;
while (!F->use_empty()) {
CallSite CS(F->use_back());
assert(CS.getCalledFunction() == F);
Instruction *Call = CS.getInstruction();
const AttrListPtr &CallPAL = CS.getAttributes();
// Add any return attributes.
if (Attributes attrs = CallPAL.getRetAttributes())
AttributesVec.push_back(AttributeWithIndex::get(0, attrs));
// Loop over the operands, inserting GEP and loads in the caller as
// appropriate.
CallSite::arg_iterator AI = CS.arg_begin();
ArgIndex = 1;
for (Function::arg_iterator I = F->arg_begin(), E = F->arg_end();
I != E; ++I, ++AI, ++ArgIndex)
if (!ArgsToPromote.count(I) && !ByValArgsToTransform.count(I)) {
Args.push_back(*AI); // Unmodified argument
if (Attributes Attrs = CallPAL.getParamAttributes(ArgIndex))
AttributesVec.push_back(AttributeWithIndex::get(Args.size(), Attrs));
} else if (ByValArgsToTransform.count(I)) {
// Emit a GEP and load for each element of the struct.
const Type *AgTy = cast<PointerType>(I->getType())->getElementType();
const StructType *STy = cast<StructType>(AgTy);
Value *Idxs[2] = {
ConstantInt::get(Type::getInt32Ty(F->getContext()), 0), 0 };
for (unsigned i = 0, e = STy->getNumElements(); i != e; ++i) {
Idxs[1] = ConstantInt::get(Type::getInt32Ty(F->getContext()), i);
Value *Idx = GetElementPtrInst::Create(*AI, Idxs, Idxs+2,
(*AI)->getName()+"."+utostr(i),
Call);
// TODO: Tell AA about the new values?
Args.push_back(new LoadInst(Idx, Idx->getName()+".val", Call));
}
} else if (!I->use_empty()) {
// Non-dead argument: insert GEPs and loads as appropriate.
ScalarizeTable &ArgIndices = ScalarizedElements[I];
// Store the Value* version of the indices in here, but declare it now
// for reuse.
std::vector<Value*> Ops;
for (ScalarizeTable::iterator SI = ArgIndices.begin(),
E = ArgIndices.end(); SI != E; ++SI) {
Value *V = *AI;
LoadInst *OrigLoad = OriginalLoads[*SI];
if (!SI->empty()) {
Ops.reserve(SI->size());
const Type *ElTy = V->getType();
for (IndicesVector::const_iterator II = SI->begin(),
IE = SI->end(); II != IE; ++II) {
// Use i32 to index structs, and i64 for others (pointers/arrays).
// This satisfies GEP constraints.
const Type *IdxTy = (ElTy->isStructTy() ?
Type::getInt32Ty(F->getContext()) :
Type::getInt64Ty(F->getContext()));
Ops.push_back(ConstantInt::get(IdxTy, *II));
// Keep track of the type we're currently indexing.
ElTy = cast<CompositeType>(ElTy)->getTypeAtIndex(*II);
}
// And create a GEP to extract those indices.
V = GetElementPtrInst::Create(V, Ops.begin(), Ops.end(),
V->getName()+".idx", Call);
Ops.clear();
AA.copyValue(OrigLoad->getOperand(0), V);
}
// Since we're replacing a load make sure we take the alignment
// of the previous load.
LoadInst *newLoad = new LoadInst(V, V->getName()+".val", Call);
newLoad->setAlignment(OrigLoad->getAlignment());
// Transfer the TBAA info too.
newLoad->setMetadata(LLVMContext::MD_tbaa,
OrigLoad->getMetadata(LLVMContext::MD_tbaa));
Args.push_back(newLoad);
AA.copyValue(OrigLoad, Args.back());
}
}
if (ExtraArgHack)
Args.push_back(Constant::getNullValue(Type::getInt32Ty(F->getContext())));
// Push any varargs arguments on the list.
for (; AI != CS.arg_end(); ++AI, ++ArgIndex) {
Args.push_back(*AI);
if (Attributes Attrs = CallPAL.getParamAttributes(ArgIndex))
AttributesVec.push_back(AttributeWithIndex::get(Args.size(), Attrs));
}
// Add any function attributes.
if (Attributes attrs = CallPAL.getFnAttributes())
AttributesVec.push_back(AttributeWithIndex::get(~0, attrs));
Instruction *New;
if (InvokeInst *II = dyn_cast<InvokeInst>(Call)) {
New = InvokeInst::Create(NF, II->getNormalDest(), II->getUnwindDest(),
Args.begin(), Args.end(), "", Call);
cast<InvokeInst>(New)->setCallingConv(CS.getCallingConv());
cast<InvokeInst>(New)->setAttributes(AttrListPtr::get(AttributesVec.begin(),
AttributesVec.end()));
} else {
New = CallInst::Create(NF, Args.begin(), Args.end(), "", Call);
cast<CallInst>(New)->setCallingConv(CS.getCallingConv());
cast<CallInst>(New)->setAttributes(AttrListPtr::get(AttributesVec.begin(),
AttributesVec.end()));
if (cast<CallInst>(Call)->isTailCall())
cast<CallInst>(New)->setTailCall();
}
Args.clear();
AttributesVec.clear();
// Update the alias analysis implementation to know that we are replacing
// the old call with a new one.
AA.replaceWithNewValue(Call, New);
// Update the callgraph to know that the callsite has been transformed.
CallGraphNode *CalleeNode = CG[Call->getParent()->getParent()];
CalleeNode->replaceCallEdge(Call, New, NF_CGN);
if (!Call->use_empty()) {
Call->replaceAllUsesWith(New);
New->takeName(Call);
}
// Finally, remove the old call from the program, reducing the use-count of
// F.
Call->eraseFromParent();
}
// Since we have now created the new function, splice the body of the old
// function right into the new function, leaving the old rotting hulk of the
// function empty.
NF->getBasicBlockList().splice(NF->begin(), F->getBasicBlockList());
// Loop over the argument list, transferring uses of the old arguments over to
// the new arguments, also transferring over the names as well.
//
for (Function::arg_iterator I = F->arg_begin(), E = F->arg_end(),
I2 = NF->arg_begin(); I != E; ++I) {
if (!ArgsToPromote.count(I) && !ByValArgsToTransform.count(I)) {
// If this is an unmodified argument, move the name and users over to the
// new version.
I->replaceAllUsesWith(I2);
I2->takeName(I);
AA.replaceWithNewValue(I, I2);
++I2;
continue;
}
if (ByValArgsToTransform.count(I)) {
// In the callee, we create an alloca, and store each of the new incoming
// arguments into the alloca.
Instruction *InsertPt = NF->begin()->begin();
// Just add all the struct element types.
const Type *AgTy = cast<PointerType>(I->getType())->getElementType();
Value *TheAlloca = new AllocaInst(AgTy, 0, "", InsertPt);
const StructType *STy = cast<StructType>(AgTy);
Value *Idxs[2] = {
ConstantInt::get(Type::getInt32Ty(F->getContext()), 0), 0 };
for (unsigned i = 0, e = STy->getNumElements(); i != e; ++i) {
Idxs[1] = ConstantInt::get(Type::getInt32Ty(F->getContext()), i);
Value *Idx =
GetElementPtrInst::Create(TheAlloca, Idxs, Idxs+2,
TheAlloca->getName()+"."+Twine(i),
InsertPt);
I2->setName(I->getName()+"."+Twine(i));
new StoreInst(I2++, Idx, InsertPt);
}
// Anything that used the arg should now use the alloca.
I->replaceAllUsesWith(TheAlloca);
TheAlloca->takeName(I);
AA.replaceWithNewValue(I, TheAlloca);
continue;
}
if (I->use_empty()) {
AA.deleteValue(I);
continue;
}
// Otherwise, if we promoted this argument, then all users are load
// instructions (or GEPs with only load users), and all loads should be
// using the new argument that we added.
ScalarizeTable &ArgIndices = ScalarizedElements[I];
while (!I->use_empty()) {
if (LoadInst *LI = dyn_cast<LoadInst>(I->use_back())) {
assert(ArgIndices.begin()->empty() &&
"Load element should sort to front!");
I2->setName(I->getName()+".val");
LI->replaceAllUsesWith(I2);
AA.replaceWithNewValue(LI, I2);
LI->eraseFromParent();
DEBUG(dbgs() << "*** Promoted load of argument '" << I->getName()
<< "' in function '" << F->getName() << "'\n");
} else {
GetElementPtrInst *GEP = cast<GetElementPtrInst>(I->use_back());
IndicesVector Operands;
Operands.reserve(GEP->getNumIndices());
for (User::op_iterator II = GEP->idx_begin(), IE = GEP->idx_end();
II != IE; ++II)
Operands.push_back(cast<ConstantInt>(*II)->getSExtValue());
// GEPs with a single 0 index can be merged with direct loads
if (Operands.size() == 1 && Operands.front() == 0)
Operands.clear();
Function::arg_iterator TheArg = I2;
for (ScalarizeTable::iterator It = ArgIndices.begin();
*It != Operands; ++It, ++TheArg) {
assert(It != ArgIndices.end() && "GEP not handled??");
}
std::string NewName = I->getName();
for (unsigned i = 0, e = Operands.size(); i != e; ++i) {
NewName += "." + utostr(Operands[i]);
}
NewName += ".val";
TheArg->setName(NewName);
DEBUG(dbgs() << "*** Promoted agg argument '" << TheArg->getName()
<< "' of function '" << NF->getName() << "'\n");
// All of the uses must be load instructions. Replace them all with
// the argument specified by ArgNo.
while (!GEP->use_empty()) {
LoadInst *L = cast<LoadInst>(GEP->use_back());
L->replaceAllUsesWith(TheArg);
AA.replaceWithNewValue(L, TheArg);
L->eraseFromParent();
}
AA.deleteValue(GEP);
GEP->eraseFromParent();
}
}
// Increment I2 past all of the arguments added for this promoted pointer.
for (unsigned i = 0, e = ArgIndices.size(); i != e; ++i)
++I2;
}
// Notify the alias analysis implementation that we inserted a new argument.
if (ExtraArgHack)
AA.copyValue(Constant::getNullValue(Type::getInt32Ty(F->getContext())),
NF->arg_begin());
// Tell the alias analysis that the old function is about to disappear.
AA.replaceWithNewValue(F, NF);
NF_CGN->stealCalledFunctionsFrom(CG[F]);
// Now that the old function is dead, delete it. If there is a dangling
// reference to the CallgraphNode, just leave the dead function around for
// someone else to nuke.
CallGraphNode *CGN = CG[F];
if (CGN->getNumReferences() == 0)
delete CG.removeFunctionFromModule(CGN);
else
F->setLinkage(Function::ExternalLinkage);
return NF_CGN;
}
bool AMDGPULowerKernelArguments::runOnFunction(Function &F) {
CallingConv::ID CC = F.getCallingConv();
if (CC != CallingConv::AMDGPU_KERNEL || F.arg_empty())
return false;
auto &TPC = getAnalysis<TargetPassConfig>();
const TargetMachine &TM = TPC.getTM<TargetMachine>();
const GCNSubtarget &ST = TM.getSubtarget<GCNSubtarget>(F);
LLVMContext &Ctx = F.getParent()->getContext();
const DataLayout &DL = F.getParent()->getDataLayout();
BasicBlock &EntryBlock = *F.begin();
IRBuilder<> Builder(&*EntryBlock.begin());
const unsigned KernArgBaseAlign = 16; // FIXME: Increase if necessary
const uint64_t BaseOffset = ST.getExplicitKernelArgOffset(F);
unsigned MaxAlign;
// FIXME: Alignment is broken broken with explicit arg offset.;
const uint64_t TotalKernArgSize = ST.getKernArgSegmentSize(F, MaxAlign);
if (TotalKernArgSize == 0)
return false;
CallInst *KernArgSegment =
Builder.CreateIntrinsic(Intrinsic::amdgcn_kernarg_segment_ptr, {}, {},
nullptr, F.getName() + ".kernarg.segment");
KernArgSegment->addAttribute(AttributeList::ReturnIndex, Attribute::NonNull);
KernArgSegment->addAttribute(AttributeList::ReturnIndex,
Attribute::getWithDereferenceableBytes(Ctx, TotalKernArgSize));
unsigned AS = KernArgSegment->getType()->getPointerAddressSpace();
uint64_t ExplicitArgOffset = 0;
for (Argument &Arg : F.args()) {
Type *ArgTy = Arg.getType();
unsigned Align = DL.getABITypeAlignment(ArgTy);
unsigned Size = DL.getTypeSizeInBits(ArgTy);
unsigned AllocSize = DL.getTypeAllocSize(ArgTy);
uint64_t EltOffset = alignTo(ExplicitArgOffset, Align) + BaseOffset;
ExplicitArgOffset = alignTo(ExplicitArgOffset, Align) + AllocSize;
if (Arg.use_empty())
continue;
if (PointerType *PT = dyn_cast<PointerType>(ArgTy)) {
// FIXME: Hack. We rely on AssertZext to be able to fold DS addressing
// modes on SI to know the high bits are 0 so pointer adds don't wrap. We
// can't represent this with range metadata because it's only allowed for
// integer types.
if ((PT->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS ||
PT->getAddressSpace() == AMDGPUAS::REGION_ADDRESS) &&
ST.getGeneration() == AMDGPUSubtarget::SOUTHERN_ISLANDS)
continue;
// FIXME: We can replace this with equivalent alias.scope/noalias
// metadata, but this appears to be a lot of work.
if (Arg.hasNoAliasAttr())
continue;
}
VectorType *VT = dyn_cast<VectorType>(ArgTy);
bool IsV3 = VT && VT->getNumElements() == 3;
bool DoShiftOpt = Size < 32 && !ArgTy->isAggregateType();
VectorType *V4Ty = nullptr;
int64_t AlignDownOffset = alignDown(EltOffset, 4);
int64_t OffsetDiff = EltOffset - AlignDownOffset;
unsigned AdjustedAlign = MinAlign(DoShiftOpt ? AlignDownOffset : EltOffset,
KernArgBaseAlign);
Value *ArgPtr;
Type *AdjustedArgTy;
if (DoShiftOpt) { // FIXME: Handle aggregate types
// Since we don't have sub-dword scalar loads, avoid doing an extload by
// loading earlier than the argument address, and extracting the relevant
// bits.
//
// Additionally widen any sub-dword load to i32 even if suitably aligned,
// so that CSE between different argument loads works easily.
ArgPtr = Builder.CreateConstInBoundsGEP1_64(
Builder.getInt8Ty(), KernArgSegment, AlignDownOffset,
Arg.getName() + ".kernarg.offset.align.down");
AdjustedArgTy = Builder.getInt32Ty();
} else {
ArgPtr = Builder.CreateConstInBoundsGEP1_64(
Builder.getInt8Ty(), KernArgSegment, EltOffset,
Arg.getName() + ".kernarg.offset");
AdjustedArgTy = ArgTy;
}
if (IsV3 && Size >= 32) {
V4Ty = VectorType::get(VT->getVectorElementType(), 4);
// Use the hack that clang uses to avoid SelectionDAG ruining v3 loads
AdjustedArgTy = V4Ty;
}
ArgPtr = Builder.CreateBitCast(ArgPtr, AdjustedArgTy->getPointerTo(AS),
ArgPtr->getName() + ".cast");
LoadInst *Load =
Builder.CreateAlignedLoad(AdjustedArgTy, ArgPtr, AdjustedAlign);
Load->setMetadata(LLVMContext::MD_invariant_load, MDNode::get(Ctx, {}));
MDBuilder MDB(Ctx);
if (isa<PointerType>(ArgTy)) {
if (Arg.hasNonNullAttr())
Load->setMetadata(LLVMContext::MD_nonnull, MDNode::get(Ctx, {}));
uint64_t DerefBytes = Arg.getDereferenceableBytes();
if (DerefBytes != 0) {
Load->setMetadata(
LLVMContext::MD_dereferenceable,
MDNode::get(Ctx,
MDB.createConstant(
ConstantInt::get(Builder.getInt64Ty(), DerefBytes))));
}
uint64_t DerefOrNullBytes = Arg.getDereferenceableOrNullBytes();
if (DerefOrNullBytes != 0) {
Load->setMetadata(
LLVMContext::MD_dereferenceable_or_null,
MDNode::get(Ctx,
MDB.createConstant(ConstantInt::get(Builder.getInt64Ty(),
DerefOrNullBytes))));
}
unsigned ParamAlign = Arg.getParamAlignment();
if (ParamAlign != 0) {
Load->setMetadata(
LLVMContext::MD_align,
MDNode::get(Ctx,
MDB.createConstant(ConstantInt::get(Builder.getInt64Ty(),
ParamAlign))));
}
}
// TODO: Convert noalias arg to !noalias
if (DoShiftOpt) {
Value *ExtractBits = OffsetDiff == 0 ?
Load : Builder.CreateLShr(Load, OffsetDiff * 8);
IntegerType *ArgIntTy = Builder.getIntNTy(Size);
Value *Trunc = Builder.CreateTrunc(ExtractBits, ArgIntTy);
Value *NewVal = Builder.CreateBitCast(Trunc, ArgTy,
Arg.getName() + ".load");
Arg.replaceAllUsesWith(NewVal);
} else if (IsV3) {
Value *Shuf = Builder.CreateShuffleVector(Load, UndefValue::get(V4Ty),
{0, 1, 2},
Arg.getName() + ".load");
Arg.replaceAllUsesWith(Shuf);
} else {
Load->setName(Arg.getName() + ".load");
Arg.replaceAllUsesWith(Load);
}
}
KernArgSegment->addAttribute(
AttributeList::ReturnIndex,
Attribute::getWithAlignment(Ctx, std::max(KernArgBaseAlign, MaxAlign)));
return true;
}
void Closure::unpack_struct(Scope<Value *> &dst,
llvm::Type *
#if LLVM_VERSION >= 37
type
#endif
,
Value *src,
IRBuilder<> *builder) {
// type, type of src should be a pointer to a struct of the type returned by build_type
int idx = 0;
LLVMContext &context = builder->getContext();
vector<string> nm = names();
for (size_t i = 0; i < nm.size(); i++) {
#if LLVM_VERSION >= 37
Value *ptr = builder->CreateConstInBoundsGEP2_32(type, src, 0, idx++);
#else
Value *ptr = builder->CreateConstInBoundsGEP2_32(src, 0, idx++);
#endif
LoadInst *load = builder->CreateLoad(ptr);
if (load->getType()->isPointerTy()) {
// Give it a unique type so that tbaa tells llvm that this can't alias anything
LLVMMDNodeArgumentType md_args[] = {MDString::get(context, nm[i])};
load->setMetadata("tbaa", MDNode::get(context, md_args));
}
dst.push(nm[i], load);
load->setName(nm[i]);
}
}