diff mbox series

[WIP,OpenMP] OpenMP metadirectives support

Message ID 8d413974-0068-3a31-6ae5-d36c1be06d06@codesourcery.com
State New
Headers show
Series [WIP,OpenMP] OpenMP metadirectives support | expand

Commit Message

Kwok Cheung Yeung July 9, 2021, 11:16 a.m. UTC 
Hello

This is a WIP implementation of metadirectives as defined in the OpenMP 5.0 
spec. I intend to add support for metadirectives as specified in OpenMP 5.1 
later (where the directive can be selected dynamically at runtime), but am 
concentrating on the static part for now. Parsing has only been implemented in 
the C frontend so far. I am especially interested in feedback regarding certain 
aspects of the implementation before I become too committed to the current design.

1) When parsing each directive variant, a vector of tokens is constructed and 
populated with the tokens for a regular equivalent pragma, along with the tokens 
for its clauses and the body. The parser routine for that pragma type is then 
called with these tokens, and the entire resulting parse tree is stored as a 
sub-tree of the metadirective tree structure.

This results in the body being parsed and stored once for each directive 
variant. I believe this is necessary because the body is parsed differently if 
there is a 'for' in the directive (using c_parser_omp_for_loop) compared to if 
there is not, plus clauses in the directive (e.g. tile, collapse) can change how 
the for loop is parsed.

As an optimisation, identical body trees could be merged together, but that can 
come later.

2) Selectors in the device set (i.e. kind, isa, arch) resolve differently 
depending on whether the program is running on a target or on the host. Since we 
don't keep multiple versions of a function for each target on the host compiler, 
resolving metadirectives with these selectors needs to be delayed until after 
LTO streaming, at which point the host or offload compiler can make the 
appropriate decision.

One negative of this is that the metadirective Gimple representation lasts 
beyond the OMP expand stage, when generally we would expect all OMP directives 
to have been expanded to something else.

3) In the OpenMP examples (version 5.0.1), section 9.7, the example 
metadirective.3.c does not work as expected.

#pragma omp declare target
void exp_pi_diff(double *d, double my_pi){
    #pragma omp metadirective \
                when( construct={target}: distribute parallel for ) \
                default( parallel for simd)
...
int main()
{
    ...
    #pragma omp target teams map(tofrom: d[0:N])
    exp_pi_diff(d,my_pi);
    ...
    exp_pi_diff(d,my_pi);

In the first call to exp_pi_diff in an '#pragma omp target' construct, the 
metadirective is expected to expand to 'distribute parallel for', but in the 
second (without the '#pragma omp target'), it should expand to 'parallel for simd'.

During OMP expansion of the 'omp target', it creates a child function that calls 
exp_pi_diff:

__attribute__((omp target entrypoint))
void main._omp_fn.0 (const struct .omp_data_t.12 & restrict .omp_data_i)
{
   ...
   <bb 4> :
   __builtin_GOMP_teams (0, 0);
   exp_pi_diff (d.13, my_pi);

This is not a problem on the offload compiler (since by definition its copy of 
exp_pi_diff must be in a 'target'), but if the host device is used, the same 
version of exp_pi_diff is called in both target and non-target contexts.

What would be the best way to solve this? Offhand, I can think of two solutions:

(a) Recursively go through all functions that can be reached via a target region 
and create clones for each, redirecting all function calls in the clones to the 
new cloned versions. Resolve the metadirectives in the clones and originals 
separately.

(b) Make the construct selector a dynamic selector when OpenMP 5.1 metadirective 
support is implemented. Keep track of the current construct list every time an 
OpenMP construct is entered or exited, and make the decision at runtime.


Thanks

Kwok

Comments

Kwok Cheung Yeung July 26, 2021, 11:38 a.m. UTC | #1 
Ping? Does anyone have any opinions on how this issue should be resolved?

On 09/07/2021 12:16 pm, Kwok Cheung Yeung wrote:
> 3) In the OpenMP examples (version 5.0.1), section 9.7, the example 
> metadirective.3.c does not work as expected.
> 
> #pragma omp declare target
> void exp_pi_diff(double *d, double my_pi){
>     #pragma omp metadirective \
>                 when( construct={target}: distribute parallel for ) \
>                 default( parallel for simd)
> ...
> int main()
> {
>     ...
>     #pragma omp target teams map(tofrom: d[0:N])
>     exp_pi_diff(d,my_pi);
>     ...
>     exp_pi_diff(d,my_pi);
> 
> In the first call to exp_pi_diff in an '#pragma omp target' construct, the 
> metadirective is expected to expand to 'distribute parallel for', but in the 
> second (without the '#pragma omp target'), it should expand to 'parallel for simd'.
> 
> During OMP expansion of the 'omp target', it creates a child function that calls 
> exp_pi_diff:
> 
> __attribute__((omp target entrypoint))
> void main._omp_fn.0 (const struct .omp_data_t.12 & restrict .omp_data_i)
> {
>    ...
>    <bb 4> :
>    __builtin_GOMP_teams (0, 0);
>    exp_pi_diff (d.13, my_pi);
> 
> This is not a problem on the offload compiler (since by definition its copy of 
> exp_pi_diff must be in a 'target'), but if the host device is used, the same 
> version of exp_pi_diff is called in both target and non-target contexts.
> 
> What would be the best way to solve this? Offhand, I can think of two solutions:
> 
> (a) Recursively go through all functions that can be reached via a target region 
> and create clones for each, redirecting all function calls in the clones to the 
> new cloned versions. Resolve the metadirectives in the clones and originals 
> separately.
> 

Maybe this could be done at the same time as when marking functions implicitly 
'declare target'? It seems a lot of work for one special case though...

> (b) Make the construct selector a dynamic selector when OpenMP 5.1 metadirective 
> support is implemented. Keep track of the current construct list every time an 
> OpenMP construct is entered or exited, and make the decision at runtime.
> 

I think this would be easier to implement at runtime (assuming that the 
infrastructure for OpenMP 5.1 was already in place) since this a host-side 
issue, but it probably goes against the intent of the specification, given that 
the 'construct' selector set appeared in the 5.0 specification before dynamic 
replacements became available.

Thanks

Kwok
Jakub Jelinek July 26, 2021, 2:29 p.m. UTC | #2 
On Fri, Jul 09, 2021 at 12:16:15PM +0100, Kwok Cheung Yeung wrote:
> This is a WIP implementation of metadirectives as defined in the OpenMP 5.0
> spec. I intend to add support for metadirectives as specified in OpenMP 5.1
> later (where the directive can be selected dynamically at runtime), but am
> concentrating on the static part for now. Parsing has only been implemented
> in the C frontend so far. I am especially interested in feedback regarding
> certain aspects of the implementation before I become too committed to the
> current design.

Note, there is a partial overlap with the attribute syntax changes, see below.
c-family/c-omp.c now has omp_directives table that should be updated for
changes like this and then c_omp_categorize_directive that returns some
information about the directives given a directive name (though, that name
can be one, two or three tokens long, consider e.g. target enter data
or cancellation point directives).

For metadirective, I think very special case are declarative directives in
them, I'd tend to sorry for them at least for now, I'm pretty sure many
cases with them are just unimplementable and will need to be restricted in
the standard, others can be implemented with lots of effort.
Whether it is e.g. metadirective guarding declare target ... end declare
target pair that would only conditionally set declare target and instead of
a single bit to find out if something is declare target or not we'd until
resolved need to compute it for all possibilities, or e.g. conditional
declare reduction/declare mapper where the name lookup for reduction or map
directives would be dependent on metadirective resolution later on, etc.
I'm afraid a total nightmare nobody has really thought about details for it.

> 1) When parsing each directive variant, a vector of tokens is constructed
> and populated with the tokens for a regular equivalent pragma, along with
> the tokens for its clauses and the body. The parser routine for that pragma
> type is then called with these tokens, and the entire resulting parse tree
> is stored as a sub-tree of the metadirective tree structure.
> 
> This results in the body being parsed and stored once for each directive
> variant. I believe this is necessary because the body is parsed differently
> if there is a 'for' in the directive (using c_parser_omp_for_loop) compared
> to if there is not, plus clauses in the directive (e.g. tile, collapse) can
> change how the for loop is parsed.
> 
> As an optimisation, identical body trees could be merged together, but that
> can come later.

I'm afraid it isn't just an optimization and we need to be as smart as
possible.  I'm not sure it is possible to parse everything many times,
consider e.g. labels in the blocks, nested function definitions, variable
definitions, etc.
While OpenMP requires that essentially the code must be valid if the
metadirective is replaced by any of those mentioned directives which rules
quite some weirdo corner cases, nothing prevents e.g. two or more
when directives to be standalone directives (which don't have any body and
so whatever comes after them should be left parsed for later as normal
statement sequence), one or more to be normal constructs that accept a
structured block and one or more to be e.g. looping constructs (simd, for,
distribute, taskloop or combined versions of those).
Even when issues with labels etc. are somehow solved (e.g. for structured
blocks we have the restriction that goto, break, continue, or switch into
a case/default label, etc. can't be used to enter or exit the structured
block which could mean some cases can be handled through renaming seen
labels in all but one bodies), most important is to sync on where parsing
should continue after the metadirective.
I think it would be nice if the metadirective parsing at least made quick
analysis on what kind of bodies the directives will want and can use the new
c-omp.c infrastructure or if needed extend it (e.g. separate the C_OMP_DIR_CONSTRUCT
category into C_OMP_DIR_CONSTRUCT and C_OMP_DIR_LOOPING_CONSTRUCT where
the latter would be used for those that expect some omp loop after it).
One option would be then to parse the body as the most restricted construct
(looping (and determine highest needed collapse and ordered), then construct,
then standalone) and be able to adjust what we parsed into what the
different constructs need, but another option is the separate parsing of
the code after the directive multiple times, but at least in the order of
most restricted to least restricted, remember where to stop and don't parse
it multiple times at least for directives that need the same thing.

> 
> 2) Selectors in the device set (i.e. kind, isa, arch) resolve differently
> depending on whether the program is running on a target or on the host.
> Since we don't keep multiple versions of a function for each target on the
> host compiler, resolving metadirectives with these selectors needs to be
> delayed until after LTO streaming, at which point the host or offload
> compiler can make the appropriate decision.

How is this different from declare variant?  For declare variant, it is true
I'm never trying to resolve it already during parsing of the call and that
probably should be changed, do a first attempt at that point.  Initially
I thought it typically will not be possible, but later clarification and
strong desire of LLVM/ICC etc. to do everything or almost everything already
during parsing suggests that it must be doable at least in some cases.
E.g. we have restrictions that requires directive on which some decision
could be dependent must appear only lexically before it or not at all, etc.
So, similarly to that, metadirective ideally should see if something is
impossible already during parsing (dunno if it should mean we wouldn't parse
the body in that case, that would mean worse diagnostics), then repeat the
checks during gimplification like declare variant is resolved there, then
repeat again after IPA.  Would be probably best if for metadirectives that
resolve to executable directives we represent it by something like a magic
IFN that is told everything needed to decide and can share as much code as
possible with the declare variant decisions.

It is true other compilers implement offloading quite differently from GCC,
by repeating all of preprocessing, parsing etc. for the offloading target,
so they can decide some metadirective/declare variant decisions earlier than
we can.  On the other side that approach has also quite some disadvantages,
it is much harder to ensure ABI compatibility between the host and offload
code if one can use #ifdefs and whatever to change layout of everything in
between.

For the checks during parsing, we'll need a different way how to track which
directives are currently active (or defer anything with construct
selectors till gimplification).  It is true that resolving that during
parsing goes against the goal to parse as many bodies together as possible,
so we need to pick one or the other.  Parsing what follows for all
standalone directives isn't a problem of course, but if the metadirective
has one when with for and another with simd, then parsing the loop just once
would be a problem if there is metadirective in the body that wants to
decide whether it is in for or simd and wants that decision be done during
parsing.

> One negative of this is that the metadirective Gimple representation lasts
> beyond the OMP expand stage, when generally we would expect all OMP
> directives to have been expanded to something else.
> 
> 3) In the OpenMP examples (version 5.0.1), section 9.7, the example
> metadirective.3.c does not work as expected.
> 
> #pragma omp declare target
> void exp_pi_diff(double *d, double my_pi){
>    #pragma omp metadirective \
>                when( construct={target}: distribute parallel for ) \
>                default( parallel for simd)
> ...
> int main()
> {
>    ...
>    #pragma omp target teams map(tofrom: d[0:N])
>    exp_pi_diff(d,my_pi);
>    ...
>    exp_pi_diff(d,my_pi);

The spec says in this case that the target construct is added to the
construct set because of the function appearing in between omp declare target
and omp end declare target, so the above is something that resolves
statically to distribute parallel for.
It is true that in OpenMP 5.1 the earlier
For functions within a declare target block, the target trait is added to the beginning of the
set as c 1 for any versions of the function that are generated for target regions so the total size
of the set is increased by 1.
has been mistakenly replaced with:
For device routines, the target trait is added to the beginning of the set as c 1 for any versions of
the procedure that are generated for target regions so the total size of the set is increased by 1.
by that has been corrected in 5.2:
C/C++:
For functions that are declared in a code region that is delimited by a declare target directive and
its paired end directive, the target trait is added to the beginning of the set as c 1 for any target
variants that result from the directive so the total size of the set is increased by one.
Fortran:
If a declare target directive appears in the specification part of a procedure or in the
specification part of a procedure interface body, the target trait is added to the beginning of the
set as c 1 for any target variants that result from the directive so the total size of the set is
increased by one.

So, it is really a static decision that can be decided already during
parsing.
> --- a/gcc/Makefile.in
> +++ b/gcc/Makefile.in
> @@ -1505,6 +1505,7 @@ OBJS = \
>  	omp-general.o \
>  	omp-low.o \
>  	omp-oacc-kernels-decompose.o \
> +        omp-expand-metadirective.o \

Spaces instead of tab.

> @@ -1312,12 +1312,14 @@ static const struct omp_pragma_def omp_pragmas[] = {
>    { "allocate", PRAGMA_OMP_ALLOCATE },
>    { "atomic", PRAGMA_OMP_ATOMIC },
>    { "barrier", PRAGMA_OMP_BARRIER },
> +  { "begin", PRAGMA_OMP_BEGIN },
>    { "cancel", PRAGMA_OMP_CANCEL },
>    { "cancellation", PRAGMA_OMP_CANCELLATION_POINT },
>    { "critical", PRAGMA_OMP_CRITICAL },
>    { "depobj", PRAGMA_OMP_DEPOBJ },
> -  { "end", PRAGMA_OMP_END_DECLARE_TARGET },
> +  { "end", PRAGMA_OMP_END },
>    { "flush", PRAGMA_OMP_FLUSH },
> +  { "metadirective", PRAGMA_OMP_METADIRECTIVE },
>    { "requires", PRAGMA_OMP_REQUIRES },
>    { "section", PRAGMA_OMP_SECTION },
>    { "sections", PRAGMA_OMP_SECTIONS },

Please update for this also the omp_directives array.

> +enum pragma_kind
> +c_pp_lookup_pragma_by_name (const char *name)
> +{
> +  const int n_omp_pragmas = sizeof (omp_pragmas) / sizeof (*omp_pragmas);
> +  const int n_omp_pragmas_simd = sizeof (omp_pragmas_simd)
> +				 / sizeof (*omp_pragmas_simd);
> +
> +  void *result = bsearch (name, omp_pragmas, n_omp_pragmas,
> +			  sizeof (*omp_pragmas),
> +			  c_pp_lookup_pragma_by_name_1);
> +  if (!result)
> +    result = bsearch (name, omp_pragmas_simd, n_omp_pragmas_simd,
> +		      sizeof (*omp_pragmas_simd),
> +		      c_pp_lookup_pragma_by_name_1);
> +
> +  if (result)
> +    {
> +      const struct omp_pragma_def *def
> +	= (const struct omp_pragma_def *) result;
> +
> +      return (enum pragma_kind) def->id;
> +    }
> +
> +  return PRAGMA_NONE;
> +}

I think this should be dropped and c_omp_categorize_directive should
be used instead of it.

Please add a function comment to show the grammar.  See e.g.
c_parser_omp_declare.

>  
> +static void
> +c_parser_omp_begin (c_parser *parser, bool *if_p)
> +{
> +  location_t loc = c_parser_peek_token (parser)->location;
> +  c_parser_consume_pragma(parser);

Space before (.

> +  if (c_parser_peek_token (parser)->type == CPP_NAME)
> +    {
> +      const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
> +
> +      if (strcmp (p, "metadirective") == 0)
> +	{
> +	  char p_name[sizeof "#pragma omp teams distribute parallel for simd"];
> +	  omp_clause_mask mask (0);
> +
> +	  c_parser_consume_token (parser);
> +	  c_parser_omp_metadirective (loc, parser, p_name, mask, NULL, if_p,
> +				      true);

metadirective, by not being itself combinable, doesn't need this p_name
and mask stuff.  That is used only for combined/composite construct when
the p_name and mask need to be computed dynamically based on the exact
parsing.  The begin metadirective vs. metadirective difference is a boolean
one that can be either passed as bool, or if the pragma token is passed
one could look at its pragma kind.

> +	  return;
> +	}
> +    }
> +
> +  error_at (loc, "expected %<begin metadirective%>");

"expected %<metadirective%>" ?  #pragma omp begin already appeared...

> +  c_parser_skip_to_pragma_eol (parser);
> +}
> +
> +static void
> +c_parser_omp_end (c_parser *parser)

Similarly with the function comment.

	Jakub
Kwok Cheung Yeung July 26, 2021, 7:28 p.m. UTC | #3 
Hello

Thanks for your reply.

On 26/07/2021 3:29 pm, Jakub Jelinek wrote:
> On Fri, Jul 09, 2021 at 12:16:15PM +0100, Kwok Cheung Yeung wrote:
>> 3) In the OpenMP examples (version 5.0.1), section 9.7, the example
>> metadirective.3.c does not work as expected.
>>
>> #pragma omp declare target
>> void exp_pi_diff(double *d, double my_pi){
>>     #pragma omp metadirective \
>>                 when( construct={target}: distribute parallel for ) \
>>                 default( parallel for simd)
>> ...
>> int main()
>> {
>>     ...
>>     #pragma omp target teams map(tofrom: d[0:N])
>>     exp_pi_diff(d,my_pi);
>>     ...
>>     exp_pi_diff(d,my_pi);
> 
> The spec says in this case that the target construct is added to the
> construct set because of the function appearing in between omp declare target
> and omp end declare target, so the above is something that resolves
> statically to distribute parallel for.
> It is true that in OpenMP 5.1 the earlier
> For functions within a declare target block, the target trait is added to the beginning of the
> set as c 1 for any versions of the function that are generated for target regions so the total size
> of the set is increased by 1.
> has been mistakenly replaced with:
> For device routines, the target trait is added to the beginning of the set as c 1 for any versions of
> the procedure that are generated for target regions so the total size of the set is increased by 1.
> by that has been corrected in 5.2:
> C/C++:
> For functions that are declared in a code region that is delimited by a declare target directive and
> its paired end directive, the target trait is added to the beginning of the set as c 1 for any target
> variants that result from the directive so the total size of the set is increased by one.
> Fortran:
> If a declare target directive appears in the specification part of a procedure or in the
> specification part of a procedure interface body, the target trait is added to the beginning of the
> set as c 1 for any target variants that result from the directive so the total size of the set is
> increased by one.
> 
> So, it is really a static decision that can be decided already during
> parsing.

In Section 1.2.2 of the OpenMP TR10 spec, 'target variant' is defined as:

A version of a device routine that can only be executed as part of a target region.

So isn't this really saying the same thing as the previous versions of the spec? 
The target trait is added to the beginning of the construct set _for any target 
variants_ that result from the directive (implying that it shouldn't be added 
for non-target variants). In this example, the same function exp_pi_diff is 
being used in both a target and non-target context, so shouldn't the 
metadirective resolve differently in the two contexts, independently of the 
function being declared in a 'declare target' block? If not, there does not seem 
to be much point in that example (in section 9.7 of the OpenMP Examples v5.0.1).

 From reading the spec, I infer that they expect the device and non-device 
versions of a function with 'declare target' to be separate, but that is not 
currently the case for GCC - on the host compiler, the same version of the 
function gets called in both target and non-target regions (though in the target 
region case, it gets called indirectly via a compiler-generated function with a 
name like main._omp_fn.0). The offload compiler gets its own streamed version, 
so there is no conflict there - by definition, its version must be in a target 
context.

Thanks,

Kwok
Jakub Jelinek July 26, 2021, 7:56 p.m. UTC | #4 
On Mon, Jul 26, 2021 at 08:28:16PM +0100, Kwok Cheung Yeung wrote:
> In Section 1.2.2 of the OpenMP TR10 spec, 'target variant' is defined as:
> 
> A version of a device routine that can only be executed as part of a target region.

Yes, that is a target variant, but I'm pretty sure we've decided that
the target construct added for declare target is actually not a dynamic
property.  So basically mostly return to the 5.0 wording with clarifications
for Fortran.  See
https://github.com/OpenMP/spec/issues/2612#issuecomment-849742988
for details.
Making the target in construct dynamic would pretty much force all the
scoring to be dynamic as well.

	Jakub
Kwok Cheung Yeung July 26, 2021, 9:19 p.m. UTC | #5 
Hello

On 26/07/2021 8:56 pm, Jakub Jelinek wrote:
> On Mon, Jul 26, 2021 at 08:28:16PM +0100, Kwok Cheung Yeung wrote:
>> In Section 1.2.2 of the OpenMP TR10 spec, 'target variant' is defined as:
>>
>> A version of a device routine that can only be executed as part of a target region.
> 
> Yes, that is a target variant, but I'm pretty sure we've decided that
> the target construct added for declare target is actually not a dynamic
> property.  So basically mostly return to the 5.0 wording with clarifications
> for Fortran.  See
> https://github.com/OpenMP/spec/issues/2612#issuecomment-849742988
> for details.
> Making the target in construct dynamic would pretty much force all the
> scoring to be dynamic as well.

In that comment, Deepak says:

So, we decided to keep the target trait static, requiring that the declare 
target directive must be explicit and that the function version must be 
different from the version of the function that may be called outside of a 
target region (with the additional clarification that whether it differs or not 
will be implementation defined).

"the function version must be different from the version of the function that 
may be called outside of a target region": This is what we do not have in GCC at 
the moment - the function versions called within and outside target regions are 
the same on the host.

"whether it differs or not will be implementation defined": So whether a 
function with 'declare target' and a metadirective involving a 'target' 
construct behaves the same or not when called from both inside and outside of a 
target region is implementation defined?

I will leave the treatment of target constructs in the selector as it is then, 
with both calls going to the same function with the metadirective resolving to 
the 'target' variant. I will try to address your other concerns later.

Thanks

Kwok
Jakub Jelinek July 26, 2021, 9:23 p.m. UTC | #6 
On Mon, Jul 26, 2021 at 10:19:35PM +0100, Kwok Cheung Yeung wrote:
> > Yes, that is a target variant, but I'm pretty sure we've decided that
> > the target construct added for declare target is actually not a dynamic
> > property.  So basically mostly return to the 5.0 wording with clarifications
> > for Fortran.  See
> > https://github.com/OpenMP/spec/issues/2612#issuecomment-849742988
> > for details.
> > Making the target in construct dynamic would pretty much force all the
> > scoring to be dynamic as well.
> 
> In that comment, Deepak says:
> 
> So, we decided to keep the target trait static, requiring that the declare
> target directive must be explicit and that the function version must be
> different from the version of the function that may be called outside of a
> target region (with the additional clarification that whether it differs or
> not will be implementation defined).
> 
> "the function version must be different from the version of the function
> that may be called outside of a target region": This is what we do not have
> in GCC at the moment - the function versions called within and outside
> target regions are the same on the host.
> 
> "whether it differs or not will be implementation defined": So whether a
> function with 'declare target' and a metadirective involving a 'target'
> construct behaves the same or not when called from both inside and outside
> of a target region is implementation defined?
> 
> I will leave the treatment of target constructs in the selector as it is
> then, with both calls going to the same function with the metadirective
> resolving to the 'target' variant. I will try to address your other concerns
> later.

I think you're right, it should differ in the host vs. target version iff
it is in explicit declare target block, my memory is weak, but let's implement
the 5.0 wording for now (and ignore the 5.1 wording later on) and only when
we'll be doing 5.2 change this (and change for both metadirective and
declare variant at that point).
Ok?

	Jakub
Kwok Cheung Yeung July 26, 2021, 9:27 p.m. UTC | #7 
On 26/07/2021 10:23 pm, Jakub Jelinek wrote:
> On Mon, Jul 26, 2021 at 10:19:35PM +0100, Kwok Cheung Yeung wrote:
>> In that comment, Deepak says:
>>
>> So, we decided to keep the target trait static, requiring that the declare
>> target directive must be explicit and that the function version must be
>> different from the version of the function that may be called outside of a
>> target region (with the additional clarification that whether it differs or
>> not will be implementation defined).
>>
>> "the function version must be different from the version of the function
>> that may be called outside of a target region": This is what we do not have
>> in GCC at the moment - the function versions called within and outside
>> target regions are the same on the host.
>>
>> "whether it differs or not will be implementation defined": So whether a
>> function with 'declare target' and a metadirective involving a 'target'
>> construct behaves the same or not when called from both inside and outside
>> of a target region is implementation defined?
>>
>> I will leave the treatment of target constructs in the selector as it is
>> then, with both calls going to the same function with the metadirective
>> resolving to the 'target' variant. I will try to address your other concerns
>> later.
> 
> I think you're right, it should differ in the host vs. target version iff
> it is in explicit declare target block, my memory is weak, but let's implement
> the 5.0 wording for now (and ignore the 5.1 wording later on) and only when
> we'll be doing 5.2 change this (and change for both metadirective and
> declare variant at that point).
> Ok?
> 

Okay, the rest of the metadirective spec is quite enough to be getting on with 
for now. :-)

Thanks

Kwok
Kwok Cheung Yeung Dec. 10, 2021, 5:27 p.m. UTC | #8 
Hello

It has been several months since I posted my WIP patch, and my current 
patch set (which I will post separately) has evolved considerably since 
then. I have added C++ and Fortran support, as well as dynamic selectors 
from the OpenMP 5.1 spec (currently only the 'user={condition(<expr>)}' 
selector is implemented, target_device is TBD).

On 26/07/2021 3:29 pm, Jakub Jelinek wrote:
> Note, there is a partial overlap with the attribute syntax changes, see below.
> c-family/c-omp.c now has omp_directives table that should be updated for
> changes like this and then c_omp_categorize_directive that returns some
> information about the directives given a directive name (though, that name
> can be one, two or three tokens long, consider e.g. target enter data
> or cancellation point directives).

I have modified the C/C++ parser code to lookup the type of the 
directive using c_omp_categorize_directive.

> For metadirective, I think very special case are declarative directives in
> them, I'd tend to sorry for them at least for now, I'm pretty sure many
> cases with them are just unimplementable and will need to be restricted in
> the standard, others can be implemented with lots of effort.
> Whether it is e.g. metadirective guarding declare target ... end declare
> target pair that would only conditionally set declare target and instead of
> a single bit to find out if something is declare target or not we'd until
> resolved need to compute it for all possibilities, or e.g. conditional
> declare reduction/declare mapper where the name lookup for reduction or map
> directives would be dependent on metadirective resolution later on, etc.
> I'm afraid a total nightmare nobody has really thought about details for it.

The parsers currently emit a sorry if a C_OMP_DIR_DECLARATIVE directive 
is encountered in a metadirective, though I am sure there are many 
remaining ways that one could break it!

>> As an optimisation, identical body trees could be merged together, but that
>> can come later.
> 
> I'm afraid it isn't just an optimization and we need to be as smart as
> possible.  I'm not sure it is possible to parse everything many times,
> consider e.g. labels in the blocks, nested function definitions, variable
> definitions, etc.
> While OpenMP requires that essentially the code must be valid if the
> metadirective is replaced by any of those mentioned directives which rules
> quite some weirdo corner cases, nothing prevents e.g. two or more
> when directives to be standalone directives (which don't have any body and
> so whatever comes after them should be left parsed for later as normal
> statement sequence), one or more to be normal constructs that accept a
> structured block and one or more to be e.g. looping constructs (simd, for,
> distribute, taskloop or combined versions of those).
> Even when issues with labels etc. are somehow solved (e.g. for structured
> blocks we have the restriction that goto, break, continue, or switch into
> a case/default label, etc. can't be used to enter or exit the structured
> block which could mean some cases can be handled through renaming seen
> labels in all but one bodies), most important is to sync on where parsing
> should continue after the metadirective.
> I think it would be nice if the metadirective parsing at least made quick
> analysis on what kind of bodies the directives will want and can use the new
> c-omp.c infrastructure or if needed extend it (e.g. separate the C_OMP_DIR_CONSTRUCT
> category into C_OMP_DIR_CONSTRUCT and C_OMP_DIR_LOOPING_CONSTRUCT where
> the latter would be used for those that expect some omp loop after it).
> One option would be then to parse the body as the most restricted construct
> (looping (and determine highest needed collapse and ordered), then construct,
> then standalone) and be able to adjust what we parsed into what the
> different constructs need, but another option is the separate parsing of
> the code after the directive multiple times, but at least in the order of
> most restricted to least restricted, remember where to stop and don't parse
> it multiple times at least for directives that need the same thing.
>

After some experimentation, I'm not sure if it is possible in the 
general case to share bodies between variants. For one thing, it 
complicates the OMP region outlining and lowering, and becomes rather 
invasive to implement in the parser. Another is the possibility of 
having metadirectives nested within metadirective bodies. e.g. Something 
of the form:

#pragma omp metadirective \
     when (cond1: dir1) \
     when (cond2: dir2)
   {
     #pragma omp metadirective \
       when (construct dir1: dirA)
       when (construct dir2: dirB)
         (body)
   }

in which case the way the inner metadirective is resolved depends on the 
outer metadirective, leading to different bodies.

In my current patch set, I have implemented a limited form of statement 
body sharing when the body is not part of an OMP directive (e.g. an 'omp 
flush' followed by the body). Variables declarations and local functions 
in the body are handled by the usual scoping rules, and labels are 
handled by declaring them as __local__ (C and C++) or by renaming 
(Fortran). I have also added assertions in the parsers to ensure that 
each variant stops parsing at the same point. Would you find this 
acceptable?

>> 2) Selectors in the device set (i.e. kind, isa, arch) resolve differently
>> depending on whether the program is running on a target or on the host.
>> Since we don't keep multiple versions of a function for each target on the
>> host compiler, resolving metadirectives with these selectors needs to be
>> delayed until after LTO streaming, at which point the host or offload
>> compiler can make the appropriate decision.
> 
> How is this different from declare variant?  For declare variant, it is true
> I'm never trying to resolve it already during parsing of the call and that
> probably should be changed, do a first attempt at that point.  Initially
> I thought it typically will not be possible, but later clarification and
> strong desire of LLVM/ICC etc. to do everything or almost everything already
> during parsing suggests that it must be doable at least in some cases.
> E.g. we have restrictions that requires directive on which some decision
> could be dependent must appear only lexically before it or not at all, etc.
> So, similarly to that, metadirective ideally should see if something is
> impossible already during parsing (dunno if it should mean we wouldn't parse
> the body in that case, that would mean worse diagnostics), then repeat the
> checks during gimplification like declare variant is resolved there, then
> repeat again after IPA.  Would be probably best if for metadirectives that
> resolve to executable directives we represent it by something like a magic
> IFN that is told everything needed to decide and can share as much code as
> possible with the declare variant decisions.
> 
> It is true other compilers implement offloading quite differently from GCC,
> by repeating all of preprocessing, parsing etc. for the offloading target,
> so they can decide some metadirective/declare variant decisions earlier than
> we can.  On the other side that approach has also quite some disadvantages,
> it is much harder to ensure ABI compatibility between the host and offload
> code if one can use #ifdefs and whatever to change layout of everything in
> between.
> 
> For the checks during parsing, we'll need a different way how to track which
> directives are currently active (or defer anything with construct
> selectors till gimplification).  It is true that resolving that during
> parsing goes against the goal to parse as many bodies together as possible,
> so we need to pick one or the other.  Parsing what follows for all
> standalone directives isn't a problem of course, but if the metadirective
> has one when with for and another with simd, then parsing the loop just once
> would be a problem if there is metadirective in the body that wants to
> decide whether it is in for or simd and wants that decision be done during
> parsing.
> 

In my current patch, I attempt to resolve metadirectives at three points 
- during parsing, during Gimplification, and just after LTO.

For Fortran only, I skipped the parser resolution for now - I originally 
wanted to reuse the code from the C/C++ front ends to resolve 
metadirectives when translating from the Fortran parse tree to tree 
form, but there are quite a few references to C-family only functions in 
it (it would need to be rewritten to be more frontend-neutral).

Thanks,

Kwok
diff mbox series

Patch

diff --git a/gcc/Makefile.in b/gcc/Makefile.in
index 1164554e6d6..28e29fab93d 100644
--- a/gcc/Makefile.in
+++ b/gcc/Makefile.in
@@ -1505,6 +1505,7 @@  OBJS = \
 	omp-general.o \
 	omp-low.o \
 	omp-oacc-kernels-decompose.o \
+        omp-expand-metadirective.o \
 	omp-simd-clone.o \
 	opt-problem.o \
 	optabs.o \
diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c
index 4f8e8e0128c..01dc1e6d9c0 100644
--- a/gcc/c-family/c-pragma.c
+++ b/gcc/c-family/c-pragma.c
@@ -1312,12 +1312,14 @@  static const struct omp_pragma_def omp_pragmas[] = {
   { "allocate", PRAGMA_OMP_ALLOCATE },
   { "atomic", PRAGMA_OMP_ATOMIC },
   { "barrier", PRAGMA_OMP_BARRIER },
+  { "begin", PRAGMA_OMP_BEGIN },
   { "cancel", PRAGMA_OMP_CANCEL },
   { "cancellation", PRAGMA_OMP_CANCELLATION_POINT },
   { "critical", PRAGMA_OMP_CRITICAL },
   { "depobj", PRAGMA_OMP_DEPOBJ },
-  { "end", PRAGMA_OMP_END_DECLARE_TARGET },
+  { "end", PRAGMA_OMP_END },
   { "flush", PRAGMA_OMP_FLUSH },
+  { "metadirective", PRAGMA_OMP_METADIRECTIVE },
   { "requires", PRAGMA_OMP_REQUIRES },
   { "section", PRAGMA_OMP_SECTION },
   { "sections", PRAGMA_OMP_SECTIONS },
@@ -1387,6 +1389,41 @@  c_pp_lookup_pragma (unsigned int id, const char **space, const char **name)
   gcc_unreachable ();
 }
 
+static int
+c_pp_lookup_pragma_by_name_1 (const void *name, const void *elem)
+{
+  const struct omp_pragma_def *pragma_def
+    = (const struct omp_pragma_def *) elem;
+
+  return strcmp ((const char *) name, pragma_def->name);
+}
+
+enum pragma_kind
+c_pp_lookup_pragma_by_name (const char *name)
+{
+  const int n_omp_pragmas = sizeof (omp_pragmas) / sizeof (*omp_pragmas);
+  const int n_omp_pragmas_simd = sizeof (omp_pragmas_simd)
+				 / sizeof (*omp_pragmas_simd);
+
+  void *result = bsearch (name, omp_pragmas, n_omp_pragmas,
+			  sizeof (*omp_pragmas),
+			  c_pp_lookup_pragma_by_name_1);
+  if (!result)
+    result = bsearch (name, omp_pragmas_simd, n_omp_pragmas_simd,
+		      sizeof (*omp_pragmas_simd),
+		      c_pp_lookup_pragma_by_name_1);
+
+  if (result)
+    {
+      const struct omp_pragma_def *def
+	= (const struct omp_pragma_def *) result;
+
+      return (enum pragma_kind) def->id;
+    }
+
+  return PRAGMA_NONE;
+}
+
 /* Front-end wrappers for pragma registration to avoid dragging
    cpplib.h in almost everywhere.  */
 
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 6c34ffa5be4..6d4698d41ba 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -45,17 +45,19 @@  enum pragma_kind {
   PRAGMA_OMP_ALLOCATE,
   PRAGMA_OMP_ATOMIC,
   PRAGMA_OMP_BARRIER,
+  PRAGMA_OMP_BEGIN,
   PRAGMA_OMP_CANCEL,
   PRAGMA_OMP_CANCELLATION_POINT,
   PRAGMA_OMP_CRITICAL,
   PRAGMA_OMP_DECLARE,
   PRAGMA_OMP_DEPOBJ,
   PRAGMA_OMP_DISTRIBUTE,
-  PRAGMA_OMP_END_DECLARE_TARGET,
+  PRAGMA_OMP_END,
   PRAGMA_OMP_FLUSH,
   PRAGMA_OMP_FOR,
   PRAGMA_OMP_LOOP,
   PRAGMA_OMP_MASTER,
+  PRAGMA_OMP_METADIRECTIVE,
   PRAGMA_OMP_ORDERED,
   PRAGMA_OMP_PARALLEL,
   PRAGMA_OMP_REQUIRES,
@@ -252,6 +254,7 @@  extern enum cpp_ttype c_lex_with_flags (tree *, location_t *, unsigned char *,
 					int);
 
 extern void c_pp_lookup_pragma (unsigned int, const char **, const char **);
+extern enum pragma_kind c_pp_lookup_pragma_by_name (const char *);
 
 extern GTY(()) tree pragma_extern_prefix;
 
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index b9930d487fd..fa807530ca7 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -1583,8 +1583,12 @@  enum pragma_context { pragma_external, pragma_struct, pragma_param,
 static bool c_parser_pragma (c_parser *, enum pragma_context, bool *);
 static void c_parser_omp_cancellation_point (c_parser *, enum pragma_context);
 static bool c_parser_omp_target (c_parser *, enum pragma_context, bool *);
+static void c_parser_omp_begin (c_parser *, bool *);
+static void c_parser_omp_end (c_parser *);
 static void c_parser_omp_end_declare_target (c_parser *);
 static void c_parser_omp_declare (c_parser *, enum pragma_context);
+static tree c_parser_omp_metadirective (location_t, c_parser *, char *,
+					omp_clause_mask, tree *, bool *, bool);
 static void c_parser_omp_requires (c_parser *);
 static bool c_parser_omp_ordered (c_parser *, enum pragma_context, bool *);
 static void c_parser_oacc_routine (c_parser *, enum pragma_context);
@@ -12402,8 +12406,12 @@  c_parser_pragma (c_parser *parser, enum pragma_context context, bool *if_p)
     case PRAGMA_OMP_TARGET:
       return c_parser_omp_target (parser, context, if_p);
 
-    case PRAGMA_OMP_END_DECLARE_TARGET:
-      c_parser_omp_end_declare_target (parser);
+    case PRAGMA_OMP_BEGIN:
+      c_parser_omp_begin (parser, if_p);
+      return false;
+
+    case PRAGMA_OMP_END:
+      c_parser_omp_end (parser);
       return false;
 
     case PRAGMA_OMP_SCAN:
@@ -18195,6 +18203,7 @@  c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code,
   location_t for_loc;
   bool tiling = false;
   bool inscan = false;
+
   vec<tree, va_gc> *for_block = make_tree_vector ();
 
   for (cl = clauses; cl; cl = OMP_CLAUSE_CHAIN (cl))
@@ -20934,6 +20943,60 @@  c_parser_omp_end_declare_target (c_parser *parser)
     current_omp_declare_target_attribute--;
 }
 
+static void
+c_parser_omp_begin (c_parser *parser, bool *if_p)
+{
+  location_t loc = c_parser_peek_token (parser)->location;
+  c_parser_consume_pragma(parser);
+  if (c_parser_peek_token (parser)->type == CPP_NAME)
+    {
+      const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+
+      if (strcmp (p, "metadirective") == 0)
+	{
+	  char p_name[sizeof "#pragma omp teams distribute parallel for simd"];
+	  omp_clause_mask mask (0);
+
+	  c_parser_consume_token (parser);
+	  c_parser_omp_metadirective (loc, parser, p_name, mask, NULL, if_p,
+				      true);
+	  return;
+	}
+    }
+
+  error_at (loc, "expected %<begin metadirective%>");
+  c_parser_skip_to_pragma_eol (parser);
+}
+
+static void
+c_parser_omp_end (c_parser *parser)
+{
+  location_t loc = c_parser_peek_token (parser)->location;
+
+  if (c_parser_peek_2nd_token (parser)->type == CPP_NAME)
+    {
+      const char *p
+	= IDENTIFIER_POINTER (c_parser_peek_2nd_token (parser)->value);
+
+      if (strcmp (p, "declare") == 0)
+	{
+	  c_parser_omp_end_declare_target (parser);
+	  return;
+	}
+      else if (strcmp (p, "metadirective") == 0)
+	{
+	  /* The pragma 'omp end metadirective' should have been consumed
+	     when processing the metadirective.  */
+	  error_at (loc, "%<#pragma omp end metadirective%> without "
+			 "corresponding %<#pragma omp begin metadirective%>");
+	}
+    }
+  else
+    error_at (loc, "expected %<end declare target%> or %<end metadirective%>");
+
+  c_parser_skip_to_pragma_eol (parser);
+}
+
 
 /* OpenMP 4.0
    #pragma omp declare reduction (reduction-id : typename-list : expression) \
@@ -21607,6 +21670,295 @@  c_parser_omp_taskloop (location_t loc, c_parser *parser,
   return ret;
 }
 
+/* OpenMP 5.0:
+
+  # pragma omp metadirective [clause[, clause]]
+
+  # pragma omp begin metadirective [clause[, clause]]
+  # pragma omp end metadirective
+*/
+
+static tree
+c_parser_omp_metadirective (location_t loc, c_parser *parser,
+			    char *p_name, omp_clause_mask mask, tree *cclauses,
+			    bool *if_p,
+			    bool begin_end_p)
+{
+  tree ret;
+  bool all_selectors_resolveable = true;
+  auto_vec<auto_vec<c_token> > directive_tokens;
+  auto_vec<tree> ctxs;
+  bool default_seen = false;
+
+  ret = make_node (OMP_METADIRECTIVE);
+  SET_EXPR_LOCATION (ret, loc);
+  TREE_TYPE (ret) = void_type_node;
+  OMP_METADIRECTIVE_CLAUSES (ret) = NULL_TREE;
+  strcat (p_name, " metadirective");
+
+  while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL))
+    {
+      if (c_parser_next_token_is_not (parser, CPP_NAME)
+	  && c_parser_next_token_is_not (parser, CPP_KEYWORD))
+	{
+	  c_parser_error (parser, "expected %<when%> or %<default%>");
+	  return NULL_TREE;
+	}
+
+      location_t match_loc = c_parser_peek_token (parser)->location;
+      const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+      c_parser_consume_token (parser);
+      bool default_p = strcmp (p, "default") == 0;
+      if (default_p)
+	{
+	  if (default_seen)
+	    {
+	      c_parser_error (parser, "there can only be one default clause");
+	      return NULL_TREE;
+	    }
+	  else
+	    default_seen = true;
+	}
+      if (strcmp (p, "when") == 0 || default_p)
+	{
+	  matching_parens parens;
+	  tree ctx = NULL_TREE;
+	  bool skip = false;
+	  if (!parens.require_open (parser))
+	    return error_mark_node;
+
+	  if (!default_p)
+	    {
+	      ctx = c_parser_omp_context_selector_specification (parser,
+								 NULL_TREE);
+	      if (ctx == error_mark_node)
+		return NULL_TREE;
+	      ctx = c_omp_check_context_selector (match_loc, ctx);
+	      if (ctx == error_mark_node)
+		return NULL_TREE;
+
+	      switch (omp_context_selector_matches (ctx, true))
+		{
+		case -1:
+		  all_selectors_resolveable = false;
+		  break;
+		case 1:
+		  break;
+		case 0:
+		  /* Remove the selector from further consideration.  */
+		  skip = true;
+		  break;
+		}
+
+	      if (c_parser_next_token_is_not (parser, CPP_COLON))
+		{
+		  c_parser_error (parser, "expected colon");
+		  return NULL_TREE;
+		}
+	      c_parser_consume_token (parser);
+	    }
+
+	  /* Read in the directive type and create a dummy pragma token for
+	     it.  */
+	  c_token *token = c_parser_peek_token (parser);
+	  if (token->type != CPP_NAME)
+	    {
+	      c_parser_error (parser, "expected directive name");
+	      return NULL_TREE;
+	    }
+
+	  location_t loc = c_parser_peek_token (parser)->location;
+	  const char *p
+	    = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+	  enum pragma_kind p_kind = c_pp_lookup_pragma_by_name (p);
+
+	  c_parser_consume_token (parser);
+	  if (p_kind == PRAGMA_NONE)
+	    {
+	      c_parser_error (parser, "unknown directive name");
+	      return NULL_TREE;
+	    }
+
+	  if (!skip)
+	    {
+	      c_token pragma_token;
+	      pragma_token.type = CPP_PRAGMA;
+	      pragma_token.location = loc;
+	      pragma_token.pragma_kind = p_kind;
+
+	      directive_tokens.safe_push (auto_vec<c_token> ());
+	      directive_tokens.last ().safe_push (pragma_token);
+
+	      ctxs.safe_push (ctx);
+	    }
+
+	  /* Read in tokens for the directive clauses.  */
+	  auto_vec<c_token> *tokens = skip ? NULL : &directive_tokens.last ();
+	  int nesting_depth = 0;
+	  while (1)
+	    {
+	      c_token *token = c_parser_peek_token (parser);
+	      switch (token->type)
+		{
+		case CPP_EOF:
+		case CPP_PRAGMA_EOL:
+		  break;
+		case CPP_OPEN_PAREN:
+		  ++nesting_depth;
+		  goto add;
+		case CPP_CLOSE_PAREN:
+		  if (nesting_depth-- == 0)
+		    break;
+		  goto add;
+		default:
+		add:
+		  if (!skip)
+		    tokens->safe_push (*token);
+		  c_parser_consume_token (parser);
+		  continue;
+		}
+	      break;
+	    }
+
+	  c_parser_consume_token (parser);
+
+	  if (!skip)
+	    {
+	      c_token eol_token;
+	      memset (&eol_token, 0, sizeof (eol_token));
+	      eol_token.type = CPP_PRAGMA_EOL;
+	      tokens->safe_push (eol_token);
+	    }
+	}
+      else {
+	c_parser_error (parser, "expected %<when%> or %<default%>");
+	return NULL_TREE;
+      }
+    }
+  c_parser_skip_to_pragma_eol (parser);
+
+  /* Add the body tokens to the tokens for each candidate directive.  */
+  int nesting_depth = 0;
+  int bracket_depth = 0;
+  while (1)
+    {
+      int i;
+      auto_vec<c_token> *tokens;
+      c_token *token = c_parser_peek_token (parser);
+      bool stop = false;
+
+      if (begin_end_p)
+	{
+	  /* Keep reading until '#pragma end metadirective' is read.  */
+	  switch (token->type)
+	  {
+	  case CPP_PRAGMA:
+	    if (token->pragma_kind == PRAGMA_OMP_END)
+	      {
+		c_token *next_token = c_parser_peek_2nd_token (parser);
+		if (next_token->type == CPP_NAME
+		    && strcmp (IDENTIFIER_POINTER (next_token->value),
+			       "metadirective") == 0)
+		  {
+		    c_parser_consume_pragma (parser);
+		    c_parser_consume_token (parser);
+		    c_parser_skip_to_pragma_eol (parser);
+		    break;
+		  }
+	      }
+	  default:
+	    FOR_EACH_VEC_ELT (directive_tokens, i, tokens)
+	      tokens->safe_push (*token);
+	    if (token->type == CPP_PRAGMA)
+	      c_parser_consume_pragma (parser);
+	    else if (token->type == CPP_PRAGMA_EOL)
+	      c_parser_skip_to_pragma_eol (parser);
+	    else
+	      c_parser_consume_token (parser);
+	    continue;
+	  }
+	  break;
+	}
+      else
+	{
+	  switch (token->type)
+	    {
+	    case CPP_EOF:
+	      break;
+	    case CPP_OPEN_BRACE:
+	      ++nesting_depth;
+	      goto add2;
+	    case CPP_CLOSE_BRACE:
+	      if (--nesting_depth == 0)
+		stop = true;
+	      goto add2;
+	    case CPP_OPEN_PAREN:
+	      ++bracket_depth;
+	      goto add2;
+	    case CPP_CLOSE_PAREN:
+	      --bracket_depth;
+	      goto add2;
+	    case CPP_SEMICOLON:
+	      if (nesting_depth == 0 && bracket_depth == 0)
+		stop = true;
+	      goto add2;
+	    default:
+	    add2:
+	      FOR_EACH_VEC_ELT (directive_tokens, i, tokens)
+		tokens->safe_push (*token);
+	      if (token->type == CPP_PRAGMA)
+		c_parser_consume_pragma (parser);
+	      else if (token->type == CPP_PRAGMA_EOL)
+		c_parser_skip_to_pragma_eol (parser);
+	      else
+		c_parser_consume_token (parser);
+	      if (stop)
+		break;
+	      continue;
+	    }
+	  break;
+	}
+    }
+
+  /* Process each candidate directive.  */
+  auto_vec<c_token> *tokens;
+  int i;
+  FOR_EACH_VEC_ELT (directive_tokens, i, tokens)
+    {
+      /* Make sure nothing tries to read past the end of the tokens.  */
+      c_token eof_token;
+      memset (&eof_token, 0, sizeof (eof_token));
+      eof_token.type = CPP_EOF;
+      tokens->safe_push (eof_token);
+      tokens->safe_push (eof_token);
+
+      unsigned int tokens_avail = parser->tokens_avail;
+      gcc_assert (parser->tokens == &parser->tokens_buf[0]);
+      parser->tokens = tokens->address ();
+      parser->tokens_avail = tokens->length ();
+
+      tree block = c_begin_compound_stmt (false);
+      c_parser_omp_construct (parser, if_p);
+      block = c_end_compound_stmt (loc, block, false);
+
+      tree variant = build_tree_list (ctxs[i], block);
+      OMP_METADIRECTIVE_CLAUSES (ret)
+	= chainon (OMP_METADIRECTIVE_CLAUSES (ret), variant);
+
+      parser->tokens = &parser->tokens_buf[0];
+      parser->tokens_avail = tokens_avail;
+    }
+
+  if (all_selectors_resolveable)
+    {
+      ret = omp_resolve_metadirective (ret);
+      gcc_assert (ret != NULL_TREE);
+    }
+  add_stmt (ret);
+
+  return ret;
+}
+
 /* Main entry point to parsing most OpenMP pragmas.  */
 
 static void
@@ -21676,6 +22028,11 @@  c_parser_omp_construct (c_parser *parser, bool *if_p)
       strcpy (p_name, "#pragma omp");
       stmt = c_parser_omp_master (loc, parser, p_name, mask, NULL, if_p);
       break;
+    case PRAGMA_OMP_METADIRECTIVE:
+      strcpy (p_name, "#pragma omp");
+      stmt = c_parser_omp_metadirective (loc, parser, p_name, mask, NULL,
+					 if_p, false);
+      break;
     case PRAGMA_OMP_PARALLEL:
       strcpy (p_name, "#pragma omp");
       stmt = c_parser_omp_parallel (loc, parser, p_name, mask, NULL, if_p);
@@ -21713,7 +22070,6 @@  c_parser_omp_construct (c_parser *parser, bool *if_p)
     gcc_assert (EXPR_LOCATION (stmt) != UNKNOWN_LOCATION);
 }
 
-
 /* OpenMP 2.5:
    # pragma omp threadprivate (variable-list) */
 
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index f3503b13a5a..e2584c18571 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -43543,6 +43543,32 @@  cp_parser_omp_end_declare_target (cp_parser *parser, cp_token *pragma_tok)
     scope_chain->omp_declare_target_attribute--;
 }
 
+static void
+cp_parser_omp_end (cp_parser *parser, cp_token *pragma_tok)
+{
+  const char *p = "";
+  if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+    {
+      tree id = cp_lexer_peek_token (parser->lexer)->u.value;
+      p = IDENTIFIER_POINTER (id);
+    }
+  if (strcmp (p, "declare") == 0)
+    {
+      cp_parser_omp_end_declare_target (parser, pragma_tok);
+      return;
+    }
+  else if (strcmp (p, "metadirective") == 0)
+    /* The pragma 'omp end metadirective' should have been consumed
+       when processing the metadirective.  */
+    error_at (pragma_tok->location,
+	      "%<#pragma omp end metadirective%> without corresponding "
+	      "%<#pragma omp begin metadirective%>");
+  else
+    error_at (pragma_tok->location,
+	      "expected %<declare target%> or %<metadirective%>");
+  cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+}
+
 /* Helper function of cp_parser_omp_declare_reduction.  Parse the combiner
    expression and optional initializer clause of
    #pragma omp declare reduction.  We store the expression(s) as
@@ -45259,8 +45285,8 @@  cp_parser_pragma (cp_parser *parser, enum pragma_context context, bool *if_p)
       pop_omp_privatization_clauses (stmt);
       return ret;
 
-    case PRAGMA_OMP_END_DECLARE_TARGET:
-      cp_parser_omp_end_declare_target (parser, pragma_tok);
+    case PRAGMA_OMP_END:
+      cp_parser_omp_end (parser, pragma_tok);
       return false;
 
     case PRAGMA_OMP_SCAN:
diff --git a/gcc/gimple-low.c b/gcc/gimple-low.c
index fa7d4de30c0..b68b0cbb06c 100644
--- a/gcc/gimple-low.c
+++ b/gcc/gimple-low.c
@@ -234,6 +234,39 @@  lower_omp_directive (gimple_stmt_iterator *gsi, struct lower_data *data)
   gsi_next (gsi);
 }
 
+/* Lower the OpenMP metadirective statement pointed by GSI.  */
+
+static void
+lower_omp_metadirective (gimple_stmt_iterator *gsi, struct lower_data *data)
+{
+  gimple *stmt = gsi_stmt (*gsi);
+  gimple *body = gimple_omp_metadirective_bodies (stmt);
+  tree succ_label = create_artificial_label (UNKNOWN_LOCATION);
+  unsigned i;
+
+  for (i = 0; i < gimple_omp_metadirective_num_clauses (stmt); i++)
+    {
+      tree label = create_artificial_label (UNKNOWN_LOCATION);
+      gimple *g = gimple_build_label (label);
+
+      gsi_insert_after (gsi, g, GSI_CONTINUE_LINKING);
+      lower_sequence (gimple_omp_body_ptr (body), data);
+      gsi_insert_seq_after (gsi, gimple_omp_body (body), GSI_CONTINUE_LINKING);
+      gsi_insert_after (gsi, gimple_build_goto (succ_label),
+			GSI_CONTINUE_LINKING);
+      gimple_omp_metadirective_set_label (stmt, i, label);
+
+      body = body->next;
+    }
+
+  gsi_insert_after (gsi, gimple_build_label (succ_label),
+		    GSI_CONTINUE_LINKING);
+  gimple_omp_metadirective_set_succ_label (stmt, succ_label);
+  gimple_omp_metadirective_set_bodies (stmt, NULL);
+
+  gsi_next (gsi);
+}
+
 
 /* Lower statement GSI.  DATA is passed through the recursion.  We try to
    track the fallthruness of statements and get rid of unreachable return
@@ -398,6 +431,12 @@  lower_stmt (gimple_stmt_iterator *gsi, struct lower_data *data)
       data->cannot_fallthru = false;
       return;
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      data->cannot_fallthru = false;
+      lower_omp_metadirective (gsi, data);
+      data->cannot_fallthru = false;
+      return;
+
     case GIMPLE_TRANSACTION:
       lower_sequence (gimple_transaction_body_ptr (
 			as_a <gtransaction *> (stmt)),
diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c
index 0ef01e6420b..e7a4ba1171b 100644
--- a/gcc/gimple-pretty-print.c
+++ b/gcc/gimple-pretty-print.c
@@ -1978,6 +1978,64 @@  dump_gimple_omp_return (pretty_printer *buffer, const gimple *gs, int spc,
     }
 }
 
+/* Dump a GIMPLE_OMP_METADIRECTIVE tuple on the pretty_printer BUFFER.  */
+
+static void
+dump_gimple_omp_metadirective (pretty_printer *buffer, const gimple *gs,
+			       int spc, dump_flags_t flags)
+{
+  if (flags & TDF_RAW)
+    {
+      dump_gimple_fmt (buffer, spc, flags, "%G <%+BODY <%S> >", gs,
+		       gimple_omp_body (gs));
+    }
+  else
+    {
+      pp_string (buffer, "#pragma omp metadirective");
+      newline_and_indent (buffer, spc + 2);
+
+      gimple *body = gimple_omp_metadirective_bodies (gs);
+      bool has_bodies_p = body != NULL;
+      unsigned num_clauses = gimple_omp_metadirective_num_clauses (gs);
+
+      for (unsigned i = 0; i < num_clauses; i++)
+	{
+	  tree selector = gimple_omp_metadirective_selector (gs, i);
+
+	  if (selector == NULL_TREE)
+	    pp_string (buffer, "default:");
+	  else
+	    {
+	      pp_string (buffer, "when (");
+	      dump_generic_node (buffer, selector, spc, flags, false);
+	      pp_string (buffer, "):");
+	    }
+
+	  if (has_bodies_p)
+	    {
+	      newline_and_indent (buffer, spc + 4);
+	      pp_left_brace (buffer);
+	      pp_newline (buffer);
+	      dump_gimple_seq (buffer, gimple_omp_body (body), spc + 6, flags);
+	      newline_and_indent (buffer, spc + 4);
+	      pp_right_brace (buffer);
+
+	      body = body->next;
+	      if (body)
+		newline_and_indent (buffer, spc + 2);
+	    }
+	  else
+	    {
+	      tree label = gimple_omp_metadirective_label (gs, i);
+	      pp_string (buffer, " ");
+	      dump_generic_node (buffer, label, spc, flags, false);
+	      if (i != num_clauses - 1)
+		newline_and_indent (buffer, spc + 2);
+	    }
+	}
+    }
+}
+
 /* Dump a GIMPLE_TRANSACTION tuple on the pretty_printer BUFFER.  */
 
 static void
@@ -2729,6 +2787,12 @@  pp_gimple_stmt_1 (pretty_printer *buffer, const gimple *gs, int spc,
 				flags);
       break;
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      dump_gimple_omp_metadirective (buffer,
+				     as_a <const gomp_metadirective *> (gs),
+				     spc, flags);
+      break;
+
     case GIMPLE_CATCH:
       dump_gimple_catch (buffer, as_a <const gcatch *> (gs), spc, flags);
       break;
diff --git a/gcc/gimple-streamer-in.c b/gcc/gimple-streamer-in.c
index 1c979f438a5..66bbc2e8e0a 100644
--- a/gcc/gimple-streamer-in.c
+++ b/gcc/gimple-streamer-in.c
@@ -151,6 +151,7 @@  input_gimple_stmt (class lto_input_block *ib, class data_in *data_in,
     case GIMPLE_COND:
     case GIMPLE_GOTO:
     case GIMPLE_DEBUG:
+    case GIMPLE_OMP_METADIRECTIVE:
       for (i = 0; i < num_ops; i++)
 	{
 	  tree *opp, op = stream_read_tree (ib, data_in);
@@ -188,6 +189,10 @@  input_gimple_stmt (class lto_input_block *ib, class data_in *data_in,
 	  else
 	    gimple_call_set_fntype (call_stmt, stream_read_tree (ib, data_in));
 	}
+      if (gomp_metadirective *metadirective_stmt
+	    = dyn_cast <gomp_metadirective*> (stmt))
+	gimple_omp_metadirective_set_succ_label (metadirective_stmt,
+						 stream_read_tree (ib, data_in));
       break;
 
     case GIMPLE_NOP:
diff --git a/gcc/gimple-streamer-out.c b/gcc/gimple-streamer-out.c
index fcbf92300d4..f0ddd6a81a3 100644
--- a/gcc/gimple-streamer-out.c
+++ b/gcc/gimple-streamer-out.c
@@ -127,6 +127,7 @@  output_gimple_stmt (struct output_block *ob, struct function *fn, gimple *stmt)
     case GIMPLE_COND:
     case GIMPLE_GOTO:
     case GIMPLE_DEBUG:
+    case GIMPLE_OMP_METADIRECTIVE:
       for (i = 0; i < gimple_num_ops (stmt); i++)
 	{
 	  tree op = gimple_op (stmt, i);
@@ -169,6 +170,8 @@  output_gimple_stmt (struct output_block *ob, struct function *fn, gimple *stmt)
 	  else
 	    stream_write_tree (ob, gimple_call_fntype (stmt), true);
 	}
+      if (gimple_code (stmt) == GIMPLE_OMP_METADIRECTIVE)
+	stream_write_tree (ob, gimple_omp_metadirective_succ_label (stmt), true);
       break;
 
     case GIMPLE_NOP:
diff --git a/gcc/gimple-walk.c b/gcc/gimple-walk.c
index e4a55f1eeb6..ad7e1c0839e 100644
--- a/gcc/gimple-walk.c
+++ b/gcc/gimple-walk.c
@@ -674,6 +674,21 @@  walk_gimple_stmt (gimple_stmt_iterator *gsi, walk_stmt_fn callback_stmt,
 	return wi->callback_result;
       break;
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      {
+	gimple *body = gimple_omp_metadirective_bodies (stmt);
+
+	while (body)
+	  {
+	    ret = walk_gimple_seq_mod (gimple_omp_body_ptr (body),
+				       callback_stmt, callback_op, wi);
+	    if (ret)
+	      return wi->callback_result;
+	    body = body->next;
+	  }
+      }
+      break;
+
     case GIMPLE_WITH_CLEANUP_EXPR:
       ret = walk_gimple_seq_mod (gimple_wce_cleanup_ptr (stmt), callback_stmt,
 			     callback_op, wi);
diff --git a/gcc/gimple.c b/gcc/gimple.c
index f1044e9c630..f94009d39f3 100644
--- a/gcc/gimple.c
+++ b/gcc/gimple.c
@@ -1234,6 +1234,28 @@  gimple_build_omp_atomic_store (tree val, enum omp_memory_order mo)
   return p;
 }
 
+/* Build a GIMPLE_OMP_METADIRECTIVE statement.  */
+
+gomp_metadirective *
+gimple_build_omp_metadirective (int clause_count)
+{
+  gomp_metadirective *p
+    = as_a <gomp_metadirective *> (gimple_alloc (GIMPLE_OMP_METADIRECTIVE,
+						 clause_count * 2));
+  gimple_omp_metadirective_set_bodies (p, NULL);
+  return p;
+}
+
+
+gomp_metadirective_body *
+gimple_build_omp_metadirective_body (gimple_seq body)
+{
+  gomp_metadirective_body *m_body = as_a <gomp_metadirective_body *>
+    (gimple_alloc (GIMPLE_OMP_METADIRECTIVE_BODY, 0));
+  gimple_omp_set_body (m_body, body);
+  return m_body;
+}
+
 /* Build a GIMPLE_TRANSACTION statement.  */
 
 gtransaction *
diff --git a/gcc/gimple.def b/gcc/gimple.def
index 0ac0cf72bfa..1da68c16a91 100644
--- a/gcc/gimple.def
+++ b/gcc/gimple.def
@@ -384,6 +384,13 @@  DEFGSCODE(GIMPLE_OMP_TEAMS, "gimple_omp_teams", GSS_OMP_PARALLEL_LAYOUT)
    CLAUSES is an OMP_CLAUSE chain holding the associated clauses.  */
 DEFGSCODE(GIMPLE_OMP_ORDERED, "gimple_omp_ordered", GSS_OMP_SINGLE_LAYOUT)
 
+/* GIMPLE_OMP_METADIRECTIVE represents #pragma omp metadirective.  */
+DEFGSCODE(GIMPLE_OMP_METADIRECTIVE, "gimple_omp_metadirective",
+	  GSS_OMP_METADIRECTIVE)
+
+DEFGSCODE(GIMPLE_OMP_METADIRECTIVE_BODY, "gimple_omp_metadirective_body",
+	  GSS_OMP_METADIRECTIVE_BODY)
+
 /* GIMPLE_PREDICT <PREDICT, OUTCOME> specifies a hint for branch prediction.
 
    PREDICT is one of the predictors from predict.def.
diff --git a/gcc/gimple.h b/gcc/gimple.h
index 91b92b4a4d1..c5288af2bd7 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -822,6 +822,29 @@  struct GTY((tag("GSS_OMP_ATOMIC_STORE_LAYOUT")))
          stmt->code == GIMPLE_OMP_RETURN.  */
 };
 
+struct GTY((tag("GSS_OMP_METADIRECTIVE_BODY")))
+  gomp_metadirective_body : public gimple_statement_omp_single_layout
+{
+    /* No extra fields; adds invariant:
+       stmt->code == GIMPLE_OMP_METADIRECTIVE_BODY.  */
+};
+
+struct GTY((tag("GSS_OMP_METADIRECTIVE")))
+  gomp_metadirective : public gimple_statement_with_ops_base
+{
+  /* [ WORD 1-7 ] : base class */
+
+  /* [ WORD 8 ]  */
+  gomp_metadirective_body *bodies;
+
+  /* [ WORD 9 ] : a label after the metadirective
+     and all the candidate bodies  */
+  tree succ_label;
+
+  /* [ WORD 10 ] : operand vector.  */
+  tree GTY((length ("%h.num_ops"))) op[1];
+};
+
 /* GIMPLE_TRANSACTION.  */
 
 /* Bits to be stored in the GIMPLE_TRANSACTION subcode.  */
@@ -1233,6 +1256,22 @@  is_a_helper <gomp_task *>::test (gimple *gs)
   return gs->code == GIMPLE_OMP_TASK;
 }
 
+template <>
+template <>
+inline bool
+is_a_helper <gomp_metadirective *>::test (gimple *gs)
+{
+  return gs->code == GIMPLE_OMP_METADIRECTIVE;
+}
+
+template <>
+template <>
+inline bool
+is_a_helper <gomp_metadirective_body *>::test (gimple *gs)
+{
+  return gs->code == GIMPLE_OMP_METADIRECTIVE_BODY;
+}
+
 template <>
 template <>
 inline bool
@@ -1475,6 +1514,22 @@  is_a_helper <const gomp_task *>::test (const gimple *gs)
   return gs->code == GIMPLE_OMP_TASK;
 }
 
+template <>
+template <>
+inline bool
+is_a_helper <const gomp_metadirective *>::test (const gimple *gs)
+{
+  return gs->code == GIMPLE_OMP_METADIRECTIVE;
+}
+
+template <>
+template <>
+inline bool
+is_a_helper <const gomp_metadirective_body *>::test (const gimple *gs)
+{
+  return gs->code == GIMPLE_OMP_METADIRECTIVE_BODY;
+}
+
 template <>
 template <>
 inline bool
@@ -1572,6 +1627,8 @@  gomp_teams *gimple_build_omp_teams (gimple_seq, tree);
 gomp_atomic_load *gimple_build_omp_atomic_load (tree, tree,
 						enum omp_memory_order);
 gomp_atomic_store *gimple_build_omp_atomic_store (tree, enum omp_memory_order);
+gomp_metadirective *gimple_build_omp_metadirective (int clause_count);
+gomp_metadirective_body *gimple_build_omp_metadirective_body (gimple_seq body);
 gtransaction *gimple_build_transaction (gimple_seq);
 extern void gimple_seq_add_stmt (gimple_seq *, gimple *);
 extern void gimple_seq_add_stmt_without_update (gimple_seq *, gimple *);
@@ -1827,6 +1884,7 @@  gimple_has_substatements (gimple *g)
     case GIMPLE_OMP_TARGET:
     case GIMPLE_OMP_TEAMS:
     case GIMPLE_OMP_CRITICAL:
+    case GIMPLE_OMP_METADIRECTIVE:
     case GIMPLE_WITH_CLEANUP_EXPR:
     case GIMPLE_TRANSACTION:
       return true;
@@ -2479,12 +2537,21 @@  gimple_ops (gimple *gs)
 }
 
 
+/* Return true if GIMPLE statement G has any operands, including any that
+   should not be processed by the SSA passes.  */
+
+static inline bool
+gimple_has_ops_1 (const gimple *g)
+{
+  return gimple_has_ops (g) || gimple_code (g) == GIMPLE_OMP_METADIRECTIVE;
+}
+
 /* Return operand I for statement GS.  */
 
 static inline tree
 gimple_op (const gimple *gs, unsigned i)
 {
-  if (gimple_has_ops (gs))
+  if (gimple_has_ops_1 (gs))
     {
       gcc_gimple_checking_assert (i < gimple_num_ops (gs));
       return gimple_ops (CONST_CAST_GIMPLE (gs))[i];
@@ -2498,7 +2565,7 @@  gimple_op (const gimple *gs, unsigned i)
 static inline tree *
 gimple_op_ptr (gimple *gs, unsigned i)
 {
-  if (gimple_has_ops (gs))
+  if (gimple_has_ops_1 (gs))
     {
       gcc_gimple_checking_assert (i < gimple_num_ops (gs));
       return gimple_ops (gs) + i;
@@ -2512,7 +2579,7 @@  gimple_op_ptr (gimple *gs, unsigned i)
 static inline void
 gimple_set_op (gimple *gs, unsigned i, tree op)
 {
-  gcc_gimple_checking_assert (gimple_has_ops (gs) && i < gimple_num_ops (gs));
+  gcc_gimple_checking_assert (gimple_has_ops_1 (gs) && i < gimple_num_ops (gs));
 
   /* Note.  It may be tempting to assert that OP matches
      is_gimple_operand, but that would be wrong.  Different tuples
@@ -6330,6 +6397,77 @@  gimple_omp_continue_set_control_use (gomp_continue *cont_stmt, tree use)
   cont_stmt->control_use = use;
 }
 
+
+static inline tree
+gimple_omp_metadirective_succ_label (const gimple *g)
+{
+  const gomp_metadirective *omp_metadirective
+    = as_a <const gomp_metadirective *> (g);
+  return omp_metadirective->succ_label;
+}
+
+
+static inline void
+gimple_omp_metadirective_set_succ_label (gimple *g, tree succ_label)
+{
+  gomp_metadirective *omp_metadirective = as_a <gomp_metadirective *> (g);
+  omp_metadirective->succ_label = succ_label;
+}
+
+
+static inline gomp_metadirective_body *
+gimple_omp_metadirective_bodies (const gimple *g)
+{
+  const gomp_metadirective *omp_metadirective
+    = as_a <const gomp_metadirective *> (g);
+  return omp_metadirective->bodies;
+}
+
+
+static inline void
+gimple_omp_metadirective_set_bodies (gimple *g,
+				     gomp_metadirective_body *bodies)
+{
+  gomp_metadirective *omp_metadirective = as_a <gomp_metadirective *> (g);
+  omp_metadirective->bodies = bodies;
+}
+
+
+static inline unsigned
+gimple_omp_metadirective_num_clauses (const gimple *g)
+{
+  return gimple_num_ops (g) / 2;
+}
+
+
+static inline tree
+gimple_omp_metadirective_selector (const gimple *g, unsigned n)
+{
+  return gimple_op (g, n * 2);
+}
+
+
+static inline void
+gimple_omp_metadirective_set_selector (gimple *g, unsigned n, tree selector)
+{
+  return gimple_set_op (g, n * 2, selector);
+}
+
+
+static inline tree
+gimple_omp_metadirective_label (const gimple *g, unsigned n)
+{
+  return gimple_op (g, n * 2 + 1);
+}
+
+
+static inline void
+gimple_omp_metadirective_set_label (gimple *g, unsigned n, tree label)
+{
+  return gimple_set_op (g, n * 2 + 1, label);
+}
+
+
 /* Return a pointer to the body for the GIMPLE_TRANSACTION statement
    TRANSACTION_STMT.  */
 
@@ -6478,6 +6616,7 @@  gimple_return_set_retval (greturn *gs, tree retval)
     case GIMPLE_OMP_RETURN:			\
     case GIMPLE_OMP_ATOMIC_LOAD:		\
     case GIMPLE_OMP_ATOMIC_STORE:		\
+    case GIMPLE_OMP_METADIRECTIVE:		\
     case GIMPLE_OMP_CONTINUE
 
 static inline bool
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 2730f225187..49f4bcb01f3 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -5646,6 +5646,7 @@  is_gimple_stmt (tree t)
     case OMP_TASKGROUP:
     case OMP_ORDERED:
     case OMP_CRITICAL:
+    case OMP_METADIRECTIVE:
     case OMP_TASK:
     case OMP_TARGET:
     case OMP_TARGET_DATA:
@@ -13785,6 +13786,49 @@  gimplify_omp_ordered (tree expr, gimple_seq body)
   return gimple_build_omp_ordered (body, OMP_ORDERED_CLAUSES (expr));
 }
 
+static enum gimplify_status
+gimplify_omp_metadirective (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
+			    bool (*gimple_test_f) (tree), fallback_t fallback)
+{
+  gomp_metadirective_body *first_body = NULL;
+  gomp_metadirective_body *prev_body = NULL;
+  auto_vec<tree> selectors;
+  unsigned i;
+  tree clause = OMP_METADIRECTIVE_CLAUSES (*expr_p);
+
+  while (clause)
+    {
+      tree selector = TREE_PURPOSE (clause);
+      tree directive = TREE_VALUE (clause);
+
+      selectors.safe_push (selector);
+      gomp_metadirective_body *body
+	= gimple_build_omp_metadirective_body (NULL);
+      gimplify_stmt (&directive, gimple_omp_body_ptr (body));
+      if (!first_body)
+	first_body = body;
+      if (prev_body)
+	{
+	  prev_body->next = body;
+	  body->prev = prev_body;
+	}
+      prev_body = body;
+
+      clause = TREE_CHAIN (clause);
+    }
+
+  gomp_metadirective *stmt
+    = gimple_build_omp_metadirective (selectors.length ());
+  gimple_omp_metadirective_set_bodies (stmt, first_body);
+  gimplify_seq_add_stmt (pre_p, stmt);
+
+  tree selector;
+  FOR_EACH_VEC_ELT (selectors, i, selector)
+    gimple_omp_metadirective_set_selector (stmt, i, selector);
+
+  return GS_ALL_DONE;
+}
+
 /* Convert the GENERIC expression tree *EXPR_P to GIMPLE.  If the
    expression produces a value to be used as an operand inside a GIMPLE
    statement, the value will be stored back in *EXPR_P.  This value will
@@ -14680,6 +14724,11 @@  gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	  ret = gimplify_omp_atomic (expr_p, pre_p);
 	  break;
 
+	case OMP_METADIRECTIVE:
+	  ret = gimplify_omp_metadirective (expr_p, pre_p, post_p,
+					    gimple_test_f, fallback);
+	  break;
+
 	case TRANSACTION_EXPR:
 	  ret = gimplify_transaction (expr_p, pre_p);
 	  break;
diff --git a/gcc/gsstruct.def b/gcc/gsstruct.def
index 8f777e2bb95..f22ac1f65cb 100644
--- a/gcc/gsstruct.def
+++ b/gcc/gsstruct.def
@@ -50,4 +50,6 @@  DEFGSSTRUCT(GSS_OMP_SINGLE_LAYOUT, gimple_statement_omp_single_layout, false)
 DEFGSSTRUCT(GSS_OMP_CONTINUE, gomp_continue, false)
 DEFGSSTRUCT(GSS_OMP_ATOMIC_LOAD, gomp_atomic_load, false)
 DEFGSSTRUCT(GSS_OMP_ATOMIC_STORE_LAYOUT, gomp_atomic_store, false)
+DEFGSSTRUCT(GSS_OMP_METADIRECTIVE, gomp_metadirective, true)
+DEFGSSTRUCT(GSS_OMP_METADIRECTIVE_BODY, gomp_metadirective_body, false)
 DEFGSSTRUCT(GSS_TRANSACTION, gtransaction, false)
diff --git a/gcc/omp-expand-metadirective.cc b/gcc/omp-expand-metadirective.cc
new file mode 100644
index 00000000000..f4620df577a
--- /dev/null
+++ b/gcc/omp-expand-metadirective.cc
@@ -0,0 +1,125 @@ 
+/* Expand an OpenMP metadirective.
+
+   Copyright (C) 2021 Free Software Foundation, Inc.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include "config.h"
+#include "system.h"
+#include "coretypes.h"
+#include "backend.h"
+#include "target.h"
+#include "tree.h"
+#include "langhooks.h"
+#include "gimple.h"
+#include "tree-pass.h"
+#include "cgraph.h"
+#include "fold-const.h"
+#include "gimplify.h"
+#include "gimple-iterator.h"
+#include "gimple-walk.h"
+#include "gomp-constants.h"
+#include "omp-general.h"
+#include "diagnostic-core.h"
+#include "tree-cfg.h"
+#include "cfganal.h"
+
+static void
+omp_expand_metadirective (function *fun, basic_block bb)
+{
+  gimple *stmt = last_stmt (bb);
+  tree selected_label = omp_resolve_metadirective (stmt);
+
+  /* This is the last chance for the metadirective to be resolved.  */
+  if (!selected_label)
+    gcc_unreachable ();
+
+  /* Delete all variant BBs except for the selected one.  */
+  calculate_dominance_info (CDI_DOMINATORS);
+  for (unsigned i = 0; i < gimple_omp_metadirective_num_clauses (stmt); i++)
+    {
+      tree label = gimple_omp_metadirective_label (stmt, i);
+      edge edge = find_edge (bb, label_to_block (fun, label));
+      if (label == selected_label)
+	edge->flags |= EDGE_FALLTHRU;
+      else
+	remove_edge_and_dominated_blocks (edge);
+    }
+
+  /* Remove the metadirective statement.  */
+  gimple_stmt_iterator gsi = gsi_last_bb (bb);
+  gsi_remove (&gsi, true);
+}
+
+namespace {
+
+const pass_data pass_data_omp_expand_metadirective =
+{
+  GIMPLE_PASS, /* type */
+  "omp_expand_metadirective", /* name */
+  OPTGROUP_OMP, /* optinfo_flags */
+  TV_NONE, /* tv_id */
+  PROP_gimple_lcf, /* properties_required */
+  0, /* properties_provided */
+  0, /* properties_destroyed */
+  0, /* todo_flags_start */
+  0, /* todo_flags_finish */
+};
+
+class pass_omp_expand_metadirective : public gimple_opt_pass
+{
+public:
+  pass_omp_expand_metadirective (gcc::context *ctxt)
+    : gimple_opt_pass (pass_data_omp_expand_metadirective, ctxt)
+  {}
+
+  /* opt_pass methods: */
+  virtual bool gate (function *)
+  {
+    return (flag_openmp);
+  }
+
+  virtual unsigned int execute (function *fun);
+}; // class pass_omp_oacc_kernels_decompose
+
+unsigned int
+pass_omp_expand_metadirective::execute (function *fun)
+{
+  basic_block bb;
+  auto_vec<basic_block> metadirective_bbs;
+
+  FOR_EACH_BB_FN (bb, fun)
+    {
+      gimple *stmt = last_stmt (bb);
+      if (stmt && is_a<gomp_metadirective *> (stmt))
+	metadirective_bbs.safe_push (bb);
+    }
+
+  for (unsigned i = 0; i < metadirective_bbs.length (); i++)
+    omp_expand_metadirective (fun, metadirective_bbs[i]);
+
+  return 0;
+}
+
+} // anon namespace
+
+
+gimple_opt_pass *
+make_pass_omp_expand_metadirective (gcc::context *ctxt)
+{
+  return new pass_omp_expand_metadirective (ctxt);
+}
diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index 0f843bad79a..2c1affb64f8 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -9866,6 +9866,22 @@  expand_omp_target (struct omp_region *region)
     }
 }
 
+static void
+expand_omp_metadirective (struct omp_region *region)
+{
+  gomp_metadirective *stmt
+    = as_a <gomp_metadirective *> (last_stmt (region->entry));
+  tree succ_label = gimple_omp_metadirective_succ_label (stmt);
+  basic_block succ_bb = label_to_block (cfun, succ_label);
+  gimple_stmt_iterator gsi = gsi_start_bb (succ_bb);
+  while (!gsi_end_p (gsi)
+	 && gimple_code (gsi_stmt (gsi)) != GIMPLE_OMP_RETURN)
+    gsi_next (&gsi);
+
+  gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_RETURN);
+  gsi_remove (&gsi, true);
+}
+
 /* Expand the parallel region tree rooted at REGION.  Expansion
    proceeds in depth-first order.  Innermost regions are expanded
    first.  This way, parallel regions that require a new function to
@@ -9952,6 +9968,10 @@  expand_omp (struct omp_region *region)
 	  expand_omp_target (region);
 	  break;
 
+	case GIMPLE_OMP_METADIRECTIVE:
+	  expand_omp_metadirective (region);
+	  break;
+
 	default:
 	  gcc_unreachable ();
 	}
@@ -10388,6 +10408,24 @@  omp_make_gimple_edges (basic_block bb, struct omp_region **region,
 	}
       break;
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      /* Create an edge to the beginning of the body of each candidate
+	 directive.  */
+      {
+	unsigned i;
+	cur_region = new_omp_region (bb, code, cur_region);
+	gimple *stmt = last_stmt (bb);
+	for (i = 0; i < gimple_omp_metadirective_num_clauses (stmt); i++)
+	  {
+	    tree dest = gimple_omp_metadirective_label (stmt, i);
+	    basic_block dest_bb = label_to_block (cfun, dest);
+	    make_edge (bb, dest_bb, 0);
+	  }
+
+	fallthru = false;
+      }
+      break;
+
     default:
       gcc_unreachable ();
     }
diff --git a/gcc/omp-general.c b/gcc/omp-general.c
index a1bb9d8d25d..4839a9849bc 100644
--- a/gcc/omp-general.c
+++ b/gcc/omp-general.c
@@ -44,6 +44,7 @@  along with GCC; see the file COPYING3.  If not see
 #include "tree-iterator.h"
 #include "data-streamer.h"
 #include "streamer-hooks.h"
+#include "tree-pretty-print.h"
 
 enum omp_requires omp_requires_mask;
 
@@ -1100,8 +1101,13 @@  omp_context_name_list_prop (tree prop)
    others need to wait until the whole TU is parsed, others need to wait until
    IPA, others until vectorization.  */
 
+#define DELAY_METADIRECTIVES_AFTER_LTO { \
+  if (metadirective_p && !(cfun->curr_properties & PROP_gimple_lomp_dev))	\
+    return -1;	\
+}
+
 int
-omp_context_selector_matches (tree ctx)
+omp_context_selector_matches (tree ctx, bool metadirective_p)
 {
   int ret = 1;
   for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
@@ -1222,6 +1228,8 @@  omp_context_selector_matches (tree ctx)
 		    const char *arch = omp_context_name_list_prop (t3);
 		    if (arch == NULL)
 		      return 0;
+		    DELAY_METADIRECTIVES_AFTER_LTO;
+
 		    int r = 0;
 		    if (targetm.omp.device_kind_arch_isa != NULL)
 		      r = targetm.omp.device_kind_arch_isa (omp_device_arch,
@@ -1340,6 +1348,8 @@  omp_context_selector_matches (tree ctx)
 			  return 0;
 			continue;
 		      }
+		    DELAY_METADIRECTIVES_AFTER_LTO;
+
 		    int r = 0;
 		    if (targetm.omp.device_kind_arch_isa != NULL)
 		      r = targetm.omp.device_kind_arch_isa (omp_device_kind,
@@ -1379,6 +1389,8 @@  omp_context_selector_matches (tree ctx)
 		    const char *isa = omp_context_name_list_prop (t3);
 		    if (isa == NULL)
 		      return 0;
+		    DELAY_METADIRECTIVES_AFTER_LTO;
+
 		    int r = 0;
 		    if (targetm.omp.device_kind_arch_isa != NULL)
 		      r = targetm.omp.device_kind_arch_isa (omp_device_isa,
@@ -1445,6 +1457,8 @@  omp_context_selector_matches (tree ctx)
   return ret;
 }
 
+#undef DELAY_METADIRECTIVES_AFTER_LTO
+
 /* Compare construct={simd} CLAUSES1 with CLAUSES2, return 0/-1/1/2 as
    in omp_context_selector_set_compare.  */
 
@@ -2459,6 +2473,161 @@  omp_lto_input_declare_variant_alt (lto_input_block *ib, cgraph_node *node,
 						 INSERT) = entryp;
 }
 
+tree
+omp_resolve_metadirective (tree metadirective)
+{
+  auto_vec <tree, 16> clauses;
+  auto_vec <widest_int, 16> scores;
+  tree clause = OMP_METADIRECTIVE_CLAUSES (metadirective);
+  tree default_variant = NULL_TREE;
+
+  while (clause)
+    {
+      tree selector = TREE_PURPOSE (clause);
+      widest_int score;
+
+      if (selector == NULL_TREE)
+	default_variant = TREE_VALUE (clause);
+      else
+	switch (omp_context_selector_matches (selector, true))
+	  {
+	  case -1:
+	    return NULL_TREE;
+	  case 1:
+	    clauses.safe_push (clause);
+	    /* TODO: Handle SIMD score?  */
+	    omp_context_compute_score (selector, &score, false);
+	    scores.safe_push (score);
+	    break;
+	  case 0:
+	    break;
+	  }
+      clause = TREE_CHAIN (clause);
+    }
+
+  /* TODO: Handle case where there is no default.  */
+  if (clauses.is_empty ())
+    {
+      if (dump_file)
+	fprintf (dump_file, "Selecting default directive variant\n");
+      return default_variant;
+    }
+
+  /* A context selector that is a strict subset of another context selector
+     has a score of zero.  */
+  tree clause1, clause2;
+  unsigned int i, j;
+  FOR_EACH_VEC_ELT (clauses, i, clause1)
+    FOR_EACH_VEC_ELT_FROM (clauses, j, clause2, i + 1)
+      {
+	int r = omp_context_selector_compare (TREE_PURPOSE (clause1),
+					      TREE_PURPOSE (clause2));
+	if (r == -1)
+	  {
+	    /* ctx1 is a strict subset of ctx2.  */
+	    scores[i] = 0;
+	    break;
+	  }
+	else if (r == 1)
+	  /* ctx2 is a strict subset of ctx1.  */
+	  scores[j] = 0;
+      }
+
+  widest_int score, highest_score = -1;
+  FOR_EACH_VEC_ELT (scores, i, score)
+    if (score > highest_score)
+      {
+	highest_score = score;
+	clause = clauses[i];
+      }
+
+  if (dump_file)
+    {
+      fprintf (dump_file, "Selecting directive variant with selector:");
+      print_generic_expr (dump_file, TREE_PURPOSE (clause));
+      fprintf (dump_file, "\n");
+    }
+  return TREE_VALUE (clause);
+}
+
+tree
+omp_resolve_metadirective (gimple *gs)
+{
+  auto_vec <tree, 16> labels;
+  auto_vec <tree, 16> selectors;
+  auto_vec <widest_int, 16> scores;
+  tree default_label = gimple_omp_metadirective_succ_label (gs);
+
+  for (unsigned i = 0; i < gimple_omp_metadirective_num_clauses (gs); i++)
+    {
+      tree selector = gimple_omp_metadirective_selector (gs, i);
+      widest_int score;
+      if (selector == NULL_TREE)
+	default_label = gimple_omp_metadirective_label (gs, i);
+      else
+	switch (omp_context_selector_matches (selector, true))
+	  {
+	  case -1:
+	    return NULL;
+	  case 1:
+	    labels.safe_push (gimple_omp_metadirective_label (gs, i));
+	    selectors.safe_push (selector);
+	    /* TODO: Handle SIMD score?.  */
+	    omp_context_compute_score (selector, &score, false);
+	    scores.safe_push (score);
+	    break;
+	  case 0:
+	    break;
+	  }
+    }
+
+  if (scores.is_empty ())
+    {
+      if (dump_file)
+	fprintf (dump_file, "Selecting default directive variant\n");
+      return default_label;
+    }
+
+  /* A context selector that is a strict subset of another context selector
+     has a score of zero.  */
+  tree ctx1, ctx2;
+  unsigned int i, j;
+  FOR_EACH_VEC_ELT (selectors, i, ctx1)
+    FOR_EACH_VEC_ELT_FROM (selectors, j, ctx2, i + 1)
+      {
+	int r = omp_context_selector_compare (ctx1, ctx2);
+	if (r == -1)
+	  {
+	    /* ctx1 is a strict subset of ctx2.  */
+	    scores[i] = 0;
+	    break;
+	  }
+	else if (r == 1)
+	  /* ctx2 is a strict subset of ctx1.  */
+	  scores[j] = 0;
+      }
+
+  unsigned highest_index = 0;
+  widest_int score, highest_score = -1;
+  FOR_EACH_VEC_ELT (scores, i, score)
+  {
+    if (score > highest_score)
+      {
+	highest_score = score;
+	highest_index = i;
+      }
+  }
+
+  if (dump_file)
+    {
+      fprintf (dump_file, "Selecting directive variant with selector:");
+      print_generic_expr (dump_file, selectors[highest_index]);
+      fprintf (dump_file, "\n");
+    }
+
+  return labels[highest_index];
+}
+
 /* Encode an oacc launch argument.  This matches the GOMP_LAUNCH_PACK
    macro on gomp-constants.h.  We do not check for overflow.  */
 
diff --git a/gcc/omp-general.h b/gcc/omp-general.h
index aa04895e16d..47cea2eae01 100644
--- a/gcc/omp-general.h
+++ b/gcc/omp-general.h
@@ -104,10 +104,12 @@  extern tree find_combined_omp_for (tree *, int *, void *);
 extern poly_uint64 omp_max_vf (void);
 extern int omp_max_simt_vf (void);
 extern int omp_constructor_traits_to_codes (tree, enum tree_code *);
-extern int omp_context_selector_matches (tree);
+extern int omp_context_selector_matches (tree, bool = false);
 extern int omp_context_selector_set_compare (const char *, tree, tree);
 extern tree omp_get_context_selector (tree, const char *, const char *);
 extern tree omp_resolve_declare_variant (tree);
+extern tree omp_resolve_metadirective (tree);
+extern tree omp_resolve_metadirective (gimple *);
 extern tree oacc_launch_pack (unsigned code, tree device, unsigned op);
 extern tree oacc_replace_fn_attrib_attr (tree attribs, tree dims);
 extern void oacc_replace_fn_attrib (tree fn, tree dims);
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index d1136d181b3..c54000cfb9d 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -179,6 +179,10 @@  struct omp_context
   /* Only used for omp target contexts.  True if an OpenMP construct other
      than teams is strictly nested in it.  */
   bool nonteams_nested_p;
+
+  /* Only used for omp metadirectives.  Links to the next shallow
+     clone of this context.  */
+  struct omp_context *next_clone;
 };
 
 static splay_tree all_contexts;
@@ -964,6 +968,7 @@  new_omp_context (gimple *stmt, omp_context *outer_ctx)
   splay_tree_insert (all_contexts, (splay_tree_key) stmt,
 		     (splay_tree_value) ctx);
   ctx->stmt = stmt;
+  ctx->next_clone = NULL;
 
   if (outer_ctx)
     {
@@ -993,6 +998,17 @@  new_omp_context (gimple *stmt, omp_context *outer_ctx)
   return ctx;
 }
 
+static omp_context *
+clone_omp_context (omp_context *ctx)
+{
+  omp_context *clone_ctx = XCNEW (omp_context);
+
+  memcpy (clone_ctx, ctx, sizeof (omp_context));
+  ctx->next_clone = clone_ctx;
+
+  return clone_ctx;
+}
+
 static gimple_seq maybe_catch_exception (gimple_seq);
 
 /* Finalize task copyfn.  */
@@ -1039,6 +1055,14 @@  delete_omp_context (splay_tree_value value)
 {
   omp_context *ctx = (omp_context *) value;
 
+  /* Delete clones.  */
+  omp_context *clone = ctx->next_clone;
+  while (clone)
+    {
+      clone = clone->next_clone;
+      XDELETE (clone);
+    }
+
   delete ctx->cb.decl_map;
 
   if (ctx->field_map)
@@ -1073,6 +1097,7 @@  delete_omp_context (splay_tree_value value)
   delete ctx->lastprivate_conditional_map;
   delete ctx->allocate_map;
 
+
   XDELETE (ctx);
 }
 
@@ -3008,6 +3033,23 @@  scan_omp_teams (gomp_teams *stmt, omp_context *outer_ctx)
     ctx->record_type = ctx->receiver_decl = NULL;
 }
 
+/* Scan an OpenMP metadirective.  */
+
+static void
+scan_omp_metadirective (gomp_metadirective *stmt, omp_context *outer_ctx)
+{
+  gomp_metadirective_body *body = gimple_omp_metadirective_bodies (stmt);
+
+  while (body)
+    {
+      gimple_seq *body_p = gimple_omp_body_ptr (body);
+      omp_context *ctx = outer_ctx ? clone_omp_context (outer_ctx) : NULL;
+      scan_omp (body_p, ctx);
+
+      body = (gomp_metadirective_body *) body->next;
+    }
+}
+
 /* Check nesting restrictions.  */
 static bool
 check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
@@ -4045,6 +4087,10 @@  scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	scan_omp_teams (as_a <gomp_teams *> (stmt), ctx);
       break;
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      scan_omp_metadirective (as_a <gomp_metadirective *> (stmt), ctx);
+      break;
+
     case GIMPLE_BIND:
       {
 	tree var;
@@ -10129,6 +10175,22 @@  lower_omp_for_lastprivate (struct omp_for_data *fd, gimple_seq *body_p,
     }
 }
 
+static void
+lower_omp_metadirective (gimple_stmt_iterator *gsi_p, omp_context *ctx)
+{
+  gimple *stmt = gsi_stmt (*gsi_p);
+  gomp_metadirective_body *body = gimple_omp_metadirective_bodies (stmt);
+  while (body)
+    {
+      gimple_seq *body_p = gimple_omp_body_ptr (body);
+      omp_context *ctx = maybe_lookup_ctx (*body_p);
+      lower_omp (body_p, ctx);
+      body = (gomp_metadirective_body *) (body->next);
+    }
+  gsi_insert_after (gsi_p, gimple_build_omp_return (true),
+		    GSI_CONTINUE_LINKING);
+}
+
 /* Callback for walk_gimple_seq.  Find #pragma omp scan statement.  */
 
 static tree
@@ -13474,6 +13536,9 @@  lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       else
 	lower_omp_teams (gsi_p, ctx);
       break;
+    case GIMPLE_OMP_METADIRECTIVE:
+      lower_omp_metadirective (gsi_p, ctx);
+      break;
     case GIMPLE_CALL:
       tree fndecl;
       call_stmt = as_a <gcall *> (stmt);
diff --git a/gcc/passes.def b/gcc/passes.def
index 945d2bc797c..9aad498f266 100644
--- a/gcc/passes.def
+++ b/gcc/passes.def
@@ -186,6 +186,7 @@  along with GCC; see the file COPYING3.  If not see
   NEXT_PASS (pass_oacc_device_lower);
   NEXT_PASS (pass_omp_device_lower);
   NEXT_PASS (pass_omp_target_link);
+  NEXT_PASS (pass_omp_expand_metadirective);
   NEXT_PASS (pass_adjust_alignment);
   NEXT_PASS (pass_all_optimizations);
   PUSH_INSERT_PASSES_WITHIN (pass_all_optimizations)
diff --git a/gcc/tree-cfg.c b/gcc/tree-cfg.c
index 02256580c98..525e945a2d7 100644
--- a/gcc/tree-cfg.c
+++ b/gcc/tree-cfg.c
@@ -1668,6 +1668,23 @@  cleanup_dead_labels (void)
 	  }
 	  break;
 
+	case GIMPLE_OMP_METADIRECTIVE:
+	  {
+	    int i;
+	    for (i = 0; i < gimple_omp_metadirective_num_clauses (stmt); i++)
+	      {
+		label = gimple_omp_metadirective_label (stmt, i);
+		new_label = main_block_label (label, label_for_bb);
+		if (new_label != label)
+		  gimple_omp_metadirective_set_label (stmt, i, new_label);
+	      }
+	    label = gimple_omp_metadirective_succ_label (stmt);
+	    new_label = main_block_label (label, label_for_bb);
+	    if (new_label != label)
+	      gimple_omp_metadirective_set_succ_label (stmt, new_label);
+	  }
+	  break;
+
 	default:
 	  break;
       }
@@ -6078,6 +6095,22 @@  gimple_redirect_edge_and_branch (edge e, basic_block dest)
 				           gimple_block_label (dest));
       break;
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      {
+	for (unsigned i = 0; i < gimple_omp_metadirective_num_clauses (stmt); i++)
+	  {
+	    tree label = gimple_omp_metadirective_label (stmt, i);
+	    if (label_to_block (cfun, label) == e->dest)
+	      gimple_omp_metadirective_set_label (stmt, i,
+						  gimple_block_label (dest));
+	  }
+	tree label = gimple_omp_metadirective_succ_label (stmt);
+	if (label_to_block (cfun, label) == e->dest)
+	  gimple_omp_metadirective_set_succ_label (stmt,
+						   gimple_block_label (dest));
+      }
+      break;
+
     default:
       /* Otherwise it must be a fallthru edge, and we don't need to
 	 do anything besides redirecting it.  */
diff --git a/gcc/tree-inline.c b/gcc/tree-inline.c
index 8f945b88c12..3a21268b93b 100644
--- a/gcc/tree-inline.c
+++ b/gcc/tree-inline.c
@@ -4534,6 +4534,12 @@  estimate_num_insns (gimple *stmt, eni_weights *weights)
       return (weights->omp_cost
               + estimate_num_insns_seq (gimple_omp_body (stmt), weights));
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      /* The actual instruction will disappear eventually, so metadirective
+	 statements have zero cost.  */
+      gcc_assert (gimple_omp_body (stmt) == NULL);
+      return 0;
+
     case GIMPLE_TRANSACTION:
       return (weights->tm_cost
 	      + estimate_num_insns_seq (gimple_transaction_body (
diff --git a/gcc/tree-pass.h b/gcc/tree-pass.h
index 15693fee150..c02dda89f6a 100644
--- a/gcc/tree-pass.h
+++ b/gcc/tree-pass.h
@@ -418,6 +418,7 @@  extern gimple_opt_pass *make_pass_lower_switch_O0 (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_lower_vector (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_lower_vector_ssa (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_omp_oacc_kernels_decompose (gcc::context *ctxt);
+extern gimple_opt_pass *make_pass_omp_expand_metadirective (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_lower_omp (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_diagnose_omp_blocks (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_expand_omp (gcc::context *ctxt);
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 0a575eb9dad..5b601a303c6 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -3626,6 +3626,34 @@  dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags,
       is_expr = false;
       break;
 
+    case OMP_METADIRECTIVE:
+      {
+	pp_string (pp, "#pragma omp metadirective");
+	newline_and_indent (pp, spc + 2);
+	pp_left_brace (pp);
+
+	tree clause = OMP_METADIRECTIVE_CLAUSES (node);
+	while (clause != NULL_TREE)
+	  {
+	    newline_and_indent (pp, spc + 4);
+	    if (TREE_PURPOSE (clause) == NULL_TREE)
+	      pp_string (pp, "default:");
+	    else
+	      {
+		pp_string (pp, "when (");
+		dump_generic_node (pp, TREE_PURPOSE (clause), spc + 4, flags,
+				   false);
+		pp_string (pp, "):");
+	      }
+	    newline_and_indent (pp, spc + 6);
+	    dump_generic_node (pp, TREE_VALUE (clause), spc + 6, flags, false);
+	    clause = TREE_CHAIN (clause);
+	  }
+	newline_and_indent (pp, spc + 2);
+	pp_right_brace (pp);
+      }
+      break;
+
     case TRANSACTION_EXPR:
       if (TRANSACTION_EXPR_OUTER (node))
 	pp_string (pp, "__transaction_atomic [[outer]]");
diff --git a/gcc/tree-ssa-operands.c b/gcc/tree-ssa-operands.c
index c15575416dd..ee26451d717 100644
--- a/gcc/tree-ssa-operands.c
+++ b/gcc/tree-ssa-operands.c
@@ -978,6 +978,9 @@  operands_scanner::parse_ssa_operands ()
       append_vuse (gimple_vop (fn));
       goto do_default;
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      break;
+
     case GIMPLE_CALL:
       /* Add call-clobbered operands, if needed.  */
       maybe_add_call_vops (as_a <gcall *> (stmt));
diff --git a/gcc/tree.def b/gcc/tree.def
index eda050bdc55..9d50c739539 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1264,6 +1264,11 @@  DEFTREECODE (OMP_TARGET_ENTER_DATA, "omp_target_enter_data", tcc_statement, 1)
    Operand 0: OMP_TARGET_EXIT_DATA_CLAUSES: List of clauses.  */
 DEFTREECODE (OMP_TARGET_EXIT_DATA, "omp_target_exit_data", tcc_statement, 1)
 
+/* OpenMP - #pragma omp metadirective [clause1 ... clauseN]
+   Operand 0: OMP_METADIRECTIVE_CLAUSES: List of selectors and directive
+	variants.  */
+DEFTREECODE (OMP_METADIRECTIVE, "omp_metadirective", tcc_statement, 1)
+
 /* OMP_ATOMIC through OMP_ATOMIC_CAPTURE_NEW must be consecutive,
    or OMP_ATOMIC_SEQ_CST needs adjusting.  */
 
diff --git a/gcc/tree.h b/gcc/tree.h
index 64612cfa368..0d74cc75ce1 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1459,6 +1459,9 @@  class auto_suppress_location_wrappers
 #define OMP_TARGET_EXIT_DATA_CLAUSES(NODE)\
   TREE_OPERAND (OMP_TARGET_EXIT_DATA_CHECK (NODE), 0)
 
+#define OMP_METADIRECTIVE_CLAUSES(NODE) \
+  TREE_OPERAND (OMP_METADIRECTIVE_CHECK (NODE), 0)
+
 #define OMP_SCAN_BODY(NODE)	TREE_OPERAND (OMP_SCAN_CHECK (NODE), 0)
 #define OMP_SCAN_CLAUSES(NODE)	TREE_OPERAND (OMP_SCAN_CHECK (NODE), 1)