CallInst *IRBuilderBase::
CreateMemSet(Value *Ptr, Value *Val, Value *Size, unsigned Align,
bool isVolatile, MDNode *TBAATag, MDNode *ScopeTag,
MDNode *NoAliasTag) {
Ptr = getCastedInt8PtrValue(Ptr);
Value *Ops[] = {Ptr, Val, Size, getInt1(isVolatile)};
Type *Tys[] = { Ptr->getType(), Size->getType() };
Module *M = BB->getParent()->getParent();
Value *TheFn = Intrinsic::getDeclaration(M, Intrinsic::memset, Tys);
CallInst *CI = createCallHelper(TheFn, Ops, this);
if (Align > 0)
cast<MemSetInst>(CI)->setDestAlignment(Align);
// Set the TBAA info if present.
if (TBAATag)
CI->setMetadata(LLVMContext::MD_tbaa, TBAATag);
if (ScopeTag)
CI->setMetadata(LLVMContext::MD_alias_scope, ScopeTag);
if (NoAliasTag)
CI->setMetadata(LLVMContext::MD_noalias, NoAliasTag);
return CI;
}
CallInst *IRBuilderBase::
CreateMemCpy(Value *Dst, Value *Src, Value *Size, unsigned Align,
bool isVolatile, MDNode *TBAATag, MDNode *TBAAStructTag,
MDNode *ScopeTag, MDNode *NoAliasTag) {
Dst = getCastedInt8PtrValue(Dst);
Src = getCastedInt8PtrValue(Src);
Value *Ops[] = { Dst, Src, Size, getInt32(Align), getInt1(isVolatile) };
Type *Tys[] = { Dst->getType(), Src->getType(), Size->getType() };
Module *M = BB->getParent()->getParent();
Value *TheFn = Intrinsic::getDeclaration(M, Intrinsic::memcpy, Tys);
CallInst *CI = createCallHelper(TheFn, Ops, this);
// Set the TBAA info if present.
if (TBAATag)
CI->setMetadata(LLVMContext::MD_tbaa, TBAATag);
// Set the TBAA Struct info if present.
if (TBAAStructTag)
CI->setMetadata(LLVMContext::MD_tbaa_struct, TBAAStructTag);
if (ScopeTag)
CI->setMetadata(LLVMContext::MD_alias_scope, ScopeTag);
if (NoAliasTag)
CI->setMetadata(LLVMContext::MD_noalias, NoAliasTag);
return CI;
}
CallInst *IRBuilderBase::
CreateMemMove(Value *Dst, unsigned DstAlign, Value *Src, unsigned SrcAlign,
Value *Size, bool isVolatile, MDNode *TBAATag, MDNode *ScopeTag,
MDNode *NoAliasTag) {
assert((DstAlign == 0 || isPowerOf2_32(DstAlign)) && "Must be 0 or a power of 2");
assert((SrcAlign == 0 || isPowerOf2_32(SrcAlign)) && "Must be 0 or a power of 2");
Dst = getCastedInt8PtrValue(Dst);
Src = getCastedInt8PtrValue(Src);
Value *Ops[] = {Dst, Src, Size, getInt1(isVolatile)};
Type *Tys[] = { Dst->getType(), Src->getType(), Size->getType() };
Module *M = BB->getParent()->getParent();
Value *TheFn = Intrinsic::getDeclaration(M, Intrinsic::memmove, Tys);
CallInst *CI = createCallHelper(TheFn, Ops, this);
auto *MMI = cast<MemMoveInst>(CI);
if (DstAlign > 0)
MMI->setDestAlignment(DstAlign);
if (SrcAlign > 0)
MMI->setSourceAlignment(SrcAlign);
// Set the TBAA info if present.
if (TBAATag)
CI->setMetadata(LLVMContext::MD_tbaa, TBAATag);
if (ScopeTag)
CI->setMetadata(LLVMContext::MD_alias_scope, ScopeTag);
if (NoAliasTag)
CI->setMetadata(LLVMContext::MD_noalias, NoAliasTag);
return CI;
}
Value *AMDGPUPromoteAlloca::getWorkitemID(IRBuilder<> &Builder, unsigned N) {
Intrinsic::ID IntrID = Intrinsic::ID::not_intrinsic;
switch (N) {
case 0:
IntrID = IsAMDGCN ? Intrinsic::amdgcn_workitem_id_x
: Intrinsic::r600_read_tidig_x;
break;
case 1:
IntrID = IsAMDGCN ? Intrinsic::amdgcn_workitem_id_y
: Intrinsic::r600_read_tidig_y;
break;
case 2:
IntrID = IsAMDGCN ? Intrinsic::amdgcn_workitem_id_z
: Intrinsic::r600_read_tidig_z;
break;
default:
llvm_unreachable("invalid dimension");
}
Function *WorkitemIdFn = Intrinsic::getDeclaration(Mod, IntrID);
CallInst *CI = Builder.CreateCall(WorkitemIdFn);
CI->setMetadata(LLVMContext::MD_range, MaxWorkGroupSizeRange);
return CI;
}
// Returns true if able to find a call instruction to mark
bool Builder::SetNamedMetaDataOnCallInstr(Instruction* inst, StringRef mdName)
{
CallInst* pCallInstr = dyn_cast<CallInst>(inst);
if (pCallInstr)
{
MDNode* N = MDNode::get(JM()->mContext, MDString::get(JM()->mContext, mdName));
pCallInstr->setMetadata(mdName, N);
return true;
}
else
{
// Follow use def chain back up
for (Use& u : inst->operands())
{
Instruction* srcInst = dyn_cast<Instruction>(u.get());
if (srcInst)
{
if (SetNamedMetaDataOnCallInstr(srcInst, mdName))
{
return true;
}
}
}
}
return false;
}
void InstrumentMemoryAccesses::instrument(Value *Pointer, Value *AccessSize,
Function *Check, Instruction &I) {
Builder->SetInsertPoint(&I);
Value *VoidPointer = Builder->CreatePointerCast(Pointer, VoidPtrTy);
CallInst *CI = Builder->CreateCall2(Check, VoidPointer, AccessSize);
// Copy debug information if it is present.
if (MDNode *MD = I.getMetadata("dbg"))
CI->setMetadata("dbg", MD);
}
CallInst *IRBuilderBase::CreateElementUnorderedAtomicMemCpy(
Value *Dst, unsigned DstAlign, Value *Src, unsigned SrcAlign, Value *Size,
uint32_t ElementSize, MDNode *TBAATag, MDNode *TBAAStructTag,
MDNode *ScopeTag, MDNode *NoAliasTag) {
assert(DstAlign >= ElementSize &&
"Pointer alignment must be at least element size");
assert(SrcAlign >= ElementSize &&
"Pointer alignment must be at least element size");
Dst = getCastedInt8PtrValue(Dst);
Src = getCastedInt8PtrValue(Src);
Value *Ops[] = {Dst, Src, Size, getInt32(ElementSize)};
Type *Tys[] = {Dst->getType(), Src->getType(), Size->getType()};
Module *M = BB->getParent()->getParent();
Value *TheFn = Intrinsic::getDeclaration(
M, Intrinsic::memcpy_element_unordered_atomic, Tys);
CallInst *CI = createCallHelper(TheFn, Ops, this);
// Set the alignment of the pointer args.
auto *AMCI = cast<AtomicMemCpyInst>(CI);
AMCI->setDestAlignment(DstAlign);
AMCI->setSourceAlignment(SrcAlign);
// Set the TBAA info if present.
if (TBAATag)
CI->setMetadata(LLVMContext::MD_tbaa, TBAATag);
// Set the TBAA Struct info if present.
if (TBAAStructTag)
CI->setMetadata(LLVMContext::MD_tbaa_struct, TBAAStructTag);
if (ScopeTag)
CI->setMetadata(LLVMContext::MD_alias_scope, ScopeTag);
if (NoAliasTag)
CI->setMetadata(LLVMContext::MD_noalias, NoAliasTag);
return CI;
}
//
// Method: visitGetElementPtrInst()
//
// Description:
// This method checks to see if the specified GEP is safe. If it cannot prove
// it safe, it then adds a run-time check for it.
//
void
InsertGEPChecks::visitGetElementPtrInst (GetElementPtrInst & GEP) {
//
// Don't insert a check if GEP only indexes into a structure and the
// user doesn't want to do structure index checking.
//
if (DisableStructChecks && indexesStructsOnly (&GEP)) {
return;
}
//
// Get the function in which the GEP instruction lives.
//
Value * PH = ConstantPointerNull::get (getVoidPtrType(GEP.getContext()));
BasicBlock::iterator InsertPt = &GEP;
++InsertPt;
Instruction * ResultPtr = castTo (&GEP,
getVoidPtrType(GEP.getContext()),
GEP.getName() + ".cast",
InsertPt);
//
// Make this an actual cast instruction; it will make it easier to update
// DSA.
//
Value * SrcPtr = castTo (GEP.getPointerOperand(),
getVoidPtrType(GEP.getContext()),
GEP.getName()+".cast",
InsertPt);
//
// Create the call to the run-time check.
//
std::vector<Value *> args(1, PH);
args.push_back (SrcPtr);
args.push_back (ResultPtr);
CallInst * CI = CallInst::Create (PoolCheckArrayUI, args, "", InsertPt);
//
// Add debugging info metadata to the run-time check.
//
if (MDNode * MD = GEP.getMetadata ("dbg"))
CI->setMetadata ("dbg", MD);
//
// Update the statistics.
//
++GEPChecks;
return;
}
CallInst *IRBuilderBase::
CreateMemSet(Value *Ptr, Value *Val, Value *Size, unsigned Align,
bool isVolatile, MDNode *TBAATag) {
Ptr = getCastedInt8PtrValue(Ptr);
Value *Ops[] = { Ptr, Val, Size, getInt32(Align), getInt1(isVolatile) };
const Type *Tys[] = { Ptr->getType(), Size->getType() };
Module *M = BB->getParent()->getParent();
Value *TheFn = Intrinsic::getDeclaration(M, Intrinsic::memset, Tys, 2);
CallInst *CI = createCallHelper(TheFn, Ops, 5, this);
// Set the TBAA info if present.
if (TBAATag)
CI->setMetadata(LLVMContext::MD_tbaa, TBAATag);
return CI;
}
/// createFastCheck - create the fast memory safety check given the old check
/// and the corresponding object and its size.
///
void ExactCheckOpt::createFastCheck(CheckInfoType* Info, CallInst *CI,
Value *ObjPtr, Value *ObjSize) {
Module &M = *CI->getParent()->getParent()->getParent();
// Get a pointer to the fast check function.
CheckInfoType *FastInfo = Info->FastVersionInfo;
Function *FastFn = FastInfo->getFunction(M);
assert(FastFn && "The fast check function should be defined.");
// Copy the old arguments to preserve extra arguments in fixed positions.
SmallVector <Value*, 8> Args(FastFn->arg_size());
assert(FastFn->arg_size() >= CI->getNumArgOperands());
for (unsigned i = 0, N = CI->getNumArgOperands(); i < N; ++i)
Args[i] = CI->getArgOperand(i);
// Set the known arguments to right values.
Args[FastInfo->PtrArgNo] = CI->getArgOperand(Info->PtrArgNo);
Args[FastInfo->ObjArgNo] = ObjPtr;
Args[FastInfo->ObjSizeArgNo] = ObjSize;
if (Info->isMemoryCheck())
Args[FastInfo->SizeArgNo] = CI->getArgOperand(Info->SizeArgNo);
else // must be a gep check
Args[FastInfo->DestPtrArgNo] = CI->getArgOperand(Info->DestPtrArgNo);
// Create the call just before the old call.
IRBuilder<> Builder(CI);
CallInst *FastCI = Builder.CreateCall(FastFn, Args);
// Copy the debug information if it is present.
if (MDNode *MD = CI->getMetadata("dbg"))
FastCI->setMetadata("dbg", MD);
if (Info->isGEPCheck())
CI->replaceAllUsesWith(FastCI);
}
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);
}
