OpenMP 5.0 metadirective

[alokmishra.besu]

Summary

This patch supports OpenMP 5.0 metadirective features.
It is implemented keeping the OpenMP 5.1 features like dynamic user condition in mind.

A new function, getBestWhenMatchForContext, is defined in llvm/Frontend/OpenMP/OMPContext.h

Currently this function return the index of the when clause with the highest score from the ones applicable in the Context.
But this function is declared with an array which can be used in OpenMP 5.1 implementation to select all the valid when clauses which can be resolved in runtime. Currently this array is set to null by default and its implementation is left for future.

[jdoerfert]

This looks close to an OpenMP 5.0 implementation. I left comments inlined.

We need tests that show how non-selected alternatives *do not* impact the program. As an example, a template instantiation inside of a non-selected alternative is not actually performed.

We also need test with ill-formed metadirectives.

clang/include/clang/Basic/DiagnosticSemaKinds.td

10493

no !. The default clause doesn't need to be in the end either.

clang/include/clang/Serialization/ASTBitCodes.h

1952

Unrelated.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10211

Can you explain this, this seems odd to me.

clang/lib/Parse/ParseOpenMP.cpp

2186

Usually you would check the return value in case we later actually propagate errors while parsing the context selector.

2193

If we give up it should be an error, I think. If we issue a warning we just pretend the colon was there afterwards.

2204

We have balanced trackers for this that also deal with the fact that there might be a ) missing. This code will probably crash.

2216

Add a TODO that nullptr should be replaced as per the use in Sema::ActOnOpenMPCall.

2237

Use a BalancedDelimiterTracker.

2254

If you warn and continue above you need to check for : here again.

2260

What is this token? We have skipped this part before so we need to validate it is what we expect it to be.

2264

Should we not go back to the original code handling "directives" instead? This looks like it is copied here.

2317

Same as below, change the order. Also, the "skipping" part is always the same, put it in a helper function or lambda.

2324

Move the smaller case first and use an early exit. That will reduce the indention of the larger case by 1.

clang/test/OpenMP/metadirective_construct.cpp

12

Since when does the construct trait work? I'm confused.

clang/test/OpenMP/metadirective_empty.cpp

1

no -fopenmp-targets please.

clang/test/OpenMP/metadirective_implementation.cpp

1

Can we run this for all configurations of the metadirective so we can actually see it will pick the right one, not the first?

llvm/include/llvm/Frontend/OpenMP/OMPContext.h

197

Leftover.

llvm/lib/Frontend/OpenMP/OMPContext.cpp

337

This looks like a clone of getBestVariantMatchForContext with an extra unused argument.

400

Leftover.

[alokmishra.besu]

I have replied to the comments and will update the code accordingly. Some of the codes are intentionally left to be update in 5.1 implementation. Will add TODO comments there.
I will revisit the implementation of getBestWhenMatchForContext and also add more test cases over the weekend and submit new code by next week.

clang/include/clang/Basic/DiagnosticSemaKinds.td

10493

OK. Will update accordingly.

clang/include/clang/Serialization/ASTBitCodes.h

1952

Only STMT_OMP_META_DIRECTIVE was added. Rest was formatted by git clang-format

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10211

Good catch. This was an experimental code for 5.1. Got committed by mistake. Will update.

clang/lib/Parse/ParseOpenMP.cpp

2186

OK. Will update accordingly.

2193

OK. Will update accordingly.

2204

Come to think of it, in case of a missing ')' this code might end up in an infinite loop. Will update accordingly.

2216

Will do

2254

Will update accordingly.

2264

Unfortunately we cannot go to the original code handling since the original code handling assumes that the directive always ends with annot_pragma_openmp_end, while here it will always end with ')'.
In specification 5.0, since we are choosing only 1 directive, the body of the while block remains the same as the original code. Only the condition of the while block changes. In specification 5.1, we will need to generate code for dynamic handling and even the body will differ as we might need to generate AST node for multiple directives. It is best if we handle this code here for easier handling of 5.1 code, than in the original code space.
I will add a TODO comment here.

2317

OK

2324

OK

clang/test/OpenMP/metadirective_empty.cpp

1

OK

clang/test/OpenMP/metadirective_implementation.cpp

1

Good catch. I will go through the implementation of getBestWhenMatchForContext to check that.

llvm/include/llvm/Frontend/OpenMP/OMPContext.h

197

Will remove.

llvm/lib/Frontend/OpenMP/OMPContext.cpp

337

I intended to keep it similar for 5.0 to be updated in 5.1 code. But anyways it seems to be giving wrong result. Will go through this again.

400

Will remove.

[alokmishra.besu]

Updated version of metadirective supporting OpenMP 5.0 features

[ABataev]

clang/include/clang/AST/StmtOpenMP.h

373

I think, metadirective should be a kind of a container for different sub-directives. The problem is that that subdirectives could be completely different, they may capture different variables, using different capture kind (by value or by reference) etc.So, you need to generate each possible sub-directive independently and store them in the meta directive node. Otherwise you won't be able to generate the code correctly.

[alokmishra.besu]

clang/include/clang/AST/StmtOpenMP.h

373

In OpenMP 5.0, we do not need to generate every sub-directive. Rather we need to select one (or none) directive and replace metadirective with it. So this is not needed.
Yes with future specifications we will need to include a list of all valid directives which need to be resolved at runtime. That is when we will need to generate and store multiple sub-directives inside the OMPMetaDirective class.

[ABataev]

clang/include/clang/AST/StmtOpenMP.h

373

I think you still need to do it even for 5.0. I don't mean y need to emit the code for every sub-directive, but to generate them in the AST. Even in the example above, when(user={condition(N>10)} is used in the template and N is a template argument, the chosen directive can be selected on the instantiation and, thus, the original templated directive has to contain all possible sub-directives.

[jdoerfert]

clang/include/clang/Serialization/ASTBitCodes.h

1952

Rest was formatted by git clang-format

Please don't format unrelated code.

clang/lib/Parse/ParseOpenMP.cpp

2264

Unfortunately we cannot go to the original code handling since the original code handling assumes that the directive always ends with annot_pragma_openmp_end, while here it will always end with ')'.

Let's add a flag to the original handling to make this possible then. Copying it is going to create more long term problems.

llvm/lib/Frontend/OpenMP/OMPContext.cpp

337

Use the getBestVariantMatchForContext interface now and for the 5.1 semantics we introduce the OrderedMatch. That can also be used in the call site declare variant handling which currently calls getBestVariantMatchForContext multiple times.

[alokmishra.besu]

clang/include/clang/Serialization/ASTBitCodes.h

1952

Removed.

clang/lib/Parse/ParseOpenMP.cpp

2264

Done.

[dreachem]

In D91944#2414364, @jdoerfert wrote:

This looks close to an OpenMP 5.0 implementation. I left comments inlined.

We need tests that show how non-selected alternatives *do not* impact the program. As an example, a template instantiation inside of a non-selected alternative is not actually performed.

We also need test with ill-formed metadirectives.

@jdoerfert I'm still trying to understand this thing regarding template instantiations. The spec says that a directive variant associated with a when clause can only affect the program if it is a dynamic replacement candidate. I had assumed this is referring to the runtime behavior of the program, and not (for instance) whether a compiler error is emitted due to a failing static assert.

Also, in order to determine what are the dynamic replacement candidates, and their order, all specified score expressions would need to be evaluated. Then, you'd need to evaluate the static context selectors, in decreasing order of their score, to find which one is the last dynamic replacement candidate. So I think template instantiations could be possible for those expressions, even if the corresponding variants aren't selected?

clang/include/clang/AST/StmtOpenMP.h

373

r

[jdoerfert]

In D91944#2457744, @dreachem wrote:

In D91944#2414364, @jdoerfert wrote:

This looks close to an OpenMP 5.0 implementation. I left comments inlined.

We need tests that show how non-selected alternatives *do not* impact the program. As an example, a template instantiation inside of a non-selected alternative is not actually performed.

We also need test with ill-formed metadirectives.

@jdoerfert I'm still trying to understand this thing regarding template instantiations.

The problem is this patch can only resolve to a single directive during parsing. Take

template<bool b>
void foo() {
  #pragma omp metadirective when(user={condition(b)}) ...
}

which is not resolvable at parsing time but it is a valid OpenMP 5.0 use of the metadirective
that needs to resolve at compile time.

The spec says that a directive variant associated with a when clause can only affect the program if it is a dynamic replacement candidate. I had assumed this is referring to the runtime behavior of the program, and not (for instance) whether a compiler error is emitted due to a failing static assert.

I always assume(d) that static_assert is something "affecting the program". It doesn't matter though because if you instantiate all the clauses eagerly you change things for the runtime as well, e.g., you can cause different template instances to be selected later on depending on earlier instantiations. So it's not "just" static_assert.

Also, in order to determine what are the dynamic replacement candidates, and their order, all specified score expressions would need to be evaluated. Then, you'd need to evaluate the static context selectors, in decreasing order of their score, to find which one is the last dynamic replacement candidate. So I think template instantiations could be possible for those expressions, even if the corresponding variants aren't selected?

Yes, we will always evaluate selectors and scores. The key is that we won't evaluate the directive variant, which is neither a selector nor a score.

I hope this makes sense.

[jdoerfert]

As discussed in the OpenMP in LLVM call on December 9th, we should go ahead with this patch and add support for the missing features, e.g., dependent selectors, afterwards. LGTM

[jdoerfert]

(Please make sure all clang and clang-tidy tests pass properly so we can merge the patch)

[ABataev]

clang/lib/Sema/TreeTransform.h

8383–8388

It would not work correctly, I would replace it with just error emission that instantiation is not supported yet.

[alokmishra.besu]

All clang test and clang-tidy pass. Ran clang-format on all code, except clang/include/clang/Serialization/ASTBitCodes.h which was formatting unrelated code as well.

[alokmishra.besu]

Emitting error instead of instantiation in clang/lib/Sema/TreeTransform.h. Added a TODO for future.

[alokmishra.besu]

Rebased with current branch to avoid patch failure

[cchen]

@alokmishra.besu, I'm trying to play/learn with your patch so I downloaded the diff and apply it to the master branch, however, I'm getting error messages everywhere if llvm::omp::OMPD_metadirective appears:

llvm-project/clang/include/clang/AST/StmtOpenMP.h:384:43: error: no member named 'OMPD_metadirective' in namespace 'llvm::omp'
                               llvm::omp::OMPD_metadirective, SourceLocation(),
                               ~~~~~~~~~~~^

Do you know why this is happening? (I've also tried clean build it but still get the error msg) Thanks!

Edit: I've been successfully compiled it after changing a computer, thanks.

[jdoerfert]

Are the test all passing?

[protze.joachim]

I found two issues with this patch regarding the default clause:

• The spec does not require a default clause. I get error: expected expression if I omit a default clause. The error is gone if I add default().
• The spec does not allow an empty default() clause, but rather expects default(nothing)or omission of the default clause. This patch accepts the empty default clause.

[jdoerfert]

In D91944#2504924, @protze.joachim wrote:

I found two issues with this patch regarding the default clause:

• The spec does not require a default clause. I get error: expected expression if I omit a default clause. The error is gone if I add default().
• The spec does not allow an empty default() clause, but rather expects default(nothing)or omission of the default clause. This patch accepts the empty default clause.

There are various issues with this patch. We will fix them afterwards. The patch provides a lot of infrastructure we need. Please provide the problematic cases in a bug.

[jdoerfert]

In D91944#2485167, @jdoerfert wrote:

Are the test all passing?

Can we fix the failing tests so we can merge this?

[alokmishra.besu]

In D91944#2505159, @jdoerfert wrote:

In D91944#2485167, @jdoerfert wrote:

Are the test all passing?

Can we fix the failing tests so we can merge this?

I will work on it this week. Will update soon.

[alokmishra.besu]

I've rebase this patch with the latest git code.
All test cases pass.
I've also applied this patch to a new git clone. It applies and builds successfully.

[jdoerfert]

There are no test cases anymore, as far as I can tell.

[alokmishra.besu]

In D91944#2540178, @jdoerfert wrote:

There are no test cases anymore, as far as I can tell.

Updated patch.

[jdoerfert]

Asssuming this passes the tests I'll merge it as is. There are unresolved review requests, e.g., custom implementation instead of balanced parenthesis trackers, the template support, etc.
but we decided we can make faster progress in tree once this is merged.

[protze.joachim]

In D91944#2458154, @jdoerfert wrote:

The problem is this patch can only resolve to a single directive during parsing. Take

template<bool b>
void foo() {
  #pragma omp metadirective when(user={condition(b)}) ...
}

which is not resolvable at parsing time but it is a valid OpenMP 5.0 use of the metadirective
that needs to resolve at compile time.

It's unfortunate that this pattern is not supported by this patch. How difficult would it be to provide this support, or would the implementation be not different to dynamic conditions?
For this templated pattern, the optimizer would throw away the branch.

[cchen]

I'm guessing the tests were not pass on buildbot but passed on the author's side is due to the assertion was disabled on the author's side.
Here is the patch for avoiding all the assertion errors and I'm able to get all the metadirective tests passed (and no regression for the existing omp tests) with this change:

diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 5ac6dd0fb55a..3313efe524f3 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -2277,8 +2277,10 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
 ///
 StmtResult
 Parser::ParseOpenMPDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx) {
-  assert(Tok.is(tok::annot_pragma_openmp) && "Not an OpenMP directive!");
   static bool ReadDirectiveWithinMetadirective = false;
+  if (!ReadDirectiveWithinMetadirective)
+    assert(Tok.isOneOf(tok::annot_pragma_openmp, tok::annot_attr_openmp) &&
+           "Not an OpenMP directive!");
   ParsingOpenMPDirectiveRAII DirScope(*this);
   ParenBraceBracketBalancer BalancerRAIIObj(*this);
   SmallVector<OMPClause *, 5> Clauses;
@@ -2288,7 +2290,7 @@ Parser::ParseOpenMPDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx) {
   unsigned ScopeFlags = Scope::FnScope | Scope::DeclScope |
                         Scope::CompoundStmtScope | Scope::OpenMPDirectiveScope;
   if (!ReadDirectiveWithinMetadirective)
-      ConsumeAnnotationToken();
+    ConsumeAnnotationToken();
   SourceLocation Loc = Tok.getLocation(), EndLoc;
   OpenMPDirectiveKind DKind = parseOpenMPDirectiveKind(*this);
   if (ReadDirectiveWithinMetadirective && DKind == OMPD_unknown) {
@@ -2331,6 +2333,7 @@ Parser::ParseOpenMPDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx) {
         parseOMPContextSelectors(Loc, TI);
         if (TI.Sets.size() == 0) {
           Diag(Tok, diag::err_omp_expected_context_selector) << "when clause";
+          TPA.Commit();
           return Directive;
         }

@@ -2339,6 +2342,7 @@ Parser::ParseOpenMPDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx) {
           ConsumeAnyToken();
         else {
           Diag(Tok, diag::err_omp_expected_colon) << "when clause";
+          TPA.Commit();
           return Directive;
         }
       }
@@ -2352,6 +2356,8 @@ Parser::ParseOpenMPDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx) {
         if (Tok.is(tok::annot_pragma_openmp_end)) {
           Diag(Tok, diag::err_omp_expected_punc)
               << getOpenMPClauseName(CKind) << 0;
+          TPA.Commit();
           return Directive;
         }
         ConsumeAnyToken();
@@ -2627,7 +2633,7 @@ Parser::ParseOpenMPDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx) {
       // ends with a ')'.
       if (ReadDirectiveWithinMetadirective && Tok.is(tok::r_paren)) {
         while (Tok.isNot(tok::annot_pragma_openmp_end))
-          ConsumeToken();
+          ConsumeAnyToken();
         break;
       }
       bool HasImplicitClause = false;

[cchen]

@alokmishra.besu do you mind if I push the patch for solving those assertions with rebase? The patch does not change the logic in your program and I've made sure that the tests could pass in debug mode.

[alokmishra.besu]

In D91944#2990234, @cchen wrote:

@alokmishra.besu do you mind if I push the patch for solving those assertions with rebase? The patch does not change the logic in your program and I've made sure that the tests could pass in debug mode.

I think your patch is valid. Thanks for it. Please go ahead.

[cchen]

Rebase and avoid assert errors

[cchen]

Remove redundant file

[cchen]

Remove debug info and spaces

[cchen]

Fix ast errors and add some tiny fixes

[cchen]

Rebase and fix warning

[cchen]

Fix flang issue

[cchen]

Fix tests for Windows

[ABataev]

clang/include/clang/AST/StmtOpenMP.h

4789

1. Make it private, please
2. Var names must start from capital letter stmt->S

clang/lib/Parse/ParseOpenMP.cpp

2239

Please, fix var names.

[cchen]

clang/include/clang/AST/StmtOpenMP.h

4789

I'm going to fix it. By the way, do you know how to fix those failing Windows tests? I tried to fix it by adding ifndef HEADER to test files and also replace var names with regex, but those Windows tests still fail. Thanks

[ABataev]

clang/include/clang/AST/StmtOpenMP.h

4789

No idea. Need a bit of context (I mean, would be good to see the test log etc.).

[cchen]

clang/include/clang/AST/StmtOpenMP.h

4789

I think the run line might be so general, I update it and see if that helps.

[ABataev]

clang/include/clang/AST/StmtOpenMP.h

4789

Just provide target triple in all test runs commands for linux and make a separate test for windows (if required), it should fix the issue with the testing.

[cchen]

Fix tests and coding style

[cchen]

I reapplied the patch since I forgot to add author's name in the commit

[phosek]

We started seeing a test failure in Clang :: OpenMP/metadirective_device_kind_codegen.c on our macOS bot after this change landed (we aren't seeing it on other systems):

Script:
--
: 'RUN: at line 1';   /opt/s/w/ir/x/w/staging/llvm_build/bin/clang -cc1 -internal-isystem /opt/s/w/ir/x/w/staging/llvm_build/lib/clang/14.0.0/include -nostdsysteminc -verify -fopenmp -x c -std=c99 -emit-llvm /opt/s/w/ir/x/w/llvm-project/clang/test/OpenMP/metadirective_device_kind_codegen.c -o - | /opt/s/w/ir/x/w/staging/llvm_build/bin/FileCheck --allow-unused-prefixes /opt/s/w/ir/x/w/llvm-project/clang/test/OpenMP/metadirective_device_kind_codegen.c
: 'RUN: at line 2';   /opt/s/w/ir/x/w/staging/llvm_build/bin/clang -cc1 -internal-isystem /opt/s/w/ir/x/w/staging/llvm_build/lib/clang/14.0.0/include -nostdsysteminc -verify -fopenmp -x c -triple x86_64-unknown-linux -emit-llvm /opt/s/w/ir/x/w/llvm-project/clang/test/OpenMP/metadirective_device_kind_codegen.c -o - | /opt/s/w/ir/x/w/staging/llvm_build/bin/FileCheck --allow-unused-prefixes /opt/s/w/ir/x/w/llvm-project/clang/test/OpenMP/metadirective_device_kind_codegen.c
: 'RUN: at line 3';   /opt/s/w/ir/x/w/staging/llvm_build/bin/clang -cc1 -internal-isystem /opt/s/w/ir/x/w/staging/llvm_build/lib/clang/14.0.0/include -nostdsysteminc -verify -fopenmp -x c -triple aarch64-unknown-linux -emit-llvm /opt/s/w/ir/x/w/llvm-project/clang/test/OpenMP/metadirective_device_kind_codegen.c -o - | /opt/s/w/ir/x/w/staging/llvm_build/bin/FileCheck --allow-unused-prefixes /opt/s/w/ir/x/w/llvm-project/clang/test/OpenMP/metadirective_device_kind_codegen.c
: 'RUN: at line 4';   /opt/s/w/ir/x/w/staging/llvm_build/bin/clang -cc1 -internal-isystem /opt/s/w/ir/x/w/staging/llvm_build/lib/clang/14.0.0/include -nostdsysteminc -verify -fopenmp -x c -triple ppc64le-unknown-linux -emit-llvm /opt/s/w/ir/x/w/llvm-project/clang/test/OpenMP/metadirective_device_kind_codegen.c -o - | /opt/s/w/ir/x/w/staging/llvm_build/bin/FileCheck --allow-unused-prefixes /opt/s/w/ir/x/w/llvm-project/clang/test/OpenMP/metadirective_device_kind_codegen.c
--
Exit Code: 1

Command Output (stderr):
--
/opt/s/w/ir/x/w/llvm-project/clang/test/OpenMP/metadirective_device_kind_codegen.c:39:17: error: CHECK-LABEL: expected string not found in input
// CHECK-LABEL: define {{.+}} void @foo()
                ^
<stdin>:1:1: note: scanning from here
; ModuleID = '/opt/s/w/ir/x/w/llvm-project/clang/test/OpenMP/metadirective_device_kind_codegen.c'
^
<stdin>:13:1: note: possible intended match here
define void @foo() #0 {
^

Input file: <stdin>
Check file: /opt/s/w/ir/x/w/llvm-project/clang/test/OpenMP/metadirective_device_kind_codegen.c

-dump-input=help explains the following input dump.

Input was:
<<<<<<
            1: ; ModuleID = '/opt/s/w/ir/x/w/llvm-project/clang/test/OpenMP/metadirective_device_kind_codegen.c' 
label:39'0     X~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ error: no match found
            2: source_filename = "/opt/s/w/ir/x/w/llvm-project/clang/test/OpenMP/metadirective_device_kind_codegen.c" 
label:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            3: target datalayout = "e-m:o-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128" 
label:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            4: target triple = "x86_64-apple-darwin19.6.0" 
label:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            5:  
label:39'0     ~
            6: %struct.ident_t = type { i32, i32, i32, i32, i8* } 
label:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            7:  
label:39'0     ~
            8: @0 = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1 
label:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            9: @1 = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @0, i32 0, i32 0) }, align 8 
label:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           10: @2 = private unnamed_addr constant %struct.ident_t { i32 0, i32 514, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @0, i32 0, i32 0) }, align 8 
label:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           11:  
label:39'0     ~
           12: ; Function Attrs: noinline nounwind optnone 
label:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           13: define void @foo() #0 { 
label:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~
label:39'1     ?                        possible intended match
           14: entry: 
label:39'0     ~~~~~~~
           15:  %0 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1) 
label:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           16:  call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @1, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*)) 
label:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           17:  call void @__kmpc_push_num_threads(%struct.ident_t* @1, i32 %0, i32 4) 
label:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           18:  call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @1, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*)) 
label:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            .
            .
            .
>>>>>>

--

Would it be possible to take a look and either quickly address the issue or revert the change?

[thakis]

(Reverted for now since mac bots have been broken for 16 hours now.)

[davezarzycki]

Also, when this lands again, please ensure that all the tests don't accidentally dump build artifacts into the source directory. Specifically, the following test was missing -o -: clang/test/OpenMP/metadirective_messages.cpp

[uweigand]

Looks like this was committed again, breaking the SystemZ build bots once again:
https://lab.llvm.org/buildbot/#/builders/94/builds/5661

[cchen]

In D91944#3009003, @uweigand wrote:

Looks like this was committed again, breaking the SystemZ build bots once again:
https://lab.llvm.org/buildbot/#/builders/94/builds/5661

I'm going to fix this issue now.

[cchen]

The SystemZ issue is due to the fact that we assumed that device(cpu) should be evaluated to true and device(gpu) should be evaluated to false in the test so the test should be fixed by specifying the triple. (https://github.com/llvm/llvm-project/commit/3679d2001c87f37101e7f20c646b21e97d8a0867)

[uweigand]

In D91944#3009868, @cchen wrote:

The SystemZ issue is due to the fact that we assumed that device(cpu) should be evaluated to true and device(gpu) should be evaluated to false in the test so the test should be fixed by specifying the triple. (https://github.com/llvm/llvm-project/commit/3679d2001c87f37101e7f20c646b21e97d8a0867)

Thanks, it looks like this fixed the problem.

clang/include/clang-c/Index.h

enum CXCursorKind {
   /** OpenMP depobj directive.
    */
   CXCursor_OMPDepobjDirective = 286,
 
   /** OpenMP scan directive.
    */
   CXCursor_OMPScanDirective = 287,
 
-  CXCursor_LastStmt = CXCursor_OMPScanDirective,
+  /** OpenMP metadirective directive.
+   */
+  CXCursor_OMPMetaDirective = 288,
+
+  CXCursor_LastStmt = CXCursor_OMPMetaDirective,
 
   /**
    * Cursor that represents the translation unit itself.
    *
    * The translation unit cursor exists primarily to act as the root
    * cursor for traversing the contents of a translation unit.
    */
   CXCursor_TranslationUnit = 300,

clang/include/clang/AST/RecursiveASTVisitor.h

 }
 
 template <typename Derived>
 bool
 RecursiveASTVisitor<Derived>::TraverseOMPLoopDirective(OMPLoopDirective *S) {
   return TraverseOMPExecutableDirective(S);
 }
 
+DEF_TRAVERSE_STMT(OMPMetaDirective,
+                  { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
 DEF_TRAVERSE_STMT(OMPParallelDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPSimdDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPForDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })

clang/include/clang/AST/StmtOpenMP.h

public:
   }
   Stmt *getRawStmt() {
     assert(hasAssociatedStmt() &&
            "Expected directive with the associated statement.");
     return Data->getRawStmt();
   }
 };
 
+/// This represents '#pragma omp metadirective' directive.
+///
+/// \code
+/// #pragma omp metadirective when(user={condition(N>10)}: parallel for)
+/// \endcode
+/// In this example directive '#pragma omp metadirective' has clauses 'when'
+/// with a dynamic user condition to check if a variable 'N > 10'
+///
+class OMPMetaDirective final : public OMPExecutableDirective {
+  friend class ASTStmtReader;
+  friend class OMPExecutableDirective;
+  Stmt *IfStmt;
+
+  OMPMetaDirective(SourceLocation StartLoc, SourceLocation EndLoc)
+      : OMPExecutableDirective(OMPMetaDirectiveClass,
+                               llvm::omp::OMPD_metadirective, StartLoc,
+                               EndLoc) {}
+  explicit OMPMetaDirective()
+      : OMPExecutableDirective(OMPMetaDirectiveClass,
+                               llvm::omp::OMPD_metadirective, SourceLocation(),
+                               SourceLocation()) {}
+
+public:
+  static OMPMetaDirective *Create(const ASTContext &C, SourceLocation StartLoc,
+                                  SourceLocation EndLoc,
+                                  ArrayRef<OMPClause *> Clauses,
+                                  Stmt *AssociatedStmt, Stmt *IfStmt);
+  static OMPMetaDirective *CreateEmpty(const ASTContext &C, unsigned NumClauses,
+                                       EmptyShell);
+
+  void setIfStmt(Stmt *stmt) { IfStmt = stmt; }
+  Stmt *getIfStmt() const { return IfStmt; }
+
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OMPMetaDirectiveClass;
+  }
+};
+
 /// This represents '#pragma omp parallel' directive.
 ///
 /// \code
 /// #pragma omp parallel private(a,b) reduction(+: c,d)
 /// \endcode
 /// In this example directive '#pragma omp parallel' has clauses 'private'
 /// with the variables 'a' and 'b' and 'reduction' with operator '+' and
 /// variables 'c' and 'd'.
public:
 
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == OMPScanDirectiveClass;
   }
 };
 
 } // end namespace clang
 
 #endif

clang/include/clang/Basic/DiagnosticSemaKinds.td

def err_omp_allocator_used_in_clauses : Error<
   "allocators used in 'uses_allocators' clause cannot appear in other "
   "data-sharing or data-mapping attribute clauses">;
 def err_omp_allocator_not_in_uses_allocators : Error<
   "allocator must be specified in the 'uses_allocators' clause">;
 def note_omp_protected_structured_block
     : Note<"jump bypasses OpenMP structured block">;
 def note_omp_exits_structured_block
     : Note<"jump exits scope of OpenMP structured block">;
+def err_omp_misplaced_default_clause : Error<
+  "misplaced default clause! Only one default clause is allowed in "
+  "metadirective in the end">;
 } // end of OpenMP category
 
 let CategoryName = "Related Result Type Issue" in {
 // Objective-C related result type compatibility
 def warn_related_result_type_compatibility_class : Warning<
   "method is expected to return an instance of its class type "
   "%diff{$, but is declared to return $|"
   ", but is declared to return different type}0,1">;

clang/include/clang/Basic/StmtNodes.td

 def SEHLeaveStmt : StmtNode<Stmt>;
 def MSDependentExistsStmt : StmtNode<Stmt>;
 
 // OpenCL Extensions.
 def AsTypeExpr : StmtNode<Expr>;
 
 // OpenMP Directives.
 def OMPExecutableDirective : StmtNode<Stmt, 1>;
+def OMPMetaDirective : StmtNode<OMPExecutableDirective>;
 def OMPLoopDirective : StmtNode<OMPExecutableDirective, 1>;
 def OMPParallelDirective : StmtNode<OMPExecutableDirective>;
 def OMPSimdDirective : StmtNode<OMPLoopDirective>;
 def OMPForDirective : StmtNode<OMPLoopDirective>;
 def OMPForSimdDirective : StmtNode<OMPLoopDirective>;
 def OMPSectionsDirective : StmtNode<OMPExecutableDirective>;
 def OMPSectionDirective : StmtNode<OMPExecutableDirective>;
 def OMPSingleDirective : StmtNode<OMPExecutableDirective>;

clang/include/clang/Sema/Sema.h

public:
 
   /// Check if the current region is an OpenMP loop region and if it is,
   /// mark loop control variable, used in \p Init for loop initialization, as
   /// private by default.
   /// \param Init First part of the for loop.
   void ActOnOpenMPLoopInitialization(SourceLocation ForLoc, Stmt *Init);
 
   // OpenMP directives and clauses.
+  /// Called on well-formed '\#pragma omp metadirective' after parsing
+  /// of the  associated statement.
+  StmtResult ActOnOpenMPMetaDirective(ArrayRef<OMPClause *> Clauses,
+                                      Stmt *AStmt, SourceLocation StartLoc,
+                                      SourceLocation EndLoc);
   /// Called on correct id-expression from the '#pragma omp
   /// threadprivate'.
   ExprResult ActOnOpenMPIdExpression(Scope *CurScope, CXXScopeSpec &ScopeSpec,
                                      const DeclarationNameInfo &Id,
                                      OpenMPDirectiveKind Kind);
   /// Called on well-formed '#pragma omp threadprivate'.
   DeclGroupPtrTy ActOnOpenMPThreadprivateDirective(
                                      SourceLocation Loc,
OMPClause *ActOnOpenMPDetachClause(Expr *Evt, SourceLocation StartLoc,
                                      SourceLocation EndLoc);
 
   OMPClause *ActOnOpenMPSimpleClause(OpenMPClauseKind Kind,
                                      unsigned Argument,
                                      SourceLocation ArgumentLoc,
                                      SourceLocation StartLoc,
                                      SourceLocation LParenLoc,
                                      SourceLocation EndLoc);
+  /// Called on well-formed 'when' clause.
+  OMPClause *ActOnOpenMPWhenClause(OMPTraitInfo &TI, SourceLocation StartLoc,
+                                   SourceLocation LParenLoc,
+                                   SourceLocation EndLoc);
   /// Called on well-formed 'default' clause.
   OMPClause *ActOnOpenMPDefaultClause(llvm::omp::DefaultKind Kind,
                                       SourceLocation KindLoc,
                                       SourceLocation StartLoc,
                                       SourceLocation LParenLoc,
                                       SourceLocation EndLoc);
   /// Called on well-formed 'proc_bind' clause.
   OMPClause *ActOnOpenMPProcBindClause(llvm::omp::ProcBindKind Kind,

clang/include/clang/Serialization/ASTBitCodes.h

enum StmtCode {
       EXPR_USER_DEFINED_LITERAL,
 
       /// A CXXStdInitializerListExpr record.
       EXPR_CXX_STD_INITIALIZER_LIST,
 
       /// A CXXBoolLiteralExpr record.
       EXPR_CXX_BOOL_LITERAL,
 
-      EXPR_CXX_NULL_PTR_LITERAL,  // CXXNullPtrLiteralExpr
-      EXPR_CXX_TYPEID_EXPR,       // CXXTypeidExpr (of expr).
-      EXPR_CXX_TYPEID_TYPE,       // CXXTypeidExpr (of type).
-      EXPR_CXX_THIS,              // CXXThisExpr
-      EXPR_CXX_THROW,             // CXXThrowExpr
-      EXPR_CXX_DEFAULT_ARG,       // CXXDefaultArgExpr
-      EXPR_CXX_DEFAULT_INIT,      // CXXDefaultInitExpr
-      EXPR_CXX_BIND_TEMPORARY,    // CXXBindTemporaryExpr
+      EXPR_CXX_NULL_PTR_LITERAL, // CXXNullPtrLiteralExpr
+      EXPR_CXX_TYPEID_EXPR,      // CXXTypeidExpr (of expr).
+      EXPR_CXX_TYPEID_TYPE,      // CXXTypeidExpr (of type).
+      EXPR_CXX_THIS,             // CXXThisExpr
+      EXPR_CXX_THROW,            // CXXThrowExpr
+      EXPR_CXX_DEFAULT_ARG,      // CXXDefaultArgExpr
+      EXPR_CXX_DEFAULT_INIT,     // CXXDefaultInitExpr
+      EXPR_CXX_BIND_TEMPORARY,   // CXXBindTemporaryExpr
 
       EXPR_CXX_SCALAR_VALUE_INIT, // CXXScalarValueInitExpr
       EXPR_CXX_NEW,               // CXXNewExpr
       EXPR_CXX_DELETE,            // CXXDeleteExpr
       EXPR_CXX_PSEUDO_DESTRUCTOR, // CXXPseudoDestructorExpr
 
-      EXPR_EXPR_WITH_CLEANUPS,    // ExprWithCleanups
+      EXPR_EXPR_WITH_CLEANUPS, // ExprWithCleanups
 
       EXPR_CXX_DEPENDENT_SCOPE_MEMBER,   // CXXDependentScopeMemberExpr
       EXPR_CXX_DEPENDENT_SCOPE_DECL_REF, // DependentScopeDeclRefExpr
       EXPR_CXX_UNRESOLVED_CONSTRUCT,     // CXXUnresolvedConstructExpr
       EXPR_CXX_UNRESOLVED_MEMBER,        // UnresolvedMemberExpr
       EXPR_CXX_UNRESOLVED_LOOKUP,        // UnresolvedLookupExpr
 
-      EXPR_CXX_EXPRESSION_TRAIT,  // ExpressionTraitExpr
-      EXPR_CXX_NOEXCEPT,          // CXXNoexceptExpr
+      EXPR_CXX_EXPRESSION_TRAIT, // ExpressionTraitExpr
+      EXPR_CXX_NOEXCEPT,         // CXXNoexceptExpr
 
-      EXPR_OPAQUE_VALUE,          // OpaqueValueExpr
-      EXPR_BINARY_CONDITIONAL_OPERATOR,  // BinaryConditionalOperator
-      EXPR_TYPE_TRAIT,            // TypeTraitExpr
-      EXPR_ARRAY_TYPE_TRAIT,      // ArrayTypeTraitIntExpr
+      EXPR_OPAQUE_VALUE,                // OpaqueValueExpr
+      EXPR_BINARY_CONDITIONAL_OPERATOR, // BinaryConditionalOperator
+      EXPR_TYPE_TRAIT,                  // TypeTraitExpr
+      EXPR_ARRAY_TYPE_TRAIT,            // ArrayTypeTraitIntExpr
 
-      EXPR_PACK_EXPANSION,        // PackExpansionExpr
-      EXPR_SIZEOF_PACK,           // SizeOfPackExpr
-      EXPR_SUBST_NON_TYPE_TEMPLATE_PARM, // SubstNonTypeTemplateParmExpr
-      EXPR_SUBST_NON_TYPE_TEMPLATE_PARM_PACK,// SubstNonTypeTemplateParmPackExpr
-      EXPR_FUNCTION_PARM_PACK,    // FunctionParmPackExpr
-      EXPR_MATERIALIZE_TEMPORARY, // MaterializeTemporaryExpr
-      EXPR_CXX_FOLD,              // CXXFoldExpr
-      EXPR_CONCEPT_SPECIALIZATION,// ConceptSpecializationExpr
-      EXPR_REQUIRES,              // RequiresExpr
+      EXPR_PACK_EXPANSION,                    // PackExpansionExpr
+      EXPR_SIZEOF_PACK,                       // SizeOfPackExpr
+      EXPR_SUBST_NON_TYPE_TEMPLATE_PARM,      // SubstNonTypeTemplateParmExpr
+      EXPR_SUBST_NON_TYPE_TEMPLATE_PARM_PACK, // SubstNonTypeTemplateParmPackExpr
+      EXPR_FUNCTION_PARM_PACK,                // FunctionParmPackExpr
+      EXPR_MATERIALIZE_TEMPORARY,             // MaterializeTemporaryExpr
+      EXPR_CXX_FOLD,                          // CXXFoldExpr
+      EXPR_CONCEPT_SPECIALIZATION,            // ConceptSpecializationExpr
+      EXPR_REQUIRES,                          // RequiresExpr
 
       // CUDA
-      EXPR_CUDA_KERNEL_CALL,       // CUDAKernelCallExpr
+      EXPR_CUDA_KERNEL_CALL, // CUDAKernelCallExpr
 
       // OpenCL
-      EXPR_ASTYPE,                 // AsTypeExpr
+      EXPR_ASTYPE, // AsTypeExpr
 
       // Microsoft
-      EXPR_CXX_PROPERTY_REF_EXPR, // MSPropertyRefExpr
+      EXPR_CXX_PROPERTY_REF_EXPR,       // MSPropertyRefExpr
       EXPR_CXX_PROPERTY_SUBSCRIPT_EXPR, // MSPropertySubscriptExpr
-      EXPR_CXX_UUIDOF_EXPR,       // CXXUuidofExpr (of expr).
-      EXPR_CXX_UUIDOF_TYPE,       // CXXUuidofExpr (of type).
-      STMT_SEH_LEAVE,             // SEHLeaveStmt
-      STMT_SEH_EXCEPT,            // SEHExceptStmt
-      STMT_SEH_FINALLY,           // SEHFinallyStmt
-      STMT_SEH_TRY,               // SEHTryStmt
+      EXPR_CXX_UUIDOF_EXPR,             // CXXUuidofExpr (of expr).
+      EXPR_CXX_UUIDOF_TYPE,             // CXXUuidofExpr (of type).
+      STMT_SEH_LEAVE,                   // SEHLeaveStmt
+      STMT_SEH_EXCEPT,                  // SEHExceptStmt
+      STMT_SEH_FINALLY,                 // SEHFinallyStmt
+      STMT_SEH_TRY,                     // SEHTryStmt
 
       // OpenMP directives
+      STMT_OMP_META_DIRECTIVE,
       STMT_OMP_PARALLEL_DIRECTIVE,
       STMT_OMP_SIMD_DIRECTIVE,
       STMT_OMP_FOR_DIRECTIVE,
       STMT_OMP_FOR_SIMD_DIRECTIVE,
       STMT_OMP_SECTIONS_DIRECTIVE,
       STMT_OMP_SECTION_DIRECTIVE,
       STMT_OMP_SINGLE_DIRECTIVE,
       STMT_OMP_MASTER_DIRECTIVE,
enum StmtCode {
       STMT_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_FOR_DIRECTIVE,
       STMT_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_FOR_SIMD_DIRECTIVE,
       STMT_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD_DIRECTIVE,
       EXPR_OMP_ARRAY_SECTION,
       EXPR_OMP_ARRAY_SHAPING,
       EXPR_OMP_ITERATOR,
 
       // ARC
-      EXPR_OBJC_BRIDGED_CAST,     // ObjCBridgedCastExpr
+      EXPR_OBJC_BRIDGED_CAST, // ObjCBridgedCastExpr
 
-      STMT_MS_DEPENDENT_EXISTS,   // MSDependentExistsStmt
-      EXPR_LAMBDA,                // LambdaExpr
+      STMT_MS_DEPENDENT_EXISTS, // MSDependentExistsStmt
+      EXPR_LAMBDA,              // LambdaExpr
       STMT_COROUTINE_BODY,
       STMT_CORETURN,
       EXPR_COAWAIT,
       EXPR_COYIELD,
       EXPR_DEPENDENT_COAWAIT,
 
       // FixedPointLiteral
       EXPR_FIXEDPOINT_LITERAL,

clang/lib/AST/OpenMPClause.cpp

const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
     break;
   default:
     break;
   }
 
   return nullptr;
 }
 
const OMPClauseWithPostUpdate *OMPClauseWithPostUpdate::get(const OMPClause *C) {
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
     break;
   default:
     break;
   }
 
   return nullptr;
 }
 

clang/lib/AST/StmtOpenMP.cpp

 }
 
 void OMPLoopDirective::setFinalsConditions(ArrayRef<Expr *> A) {
   assert(A.size() == getCollapsedNumber() &&
          "Number of finals conditions is not the same as the collapsed number");
   llvm::copy(A, getFinalsConditions().begin());
 }
 
+OMPMetaDirective *OMPMetaDirective::Create(const ASTContext &C,
+                                           SourceLocation StartLoc,
+                                           SourceLocation EndLoc,
+                                           ArrayRef<OMPClause *> Clauses,
+                                           Stmt *AssociatedStmt, Stmt *IfStmt) {
+  auto *Dir = createDirective<OMPMetaDirective>(
+      C, Clauses, AssociatedStmt, /*NumChildren=*/1, StartLoc, EndLoc);
+  Dir->setIfStmt(IfStmt);
+  return Dir;
+}
+
+OMPMetaDirective *OMPMetaDirective::CreateEmpty(const ASTContext &C,
+                                                unsigned NumClauses,
+                                                EmptyShell) {
+  return createEmptyDirective<OMPMetaDirective>(C, NumClauses,
+                                                /*HasAssociatedStmt=*/true,
+                                                /*NumChildren=*/1);
+}
+
 OMPParallelDirective *OMPParallelDirective::Create(
     const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
     ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt, Expr *TaskRedRef,
     bool HasCancel) {
   auto *Dir = createDirective<OMPParallelDirective>(
       C, Clauses, AssociatedStmt, /*NumChildren=*/1, StartLoc, EndLoc);
   Dir->setTaskReductionRefExpr(TaskRedRef);
   Dir->setHasCancel(HasCancel);

clang/lib/AST/StmtPrinter.cpp

if (Clause && !Clause->isImplicit()) {
       OS << ' ';
       Printer.Visit(Clause);
     }
   OS << NL;
   if (!ForceNoStmt && S->hasAssociatedStmt())
     PrintStmt(S->getRawStmt());
 }
 
+void StmtPrinter::VisitOMPMetaDirective(OMPMetaDirective *Node) {
+  Indent() << "#pragma omp metadirective";
+  PrintOMPExecutableDirective(Node);
+}
+
 void StmtPrinter::VisitOMPParallelDirective(OMPParallelDirective *Node) {
   Indent() << "#pragma omp parallel";
   PrintOMPExecutableDirective(Node);
 }
 
 void StmtPrinter::VisitOMPSimdDirective(OMPSimdDirective *Node) {
   Indent() << "#pragma omp simd";
   PrintOMPExecutableDirective(Node);

clang/lib/AST/StmtProfile.cpp

StmtProfiler::VisitOMPExecutableDirective(const OMPExecutableDirective *S) {
   OMPClauseProfiler P(this);
   ArrayRef<OMPClause *> Clauses = S->clauses();
   for (ArrayRef<OMPClause *>::iterator I = Clauses.begin(), E = Clauses.end();
        I != E; ++I)
     if (*I)
       P.Visit(*I);
 }
 
+void StmtProfiler::VisitOMPMetaDirective(const OMPMetaDirective *S) {
+  VisitOMPExecutableDirective(S);
+}
+
 void StmtProfiler::VisitOMPLoopDirective(const OMPLoopDirective *S) {
   VisitOMPExecutableDirective(S);
 }
 
 void StmtProfiler::VisitOMPParallelDirective(const OMPParallelDirective *S) {
   VisitOMPExecutableDirective(S);
 }
 

clang/lib/Basic/OpenMPKinds.cpp

#include "clang/Basic/OpenMPKinds.def"
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
     break;
   default:
     break;
   }
   llvm_unreachable("Invalid OpenMP simple clause kind");
 }
 
 const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind,
#include "clang/Basic/OpenMPKinds.def"
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
     break;
   default:
     break;
   }
   llvm_unreachable("Invalid OpenMP simple clause kind");
 }
 
 bool clang::isOpenMPLoopDirective(OpenMPDirectiveKind DKind) {
return Kind == OMPD_distribute_parallel_for ||
          Kind == OMPD_target_teams_distribute_parallel_for_simd;
 }
 
 void clang::getOpenMPCaptureRegions(
     SmallVectorImpl<OpenMPDirectiveKind> &CaptureRegions,
     OpenMPDirectiveKind DKind) {
   assert(unsigned(DKind) < llvm::omp::Directive_enumSize);
   switch (DKind) {
+  case OMPD_metadirective:
+    CaptureRegions.push_back(OMPD_metadirective);
+    break;
   case OMPD_parallel:
   case OMPD_parallel_for:
   case OMPD_parallel_for_simd:
   case OMPD_parallel_master:
   case OMPD_parallel_sections:
   case OMPD_distribute_parallel_for:
   case OMPD_distribute_parallel_for_simd:
     CaptureRegions.push_back(OMPD_parallel);

clang/lib/CodeGen/CGOpenMPRuntime.cpp

emitNumThreadsForTargetDirective(CodeGenFunction &CGF,
   case OMPD_taskloop:
   case OMPD_taskloop_simd:
   case OMPD_master_taskloop:
   case OMPD_master_taskloop_simd:
   case OMPD_parallel_master_taskloop:
   case OMPD_parallel_master_taskloop_simd:
   case OMPD_requires:
   case OMPD_unknown:
+  case OMPD_metadirective:
     break;
   default:
     break;
   }
   llvm_unreachable("Unsupported directive kind.");
 }
 
 namespace {
if (const auto *NestedDir =
     case OMPD_declare_mapper:
     case OMPD_taskloop:
     case OMPD_taskloop_simd:
     case OMPD_master_taskloop:
     case OMPD_master_taskloop_simd:
     case OMPD_parallel_master_taskloop:
     case OMPD_parallel_master_taskloop_simd:
     case OMPD_requires:
+    case OMPD_metadirective:
     case OMPD_unknown:
     default:
       llvm_unreachable("Unexpected directive.");
     }
   }
 
   return nullptr;
 }
void CGOpenMPRuntime::emitTargetCall(
   }
 }
 
 void CGOpenMPRuntime::scanForTargetRegionsFunctions(const Stmt *S,
                                                     StringRef ParentName) {
   if (!S)
     return;
 
+  if (isa<OMPMetaDirective>(S)) {
+    const auto &M = *cast<OMPMetaDirective>(S);
+    scanForTargetRegionsFunctions(M.getIfStmt(), ParentName);
+    return;
+  }
+
   // Codegen OMP target directives that offload compute to the device.
   bool RequiresDeviceCodegen =
       isa<OMPExecutableDirective>(S) &&
       isOpenMPTargetExecutionDirective(
           cast<OMPExecutableDirective>(S)->getDirectiveKind());
 
   if (RequiresDeviceCodegen) {
     const auto &E = *cast<OMPExecutableDirective>(S);
auto &&ThenGen = [this, &D, Device, &InputInfo, &MapTypesArray,
     case OMPD_target_teams_distribute_simd:
     case OMPD_target_teams_distribute_parallel_for:
     case OMPD_target_teams_distribute_parallel_for_simd:
     case OMPD_target_teams:
     case OMPD_target_parallel:
     case OMPD_target_parallel_for:
     case OMPD_target_parallel_for_simd:
     case OMPD_requires:
+    case OMPD_metadirective:
     case OMPD_unknown:
     default:
       llvm_unreachable("Unexpected standalone target data directive.");
       break;
     }
     CGF.EmitRuntimeCall(
         OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(), RTLFn),
         OffloadingArgs);

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

if (const auto *NestedDir =
     case OMPD_declare_mapper:
     case OMPD_taskloop:
     case OMPD_taskloop_simd:
     case OMPD_master_taskloop:
     case OMPD_master_taskloop_simd:
     case OMPD_parallel_master_taskloop:
     case OMPD_parallel_master_taskloop_simd:
     case OMPD_requires:
+    case OMPD_metadirective:
     case OMPD_unknown:
     default:
       llvm_unreachable("Unexpected directive.");
     }
   }
 
   return false;
 }
static bool supportsSPMDExecutionMode(ASTContext &Ctx,
   case OMPD_declare_mapper:
   case OMPD_taskloop:
   case OMPD_taskloop_simd:
   case OMPD_master_taskloop:
   case OMPD_master_taskloop_simd:
   case OMPD_parallel_master_taskloop:
   case OMPD_parallel_master_taskloop_simd:
   case OMPD_requires:
+  case OMPD_metadirective:
   case OMPD_unknown:
   default:
     break;
   }
   llvm_unreachable(
       "Unknown programming model for OpenMP directive on NVPTX target.");
 }
 
if (const auto *NestedDir =
     case OMPD_declare_mapper:
     case OMPD_taskloop:
     case OMPD_taskloop_simd:
     case OMPD_master_taskloop:
     case OMPD_master_taskloop_simd:
     case OMPD_parallel_master_taskloop:
     case OMPD_parallel_master_taskloop_simd:
     case OMPD_requires:
+    case OMPD_metadirective:
     case OMPD_unknown:
     default:
       llvm_unreachable("Unexpected directive.");
     }
   }
 
   return false;
 }
static bool supportsLightweightRuntime(ASTContext &Ctx,
   case OMPD_declare_mapper:
   case OMPD_taskloop:
   case OMPD_taskloop_simd:
   case OMPD_master_taskloop:
   case OMPD_master_taskloop_simd:
   case OMPD_parallel_master_taskloop:
   case OMPD_parallel_master_taskloop_simd:
   case OMPD_requires:
+  case OMPD_metadirective:
   case OMPD_unknown:
   default:
     break;
   }
   llvm_unreachable(
       "Unknown programming model for OpenMP directive on NVPTX target.");
 }
 

clang/lib/CodeGen/CGStmt.cpp

case Stmt::CXXTryStmtClass:
     EmitCXXTryStmt(cast<CXXTryStmt>(*S));
     break;
   case Stmt::CXXForRangeStmtClass:
     EmitCXXForRangeStmt(cast<CXXForRangeStmt>(*S), Attrs);
     break;
   case Stmt::SEHTryStmtClass:
     EmitSEHTryStmt(cast<SEHTryStmt>(*S));
     break;
+  case Stmt::OMPMetaDirectiveClass:
+    EmitOMPMetaDirective(cast<OMPMetaDirective>(*S));
+    break;
   case Stmt::OMPParallelDirectiveClass:
     EmitOMPParallelDirective(cast<OMPParallelDirective>(*S));
     break;
   case Stmt::OMPSimdDirectiveClass:
     EmitOMPSimdDirective(cast<OMPSimdDirective>(*S));
     break;
   case Stmt::OMPForDirectiveClass:
     EmitOMPForDirective(cast<OMPForDirective>(*S));

clang/lib/CodeGen/CGStmtOpenMP.cpp

static void emitOMPAtomicExpr(CodeGenFunction &CGF, OpenMPClauseKind Kind,
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
   default:
     llvm_unreachable("Clause is not allowed in 'omp atomic'.");
   }
 }
 
 void CodeGenFunction::EmitOMPAtomicDirective(const OMPAtomicDirective &S) {
   llvm::AtomicOrdering AO = llvm::AtomicOrdering::Monotonic;
   bool MemOrderingSpecified = false;
auto &&CodeGen = [this, &S](CodeGenFunction &CGF, PrePostActionTy &Action) {
     EmitOMPTaskLoopBasedDirective(S);
   };
   auto LPCRegion =
       CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S);
   OMPLexicalScope Scope(*this, S);
   CGM.getOpenMPRuntime().emitMasterRegion(*this, CodeGen, S.getBeginLoc());
 }
 
+void CodeGenFunction::EmitOMPMetaDirective(const OMPMetaDirective &S) {
+  EmitStmt(S.getIfStmt());
+}
+
 void CodeGenFunction::EmitOMPParallelMasterTaskLoopDirective(
     const OMPParallelMasterTaskLoopDirective &S) {
   auto &&CodeGen = [this, &S](CodeGenFunction &CGF, PrePostActionTy &Action) {
     auto &&TaskLoopCodeGen = [&S](CodeGenFunction &CGF,
                                   PrePostActionTy &Action) {
       Action.Enter(CGF);
       CGF.EmitOMPTaskLoopBasedDirective(S);
     };

clang/lib/CodeGen/CodeGenFunction.h

OMPTargetDataInfo(Address BasePointersArray, Address PointersArray,
         : BasePointersArray(BasePointersArray), PointersArray(PointersArray),
           SizesArray(SizesArray), MappersArray(MappersArray),
           NumberOfTargetItems(NumberOfTargetItems) {}
   };
   void EmitOMPTargetTaskBasedDirective(const OMPExecutableDirective &S,
                                        const RegionCodeGenTy &BodyGen,
                                        OMPTargetDataInfo &InputInfo);
 
+  void EmitOMPMetaDirective(const OMPMetaDirective &S);
   void EmitOMPParallelDirective(const OMPParallelDirective &S);
   void EmitOMPSimdDirective(const OMPSimdDirective &S);
   void EmitOMPForDirective(const OMPForDirective &S);
   void EmitOMPForSimdDirective(const OMPForSimdDirective &S);
   void EmitOMPSectionsDirective(const OMPSectionsDirective &S);
   void EmitOMPSectionDirective(const OMPSectionDirective &S);
   void EmitOMPSingleDirective(const OMPSingleDirective &S);
   void EmitOMPMasterDirective(const OMPMasterDirective &S);

clang/lib/Parse/ParseOpenMP.cpp

Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
   case OMPD_teams_distribute_simd:
   case OMPD_teams_distribute_parallel_for_simd:
   case OMPD_teams_distribute_parallel_for:
   case OMPD_target_teams:
   case OMPD_target_teams_distribute:
   case OMPD_target_teams_distribute_parallel_for:
   case OMPD_target_teams_distribute_parallel_for_simd:
   case OMPD_target_teams_distribute_simd:
+  case OMPD_metadirective:
     Diag(Tok, diag::err_omp_unexpected_directive)
         << 1 << getOpenMPDirectiveName(DKind);
     break;
   default:
     break;
   }
   while (Tok.isNot(tok::annot_pragma_openmp_end))
     ConsumeAnyToken();
Parser::ParseOpenMPDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx) {
   unsigned ScopeFlags = Scope::FnScope | Scope::DeclScope |
                         Scope::CompoundStmtScope | Scope::OpenMPDirectiveScope;
   SourceLocation Loc = ConsumeAnnotationToken(), EndLoc;
   OpenMPDirectiveKind DKind = parseOpenMPDirectiveKind(*this);
   OpenMPDirectiveKind CancelRegion = OMPD_unknown;
   // Name of critical directive.
   DeclarationNameInfo DirName;
   StmtResult Directive = StmtError();
+
   bool HasAssociatedStatement = true;
 
   switch (DKind) {
+  case OMPD_metadirective: {
+    ConsumeToken();
+    SmallVector<VariantMatchInfo, 4> VMIs;
+    OpenMPDirectiveKind DirKind = OMPD_unknown;
+    StmtResult AssociatedStmt;
+    TentativeParsingAction TPA(*this);
+    ASTContext &ASTContext = Actions.getASTContext();
+
+    BalancedDelimiterTracker T(*this, tok::l_paren,
+                               tok::annot_pragma_openmp_end);
+    while (Tok.isNot(tok::annot_pragma_openmp_end)) {
+      OpenMPClauseKind CKind = Tok.isAnnotation()
+                                   ? OMPC_unknown
+                                   : getOpenMPClauseKind(PP.getSpelling(Tok));
+      SourceLocation Loc = ConsumeToken();
+
+      // Parse '('.
+      if (T.expectAndConsume(diag::err_expected_lparen_after,
+                             getOpenMPClauseName(CKind).data()))
+        return Directive;
+
+      OMPTraitInfo &TI = Actions.getASTContext().getNewOMPTraitInfo();
+      if (CKind == OMPC_when) {
+        // parse and get condition expression to pass to the When clause
+        parseOMPContextSelectors(Loc, TI);
+
+        // Parse ':'
+        if (Tok.is(tok::colon))
+          ConsumeAnyToken();
+        else {
+          Diag(Tok, diag::warn_pragma_expected_colon) << "when clause";
+          return Directive;
+        }
+      }
+      // Skip Directive
+      int paren = 0;
+      while (Tok.isNot(tok::r_paren) || paren != 0) {
+        if (Tok.is(tok::l_paren))
+          paren++;
+        if (Tok.is(tok::r_paren))
+          paren--;
+        ConsumeAnyToken();
+      }
+      // Parse ')'
+      if (Tok.is(tok::r_paren))
+        T.consumeClose();
+
+      VariantMatchInfo VMI;
+      TI.getAsVariantMatchInfo(ASTContext, VMI);
+
+      VMIs.push_back(VMI);
+    }
+
+    TPA.Revert();
+    TargetOMPContext OMPCtx(ASTContext, nullptr, nullptr);
+    int BestIdx = getBestWhenMatchForContext(VMIs, OMPCtx);
+
+    int idx = 0;
+    while (Tok.isNot(tok::annot_pragma_openmp_end)) {
+      if (idx++ != BestIdx) { // Ignore this clause
+        ConsumeToken();       // Consume clause name
+        T.consumeOpen();      // Consume '('
+        int paren = 0;
+        // Skip everything inside the clause
+        while (Tok.isNot(tok::r_paren) || paren != 0) {
+          if (Tok.is(tok::l_paren))
+            paren++;
+          if (Tok.is(tok::r_paren))
+            paren--;
+          ConsumeAnyToken();
+        }
+        // Parse ')'
+        if (Tok.is(tok::r_paren))
+          T.consumeClose();
+        continue;
+      }
+
+      OpenMPClauseKind CKind = Tok.isAnnotation()
+                                   ? OMPC_unknown
+                                   : getOpenMPClauseKind(PP.getSpelling(Tok));
+      SourceLocation Loc = ConsumeToken();
+
+      // Parse '('.
+      T.consumeOpen();
+
+      // Skip ContextSelectors for when clause
+      if (CKind == OMPC_when) {
+        OMPTraitInfo &TI = Actions.getASTContext().getNewOMPTraitInfo();
+        // parse and skip ContextSelectors
+        parseOMPContextSelectors(Loc, TI);
+
+        // Parse ':'
+        ConsumeAnyToken();
+      }
+
+      // Parse Directive
+      if (Tok.isNot(tok::r_paren)) {
+        DirKind = parseOpenMPDirectiveKind(*this);
+        ConsumeToken();
+        if (DirKind != OMPD_unknown) {
+          ParseScope OMPDirectiveScope(this, ScopeFlags);
+          Actions.StartOpenMPDSABlock(DirKind, DirName, Actions.getCurScope(),
+                                      Loc);
+          int paren = 0;
+          while (Tok.isNot(tok::r_paren) || paren != 0) {
+            if (Tok.is(tok::l_paren))
+              paren++;
+            if (Tok.is(tok::r_paren))
+              paren--;
+            OpenMPClauseKind CKind =
+                Tok.isAnnotation() ? OMPC_unknown
+                                   : getOpenMPClauseKind(PP.getSpelling(Tok));
+            Actions.StartOpenMPClause(CKind);
+            OMPClause *DClause = ParseOpenMPClause(
+                DirKind, CKind, !FirstClauses[(unsigned)CKind].getInt());
+            FirstClauses[(unsigned)CKind].setInt(true);
+            if (DClause) {
+              FirstClauses[(unsigned)CKind].setPointer(DClause);
+              Clauses.push_back(DClause);
+            }
+            // Skip ',' if any.
+            if (Tok.is(tok::comma))
+              ConsumeToken();
+            Actions.EndOpenMPClause();
+          }
+          while (Tok.isNot(tok::annot_pragma_openmp_end))
+            ConsumeAnyToken();
+          // End location of the directive.
+          EndLoc = Tok.getLocation();
+          // Consume final annot_pragma_openmp_end.
+          ConsumeAnnotationToken();
+
+          Actions.ActOnOpenMPRegionStart(DirKind, getCurScope());
+          ParsingOpenMPDirectiveRAII NormalScope(*this, /*Value=*/false);
+          ParseScope InnerStmtScope(this, Scope::DeclScope,
+                                    getLangOpts().C99 ||
+                                        getLangOpts().CPlusPlus,
+                                    Tok.is(tok::l_brace));
+          StmtResult AStmt = ParseStatement();
+          InnerStmtScope.Exit();
+          AssociatedStmt = (Sema::CompoundScopeRAII(Actions), AStmt);
+          AssociatedStmt =
+              Actions.ActOnOpenMPRegionEnd(AssociatedStmt, Clauses);
+          Directive = Actions.ActOnOpenMPExecutableDirective(
+              DirKind, DirName, CancelRegion, Clauses, AssociatedStmt.get(),
+              Loc, EndLoc);
+          // Exit scope.
+          Actions.EndOpenMPDSABlock(Directive.get());
+          OMPDirectiveScope.Exit();
+        } else {
+          // Unknown Directive encountered. Skip directive.
+          Diag(Tok, diag::err_omp_unknown_directive);
+          while (Tok.isNot(tok::annot_pragma_openmp_end))
+            ConsumeAnyToken();
+          ConsumeAnnotationToken();
+        }
+      } else {
+        // Empty Directive. Skip in OpenMP 5.0.
+        // Generate nothing directive from OpenMP 5.1.
+        while (Tok.isNot(tok::annot_pragma_openmp_end))
+          ConsumeAnyToken();
+        ConsumeAnnotationToken();
+      }
+      break;
+    }
+    break;
+  }
   case OMPD_threadprivate: {
     // FIXME: Should this be permitted in C++?
     if ((StmtCtx & ParsedStmtContext::AllowDeclarationsInC) ==
         ParsedStmtContext()) {
       Diag(Tok, diag::err_omp_immediate_directive)
           << getOpenMPDirectiveName(DKind) << 0;
     }
     ConsumeToken();
case OMPC_uses_allocators:
     break;
   case OMPC_device_type:
   case OMPC_unknown:
     skipUntilPragmaOpenMPEnd(DKind);
     break;
   case OMPC_threadprivate:
   case OMPC_uniform:
   case OMPC_match:
+  case OMPC_when:
     if (!WrongDirective)
       Diag(Tok, diag::err_omp_unexpected_clause)
           << getOpenMPClauseName(CKind) << getOpenMPDirectiveName(DKind);
     SkipUntil(tok::comma, tok::annot_pragma_openmp_end, StopBeforeMatch);
     break;
   default:
     break;
   }

clang/lib/Sema/SemaExceptionSpec.cpp

CanThrowResult Sema::canThrow(const Stmt *S) {
   case Stmt::OMPTaskLoopSimdDirectiveClass:
   case Stmt::OMPTaskwaitDirectiveClass:
   case Stmt::OMPTaskyieldDirectiveClass:
   case Stmt::OMPTeamsDirectiveClass:
   case Stmt::OMPTeamsDistributeDirectiveClass:
   case Stmt::OMPTeamsDistributeParallelForDirectiveClass:
   case Stmt::OMPTeamsDistributeParallelForSimdDirectiveClass:
   case Stmt::OMPTeamsDistributeSimdDirectiveClass:
+  case Stmt::OMPMetaDirectiveClass:
   case Stmt::ReturnStmtClass:
   case Stmt::SEHExceptStmtClass:
   case Stmt::SEHFinallyStmtClass:
   case Stmt::SEHLeaveStmtClass:
   case Stmt::SEHTryStmtClass:
   case Stmt::SwitchStmtClass:
   case Stmt::WhileStmtClass:
     return canSubStmtsThrow(*this, S);

clang/lib/Sema/SemaOpenMP.cpp

if (isOpenMPTargetExecutionDirective(S->getCurrentDirective())) {
         Visit(E);
     }
   }
 };
 } // namespace
 
 void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
   switch (DKind) {
+  case OMPD_metadirective:
   case OMPD_parallel:
   case OMPD_parallel_for:
   case OMPD_parallel_for_simd:
   case OMPD_parallel_sections:
   case OMPD_parallel_master:
   case OMPD_teams:
   case OMPD_teams_distribute:
   case OMPD_teams_distribute_simd: {
case OMPD_target_teams_distribute_parallel_for_simd:
     break;
   case OMPD_target_teams_distribute_simd:
     Res = ActOnOpenMPTargetTeamsDistributeSimdDirective(
         ClausesWithImplicit, AStmt, StartLoc, EndLoc, VarsWithInheritedDSA);
     AllowedNameModifiers.push_back(OMPD_target);
     if (LangOpts.OpenMP >= 50)
       AllowedNameModifiers.push_back(OMPD_simd);
     break;
+  case OMPD_metadirective:
   case OMPD_declare_target:
   case OMPD_end_declare_target:
   case OMPD_threadprivate:
   case OMPD_allocate:
   case OMPD_declare_reduction:
   case OMPD_declare_mapper:
   case OMPD_declare_simd:
   case OMPD_requires:
for (OMPClause *C : Clauses) {
       case OMPC_unknown:
       case OMPC_unified_address:
       case OMPC_unified_shared_memory:
       case OMPC_reverse_offload:
       case OMPC_dynamic_allocators:
       case OMPC_atomic_default_mem_order:
       case OMPC_device_type:
       case OMPC_match:
+      case OMPC_when:
       default:
         llvm_unreachable("Unexpected clause");
       }
       for (Stmt *CC : C->children()) {
         if (CC)
           DSAChecker.Visit(CC);
       }
     }
OMPClause *Sema::ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind, Expr *Expr,
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_destroy:
   case OMPC_inclusive:
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
   default:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 // An OpenMP directive such as 'target parallel' has two captured regions:
 // for the 'target' and 'parallel' respectively.  This function returns
case OMPC_if:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_teams_distribute:
     case OMPD_requires:
+    case OMPD_metadirective:
       llvm_unreachable("Unexpected OpenMP directive with if-clause");
     case OMPD_unknown:
     default:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_num_threads:
     switch (DKind) {
case OMPC_num_threads:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_distribute_simd:
     case OMPD_teams_distribute:
     case OMPD_teams_distribute_simd:
     case OMPD_requires:
+    case OMPD_metadirective:
       llvm_unreachable("Unexpected OpenMP directive with num_threads-clause");
     case OMPD_unknown:
     default:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_num_teams:
     switch (DKind) {
case OMPC_num_teams:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_distribute_simd:
     case OMPD_requires:
+    case OMPD_metadirective:
       llvm_unreachable("Unexpected OpenMP directive with num_teams-clause");
     case OMPD_unknown:
     default:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_thread_limit:
     switch (DKind) {
case OMPC_thread_limit:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_distribute_simd:
     case OMPD_requires:
+    case OMPD_metadirective:
       llvm_unreachable("Unexpected OpenMP directive with thread_limit-clause");
     case OMPD_unknown:
     default:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_schedule:
     switch (DKind) {
case OMPC_schedule:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_distribute_simd:
     case OMPD_target_teams:
     case OMPD_requires:
+    case OMPD_metadirective:
       llvm_unreachable("Unexpected OpenMP directive with schedule clause");
     case OMPD_unknown:
     default:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_dist_schedule:
     switch (DKind) {
case OMPC_dist_schedule:
     case OMPD_single:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_target_teams:
     case OMPD_requires:
+    case OMPD_metadirective:
       llvm_unreachable("Unexpected OpenMP directive with schedule clause");
     case OMPD_unknown:
     default:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_device:
     switch (DKind) {
case OMPC_device:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_distribute_simd:
     case OMPD_requires:
+    case OMPD_metadirective:
       llvm_unreachable("Unexpected OpenMP directive with num_teams-clause");
     case OMPD_unknown:
     default:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_grainsize:
   case OMPC_num_tasks:
case OMPC_priority:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_distribute_simd:
     case OMPD_requires:
+    case OMPD_metadirective:
       llvm_unreachable("Unexpected OpenMP directive with grainsize-clause");
     case OMPD_unknown:
     default:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
+  case OMPC_when:
+    if (DKind == OMPD_metadirective) {
+      CaptureRegion = OMPD_metadirective;
+    } else if (DKind == OMPD_unknown) {
+      llvm_unreachable("Unknown OpenMP directive");
+    } else {
+      llvm_unreachable("Unexpected OpenMP directive with when clause");
+    }
+    break;
   case OMPC_firstprivate:
   case OMPC_lastprivate:
   case OMPC_reduction:
   case OMPC_task_reduction:
   case OMPC_in_reduction:
   case OMPC_linear:
   case OMPC_default:
   case OMPC_proc_bind:
OMPClause *Sema::ActOnOpenMPSimpleClause(
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
   default:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 static std::string
 getListOfPossibleValues(OpenMPClauseKind K, unsigned First, unsigned Last,
OMPClause *Sema::ActOnOpenMPSingleExprWithArgClause(
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
   default:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 static bool checkScheduleModifiers(Sema &S, OpenMPScheduleClauseModifier M1,
                                    OpenMPScheduleClauseModifier M2,
OMPClause *Sema::ActOnOpenMPClause(OpenMPClauseKind Kind,
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
   default:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 OMPClause *Sema::ActOnOpenMPNowaitClause(SourceLocation StartLoc,
                                          SourceLocation EndLoc) {
OMPClause *Sema::ActOnOpenMPVarListClause(
   case OMPC_dynamic_allocators:
   case OMPC_atomic_default_mem_order:
   case OMPC_device_type:
   case OMPC_match:
   case OMPC_order:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_uses_allocators:
+  case OMPC_when:
   default:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 ExprResult Sema::getOpenMPCapturedExpr(VarDecl *Capture, ExprValueKind VK,
                                        ExprObjectKind OK, SourceLocation Loc) {

clang/lib/Sema/TreeTransform.h

StmtResult TreeTransform<Derived>::TransformOMPExecutableDirective(
 
   return getDerived().RebuildOMPExecutableDirective(
       D->getDirectiveKind(), DirName, CancelRegion, TClauses,
       AssociatedStmt.get(), D->getBeginLoc(), D->getEndLoc());
 }
 
 template <typename Derived>
 StmtResult
+TreeTransform<Derived>::TransformOMPMetaDirective(OMPMetaDirective *D) {
+  DeclarationNameInfo DirName;
+  getDerived().getSema().StartOpenMPDSABlock(OMPD_metadirective, DirName,
+                                             nullptr, D->getBeginLoc());
+  StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
+  getDerived().getSema().EndOpenMPDSABlock(Res.get());
+  return Res;
+}
+
+template <typename Derived>
+StmtResult
 TreeTransform<Derived>::TransformOMPParallelDirective(OMPParallelDirective *D) {
   DeclarationNameInfo DirName;
   getDerived().getSema().StartOpenMPDSABlock(OMPD_parallel, DirName, nullptr,
                                              D->getBeginLoc());
   StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
   getDerived().getSema().EndOpenMPDSABlock(Res.get());
   return Res;
 }

clang/lib/Serialization/ASTReaderStmt.cpp

 
 void ASTStmtReader::VisitOMPLoopDirective(OMPLoopDirective *D) {
   VisitStmt(D);
   // Field CollapsedNum was read in ReadStmtFromStream.
   Record.skipInts(1);
   VisitOMPExecutableDirective(D);
 }
 
+void ASTStmtReader::VisitOMPMetaDirective(OMPMetaDirective *D) {
+  VisitStmt(D);
+  // The NumClauses field was read in ReadStmtFromStream.
+  Record.skipInts(1);
+  VisitOMPExecutableDirective(D);
+}
+
 void ASTStmtReader::VisitOMPParallelDirective(OMPParallelDirective *D) {
   VisitStmt(D);
   VisitOMPExecutableDirective(D);
   D->setHasCancel(Record.readBool());
 }
 
 void ASTStmtReader::VisitOMPSimdDirective(OMPSimdDirective *D) {
   VisitOMPLoopDirective(D);
while (true) {
 
     case STMT_MS_DEPENDENT_EXISTS:
       S = new (Context) MSDependentExistsStmt(SourceLocation(), true,
                                               NestedNameSpecifierLoc(),
                                               DeclarationNameInfo(),
                                               nullptr);
       break;
 
+    case STMT_OMP_META_DIRECTIVE:
+      S = OMPMetaDirective::CreateEmpty(
+          Context, Record[ASTStmtReader::NumStmtFields], Empty);
+      break;
+
     case STMT_OMP_PARALLEL_DIRECTIVE:
       S =
         OMPParallelDirective::CreateEmpty(Context,
                                           Record[ASTStmtReader::NumStmtFields],
                                           Empty);
       break;
 
     case STMT_OMP_SIMD_DIRECTIVE: {

clang/lib/Serialization/ASTWriterStmt.cpp

 //===----------------------------------------------------------------------===//
 
 void ASTStmtWriter::VisitOMPExecutableDirective(OMPExecutableDirective *E) {
   Record.writeOMPChildren(E->Data);
   Record.AddSourceLocation(E->getBeginLoc());
   Record.AddSourceLocation(E->getEndLoc());
 }
 
+void ASTStmtWriter::VisitOMPMetaDirective(OMPMetaDirective *D) {
+  VisitStmt(D);
+  Record.push_back(D->getNumClauses());
+  VisitOMPExecutableDirective(D);
+  Code = serialization::STMT_OMP_META_DIRECTIVE;
+}
+
 void ASTStmtWriter::VisitOMPLoopDirective(OMPLoopDirective *D) {
   VisitStmt(D);
   Record.writeUInt32(D->getCollapsedNumber());
   VisitOMPExecutableDirective(D);
 }
 
 void ASTStmtWriter::VisitOMPParallelDirective(OMPParallelDirective *D) {
   VisitStmt(D);

clang/lib/StaticAnalyzer/Core/ExprEngine.cpp

switch (S->getStmtClass()) {
     case Stmt::OMPTeamsDistributeSimdDirectiveClass:
     case Stmt::OMPTeamsDistributeParallelForSimdDirectiveClass:
     case Stmt::OMPTeamsDistributeParallelForDirectiveClass:
     case Stmt::OMPTargetTeamsDirectiveClass:
     case Stmt::OMPTargetTeamsDistributeDirectiveClass:
     case Stmt::OMPTargetTeamsDistributeParallelForDirectiveClass:
     case Stmt::OMPTargetTeamsDistributeParallelForSimdDirectiveClass:
     case Stmt::OMPTargetTeamsDistributeSimdDirectiveClass:
+    case Stmt::OMPMetaDirectiveClass:
     case Stmt::CapturedStmtClass: {
       const ExplodedNode *node = Bldr.generateSink(S, Pred, Pred->getState());
       Engine.addAbortedBlock(node, currBldrCtx->getBlock());
       break;
     }
 
     case Stmt::ParenExprClass:
       llvm_unreachable("ParenExprs already handled.");

clang/test/OpenMP/metadirective_ast_print.cpp

+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -ast-print %s -o - | FileCheck %s
+// expected-no-diagnostics
+#define N 10000
+int main() {
+  int v1[N], v2[N], v3[N];
+  for (int i = 0; i < N; i++) {
+    v1[i] = (i + 1);
+    v2[i] = -(i + 1);
+  }
+#pragma omp metadirective           \
+    when(device = {arch("nvptx64")} \
+         : target teams distribute parallel for map(v1 [0:N], v2 [0:N], v3 [0:N])) default()
+  for (int i = 0; i < N; i++) {
+    v3[i] = v1[i] * v2[i];
+  }
+  return 0;
+}
+// CHECK:     int v1[10000], v2[10000], v3[10000];
+// CHECK:     for (int i = 0; i < 10000; i++) {
+// CHECK:         v1[i] = (i + 1);
+// CHECK:         v2[i] = -(i + 1);
+// CHECK:     }
+// CHECK:     #pragma omp target teams distribute parallel for map(tofrom: v1[0:10000],v2[0:10000],v3[0:10000])
+// CHECK:         for (int i = 0; i < 10000; i++) {
+// CHECK:             v3[i] = v1[i] * v2[i];
+// CHECK:         }

clang/test/OpenMP/metadirective_codegen.cpp

+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+#define N 10000
+int func() {
+  int v1[N], v2[N], v3[N];
+  for (int i = 0; i < N; i++) {
+    v1[i] = (i + 1);
+    v2[i] = -(i + 1);
+  }
+
+#pragma omp metadirective           \
+    when(device = {arch("nvptx64")} \
+         : target teams distribute parallel for map(v1 [0:N], v2 [0:N], v3 [0:N])) default()
+  for (int i = 0; i < N; i++) {
+    v3[i] = v1[i] * v2[i];
+  }
+
+  return 0;
+}
+// CHECK: %arrayidx4 = getelementptr inbounds [10000 x i32], [10000 x i32]* %v3, i64 0, i64 0
+// CHECK: %arrayidx5 = getelementptr inbounds [10000 x i32], [10000 x i32]* %v1, i64 0, i64 0
+// CHECK: %arrayidx6 = getelementptr inbounds [10000 x i32], [10000 x i32]* %v2, i64 0, i64 0
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_baseptrs, i32 0, i32 0
+// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to [10000 x i32]**
+// CHECK: store [10000 x i32]* %v3, [10000 x i32]** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 0
+// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to i32**
+// CHECK: store i32* %arrayidx4, i32** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_mappers, i64 0, i64 0
+// CHECK: store i8* null, i8** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_baseptrs, i32 0, i32 1
+// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to [10000 x i32]**
+// CHECK: store [10000 x i32]* %v1, [10000 x i32]** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 1
+// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to i32**
+// CHECK: store i32* %arrayidx5, i32** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_mappers, i64 0, i64 1
+// CHECK: store i8* null, i8** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_baseptrs, i32 0, i32 2
+// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to [10000 x i32]**
+// CHECK: store [10000 x i32]* %v2, [10000 x i32]** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 2
+// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to i32**
+// CHECK: store i32* %arrayidx6, i32** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_mappers, i64 0, i64 2
+// CHECK: store i8* null, i8** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_baseptrs, i32 0, i32 0
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 0
+// CHECK: call void @__kmpc_push_target_tripcount(%struct.ident_t* @3, i64 -1, i64 10000)
+// CHECK: %{{[0-9]*}} = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @3, i64 -1, i8* [[OMP_OUTLINED:@.+]].region_id, i32 3, i8** %{{[0-9]*}}, i8** %{{[0-9]*}}, i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 0, i32 0)
+// CHECK: %{{[0-9]*}} = icmp ne i32 %{{[0-9]*}}, 0
+// CHECK: br i1 %{{[0-9]*}}, label %omp_offload.failed, label %omp_offload.cont
+// CHECK: omp_offload.failed:
+// CHECK: call void [[OMP_OUTLINED:@.+]]
+// CHECK: br label %omp_offload.cont
+// CHECK: omp_offload.cont:
+// CHECK: ret i32 0

clang/test/OpenMP/metadirective_construct.cpp

+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+#define N 1000
+
+#pragma omp declare target
+void func1() {
+  // Test to check the case where metadirective is called within a target construct.
+  // Here the directive for 'distribute parallel for' will be generated
+#pragma omp metadirective       \
+    when(construct = {"target"} \
+         : distribute parallel for) default()
+  for (int i = 0; i < N; i++)
+    ;
+}
+#pragma omp end declare target
+
+void func() {
+#pragma omp target teams
+  func1();
+}
+
+// CHECK: define void @_Z5func1v() #0 {
+// CHECK: entry:
+// CHECK:   %.omp.iv = alloca i32, align 4
+// CHECK:   %tmp = alloca i32, align 4
+// CHECK:   %.omp.comb.lb = alloca i32, align 4
+// CHECK:   %.omp.comb.ub = alloca i32, align 4
+// CHECK:   %.omp.stride = alloca i32, align 4
+// CHECK:   %.omp.is_last = alloca i32, align 4
+// CHECK:   %i = alloca i32, align 4
+// CHECK:   %{{[0-9]*}} = call i32 @__kmpc_global_thread_num(%struct.ident_t* @2)
+// CHECK:   store i32 0, i32* %.omp.comb.lb, align 4
+// CHECK:   store i32 999, i32* %.omp.comb.ub, align 4
+// CHECK:   store i32 1, i32* %.omp.stride, align 4
+// CHECK:   store i32 0, i32* %.omp.is_last, align 4
+// CHECK:   call void @__kmpc_for_static_init_4(%struct.ident_t* @1, i32 %{{[0-9]*}}, i32 92, i32* %.omp.is_last, i32* %.omp.comb.lb, i32* %.omp.comb.ub, i32* %.omp.stride, i32 1, i32 1)
+// CHECK:   %{{[0-9]*}} = load i32, i32* %.omp.comb.ub, align 4
+// CHECK:   %cmp = icmp sgt i32 %{{[0-9]*}}, 999
+// CHECK:   br i1 %cmp, label %cond.true, label %cond.false
+// CHECK: cond.true:
+// CHECK:   br label %cond.end
+// CHECK: cond.false:
+// CHECK:   %{{[0-9]*}} = load i32, i32* %.omp.comb.ub, align 4
+// CHECK:   br label %cond.end
+// CHECK: cond.end:
+// CHECK:   %cond = phi i32 [ 999, %cond.true ], [ %{{[0-9]*}}, %cond.false ]
+// CHECK:   store i32 %cond, i32* %.omp.comb.ub, align 4
+// CHECK:   %{{[0-9]*}} = load i32, i32* %.omp.comb.lb, align 4
+// CHECK:   store i32 %{{[0-9]*}}, i32* %.omp.iv, align 4
+// CHECK:   br label %omp.inner.for.cond
+// CHECK: omp.inner.for.cond:
+// CHECK:   %{{[0-9]*}} = load i32, i32* %.omp.iv, align 4
+// CHECK:   %{{[0-9]*}} = load i32, i32* %.omp.comb.ub, align 4
+// CHECK:   %cmp1 = icmp sle i32 %{{[0-9]*}}, %{{[0-9]*}}
+// CHECK:   br i1 %cmp1, label %omp.inner.for.body, label %omp.inner.for.end
+// CHECK: omp.inner.for.body:
+// CHECK:   %{{[0-9]*}} = load i32, i32* %.omp.comb.lb, align 4
+// CHECK:   %{{[0-9]*}} = zext i32 %{{[0-9]*}} to i64
+// CHECK:   %{{[0-9]*}} = load i32, i32* %.omp.comb.ub, align 4
+// CHECK:   %{{[0-9]*}} = zext i32 %{{[0-9]*}} to i64
+// CHECK:   call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @2, i32 2, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i64, i64)* [[OMP_OUTLINED:@.+]] to void (i32*, i32*, ...)*), i64 %{{[0-9]*}}, i64 %{{[0-9]*}})
+// CHECK:   br label %omp.inner.for.inc
+// CHECK: omp.inner.for.inc:
+// CHECK:   %{{[0-9]*}} = load i32, i32* %.omp.iv, align 4
+// CHECK:   %{{[0-9]*}} = load i32, i32* %.omp.stride, align 4
+// CHECK:   %add = add nsw i32 %{{[0-9]*}}, %{{[0-9]*}}
+// CHECK:   store i32 %add, i32* %.omp.iv, align 4
+// CHECK:   br label %omp.inner.for.cond
+// CHECK: omp.inner.for.end:
+// CHECK:   br label %omp.loop.exit
+// CHECK: omp.loop.exit:
+// CHECK:   call void @__kmpc_for_static_fini(%struct.ident_t* @1, i32 %{{[0-9]*}})
+// CHECK:   ret void
+// CHECK: }
+
+// Declaration of funtion handling distribute parallel for
+// CHECK: define internal void [[OMP_OUTLINED:@.+]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i64 %.previous.lb., i64 %.previous.ub.)
+
+// CHECK: define void @_Z4funcv() #0 {
+// CHECK: entry:
+// CHECK:   %{{[0-9]*}} = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @2, i64 -1, i8* [[OMP_OFFLOADING:@.+]].region_id
+// CHECK:   %{{[0-9]*}} = icmp ne i32 %{{[0-9]*}}, 0
+// CHECK:   br i1 %{{[0-9]*}}, label %omp_offload.failed, label %omp_offload.cont
+// CHECK: omp_offload.failed:
+// CHECK:   call void [[OMP_OFFLOADING:@.+]]() #1
+// CHECK:   br label %omp_offload.cont
+// CHECK: omp_offload.cont:
+// CHECK:   ret void
+// CHECK: }
+
+// CHECK: define internal void [[OMP_OFFLOADING:@.+]]() #2 {
+// CHECK: entry:
+// CHECK:   call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @2, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*))
+// CHECK:   ret void
+// CHECK: }
+
+// CHECK: define internal void @.omp_outlined..1(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #2 {
+// CHECK: entry:
+// CHECK:   %.global_tid..addr = alloca i32*, align 8
+// CHECK:   %.bound_tid..addr = alloca i32*, align 8
+// CHECK:   store i32* %.global_tid., i32** %.global_tid..addr, align 8
+// CHECK:   store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
+// CHECK:   call void @_Z5func1v()
+// CHECK:   ret void
+// CHECK: }

clang/test/OpenMP/metadirective_empty.cpp

+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+#define N 1000
+void func() {
+  // Test where a valid when clause contains empty directive.
+  // The directive will be ignored and code for a serial for loop will be generated.
+#pragma omp metadirective when(implementation = {vendor(llvm)} \
+                               :) default(parallel for)
+  for (int i = 0; i < N; i++)
+    ;
+}
+
+// CHECK: define void @_Z4funcv() #0 {
+// CHECK: entry:
+// CHECK:   %i = alloca i32, align 4
+// CHECK:   store i32 0, i32* %i, align 4
+// CHECK:   br label %for.cond
+// CHECK: for.cond:
+// CHECK:   %0 = load i32, i32* %i, align 4
+// CHECK:   %cmp = icmp slt i32 %0, 1000
+// CHECK:   br i1 %cmp, label %for.body, label %for.end
+// CHECK: for.body:
+// CHECK:   br label %for.inc
+// CHECK: for.inc:
+// CHECK:   %1 = load i32, i32* %i, align 4
+// CHECK:   %inc = add nsw i32 %1, 1
+// CHECK:   store i32 %inc, i32* %i, align 4
+// CHECK:   br label %for.cond, !llvm.loop !2
+// CHECK: for.end:
+// CHECK:   ret void
+// CHECK: }

clang/test/OpenMP/metadirective_implementation.cpp

+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+void func() {
+#pragma omp target
+  // Test to check implementation and device.
+  // Since we are building with llvm for nvptx64 device, 512 teams with thread_limit of 32 are generated.
+#pragma omp metadirective                                            \
+    when(implementation = {vendor(llvm)}, device = {arch(nvptx64)}   \
+         : teams num_teams(512) thread_limit(32))                    \
+        when(implementation = {vendor(amd)}, device = {arch(amdgcn)} \
+             : teams num_teams(512) thread_limit(64)) default(teams)
+#pragma omp distribute parallel for
+  for (int i = 0; i < 10000; i++)
+    ;
+}
+
+// CHECK: define void @_Z4funcv() #0 {
+// CHECK: entry:
+// CHECK:   %tmp = alloca i32, align 4
+// CHECK:   call void @__kmpc_push_target_tripcount
+// CHECK:   %0 = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @3, i64 -1, i8* [[OMP_OUTLINED:@.+]].region_id,
+// CHECK:   %1 = icmp ne i32 %0, 0
+// CHECK:   br i1 %1, label %omp_offload.failed, label %omp_offload.cont
+// CHECK: omp_offload.failed:                               ; preds = %entry
+// CHECK:   call void [[OMP_OUTLINED:@.+]] #2
+// CHECK:   br label %omp_offload.cont
+// CHECK: omp_offload.cont:                                 ; preds = %omp_offload.failed, %entry
+// CHECK:   ret void
+// CHECK: }
+// CHECK: define internal void [[OMP_OUTLINED:@.+]] #1 {
+// CHECK: entry:
+// CHECK:   %0 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @3)
+// CHECK:   call void @__kmpc_push_num_teams(%struct.ident_t* @3, i32 %0, i32 512, i32 32)
+// CHECK:   call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @3, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OMP_OUTLINED:@.+]] to void (i32*, i32*, ...)*))
+// CHECK:   ret void
+// CHECK: }

clang/tools/libclang/CIndex.cpp

CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) {
   case CXCursor_ObjCSynthesizeDecl:
     return cxstring::createRef("ObjCSynthesizeDecl");
   case CXCursor_ObjCDynamicDecl:
     return cxstring::createRef("ObjCDynamicDecl");
   case CXCursor_CXXAccessSpecifier:
     return cxstring::createRef("CXXAccessSpecifier");
   case CXCursor_ModuleImportDecl:
     return cxstring::createRef("ModuleImport");
+  case CXCursor_OMPMetaDirective:
+    return cxstring::createRef("OMPMetaDirective");
   case CXCursor_OMPParallelDirective:
     return cxstring::createRef("OMPParallelDirective");
   case CXCursor_OMPSimdDirective:
     return cxstring::createRef("OMPSimdDirective");
   case CXCursor_OMPForDirective:
     return cxstring::createRef("OMPForDirective");
   case CXCursor_OMPForSimdDirective:
     return cxstring::createRef("OMPForSimdDirective");

clang/tools/libclang/CXCursor.cpp

case Stmt::ObjCMessageExprClass: {
     }
     CXCursor C = {K, 0, {Parent, S, TU}};
     return getSelectorIdentifierCursor(SelectorIdIndex, C);
   }
 
   case Stmt::MSDependentExistsStmtClass:
     K = CXCursor_UnexposedStmt;
     break;
+  case Stmt::OMPMetaDirectiveClass:
+    K = CXCursor_OMPMetaDirective;
+    break;
   case Stmt::OMPParallelDirectiveClass:
     K = CXCursor_OMPParallelDirective;
     break;
   case Stmt::OMPSimdDirectiveClass:
     K = CXCursor_OMPSimdDirective;
     break;
   case Stmt::OMPForDirectiveClass:
     K = CXCursor_OMPForDirective;

llvm/include/llvm/Frontend/OpenMP/OMP.td

 def OMPC_SimdLen : Clause<"simdlen"> {
   let clangClass = "OMPSimdlenClause";
   let flangClassValue = "ScalarIntConstantExpr";
 }
 def OMPC_Collapse : Clause<"collapse"> {
   let clangClass = "OMPCollapseClause";
   let flangClassValue = "ScalarIntConstantExpr";
 }
+def OMPC_When: Clause<"when"> {}
 def OMPC_Default : Clause<"default"> {
   let clangClass = "OMPDefaultClause";
   let flangClass = "OmpDefaultClause";
 }
 def OMPC_Private : Clause<"private"> {
   let clangClass = "OMPPrivateClause";
   let flangClassValue = "OmpObjectList";
 }

 }
 def OMPC_Inbranch : Clause<"inbranch"> {}
 def OMPC_Notinbranch : Clause<"notinbranch"> {}
 
 //===----------------------------------------------------------------------===//
 // Definition of OpenMP directives
 //===----------------------------------------------------------------------===//
 
+def OMP_Metadirective : Directive<"metadirective"> {
+  let allowedClauses = [
+    VersionedClause<OMPC_When>
+  ];
+  let allowedOnceClauses = [
+    VersionedClause<OMPC_Default>
+  ];
+}
 def OMP_ThreadPrivate : Directive<"threadprivate"> {}
 def OMP_Parallel : Directive<"parallel"> {
   let allowedClauses = [
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_Shared>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Copyin>,

llvm/include/llvm/Frontend/OpenMP/OMPContext.h

 /// Return true if \p VMI is applicable in \p Ctx, that is, all traits required
 /// by \p VMI are available in the OpenMP context \p Ctx. If \p DeviceSetOnly is
 /// true, only the device selector set, if present, are checked. Note that we
 /// still honor extension traits provided by the user.
 bool isVariantApplicableInContext(const VariantMatchInfo &VMI,
                                   const OMPContext &Ctx,
                                   bool DeviceSetOnly = false);
 
+/// Return the index (into \p VMIs) of the When clause with the highest score
+/// from the ones applicble in \p Ctx.
+/// In OpenMP 5.1 set OrderedMatch to those conditions which need runtime
+/// resolution.
+int getBestWhenMatchForContext(
+    const SmallVectorImpl<VariantMatchInfo> &VMIs, const OMPContext &Ctx,
+    SmallVectorImpl<unsigned> *OrderedMatch = nullptr);
+/// Return the index (into \p VMIs) of the When clause with the highest score
+/// from the ones applicble in \p Ctx.
+// int getBestWhenMatchForContext(const SmallVectorImpl<VariantMatchInfo> &VMIs,
+//                               const OMPContext &Ctx);
 /// Return the index (into \p VMIs) of the variant with the highest score
 /// from the ones applicble in \p Ctx. See llvm::isVariantApplicableInContext.
 int getBestVariantMatchForContext(const SmallVectorImpl<VariantMatchInfo> &VMIs,
                                   const OMPContext &Ctx);
 
 } // namespace omp
 
 template <> struct DenseMapInfo<omp::TraitProperty> {

llvm/include/llvm/Frontend/OpenMP/OMPKinds.def

 __OMP_CLAUSE(exclusive, OMPExclusiveClause)
 __OMP_CLAUSE(uses_allocators, OMPUsesAllocatorsClause)
 __OMP_CLAUSE(affinity, OMPAffinityClause)
 __OMP_CLAUSE(use_device_addr, OMPUseDeviceAddrClause)
 
 __OMP_CLAUSE_NO_CLASS(uniform)
 __OMP_CLAUSE_NO_CLASS(device_type)
 __OMP_CLAUSE_NO_CLASS(match)
+__OMP_CLAUSE_NO_CLASS(when)
 
 __OMP_IMPLICIT_CLAUSE_CLASS(depobj, "depobj", OMPDepobjClause)
 __OMP_IMPLICIT_CLAUSE_CLASS(flush, "flush", OMPFlushClause)
 
 __OMP_IMPLICIT_CLAUSE_NO_CLASS(threadprivate, "threadprivate or thread local")
 __OMP_IMPLICIT_CLAUSE_NO_CLASS(unknown, "unknown")
 
 #undef __OMP_IMPLICIT_CLAUSE_NO_CLASS

llvm/lib/Frontend/OpenMP/OMPContext.cpp

for (TraitProperty Property : VMI.ConstructTraits) {
     Score += (1ULL << ConstructMatches[ConstructIdx++]);
   }
 
   LLVM_DEBUG(dbgs() << "[" << DEBUG_TYPE << "] Variant has a score of " << Score
                     << "\n");
   return Score;
 }
 
+int llvm::omp::getBestWhenMatchForContext(
+    const SmallVectorImpl<VariantMatchInfo> &VMIs, const OMPContext &Ctx,
+    SmallVectorImpl<unsigned> *OrderedMatch) {
+
+  APInt BestScore(64, 0);
+  int BestVMIIdx = -1;
+  const VariantMatchInfo *BestVMI = nullptr;
+
+  for (unsigned u = 0, e = VMIs.size(); u < e; ++u) {
+    const VariantMatchInfo &VMI = VMIs[u];
+
+    SmallVector<unsigned, 8> ConstructMatches;
+    // Check if its clearly not the best.
+    APInt Score = getVariantMatchScore(VMI, Ctx, ConstructMatches);
+    if (Score.ult(BestScore))
+      continue;
+    // Equal score need subset checks.
+    if (Score.eq(BestScore)) {
+      // Strict subset are never best.
+      if (isStrictSubset(VMI, *BestVMI))
+        continue;
+      // Same score and the current best is no strict subset so we keep it.
+      if (!isStrictSubset(*BestVMI, VMI))
+        continue;
+    }
+    // New best found.
+    BestVMI = &VMI;
+    BestVMIIdx = u;
+    BestScore = Score;
+  }
+
+  return BestVMIIdx;
+}
+
+/*int llvm::omp::getBestWhenMatchForContext(
+    const SmallVectorImpl<VariantMatchInfo> &VMIs, const OMPContext &Ctx) {
+
+  APInt BestScore(64, 0);
+  int BestVMIIdx = -1;
+  const VariantMatchInfo *BestVMI = nullptr;
+
+  for (unsigned u = 0, e = VMIs.size(); u < e; ++u) {
+    const VariantMatchInfo &VMI = VMIs[u];
+
+    SmallVector<unsigned, 8> ConstructMatches;
+    // Check if its clearly not the best.
+    APInt Score = getVariantMatchScore(VMI, Ctx, ConstructMatches);
+    if (Score.ult(BestScore))
+      continue;
+    // Equal score need subset checks.
+    if (Score.eq(BestScore)) {
+      // Strict subset are never best.
+      if (isStrictSubset(VMI, *BestVMI))
+        continue;
+      // Same score and the current best is no strict subset so we keep it.
+      if (!isStrictSubset(*BestVMI, VMI))
+        continue;
+    }
+    // New best found.
+    BestVMI = &VMI;
+    BestVMIIdx = u;
+    BestScore = Score;
+  }
+
+  return BestVMIIdx;
+}*/
+
 int llvm::omp::getBestVariantMatchForContext(
     const SmallVectorImpl<VariantMatchInfo> &VMIs, const OMPContext &Ctx) {
 
   APInt BestScore(64, 0);
   int BestVMIIdx = -1;
   const VariantMatchInfo *BestVMI = nullptr;
 
   for (unsigned u = 0, e = VMIs.size(); u < e; ++u) {