RValue CodeGenFunction::EmitAtomicExpr(AtomicExpr *E, llvm::Value *Dest) {
QualType AtomicTy = E->getPtr()->getType()->getPointeeType();
QualType MemTy = AtomicTy;
if (const AtomicType *AT = AtomicTy->getAs<AtomicType>())
MemTy = AT->getValueType();
CharUnits sizeChars = getContext().getTypeSizeInChars(AtomicTy);
uint64_t Size = sizeChars.getQuantity();
CharUnits alignChars = getContext().getTypeAlignInChars(AtomicTy);
unsigned Align = alignChars.getQuantity();
unsigned MaxInlineWidthInBits =
getTarget().getMaxAtomicInlineWidth();
bool UseLibcall = (Size != Align ||
getContext().toBits(sizeChars) > MaxInlineWidthInBits);
llvm::Value *Ptr, *Order, *OrderFail = 0, *Val1 = 0, *Val2 = 0;
Ptr = EmitScalarExpr(E->getPtr());
if (E->getOp() == AtomicExpr::AO__c11_atomic_init) {
assert(!Dest && "Init does not return a value");
LValue lvalue = LValue::MakeAddr(Ptr, AtomicTy, alignChars, getContext());
EmitAtomicInit(E->getVal1(), lvalue);
return RValue::get(0);
}
Order = EmitScalarExpr(E->getOrder());
switch (E->getOp()) {
case AtomicExpr::AO__c11_atomic_init:
llvm_unreachable("Already handled!");
case AtomicExpr::AO__c11_atomic_load:
case AtomicExpr::AO__atomic_load_n:
break;
case AtomicExpr::AO__atomic_load:
Dest = EmitScalarExpr(E->getVal1());
break;
case AtomicExpr::AO__atomic_store:
Val1 = EmitScalarExpr(E->getVal1());
break;
case AtomicExpr::AO__atomic_exchange:
Val1 = EmitScalarExpr(E->getVal1());
Dest = EmitScalarExpr(E->getVal2());
break;
case AtomicExpr::AO__c11_atomic_compare_exchange_strong:
case AtomicExpr::AO__c11_atomic_compare_exchange_weak:
case AtomicExpr::AO__atomic_compare_exchange_n:
case AtomicExpr::AO__atomic_compare_exchange:
Val1 = EmitScalarExpr(E->getVal1());
if (E->getOp() == AtomicExpr::AO__atomic_compare_exchange)
Val2 = EmitScalarExpr(E->getVal2());
else
Val2 = EmitValToTemp(*this, E->getVal2());
OrderFail = EmitScalarExpr(E->getOrderFail());
// Evaluate and discard the 'weak' argument.
if (E->getNumSubExprs() == 6)
EmitScalarExpr(E->getWeak());
break;
case AtomicExpr::AO__c11_atomic_fetch_add:
case AtomicExpr::AO__c11_atomic_fetch_sub:
if (MemTy->isPointerType()) {
// For pointer arithmetic, we're required to do a bit of math:
// adding 1 to an int* is not the same as adding 1 to a uintptr_t.
// ... but only for the C11 builtins. The GNU builtins expect the
// user to multiply by sizeof(T).
QualType Val1Ty = E->getVal1()->getType();
llvm::Value *Val1Scalar = EmitScalarExpr(E->getVal1());
CharUnits PointeeIncAmt =
getContext().getTypeSizeInChars(MemTy->getPointeeType());
Val1Scalar = Builder.CreateMul(Val1Scalar, CGM.getSize(PointeeIncAmt));
Val1 = CreateMemTemp(Val1Ty, ".atomictmp");
EmitStoreOfScalar(Val1Scalar, MakeAddrLValue(Val1, Val1Ty));
break;
}
// Fall through.
case AtomicExpr::AO__atomic_fetch_add:
case AtomicExpr::AO__atomic_fetch_sub:
case AtomicExpr::AO__atomic_add_fetch:
case AtomicExpr::AO__atomic_sub_fetch:
case AtomicExpr::AO__c11_atomic_store:
case AtomicExpr::AO__c11_atomic_exchange:
case AtomicExpr::AO__atomic_store_n:
case AtomicExpr::AO__atomic_exchange_n:
case AtomicExpr::AO__c11_atomic_fetch_and:
case AtomicExpr::AO__c11_atomic_fetch_or:
case AtomicExpr::AO__c11_atomic_fetch_xor:
case AtomicExpr::AO__atomic_fetch_and:
case AtomicExpr::AO__atomic_fetch_or:
case AtomicExpr::AO__atomic_fetch_xor:
case AtomicExpr::AO__atomic_fetch_nand:
case AtomicExpr::AO__atomic_and_fetch:
case AtomicExpr::AO__atomic_or_fetch:
case AtomicExpr::AO__atomic_xor_fetch:
case AtomicExpr::AO__atomic_nand_fetch:
Val1 = EmitValToTemp(*this, E->getVal1());
break;
}
if (!E->getType()->isVoidType() && !Dest)
Dest = CreateMemTemp(E->getType(), ".atomicdst");
// Use a library call. See: http://gcc.gnu.org/wiki/Atomic/GCCMM/LIbrary .
if (UseLibcall) {
bool UseOptimizedLibcall = false;
switch (E->getOp()) {
case AtomicExpr::AO__c11_atomic_fetch_add:
case AtomicExpr::AO__atomic_fetch_add:
case AtomicExpr::AO__c11_atomic_fetch_and:
case AtomicExpr::AO__atomic_fetch_and:
case AtomicExpr::AO__c11_atomic_fetch_or:
case AtomicExpr::AO__atomic_fetch_or:
case AtomicExpr::AO__c11_atomic_fetch_sub:
case AtomicExpr::AO__atomic_fetch_sub:
case AtomicExpr::AO__c11_atomic_fetch_xor:
case AtomicExpr::AO__atomic_fetch_xor:
// For these, only library calls for certain sizes exist.
UseOptimizedLibcall = true;
break;
default:
// Only use optimized library calls for sizes for which they exist.
if (Size == 1 || Size == 2 || Size == 4 || Size == 8)
UseOptimizedLibcall = true;
break;
}
CallArgList Args;
if (!UseOptimizedLibcall) {
// For non-optimized library calls, the size is the first parameter
Args.add(RValue::get(llvm::ConstantInt::get(SizeTy, Size)),
getContext().getSizeType());
}
// Atomic address is the first or second parameter
Args.add(RValue::get(EmitCastToVoidPtr(Ptr)),
getContext().VoidPtrTy);
std::string LibCallName;
QualType RetTy;
bool HaveRetTy = false;
switch (E->getOp()) {
// There is only one libcall for compare an exchange, because there is no
// optimisation benefit possible from a libcall version of a weak compare
// and exchange.
// bool __atomic_compare_exchange(size_t size, void *mem, void *expected,
// void *desired, int success, int failure)
// bool __atomic_compare_exchange_N(T *mem, T *expected, T desired,
// int success, int failure)
case AtomicExpr::AO__c11_atomic_compare_exchange_weak:
case AtomicExpr::AO__c11_atomic_compare_exchange_strong:
case AtomicExpr::AO__atomic_compare_exchange:
case AtomicExpr::AO__atomic_compare_exchange_n:
LibCallName = "__atomic_compare_exchange";
RetTy = getContext().BoolTy;
HaveRetTy = true;
Args.add(RValue::get(EmitCastToVoidPtr(Val1)),
getContext().VoidPtrTy);
AddDirectArgument(*this, Args, UseOptimizedLibcall, Val2, MemTy);
Args.add(RValue::get(Order),
getContext().IntTy);
Order = OrderFail;
break;
// void __atomic_exchange(size_t size, void *mem, void *val, void *return,
// int order)
// T __atomic_exchange_N(T *mem, T val, int order)
case AtomicExpr::AO__c11_atomic_exchange:
case AtomicExpr::AO__atomic_exchange_n:
case AtomicExpr::AO__atomic_exchange:
LibCallName = "__atomic_exchange";
AddDirectArgument(*this, Args, UseOptimizedLibcall, Val1, MemTy);
break;
// void __atomic_store(size_t size, void *mem, void *val, int order)
// void __atomic_store_N(T *mem, T val, int order)
case AtomicExpr::AO__c11_atomic_store:
case AtomicExpr::AO__atomic_store:
case AtomicExpr::AO__atomic_store_n:
LibCallName = "__atomic_store";
RetTy = getContext().VoidTy;
HaveRetTy = true;
AddDirectArgument(*this, Args, UseOptimizedLibcall, Val1, MemTy);
break;
// void __atomic_load(size_t size, void *mem, void *return, int order)
// T __atomic_load_N(T *mem, int order)
case AtomicExpr::AO__c11_atomic_load:
case AtomicExpr::AO__atomic_load:
case AtomicExpr::AO__atomic_load_n:
LibCallName = "__atomic_load";
break;
// T __atomic_fetch_add_N(T *mem, T val, int order)
case AtomicExpr::AO__c11_atomic_fetch_add:
case AtomicExpr::AO__atomic_fetch_add:
LibCallName = "__atomic_fetch_add";
AddDirectArgument(*this, Args, UseOptimizedLibcall, Val1, MemTy);
break;
// T __atomic_fetch_and_N(T *mem, T val, int order)
case AtomicExpr::AO__c11_atomic_fetch_and:
case AtomicExpr::AO__atomic_fetch_and:
LibCallName = "__atomic_fetch_and";
AddDirectArgument(*this, Args, UseOptimizedLibcall, Val1, MemTy);
break;
// T __atomic_fetch_or_N(T *mem, T val, int order)
case AtomicExpr::AO__c11_atomic_fetch_or:
case AtomicExpr::AO__atomic_fetch_or:
LibCallName = "__atomic_fetch_or";
AddDirectArgument(*this, Args, UseOptimizedLibcall, Val1, MemTy);
break;
// T __atomic_fetch_sub_N(T *mem, T val, int order)
case AtomicExpr::AO__c11_atomic_fetch_sub:
case AtomicExpr::AO__atomic_fetch_sub:
LibCallName = "__atomic_fetch_sub";
AddDirectArgument(*this, Args, UseOptimizedLibcall, Val1, MemTy);
break;
// T __atomic_fetch_xor_N(T *mem, T val, int order)
case AtomicExpr::AO__c11_atomic_fetch_xor:
case AtomicExpr::AO__atomic_fetch_xor:
LibCallName = "__atomic_fetch_xor";
AddDirectArgument(*this, Args, UseOptimizedLibcall, Val1, MemTy);
break;
default: return EmitUnsupportedRValue(E, "atomic library call");
}
// Optimized functions have the size in their name.
if (UseOptimizedLibcall)
LibCallName += "_" + llvm::utostr(Size);
// By default, assume we return a value of the atomic type.
if (!HaveRetTy) {
if (UseOptimizedLibcall) {
// Value is returned directly.
RetTy = MemTy;
} else {
// Value is returned through parameter before the order.
RetTy = getContext().VoidTy;
Args.add(RValue::get(EmitCastToVoidPtr(Dest)),
getContext().VoidPtrTy);
}
}
// order is always the last parameter
Args.add(RValue::get(Order),
getContext().IntTy);
const CGFunctionInfo &FuncInfo =
CGM.getTypes().arrangeFreeFunctionCall(RetTy, Args,
FunctionType::ExtInfo(), RequiredArgs::All);
llvm::FunctionType *FTy = CGM.getTypes().GetFunctionType(FuncInfo);
llvm::Constant *Func = CGM.CreateRuntimeFunction(FTy, LibCallName);
RValue Res = EmitCall(FuncInfo, Func, ReturnValueSlot(), Args);
if (!RetTy->isVoidType())
return Res;
if (E->getType()->isVoidType())
return RValue::get(0);
return convertTempToRValue(Dest, E->getType());
}
bool IsStore = E->getOp() == AtomicExpr::AO__c11_atomic_store ||
E->getOp() == AtomicExpr::AO__atomic_store ||
E->getOp() == AtomicExpr::AO__atomic_store_n;
bool IsLoad = E->getOp() == AtomicExpr::AO__c11_atomic_load ||
E->getOp() == AtomicExpr::AO__atomic_load ||
E->getOp() == AtomicExpr::AO__atomic_load_n;
llvm::Type *IPtrTy =
llvm::IntegerType::get(getLLVMContext(), Size * 8)->getPointerTo();
llvm::Value *OrigDest = Dest;
Ptr = Builder.CreateBitCast(Ptr, IPtrTy);
if (Val1) Val1 = Builder.CreateBitCast(Val1, IPtrTy);
if (Val2) Val2 = Builder.CreateBitCast(Val2, IPtrTy);
if (Dest && !E->isCmpXChg()) Dest = Builder.CreateBitCast(Dest, IPtrTy);
if (isa<llvm::ConstantInt>(Order)) {
int ord = cast<llvm::ConstantInt>(Order)->getZExtValue();
switch (ord) {
case AO_ABI_memory_order_relaxed:
EmitAtomicOp(*this, E, Dest, Ptr, Val1, Val2, Size, Align,
llvm::Monotonic);
break;
case AO_ABI_memory_order_consume:
case AO_ABI_memory_order_acquire:
if (IsStore)
break; // Avoid crashing on code with undefined behavior
EmitAtomicOp(*this, E, Dest, Ptr, Val1, Val2, Size, Align,
llvm::Acquire);
break;
case AO_ABI_memory_order_release:
if (IsLoad)
break; // Avoid crashing on code with undefined behavior
EmitAtomicOp(*this, E, Dest, Ptr, Val1, Val2, Size, Align,
llvm::Release);
break;
case AO_ABI_memory_order_acq_rel:
if (IsLoad || IsStore)
break; // Avoid crashing on code with undefined behavior
EmitAtomicOp(*this, E, Dest, Ptr, Val1, Val2, Size, Align,
llvm::AcquireRelease);
break;
case AO_ABI_memory_order_seq_cst:
EmitAtomicOp(*this, E, Dest, Ptr, Val1, Val2, Size, Align,
llvm::SequentiallyConsistent);
break;
default: // invalid order
// We should not ever get here normally, but it's hard to
// enforce that in general.
break;
}
if (E->getType()->isVoidType())
return RValue::get(0);
return convertTempToRValue(OrigDest, E->getType());
}
// Long case, when Order isn't obviously constant.
// Create all the relevant BB's
llvm::BasicBlock *MonotonicBB = 0, *AcquireBB = 0, *ReleaseBB = 0,
*AcqRelBB = 0, *SeqCstBB = 0;
MonotonicBB = createBasicBlock("monotonic", CurFn);
if (!IsStore)
AcquireBB = createBasicBlock("acquire", CurFn);
if (!IsLoad)
ReleaseBB = createBasicBlock("release", CurFn);
if (!IsLoad && !IsStore)
AcqRelBB = createBasicBlock("acqrel", CurFn);
SeqCstBB = createBasicBlock("seqcst", CurFn);
llvm::BasicBlock *ContBB = createBasicBlock("atomic.continue", CurFn);
// Create the switch for the split
// MonotonicBB is arbitrarily chosen as the default case; in practice, this
// doesn't matter unless someone is crazy enough to use something that
// doesn't fold to a constant for the ordering.
Order = Builder.CreateIntCast(Order, Builder.getInt32Ty(), false);
llvm::SwitchInst *SI = Builder.CreateSwitch(Order, MonotonicBB);
// Emit all the different atomics
Builder.SetInsertPoint(MonotonicBB);
EmitAtomicOp(*this, E, Dest, Ptr, Val1, Val2, Size, Align,
llvm::Monotonic);
Builder.CreateBr(ContBB);
if (!IsStore) {
Builder.SetInsertPoint(AcquireBB);
EmitAtomicOp(*this, E, Dest, Ptr, Val1, Val2, Size, Align,
llvm::Acquire);
Builder.CreateBr(ContBB);
SI->addCase(Builder.getInt32(1), AcquireBB);
SI->addCase(Builder.getInt32(2), AcquireBB);
}
if (!IsLoad) {
Builder.SetInsertPoint(ReleaseBB);
EmitAtomicOp(*this, E, Dest, Ptr, Val1, Val2, Size, Align,
llvm::Release);
Builder.CreateBr(ContBB);
SI->addCase(Builder.getInt32(3), ReleaseBB);
}
if (!IsLoad && !IsStore) {
Builder.SetInsertPoint(AcqRelBB);
EmitAtomicOp(*this, E, Dest, Ptr, Val1, Val2, Size, Align,
llvm::AcquireRelease);
Builder.CreateBr(ContBB);
SI->addCase(Builder.getInt32(4), AcqRelBB);
}
Builder.SetInsertPoint(SeqCstBB);
EmitAtomicOp(*this, E, Dest, Ptr, Val1, Val2, Size, Align,
llvm::SequentiallyConsistent);
Builder.CreateBr(ContBB);
SI->addCase(Builder.getInt32(5), SeqCstBB);
// Cleanup and return
Builder.SetInsertPoint(ContBB);
if (E->getType()->isVoidType())
return RValue::get(0);
return convertTempToRValue(OrigDest, E->getType());
}
/// Emit a load from an l-value of atomic type. Note that the r-value
/// we produce is an r-value of the atomic *value* type.
RValue CodeGenFunction::EmitAtomicLoad(LValue src, AggValueSlot resultSlot) {
AtomicInfo atomics(*this, src);
// Check whether we should use a library call.
if (atomics.shouldUseLibcall()) {
llvm::Value *tempAddr;
if (resultSlot.isValueOfAtomic()) {
assert(atomics.getEvaluationKind() == TEK_Aggregate);
tempAddr = resultSlot.getPaddedAtomicAddr();
} else if (!resultSlot.isIgnored() && !atomics.hasPadding()) {
assert(atomics.getEvaluationKind() == TEK_Aggregate);
tempAddr = resultSlot.getAddr();
} else {
tempAddr = CreateMemTemp(atomics.getAtomicType(), "atomic-load-temp");
}
// void __atomic_load(size_t size, void *mem, void *return, int order);
CallArgList args;
args.add(RValue::get(atomics.getAtomicSizeValue()),
getContext().getSizeType());
args.add(RValue::get(EmitCastToVoidPtr(src.getAddress())),
getContext().VoidPtrTy);
args.add(RValue::get(EmitCastToVoidPtr(tempAddr)),
getContext().VoidPtrTy);
args.add(RValue::get(llvm::ConstantInt::get(IntTy,
AO_ABI_memory_order_seq_cst)),
getContext().IntTy);
emitAtomicLibcall(*this, "__atomic_load", getContext().VoidTy, args);
// Produce the r-value.
return atomics.convertTempToRValue(tempAddr, resultSlot);
}
// Okay, we're doing this natively.
llvm::Value *addr = atomics.emitCastToAtomicIntPointer(src.getAddress());
llvm::LoadInst *load = Builder.CreateLoad(addr, "atomic-load");
load->setAtomic(llvm::SequentiallyConsistent);
// Other decoration.
load->setAlignment(src.getAlignment().getQuantity());
if (src.isVolatileQualified())
load->setVolatile(true);
if (src.getTBAAInfo())
CGM.DecorateInstruction(load, src.getTBAAInfo());
// Okay, turn that back into the original value type.
QualType valueType = atomics.getValueType();
llvm::Value *result = load;
// If we're ignoring an aggregate return, don't do anything.
if (atomics.getEvaluationKind() == TEK_Aggregate && resultSlot.isIgnored())
return RValue::getAggregate(0, false);
// The easiest way to do this this is to go through memory, but we
// try not to in some easy cases.
if (atomics.getEvaluationKind() == TEK_Scalar && !atomics.hasPadding()) {
llvm::Type *resultTy = CGM.getTypes().ConvertTypeForMem(valueType);
if (isa<llvm::IntegerType>(resultTy)) {
assert(result->getType() == resultTy);
result = EmitFromMemory(result, valueType);
} else if (isa<llvm::PointerType>(resultTy)) {
result = Builder.CreateIntToPtr(result, resultTy);
} else {
result = Builder.CreateBitCast(result, resultTy);
}
return RValue::get(result);
}
// Create a temporary. This needs to be big enough to hold the
// atomic integer.
llvm::Value *temp;
bool tempIsVolatile = false;
CharUnits tempAlignment;
if (atomics.getEvaluationKind() == TEK_Aggregate &&
(!atomics.hasPadding() || resultSlot.isValueOfAtomic())) {
assert(!resultSlot.isIgnored());
if (resultSlot.isValueOfAtomic()) {
temp = resultSlot.getPaddedAtomicAddr();
tempAlignment = atomics.getAtomicAlignment();
} else {
temp = resultSlot.getAddr();
tempAlignment = atomics.getValueAlignment();
}
tempIsVolatile = resultSlot.isVolatile();
} else {
temp = CreateMemTemp(atomics.getAtomicType(), "atomic-load-temp");
tempAlignment = atomics.getAtomicAlignment();
}
// Slam the integer into the temporary.
llvm::Value *castTemp = atomics.emitCastToAtomicIntPointer(temp);
Builder.CreateAlignedStore(result, castTemp, tempAlignment.getQuantity())
->setVolatile(tempIsVolatile);
return atomics.convertTempToRValue(temp, resultSlot);
}
void CodeGenFunction::GenerateThunk(llvm::Function *Fn,
const CGFunctionInfo &FnInfo,
GlobalDecl GD, const ThunkInfo &Thunk) {
const CXXMethodDecl *MD = cast<CXXMethodDecl>(GD.getDecl());
const FunctionProtoType *FPT = MD->getType()->getAs<FunctionProtoType>();
QualType ResultType = FPT->getResultType();
QualType ThisType = MD->getThisType(getContext());
FunctionArgList FunctionArgs;
// FIXME: It would be nice if more of this code could be shared with
// CodeGenFunction::GenerateCode.
// Create the implicit 'this' parameter declaration.
CurGD = GD;
CGM.getCXXABI().BuildInstanceFunctionParams(*this, ResultType, FunctionArgs);
// Add the rest of the parameters.
for (FunctionDecl::param_const_iterator I = MD->param_begin(),
E = MD->param_end(); I != E; ++I) {
ParmVarDecl *Param = *I;
FunctionArgs.push_back(Param);
}
// Initialize debug info if needed.
maybeInitializeDebugInfo();
StartFunction(GlobalDecl(), ResultType, Fn, FnInfo, FunctionArgs,
SourceLocation());
CGM.getCXXABI().EmitInstanceFunctionProlog(*this);
CXXThisValue = CXXABIThisValue;
// Adjust the 'this' pointer if necessary.
llvm::Value *AdjustedThisPtr =
PerformTypeAdjustment(*this, LoadCXXThis(),
Thunk.This.NonVirtual,
Thunk.This.VCallOffsetOffset,
/*IsReturnAdjustment*/false);
CallArgList CallArgs;
// Add our adjusted 'this' pointer.
CallArgs.add(RValue::get(AdjustedThisPtr), ThisType);
// Add the rest of the parameters.
for (FunctionDecl::param_const_iterator I = MD->param_begin(),
E = MD->param_end(); I != E; ++I) {
ParmVarDecl *param = *I;
EmitDelegateCallArg(CallArgs, param);
}
// Get our callee.
llvm::Type *Ty =
CGM.getTypes().GetFunctionType(CGM.getTypes().arrangeGlobalDeclaration(GD));
llvm::Value *Callee = CGM.GetAddrOfFunction(GD, Ty, /*ForVTable=*/true);
#ifndef NDEBUG
const CGFunctionInfo &CallFnInfo =
CGM.getTypes().arrangeCXXMethodCall(CallArgs, FPT,
RequiredArgs::forPrototypePlus(FPT, 1));
assert(CallFnInfo.getRegParm() == FnInfo.getRegParm() &&
CallFnInfo.isNoReturn() == FnInfo.isNoReturn() &&
CallFnInfo.getCallingConvention() == FnInfo.getCallingConvention());
assert(isa<CXXDestructorDecl>(MD) || // ignore dtor return types
similar(CallFnInfo.getReturnInfo(), CallFnInfo.getReturnType(),
FnInfo.getReturnInfo(), FnInfo.getReturnType()));
assert(CallFnInfo.arg_size() == FnInfo.arg_size());
for (unsigned i = 0, e = FnInfo.arg_size(); i != e; ++i)
assert(similar(CallFnInfo.arg_begin()[i].info,
CallFnInfo.arg_begin()[i].type,
FnInfo.arg_begin()[i].info, FnInfo.arg_begin()[i].type));
#endif
// Determine whether we have a return value slot to use.
ReturnValueSlot Slot;
if (!ResultType->isVoidType() &&
FnInfo.getReturnInfo().getKind() == ABIArgInfo::Indirect &&
!hasScalarEvaluationKind(CurFnInfo->getReturnType()))
Slot = ReturnValueSlot(ReturnValue, ResultType.isVolatileQualified());
// Now emit our call.
RValue RV = EmitCall(FnInfo, Callee, Slot, CallArgs, MD);
if (!Thunk.Return.isEmpty())
RV = PerformReturnAdjustment(*this, ResultType, RV, Thunk);
if (!ResultType->isVoidType() && Slot.isNull())
CGM.getCXXABI().EmitReturnFromThunk(*this, RV, ResultType);
// Disable the final ARC autorelease.
AutoreleaseResult = false;
FinishFunction();
// Set the right linkage.
CGM.setFunctionLinkage(GD, Fn);
// Set the right visibility.
setThunkVisibility(CGM, MD, Thunk, Fn);
}
void CodeGenFunction::EmitCallAndReturnForThunk(GlobalDecl GD,
llvm::Value *Callee,
const ThunkInfo *Thunk) {
assert(isa<CXXMethodDecl>(CurGD.getDecl()) &&
"Please use a new CGF for this thunk");
const CXXMethodDecl *MD = cast<CXXMethodDecl>(GD.getDecl());
// Adjust the 'this' pointer if necessary
llvm::Value *AdjustedThisPtr = Thunk ? CGM.getCXXABI().performThisAdjustment(
*this, LoadCXXThis(), Thunk->This)
: LoadCXXThis();
// Start building CallArgs.
CallArgList CallArgs;
QualType ThisType = MD->getThisType(getContext());
CallArgs.add(RValue::get(AdjustedThisPtr), ThisType);
if (isa<CXXDestructorDecl>(MD))
CGM.getCXXABI().adjustCallArgsForDestructorThunk(*this, GD, CallArgs);
// Add the rest of the arguments.
for (FunctionDecl::param_const_iterator I = MD->param_begin(),
E = MD->param_end(); I != E; ++I)
EmitDelegateCallArg(CallArgs, *I, (*I)->getLocStart());
const FunctionProtoType *FPT = MD->getType()->getAs<FunctionProtoType>();
#ifndef NDEBUG
const CGFunctionInfo &CallFnInfo =
CGM.getTypes().arrangeCXXMethodCall(CallArgs, FPT,
RequiredArgs::forPrototypePlus(FPT, 1));
assert(CallFnInfo.getRegParm() == CurFnInfo->getRegParm() &&
CallFnInfo.isNoReturn() == CurFnInfo->isNoReturn() &&
CallFnInfo.getCallingConvention() == CurFnInfo->getCallingConvention());
assert(isa<CXXDestructorDecl>(MD) || // ignore dtor return types
similar(CallFnInfo.getReturnInfo(), CallFnInfo.getReturnType(),
CurFnInfo->getReturnInfo(), CurFnInfo->getReturnType()));
assert(CallFnInfo.arg_size() == CurFnInfo->arg_size());
for (unsigned i = 0, e = CurFnInfo->arg_size(); i != e; ++i)
assert(similar(CallFnInfo.arg_begin()[i].info,
CallFnInfo.arg_begin()[i].type,
CurFnInfo->arg_begin()[i].info,
CurFnInfo->arg_begin()[i].type));
#endif
// Determine whether we have a return value slot to use.
QualType ResultType =
CGM.getCXXABI().HasThisReturn(GD) ? ThisType : FPT->getResultType();
ReturnValueSlot Slot;
if (!ResultType->isVoidType() &&
CurFnInfo->getReturnInfo().getKind() == ABIArgInfo::Indirect &&
!hasScalarEvaluationKind(CurFnInfo->getReturnType()))
Slot = ReturnValueSlot(ReturnValue, ResultType.isVolatileQualified());
// Now emit our call.
RValue RV = EmitCall(*CurFnInfo, Callee, Slot, CallArgs, MD);
// Consider return adjustment if we have ThunkInfo.
if (Thunk && !Thunk->Return.isEmpty())
RV = PerformReturnAdjustment(*this, ResultType, RV, *Thunk);
// Emit return.
if (!ResultType->isVoidType() && Slot.isNull())
CGM.getCXXABI().EmitReturnFromThunk(*this, RV, ResultType);
// Disable the final ARC autorelease.
AutoreleaseResult = false;
FinishFunction();
}
void CodeGenFunction::EmitCallAndReturnForThunk(llvm::Value *Callee,
const ThunkInfo *Thunk) {
assert(isa<CXXMethodDecl>(CurGD.getDecl()) &&
"Please use a new CGF for this thunk");
const CXXMethodDecl *MD = cast<CXXMethodDecl>(CurGD.getDecl());
// Adjust the 'this' pointer if necessary
llvm::Value *AdjustedThisPtr =
Thunk ? CGM.getCXXABI().performThisAdjustment(
*this, LoadCXXThisAddress(), Thunk->This)
: LoadCXXThis();
if (CurFnInfo->usesInAlloca()) {
// We don't handle return adjusting thunks, because they require us to call
// the copy constructor. For now, fall through and pretend the return
// adjustment was empty so we don't crash.
if (Thunk && !Thunk->Return.isEmpty()) {
CGM.ErrorUnsupported(
MD, "non-trivial argument copy for return-adjusting thunk");
}
EmitMustTailThunk(MD, AdjustedThisPtr, Callee);
return;
}
// Start building CallArgs.
CallArgList CallArgs;
QualType ThisType = MD->getThisType(getContext());
CallArgs.add(RValue::get(AdjustedThisPtr), ThisType);
if (isa<CXXDestructorDecl>(MD))
CGM.getCXXABI().adjustCallArgsForDestructorThunk(*this, CurGD, CallArgs);
// Add the rest of the arguments.
for (const ParmVarDecl *PD : MD->parameters())
EmitDelegateCallArg(CallArgs, PD, PD->getLocStart());
const FunctionProtoType *FPT = MD->getType()->getAs<FunctionProtoType>();
#ifndef NDEBUG
const CGFunctionInfo &CallFnInfo = CGM.getTypes().arrangeCXXMethodCall(
CallArgs, FPT, RequiredArgs::forPrototypePlus(FPT, 1, MD));
assert(CallFnInfo.getRegParm() == CurFnInfo->getRegParm() &&
CallFnInfo.isNoReturn() == CurFnInfo->isNoReturn() &&
CallFnInfo.getCallingConvention() == CurFnInfo->getCallingConvention());
assert(isa<CXXDestructorDecl>(MD) || // ignore dtor return types
similar(CallFnInfo.getReturnInfo(), CallFnInfo.getReturnType(),
CurFnInfo->getReturnInfo(), CurFnInfo->getReturnType()));
assert(CallFnInfo.arg_size() == CurFnInfo->arg_size());
for (unsigned i = 0, e = CurFnInfo->arg_size(); i != e; ++i)
assert(similar(CallFnInfo.arg_begin()[i].info,
CallFnInfo.arg_begin()[i].type,
CurFnInfo->arg_begin()[i].info,
CurFnInfo->arg_begin()[i].type));
#endif
// Determine whether we have a return value slot to use.
QualType ResultType = CGM.getCXXABI().HasThisReturn(CurGD)
? ThisType
: CGM.getCXXABI().hasMostDerivedReturn(CurGD)
? CGM.getContext().VoidPtrTy
: FPT->getReturnType();
ReturnValueSlot Slot;
if (!ResultType->isVoidType() &&
CurFnInfo->getReturnInfo().getKind() == ABIArgInfo::Indirect &&
!hasScalarEvaluationKind(CurFnInfo->getReturnType()))
Slot = ReturnValueSlot(ReturnValue, ResultType.isVolatileQualified());
// Now emit our call.
llvm::Instruction *CallOrInvoke;
RValue RV = EmitCall(*CurFnInfo, Callee, Slot, CallArgs, MD, &CallOrInvoke);
// Consider return adjustment if we have ThunkInfo.
if (Thunk && !Thunk->Return.isEmpty())
RV = PerformReturnAdjustment(*this, ResultType, RV, *Thunk);
else if (llvm::CallInst* Call = dyn_cast<llvm::CallInst>(CallOrInvoke))
Call->setTailCallKind(llvm::CallInst::TCK_Tail);
// Emit return.
if (!ResultType->isVoidType() && Slot.isNull())
CGM.getCXXABI().EmitReturnFromThunk(*this, RV, ResultType);
// Disable the final ARC autorelease.
AutoreleaseResult = false;
FinishThunk();
}
// CUDA 9.0+ uses new way to launch kernels. Parameters are packed in a local
// array and kernels are launched using cudaLaunchKernel().
void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
FunctionArgList &Args) {
// Build the shadow stack entry at the very start of the function.
// Calculate amount of space we will need for all arguments. If we have no
// args, allocate a single pointer so we still have a valid pointer to the
// argument array that we can pass to runtime, even if it will be unused.
Address KernelArgs = CGF.CreateTempAlloca(
VoidPtrTy, CharUnits::fromQuantity(16), "kernel_args",
llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size())));
// Store pointers to the arguments in a locally allocated launch_args.
for (unsigned i = 0; i < Args.size(); ++i) {
llvm::Value* VarPtr = CGF.GetAddrOfLocalVar(Args[i]).getPointer();
llvm::Value *VoidVarPtr = CGF.Builder.CreatePointerCast(VarPtr, VoidPtrTy);
CGF.Builder.CreateDefaultAlignedStore(
VoidVarPtr, CGF.Builder.CreateConstGEP1_32(KernelArgs.getPointer(), i));
}
llvm::BasicBlock *EndBlock = CGF.createBasicBlock("setup.end");
// Lookup cudaLaunchKernel function.
// cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
// void **args, size_t sharedMem,
// cudaStream_t stream);
TranslationUnitDecl *TUDecl = CGM.getContext().getTranslationUnitDecl();
DeclContext *DC = TranslationUnitDecl::castToDeclContext(TUDecl);
IdentifierInfo &cudaLaunchKernelII =
CGM.getContext().Idents.get("cudaLaunchKernel");
FunctionDecl *cudaLaunchKernelFD = nullptr;
for (const auto &Result : DC->lookup(&cudaLaunchKernelII)) {
if (FunctionDecl *FD = dyn_cast<FunctionDecl>(Result))
cudaLaunchKernelFD = FD;
}
if (cudaLaunchKernelFD == nullptr) {
CGM.Error(CGF.CurFuncDecl->getLocation(),
"Can't find declaration for cudaLaunchKernel()");
return;
}
// Create temporary dim3 grid_dim, block_dim.
ParmVarDecl *GridDimParam = cudaLaunchKernelFD->getParamDecl(1);
QualType Dim3Ty = GridDimParam->getType();
Address GridDim =
CGF.CreateMemTemp(Dim3Ty, CharUnits::fromQuantity(8), "grid_dim");
Address BlockDim =
CGF.CreateMemTemp(Dim3Ty, CharUnits::fromQuantity(8), "block_dim");
Address ShmemSize =
CGF.CreateTempAlloca(SizeTy, CGM.getSizeAlign(), "shmem_size");
Address Stream =
CGF.CreateTempAlloca(VoidPtrTy, CGM.getPointerAlign(), "stream");
llvm::FunctionCallee cudaPopConfigFn = CGM.CreateRuntimeFunction(
llvm::FunctionType::get(IntTy,
{/*gridDim=*/GridDim.getType(),
/*blockDim=*/BlockDim.getType(),
/*ShmemSize=*/ShmemSize.getType(),
/*Stream=*/Stream.getType()},
/*isVarArg=*/false),
"__cudaPopCallConfiguration");
CGF.EmitRuntimeCallOrInvoke(cudaPopConfigFn,
{GridDim.getPointer(), BlockDim.getPointer(),
ShmemSize.getPointer(), Stream.getPointer()});
// Emit the call to cudaLaunch
llvm::Value *Kernel = CGF.Builder.CreatePointerCast(CGF.CurFn, VoidPtrTy);
CallArgList LaunchKernelArgs;
LaunchKernelArgs.add(RValue::get(Kernel),
cudaLaunchKernelFD->getParamDecl(0)->getType());
LaunchKernelArgs.add(RValue::getAggregate(GridDim), Dim3Ty);
LaunchKernelArgs.add(RValue::getAggregate(BlockDim), Dim3Ty);
LaunchKernelArgs.add(RValue::get(KernelArgs.getPointer()),
cudaLaunchKernelFD->getParamDecl(3)->getType());
LaunchKernelArgs.add(RValue::get(CGF.Builder.CreateLoad(ShmemSize)),
cudaLaunchKernelFD->getParamDecl(4)->getType());
LaunchKernelArgs.add(RValue::get(CGF.Builder.CreateLoad(Stream)),
cudaLaunchKernelFD->getParamDecl(5)->getType());
QualType QT = cudaLaunchKernelFD->getType();
QualType CQT = QT.getCanonicalType();
llvm::Type *Ty = CGM.getTypes().ConvertType(CQT);
llvm::FunctionType *FTy = dyn_cast<llvm::FunctionType>(Ty);
const CGFunctionInfo &FI =
CGM.getTypes().arrangeFunctionDeclaration(cudaLaunchKernelFD);
llvm::FunctionCallee cudaLaunchKernelFn =
CGM.CreateRuntimeFunction(FTy, "cudaLaunchKernel");
CGF.EmitCall(FI, CGCallee::forDirect(cudaLaunchKernelFn), ReturnValueSlot(),
LaunchKernelArgs);
CGF.EmitBranch(EndBlock);
CGF.EmitBlock(EndBlock);
}
/// Emit a store to an l-value of atomic type.
///
/// Note that the r-value is expected to be an r-value *of the atomic
/// type*; this means that for aggregate r-values, it should include
/// storage for any padding that was necessary.
void CodeGenFunction::EmitAtomicStore(RValue rvalue, LValue dest,
bool isInit) {
// If this is an aggregate r-value, it should agree in type except
// maybe for address-space qualification.
assert(!rvalue.isAggregate() ||
rvalue.getAggregateAddr()->getType()->getPointerElementType()
== dest.getAddress()->getType()->getPointerElementType());
AtomicInfo atomics(*this, dest);
// If this is an initialization, just put the value there normally.
if (isInit) {
atomics.emitCopyIntoMemory(rvalue, dest);
return;
}
// Check whether we should use a library call.
if (atomics.shouldUseLibcall()) {
// Produce a source address.
llvm::Value *srcAddr = atomics.materializeRValue(rvalue);
// void __atomic_store(size_t size, void *mem, void *val, int order)
CallArgList args;
args.add(RValue::get(atomics.getAtomicSizeValue()),
getContext().getSizeType());
args.add(RValue::get(EmitCastToVoidPtr(dest.getAddress())),
getContext().VoidPtrTy);
args.add(RValue::get(EmitCastToVoidPtr(srcAddr)),
getContext().VoidPtrTy);
args.add(RValue::get(llvm::ConstantInt::get(IntTy,
AO_ABI_memory_order_seq_cst)),
getContext().IntTy);
emitAtomicLibcall(*this, "__atomic_store", getContext().VoidTy, args);
return;
}
// Okay, we're doing this natively.
llvm::Value *intValue;
// If we've got a scalar value of the right size, try to avoid going
// through memory.
if (rvalue.isScalar() && !atomics.hasPadding()) {
llvm::Value *value = rvalue.getScalarVal();
if (isa<llvm::IntegerType>(value->getType())) {
intValue = value;
} else {
llvm::IntegerType *inputIntTy =
llvm::IntegerType::get(getLLVMContext(), atomics.getValueSizeInBits());
if (isa<llvm::PointerType>(value->getType())) {
intValue = Builder.CreatePtrToInt(value, inputIntTy);
} else {
intValue = Builder.CreateBitCast(value, inputIntTy);
}
}
// Otherwise, we need to go through memory.
} else {
// Put the r-value in memory.
llvm::Value *addr = atomics.materializeRValue(rvalue);
// Cast the temporary to the atomic int type and pull a value out.
addr = atomics.emitCastToAtomicIntPointer(addr);
intValue = Builder.CreateAlignedLoad(addr,
atomics.getAtomicAlignment().getQuantity());
}
// Do the atomic store.
llvm::Value *addr = atomics.emitCastToAtomicIntPointer(dest.getAddress());
llvm::StoreInst *store = Builder.CreateStore(intValue, addr);
// Initializations don't need to be atomic.
if (!isInit) store->setAtomic(llvm::SequentiallyConsistent);
// Other decoration.
store->setAlignment(dest.getAlignment().getQuantity());
if (dest.isVolatileQualified())
store->setVolatile(true);
if (dest.getTBAAInfo())
CGM.DecorateInstruction(store, dest.getTBAAInfo());
}
