[Bug middle-end/112779] New: [OpenMP] Support omp Metadirectives

[Bug middle-end/112779] New: [OpenMP] Support omp Metadirectives

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=112779

            Bug ID: 112779
           Summary: [OpenMP] Support omp Metadirectives
           Product: gcc
           Version: 14.0
            Status: UNCONFIRMED
          Severity: normal
          Priority: P3
         Component: middle-end
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
  Target Milestone: ---

There is a rather complete support for metadirectives on the OG13 branch, i.e.
devel/omp/gcc-13 branch.

Several patches have been already posted.

* * *

Known issues with those patches:

(A) OpenMP 5.2's renaming of clause 'default' is now (also) 'otherwise'

(B) -------------------------------------------
ICE (segfault) for "kind: nohost" in default()

og12-offload/testlogs-2023-05-04 shows several 'internal compiler error:
Segmentation fault' for the 'default' clause of the metadirectives

→ sollve_vv's test_metadirective_target_device{,_kind{,_any}}.c testcases.

The problem for one test case at least is in omp-general.cc's omp_dynamic_cond
:

'kind_sel' = {purpose = "kind", value = "{ purpose: "nohost", value: NULL}" } -
and accessing TREE_VALUE (TREE_VALUE (kind_sel)).

That's for the following code:



      tree kind_sel = omp_get_context_selector (ctx, "target_device", "kind");
      if (kind_sel)
        {
          const char *str
            = (TREE_VALUE (TREE_VALUE (kind_sel))
               ? TREE_STRING_POINTER (TREE_VALUE (TREE_VALUE (kind_sel)))
               : IDENTIFIER_POINTER (TREE_PURPOSE (TREE_VALUE (kind_sel))));
          kind =  build_string_literal (strlen (str) + 1, str);
        } 



I wonder why that's not already handled in gcc/omp-general.cc's
omp_context_selector_matches (which has some code) – but it might indeed only
be available at run time?!?


(C) ------------------------------------------------------

wonder whether libgomp/target.c's GOMP_evaluate_target_device lacks a check for
kind == "nohost" — I only see "host" (for the host) and "gpu" (for the GPU) and
the generic "any".


(D) ---------------------------------------------------------

Fortran's DO with do-end-label{}

Fortran permits loops with label instead of a simple END DO, example:



DO 123 i=1,5
  DO 123 j=1,5
123 CONTINUE  ! or '123 END DO' — Note the *shared* end-do-label (which
invalid/deleted since F2018, before deprecated but valid)

Such code is not handled with metadirectives as already indicated at
gcc/fortran/parse.cc's parse_omp_metadirective_body:





        case_omp_do:
          st = parse_omp_do (clause->stmt);
          /* TODO: Does st == ST_IMPLIED_ENDDO need special handling?  */
          break; 

The answer seems to be yes — failing testcase is the following ("Error: END DO
statement expected"):



implicit none
integer :: i, j, psi(5,5)!$omp  metadirective  &
!$omp&    when(user={condition(.false. )}: target teams  &
!$omp&         distribute parallel do simd collapse(2))  &
!$omp&    when(user={condition(.false. )}: target teams  &
!$omp&         distribute parallel do)  &
!$omp&    default(target teams loop collapse(2))
      DO 50 I=1,5
!$omp  metadirective  &
!$omp& when(user={condition(.false. )}: simd)               
      DO 51 J=1,5
        PSI(j,i) = j
   51 CONTINUE
   50 CONTINUE
end 


(E) -------------------------------------------------------

internal compiler error: in c_parser_omp_metadirective, at c/c-parser.cc:26565
   11 |  #pragma omp metadirective when (user = { condition (USE_GPU == 1) } :
target enter data map(alloc : number[ : SIZE]))

for:

#include <stdio.h>
#include <omp.h>

int main(int argc, char ** argv){

    const int SIZE = 10; 
    int USE_GPU = 1;
    double number[SIZE];
    double *number_d;
 #pragma omp metadirective when (user = { condition (USE_GPU == 1) } : target
enter data map(alloc : number[ : SIZE]))
    if (USE_GPU)
        number_d = (double *)omp_get_mapped_ptr(number,
omp_get_default_device());
    else
        number_d = number;
    printf("number_d = %p    number= %p\n", number_d, number);

    return 0;
}


(F) ----------------------------------------------------------------

For C/C++, begin/end metadirective is not handled - it is for Fortran, where it
is much more useful.

Note: It is less useful that it sounds. From an internal bug tracker:

This was a deliberate design decision. From the OpenMP 5.0 spec (2.3.4):

"The begin metadirective directive behaves identically to the metadirective
directive, except that the directive syntax for the specified directive
variants must accept a paired end directive."

so having 'target enter data' in a 'begin metadirective' is invalid.

The only OpenMP directive supported in GCC that takes an end directive in C/C++
is 'declare target' (is this still true?), and we have already said that we
would not support declarative constructs in metadirectives. So the 'begin/end
metadirective' support was left out in C/C++.

(Invalid) testcase:

#pragma omp begin metadirective  when (user = { condition (USE_GPU == 1) } )
#pragma omp  target enter data map(alloc : number[ : SIZE])
    ;   
#pragma omp end metadirective

I get
$ gcc -O3 -fopenmp -foffload="-march=gfx90a -lm" -lm  ./test_meta_user.c
./test_meta_user.c: In function 'main':
./test_meta_user.c:11:19: error: expected 'declare target' or 'assumes' before
'metadirective'
   11 | #pragma omp begin metadirective  when (user = { condition (USE_GPU ==
1) } )
      |                   ^~~~~~~~~~~~~
./test_meta_user.c:14:17: error: expected 'declare' or 'assumes' before
'metadirective'
   14 | #pragma omp end metadirective
      |                 ^~~~~~~~~~~~~
----------------

Note:
In C/C++, only delimited directives have a begin/end + and only 'begin/end
assumes' seems to be useful.

Details:
Quoting the OpenMP spec (git version; 5.0 and later have similar wording):

"The restrictions to begin metadirective are as follows:
• Any directive-variant that is specified by a when or otherwise clause must be
a directive that has a paired end directive or must be the nothing directive."

(This applies to both Fortran and C/C++, but while in Fortran any structured
block can have an end directive, only a few directives in C/C++ can have an end
directive..)

In case of "target enter data", it is clear that this does not have any paired
end directive - neither in C/C++ nor in Fortran.


Otherwise, OpenMP 5.0 and 5.1 makes it very simple to search for '#pragma omp
end'. This turns up the following directives that have an end-directive in
C/C++:

#pragma omp begin/end metadirective  - OMP 5.0
#pragma omp (end) declare target  - OMP 5.0
#pragma omp begin/end declare target  - OMP 5.1
#pragma omp begin/end declare variant  - OMP 5.1
#pragma omp begin/end assumes  - OMP 5.1


Additionally, there is - quoting 5.1:

"Restrictions to metadirectives are as follows:
• The directive variant appearing in a when or default clause must not specify
a metadirective, begin metadirective, or end metadirective directive.
C / C++
• The directive variant that appears in a when or default clause must not
specify a begin declare variant or end declare variant"


Thus, that leaves for 5.0 only "omp begin declare target" and for/since 5.1
additionally "omp begin assumes".


In 5.2 and later, the restrictions moved to the when and otherwise (=
"default") clauses:

  "directive-variant must not specify a metadirective."
  "directive-variant must not specify a begin declare variant directive"


Additionally, there is:

• If the list of dynamic replacement candidates has multiple items then all
items must be executable directives.

Which in turn disallows declarative directives unless there is a single dynamic
replacement candidate - implying that the dynamic context selector does not
need to be evaluated in that case.


(G) ------------------------------------------------------------

ICE on invalid:

Regarding the following testcase, I believe the following is invalid:

    !$OMP metadirective &
    !$OMP   when ( user = { condition ( UseDevice ) } &
    !$OMP     : end target teams distribute parallel do simd ) &
    !$OMP   default ( end parallel do simd ) 

Namely, neither "end target teams ..." nor "end parallel ..." is a valid
directive-specification — and the latter starts with a valid directive-name –
and "end" isn't one nor is 'end target' nor ...

Hence, I believe that OpenMP requires for directives that require an
end-directive to use metadirective's delimited form, i.e. 'omp begin
metadirective" ... "omp end metadirective".

----------
program OpenMP_Metadirective_WrongEnd_Test
  integer :: &
    iV, jV, kV
  integer, dimension ( 3 ) :: &
    lV, uV
  logical :: &
    UseDevice

    UseDevice = .False.

    !$OMP metadirective &
    !$OMP   when ( user = { condition ( UseDevice ) } &
    !$OMP     : target teams distribute parallel do simd collapse ( 3 ) &
    !$OMP         private ( iaVS ) ) &
    !$OMP   default ( parallel do simd collapse ( 3 ) private ( iaVS ) )
    do kV = lV ( 3 ), uV ( 3 ) 
      do jV = lV ( 2 ), uV ( 2 ) 
        do iV = lV ( 1 ), uV ( 1 ) 


        end do
      end do
    end do
    !$OMP metadirective &
    !$OMP   when ( user = { condition ( UseDevice ) } &
    !$OMP     : end target teams distribute parallel do simd ) &
    !$OMP   default ( end parallel do simd ) 
end program
----------


(H) ---------------------------------------------

The following gives also an ICE:

program test
  implicit none  integer, parameter :: N = 100
  integer :: x(N), y(N), z(N)  block
    integer :: k
    !$omp metadirective &
                !$omp& when(device={arch("nvptx")}: teams loop) &
                !$omp& default(parallel loop)
    do i = 1, N
          z(i) = x(i) * y(i)
    enddo
   end block
end