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

["burnus at gcc dot gnu.org" <gcc-bugzilla@gcc.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.