Decl *Sema::ActOnEntityDecl(ASTContext &C, const QualType &T, SourceLocation IDLoc, const IdentifierInfo *IDInfo) { auto Quals = T.getQualifiers(); if(Quals.hasAttributeSpec(Qualifiers::AS_external)) return ActOnExternalEntityDecl(C, T, IDLoc, IDInfo); else if(Quals.hasAttributeSpec(Qualifiers::AS_intrinsic)) return ActOnIntrinsicEntityDecl(C, T, IDLoc, IDInfo); if (auto Prev = LookupIdentifier(IDInfo)) { FunctionDecl *FD = dyn_cast<FunctionDecl>(Prev); if(auto VD = dyn_cast<VarDecl>(Prev)) { if(VD->isArgument() && VD->getType().isNull()) { VD->setType(T); return VD; } else if(VD->isFunctionResult()) FD = CurrentContextAsFunction(); } if(FD && (FD->isNormalFunction() || FD->isExternal())) { if(FD->getType().isNull() || FD->getType()->isVoidType()) { SetFunctionType(FD, T, IDLoc, SourceRange()); //Fixme: proper loc and range return FD; } else { Diags.Report(IDLoc, diag::err_func_return_type_already_specified) << IDInfo; return nullptr; } } Diags.Report(IDLoc, diag::err_redefinition) << IDInfo; Diags.Report(Prev->getLocation(), diag::note_previous_definition); return nullptr; } VarDecl *VD = VarDecl::Create(C, CurContext, IDLoc, IDInfo, T); CurContext->addDecl(VD); if(!T.isNull()) { auto SubT = T; if(T->isArrayType()) { CheckArrayTypeDeclarationCompability(T->asArrayType(), VD); SubT = T->asArrayType()->getElementType(); VD->MarkUsedAsVariable(IDLoc); } else if(SubT->isCharacterType()) CheckCharacterLengthDeclarationCompability(SubT, VD); } return VD; }
static CallExpr *create_call_once_lambda_call(ASTContext &C, ASTMaker M, const ParmVarDecl *Callback, CXXRecordDecl *CallbackDecl, ArrayRef<Expr *> CallArgs) { assert(CallbackDecl != nullptr); assert(CallbackDecl->isLambda()); FunctionDecl *callOperatorDecl = CallbackDecl->getLambdaCallOperator(); assert(callOperatorDecl != nullptr); DeclRefExpr *callOperatorDeclRef = DeclRefExpr::Create(/* Ctx =*/ C, /* QualifierLoc =*/ NestedNameSpecifierLoc(), /* TemplateKWLoc =*/ SourceLocation(), const_cast<FunctionDecl *>(callOperatorDecl), /* RefersToEnclosingVariableOrCapture=*/ false, /* NameLoc =*/ SourceLocation(), /* T =*/ callOperatorDecl->getType(), /* VK =*/ VK_LValue); return new (C) CXXOperatorCallExpr(/*AstContext=*/C, OO_Call, callOperatorDeclRef, /*args=*/CallArgs, /*QualType=*/C.VoidTy, /*ExprValueType=*/VK_RValue, /*SourceLocation=*/SourceLocation(), FPOptions()); }
// Here is the test Eval function specialization. Here the CallExpr to the // function is created. CallExpr* EvaluateTSynthesizer::BuildEvalCallExpr(const QualType InstTy, Expr* SubTree, ASTOwningVector<Expr*>& CallArgs) { // Set up new context for the new FunctionDecl DeclContext* PrevContext = m_Sema->CurContext; m_Sema->CurContext = m_EvalDecl->getDeclContext(); // Create template arguments Sema::InstantiatingTemplate Inst(*m_Sema, m_NoSLoc, m_EvalDecl); TemplateArgument Arg(InstTy); TemplateArgumentList TemplateArgs(TemplateArgumentList::OnStack, &Arg, 1U); // Substitute the declaration of the templated function, with the // specified template argument Decl* D = m_Sema->SubstDecl(m_EvalDecl, m_EvalDecl->getDeclContext(), MultiLevelTemplateArgumentList(TemplateArgs)); FunctionDecl* Fn = dyn_cast<FunctionDecl>(D); // Creates new body of the substituted declaration m_Sema->InstantiateFunctionDefinition(Fn->getLocation(), Fn, true, true); m_Sema->CurContext = PrevContext; const FunctionProtoType* FPT = Fn->getType()->getAs<FunctionProtoType>(); FunctionProtoType::ExtProtoInfo EPI = FPT->getExtProtoInfo(); QualType FnTy = m_Context->getFunctionType(Fn->getResultType(), FPT->arg_type_begin(), FPT->getNumArgs(), EPI); DeclRefExpr* DRE = m_Sema->BuildDeclRefExpr(Fn, FnTy, VK_RValue, m_NoSLoc ).takeAs<DeclRefExpr>(); // TODO: Figure out a way to avoid passing in wrong source locations // of the symbol being replaced. This is important when we calculate the // size of the memory buffers and may lead to creation of wrong wrappers. Scope* S = m_Sema->getScopeForContext(m_Sema->CurContext); CallExpr* EvalCall = m_Sema->ActOnCallExpr(S, DRE, SubTree->getLocStart(), move_arg(CallArgs), SubTree->getLocEnd() ).takeAs<CallExpr>(); assert (EvalCall && "Cannot create call to Eval"); return EvalCall; }
ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc, MultiExprArg ExecConfig, SourceLocation GGGLoc) { FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl(); if (!ConfigDecl) return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use) << "cudaConfigureCall"); QualType ConfigQTy = ConfigDecl->getType(); DeclRefExpr *ConfigDR = new (Context) DeclRefExpr(ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc); MarkFunctionReferenced(LLLLoc, ConfigDecl); return ActOnCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr, /*IsExecConfig=*/true); }
void ExprEngine::VisitCXXNewExpr(const CXXNewExpr *CNE, ExplodedNode *Pred, ExplodedNodeSet &Dst) { // FIXME: Much of this should eventually migrate to CXXAllocatorCall. // Also, we need to decide how allocators actually work -- they're not // really part of the CXXNewExpr because they happen BEFORE the // CXXConstructExpr subexpression. See PR12014 for some discussion. StmtNodeBuilder Bldr(Pred, Dst, *currBldrCtx); unsigned blockCount = currBldrCtx->blockCount(); const LocationContext *LCtx = Pred->getLocationContext(); DefinedOrUnknownSVal symVal = svalBuilder.conjureSymbolVal(0, CNE, LCtx, CNE->getType(), blockCount); ProgramStateRef State = Pred->getState(); CallEventManager &CEMgr = getStateManager().getCallEventManager(); CallEventRef<CXXAllocatorCall> Call = CEMgr.getCXXAllocatorCall(CNE, State, LCtx); // Invalidate placement args. // FIXME: Once we figure out how we want allocators to work, // we should be using the usual pre-/(default-)eval-/post-call checks here. State = Call->invalidateRegions(blockCount); // If we're compiling with exceptions enabled, and this allocation function // is not declared as non-throwing, failures /must/ be signalled by // exceptions, and thus the return value will never be NULL. // C++11 [basic.stc.dynamic.allocation]p3. FunctionDecl *FD = CNE->getOperatorNew(); if (FD && getContext().getLangOpts().CXXExceptions) { QualType Ty = FD->getType(); if (const FunctionProtoType *ProtoType = Ty->getAs<FunctionProtoType>()) if (!ProtoType->isNothrow(getContext())) State = State->assume(symVal, true); } if (CNE->isArray()) { // FIXME: allocating an array requires simulating the constructors. // For now, just return a symbolicated region. const MemRegion *NewReg = cast<loc::MemRegionVal>(symVal).getRegion(); QualType ObjTy = CNE->getType()->getAs<PointerType>()->getPointeeType(); const ElementRegion *EleReg = getStoreManager().GetElementZeroRegion(NewReg, ObjTy); State = State->BindExpr(CNE, Pred->getLocationContext(), loc::MemRegionVal(EleReg)); Bldr.generateNode(CNE, Pred, State); return; } // FIXME: Once we have proper support for CXXConstructExprs inside // CXXNewExpr, we need to make sure that the constructed object is not // immediately invalidated here. (The placement call should happen before // the constructor call anyway.) if (FD && FD->isReservedGlobalPlacementOperator()) { // Non-array placement new should always return the placement location. SVal PlacementLoc = State->getSVal(CNE->getPlacementArg(0), LCtx); SVal Result = svalBuilder.evalCast(PlacementLoc, CNE->getType(), CNE->getPlacementArg(0)->getType()); State = State->BindExpr(CNE, LCtx, Result); } else { State = State->BindExpr(CNE, LCtx, symVal); } // If the type is not a record, we won't have a CXXConstructExpr as an // initializer. Copy the value over. if (const Expr *Init = CNE->getInitializer()) { if (!isa<CXXConstructExpr>(Init)) { QualType ObjTy = CNE->getType()->getAs<PointerType>()->getPointeeType(); (void)ObjTy; assert(!ObjTy->isRecordType()); SVal Location = State->getSVal(CNE, LCtx); if (isa<Loc>(Location)) State = State->bindLoc(cast<Loc>(Location), State->getSVal(Init, LCtx)); } } Bldr.generateNode(CNE, Pred, State); }
void ExprEngine::VisitCXXNewExpr(const CXXNewExpr *CNE, ExplodedNode *Pred, ExplodedNodeSet &Dst) { // FIXME: Much of this should eventually migrate to CXXAllocatorCall. // Also, we need to decide how allocators actually work -- they're not // really part of the CXXNewExpr because they happen BEFORE the // CXXConstructExpr subexpression. See PR12014 for some discussion. unsigned blockCount = currBldrCtx->blockCount(); const LocationContext *LCtx = Pred->getLocationContext(); DefinedOrUnknownSVal symVal = UnknownVal(); FunctionDecl *FD = CNE->getOperatorNew(); bool IsStandardGlobalOpNewFunction = false; if (FD && !isa<CXXMethodDecl>(FD) && !FD->isVariadic()) { if (FD->getNumParams() == 2) { QualType T = FD->getParamDecl(1)->getType(); if (const IdentifierInfo *II = T.getBaseTypeIdentifier()) // NoThrow placement new behaves as a standard new. IsStandardGlobalOpNewFunction = II->getName().equals("nothrow_t"); } else // Placement forms are considered non-standard. IsStandardGlobalOpNewFunction = (FD->getNumParams() == 1); } // We assume all standard global 'operator new' functions allocate memory in // heap. We realize this is an approximation that might not correctly model // a custom global allocator. if (IsStandardGlobalOpNewFunction) symVal = svalBuilder.getConjuredHeapSymbolVal(CNE, LCtx, blockCount); else symVal = svalBuilder.conjureSymbolVal(nullptr, CNE, LCtx, CNE->getType(), blockCount); ProgramStateRef State = Pred->getState(); CallEventManager &CEMgr = getStateManager().getCallEventManager(); CallEventRef<CXXAllocatorCall> Call = CEMgr.getCXXAllocatorCall(CNE, State, LCtx); // Invalidate placement args. // FIXME: Once we figure out how we want allocators to work, // we should be using the usual pre-/(default-)eval-/post-call checks here. State = Call->invalidateRegions(blockCount); if (!State) return; // If this allocation function is not declared as non-throwing, failures // /must/ be signalled by exceptions, and thus the return value will never be // NULL. -fno-exceptions does not influence this semantics. // FIXME: GCC has a -fcheck-new option, which forces it to consider the case // where new can return NULL. If we end up supporting that option, we can // consider adding a check for it here. // C++11 [basic.stc.dynamic.allocation]p3. if (FD) { QualType Ty = FD->getType(); if (const FunctionProtoType *ProtoType = Ty->getAs<FunctionProtoType>()) if (!ProtoType->isNothrow(getContext())) State = State->assume(symVal, true); } StmtNodeBuilder Bldr(Pred, Dst, *currBldrCtx); if (CNE->isArray()) { // FIXME: allocating an array requires simulating the constructors. // For now, just return a symbolicated region. const MemRegion *NewReg = symVal.castAs<loc::MemRegionVal>().getRegion(); QualType ObjTy = CNE->getType()->getAs<PointerType>()->getPointeeType(); const ElementRegion *EleReg = getStoreManager().GetElementZeroRegion(NewReg, ObjTy); State = State->BindExpr(CNE, Pred->getLocationContext(), loc::MemRegionVal(EleReg)); Bldr.generateNode(CNE, Pred, State); return; } // FIXME: Once we have proper support for CXXConstructExprs inside // CXXNewExpr, we need to make sure that the constructed object is not // immediately invalidated here. (The placement call should happen before // the constructor call anyway.) SVal Result = symVal; if (FD && FD->isReservedGlobalPlacementOperator()) { // Non-array placement new should always return the placement location. SVal PlacementLoc = State->getSVal(CNE->getPlacementArg(0), LCtx); Result = svalBuilder.evalCast(PlacementLoc, CNE->getType(), CNE->getPlacementArg(0)->getType()); } // Bind the address of the object, then check to see if we cached out. State = State->BindExpr(CNE, LCtx, Result); ExplodedNode *NewN = Bldr.generateNode(CNE, Pred, State); if (!NewN) return; // If the type is not a record, we won't have a CXXConstructExpr as an // initializer. Copy the value over. if (const Expr *Init = CNE->getInitializer()) { if (!isa<CXXConstructExpr>(Init)) { assert(Bldr.getResults().size() == 1); Bldr.takeNodes(NewN); evalBind(Dst, CNE, NewN, Result, State->getSVal(Init, LCtx), /*FirstInit=*/IsStandardGlobalOpNewFunction); } } }
llvm::Value *CodeGenFunction::EmitCXXNewExpr(const CXXNewExpr *E) { QualType AllocType = E->getAllocatedType(); FunctionDecl *NewFD = E->getOperatorNew(); const FunctionProtoType *NewFTy = NewFD->getType()->getAs<FunctionProtoType>(); CallArgList NewArgs; // The allocation size is the first argument. QualType SizeTy = getContext().getSizeType(); llvm::Value *NumElements = 0; llvm::Value *AllocSize = EmitCXXNewAllocSize(*this, E, NumElements); NewArgs.push_back(std::make_pair(RValue::get(AllocSize), SizeTy)); // Emit the rest of the arguments. // FIXME: Ideally, this should just use EmitCallArgs. CXXNewExpr::const_arg_iterator NewArg = E->placement_arg_begin(); // First, use the types from the function type. // We start at 1 here because the first argument (the allocation size) // has already been emitted. for (unsigned i = 1, e = NewFTy->getNumArgs(); i != e; ++i, ++NewArg) { QualType ArgType = NewFTy->getArgType(i); assert(getContext().getCanonicalType(ArgType.getNonReferenceType()). getTypePtr() == getContext().getCanonicalType(NewArg->getType()).getTypePtr() && "type mismatch in call argument!"); NewArgs.push_back(std::make_pair(EmitCallArg(*NewArg, ArgType), ArgType)); } // Either we've emitted all the call args, or we have a call to a // variadic function. assert((NewArg == E->placement_arg_end() || NewFTy->isVariadic()) && "Extra arguments in non-variadic function!"); // If we still have any arguments, emit them using the type of the argument. for (CXXNewExpr::const_arg_iterator NewArgEnd = E->placement_arg_end(); NewArg != NewArgEnd; ++NewArg) { QualType ArgType = NewArg->getType(); NewArgs.push_back(std::make_pair(EmitCallArg(*NewArg, ArgType), ArgType)); } // Emit the call to new. RValue RV = EmitCall(CGM.getTypes().getFunctionInfo(NewFTy->getResultType(), NewArgs), CGM.GetAddrOfFunction(NewFD), NewArgs, NewFD); // If an allocation function is declared with an empty exception specification // it returns null to indicate failure to allocate storage. [expr.new]p13. // (We don't need to check for null when there's no new initializer and // we're allocating a POD type). bool NullCheckResult = NewFTy->hasEmptyExceptionSpec() && !(AllocType->isPODType() && !E->hasInitializer()); llvm::BasicBlock *NewNull = 0; llvm::BasicBlock *NewNotNull = 0; llvm::BasicBlock *NewEnd = 0; llvm::Value *NewPtr = RV.getScalarVal(); if (NullCheckResult) { NewNull = createBasicBlock("new.null"); NewNotNull = createBasicBlock("new.notnull"); NewEnd = createBasicBlock("new.end"); llvm::Value *IsNull = Builder.CreateICmpEQ(NewPtr, llvm::Constant::getNullValue(NewPtr->getType()), "isnull"); Builder.CreateCondBr(IsNull, NewNull, NewNotNull); EmitBlock(NewNotNull); } if (uint64_t CookiePadding = CalculateCookiePadding(getContext(), E)) { uint64_t CookieOffset = CookiePadding - getContext().getTypeSize(SizeTy) / 8; llvm::Value *NumElementsPtr = Builder.CreateConstInBoundsGEP1_64(NewPtr, CookieOffset); NumElementsPtr = Builder.CreateBitCast(NumElementsPtr, ConvertType(SizeTy)->getPointerTo()); Builder.CreateStore(NumElements, NumElementsPtr); // Now add the padding to the new ptr. NewPtr = Builder.CreateConstInBoundsGEP1_64(NewPtr, CookiePadding); } NewPtr = Builder.CreateBitCast(NewPtr, ConvertType(E->getType())); EmitNewInitializer(*this, E, NewPtr, NumElements); if (NullCheckResult) { Builder.CreateBr(NewEnd); NewNotNull = Builder.GetInsertBlock(); EmitBlock(NewNull); Builder.CreateBr(NewEnd); EmitBlock(NewEnd); llvm::PHINode *PHI = Builder.CreatePHI(NewPtr->getType()); PHI->reserveOperandSpace(2); PHI->addIncoming(NewPtr, NewNotNull); PHI->addIncoming(llvm::Constant::getNullValue(NewPtr->getType()), NewNull); NewPtr = PHI; } return NewPtr; }
bool CoroutineStmtBuilder::makeNewAndDeleteExpr() { // Form and check allocation and deallocation calls. assert(!IsPromiseDependentType && "cannot make statement while the promise type is dependent"); QualType PromiseType = Fn.CoroutinePromise->getType(); if (S.RequireCompleteType(Loc, PromiseType, diag::err_incomplete_type)) return false; const bool RequiresNoThrowAlloc = ReturnStmtOnAllocFailure != nullptr; // FIXME: Add support for stateful allocators. FunctionDecl *OperatorNew = nullptr; FunctionDecl *OperatorDelete = nullptr; FunctionDecl *UnusedResult = nullptr; bool PassAlignment = false; SmallVector<Expr *, 1> PlacementArgs; S.FindAllocationFunctions(Loc, SourceRange(), /*UseGlobal*/ false, PromiseType, /*isArray*/ false, PassAlignment, PlacementArgs, OperatorNew, UnusedResult); bool IsGlobalOverload = OperatorNew && !isa<CXXRecordDecl>(OperatorNew->getDeclContext()); // If we didn't find a class-local new declaration and non-throwing new // was is required then we need to lookup the non-throwing global operator // instead. if (RequiresNoThrowAlloc && (!OperatorNew || IsGlobalOverload)) { auto *StdNoThrow = buildStdNoThrowDeclRef(S, Loc); if (!StdNoThrow) return false; PlacementArgs = {StdNoThrow}; OperatorNew = nullptr; S.FindAllocationFunctions(Loc, SourceRange(), /*UseGlobal*/ true, PromiseType, /*isArray*/ false, PassAlignment, PlacementArgs, OperatorNew, UnusedResult); } assert(OperatorNew && "expected definition of operator new to be found"); if (RequiresNoThrowAlloc) { const auto *FT = OperatorNew->getType()->getAs<FunctionProtoType>(); if (!FT->isNothrow(S.Context, /*ResultIfDependent*/ false)) { S.Diag(OperatorNew->getLocation(), diag::err_coroutine_promise_new_requires_nothrow) << OperatorNew; S.Diag(Loc, diag::note_coroutine_promise_call_implicitly_required) << OperatorNew; return false; } } if ((OperatorDelete = findDeleteForPromise(S, Loc, PromiseType)) == nullptr) return false; Expr *FramePtr = buildBuiltinCall(S, Loc, Builtin::BI__builtin_coro_frame, {}); Expr *FrameSize = buildBuiltinCall(S, Loc, Builtin::BI__builtin_coro_size, {}); // Make new call. ExprResult NewRef = S.BuildDeclRefExpr(OperatorNew, OperatorNew->getType(), VK_LValue, Loc); if (NewRef.isInvalid()) return false; SmallVector<Expr *, 2> NewArgs(1, FrameSize); for (auto Arg : PlacementArgs) NewArgs.push_back(Arg); ExprResult NewExpr = S.ActOnCallExpr(S.getCurScope(), NewRef.get(), Loc, NewArgs, Loc); NewExpr = S.ActOnFinishFullExpr(NewExpr.get()); if (NewExpr.isInvalid()) return false; // Make delete call. QualType OpDeleteQualType = OperatorDelete->getType(); ExprResult DeleteRef = S.BuildDeclRefExpr(OperatorDelete, OpDeleteQualType, VK_LValue, Loc); if (DeleteRef.isInvalid()) return false; Expr *CoroFree = buildBuiltinCall(S, Loc, Builtin::BI__builtin_coro_free, {FramePtr}); SmallVector<Expr *, 2> DeleteArgs{CoroFree}; // Check if we need to pass the size. const auto *OpDeleteType = OpDeleteQualType.getTypePtr()->getAs<FunctionProtoType>(); if (OpDeleteType->getNumParams() > 1) DeleteArgs.push_back(FrameSize); ExprResult DeleteExpr = S.ActOnCallExpr(S.getCurScope(), DeleteRef.get(), Loc, DeleteArgs, Loc); DeleteExpr = S.ActOnFinishFullExpr(DeleteExpr.get()); if (DeleteExpr.isInvalid()) return false; this->Allocate = NewExpr.get(); this->Deallocate = DeleteExpr.get(); return true; }
// 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); }