DominatingValue<RValue>::saved_type
DominatingValue<RValue>::saved_type::save(CodeGenFunction &CGF, RValue rv) {
if (rv.isScalar()) {
llvm::Value *V = rv.getScalarVal();
// These automatically dominate and don't need to be saved.
if (!DominatingLLVMValue::needsSaving(V))
return saved_type(V, ScalarLiteral);
// Everything else needs an alloca.
llvm::Value *addr = CGF.CreateTempAlloca(V->getType(), "saved-rvalue");
CGF.Builder.CreateStore(V, addr);
return saved_type(addr, ScalarAddress);
}
if (rv.isComplex()) {
CodeGenFunction::ComplexPairTy V = rv.getComplexVal();
llvm::Type *ComplexTy =
llvm::StructType::get(V.first->getType(), V.second->getType(),
(void*) nullptr);
llvm::Value *addr = CGF.CreateTempAlloca(ComplexTy, "saved-complex");
CGF.Builder.CreateStore(V.first, CGF.Builder.CreateStructGEP(addr, 0));
CGF.Builder.CreateStore(V.second, CGF.Builder.CreateStructGEP(addr, 1));
return saved_type(addr, ComplexAddress);
}
assert(rv.isAggregate());
llvm::Value *V = rv.getAggregateAddr(); // TODO: volatile?
if (!DominatingLLVMValue::needsSaving(V))
return saved_type(V, AggregateLiteral);
llvm::Value *addr = CGF.CreateTempAlloca(V->getType(), "saved-rvalue");
CGF.Builder.CreateStore(V, addr);
return saved_type(addr, AggregateAddress);
}
/// The given cleanup block is changing activation state. Configure a
/// cleanup variable if necessary.
///
/// It would be good if we had some way of determining if there were
/// extra uses *after* the change-over point.
static void SetupCleanupBlockActivation(CodeGenFunction &CGF,
EHScopeStack::stable_iterator C,
ForActivation_t kind,
llvm::Instruction *dominatingIP) {
EHCleanupScope &Scope = cast<EHCleanupScope>(*CGF.EHStack.find(C));
// We always need the flag if we're activating the cleanup in a
// conditional context, because we have to assume that the current
// location doesn't necessarily dominate the cleanup's code.
bool isActivatedInConditional =
(kind == ForActivation && CGF.isInConditionalBranch());
bool needFlag = false;
// Calculate whether the cleanup was used:
// - as a normal cleanup
if (Scope.isNormalCleanup() &&
(isActivatedInConditional || IsUsedAsNormalCleanup(CGF.EHStack, C))) {
Scope.setTestFlagInNormalCleanup();
needFlag = true;
}
// - as an EH cleanup
if (Scope.isEHCleanup() &&
(isActivatedInConditional || IsUsedAsEHCleanup(CGF.EHStack, C))) {
Scope.setTestFlagInEHCleanup();
needFlag = true;
}
// If it hasn't yet been used as either, we're done.
if (!needFlag) return;
Address var = Scope.getActiveFlag();
if (!var.isValid()) {
var = CGF.CreateTempAlloca(CGF.Builder.getInt1Ty(), CharUnits::One(),
"cleanup.isactive");
Scope.setActiveFlag(var);
assert(dominatingIP && "no existing variable and no dominating IP!");
// Initialize to true or false depending on whether it was
// active up to this point.
llvm::Constant *value = CGF.Builder.getInt1(kind == ForDeactivation);
// If we're in a conditional block, ignore the dominating IP and
// use the outermost conditional branch.
if (CGF.isInConditionalBranch()) {
CGF.setBeforeOutermostConditional(value, var);
} else {
createStoreInstBefore(value, var, dominatingIP);
}
}
CGF.Builder.CreateStore(CGF.Builder.getInt1(kind == ForActivation), var);
}
// CopyObject - Utility to copy an object. Calls copy constructor as necessary.
// N is casted to the right type.
static void CopyObject(CodeGenFunction &CGF, QualType ObjectType,
bool WasPointer, bool WasPointerReference,
llvm::Value *E, llvm::Value *N) {
// Store the throw exception in the exception object.
if (WasPointer || !CGF.hasAggregateLLVMType(ObjectType)) {
llvm::Value *Value = E;
if (!WasPointer)
Value = CGF.Builder.CreateLoad(Value);
const llvm::Type *ValuePtrTy = Value->getType()->getPointerTo(0);
if (WasPointerReference) {
llvm::Value *Tmp = CGF.CreateTempAlloca(Value->getType(), "catch.param");
CGF.Builder.CreateStore(Value, Tmp);
Value = Tmp;
ValuePtrTy = Value->getType()->getPointerTo(0);
}
N = CGF.Builder.CreateBitCast(N, ValuePtrTy);
CGF.Builder.CreateStore(Value, N);
} else {
const llvm::Type *Ty = CGF.ConvertType(ObjectType)->getPointerTo(0);
const CXXRecordDecl *RD;
RD = cast<CXXRecordDecl>(ObjectType->getAs<RecordType>()->getDecl());
llvm::Value *This = CGF.Builder.CreateBitCast(N, Ty);
if (RD->hasTrivialCopyConstructor()) {
CGF.EmitAggregateCopy(This, E, ObjectType);
} else if (CXXConstructorDecl *CopyCtor
= RD->getCopyConstructor(CGF.getContext(), 0)) {
llvm::Value *Src = E;
// Stolen from EmitClassAggrMemberwiseCopy
llvm::Value *Callee = CGF.CGM.GetAddrOfCXXConstructor(CopyCtor,
Ctor_Complete);
CallArgList CallArgs;
CallArgs.push_back(std::make_pair(RValue::get(This),
CopyCtor->getThisType(CGF.getContext())));
// Push the Src ptr.
CallArgs.push_back(std::make_pair(RValue::get(Src),
CopyCtor->getParamDecl(0)->getType()));
const FunctionProtoType *FPT
= CopyCtor->getType()->getAs<FunctionProtoType>();
CGF.EmitCall(CGF.CGM.getTypes().getFunctionInfo(CallArgs, FPT),
Callee, ReturnValueSlot(), CallArgs, CopyCtor);
} else
llvm_unreachable("uncopyable object");
}
}
DominatingValue<RValue>::saved_type
DominatingValue<RValue>::saved_type::save(CodeGenFunction &CGF, RValue rv) {
if (rv.isScalar()) {
llvm::Value *V = rv.getScalarVal();
// These automatically dominate and don't need to be saved.
if (!DominatingLLVMValue::needsSaving(V))
return saved_type(V, ScalarLiteral);
// Everything else needs an alloca.
Address addr =
CGF.CreateDefaultAlignTempAlloca(V->getType(), "saved-rvalue");
CGF.Builder.CreateStore(V, addr);
return saved_type(addr.getPointer(), ScalarAddress);
}
if (rv.isComplex()) {
CodeGenFunction::ComplexPairTy V = rv.getComplexVal();
llvm::Type *ComplexTy =
llvm::StructType::get(V.first->getType(), V.second->getType(),
(void*) nullptr);
Address addr = CGF.CreateDefaultAlignTempAlloca(ComplexTy, "saved-complex");
CGF.Builder.CreateStore(V.first,
CGF.Builder.CreateStructGEP(addr, 0, CharUnits()));
CharUnits offset = CharUnits::fromQuantity(
CGF.CGM.getDataLayout().getTypeAllocSize(V.first->getType()));
CGF.Builder.CreateStore(V.second,
CGF.Builder.CreateStructGEP(addr, 1, offset));
return saved_type(addr.getPointer(), ComplexAddress);
}
assert(rv.isAggregate());
Address V = rv.getAggregateAddress(); // TODO: volatile?
if (!DominatingLLVMValue::needsSaving(V.getPointer()))
return saved_type(V.getPointer(), AggregateLiteral,
V.getAlignment().getQuantity());
Address addr =
CGF.CreateTempAlloca(V.getType(), CGF.getPointerAlign(), "saved-rvalue");
CGF.Builder.CreateStore(V.getPointer(), addr);
return saved_type(addr.getPointer(), AggregateAddress,
V.getAlignment().getQuantity());
}
/// CreateCoercedStore - Create a store to \arg DstPtr from \arg Src,
/// where the source and destination may have different types.
///
/// This safely handles the case when the src type is larger than the
/// destination type; the upper bits of the src will be lost.
static void CreateCoercedStore(llvm::Value *Src,
llvm::Value *DstPtr,
bool DstIsVolatile,
CodeGenFunction &CGF) {
const llvm::Type *SrcTy = Src->getType();
const llvm::Type *DstTy =
cast<llvm::PointerType>(DstPtr->getType())->getElementType();
uint64_t SrcSize = CGF.CGM.getTargetData().getTypeAllocSize(SrcTy);
uint64_t DstSize = CGF.CGM.getTargetData().getTypeAllocSize(DstTy);
// If store is legal, just bitcast the src pointer.
if (SrcSize <= DstSize) {
llvm::Value *Casted =
CGF.Builder.CreateBitCast(DstPtr, llvm::PointerType::getUnqual(SrcTy));
// FIXME: Use better alignment / avoid requiring aligned store.
CGF.Builder.CreateStore(Src, Casted, DstIsVolatile)->setAlignment(1);
} else {
// Otherwise do coercion through memory. This is stupid, but
// simple.
// Generally SrcSize is never greater than DstSize, since this means we are
// losing bits. However, this can happen in cases where the structure has
// additional padding, for example due to a user specified alignment.
//
// FIXME: Assert that we aren't truncating non-padding bits when have access
// to that information.
llvm::Value *Tmp = CGF.CreateTempAlloca(SrcTy);
CGF.Builder.CreateStore(Src, Tmp);
llvm::Value *Casted =
CGF.Builder.CreateBitCast(Tmp, llvm::PointerType::getUnqual(DstTy));
llvm::LoadInst *Load = CGF.Builder.CreateLoad(Casted);
// FIXME: Use better alignment / avoid requiring aligned load.
Load->setAlignment(1);
CGF.Builder.CreateStore(Load, DstPtr, DstIsVolatile);
}
}
// CopyObject - Utility to copy an object. Calls copy constructor as necessary.
// DestPtr is casted to the right type.
static void CopyObject(CodeGenFunction &CGF, const Expr *E,
llvm::Value *DestPtr, llvm::Value *ExceptionPtrPtr) {
QualType ObjectType = E->getType();
// Store the throw exception in the exception object.
if (!CGF.hasAggregateLLVMType(ObjectType)) {
llvm::Value *Value = CGF.EmitScalarExpr(E);
const llvm::Type *ValuePtrTy = Value->getType()->getPointerTo();
CGF.Builder.CreateStore(Value,
CGF.Builder.CreateBitCast(DestPtr, ValuePtrTy));
} else {
const llvm::Type *Ty = CGF.ConvertType(ObjectType)->getPointerTo();
const CXXRecordDecl *RD =
cast<CXXRecordDecl>(ObjectType->getAs<RecordType>()->getDecl());
llvm::Value *This = CGF.Builder.CreateBitCast(DestPtr, Ty);
if (RD->hasTrivialCopyConstructor()) {
CGF.EmitAggExpr(E, This, false);
} else if (CXXConstructorDecl *CopyCtor
= RD->getCopyConstructor(CGF.getContext(), 0)) {
llvm::Value *CondPtr = 0;
if (CGF.Exceptions) {
CodeGenFunction::EHCleanupBlock Cleanup(CGF);
llvm::Constant *FreeExceptionFn = getFreeExceptionFn(CGF);
llvm::BasicBlock *CondBlock = CGF.createBasicBlock("cond.free");
llvm::BasicBlock *Cont = CGF.createBasicBlock("cont");
CondPtr = CGF.CreateTempAlloca(llvm::Type::getInt1Ty(CGF.getLLVMContext()),
"doEHfree");
CGF.Builder.CreateCondBr(CGF.Builder.CreateLoad(CondPtr),
CondBlock, Cont);
CGF.EmitBlock(CondBlock);
// Load the exception pointer.
llvm::Value *ExceptionPtr = CGF.Builder.CreateLoad(ExceptionPtrPtr);
CGF.Builder.CreateCall(FreeExceptionFn, ExceptionPtr);
CGF.EmitBlock(Cont);
}
if (CondPtr)
CGF.Builder.CreateStore(llvm::ConstantInt::getTrue(CGF.getLLVMContext()),
CondPtr);
llvm::Value *Src = CGF.EmitLValue(E).getAddress();
if (CondPtr)
CGF.Builder.CreateStore(llvm::ConstantInt::getFalse(CGF.getLLVMContext()),
CondPtr);
llvm::BasicBlock *TerminateHandler = CGF.getTerminateHandler();
llvm::BasicBlock *PrevLandingPad = CGF.getInvokeDest();
CGF.setInvokeDest(TerminateHandler);
// Stolen from EmitClassAggrMemberwiseCopy
llvm::Value *Callee = CGF.CGM.GetAddrOfCXXConstructor(CopyCtor,
Ctor_Complete);
CallArgList CallArgs;
CallArgs.push_back(std::make_pair(RValue::get(This),
CopyCtor->getThisType(CGF.getContext())));
// Push the Src ptr.
CallArgs.push_back(std::make_pair(RValue::get(Src),
CopyCtor->getParamDecl(0)->getType()));
const FunctionProtoType *FPT
= CopyCtor->getType()->getAs<FunctionProtoType>();
CGF.EmitCall(CGF.CGM.getTypes().getFunctionInfo(CallArgs, FPT),
Callee, ReturnValueSlot(), CallArgs, CopyCtor);
CGF.setInvokeDest(PrevLandingPad);
} else
llvm_unreachable("uncopyable object");
}
}
// 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);
}
