diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h --- a/clang/lib/CodeGen/CGOpenMPRuntime.h +++ b/clang/lib/CodeGen/CGOpenMPRuntime.h @@ -1012,6 +1012,14 @@ const RegionCodeGenTy &MasterOpGen, SourceLocation Loc); + /// Emits a masked region. + /// \param MaskedOpGen Generator for the statement associated with the given + /// masked region. + virtual void emitMaskedRegion(CodeGenFunction &CGF, + const RegionCodeGenTy &MaskedOpGen, + SourceLocation Loc, + const Expr *Filter = nullptr); + /// Emits code for a taskyield directive. virtual void emitTaskyieldCall(CodeGenFunction &CGF, SourceLocation Loc); @@ -1984,6 +1992,17 @@ const RegionCodeGenTy &MasterOpGen, SourceLocation Loc) override; + /// Emits a masked region. + /// \param MaskedOpGen Generator for the statement associated with the given + /// masked region. + void emitMaskedRegion(CodeGenFunction &CGF, + const RegionCodeGenTy &MaskedOpGen, SourceLocation Loc, + const Expr *Filter = nullptr) override; + + /// Emits a masked region. + /// \param MaskedOpGen Generator for the statement associated with the given + /// masked region. + /// Emits code for a taskyield directive. void emitTaskyieldCall(CodeGenFunction &CGF, SourceLocation Loc) override; diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -2279,6 +2279,35 @@ Action.Done(CGF); } +void CGOpenMPRuntime::emitMaskedRegion(CodeGenFunction &CGF, + const RegionCodeGenTy &MaskedOpGen, + SourceLocation Loc, const Expr *Filter) { + if (!CGF.HaveInsertPoint()) + return; + // if(__kmpc_masked(ident_t *, gtid, filter)) { + // MaskedOpGen(); + // __kmpc_end_masked(iden_t *, gtid); + // } + // Prepare arguments and build a call to __kmpc_masked + llvm::Value *FilterVal = Filter + ? CGF.EmitScalarExpr(Filter, CGF.Int32Ty) + : llvm::ConstantInt::get(CGM.Int32Ty, /*V=*/0); + llvm::Value *Args[] = {emitUpdateLocation(CGF, Loc), getThreadID(CGF, Loc), + FilterVal}; + llvm::Value *ArgsEnd[] = {emitUpdateLocation(CGF, Loc), + getThreadID(CGF, Loc)}; + CommonActionTy Action(OMPBuilder.getOrCreateRuntimeFunction( + CGM.getModule(), OMPRTL___kmpc_masked), + Args, + OMPBuilder.getOrCreateRuntimeFunction( + CGM.getModule(), OMPRTL___kmpc_end_masked), + ArgsEnd, + /*Conditional=*/true); + MaskedOpGen.setAction(Action); + emitInlinedDirective(CGF, OMPD_masked, MaskedOpGen); + Action.Done(CGF); +} + void CGOpenMPRuntime::emitTaskyieldCall(CodeGenFunction &CGF, SourceLocation Loc) { if (!CGF.HaveInsertPoint()) @@ -6232,7 +6261,8 @@ return; InlinedOpenMPRegionRAII Region(CGF, CodeGen, InnerKind, HasCancel, InnerKind != OMPD_critical && - InnerKind != OMPD_master); + InnerKind != OMPD_master && + InnerKind != OMPD_masked); CGF.CapturedStmtInfo->EmitBody(CGF, /*S=*/nullptr); } @@ -12596,6 +12626,13 @@ llvm_unreachable("Not supported in SIMD-only mode"); } +void CGOpenMPSIMDRuntime::emitMaskedRegion(CodeGenFunction &CGF, + const RegionCodeGenTy &MasterOpGen, + SourceLocation Loc, + const Expr *Filter) { + llvm_unreachable("Not supported in SIMD-only mode"); +} + void CGOpenMPSIMDRuntime::emitTaskyieldCall(CodeGenFunction &CGF, SourceLocation Loc) { llvm_unreachable("Not supported in SIMD-only mode"); diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp --- a/clang/lib/CodeGen/CGStmt.cpp +++ b/clang/lib/CodeGen/CGStmt.cpp @@ -382,7 +382,7 @@ llvm_unreachable("Dispatch directive not supported yet."); break; case Stmt::OMPMaskedDirectiveClass: - llvm_unreachable("Masked directive not supported yet."); + EmitOMPMaskedDirective(cast(*S)); break; } } diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -3844,6 +3844,56 @@ emitMaster(*this, S); } +static void emitMasked(CodeGenFunction &CGF, const OMPExecutableDirective &S) { + auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) { + Action.Enter(CGF); + CGF.EmitStmt(S.getRawStmt()); + }; + Expr *Filter = nullptr; + if (const auto *FilterClause = S.getSingleClause()) { + Filter = FilterClause->getThreadID(); + } + CGF.CGM.getOpenMPRuntime().emitMaskedRegion(CGF, CodeGen, S.getBeginLoc(), + Filter); +} + +void CodeGenFunction::EmitOMPMaskedDirective(const OMPMaskedDirective &S) { + if (CGM.getLangOpts().OpenMPIRBuilder) { + llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder(); + using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy; + + const Stmt *MaskedRegionBodyStmt = S.getAssociatedStmt(); + const Expr *Filter = nullptr; + if (const auto *FilterClause = S.getSingleClause()) + Filter = FilterClause->getThreadID(); + llvm::Value *FilterVal = Filter + ? EmitScalarExpr(Filter, CGM.Int32Ty) + : llvm::ConstantInt::get(CGM.Int32Ty, /*V=*/0); + + auto FiniCB = [this](InsertPointTy IP) { + OMPBuilderCBHelpers::FinalizeOMPRegion(*this, IP); + }; + + auto BodyGenCB = [MaskedRegionBodyStmt, this](InsertPointTy AllocaIP, + InsertPointTy CodeGenIP, + llvm::BasicBlock &FiniBB) { + OMPBuilderCBHelpers::InlinedRegionBodyRAII IRB(*this, AllocaIP, FiniBB); + OMPBuilderCBHelpers::EmitOMPRegionBody(*this, MaskedRegionBodyStmt, + CodeGenIP, FiniBB); + }; + + LexicalScope Scope(*this, S.getSourceRange()); + EmitStopPoint(&S); + Builder.restoreIP( + OMPBuilder.createMasked(Builder, BodyGenCB, FiniCB, FilterVal)); + + return; + } + LexicalScope Scope(*this, S.getSourceRange()); + EmitStopPoint(&S); + emitMasked(*this, S); +} + void CodeGenFunction::EmitOMPCriticalDirective(const OMPCriticalDirective &S) { if (CGM.getLangOpts().OpenMPIRBuilder) { llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder(); @@ -6930,7 +6980,8 @@ if (D.getDirectiveKind() == OMPD_atomic || D.getDirectiveKind() == OMPD_critical || D.getDirectiveKind() == OMPD_section || - D.getDirectiveKind() == OMPD_master) { + D.getDirectiveKind() == OMPD_master || + D.getDirectiveKind() == OMPD_masked) { EmitStmt(D.getAssociatedStmt()); } else { auto LPCRegion = diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h --- a/clang/lib/CodeGen/CodeGenFunction.h +++ b/clang/lib/CodeGen/CodeGenFunction.h @@ -3421,6 +3421,7 @@ void EmitOMPSectionDirective(const OMPSectionDirective &S); void EmitOMPSingleDirective(const OMPSingleDirective &S); void EmitOMPMasterDirective(const OMPMasterDirective &S); + void EmitOMPMaskedDirective(const OMPMaskedDirective &S); void EmitOMPCriticalDirective(const OMPCriticalDirective &S); void EmitOMPParallelForDirective(const OMPParallelForDirective &S); void EmitOMPParallelForSimdDirective(const OMPParallelForSimdDirective &S); diff --git a/clang/test/OpenMP/masked_codegen.cpp b/clang/test/OpenMP/masked_codegen.cpp new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/masked_codegen.cpp @@ -0,0 +1,102 @@ +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=51 -x c++ -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s --check-prefixes=ALL,NORMAL +// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -x c++ -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefixes=ALL,NORMAL +// RUN: %clang_cc1 -verify -triple x86_64-apple-darwin10 -fopenmp -fopenmp-version=51 -fexceptions -fcxx-exceptions -debug-info-kind=line-tables-only -x c++ -emit-llvm %s -o - | FileCheck %s --check-prefix=TERM_DEBUG +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=51 -fopenmp-enable-irbuilder -x c++ -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s --check-prefixes=ALL,IRBUILDER +// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-enable-irbuilder -x c++ -std=c++11 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-enable-irbuilder -x c++ -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefixes=ALL,IRBUILDER + +// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -verify -triple x86_64-apple-darwin10 -fopenmp-simd -fopenmp-version=51 -fexceptions -fcxx-exceptions -debug-info-kind=line-tables-only -x c++ -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// SIMD-ONLY0-NOT: {{__kmpc|__tgt}} +// expected-no-diagnostics +#ifndef HEADER +#define HEADER + +// ALL: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* } + +// ALL: define {{.*}}void [[FOO:@.+]]() + +void foo() { extern void mayThrow(); mayThrow(); } + +// ALL-LABEL: @main +// TERM_DEBUG-LABEL: @main +int main() { + // ALL: [[A_ADDR:%.+]] = alloca i8 + char a; + +// ALL: [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEFAULT_LOC:@.+]]) +// ALL: [[RES:%.+]] = call {{.*}}i32 @__kmpc_masked([[IDENT_T_TY]]* [[DEFAULT_LOC]], i32 [[GTID]], i32 0) +// ALL-NEXT: [[IS_MASKED:%.+]] = icmp ne i32 [[RES]], 0 +// ALL-NEXT: br i1 [[IS_MASKED]], label {{%?}}[[THEN:.+]], label {{%?}}[[EXIT:.+]] +// ALL: [[THEN]] +// ALL-NEXT: store i8 2, i8* [[A_ADDR]] +// ALL-NEXT: call {{.*}}void @__kmpc_end_masked([[IDENT_T_TY]]* [[DEFAULT_LOC]], i32 [[GTID]]) +// ALL-NEXT: br label {{%?}}[[EXIT]] +// ALL: [[EXIT]] +#pragma omp masked + a = 2; +// IRBUILDER: [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEFAULT_LOC:@.+]]) +// ALL: [[RES:%.+]] = call {{.*}}i32 @__kmpc_masked([[IDENT_T_TY]]* [[DEFAULT_LOC]], i32 [[GTID]], i32 0) +// ALL-NEXT: [[IS_MASKED:%.+]] = icmp ne i32 [[RES]], 0 +// ALL-NEXT: br i1 [[IS_MASKED]], label {{%?}}[[THEN:.+]], label {{%?}}[[EXIT:.+]] +// ALL: [[THEN]] +// IRBUILDER-NEXT: call {{.*}}void [[FOO]]() +// NORMAL-NEXT: invoke {{.*}}void [[FOO]]() +// ALL: call {{.*}}void @__kmpc_end_masked([[IDENT_T_TY]]* [[DEFAULT_LOC]], i32 [[GTID]]) +// ALL-NEXT: br label {{%?}}[[EXIT]] +// ALL: [[EXIT]] +#pragma omp masked + foo(); + // ALL-NOT: call i32 @__kmpc_masked + // ALL-NOT: call void @__kmpc_end_masked + return a; +} + +// ALL-LABEL: lambda_masked +// TERM_DEBUG-LABEL: lambda_masked +void lambda_masked(int a, int b) { + auto l = [=]() { +#pragma omp masked + { + // ALL: call i32 @__kmpc_masked( + int c = a + b; + } + }; + + l(); + + auto l1 = [=]() { +#pragma omp parallel +#pragma omp masked + { + // ALL: call i32 @__kmpc_masked( + int c = a + b; + } + }; + + l1(); +} + +// ALL-LABEL: parallel_master +// TERM_DEBUG-LABEL: parallel_master +void parallel_master() { +#pragma omp parallel +#pragma omp masked + // TERM_DEBUG-NOT: __kmpc_global_thread_num + // TERM_DEBUG: call i32 @__kmpc_masked({{.+}}), !dbg [[DBG_LOC_START:![0-9]+]] + // TERM_DEBUG: invoke void {{.*}}foo{{.*}}() + // TERM_DEBUG: unwind label %[[TERM_LPAD:.+]], + // TERM_DEBUG-NOT: __kmpc_global_thread_num + // TERM_DEBUG: call void @__kmpc_end_masked({{.+}}), !dbg [[DBG_LOC_END:![0-9]+]] + // TERM_DEBUG: [[TERM_LPAD]] + // TERM_DEBUG: call void @__clang_call_terminate + // TERM_DEBUG: unreachable + foo(); +} +// TERM_DEBUG-DAG: [[DBG_LOC_START]] = !DILocation(line: [[@LINE-12]], +// TERM_DEBUG-DAG: [[DBG_LOC_END]] = !DILocation(line: [[@LINE-3]], + +#endif diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h --- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h @@ -629,6 +629,17 @@ BodyGenCallbackTy BodyGenCB, FinalizeCallbackTy FiniCB); + /// Generator for '#omp masked' + /// + /// \param Loc The insert and source location description. + /// \param BodyGenCB Callback that will generate the region code. + /// \param FiniCB Callback to finialize variable copies. + /// + /// \returns The insertion position *after* the master. + InsertPointTy createMasked(const LocationDescription &Loc, + BodyGenCallbackTy BodyGenCB, + FinalizeCallbackTy FiniCB, Value *Filter); + /// Generator for '#omp critical' /// /// \param Loc The insert and source location description. diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def --- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def +++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def @@ -237,6 +237,8 @@ __OMP_RTL(__kmpc_master, false, Int32, IdentPtr, Int32) __OMP_RTL(__kmpc_end_master, false, Void, IdentPtr, Int32) +__OMP_RTL(__kmpc_masked, false, Int32, IdentPtr, Int32, Int32) +__OMP_RTL(__kmpc_end_masked, false, Void, IdentPtr, Int32) __OMP_RTL(__kmpc_critical, false, Void, IdentPtr, Int32, KmpCriticalNamePtrTy) __OMP_RTL(__kmpc_critical_with_hint, false, Void, IdentPtr, Int32, KmpCriticalNamePtrTy, Int32) @@ -640,6 +642,10 @@ ParamAttrs(ReadOnlyPtrAttrs)) __OMP_RTL_ATTRS(__kmpc_end_master, InaccessibleArgOnlyAttrs, AttributeSet(), ParamAttrs(ReadOnlyPtrAttrs)) +__OMP_RTL_ATTRS(__kmpc_masked, InaccessibleArgOnlyAttrs, AttributeSet(), + ParamAttrs(ReadOnlyPtrAttrs)) +__OMP_RTL_ATTRS(__kmpc_end_masked, InaccessibleArgOnlyAttrs, AttributeSet(), + ParamAttrs(ReadOnlyPtrAttrs)) __OMP_RTL_ATTRS(__kmpc_critical, BarrierAttrs, AttributeSet(), ParamAttrs(ReadOnlyPtrAttrs, AttributeSet(), AttributeSet())) __OMP_RTL_ATTRS(__kmpc_critical_with_hint, BarrierAttrs, AttributeSet(), diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -901,6 +901,30 @@ /*Conditional*/ true, /*hasFinalize*/ true); } +OpenMPIRBuilder::InsertPointTy +OpenMPIRBuilder::createMasked(const LocationDescription &Loc, + BodyGenCallbackTy BodyGenCB, + FinalizeCallbackTy FiniCB, Value *Filter) { + if (!updateToLocation(Loc)) + return Loc.IP; + + Directive OMPD = Directive::OMPD_masked; + Constant *SrcLocStr = getOrCreateSrcLocStr(Loc); + Value *Ident = getOrCreateIdent(SrcLocStr); + Value *ThreadId = getOrCreateThreadID(Ident); + Value *Args[] = {Ident, ThreadId, Filter}; + Value *ArgsEnd[] = {Ident, ThreadId}; + + Function *EntryRTLFn = getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_masked); + Instruction *EntryCall = Builder.CreateCall(EntryRTLFn, Args); + + Function *ExitRTLFn = getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_end_masked); + Instruction *ExitCall = Builder.CreateCall(ExitRTLFn, ArgsEnd); + + return EmitOMPInlinedRegion(OMPD, EntryCall, ExitCall, BodyGenCB, FiniCB, + /*Conditional*/ true, /*hasFinalize*/ true); +} + CanonicalLoopInfo *OpenMPIRBuilder::createLoopSkeleton( DebugLoc DL, Value *TripCount, Function *F, BasicBlock *PreInsertBefore, BasicBlock *PostInsertBefore, const Twine &Name) { diff --git a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp --- a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp +++ b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp @@ -1788,6 +1788,88 @@ EXPECT_EQ(MasterEndCI->getArgOperand(1), MasterEntryCI->getArgOperand(1)); } +TEST_F(OpenMPIRBuilderTest, MaskedDirective) { + using InsertPointTy = OpenMPIRBuilder::InsertPointTy; + OpenMPIRBuilder OMPBuilder(*M); + OMPBuilder.initialize(); + F->setName("func"); + IRBuilder<> Builder(BB); + + OpenMPIRBuilder::LocationDescription Loc({Builder.saveIP(), DL}); + + AllocaInst *PrivAI = nullptr; + + BasicBlock *EntryBB = nullptr; + BasicBlock *ExitBB = nullptr; + BasicBlock *ThenBB = nullptr; + + auto BodyGenCB = [&](InsertPointTy AllocaIP, InsertPointTy CodeGenIP, + BasicBlock &FiniBB) { + if (AllocaIP.isSet()) + Builder.restoreIP(AllocaIP); + else + Builder.SetInsertPoint(&*(F->getEntryBlock().getFirstInsertionPt())); + PrivAI = Builder.CreateAlloca(F->arg_begin()->getType()); + Builder.CreateStore(F->arg_begin(), PrivAI); + + llvm::BasicBlock *CodeGenIPBB = CodeGenIP.getBlock(); + llvm::Instruction *CodeGenIPInst = &*CodeGenIP.getPoint(); + EXPECT_EQ(CodeGenIPBB->getTerminator(), CodeGenIPInst); + + Builder.restoreIP(CodeGenIP); + + // collect some info for checks later + ExitBB = FiniBB.getUniqueSuccessor(); + ThenBB = Builder.GetInsertBlock(); + EntryBB = ThenBB->getUniquePredecessor(); + + // simple instructions for body + Value *PrivLoad = + Builder.CreateLoad(PrivAI->getAllocatedType(), PrivAI, "local.use"); + Builder.CreateICmpNE(F->arg_begin(), PrivLoad); + }; + + auto FiniCB = [&](InsertPointTy IP) { + BasicBlock *IPBB = IP.getBlock(); + EXPECT_NE(IPBB->end(), IP.getPoint()); + }; + + Constant *Filter = ConstantInt::get(Type::getInt32Ty(M->getContext()), 0); + Builder.restoreIP( + OMPBuilder.createMasked(Builder, BodyGenCB, FiniCB, Filter)); + Value *EntryBBTI = EntryBB->getTerminator(); + EXPECT_NE(EntryBBTI, nullptr); + EXPECT_TRUE(isa(EntryBBTI)); + BranchInst *EntryBr = cast(EntryBB->getTerminator()); + EXPECT_TRUE(EntryBr->isConditional()); + EXPECT_EQ(EntryBr->getSuccessor(0), ThenBB); + EXPECT_EQ(ThenBB->getUniqueSuccessor(), ExitBB); + EXPECT_EQ(EntryBr->getSuccessor(1), ExitBB); + + CmpInst *CondInst = cast(EntryBr->getCondition()); + EXPECT_TRUE(isa(CondInst->getOperand(0))); + + CallInst *MaskedEntryCI = cast(CondInst->getOperand(0)); + EXPECT_EQ(MaskedEntryCI->getNumArgOperands(), 3U); + EXPECT_EQ(MaskedEntryCI->getCalledFunction()->getName(), "__kmpc_masked"); + EXPECT_TRUE(isa(MaskedEntryCI->getArgOperand(0))); + + CallInst *MaskedEndCI = nullptr; + for (auto &FI : *ThenBB) { + Instruction *cur = &FI; + if (isa(cur)) { + MaskedEndCI = cast(cur); + if (MaskedEndCI->getCalledFunction()->getName() == "__kmpc_end_masked") + break; + MaskedEndCI = nullptr; + } + } + EXPECT_NE(MaskedEndCI, nullptr); + EXPECT_EQ(MaskedEndCI->getNumArgOperands(), 2U); + EXPECT_TRUE(isa(MaskedEndCI->getArgOperand(0))); + EXPECT_EQ(MaskedEndCI->getArgOperand(1), MaskedEntryCI->getArgOperand(1)); +} + TEST_F(OpenMPIRBuilderTest, CriticalDirective) { using InsertPointTy = OpenMPIRBuilder::InsertPointTy; OpenMPIRBuilder OMPBuilder(*M);