[OpenMP][5.1] Fix parallel masked is ignored #59939

[randreshg]

Summary

Code generation support for 'parallel masked' directive.

The EmitOMPParallelMaskedDirective was implemented.
In addition, the appropiate device functions were added.

Fix #59939.

[tianshilei1992]

Can you add a test please?

[tianshilei1992]

clang/lib/CodeGen/CGStmt.cpp

431

also point out which directive is not supported.

[josemonsalve2]

I’m confused that Shilei said in the issue in GitHub that this is supported already. But front end is just throwing an error message.

Also. Should the same CGM.ErrorUnsupported be applied to all the llvm_unreachable in that switch

[tianshilei1992]

In D143527#4111295, @josemonsalve2 wrote:

I’m confused that Shilei said in the issue in GitHub that this is supported already. But front end is just throwing an error message.

Our Sema can support it but CodeGen doesn't. Here we want to error it out instead of "silently" pass if assertion is disabled.

Also. Should the same CGM.ErrorUnsupported be applied to all the llvm_unreachable in that switch

That might be a good idea.

[tianshilei1992]

clang/lib/CodeGen/CGStmt.cpp

431

We might want a more detailed description, like the code gen of xxx is not supported.

[josemonsalve2]

Thanks for the clarification Shilei

[randreshg]

clang/lib/CodeGen/CGStmt.cpp

431

The CGM.ErrorUnsupported function already prints out that information. I can create another function that explicitly errors out the message you suggest. WDYT?

[tianshilei1992]

clang/lib/CodeGen/CGStmt.cpp

431

Yeah, that's what I was referring to. It only says "this directive". I'd prefer explicit message.

[jdoerfert]

This is not a fix. It's "progress" in the sense that we should not have unreachable in the first place but proper errors. That said, it still doesn't solve the issue and certainly doesn't fix https://github.com/llvm/llvm-project/issues/59939. See my comments and use https://godbolt.org/z/q7MaKbarv as test for this.

clang/lib/CodeGen/CGStmt.cpp

248

... here. And see what happens.

430

Move this line...

[randreshg]

clang/lib/CodeGen/CGStmt.cpp

248

The front-end crashes in the cast to OMPParallelMasterDirective. I guess the best solution would be to implement EmitOMPParallelMaskedDirective function. WDYT?

[randreshg]

The parallel masked directive is emitted as a combined directive that consists of two implicit directives: 'parallel' with 'masked' directive.

[jdoerfert]

clang/test/OpenMP/parallel_masked.cpp

2

add the flag to include generated functions.

[josemonsalve2]

Just a suggestion, I think this works, although it may be best to use OpenMPIRBuilder.

clang/lib/CodeGen/CGStmtOpenMP.cpp

4492

I noticed that in EmitOMPParallelDirective and EmitOMPMaskedDirective there is an option to use OpenMPIRBuilder, instead of the traditional front end infrastructure.

Would it make sense to use this instead?

[randreshg]

Generated functions included in the test

[randreshg]

clang/lib/CodeGen/CGStmtOpenMP.cpp

4492

what do you think @jdoerfert ?

[jdoerfert]

clang/lib/CodeGen/CGStmtOpenMP.cpp

4492

I think we need to move it at some point. Didn't we move some of these already?

clang/test/OpenMP/parallel_masked.cpp

36

Did you try to run this? I would guess it fails to link because this function is missing in the device runtime.

[randreshg]

clang/lib/CodeGen/CGStmtOpenMP.cpp

4492

I'll check and come back to you if any questions

clang/test/OpenMP/parallel_masked.cpp

36

Yes, I did, and it worked. When you say it would fail to link in the device runtime, would that be the case where the parallel masked region is inside a target region? If that is what you mean, I haven't tested that case yet.

[jdoerfert]

clang/test/OpenMP/parallel_masked.cpp

36

Yes, that is what I meant. Sorry. We can do that in a follow up or in this patch.

[randreshg]

Adding masked support to the device

[randreshg]

clang/lib/CodeGen/CGStmtOpenMP.cpp

4492

I took a look at this and I think it would be better if we do it in a follow-up patch. The way I see it, in that patch, we can use OpenMPIRBuilder not only for EmitOMPParallelMasked but also for any other directive that uses the emitCommonOMPParallelDirective function, e.g, parallel for or parallel sections. WDYT? @josemonsalve2 @jdoerfert

[josemonsalve2]

clang/lib/CodeGen/CGStmtOpenMP.cpp

4492

Fine by me. The patch looks good to me but let’s wait for the big boss to respond.

[jdoerfert]

LG, minor nits, see below.

clang/test/OpenMP/parallel_masked.cpp

12

There is no test for the filter clause.

openmp/libomptarget/DeviceRTL/include/Interface.h

263

Style

openmp/libomptarget/DeviceRTL/src/Synchronization.cpp

523

In a follow up we should finally fix/cleanup the TId stuff. We still pass in the "level 0/1" TId and hope for the best. We could even use that if we knew we were in Lvl 0/1, that is, no nested parallel.
However, we generally don't. So we either generate better code to determine the TId we pass in, or have a special function here that uses TId for LVL 0/1 and the constant 0 otherwise. This could actually improve performance for a lot of the APIs as we had to replace the TId that is passed in all over the place to make it correct, but we don't always optimize the omp_get_thread_num out, I'd assume. Maybe we do and this is a non-issue, though we should check.

525

Style: Filter

[randreshg]

Test with filter clause added

[randreshg]

Test fixed

[randreshg]

Test updated.

[randreshg]

Test updated.
This should work 🤞

clang/lib/CodeGen/CGStmt.cpp

case Stmt::OMPCriticalDirectiveClass:
     EmitOMPCriticalDirective(cast<OMPCriticalDirective>(*S));
     break;
   case Stmt::OMPParallelForDirectiveClass:
     EmitOMPParallelForDirective(cast<OMPParallelForDirective>(*S));
     break;
   case Stmt::OMPParallelForSimdDirectiveClass:
     EmitOMPParallelForSimdDirective(cast<OMPParallelForSimdDirective>(*S));
     break;
   case Stmt::OMPParallelMasterDirectiveClass:
     EmitOMPParallelMasterDirective(cast<OMPParallelMasterDirective>(*S));
     break;
   case Stmt::OMPParallelSectionsDirectiveClass:
     EmitOMPParallelSectionsDirective(cast<OMPParallelSectionsDirective>(*S));
     break;
   case Stmt::OMPTaskDirectiveClass:
     EmitOMPTaskDirective(cast<OMPTaskDirective>(*S));
     break;
case Stmt::OMPTargetTeamsGenericLoopDirectiveClass:
     llvm_unreachable("target teams loop directive not supported yet.");
     break;
   case Stmt::OMPParallelGenericLoopDirectiveClass:
     llvm_unreachable("parallel loop directive not supported yet.");
     break;
   case Stmt::OMPTargetParallelGenericLoopDirectiveClass:
     llvm_unreachable("target parallel loop directive not supported yet.");
     break;
   case Stmt::OMPParallelMaskedDirectiveClass:
-    llvm_unreachable("parallel masked directive not supported yet.");
+    EmitOMPParallelMaskedDirective(cast<OMPParallelMaskedDirective>(*S));
     break;
   }
 }
 
 bool CodeGenFunction::EmitSimpleStmt(const Stmt *S,
                                      ArrayRef<const Attr *> Attrs) {
   switch (S->getStmtClass()) {
   default:

clang/lib/CodeGen/CGStmtOpenMP.cpp

emitCommonOMPParallelDirective(*this, S, OMPD_master, CodeGen,
                                    emitEmptyBoundParameters);
     emitPostUpdateForReductionClause(*this, S,
                                      [](CodeGenFunction &) { return nullptr; });
   }
   // Check for outer lastprivate conditional update.
   checkForLastprivateConditionalUpdate(*this, S);
 }
 
+void CodeGenFunction::EmitOMPParallelMaskedDirective(
+    const OMPParallelMaskedDirective &S) {
+  // Emit directive as a combined directive that consists of two implicit
+  // directives: 'parallel' with 'masked' directive.
+  auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
+    Action.Enter(CGF);
+    OMPPrivateScope PrivateScope(CGF);
+    emitOMPCopyinClause(CGF, S);
+    (void)CGF.EmitOMPFirstprivateClause(S, PrivateScope);
+    CGF.EmitOMPPrivateClause(S, PrivateScope);
+    CGF.EmitOMPReductionClauseInit(S, PrivateScope);
+    (void)PrivateScope.Privatize();
+    emitMasked(CGF, S);
+    CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_parallel);
+  };
+  {
+    auto LPCRegion =
+        CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S);
+    emitCommonOMPParallelDirective(*this, S, OMPD_masked, CodeGen,
+                                   emitEmptyBoundParameters);
+    emitPostUpdateForReductionClause(*this, S,
+                                     [](CodeGenFunction &) { return nullptr; });
+  }
+  // Check for outer lastprivate conditional update.
+  checkForLastprivateConditionalUpdate(*this, S);
+}
+
 void CodeGenFunction::EmitOMPParallelSectionsDirective(
     const OMPParallelSectionsDirective &S) {
   // Emit directive as a combined directive that consists of two implicit
   // directives: 'parallel' with 'sections' directive.
   auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
     Action.Enter(CGF);
     emitOMPCopyinClause(CGF, S);
     CGF.EmitSections(S);

clang/lib/CodeGen/CodeGenFunction.h

   void EmitOMPTargetTeamsDistributeParallelForDirective(
       const OMPTargetTeamsDistributeParallelForDirective &S);
   void EmitOMPTargetTeamsDistributeParallelForSimdDirective(
       const OMPTargetTeamsDistributeParallelForSimdDirective &S);
   void EmitOMPTargetTeamsDistributeSimdDirective(
       const OMPTargetTeamsDistributeSimdDirective &S);
   void EmitOMPGenericLoopDirective(const OMPGenericLoopDirective &S);
   void EmitOMPInteropDirective(const OMPInteropDirective &S);
+  void EmitOMPParallelMaskedDirective(const OMPParallelMaskedDirective &S);
 
   /// Emit device code for the target directive.
   static void EmitOMPTargetDeviceFunction(CodeGenModule &CGM,
                                           StringRef ParentName,
                                           const OMPTargetDirective &S);
   static void
   EmitOMPTargetParallelDeviceFunction(CodeGenModule &CGM, StringRef ParentName,
                                       const OMPTargetParallelDirective &S);

clang/lib/Parse/ParseOpenMP.cpp

 ///         taskloop' | 'parallel master taskloop simd' | 'distribute' | 'target
 ///         enter data' | 'target exit data' | 'target parallel' | 'target
 ///         parallel for' | 'target update' | 'distribute parallel for' |
 ///         'distribute paralle for simd' | 'distribute simd' | 'target parallel
 ///         for simd' | 'target simd' | 'teams distribute' | 'teams distribute
 ///         simd' | 'teams distribute parallel for simd' | 'teams distribute
 ///         parallel for' | 'target teams' | 'target teams distribute' | 'target
 ///         teams distribute parallel for' | 'target teams distribute parallel
-///         for simd' | 'target teams distribute simd' | 'masked' {clause}
-///         annot_pragma_openmp_end
+///         for simd' | 'target teams distribute simd' | 'masked' |
+///         'parallel masked' {clause} annot_pragma_openmp_end
 ///
 StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
     ParsedStmtContext StmtCtx, bool ReadDirectiveWithinMetadirective) {
   if (!ReadDirectiveWithinMetadirective)
     assert(Tok.isOneOf(tok::annot_pragma_openmp, tok::annot_attr_openmp) &&
            "Not an OpenMP directive!");
   ParsingOpenMPDirectiveRAII DirScope(*this);
   ParenBraceBracketBalancer BalancerRAIIObj(*this);

clang/test/OpenMP/parallel_masked.cpp

+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --prefix-filecheck-ir-name _
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -fopenmp-version=52 -x c -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+void foo();
+
+void masked() {
+    #pragma omp parallel masked
+    {
+        foo();
+    }
+}
+
+void maskedFilter() {
+    const int tid = 1;
+    #pragma omp parallel masked filter(tid)
+    {
+        foo();
+    }
+}
+
+void master() {
+    #pragma omp parallel master
+    {
+        foo();
+    }
+}
+// CHECK-LABEL: define {{[^@]+}}@masked
+// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1:[0-9]+]], i32 0, ptr @.omp_outlined.)
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@.omp_outlined.
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
+// CHECK-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_masked(ptr @[[GLOB1]], i32 [[TMP1]], i32 0)
+// CHECK-NEXT:    [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0
+// CHECK-NEXT:    br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]]
+// CHECK:       omp_if.then:
+// CHECK-NEXT:    call void (...) @foo()
+// CHECK-NEXT:    call void @__kmpc_end_masked(ptr @[[GLOB1]], i32 [[TMP1]])
+// CHECK-NEXT:    br label [[OMP_IF_END]]
+// CHECK:       omp_if.end:
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@maskedFilter
+// CHECK-SAME: () #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[TID:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    store i32 1, ptr [[TID]], align 4
+// CHECK-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @.omp_outlined..1)
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@.omp_outlined..1
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
+// CHECK-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_masked(ptr @[[GLOB1]], i32 [[TMP1]], i32 1)
+// CHECK-NEXT:    [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0
+// CHECK-NEXT:    br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]]
+// CHECK:       omp_if.then:
+// CHECK-NEXT:    call void (...) @foo()
+// CHECK-NEXT:    call void @__kmpc_end_masked(ptr @[[GLOB1]], i32 [[TMP1]])
+// CHECK-NEXT:    br label [[OMP_IF_END]]
+// CHECK:       omp_if.end:
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@master
+// CHECK-SAME: () #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @.omp_outlined..2)
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@.omp_outlined..2
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
+// CHECK-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_master(ptr @[[GLOB1]], i32 [[TMP1]])
+// CHECK-NEXT:    [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0
+// CHECK-NEXT:    br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]]
+// CHECK:       omp_if.then:
+// CHECK-NEXT:    call void (...) @foo()
+// CHECK-NEXT:    call void @__kmpc_end_master(ptr @[[GLOB1]], i32 [[TMP1]])
+// CHECK-NEXT:    br label [[OMP_IF_END]]
+// CHECK:       omp_if.end:
+// CHECK-NEXT:    ret void
+//

clang/test/OpenMP/parallel_masked_target.cpp

+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --prefix-filecheck-ir-name _
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -fopenmp-version=52 -fopenmp-targets=nvptx64 -offload-device-only -x c -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+void foo();
+
+void masked() {
+    #pragma target
+    #pragma omp parallel masked
+    {
+        foo();
+    }
+}
+
+void maskedFilter() {
+    const int tid = 1;
+    #pragma target
+    #pragma omp parallel masked filter(tid)
+    {
+        foo();
+    }
+}
+
+void master() {
+    #pragma target
+    #pragma omp parallel master
+    {
+        foo();
+    }
+}
+// CHECK-LABEL: define {{[^@]+}}@masked
+// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1:[0-9]+]], i32 0, ptr @.omp_outlined.)
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@.omp_outlined.
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
+// CHECK-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_masked(ptr @[[GLOB1]], i32 [[TMP1]], i32 0)
+// CHECK-NEXT:    [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0
+// CHECK-NEXT:    br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]]
+// CHECK:       omp_if.then:
+// CHECK-NEXT:    call void (...) @foo()
+// CHECK-NEXT:    call void @__kmpc_end_masked(ptr @[[GLOB1]], i32 [[TMP1]])
+// CHECK-NEXT:    br label [[OMP_IF_END]]
+// CHECK:       omp_if.end:
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@maskedFilter
+// CHECK-SAME: () #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[TID:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    store i32 1, ptr [[TID]], align 4
+// CHECK-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @.omp_outlined..1)
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@.omp_outlined..1
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
+// CHECK-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_masked(ptr @[[GLOB1]], i32 [[TMP1]], i32 1)
+// CHECK-NEXT:    [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0
+// CHECK-NEXT:    br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]]
+// CHECK:       omp_if.then:
+// CHECK-NEXT:    call void (...) @foo()
+// CHECK-NEXT:    call void @__kmpc_end_masked(ptr @[[GLOB1]], i32 [[TMP1]])
+// CHECK-NEXT:    br label [[OMP_IF_END]]
+// CHECK:       omp_if.end:
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@master
+// CHECK-SAME: () #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @.omp_outlined..2)
+// CHECK-NEXT:    ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@.omp_outlined..2
+// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
+// CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
+// CHECK-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_master(ptr @[[GLOB1]], i32 [[TMP1]])
+// CHECK-NEXT:    [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0
+// CHECK-NEXT:    br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]]
+// CHECK:       omp_if.then:
+// CHECK-NEXT:    call void (...) @foo()
+// CHECK-NEXT:    call void @__kmpc_end_master(ptr @[[GLOB1]], i32 [[TMP1]])
+// CHECK-NEXT:    br label [[OMP_IF_END]]
+// CHECK:       omp_if.end:
+// CHECK-NEXT:    ret void
+//

openmp/libomptarget/DeviceRTL/include/Interface.h

 void __kmpc_barrier_simple_spmd(IdentTy *Loc_ref, int32_t TId);
 
 void __kmpc_barrier_simple_generic(IdentTy *Loc_ref, int32_t TId);
 
 int32_t __kmpc_master(IdentTy *Loc, int32_t TId);
 
 void __kmpc_end_master(IdentTy *Loc, int32_t TId);
 
+int32_t __kmpc_masked(IdentTy *Loc, int32_t TId, int32_t Filter);
+
+void __kmpc_end_masked(IdentTy *Loc, int32_t TId);
+
 int32_t __kmpc_single(IdentTy *Loc, int32_t TId);
 
 void __kmpc_end_single(IdentTy *Loc, int32_t TId);
 
 void __kmpc_flush(IdentTy *Loc);
 
 uint64_t __kmpc_warp_active_thread_mask(void);
 

openmp/libomptarget/DeviceRTL/src/Synchronization.cpp

 
 int32_t __kmpc_master(IdentTy *Loc, int32_t TId) {
   FunctionTracingRAII();
   return omp_get_thread_num() == 0;
 }
 
 void __kmpc_end_master(IdentTy *Loc, int32_t TId) { FunctionTracingRAII(); }
 
+int32_t __kmpc_masked(IdentTy *Loc, int32_t TId, int32_t Filter) {
+  FunctionTracingRAII();
+  return omp_get_thread_num() == Filter;
+}
+
+void __kmpc_end_masked(IdentTy *Loc, int32_t TId) { FunctionTracingRAII(); }
+
 int32_t __kmpc_single(IdentTy *Loc, int32_t TId) {
   FunctionTracingRAII();
   return __kmpc_master(Loc, TId);
 }
 
 void __kmpc_end_single(IdentTy *Loc, int32_t TId) {
   FunctionTracingRAII();
   // The barrier is explicitly called.