Re: [PATCH v7 2/5] OpenMP/OpenACC: Rework clause expansion and nested struct handling

[Tobias Burnus <tobias@codesourcery.com>]

Hi Julian,

first round of comments - I think I need a second pass as the patch is
long and complex. The depth of review also was decreasing, hence, I
assume I will spot things in later parts of the compiler.

In any case, I think the patch is a huge leap forward and very useful!

Contrary to previous iterations of the patch series, it looks as if
{1,2}/5 could be sensibly applied before {3,4,5}/5. However, those are
smaller or tiny, it should be way quicker to review them!

* * *

CROSS REFS

This patch is part of the patch set
"[PATCH v7 0/5] OpenMP/OpenACC: map clause and OMP gimplify rework"
https://gcc.gnu.org/pipermail/gcc-patches/2023-August/627895.html

This patch set in turn is required for the follow-up patch set
https://gcc.gnu.org/pipermail/gcc-patches/2023-September/629363.html
'[PATCH 0/8] OpenMP: lvalue parsing and "declare mapper" support'

and it looks as if it would be very useful for some other WIP patches
like map+iterator.

* * *

The patch 1/5 is just an indentation patch to make the following patch
smaller - and it has been approved multiple times.

This patch can be found at
   https://gcc.gnu.org/pipermail/gcc-patches/2023-August/627897.html
in case someone wants to read the full patch or Julian's patch-intro
comments.

* * *

I do note that all 5 patches of this patch set apply cleanly with some
fuzzy line matching, but:

WARNING: at least for this patch (2/5), code gets wrongly applied
in (c-)parser.cc - see below.

On 19.08.23 00:47, Julian Brown wrote:
> This patch reworks clause expansion in the C, C++ and (to a lesser
> extent) Fortran front ends for OpenMP and OpenACC mapping nodes used in
> GPU offloading support.
>
> At present a single clause may be turned into several mapping nodes,
> or have its mapping type changed, in several places scattered through
> the front- and middle-end.  The analysis relating to which particular
> transformations are needed for some given expression has become quite hard
> to follow.  Briefly, we manipulate clause types in the following places:

[...]

[As the patch is very large, I copied the interesting bits out of the patch
and comment on them - I hope that preserves enough context to be useful.]


   tree c_omp_address_inspector::get_address ()


This class member function is unused. It returns 'orig' and that member variable
is used a couple of times in the member functions - but not this function.

I think it can be removed. (I also find the name slightly misleading.)

* * *


+c_omp_address_inspector::maybe_zero_length_array_section (tree clause)
+{
+  switch (OMP_CLAUSE_MAP_KIND (clause))
+    {
+    case GOMP_MAP_ALLOC:
+    case GOMP_MAP_IF_PRESENT:
+    case GOMP_MAP_TO:
+    case GOMP_MAP_FROM:
+    case GOMP_MAP_TOFROM:
+    case GOMP_MAP_ALWAYS_TO:
+    case GOMP_MAP_ALWAYS_FROM:
+    case GOMP_MAP_ALWAYS_TOFROM:
+    case GOMP_MAP_RELEASE:
+    case GOMP_MAP_DELETE:
+    case GOMP_MAP_FORCE_TO:
+    case GOMP_MAP_FORCE_FROM:
+    case GOMP_MAP_FORCE_TOFROM:
+    case GOMP_MAP_FORCE_PRESENT:

Shouldn't this also include:

GOMP_MAP_ALWAYS_PRESENT_TO:
GOMP_MAP_ALWAYS_PRESENT_FROM:
GOMP_MAP_ALWAYS_PRESENT_TOFROM:
GOMP_MAP_PRESENT_ALLOC:
GOMP_MAP_PRESENT_TO:

?

* * *

+omp_expand_access_chain (tree c, tree expr, vec<omp_addr_token *> &addr_tokens,
+                        unsigned *idx)
+{
+  using namespace omp_addr_tokenizer;
+  location_t loc = OMP_CLAUSE_LOCATION (c);
+  unsigned i = *idx;
+  tree c2 = NULL_TREE;
+  gomp_map_kind kind
+    = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM
+       || (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+          && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM
+              || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DELETE
+              || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_RELEASE
+              || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_FROM
+              || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_FROM
+              || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PRESENT_FROM)))
+      ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH;

GOMP_MAP_ALWAYS_PRESENT_FROM  missing?

And I personally find it more readable to have
   if (...
       ...
       ...)
     kind = GOMP_MAP_DETACH;
   else
     kind = GOMP_MAP_ATTACH;

as in the assignment above, the lhs is in this case really far from the rhs.


* * *

I do see in omp_expand_access_chain, which is shared by C and C++, that the following
accesses are handled:

   ACCESS_POINTER,
   ACCESS_POINTER_OFFSET,
   ACCESS_INDEXED_ARRAY,

But not:
   ACCESS_REF_TO_POINTER,
   ACCESS_REF_TO_POINTER_OFFSET,
   ACCESS_INDEXED_REF_TO_ARRAY

Do those need to be handled as well?

SIDENOTE: As Fortran allocatable components, reference-type variables in a C++ class ('class', 'struct')
are automatically be copied when doing 'map(var)' - contrary to pointer members that need to be explicitly
be mapped (either at mapping side or via a mapper). (And contrary to Fortran allocatables, reference-type members
are always "there" - not like with allocatables can change from unallocated to allocated - with all compilations
of polymorphism, array sizes, string lengths, ...).

* * *

+tree
+c_omp_address_inspector::expand_array_base (tree c,
...
+  bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP;
+  bool implicit_p = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+                    && OMP_CLAUSE_MAP_IMPLICIT (c));
...
+  if (!openmp
+      && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP

How about using 'map_p' (twice), given that it is already
defined?

...

+       if (target)
+         {
+           c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+           if (target
+               && !ref_to_ptr

Remove 'target &&' from the second condition.


* * *

+++ b/gcc/c/c-parser.cc
@@ -17659,7 +17659,8 @@ c_parser_omp_clause_detach (c_parser *parser, tree list)

  static tree
  c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
-                          const char *where, bool finish_p = true)
+                          const char *where, bool finish_p = true,
+                          bool target = false)

Here but also in gcc/c-family/c-omp.cc - I wonder whether this should be target_p for
consistency with most other boolean variables. The code is not really consistent,
but given that there is 'finish_p' here and map_p etc. in the other file, it seems
to make sense to use 'target_p.

* * *

If I try the example at
   https://discourse.llvm.org/t/how-do-i-map-an-array-of-pointers-using-array-notation/53640/3
it fails to compile with:

foo.c:11:34: error: ‘#pragma omp target enter data’ with map-type other than ‘to’, ‘tofrom’ or ‘alloc’ on ‘map’ clause
    11 | #pragma omp target enter data map(to: a[i][0:1])


The problem is that fuzzy 'patch' mis-apply the patch for c-parser.cc's

@@ -22106,6 +22111,7 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
           case GOMP_MAP_FIRSTPRIVATE_POINTER:
           case GOMP_MAP_ALWAYS_POINTER:
           case GOMP_MAP_ATTACH_DETACH:
+         case GOMP_MAP_ATTACH:
             break;
           default:
             map_seen |= 1;
@@ -22215,6 +22221,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
           case GOMP_MAP_FIRSTPRIVATE_POINTER:
           case GOMP_MAP_ALWAYS_POINTER:
           case GOMP_MAP_ATTACH_DETACH:
+         case GOMP_MAP_DETACH:
             break;
           default:
             map_seen |= 1;

Namely after applying it, the result is:


@@ -23329,0 +23335 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
+         case GOMP_MAP_ATTACH:
@@ -23494,0 +23501 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
+         case GOMP_MAP_DETACH:

That is, the new code is added
   to c_parser_omp_target_data       and c_parser_omp_target_enter_data
instead of
   to c_parser_omp_target_enter_data and c_parser_omp_target_exit_data.

* * *

Likewise issue for C++'s cp/parser.cc.

* * *

The first question is whether such code is also needed for 'target' and 'target data'
besides 'target enter/exit data'. It looks as if omp_expand_access_chain should also
apply here.

* * *

And trying the following snippet:

void f(){
   int **a,i,j,n;

   #pragma omp target enter data map(a[i][j:n])
   #pragma omp target exit data map(a[i][j:n])

   #pragma omp target data map(a[i][j:n])
   {
   #pragma omp target map(a[i][j:n])
      ;
   }
}

Results - when putting the enums to enter/exit data
instead of to target data / target enter data:


g++ fails, unless I messed up, with:

fiof.c:4:36: error: ‘#pragma omp target exit data’ with map-type other than ‘from’, ‘tofrom’, ‘release’ or ‘delete’ on ‘map’ clause
     4 |   #pragma omp target exit data map(a[i][j:n])
       |                                    ^
fiof.c:5:31: error: ‘#pragma omp target data’ with map-type other than ‘to’, ‘from’, ‘tofrom’ or ‘alloc’ on ‘map’ clause
     5 |   #pragma omp target data map(a[i][j:n])


And gcc fails with:

fiof.c:3:36: error: ‘a’ appears more than once in map clauses
     3 |   #pragma omp target enter data map(a[i][j:n])
       |                                    ^
fiof.c:4:35: error: ‘a’ appears more than once in map clauses
     4 |   #pragma omp target exit data map(a[i][j:n])
       |                                   ^
fiof.c:5:30: error: ‘a’ appears more than once in map clauses
     5 |   #pragma omp target data map(a[i][j:n])
       |                              ^
fiof.c:7:25: error: ‘a’ appears more than once in map clauses
     7 |   #pragma omp target map(a[i][j:n])

(Plus all '... data' ones: '... must contain at least one ...')


Can you fix this - and add an example in the spirit of the LLVM example above,
except that it should test all four 'target' variants.

Thanks,

Tobias

-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955