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