[Bug middle-end/103416] New: [12 Regression][OpenMP] Bogus firstprivate(n) map(to:n [len: 4][implicit])

[Bug middle-end/103416] New: [12 Regression][OpenMP] Bogus firstprivate(n) map(to:n [len: 4][implicit])

[burnus at gcc dot gnu.org]

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

            Bug ID: 103416
           Summary: [12 Regression][OpenMP] Bogus firstprivate(n) map(to:n
                    [len: 4][implicit])
           Product: gcc
           Version: 12.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: cltang at gcc dot gnu.org, jakub at gcc dot gnu.org
  Target Milestone: ---

Side remark: PR80330 (suggested as possible duplicate) looks a bit similar from
the subject line.

Caused by commit r12-5194-gb7e20480630e3eeb9eed8b3941da3b3f0c22c969
Author: Chung-Lin Tang <cltang@codesourcery.com>
Date:   Fri Nov 12 20:29:00 2021 +0800

    openmp: Relax handling of implicit map vs. existing device mappings

Works when undoing the omp-low.c + gimplify.c bits.

I am not 100% sure whether there is some other issue - and this patch just
exposes it or whether it is the cause. However, the dump looks in any case
wrong/odd.

For some reasons, replacing '8' (double-precision complex) by '4'
(single-precision complex) makes it pass. The ptx assembler looks fine at a
glance, but I do observe that there are other DCmode related issues: PR102429
(nvptx, SIMT_XCHG_BFLY)

 * * *

In any case, running
https://github.com/TApplencourt/OvO/blob/master/test_src/fortran/hierarchical_parallelism/memcopy-double_complex/target_parallel_do.F90
or the program below fails with nvptx offloading with:

  libgomp: cuCtxSynchronize error: misaligned address
  libgomp: cuMemFree_v2 error: misaligned address
  libgomp: device finalization failed

PROGRAM target_parallel_do
  implicit none
  INTEGER :: i0, N
  COMPLEX(8) :: scalar
  COMPLEX(8), ALLOCATABLE :: src(:)
  N = 1
  ALLOCATE(src(1), stat=i0)
  !$OMP TARGET PARALLEL DO map(to: src) map(from: scalar) private(i0)
  DO i0 = 1, N
    scalar%re = src(i0)%re
  END DO
END PROGRAM target_parallel_do

 * * *

When playing around using not the combined directive but

  !$OMP TARGET map(to: src) map(from: scalar)
  !$OMP PARALLEL DO private(i0)
  DO i0 = 1, N
    scalar%re = src(i0)%re
  END DO

it started to work. Looking at the dump of the failing version:

          D.4246 = n;
#pragma omp target map(to:D.4246 [len: 4][implicit]) firstprivate(n) ...
  ...
  #pragma omp parallel firstprivate(D.4246) firstprivate(n) ...


while the second, working version has:

#pragma omp target firstprivate(n)
...
D.4246 = n;
....
  #pragma omp parallel firstprivate(D.4246)

 * * *

In any case, it is very odd to have both a 'map(to:' and a 'firstprivate' for
the same variable. Especially, if it causes wrong results/a crash.

[Bug middle-end/103416] [12 Regression][OpenMP] Bogus firstprivate(n) map(to:n [len: 4][implicit])

[cltang at gcc dot gnu.org]

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

--- Comment #1 from Chung-Lin Tang <cltang at gcc dot gnu.org> ---
Can you see if adding this patch:
https://gcc.gnu.org/pipermail/gcc-patches/2021-November/583279.html

fixes this problem? If so, then it should be another occurrence of PR90030

[Bug middle-end/103416] [12 Regression][OpenMP] Bogus firstprivate(n) map(to:n [len: 4][implicit])

[rguenth at gcc dot gnu.org]

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

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

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

[Bug middle-end/103416] [12 Regression][OpenMP] Bogus firstprivate(n) map(to:n [len: 4][implicit])

[burnus at gcc dot gnu.org]

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

--- Comment #2 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Chung-Lin Tang from comment #1)
> Can you see if adding this patch:
> https://gcc.gnu.org/pipermail/gcc-patches/2021-November/583279.html
> fixes this problem? If so, then it should be another occurrence of PR90030

Yes and no.

That patch FIXES the issue
  libgomp: cuCtxSynchronize error: misaligned address
  libgomp: cuMemFree_v2 error: misaligned address
  libgomp: device finalization failed


BUT I still see
  #pragma omp target map(to:D.4246 [len: 4][implicit]) ...

and when adding an explicit firstprivate(n) I see

 #pragma omp target firstprivate(n) map(to:D.4246 [len: 4][implicit]) ...

which looks wrong. I understand that implicit mapping tries to solve the
problem of explicitly mapping an array section via 'omp target (enter) data' -
and then implicitly mapping the whole array. — But I think it does not make
sense to add this implicit mapping if there is an implicit or explicit
'firstprivate' for that variable. – That's just generates pointless code,
obfuscates the dump, adds an overhead to libgomp, ...


Actually, it does not only apply to 'firstprivate' - the same also can be
caused for 'tofrom' vs. 'to' as in

#pragma omp target ... map(to:D.4217 [len: 4][implicit]) map(tofrom:n [len:
4][implicit])


for the following code:

PROGRAM target_parallel_do
  implicit none
  INTEGER :: i0, N
  COMPLEX(8) :: scalar
  N = 1
  !$OMP TARGET PARALLEL do map(from: scalar) private(i0) defaultmap(tofrom)
  DO i0 = 1, N
    scalar%re = n
  END DO
  !$omp end target parallel do
END PROGRAM target_parallel_do


Admittedly, I have not yet managed to construct something which causes an
observable misbehavior.

[Bug middle-end/103416] [12 Regression][OpenMP] Bogus firstprivate(n) map(to:n [len: 4][implicit])

[burnus at gcc dot gnu.org]

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

--- Comment #3 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Okay, the
  map(to:D.4217 [len: 4][implicit]) map(tofrom:n [len: 4][implicit])
issue is not new – only the '[implicit]' + the misaligned address one (fixed by
the patch from comment 1).

 * * *

Thus regression → comment 1

***

And extra map(to:D.4217) – this occurs already with GCC 11 and is, thus, not a
regression.
(Only '[implicit]' is new.)

This seems to be due to a scope issue. In gimplify_omp_for

(gdb) p debug(for_stmt)
#pragma omp for nowait
for (i0 = 1; i0 <= D.4217; i0 = i0 + 1)
      REALPART_EXPR <scalar> = (real(kind=8)) n;

which then calls at some point
  omp_notice_variable (
debug_tree(decl) which ends up with
7698                            nflags |= GOVD_MAP | GOVD_MAP_TO_ONLY;

The problem is that here
        integer(kind=4) D.3933;
is generated in the parent scope of '#omp target' instead of in the parent
scope of '#omp parallel do'.

Thus, there are two questions:
* Why is 'map(to:' and not 'firstprivate' used?
* Why is the var generated in the parent scope of 'omp target' instead of
inside 'omp target'?

[Bug middle-end/103416] [12 Regression][OpenMP] Bogus firstprivate(n) map(to:n [len: 4][implicit])

[burnus at gcc dot gnu.org]

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

--- Comment #4 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Created attachment 51872
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=51872&action=edit
RFC Patch to avoid the pointless evaluation, see comment 4

(In reply to Tobias Burnus from comment #3)

> * Why is the var generated in the parent scope of 'omp target' instead of
> inside 'omp target'?

The problem is a forced evaluation of the array bounds, which I regard as
pointless if the variable is just a plain variable - no array ref, not struct
ref no ...

Cf. attachment. (The question is when 'force=true' is needed and whether the
DECL_P check is the right one or whether more or less should be permitted.)

This is indeed the same as issue as PR80330 (8...)

 * * *

The
  libgomp: cuCtxSynchronize error: misaligned address
is a regression – see comment 1 for a patch which fixes it. This is PR90030
(9...)


 * * *

> * Why is 'map(to:' and not 'firstprivate' used?

Because of:

gfc_omp_predetermined_mapping (tree decl)
{
  if (DECL_ARTIFICIAL (decl)
      && ! GFC_DECL_RESULT (decl)
      && ! (DECL_LANG_SPECIFIC (decl)
            && GFC_DECL_SAVED_DESCRIPTOR (decl)))
    return OMP_CLAUSE_DEFAULTMAP_TO;

I wonder whether OMP_DEFAULTMAP_FIRSTPRIVATE  wouldn't make more sense in this
case – at least for gfc_omp_scalar_target_p ?

Which is also related to PR80330 (8...)

[Bug middle-end/103416] [12 Regression][OpenMP] Bogus firstprivate(n) map(to:n [len: 4][implicit])

[burnus at gcc dot gnu.org]

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

--- Comment #5 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #4)
> Created attachment 51872 [details]
> RFC Patch to avoid the pointless evaluation, see comment 4

The default was supposed to be 'false' - to be overridden where needed.

Otherwise, the 'false' has to added in gfc_trans_omp_do under
  /* Evaluate all the expressions in the iterator.  */

[Bug middle-end/103416] [OpenMP] Double mapping via firstprivate(n) map(to:n [len: 4][implicit])

[burnus at gcc dot gnu.org]

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

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

           What    |Removed                     |Added
----------------------------------------------------------------------------
            Summary|[12 Regression][OpenMP]     |[OpenMP] Double mapping via
                   |Bogus firstprivate(n)       |firstprivate(n) map(to:n
                   |map(to:n [len:              |[len: 4][implicit])
                   |4][implicit])               |
           Keywords|wrong-code                  |missed-optimization

--- Comment #6 from Tobias Burnus <burnus at gcc dot gnu.org> ---
I changed the bug subject as the "12 Regression" issue has been fixed (cf.
below,
committed with reference to PR90030 not to this PR) – but the code-improvement
part still remains to be done.

 * * *

(In reply to Chung-Lin Tang from comment #1)
> Can you see if adding this patch:
> https://gcc.gnu.org/pipermail/gcc-patches/2021-November/583279.html
> fixes this problem? If so, then it should be another occurrence of PR90030

This patch is now committed:

commit r12-5706-g1ac7a8c9e4798d352eb8c64905dd38086af4e1cd
Author: Chung-Lin Tang <cltang@codesourcery.com>
Date:   Thu Dec 2 18:24:03 2021 +0800

    fortran: OpenMP/OpenACC array mapping alignment fix (PR90030)

* * *

TODO:  LOOP VAR double mapping – cf. PR80330, i.e.
  map(to:D.4217 [len: 4][implicit]) map(tofrom:n [len: 4][implicit])
where the loop variable is mapped twice.

See Comment 4 for what remains to be done and a draft patch.

[Bug fortran/103416] [OpenMP] Double mapping via firstprivate(n) map(to:n [len: 4][implicit]) of loop bounds

[jakub at gcc dot gnu.org]

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

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

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

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

[Bug fortran/103416] [OpenMP] Double mapping via firstprivate(n) map(to:n [len: 4][implicit]) of loop bounds

[rguenth at gcc dot gnu.org]

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

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

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

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