[Bug middle-end/113867] New: [14 Regression][OpenMP] Wrong code with mapping pointers in structs

[Bug middle-end/113867] New: [14 Regression][OpenMP] Wrong code with mapping pointers in structs

[burnus at gcc dot gnu.org]

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

            Bug ID: 113867
           Summary: [14 Regression][OpenMP] Wrong code with mapping
                    pointers in structs
           Product: gcc
           Version: 14.0
            Status: UNCONFIRMED
          Keywords: openmp, wrong-code
          Severity: normal
          Priority: P3
         Component: middle-end
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
                CC: jakub at gcc dot gnu.org
  Target Milestone: ---

Created attachment 57381
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=57381&action=edit
Testcase, compile with 'gcc -fopenmp' and run with an offload device

Split off from PR113724 which mainly about an ICE.

The attached programs work with GCC 13 but fails with mainline.
(Requires actual offloading; tried here with nvptx.)

Probably due to Julian's mapping patches.


With mainline, it fails for 'g()' when executing

   omp target data map(tofrom: s2.p[:100])

(i.e. GOMP_target_data_ext → gomp_copy_host2dev → gomp_device_copy) with

  libgomp: cuMemGetAddressRange_v2 error: named symbol not found
  libgomp: Copying of host object [0x118c500..0x118c690) to dev object
[0x7f7e721cae00..0x7f7e721caf90) failed

It works using a separate target enter/exit data (i.e. for 'f()').


The mainline dump shows:

map(struct:s2 [len: 1]) map(alloc:s2.p [len: 0]) map(tofrom:*_2 [len: 400])
map(attach:s2.p [bias: 0])

I somehow hadn't expected the

   map(struct:s2 [len: 1]) map(alloc:s2.p [len: 0])

which might or might not be the issue. As it works with 'f()' (i.e. enter/exit
data), it might be a red herring (or not).

[Bug middle-end/113867] [14 Regression][OpenMP] Wrong code with mapping pointers in structs

[burnus at gcc dot gnu.org]

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

--- Comment #1 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Created attachment 57382
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=57382&action=edit
Fortran testcase, kind of, as pointer + pointee mapping cannot be split
(working)

For completeness, a Fortran testcase.
(This Testcase works on GCC 13 and mainline.)

As in Fortran, 'map(ptr, dt%ptr)' will always attempt to map the pointer and
the pointer, it is not possible to split
   map(s2.p[:N])  // map the pointee
   map(s2.p, s2.p[:0]) // map the pointer and try pointer attachment
as in C/C++. And using 'map(s.p)' will prevent a later inner 'map(s.a,...)' as
's' is already partially mapped. Hence, the aux 'ptr' is used, but that kind of
defeats the purpose of this testcase.

[Bug middle-end/113867] [14 Regression][OpenMP] Wrong code with mapping pointers in structs

[rguenth at gcc dot gnu.org]

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

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

           What    |Removed                     |Added
----------------------------------------------------------------------------
   Target Milestone|---                         |14.0

[Bug libgomp/113867] [14 Regression][OpenMP] Wrong code with mapping pointers in structs

[burnus at gcc dot gnu.org]

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

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

           What    |Removed                     |Added
----------------------------------------------------------------------------
           Assignee|unassigned at gcc dot gnu.org      |burnus at gcc dot gnu.org
          Component|middle-end                  |libgomp

--- Comment #2 from Tobias Burnus <burnus at gcc dot gnu.org> ---
The problem here is in libgomp's gomp_map_vars_internal:

                /* Fallthrough.  */
              case GOMP_MAP_STRUCT:
                first = i + 1;
                last = i + sizes[i];
                cur_node.host_start = (uintptr_t) hostaddrs[i];
                cur_node.host_end = (uintptr_t) hostaddrs[last]
                                    + sizes[last];
                if (tgt->list[first].key != NULL)
                  continue;
                if (sizes[last] == 0)
                  cur_node.host_end++;
                n = splay_tree_lookup (mem_map, &cur_node);
                if (sizes[last] == 0)
                  cur_node.host_end--;
                if (n == NULL && cur_node.host_start == cur_node.host_end)
                  {
                    gomp_mutex_unlock (&devicep->lock);
                    gomp_fatal ("Struct pointer member not mapped (%p)",
                                (void*) hostaddrs[first]);
                  }
                if (n == NULL)
...
                    field_tgt_base = (uintptr_t) hostaddrs[first];
...
                    field_tgt_clear = last;

here: n == NULL and cur_node.host_end - cur_node.host_start = 8 [i.e.
sizeof(void*)?!]:

For i=1, there is no action to be taken due to the
GOMP_MAP_ZERO_LEN_ARRAY_SECTION.

And for i=2,

                if (field_tgt_clear != FIELD_TGT_EMPTY)
                  {
                    k->tgt_offset = k->host_start - field_tgt_base
                                    + field_tgt_offset;

Here, k->tgt_offset = hostaddr of the struct but we are no longer mapping a
struct here. - Clearly, resetting was forgotten ...

[Bug libgomp/113867] [14 Regression][OpenMP] Wrong code with mapping pointers in structs

[burnus at gcc dot gnu.org]

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

--- Comment #3 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Created attachment 57398
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=57398&action=edit
Patch - handling the libgomp issue

Possible patch - lightly tested. This fixes the issue in libgomp.

While always correct and possibly avoiding other corner cases (if there are
any), an alternative approach is to not create those 'struct: s [len: 1]' +
'alloc:s2.p [len: 0]'.

RFC:
- Is there a reason why we want to have the struct in such a case?
  (GCC <= 13 doesn't create this struct)
- Do we want to have this libgomp change even when not generating the struct
  as extra safety net?

[Bug libgomp/113867] [14 Regression][OpenMP] Wrong code with mapping pointers in structs

[burnus at gcc dot gnu.org]

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

--- Comment #4 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Also not handled:
  struct s { int *p; } s1;
  ...
  #pragma omp target map(s1.p[:N])
    p[0] = p[N-1] = 99;

Here, the pointer attachment is missing.  See also PR113724 's attachment 57407
for a testcase for this and (some) other issues.

TODO:
- Fix the extra struct issue (→ this patch or other solution)
- Fix the missing attachment issue (this comment's example)
- Audit whether other changes are required.

[Bug libgomp/113867] [14 Regression][OpenMP] Wrong code with mapping pointers in structs

[burnus at gcc dot gnu.org]

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

--- Comment #5 from Tobias Burnus <burnus at gcc dot gnu.org> ---
For:

  int *q;
  #pragma omp target device(y5) map(q, q[:5])

GCC currently generates:
  map(tofrom:q [len: 8]) map(tofrom:*q.4_1 [len: 20]) map(attach:q [bias: 0])

Expected:
  'alloc:' instead of 'attach:'
or even:
  map(tofrom:*q [len: 20]) map(firstprivate:q [pointer assign, bias: 0])

In any case, the first 'tofrom' is pointless!


NOTE: GCC 13 shows:
  error: 'q' appears both in data and map clauses

 * * *

For
  #pragma omp target map(s.p[:5])

GCC should do:
  map(tofrom:s [len: 24][implicit]) map(tofrom:*_5 [len: 16])
    map(attach:s.p [bias: 0])

But (regression!) it does:
  map(struct:s [len: 1]) map(alloc:s.p [len: 8]) map(tofrom:*_5 [len: 16])
    map(attach:s.p [bias: 0])

Solution:

--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -12381,3 +12381,4 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
              if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
-                 || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+                 || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH
+                 || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
                break;


However, unless I messed up, this will cause tons of ICE(segfault).

[Bug libgomp/113867] [14 Regression][OpenMP] Wrong code with mapping pointers in structs

[law at gcc dot gnu.org]

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

Jeffrey A. Law <law at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
           Priority|P3                          |P1
                 CC|                            |law at gcc dot gnu.org

[Bug libgomp/113867] [14 Regression][OpenMP] Wrong code with mapping pointers in structs

[rguenth at gcc dot gnu.org]

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

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

           What    |Removed                     |Added
----------------------------------------------------------------------------
             Status|UNCONFIRMED                 |ASSIGNED
     Ever confirmed|0                           |1
   Last reconfirmed|                            |2024-04-03

[Bug libgomp/113867] [14 Regression][OpenMP] Wrong code with mapping pointers in structs

[jakub at gcc dot gnu.org]

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

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

           What    |Removed                     |Added
----------------------------------------------------------------------------
           Priority|P1                          |P2

--- Comment #6 from Jakub Jelinek <jakub at gcc dot gnu.org> ---
Per IRC discussions downgrading this to P2, it isn't something used widely in
real-world code and while we can get some partial fix soon, a full fix will
take likely longer.  If it is fixed relatively soon in stage1, we can consider
backporting the non-risky parts of the fixes to 14.2.

[Bug libgomp/113867] [14/15 Regression][OpenMP] Wrong code with mapping pointers in structs

[rguenth at gcc dot gnu.org]

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

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

           What    |Removed                     |Added
----------------------------------------------------------------------------
   Target Milestone|14.0                        |14.2

--- Comment #7 from Richard Biener <rguenth at gcc dot gnu.org> ---
GCC 14.1 is being released, retargeting bugs to GCC 14.2.