[Bug fortran/104696] New: [12 Regression][OpenMP] Implicit mapping breaks struct mapping

[Bug fortran/104696] New: [12 Regression][OpenMP] Implicit mapping breaks struct mapping

[burnus at gcc dot gnu.org]

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

            Bug ID: 104696
           Summary: [12 Regression][OpenMP] Implicit mapping breaks struct
                    mapping
           Product: gcc
           Version: 12.0
            Status: UNCONFIRMED
          Keywords: openmp, wrong-code
          Severity: normal
          Priority: P3
         Component: fortran
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
                CC: cltang at gcc dot gnu.org, jakub at gcc dot gnu.org
  Target Milestone: ---

Created attachment 52519
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=52519&action=edit
implicit.f90 - compile with -fopenmp and run with a non-shared-memory device.

In my understanding, the following is valid and should work.

However, it fails – and I bet it is due to the implicit mapping of 'var3'. The
code uses:

!$omp target map(to: var3%R(2)%d)
 ...

Printing on the host the  'loc (var3%R(2)%d)'  shows:
         21516C0
    7F51E5E001D8
    STOP 11

As var%R(2)%d  (or in the dump 'var3.r[1].d.data')  is now in device address
space,
accessing it after the target region crashes the program.


implicit.f90.005t.original:
  #pragma omp target
     map(to:var3.r[1].d [len: 88])
     map(to:*(struct t2[0:] *) var3.r[1].d.data [len: D.4243 * 4])
     map(always_pointer:(struct t2[0:] *) var3.r[1].d.data [pointer assign,
bias: 0])

implicit.f90.006t.gimple:
  #pragma omp target num_teams(1) thread_limit(0)
     map(tofrom:var3 [len: 440][implicit])
     map(to:var3.r[1].d [len: 88])
     map(to:MEM <struct t2[0:]> [(struct t2[0:] *)_9] [len: _8])
     map(always_pointer:var3.r[1].d.data [pointer assign, bias: 0])

[Bug fortran/104696] [12 Regression][OpenMP] Implicit mapping breaks struct mapping

[rguenth at gcc dot gnu.org]

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

Richard Biener <rguenth at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
   Target Milestone|---                         |12.0

[Bug fortran/104696] [OpenMP] Implicit mapping breaks struct mapping

[burnus at gcc dot gnu.org]

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

Tobias Burnus <burnus at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
            Summary|[12 Regression][OpenMP]     |[OpenMP] Implicit mapping
                   |Implicit mapping breaks     |breaks struct mapping
                   |struct mapping              |
           Keywords|                            |missed-optimization

--- Comment #1 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Looking closer at it, I no longer think it is a regression (to be checked if
deemed important).

But it looks as if there are two problems - one wrong-code one and one
missed-optimization one.

Namely, I think reason for both issues is that 
     map(to:var3.r[1].d [len: 88]) ...
is not turned into
     map(struct:var3 [len: 1]) map(to:var3.r[1].d [len: 88])
but into 'to' + pointer assign. This does not even work when using the 'always'
modifier.

('struct:' appears when using 'Q' instead of 'R(2)'.)

That that the struct not detected seems to be because array and component refs
are mixed – hiding that the var is memory wise in the same struct.

I believe with that fixed, it would work correctly.

  * * *

For C/C++, it "works".

But: it still does not detect that the member is part of the whole struct - and
allocates pointlessly too much memory. To illustrate this, I added a large
'arr' element.

Otherwise, that the reason that it works in C/C++ is that ATTACH and not
pointer assign is used.

Namely:
------------ test.c--------------
struct s { int *d; };
struct s2 { struct s r[5], q, arr[1024][1024]; };

int main () {
  struct s2 x;
  x.q.d = __builtin_malloc(sizeof(int));
  x.r[1].d = __builtin_malloc(sizeof(int));

  #pragma omp target map(tofrom: x.q.d[:1])
   *x.q.d = 2;
  #pragma omp target map(tofrom: x.r[1].d[:1])
   *x.r[1].d = 3;

  __builtin_printf("%d, %d\n", *x.q.d, *x.r[1].d);
  return 0;
}
------------ test.c--------------

gives:

  #pragma omp target num_teams(1) thread_limit(0)
     map(tofrom:x [len: 8388656][implicit]) map(tofrom:*_3 [len: 4])
     map(attach:x.q.d [bias: 0])

  #pragma omp target num_teams(1) thread_limit(0)
     map(tofrom:x [len: 8388656][implicit])
     map(tofrom:*_4 [len: 4]) map(attach:x.r[1].d [bias: 0])

[Bug fortran/104696] [OpenMP] Implicit mapping breaks struct mapping

[burnus at gcc dot gnu.org]

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

--- Comment #2 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #1)
> Namely:
> ------------ test.c--------------
> struct s { int *d; };

It makes more sense to use 'int d;' to match Fortran. Doing so yields in the
gimple dump:

#pragma omp target num_teams(1) thread_limit(0)
  map(struct:x [len: 1]) map(tofrom:x.q.d [len: 4])
#pragma omp target num_teams(1) thread_limit(0)
  map(tofrom:x [len: 4194328][implicit]) map(tofrom:x.r[1].d [len: 4])

Thus, C and Fortran show the same issue.

[Bug fortran/104696] [OpenMP] Implicit mapping breaks struct mapping

[cltang at gcc dot gnu.org]

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

--- Comment #3 from Chung-Lin Tang <cltang at gcc dot gnu.org> ---
The problem, at a low-level, is that the Fortran FE is using always_pointer,
while C/C++ front-ends generate attach_detach (and turns in attach after
gimplify).

attach and always_pointer both modify the GPU-side field into the correct GPU
pointer, but only attach will "clean-up" at the end, restoring the host-side
pointer contents, before the full implicit struct mapping copies back to the
host (due to tofrom map-kind).

Since Fortran array descriptors are like little structs, I think quickest fix
is to make the Fortran FE use 'attach_detach' for the array data pointer field.
That should immediately make things work.

A more elaborate way is to try to turn off the creation of the implicit mapping
for the whole array descriptor, which should be a killer of performance. (not
sure when this started, although prior changes did put implicit mappings to the
very front of the sequence)

[Bug fortran/104696] [OpenMP] Implicit mapping breaks struct mapping

[burnus at gcc dot gnu.org]

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

--- Comment #4 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Chung-Lin Tang from comment #3)
> The problem, at a low-level, is that the Fortran FE is using always_pointer,
> while C/C++ front-ends generate attach_detach (and turns in attach after
> gimplify).

Actually, if one modifies the C testcase, as alluded in previous comment 2, one
runs into the same issue. Namely, for:

------------ test.c--------------
struct s { int d; };
struct s2 { struct s r[5], q, arr[1024][1024]; };

int main () {
  struct s2 x;
  #pragma omp target map(tofrom: x.q.d)
   x.q.d = 2;
  #pragma omp target map(tofrom: x.r[1].d)
   *.r[1].d = 3;
  __builtin_printf("%d, %d\n", x.q.d, x.r[1].d);
  return 0;
}
------------ test.c--------------

The expect result is 'struct: x' + 'tofrom: x.r[1].d' - and as 'struct:x'
exists, no implicit mapping of 'x' happens.

That happens for 'x.q.d':

  map(struct:x [len: 1]) map(tofrom:x.q.d [len: 4])

but not for 'x.r[1].d':

  map(tofrom:x [len: 4194328][implicit]) map(tofrom:x.r[1].d [len: 4])

Thus the code which constructs 'struct: x' has trouble for 'component ref +
array ref + component ref'.

[Bug middle-end/104696] [OpenMP] component/array-ref/component (x.r[1].d) should use 'x' for GOMP_MAP_STRUCT (not yield 'x.r[1]' for nonptr 'x.r')

[burnus at gcc dot gnu.org]

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

Tobias Burnus <burnus at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
          Component|fortran                     |middle-end
            Summary|[OpenMP] Implicit mapping   |[OpenMP]
                   |breaks struct mapping       |component/array-ref/compone
                   |                            |nt (x.r[1].d) should use
                   |                            |'x' for GOMP_MAP_STRUCT
                   |                            |(not yield 'x.r[1]' for
                   |                            |nonptr 'x.r')

--- Comment #5 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #4)
>   #pragma omp target map(tofrom: x.r[1].d)
>    *.r[1].d = 3;

s/*/x/

The problem is that in gimplify.cc's gimplify_scan_omp_clauses:

                  tree base
                    = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ref,
                                               &bitpos1, &offset1,
                                               &tree_offset1);
                  bool do_map_struct = (base == decl && !tree_offset1);

Here, 'base' == 'x' but 'decl' is 'x.r[1]' - while for 'x.q.d' it is 'x' (==
base).

The comp refs are removed as follows. (It seems as if some additional ARRAY_REF
checking is needed for non-pointer components.)

              else if (TREE_CODE (decl) == COMPONENT_REF
                       && (OMP_CLAUSE_MAP_KIND (c)
                           != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
                { 
                  component_ref_p = true; 
                  while (TREE_CODE (decl) == COMPONENT_REF)
                    decl = TREE_OPERAND (decl, 0);
                  if (TREE_CODE (decl) == INDIRECT_REF
                      && DECL_P (TREE_OPERAND (decl, 0))
                      && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
                          == REFERENCE_TYPE))
                    decl = TREE_OPERAND (decl, 0);
                }

Probably the same for '!DECL_P (decl)' in the previous 'if' branch.

[Bug middle-end/104696] [OpenMP] component/array-ref/component (x.r[1].d) should use 'x' for GOMP_MAP_STRUCT (not yield 'x.r[1]' for nonptr 'x.r')

[jakub at gcc dot gnu.org]

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

Jakub Jelinek <jakub at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
   Target Milestone|12.0                        |12.2

--- Comment #6 from Jakub Jelinek <jakub at gcc dot gnu.org> ---
GCC 12.1 is being released, retargeting bugs to GCC 12.2.

[Bug middle-end/104696] [OpenMP] component/array-ref/component (x.r[1].d) should use 'x' for GOMP_MAP_STRUCT (not yield 'x.r[1]' for nonptr 'x.r')

[rguenth at gcc dot gnu.org]

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

Richard Biener <rguenth at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
   Target Milestone|12.3                        |12.4

--- Comment #8 from Richard Biener <rguenth at gcc dot gnu.org> ---
GCC 12.3 is being released, retargeting bugs to GCC 12.4.