[clang] e0c2125 - [OpenMP] Added codegen for masked directive
via cfe-commits
cfe-commits at lists.llvm.org
Thu Apr 15 10:55:28 PDT 2021
Author: cchen
Date: 2021-04-15T12:55:07-05:00
New Revision: e0c2125d1d1e72039b8e071d468d9f740c7dbfbd
URL: https://github.com/llvm/llvm-project/commit/e0c2125d1d1e72039b8e071d468d9f740c7dbfbd
DIFF: https://github.com/llvm/llvm-project/commit/e0c2125d1d1e72039b8e071d468d9f740c7dbfbd.diff
LOG: [OpenMP] Added codegen for masked directive
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D100514
Added:
clang/test/OpenMP/masked_codegen.cpp
Modified:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/CodeGen/CGOpenMPRuntime.h
clang/lib/CodeGen/CGStmt.cpp
clang/lib/CodeGen/CGStmtOpenMP.cpp
clang/lib/CodeGen/CodeGenFunction.h
llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 37146989e2522..18ad5066fbba2 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -2279,6 +2279,35 @@ void CGOpenMPRuntime::emitMasterRegion(CodeGenFunction &CGF,
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 @@ void CGOpenMPRuntime::emitInlinedDirective(CodeGenFunction &CGF,
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 @@ void CGOpenMPSIMDRuntime::emitMasterRegion(CodeGenFunction &CGF,
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/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h
index 541904aa3988b..c35202a513711 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -1012,6 +1012,14 @@ class CGOpenMPRuntime {
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 @@ class CGOpenMPSIMDRuntime final : public CGOpenMPRuntime {
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/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp
index daecb44abbd89..c2b64c899eeb4 100644
--- a/clang/lib/CodeGen/CGStmt.cpp
+++ b/clang/lib/CodeGen/CGStmt.cpp
@@ -382,7 +382,7 @@ void CodeGenFunction::EmitStmt(const Stmt *S, ArrayRef<const Attr *> Attrs) {
llvm_unreachable("Dispatch directive not supported yet.");
break;
case Stmt::OMPMaskedDirectiveClass:
- llvm_unreachable("Masked directive not supported yet.");
+ EmitOMPMaskedDirective(cast<OMPMaskedDirective>(*S));
break;
}
}
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 4b93fddab5235..827102ba7d7a2 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -3844,6 +3844,55 @@ void CodeGenFunction::EmitOMPMasterDirective(const OMPMasterDirective &S) {
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<OMPFilterClause>())
+ 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<OMPFilterClause>())
+ 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 +6979,8 @@ void CodeGenFunction::EmitSimpleOMPExecutableDirective(
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
index ceb161f5a8cb1..0b3da9e088d09 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -3421,6 +3421,7 @@ class CodeGenFunction : public CodeGenTypeCache {
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
index 0000000000000..97cb0375edb97
--- /dev/null
+++ b/clang/test/OpenMP/masked_codegen.cpp
@@ -0,0 +1,143 @@
+// 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 2)
+// 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 filter(2)
+ foo();
+// ALL: store i32 9, i32* [[X:.+]],
+// ALL: [[X_VAL:%.+]] = load i32, i32* [[X]]
+// 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 [[X_VAL]])
+// 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]]
+ int x = 9;
+#pragma omp masked filter(x)
+ 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 filter(1)
+ {
+ // ALL: call i32 @__kmpc_masked(
+ int c = a + b;
+ }
+ };
+
+ l1();
+
+ int y = 1;
+ auto l2 = [=](int yy) {
+#pragma omp parallel
+#pragma omp masked filter(yy)
+ {
+ // ALL: call i32 @__kmpc_masked(
+ int c = a + b;
+ }
+ };
+
+ l2(y);
+}
+
+// ALL-LABEL: parallel_masked
+// TERM_DEBUG-LABEL: parallel_masked
+void parallel_masked() {
+#pragma omp parallel
+#pragma omp masked filter(1)
+ // 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();
+
+ int x;
+#pragma omp parallel
+#pragma omp masked filter(x)
+ // 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
index 997c078d48bc6..5a4b406649409 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -629,6 +629,17 @@ class OpenMPIRBuilder {
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
index 533a10f047896..5b403e00df16c 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -237,6 +237,8 @@ __OMP_RTL(omp_set_max_active_levels, false, Void, Int32)
__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 @@ __OMP_RTL_ATTRS(__kmpc_master, InaccessibleArgOnlyAttrs, AttributeSet(),
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
index 26f5901d3440c..ec9ecced08e74 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -901,6 +901,30 @@ OpenMPIRBuilder::createMaster(const LocationDescription &Loc,
/*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
index 344c66c3165b5..da81367180492 100644
--- a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
+++ b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
@@ -1788,6 +1788,88 @@ TEST_F(OpenMPIRBuilderTest, MasterDirective) {
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<BranchInst>(EntryBBTI));
+ BranchInst *EntryBr = cast<BranchInst>(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<CmpInst>(EntryBr->getCondition());
+ EXPECT_TRUE(isa<CallInst>(CondInst->getOperand(0)));
+
+ CallInst *MaskedEntryCI = cast<CallInst>(CondInst->getOperand(0));
+ EXPECT_EQ(MaskedEntryCI->getNumArgOperands(), 3U);
+ EXPECT_EQ(MaskedEntryCI->getCalledFunction()->getName(), "__kmpc_masked");
+ EXPECT_TRUE(isa<GlobalVariable>(MaskedEntryCI->getArgOperand(0)));
+
+ CallInst *MaskedEndCI = nullptr;
+ for (auto &FI : *ThenBB) {
+ Instruction *cur = &FI;
+ if (isa<CallInst>(cur)) {
+ MaskedEndCI = cast<CallInst>(cur);
+ if (MaskedEndCI->getCalledFunction()->getName() == "__kmpc_end_masked")
+ break;
+ MaskedEndCI = nullptr;
+ }
+ }
+ EXPECT_NE(MaskedEndCI, nullptr);
+ EXPECT_EQ(MaskedEndCI->getNumArgOperands(), 2U);
+ EXPECT_TRUE(isa<GlobalVariable>(MaskedEndCI->getArgOperand(0)));
+ EXPECT_EQ(MaskedEndCI->getArgOperand(1), MaskedEntryCI->getArgOperand(1));
+}
+
TEST_F(OpenMPIRBuilderTest, CriticalDirective) {
using InsertPointTy = OpenMPIRBuilder::InsertPointTy;
OpenMPIRBuilder OMPBuilder(*M);
More information about the cfe-commits
mailing list