Home - News ( United States | United Kingdom | Italy | Germany ) - Football scores

Showing content from https://clang.llvm.org/doxygen/CGOpenMPRuntimeGPU_8cpp_source.html below:

clang: lib/CodeGen/CGOpenMPRuntimeGPU.cpp Source File

;

CodeGen;

llvm::omp;

llvm::FunctionCallee EnterCallee =

;

llvm::FunctionCallee ExitCallee =

;

llvm::BasicBlock *ContBlock =

;

NVPTXActionTy(llvm::FunctionCallee EnterCallee,

llvm::FunctionCallee ExitCallee,

: EnterCallee(EnterCallee), EnterArgs(EnterArgs), ExitCallee(ExitCallee),

llvm::Value *CallBool = CGF.

.CreateIsNotNull(EnterRes);

CGF.

.CreateCondBr(CallBool, ThenBlock, ContBlock);

ExecutionRuntimeModesRAII {

: ExecMode(ExecMode) {

SavedExecMode = ExecMode;

~ExecutionRuntimeModesRAII() { ExecMode = SavedExecMode; }

(

*ASE = dyn_cast<ArraySubscriptExpr>(RefExpr)) {

*

= ASE->getBase()->IgnoreParenImpCasts();

(

*TempASE = dyn_cast<ArraySubscriptExpr>(

))

= TempASE->getBase()->IgnoreParenImpCasts();

}

(

*OASE = dyn_cast<ArraySectionExpr>(RefExpr)) {

*

= OASE->getBase()->IgnoreParenImpCasts();

(

*TempOASE = dyn_cast<ArraySectionExpr>(

))

= TempOASE->getBase()->IgnoreParenImpCasts();

(

*TempASE = dyn_cast<ArraySubscriptExpr>(

))

= TempASE->getBase()->IgnoreParenImpCasts();

(

*DE = dyn_cast<DeclRefExpr>(RefExpr))

cast<ValueDecl>(DE->getDecl()->getCanonicalDecl());

*ME = cast<MemberExpr>(RefExpr);

cast<ValueDecl>(ME->getMemberDecl()->getCanonicalDecl());

*buildRecordForGlobalizedVars(

llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>

(EscapedDecls.empty() && EscapedDeclsForTeams.empty())

GlobalizedVars.emplace_back(

.getDeclAlign(

),

);

(

*

: EscapedDeclsForTeams)

GlobalizedVars.emplace_back(

.getDeclAlign(

),

);

*GlobalizedRD =

.buildImplicitRecord(

);

EscapedDeclsForTeams.begin(), EscapedDeclsForTeams.end());

(

&Pair : GlobalizedVars) {

=

.getPointerType(

.getNonReferenceType());

(SingleEscaped.count(VD)) {

llvm::APInt ArraySize(32, BufSize);

=

.getConstantArrayType(

, ArraySize,

,

ArraySizeModifier::Normal, 0);

llvm::APInt Align(32, Pair.first.getQuantity());

->addAttr(AlignedAttr::CreateImplicit(

.getIntTypeForBitwidth(32,

0),

{}, AlignedAttr::GNU_aligned));

GlobalizedRD->

(Field);

MappedDeclsFields.try_emplace(VD, Field);

CheckVarsEscapingDeclContext final

llvm::SetVector<const ValueDecl *> EscapedDecls;

llvm::SetVector<const ValueDecl *> EscapedVariableLengthDecls;

llvm::SetVector<const ValueDecl *> DelayedVariableLengthDecls;

llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;

AllEscaped =

;

IsForCombinedParallelRegion =

;

markAsEscaped(

*VD) {

(!isa<VarDecl>(VD) ||

OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))

IsCaptured =

;

(

*CSI = CGF.CapturedStmtInfo) {

(

*FD = CSI->lookup(cast<VarDecl>(VD))) {

(!IsForCombinedParallelRegion) {

*

= FD->getAttr<OMPCaptureKindAttr>();

(((

->getCaptureKind() != OMPC_map) &&

((

->getCaptureKind() == OMPC_map) &&

!FD->getType()->isAnyPointerType()))

(!FD->getType()->isReferenceType()) {

);

EscapedParameters.insert(VD);

}

(!IsForCombinedParallelRegion) {

((!CGF.CapturedStmtInfo ||

(IsForCombinedParallelRegion && CGF.CapturedStmtInfo)) &&

EscapedVariableLengthDecls.insert(VD);

DelayedVariableLengthDecls.insert(VD);

EscapedDecls.insert(VD);

VisitValueDecl(

*VD) {

(

*VarD = dyn_cast<VarDecl>(VD)) {

(!isa<ParmVarDecl>(VarD) && VarD->hasInit()) {

SavedAllEscaped = AllEscaped;

(VarD->getInit());

AllEscaped = SavedAllEscaped;

IsCombinedParallelRegion) {

(

.capturesVariable() && !

.capturesVariableByCopy()) {

SavedIsForCombinedParallelRegion = IsForCombinedParallelRegion;

(IsCombinedParallelRegion) {

IsForCombinedParallelRegion =

;

->getClauseKind() == OMPC_reduction ||

->getClauseKind() == OMPC_linear ||

->getClauseKind() == OMPC_private)

(

*PC = dyn_cast<OMPFirstprivateClause>(

))

Vars = PC->getVarRefs();

(

*PC = dyn_cast<OMPLastprivateClause>(

))

Vars = PC->getVarRefs();

llvm_unreachable(

);

(

*

: Vars) {

IsForCombinedParallelRegion =

;

(IsForCombinedParallelRegion)

(isa<OMPCapturedExprDecl>(VD))

IsForCombinedParallelRegion = SavedIsForCombinedParallelRegion;

buildRecordForGlobalizedVars(

IsInTTDRegion) {

assert(!GlobalizedRD &&

);

WarpSize = CGF.getTarget().getGridValue().GV_Warp_Size;

EscapedDeclsForTeams = EscapedDecls.getArrayRef();

EscapedDeclsForParallel = EscapedDecls.getArrayRef();

GlobalizedRD = ::buildRecordForGlobalizedVars(

CGF.getContext(), EscapedDeclsForParallel, EscapedDeclsForTeams,

MappedDeclsFields, WarpSize);

: CGF(CGF), EscapedDecls(TeamsReductions.begin(), TeamsReductions.end()) {

~CheckVarsEscapingDeclContext() =

;

VisitDeclStmt(

*S) {

(

*

: S->decls())

(

*VD = dyn_cast_or_null<ValueDecl>(

))

(!

->hasAssociatedStmt())

dyn_cast_or_null<CapturedStmt>(

->getAssociatedStmt())) {

(CaptureRegions.size() == 1 && CaptureRegions.back() == OMPD_unknown) {

VisitStmt(S->getCapturedStmt());

VisitOpenMPCapturedStmt(

CaptureRegions.back() == OMPD_parallel &&

(

.capturesVariable() && !

.capturesVariableByCopy()) {

(isa<OMPCapturedExprDecl>(VD))

(

.capturesVariable()) {

(

->isInitCapture(&

) || isa<OMPCapturedExprDecl>(VD))

*VD =

.getVariable();

VisitCallExpr(

*

) {

(

*Arg :

->arguments()) {

(Arg->isLValue()) {

SavedAllEscaped = AllEscaped;

AllEscaped = SavedAllEscaped;

(isa<OMPCapturedExprDecl>(VD))

(

->getOpcode() == UO_AddrOf) {

SavedAllEscaped = AllEscaped;

(

->getSubExpr());

AllEscaped = SavedAllEscaped;

(

->getSubExpr());

(

->getCastKind() == CK_ArrayToPointerDecay) {

SavedAllEscaped = AllEscaped;

(

->getSubExpr());

AllEscaped = SavedAllEscaped;

(

->getSubExpr());

VisitExpr(

*

) {

SavedAllEscaped = AllEscaped;

AllEscaped = SavedAllEscaped;

VisitStmt(

*S) {

(

*Child : S->children())

*getGlobalizedRecord(

IsInTTDRegion) {

buildRecordForGlobalizedVars(IsInTTDRegion);

assert(GlobalizedRD &&

);

MappedDeclsFields.lookup(VD);

EscapedDecls.getArrayRef();

llvm::SmallPtrSetImpl<const Decl *> &getEscapedParameters()

{

EscapedParameters;

EscapedVariableLengthDecls.getArrayRef();

DelayedVariableLengthDecls.getArrayRef();

CGOpenMPRuntimeGPU::getExecutionMode()

{

CurrentExecutionMode;

CGOpenMPRuntimeGPU::getDataSharingMode()

{

CurrentDataSharingMode;

*CS =

.getInnermostCapturedStmt();

CS->getCapturedStmt()->IgnoreContainers(

);

(

*NestedDir =

dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {

(

.getDirectiveKind()) {

(DKind == OMPD_teams) {

Body = NestedDir->getInnermostCapturedStmt()->IgnoreContainers(

(

*NND =

dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {

DKind = NND->getDirectiveKind();

OMPD_target_teams:

OMPD_target_simd:

OMPD_target_parallel:

OMPD_target_parallel_for:

OMPD_target_parallel_for_simd:

OMPD_target_teams_distribute:

OMPD_target_teams_distribute_simd:

OMPD_target_teams_distribute_parallel_for:

OMPD_target_teams_distribute_parallel_for_simd:

OMPD_parallel_for:

OMPD_parallel_master:

OMPD_parallel_sections:

OMPD_parallel_for_simd:

OMPD_cancellation_point:

OMPD_threadprivate:

OMPD_target_data:

OMPD_target_exit_data:

OMPD_target_enter_data:

OMPD_distribute:

OMPD_distribute_simd:

OMPD_distribute_parallel_for:

OMPD_distribute_parallel_for_simd:

OMPD_teams_distribute:

OMPD_teams_distribute_simd:

OMPD_teams_distribute_parallel_for:

OMPD_teams_distribute_parallel_for_simd:

OMPD_target_update:

OMPD_declare_simd:

OMPD_declare_variant:

OMPD_begin_declare_variant:

OMPD_end_declare_variant:

OMPD_declare_target:

OMPD_end_declare_target:

OMPD_declare_reduction:

OMPD_declare_mapper:

OMPD_taskloop_simd:

OMPD_master_taskloop:

OMPD_master_taskloop_simd:

OMPD_parallel_master_taskloop:

OMPD_parallel_master_taskloop_simd:

llvm_unreachable(

);

(DirectiveKind) {

OMPD_target_teams:

OMPD_target_parallel_loop:

OMPD_target_parallel:

OMPD_target_parallel_for:

OMPD_target_parallel_for_simd:

OMPD_target_teams_distribute_parallel_for:

OMPD_target_teams_distribute_parallel_for_simd:

OMPD_target_simd:

OMPD_target_teams_distribute_simd:

OMPD_target_teams_distribute:

OMPD_target_teams_loop:

(

*TTLD = dyn_cast<OMPTargetTeamsGenericLoopDirective>(&

))

TTLD->canBeParallelFor();

OMPD_parallel_for:

OMPD_parallel_master:

OMPD_parallel_sections:

OMPD_parallel_for_simd:

OMPD_cancellation_point:

OMPD_threadprivate:

OMPD_target_data:

OMPD_target_exit_data:

OMPD_target_enter_data:

OMPD_distribute:

OMPD_distribute_simd:

OMPD_distribute_parallel_for:

OMPD_distribute_parallel_for_simd:

OMPD_teams_distribute:

OMPD_teams_distribute_simd:

OMPD_teams_distribute_parallel_for:

OMPD_teams_distribute_parallel_for_simd:

OMPD_target_update:

OMPD_declare_simd:

OMPD_declare_variant:

OMPD_begin_declare_variant:

OMPD_end_declare_variant:

OMPD_declare_target:

OMPD_end_declare_target:

OMPD_declare_reduction:

OMPD_declare_mapper:

OMPD_taskloop_simd:

OMPD_master_taskloop:

OMPD_master_taskloop_simd:

OMPD_parallel_master_taskloop:

OMPD_parallel_master_taskloop_simd:

);

StringRef ParentName,

llvm::Function *&OutlinedFn,

llvm::Constant *&OutlinedFnID,

ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode,

);

EntryFunctionState EST;

WrapperFunctionsMap.clear();

[[maybe_unused]]

IsBareKernel =

.getSingleClause<

>();

assert(!IsBareKernel &&

);

CGOpenMPRuntimeGPU::EntryFunctionState &EST;

NVPTXPrePostActionTy(CGOpenMPRuntimeGPU::EntryFunctionState &EST,

: EST(EST),

(

) {}

RT.emitKernelInit(

, CGF, EST,

);

RT.setLocThreadIdInsertPt(CGF,

);

RT.emitKernelDeinit(CGF, EST,

);

IsInTTDRegion =

;

IsOffloadEntry, CodeGen);

IsInTTDRegion =

;

EntryFunctionState &EST,

IsSPMD) {

llvm::OpenMPIRBuilder::TargetKernelDefaultAttrs Attrs;

IsSPMD ? llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD

: llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_GENERIC;

Bld.restoreIP(

.createTargetInit(Bld, Attrs));

emitGenericVarsProlog(CGF, EST.Loc);

EntryFunctionState &EST,

emitGenericVarsEpilog(CGF);

, RecordDecl::TagKind::Union);

(

*TeamReductionRec : TeamsReductions) {

RecTy =

.getRecordType(TeamReductionRec);

StaticTy =

.getRecordType(StaticRD);

llvm::Type *LLVMReductionsBufferTy =

TeamsReductions.empty()

: DL.getTypeAllocSize(LLVMReductionsBufferTy).getFixedValue();

.createTargetDeinit(Bld, ReductionDataSize,

.getLangOpts().OpenMPCUDAReductionBufNum);

TeamsReductions.clear();

StringRef ParentName,

llvm::Function *&OutlinedFn,

llvm::Constant *&OutlinedFnID,

ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode,

);

EntryFunctionState EST;

CGOpenMPRuntimeGPU::EntryFunctionState &EST;

DataSharingMode Mode;

CGOpenMPRuntimeGPU::EntryFunctionState &EST,

: RT(RT), EST(EST), IsBareKernel(IsBareKernel),

Mode(RT.CurrentDataSharingMode),

(

) {}

RT.CurrentDataSharingMode = DataSharingMode::DS_CUDA;

RT.emitKernelInit(

, CGF, EST,

);

RT.setLocThreadIdInsertPt(CGF,

);

RT.CurrentDataSharingMode = Mode;

RT.clearLocThreadIdInsertPt(CGF);

RT.emitKernelDeinit(CGF, EST,

);

} Action(*

, EST, IsBareKernel,

);

IsInTTDRegion =

;

IsOffloadEntry, CodeGen);

IsInTTDRegion =

;

CGOpenMPRuntimeGPU::emitTargetOutlinedFunction(

llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,

assert(!ParentName.empty() &&

);

(Mode || IsBareKernel)

emitSPMDKernel(

, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,

emitNonSPMDKernel(

, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,

llvm::OpenMPIRBuilderConfig Config(

llvm_unreachable(

);

);

);

);

);

);

ProcBindKind ProcBind,

llvm::Value *NumThreads,

*NumTeams,

*ThreadLimit,

PrevIsInTTDRegion = IsInTTDRegion;

IsInTTDRegion =

;

CGF,

, ThreadIDVar, InnermostKind, CodeGen));

IsInTTDRegion = PrevIsInTTDRegion;

llvm::Function *WrapperFun =

createParallelDataSharingWrapper(OutlinedFun,

);

WrapperFunctionsMap[OutlinedFun] = WrapperFun;

);

.getInnermostCapturedStmt()->getCapturedStmt()->IgnoreContainers(

Dir = dyn_cast_or_null<OMPExecutableDirective>(S);

(

*

:

->getVarRefs())

);

(

*

:

->privates())

llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;

(!LastPrivatesReductions.empty()) {

GlobalizedRD = ::buildRecordForGlobalizedVars(

.

(), {}, LastPrivatesReductions, MappedDeclsFields,

}

(!LastPrivatesReductions.empty()) {

assert(!TeamAndReductions.first &&

);

TeamAndReductions.first =

.getCapturedStmt(OMPD_teams)->getCapturedDecl();

std::swap(TeamAndReductions.second, LastPrivatesReductions);

llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>

NVPTXPrePostActionTy(

llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>

:

(

), GlobalizedRD(GlobalizedRD),

MappedDeclsFields(MappedDeclsFields) {}

I = Rt.FunctionGlobalizedDecls.try_emplace(CGF.

).first;

I->getSecond().MappedParams =

std::make_unique<CodeGenFunction::OMPMapVars>();

DeclToAddrMapTy &

= I->getSecond().LocalVarData;

(

&Pair : MappedDeclsFields) {

assert(Pair.getFirst()->isCanonicalDecl() &&

);

.insert(std::make_pair(Pair.getFirst(), MappedVarData()));

Rt.emitGenericVarsProlog(CGF,

);

.emitGenericVarsEpilog(CGF);

} Action(

, GlobalizedRD, MappedDeclsFields);

CGF,

, ThreadIDVar, InnermostKind, CodeGen);

I = FunctionGlobalizedDecls.find(CGF.

);

(I == FunctionGlobalizedDecls.end())

(

&Rec : I->getSecond().LocalVarData) {

*VD = cast<VarDecl>(Rec.first);

EscapedParam = I->getSecond().EscapedParameters.count(Rec.first);

llvm::Value *ParValue;

llvm::CallBase *VoidPtr =

VoidPtr->addRetAttr(llvm::Attribute::get(

VoidPtr, Bld.getPtrTy(0), VD->

() +

);

Rec.second.PrivateAddr = VarAddr.

();

Rec.second.GlobalizedVal = VoidPtr;

I->getSecond().MappedParams->setVarAddr(CGF, VD, VarAddr.

());

VoidPtr->setDebugLoc(DI->SourceLocToDebugLoc(VD->

()));

(

*ValueD : I->getSecond().EscapedVariableLengthDecls) {

*VD = cast<VarDecl>(ValueD);

std::pair<llvm::Value *, llvm::Value *> AddrSizePair =

I->getSecond().EscapedVariableLengthDeclsAddrs.emplace_back(AddrSizePair);

I->getSecond().MappedParams->setVarAddr(CGF, VD,

.getAddress());

I->getSecond().MappedParams->apply(CGF);

I = FunctionGlobalizedDecls.find(CGF.

);

(I == FunctionGlobalizedDecls.end())

llvm::is_contained(I->getSecond().DelayedVariableLengthDecls, VD);

std::pair<llvm::Value *, llvm::Value *>

Size = Bld.CreateNUWAdd(

llvm::Value *AlignVal =

Size = Bld.CreateUDiv(Size, AlignVal);

Size = Bld.CreateNUWMul(Size, AlignVal);

llvm::Value *AllocArgs[] = {Size};

llvm::CallBase *VoidPtr =

VoidPtr->addRetAttr(llvm::Attribute::get(

std::make_pair(VoidPtr, Size);

std::pair<llvm::Value *, llvm::Value *> &AddrSizePair) {

{AddrSizePair.first, AddrSizePair.second});

CGOpenMPRuntimeGPU::emitGenericVarsEpilog(

&CGF) {

I = FunctionGlobalizedDecls.find(CGF.

);

(I != FunctionGlobalizedDecls.end()) {

(

&AddrSizePair :

llvm::reverse(I->getSecond().EscapedVariableLengthDeclsAddrs)) {

{AddrSizePair.first, AddrSizePair.second});

(

&Rec : llvm::reverse(I->getSecond().LocalVarData)) {

*VD = cast<VarDecl>(Rec.first);

I->getSecond().MappedParams->restore(CGF);

llvm::Value *FreeArgs[] = {Rec.second.GlobalizedVal,

llvm::Function *OutlinedFn,

OutlinedFnArgs.push_back(llvm::ConstantPointerNull::get(

.

));

OutlinedFnArgs.push_back(ZeroAddr.getPointer());

OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end());

llvm::Function *OutlinedFn,

llvm::Value *NumThreads) {

&&ParallelGen = [

,

, OutlinedFn, CapturedVars, IfCond,

llvm::Value *NumThreadsVal = NumThreads;

llvm::Function *WFn = WrapperFunctionsMap[OutlinedFn];

llvm::Value *ID = llvm::ConstantPointerNull::get(

.

);

ID = Bld.CreateBitOrPointerCast(WFn,

.

);

llvm::Value *FnPtr = Bld.CreateBitOrPointerCast(OutlinedFn,

.

);

llvm::ArrayType::get(

.

, CapturedVars.size()),

);

(!CapturedVars.empty()) {

(llvm::Value *

: CapturedVars) {

(

->getType()->isIntegerTy())

PtrV = Bld.CreateIntToPtr(

, CGF.

);

llvm::Value *IfCondVal =

;

IfCondVal = llvm::ConstantInt::get(CGF.

, 1);

NumThreadsVal = llvm::ConstantInt::get(CGF.

, -1);

NumThreadsVal = Bld.CreateZExtOrTrunc(NumThreadsVal, CGF.

),

assert(IfCondVal &&

);

llvm::Value *Args[] = {

llvm::ConstantInt::get(CGF.

, -1),

Bld.CreateBitOrPointerCast(CapturedVarsAddrs.

(CGF),

llvm::ConstantInt::get(

.

, CapturedVars.size())};

llvm::Value *Args[] = {

llvm::ConstantPointerNull::get(

llvm::ConstantInt::get(CGF.

,

0,

)};

.

(), OMPRTL___kmpc_barrier_simple_spmd),

.

(), OMPRTL___kmpc_warp_active_thread_mask));

llvm::Value *ThreadID = RT.getGPUThreadID(CGF);

llvm::Value *TeamWidth = RT.getGPUNumThreads(CGF);

llvm::Value *CmpLoopBound = CGF.

.CreateICmpSLT(CounterVal, TeamWidth);

CGF.

.CreateCondBr(CmpLoopBound, TestBB, ExitBB);

llvm::Value *CmpThreadToCounter =

CGF.

.CreateICmpEQ(ThreadID, CounterVal);

CGF.

.CreateCondBr(CmpThreadToCounter, BodyBB, SyncBB);

llvm::Value *IncCounterVal =

CGF.

.CreateNSWAdd(CounterVal, CGF.

.getInt32(1));

);

);

(ValTy == CastTy)

CGF.

.CreateBitCast(Val, LLVMCastTy);

CGF.

.CreateIntCast(Val, LLVMCastTy,

(Options.SimpleReduction) {

assert(!TeamsReduction && !ParallelReduction &&

);

(void)ParallelReduction;

ReductionOps, Options);

llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> VarFieldMap;

(

*DRE : Privates) {

PrivatesReductions[Cnt] = cast<DeclRefExpr>(DRE)->getDecl();

*ReductionRec = ::buildRecordForGlobalizedVars(

.

(), PrivatesReductions, {}, VarFieldMap, 1);

TeamsReductions.push_back(ReductionRec);

InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;

InsertPointTy CodeGenIP(CGF.

.GetInsertBlock(),

CGF.

.GetInsertPoint());

llvm::OpenMPIRBuilder::LocationDescription OmpLoc(

llvm::Type *ElementType;

llvm::Value *PrivateVariable;

llvm::OpenMPIRBuilder::ReductionGenAtomicCBTy AtomicReductionGen =

;

*RHSVar =

cast<VarDecl>(cast<DeclRefExpr>(RHSExprs[Idx])->getDecl());

*LHSVar =

cast<VarDecl>(cast<DeclRefExpr>(LHSExprs[Idx])->getDecl());

llvm::OpenMPIRBuilder::EvalKind EvalKind;

EvalKind = llvm::OpenMPIRBuilder::EvalKind::Scalar;

EvalKind = llvm::OpenMPIRBuilder::EvalKind::Complex;

EvalKind = llvm::OpenMPIRBuilder::EvalKind::Aggregate;

ReductionGen = [&](InsertPointTy CodeGenIP,

I,

llvm::Value **LHSPtr, llvm::Value **RHSPtr,

llvm::Function *NewFunc) {

CGF.

.restoreIP(CodeGenIP);

*CurFn = CGF.

;

CGF.

= NewFunc;

cast<VarDecl>(cast<DeclRefExpr>(LHSExprs[I])->getDecl()))

.emitRawPointer(CGF);

cast<VarDecl>(cast<DeclRefExpr>(RHSExprs[I])->getDecl()))

.emitRawPointer(CGF);

cast<DeclRefExpr>(LHSExprs[I]),

cast<DeclRefExpr>(RHSExprs[I]));

InsertPointTy(CGF.

.GetInsertBlock(),

CGF.

.GetInsertPoint());

ReductionInfos.emplace_back(llvm::OpenMPIRBuilder::ReductionInfo(

ElementType,

, PrivateVariable, EvalKind,

, ReductionGen, AtomicReductionGen));

llvm::OpenMPIRBuilder::InsertPointTy AfterIP =

OmpLoc, AllocaIP, CodeGenIP, ReductionInfos,

, TeamsReduction,

DistributeReduction, llvm::OpenMPIRBuilder::ReductionGenCBKind::Clang,

.getLangOpts().OpenMPCUDAReductionBufNum, RTLoc));

CGF.

.restoreIP(AfterIP);

*NativeParam)

{

*NonQualTy = QC.

(ArgType);

(

*

= FD->

<OMPCaptureKindAttr>()) {

(

->getCaptureKind() == OMPC_map) {

{ NVPTX_local_addr = 5 };

(isa<ImplicitParamDecl>(NativeParam))

,

,

);

*TargetParam)

{

assert(NativeParam != TargetParam &&

);

*NonQualTy = QC.

(NativeParamType);

NativePointeeAddrSpace =

llvm::PointerType::get(CGF.

(), NativePointeeAddrSpace));

NativeParamAddr;

TargetArgs.reserve(Args.size());

*FnType = OutlinedFn.getFunctionType();

(

I = 0,

= Args.size(); I <

; ++I) {

(FnType->isVarArg() && FnType->getNumParams() <= I) {

TargetArgs.append(std::next(Args.begin(), I), Args.end());

llvm::Type *TargetType = FnType->getParamType(I);

llvm::Value *NativeArg = Args[I];

(!TargetType->isPointerTy()) {

TargetArgs.emplace_back(NativeArg);

TargetArgs.emplace_back(

llvm::Function *CGOpenMPRuntimeGPU::createParallelDataSharingWrapper(

&CS = *

.getCapturedStmt(OMPD_parallel);

WrapperArgs.emplace_back(&ParallelLevelArg);

WrapperArgs.emplace_back(&WrapperArg);

*Fn = llvm::Function::Create(

Twine(OutlinedParallelFn->getName(),

), &

.

());

Fn->addFnAttr(llvm::Attribute::NoInline);

Fn->setLinkage(llvm::GlobalValue::InternalLinkage);

Fn->setDoesNotRecurse();

*RD = CS.getCapturedRecordDecl();

CurField = RD->field_begin();

CI = CS.capture_begin();

llvm::Value *GlobalArgsPtr = GlobalArgs.

();

llvm::Value *DataSharingArgs[] = {GlobalArgsPtr};

.

(), OMPRTL___kmpc_get_shared_variables),

(CS.capture_size() > 0 ||

Src, Bld.getPtrTy(0), CGF.

);

cast<OMPLoopDirective>(

).getLowerBoundVariable()->getExprLoc());

Args.emplace_back(LB);

cast<OMPLoopDirective>(

).getUpperBoundVariable()->getExprLoc());

Args.emplace_back(UB);

(CS.capture_size() > 0) {

(

I = 0,

= CS.capture_size(); I <

; ++I, ++CI, ++CurField) {

ElemTy = CurField->getType();

(CI->capturesVariableByCopy() &&

!CI->getCapturedVar()->getType()->isAnyPointerType()) {

Args.emplace_back(Arg);

assert(

&&

);

assert(FunctionGlobalizedDecls.count(CGF.

) == 0 &&

);

assert((!TeamAndReductions.first || TeamAndReductions.first ==

) &&

);

*Body =

;

NeedToDelayGlobalization =

;

(

*FD = dyn_cast<FunctionDecl>(

)) {

Body = FD->getBody();

}

(

*BD = dyn_cast<BlockDecl>(

)) {

Body = BD->getBody();

}

(

*CD = dyn_cast<CapturedDecl>(

)) {

Body = CD->getBody();

(NeedToDelayGlobalization &&

CheckVarsEscapingDeclContext VarChecker(CGF, TeamAndReductions.second);

VarChecker.Visit(Body);

VarChecker.getGlobalizedRecord(IsInTTDRegion);

TeamAndReductions.first =

;

TeamAndReductions.second.clear();

VarChecker.getEscapedVariableLengthDecls();

VarChecker.getDelayedVariableLengthDecls();

(!GlobalizedVarsRecord && EscapedVariableLengthDecls.empty() &&

DelayedVariableLengthDecls.empty())

I = FunctionGlobalizedDecls.try_emplace(CGF.

).first;

I->getSecond().MappedParams =

std::make_unique<CodeGenFunction::OMPMapVars>();

I->getSecond().EscapedParameters.insert(

VarChecker.getEscapedParameters().begin(),

VarChecker.getEscapedParameters().end());

I->getSecond().EscapedVariableLengthDecls.append(

EscapedVariableLengthDecls.begin(), EscapedVariableLengthDecls.end());

I->getSecond().DelayedVariableLengthDecls.append(

DelayedVariableLengthDecls.begin(), DelayedVariableLengthDecls.end());

DeclToAddrMapTy &

= I->getSecond().LocalVarData;

(

*VD : VarChecker.getEscapedDecls()) {

.insert(std::make_pair(VD, MappedVarData()));

(!NeedToDelayGlobalization) {

GlobalizationScope() =

;

.emitGenericVarsEpilog(CGF);

(VD && VD->

<OMPAllocateDeclAttr>()) {

*A = VD->

<OMPAllocateDeclAttr>();

(A->getAllocatorType()) {

OMPAllocateDeclAttr::OMPNullMemAlloc:

OMPAllocateDeclAttr::OMPDefaultMemAlloc:

OMPAllocateDeclAttr::OMPHighBWMemAlloc:

OMPAllocateDeclAttr::OMPLowLatMemAlloc:

OMPAllocateDeclAttr::OMPThreadMemAlloc:

OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:

OMPAllocateDeclAttr::OMPConstMemAlloc:

OMPAllocateDeclAttr::OMPPTeamMemAlloc:

OMPAllocateDeclAttr::OMPLargeCapMemAlloc:

OMPAllocateDeclAttr::OMPCGroupMemAlloc:

*GV =

llvm::GlobalVariable(

llvm::GlobalValue::InternalLinkage, llvm::PoisonValue::get(VarTy),

, llvm::GlobalValue::NotThreadLocal,

I = FunctionGlobalizedDecls.find(CGF.

);

(I == FunctionGlobalizedDecls.end())

VDI = I->getSecond().LocalVarData.find(VD);

(VDI != I->getSecond().LocalVarData.end())

VDI->second.PrivateAddr;

VDI = I->getSecond().LocalVarData.find(

cast<VarDecl>(cast<DeclRefExpr>(IT->getRef())->getDecl())

->getCanonicalDecl());

(VDI != I->getSecond().LocalVarData.end())

VDI->second.PrivateAddr;

FunctionGlobalizedDecls.erase(CGF.

);

llvm::Value *&Chunk)

{

ScheduleKind = OMPC_DIST_SCHEDULE_static;

RT.getGPUNumThreads(CGF),

S.getIterationVariable()->getType(), S.getBeginLoc());

CGF, S, ScheduleKind, Chunk);

*&ChunkExpr)

{

ScheduleKind = OMPC_SCHEDULE_static;

llvm::APInt ChunkSize(32, 1);

);

(!

.capturesVariable())

*VD =

.getCapturedVar();

*RD = VD->

()

(!RD || !RD->isLambda())

llvm::DenseMap<const ValueDecl *, FieldDecl *> Captures;

RD->getCaptureFields(Captures, ThisCapture);

*VD = LC.getCapturedVar();

It = Captures.find(VD);

assert(It != Captures.end() &&

);

(!VD || !VD->

<OMPAllocateDeclAttr>())

*A = VD->

<OMPAllocateDeclAttr>();

(A->getAllocatorType()) {

OMPAllocateDeclAttr::OMPNullMemAlloc:

OMPAllocateDeclAttr::OMPDefaultMemAlloc:

OMPAllocateDeclAttr::OMPLargeCapMemAlloc:

OMPAllocateDeclAttr::OMPCGroupMemAlloc:

OMPAllocateDeclAttr::OMPHighBWMemAlloc:

OMPAllocateDeclAttr::OMPLowLatMemAlloc:

OMPAllocateDeclAttr::OMPThreadMemAlloc:

OMPAllocateDeclAttr::OMPConstMemAlloc:

OMPAllocateDeclAttr::OMPPTeamMemAlloc:

OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:

llvm_unreachable(

);

(Feature.getValue()) {

(

*Clause :

->clauselists()) {

(Clause->getClauseKind() == OMPC_unified_shared_memory) {

llvm::raw_svector_ostream Out(Buffer);

<<

;

.

(Clause->getBeginLoc(), Out.str());

llvm_unreachable(

);

*LocSize =

;

llvm::Function *F = M->getFunction(LocSize);

F = llvm::Function::Create(llvm::FunctionType::get(CGF.

, {},

),

llvm::GlobalVariable::ExternalLinkage, LocSize,

Bld.CreateCall(F, {},

);

.

(), OMPRTL___kmpc_get_hardware_thread_id_in_block),

static void getTeamsReductionVars(ASTContext &Ctx, const OMPExecutableDirective &D, llvm::SmallVectorImpl< const ValueDecl * > &Vars)

Get list of reduction variables from the teams ... directives.

static llvm::Value * castValueToType(CodeGenFunction &CGF, llvm::Value *Val, QualType ValTy, QualType CastTy, SourceLocation Loc)

Cast value to the specified type.

static void getDistributeLastprivateVars(ASTContext &Ctx, const OMPExecutableDirective &D, llvm::SmallVectorImpl< const ValueDecl * > &Vars)

Get list of lastprivate variables from the teams distribute ... or teams {distribute ....

static bool hasNestedSPMDDirective(ASTContext &Ctx, const OMPExecutableDirective &D)

Check for inner (nested) SPMD construct, if any.

static bool supportsSPMDExecutionMode(ASTContext &Ctx, const OMPExecutableDirective &D)

static OffloadArch getOffloadArch(CodeGenModule &CGM)

This file defines OpenMP nodes for declarative directives.

This file defines OpenMP AST classes for clauses.

static std::pair< ValueDecl *, bool > getPrivateItem(Sema &S, Expr *&RefExpr, SourceLocation &ELoc, SourceRange &ERange, bool AllowArraySection=false, StringRef DiagType="")

This file defines OpenMP AST classes for executable directives and clauses.

Holds long-lived AST nodes (such as types and decls) that can be referred to throughout the semantic ...

QualType getPointerType(QualType T) const

Return the uniqued reference to the type for a pointer to the specified type.

QualType getUIntPtrType() const

Return a type compatible with "uintptr_t" (C99 7.18.1.4), as defined by the target.

QualType getIntTypeForBitwidth(unsigned DestWidth, unsigned Signed) const

getIntTypeForBitwidth - sets integer QualTy according to specified details: bitwidth,...

CanQualType getSizeType() const

Return the unique type for "size_t" (C99 7.17), defined in <stddef.h>.

CharUnits getDeclAlign(const Decl *D, bool ForAlignof=false) const

Return a conservative estimate of the alignment of the specified decl D.

CharUnits getTypeSizeInChars(QualType T) const

Return the size of the specified (complete) type T, in characters.

const TargetInfo & getTargetInfo() const

QualType getAddrSpaceQualType(QualType T, LangAS AddressSpace) const

Return the uniqued reference to the type for an address space qualified type with the specified type ...

unsigned getTargetAddressSpace(LangAS AS) const

Attr - This represents one attribute.

A class which contains all the information about a particular captured value.

BlockExpr - Adaptor class for mixing a BlockDecl with expressions.

CallExpr - Represents a function call (C99 6.5.2.2, C++ [expr.call]).

Describes the capture of either a variable, or 'this', or variable-length array type.

This captures a statement into a function.

bool capturesVariable(const VarDecl *Var) const

True if this variable has been captured.

CharUnits - This is an opaque type for sizes expressed in character units.

bool isZero() const

isZero - Test whether the quantity equals zero.

llvm::Align getAsAlign() const

getAsAlign - Returns Quantity as a valid llvm::Align, Beware llvm::Align assumes power of two 8-bit b...

QuantityType getQuantity() const

getQuantity - Get the raw integer representation of this quantity.

Like RawAddress, an abstract representation of an aligned address, but the pointer contained in this ...

llvm::Value * emitRawPointer(CodeGenFunction &CGF) const

Return the pointer contained in this class after authenticating it and adding offset to it if necessa...

Address withElementType(llvm::Type *ElemTy) const

Return address with different element type, but same pointer and alignment.

llvm::StoreInst * CreateStore(llvm::Value *Val, Address Addr, bool IsVolatile=false)

Address CreatePointerBitCastOrAddrSpaceCast(Address Addr, llvm::Type *Ty, llvm::Type *ElementTy, const llvm::Twine &Name="")

Address CreateConstArrayGEP(Address Addr, uint64_t Index, const llvm::Twine &Name="")

Given addr = [n x T]* ... produce name = getelementptr inbounds addr, i64 0, i64 index where i64 is a...

Address CreateConstInBoundsGEP(Address Addr, uint64_t Index, const llvm::Twine &Name="")

Given addr = T* ... produce name = getelementptr inbounds addr, i64 index where i64 is actually the t...

CGFunctionInfo - Class to encapsulate the information about a function definition.

void emitParallelCall(CodeGenFunction &CGF, SourceLocation Loc, llvm::Function *OutlinedFn, ArrayRef< llvm::Value * > CapturedVars, const Expr *IfCond, llvm::Value *NumThreads) override

Emits code for parallel or serial call of the OutlinedFn with variables captured in a record which ad...

llvm::Function * emitTeamsOutlinedFunction(CodeGenFunction &CGF, const OMPExecutableDirective &D, const VarDecl *ThreadIDVar, OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) override

Emits inlined function for the specified OpenMP teams.

void emitProcBindClause(CodeGenFunction &CGF, llvm::omp::ProcBindKind ProcBind, SourceLocation Loc) override

Emit call to void __kmpc_push_proc_bind(ident_t *loc, kmp_int32 global_tid, int proc_bind) to generat...

void emitReduction(CodeGenFunction &CGF, SourceLocation Loc, ArrayRef< const Expr * > Privates, ArrayRef< const Expr * > LHSExprs, ArrayRef< const Expr * > RHSExprs, ArrayRef< const Expr * > ReductionOps, ReductionOptionsTy Options) override

Emit a code for reduction clause.

DataSharingMode

Target codegen is specialized based on two data-sharing modes: CUDA, in which the local variables are...

@ DS_CUDA

CUDA data sharing mode.

@ DS_Generic

Generic data-sharing mode.

void getDefaultDistScheduleAndChunk(CodeGenFunction &CGF, const OMPLoopDirective &S, OpenMPDistScheduleClauseKind &ScheduleKind, llvm::Value *&Chunk) const override

Choose a default value for the dist_schedule clause.

Address getAddressOfLocalVariable(CodeGenFunction &CGF, const VarDecl *VD) override

Gets the OpenMP-specific address of the local variable.

void emitFunctionProlog(CodeGenFunction &CGF, const Decl *D) override

Emits OpenMP-specific function prolog.

void getDefaultScheduleAndChunk(CodeGenFunction &CGF, const OMPLoopDirective &S, OpenMPScheduleClauseKind &ScheduleKind, const Expr *&ChunkExpr) const override

Choose a default value for the schedule clause.

void emitNumTeamsClause(CodeGenFunction &CGF, const Expr *NumTeams, const Expr *ThreadLimit, SourceLocation Loc) override

This function ought to emit, in the general case, a call to.

void emitCriticalRegion(CodeGenFunction &CGF, StringRef CriticalName, const RegionCodeGenTy &CriticalOpGen, SourceLocation Loc, const Expr *Hint=nullptr) override

Emits a critical region.

void emitTeamsCall(CodeGenFunction &CGF, const OMPExecutableDirective &D, SourceLocation Loc, llvm::Function *OutlinedFn, ArrayRef< llvm::Value * > CapturedVars) override

Emits code for teams call of the OutlinedFn with variables captured in a record which address is stor...

bool hasAllocateAttributeForGlobalVar(const VarDecl *VD, LangAS &AS) override

Checks if the variable has associated OMPAllocateDeclAttr attribute with the predefined allocator and...

void getKmpcFreeShared(CodeGenFunction &CGF, const std::pair< llvm::Value *, llvm::Value * > &AddrSizePair) override

Get call to __kmpc_free_shared.

CGOpenMPRuntimeGPU(CodeGenModule &CGM)

llvm::Function * emitParallelOutlinedFunction(CodeGenFunction &CGF, const OMPExecutableDirective &D, const VarDecl *ThreadIDVar, OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) override

Emits inlined function for the specified OpenMP parallel.

void functionFinished(CodeGenFunction &CGF) override

Cleans up references to the objects in finished function.

llvm::Value * getGPUThreadID(CodeGenFunction &CGF)

Get the id of the current thread on the GPU.

void processRequiresDirective(const OMPRequiresDecl *D) override

Perform check on requires decl to ensure that target architecture supports unified addressing.

bool isDelayedVariableLengthDecl(CodeGenFunction &CGF, const VarDecl *VD) const override

Declare generalized virtual functions which need to be defined by all specializations of OpenMPGPURun...

void emitOutlinedFunctionCall(CodeGenFunction &CGF, SourceLocation Loc, llvm::FunctionCallee OutlinedFn, ArrayRef< llvm::Value * > Args={}) const override

Emits call of the outlined function with the provided arguments, translating these arguments to corre...

Address getParameterAddress(CodeGenFunction &CGF, const VarDecl *NativeParam, const VarDecl *TargetParam) const override

Gets the address of the native argument basing on the address of the target-specific parameter.

ExecutionMode

Defines the execution mode.

@ EM_NonSPMD

Non-SPMD execution mode (1 master thread, others are workers).

@ EM_Unknown

Unknown execution mode (orphaned directive).

@ EM_SPMD

SPMD execution mode (all threads are worker threads).

void emitBarrierCall(CodeGenFunction &CGF, SourceLocation Loc, OpenMPDirectiveKind Kind, bool EmitChecks=true, bool ForceSimpleCall=false) override

Emit an implicit/explicit barrier for OpenMP threads.

llvm::Value * getGPUNumThreads(CodeGenFunction &CGF)

Get the maximum number of threads in a block of the GPU.

const VarDecl * translateParameter(const FieldDecl *FD, const VarDecl *NativeParam) const override

Translates the native parameter of outlined function if this is required for target.

std::pair< llvm::Value *, llvm::Value * > getKmpcAllocShared(CodeGenFunction &CGF, const VarDecl *VD) override

Get call to __kmpc_alloc_shared.

bool isGPU() const override

Returns true if the current target is a GPU.

void emitNumThreadsClause(CodeGenFunction &CGF, llvm::Value *NumThreads, SourceLocation Loc) override

Emits call to void __kmpc_push_num_threads(ident_t *loc, kmp_int32 global_tid, kmp_int32 num_threads)...

void adjustTargetSpecificDataForLambdas(CodeGenFunction &CGF, const OMPExecutableDirective &D) const override

Adjust some parameters for the target-based directives, like addresses of the variables captured by r...

virtual Address emitThreadIDAddress(CodeGenFunction &CGF, SourceLocation Loc)

Emits address of the word in a memory where current thread id is stored.

static const Stmt * getSingleCompoundChild(ASTContext &Ctx, const Stmt *Body)

Checks if the Body is the CompoundStmt and returns its child statement iff there is only one that is ...

llvm::Value * emitUpdateLocation(CodeGenFunction &CGF, SourceLocation Loc, unsigned Flags=0, bool EmitLoc=false)

Emits object of ident_t type with info for source location.

virtual void functionFinished(CodeGenFunction &CGF)

Cleans up references to the objects in finished function.

virtual llvm::Function * emitTeamsOutlinedFunction(CodeGenFunction &CGF, const OMPExecutableDirective &D, const VarDecl *ThreadIDVar, OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen)

Emits outlined function for the specified OpenMP teams directive D.

llvm::OpenMPIRBuilder OMPBuilder

An OpenMP-IR-Builder instance.

virtual void emitTargetOutlinedFunctionHelper(const OMPExecutableDirective &D, StringRef ParentName, llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID, bool IsOffloadEntry, const RegionCodeGenTy &CodeGen)

Helper to emit outlined function for 'target' directive.

bool hasRequiresUnifiedSharedMemory() const

Return whether the unified_shared_memory has been specified.

virtual void processRequiresDirective(const OMPRequiresDecl *D)

Perform check on requires decl to ensure that target architecture supports unified addressing.

llvm::Value * getThreadID(CodeGenFunction &CGF, SourceLocation Loc)

Gets thread id value for the current thread.

void clearLocThreadIdInsertPt(CodeGenFunction &CGF)

void computeMinAndMaxThreadsAndTeams(const OMPExecutableDirective &D, CodeGenFunction &CGF, llvm::OpenMPIRBuilder::TargetKernelDefaultAttrs &Attrs)

Helper to determine the min/max number of threads/teams for D.

static unsigned getDefaultFlagsForBarriers(OpenMPDirectiveKind Kind)

Returns default flags for the barriers depending on the directive, for which this barier is going to ...

virtual llvm::Function * emitParallelOutlinedFunction(CodeGenFunction &CGF, const OMPExecutableDirective &D, const VarDecl *ThreadIDVar, OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen)

Emits outlined function for the specified OpenMP parallel directive D.

virtual void getDefaultDistScheduleAndChunk(CodeGenFunction &CGF, const OMPLoopDirective &S, OpenMPDistScheduleClauseKind &ScheduleKind, llvm::Value *&Chunk) const

Choose default schedule type and chunk value for the dist_schedule clause.

llvm::Type * getIdentTyPointerTy()

Returns pointer to ident_t type.

void emitSingleReductionCombiner(CodeGenFunction &CGF, const Expr *ReductionOp, const Expr *PrivateRef, const DeclRefExpr *LHS, const DeclRefExpr *RHS)

Emits single reduction combiner.

llvm::OpenMPIRBuilder & getOMPBuilder()

virtual void emitCriticalRegion(CodeGenFunction &CGF, StringRef CriticalName, const RegionCodeGenTy &CriticalOpGen, SourceLocation Loc, const Expr *Hint=nullptr)

Emits a critical region.

virtual void emitOutlinedFunctionCall(CodeGenFunction &CGF, SourceLocation Loc, llvm::FunctionCallee OutlinedFn, ArrayRef< llvm::Value * > Args={}) const

Emits call of the outlined function with the provided arguments, translating these arguments to corre...

virtual void emitReduction(CodeGenFunction &CGF, SourceLocation Loc, ArrayRef< const Expr * > Privates, ArrayRef< const Expr * > LHSExprs, ArrayRef< const Expr * > RHSExprs, ArrayRef< const Expr * > ReductionOps, ReductionOptionsTy Options)

Emit a code for reduction clause.

CapturedRegionKind getKind() const

bool isCXXThisExprCaptured() const

The scope used to remap some variables as private in the OpenMP loop body (or other captured region e...

CodeGenFunction - This class organizes the per-function state that is used while generating LLVM code...

void FinishFunction(SourceLocation EndLoc=SourceLocation())

FinishFunction - Complete IR generation of the current function.

static TypeEvaluationKind getEvaluationKind(QualType T)

getEvaluationKind - Return the TypeEvaluationKind of QualType T.

CGCapturedStmtInfo * CapturedStmtInfo

LValue MakeNaturalAlignPointeeRawAddrLValue(llvm::Value *V, QualType T)

Same as MakeNaturalAlignPointeeAddrLValue except that the pointer is known to be unsigned.

Address EmitLoadOfPointer(Address Ptr, const PointerType *PtrTy, LValueBaseInfo *BaseInfo=nullptr, TBAAAccessInfo *TBAAInfo=nullptr)

Load a pointer with type PtrTy stored at address Ptr.

RawAddress CreateDefaultAlignTempAlloca(llvm::Type *Ty, const Twine &Name="tmp")

CreateDefaultAlignedTempAlloca - This creates an alloca with the default ABI alignment of the given L...

llvm::BasicBlock * createBasicBlock(const Twine &name="", llvm::Function *parent=nullptr, llvm::BasicBlock *before=nullptr)

createBasicBlock - Create an LLVM basic block.

LValue EmitLValueForFieldInitialization(LValue Base, const FieldDecl *Field)

EmitLValueForFieldInitialization - Like EmitLValueForField, except that if the Field is a reference,...

void EmitBlock(llvm::BasicBlock *BB, bool IsFinished=false)

EmitBlock - Emit the given block.

llvm::Type * ConvertTypeForMem(QualType T)

llvm::AssertingVH< llvm::Instruction > AllocaInsertPt

AllocaInsertPoint - This is an instruction in the entry block before which we prefer to insert alloca...

RawAddress CreateMemTemp(QualType T, const Twine &Name="tmp", RawAddress *Alloca=nullptr)

CreateMemTemp - Create a temporary memory object of the given type, with appropriate alignmen and cas...

const TargetInfo & getTarget() const

llvm::DebugLoc SourceLocToDebugLoc(SourceLocation Location)

Converts Location to a DebugLoc, if debug information is enabled.

llvm::Value * getTypeSize(QualType Ty)

Returns calculated size of the specified type.

void StartFunction(GlobalDecl GD, QualType RetTy, llvm::Function *Fn, const CGFunctionInfo &FnInfo, const FunctionArgList &Args, SourceLocation Loc=SourceLocation(), SourceLocation StartLoc=SourceLocation())

Emit code for the start of a function.

bool HaveInsertPoint() const

HaveInsertPoint - True if an insertion point is defined.

CGDebugInfo * getDebugInfo()

void EmitBranch(llvm::BasicBlock *Block)

EmitBranch - Emit a branch to the specified basic block from the current insert block,...

ASTContext & getContext() const

llvm::Value * EmitLoadOfScalar(Address Addr, bool Volatile, QualType Ty, SourceLocation Loc, AlignmentSource Source=AlignmentSource::Type, bool isNontemporal=false)

EmitLoadOfScalar - Load a scalar value from an address, taking care to appropriately convert from the...

llvm::CallInst * EmitRuntimeCall(llvm::FunctionCallee callee, const Twine &name="")

CodeGenTypes & getTypes() const

llvm::Value * EvaluateExprAsBool(const Expr *E)

EvaluateExprAsBool - Perform the usual unary conversions on the specified expression and compare the ...

LValue MakeAddrLValue(Address Addr, QualType T, AlignmentSource Source=AlignmentSource::Type)

llvm::Value * LoadCXXThis()

LoadCXXThis - Load the value of 'this'.

LValue EmitLoadOfReferenceLValue(LValue RefLVal)

Address GetAddrOfLocalVar(const VarDecl *VD)

GetAddrOfLocalVar - Return the address of a local variable.

llvm::Value * EmitScalarConversion(llvm::Value *Src, QualType SrcTy, QualType DstTy, SourceLocation Loc)

Emit a conversion from the specified type to the specified destination type, both of which are LLVM s...

llvm::LLVMContext & getLLVMContext()

void EmitStoreOfScalar(llvm::Value *Value, Address Addr, bool Volatile, QualType Ty, AlignmentSource Source=AlignmentSource::Type, bool isInit=false, bool isNontemporal=false)

EmitStoreOfScalar - Store a scalar value to an address, taking care to appropriately convert from the...

This class organizes the cross-function state that is used while generating LLVM code.

void SetInternalFunctionAttributes(GlobalDecl GD, llvm::Function *F, const CGFunctionInfo &FI)

Set the attributes on the LLVM function for the given decl and function info.

llvm::Module & getModule() const

const LangOptions & getLangOpts() const

CodeGenTypes & getTypes()

const TargetInfo & getTarget() const

void Error(SourceLocation loc, StringRef error)

Emit a general error that something can't be done.

CGOpenMPRuntime & getOpenMPRuntime()

Return a reference to the configured OpenMP runtime.

ASTContext & getContext() const

llvm::LLVMContext & getLLVMContext()

llvm::FunctionType * GetFunctionType(const CGFunctionInfo &Info)

GetFunctionType - Get the LLVM function type for.

const CGFunctionInfo & arrangeBuiltinFunctionDeclaration(QualType resultType, const FunctionArgList &args)

A builtin function is a freestanding function using the default C conventions.

unsigned getTargetAddressSpace(QualType T) const

llvm::Type * ConvertTypeForMem(QualType T)

ConvertTypeForMem - Convert type T into a llvm::Type.

Information for lazily generating a cleanup.

FunctionArgList - Type for representing both the decl and type of parameters to a function.

LValue - This represents an lvalue references.

Address getAddress() const

A basic class for pre|post-action for advanced codegen sequence for OpenMP region.

An abstract representation of an aligned address.

llvm::Value * getPointer() const

Class provides a way to call simple version of codegen for OpenMP region, or an advanced with possibl...

void setAction(PrePostActionTy &Action) const

ConstStmtVisitor - This class implements a simple visitor for Stmt subclasses.

DeclContext - This is used only as base class of specific decl types that can act as declaration cont...

void addDecl(Decl *D)

Add the declaration D into this context.

A reference to a declared variable, function, enum, etc.

DeclStmt - Adaptor class for mixing declarations with statements and expressions.

Decl - This represents one declaration (or definition), e.g.

attr_iterator attr_end() const

bool isCanonicalDecl() const

Whether this particular Decl is a canonical one.

attr_iterator attr_begin() const

SourceLocation getLocation() const

DeclContext * getDeclContext()

SourceLocation getBeginLoc() const LLVM_READONLY

virtual Decl * getCanonicalDecl()

Retrieves the "canonical" declaration of the given declaration.

SourceLocation getBeginLoc() const LLVM_READONLY

This represents one expression.

Expr * IgnoreParenImpCasts() LLVM_READONLY

Skip past any parentheses and implicit casts which might surround this expression until reaching a fi...

Expr * IgnoreParens() LLVM_READONLY

Skip past any parentheses which might surround this expression until reaching a fixed point.

bool isLValue() const

isLValue - True if this expression is an "l-value" according to the rules of the current language.

Represents a member of a struct/union/class.

static FieldDecl * Create(const ASTContext &C, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, const IdentifierInfo *Id, QualType T, TypeSourceInfo *TInfo, Expr *BW, bool Mutable, InClassInitStyle InitStyle)

GlobalDecl - represents a global declaration.

ImplicitCastExpr - Allows us to explicitly represent implicit type conversions, which have no direct ...

static ImplicitParamDecl * Create(ASTContext &C, DeclContext *DC, SourceLocation IdLoc, IdentifierInfo *Id, QualType T, ImplicitParamKind ParamKind)

Create implicit parameter.

static IntegerLiteral * Create(const ASTContext &C, const llvm::APInt &V, QualType type, SourceLocation l)

Returns a new integer literal with value 'V' and type 'type'.

Describes the capture of a variable or of this, or of a C++1y init-capture.

A C++ lambda expression, which produces a function object (of unspecified type) that can be invoked l...

std::string OMPHostIRFile

Name of the IR file that contains the result of the OpenMP target host code generation.

IdentifierInfo * getIdentifier() const

Get the identifier that names this declaration, if there is one.

StringRef getName() const

Get the name of identifier for this declaration as a StringRef.

This is a basic class for representing single OpenMP clause.

This is a basic class for representing single OpenMP executable directive.

OpenMPDirectiveKind getDirectiveKind() const

static llvm::iterator_range< specific_clause_iterator< SpecificClause > > getClausesOfKind(ArrayRef< OMPClause * > Clauses)

This represents clause 'lastprivate' in the '#pragma omp ...' directives.

This is a common base class for loop directives ('omp simd', 'omp for', 'omp for simd' etc....

This represents clause 'reduction' in the '#pragma omp ...' directives.

This represents '#pragma omp requires...' directive.

This represents 'ompx_bare' clause in the '#pragma omp target teams ...' directive.

static ParmVarDecl * Create(ASTContext &C, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, const IdentifierInfo *Id, QualType T, TypeSourceInfo *TInfo, StorageClass S, Expr *DefArg)

PointerType - C99 6.7.5.1 - Pointer Declarators.

A (possibly-)qualified type.

LangAS getAddressSpace() const

Return the address space of this type.

QualType getNonReferenceType() const

If Type is a reference type (e.g., const int&), returns the type that the reference refers to ("const...

QualType getCanonicalType() const

A qualifier set is used to build a set of qualifiers.

const Type * strip(QualType type)

Collect any qualifiers on the given type and return an unqualified type.

QualType apply(const ASTContext &Context, QualType QT) const

Apply the collected qualifiers to the given type.

void addAddressSpace(LangAS space)

Represents a struct/union/class.

virtual void completeDefinition()

Note that the definition of this type is now complete.

Scope - A scope is a transient data structure that is used while parsing the program.

Encodes a location in the source.

RetTy Visit(PTR(Stmt) S, ParamTys... P)

Stmt - This represents one statement.

void startDefinition()

Starts the definition of this tag declaration.

unsigned getNewAlign() const

Return the largest alignment for which a suitably-sized allocation with '::operator new(size_t)' is g...

TargetOptions & getTargetOpts() const

Retrieve the target options.

virtual const llvm::omp::GV & getGridValue() const

virtual bool hasFeature(StringRef Feature) const

Determine whether the given target has the given feature.

llvm::StringMap< bool > FeatureMap

The map of which features have been enabled disabled based on the command line.

The base class of the type hierarchy.

CXXRecordDecl * getAsCXXRecordDecl() const

Retrieves the CXXRecordDecl that this type refers to, either because the type is a RecordType or beca...

bool isIntegerType() const

isIntegerType() does not include complex integers (a GCC extension).

bool isReferenceType() const

QualType getPointeeType() const

If this is a pointer, ObjC object pointer, or block pointer, this returns the respective pointee.

bool isLValueReferenceType() const

bool hasSignedIntegerRepresentation() const

Determine whether this type has an signed integer representation of some sort, e.g....

bool isVariablyModifiedType() const

Whether this type is a variably-modified type (C99 6.7.5).

UnaryOperator - This represents the unary-expression's (except sizeof and alignof),...

Represent the declaration of a variable (in which case it is an lvalue) a function (in which case it ...

bool isInitCapture() const

Whether this variable is the implicit variable for a lambda init-capture.

Represents a variable declaration or definition.

VarDecl * getCanonicalDecl() override

Retrieves the "canonical" declaration of the given declaration.

bool isInitCapture() const

Whether this variable is the implicit variable for a lambda init-capture.

specific_attr_iterator - Iterates over a subrange of an AttrVec, only providing attributes that are o...

@ Type

The l-value was considered opaque, so the alignment was determined from a type.

@ Decl

The l-value was an access to a declared entity or something equivalently strong, like the address of ...

The JSON file list parser is used to communicate input to InstallAPI.

llvm::omp::Directive OpenMPDirectiveKind

OpenMP directives.

@ ICIS_NoInit

No in-class initializer.

bool isOpenMPDistributeDirective(OpenMPDirectiveKind DKind)

Checks if the specified directive is a distribute directive.

@ LCK_ByRef

Capturing by reference.

@ Private

'private' clause, allowed on 'parallel', 'serial', 'loop', 'parallel loop', and 'serial loop' constru...

bool isOpenMPParallelDirective(OpenMPDirectiveKind DKind)

Checks if the specified directive is a parallel-kind directive.

bool isOpenMPPrivate(OpenMPClauseKind Kind)

Checks if the specified clause is one of private clauses like 'private', 'firstprivate',...

OpenMPDistScheduleClauseKind

OpenMP attributes for 'dist_schedule' clause.

bool isOpenMPTargetExecutionDirective(OpenMPDirectiveKind DKind)

Checks if the specified directive is a target code offload directive.

bool isOpenMPTeamsDirective(OpenMPDirectiveKind DKind)

Checks if the specified directive is a teams-kind directive.

OffloadArch StringToOffloadArch(llvm::StringRef S)

bool isOpenMPLoopBoundSharingDirective(OpenMPDirectiveKind Kind)

Checks if the specified directive kind is one of the composite or combined directives that need loop ...

LangAS

Defines the address space values used by the address space qualifier of QualType.

const char * OffloadArchToString(OffloadArch A)

void getOpenMPCaptureRegions(llvm::SmallVectorImpl< OpenMPDirectiveKind > &CaptureRegions, OpenMPDirectiveKind DKind)

Return the captured regions of an OpenMP directive.

LangAS getLangASFromTargetAS(unsigned TargetAS)

@ CXXThis

Parameter for C++ 'this' argument.

@ Other

Other implicit parameter.

OpenMPScheduleClauseKind

OpenMP attributes for 'schedule' clause.

llvm::PointerType * VoidPtrTy

llvm::IntegerType * SizeTy

llvm::PointerType * VoidPtrPtrTy

llvm::IntegerType * Int32Ty

llvm::PointerType * Int8PtrTy

RetroSearch is an open source project built by @garambo | Open a GitHub Issue

Search and Browse the WWW like it's 1997 | Search results from DuckDuckGo

HTML: 3.2 | Encoding: UTF-8 | Version: 0.7.4