[Bug fortran/108558] New: OpenMP/Fortran 'has_device_addr' clause getting lost?

[Bug fortran/108558] New: OpenMP/Fortran 'has_device_addr' clause getting lost?

[tschwinge at gcc dot gnu.org]

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

            Bug ID: 108558
           Summary: OpenMP/Fortran 'has_device_addr' clause getting lost?
           Product: gcc
           Version: 13.0
            Status: UNCONFIRMED
          Keywords: openmp
          Severity: normal
          Priority: P3
         Component: fortran
          Assignee: unassigned at gcc dot gnu.org
          Reporter: tschwinge at gcc dot gnu.org
                CC: burnus at gcc dot gnu.org, jakub at gcc dot gnu.org
  Target Milestone: ---

It's certainly possible that I'm doing something wrong here (first use of
OpenMP 'has_device_addr' clause), but please consider:

    subroutine vectorAdd(a, b, M)
      implicit none
      integer(4)::a(M), b(M)
      integer:: i, M
      !$omp target teams distribute parallel do has_device_addr(a, b)
      do i = 1, M
        b(i) = a(i) + b(i)
      end do
    end subroutine

(..., which is called from a '!$omp target data use_device_addr(a, b)' inside a
'!$omp target data map(tofrom:a(1:M), b(1:M))'.)

If I 'diff' the '-fopenmp -fdump-tree-all' without vs. with the
'has_device_addr(a, b)' clause, I -- unexpectedly -- get no differences (aside
from minor ones due to what seems to be different order of compiler
temporaries):

'pr.f90.005t.original':

    #pragma omp target
    [...]

(Decomposed combined construct.  Is that perhaps where the problem lies?)

'pr.f90.006t.gimple':

    #pragma omp target num_teams(0) thread_limit(0) firstprivate(m)
map(tofrom:*b [len: D.4283][implicit]) map(alloc:b [pointer assign, bias: 0])
map(tofrom:*a [len: D.4280][implicit]) map(alloc:a [pointer assign, bias: 0])
    [...]

That is, 'map' instead of 'has_device_addr'.


In contrast, for the translated C code:

    void vectorAdd(int *a, int *b, int M)
    {
      #pragma omp target teams distribute parallel for has_device_addr(a, b)
      for (int i = 1; i < M; ++i)
        b[i] = a[i] + b[i];
    }

..., I see the expected 'diff' of 'pr.c.005t.original':

    -  #pragma omp target
    +  #pragma omp target has_device_addr(a) has_device_addr(b)

..., and 'diff' of 'pr.c.006t.gimple':

    -  #pragma omp target num_teams(0) thread_limit(0) firstprivate(M)
map(alloc:MEM[(char *)b] [len: 0]) map(firstprivate:b [pointer assign, bias:
0]) map(alloc:MEM[(char *)a] [len: 0]) map(firstprivate:a [pointer assign,
bias: 0])
    +  #pragma omp target num_teams(0) thread_limit(0) has_device_addr(a)
has_device_addr(b) firstprivate(M)

(Have not examined that one any further.)


Cross-checking with corresponding OpenACC/Fortran 'deviceptr' clause ('!$acc
parallel loop deviceptr(a, b)'), that seems to work as expected (from a quick
look, not futher examined):

'pr.f90.005t.original':

    -    #pragma acc parallel
    +    #pragma acc parallel map(force_deviceptr:*a) map(alloc:a [pointer
assign, bias: 0]) map(force_deviceptr:*b) map(alloc:b [pointer assign, bias:
0])

'pr.f90.006t.gimple':

    -    #pragma omp target oacc_parallel firstprivate(D.4291) map(tofrom:*b
[len: D.4298]) map(alloc:b [pointer assign, bias: 0]) map(tofrom:*a [len:
D.4295]) map(alloc:a [pointer assign, bias: 0])
    +    a.8_9 = a;
    +    b.9_10 = b;
    +    #pragma omp target oacc_parallel map(force_deviceptr:(*a.8_9) [len:
D.4298]) map(alloc:a [pointer assign, bias: 0]) map(force_deviceptr:(*b.9_10)
[len: D.4295]) map(alloc:b [pointer assign, bias: 0]) firstprivate(D.4291)


That's with GCC based on fairly recent commit
de99049f6fe5341024d4d939ac50d063280f90db (2023-01-11).

[Bug fortran/108558] OpenMP/Fortran 'has_device_addr' clause getting lost?

[burnus at gcc dot gnu.org]

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

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

           What    |Removed                     |Added
----------------------------------------------------------------------------
   Last reconfirmed|                            |2023-01-26
     Ever confirmed|0                           |1
             Status|UNCONFIRMED                 |NEW

--- Comment #1 from Tobias Burnus <burnus at gcc dot gnu.org> ---
I bet that this is a problem of 'gfc_split_omp_clauses':

--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -6352,4 +6352,6 @@ gfc_split_omp_clauses (gfc_code *code,
          clausesa[GFC_OMP_SPLIT_TARGET].lists[OMP_LIST_IS_DEVICE_PTR]
            = code->ext.omp_clauses->lists[OMP_LIST_IS_DEVICE_PTR];
+         clausesa[GFC_OMP_SPLIT_TARGET].lists[OMP_LIST_HAS_DEVICE_ADDR]
+           = code->ext.omp_clauses->lists[OMP_LIST_HAS_DEVICE_ADDR];
          clausesa[GFC_OMP_SPLIT_TARGET].device
            = code->ext.omp_clauses->device;

[Bug fortran/108558] OpenMP/Fortran 'has_device_addr' clause getting lost?

[tschwinge at gcc dot gnu.org]

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

--- Comment #2 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #1)
> I bet that this is a problem of 'gfc_split_omp_clauses': [...]

Heh, so indeed as I suspected:

(In reply to myself from comment #0)
> (Decomposed combined construct.  Is that perhaps where the problem lies?)

:-)

With your patch (thanks!) applied, I do get what I suspect are the expected
changes:

'pr.f90.005t.original':

    -  #pragma omp target
    +  #pragma omp target has_device_addr(a) has_device_addr(b)

'pr.f90.006t.gimple':

    -  #pragma omp target num_teams(0) thread_limit(0) firstprivate(m)
map(tofrom:*b [len: D.4283][implicit]) map(alloc:b [pointer assign, bias: 0])
map(tofrom:*a [len: D.4280][implicit]) map(alloc:a [pointer assign, bias: 0])
    +  #pragma omp target num_teams(0) thread_limit(0) has_device_addr(a)
has_device_addr(b) firstprivate(D.4283) firstprivate(D.4280) firstprivate(m)

..., and my original test case behaves as expected; OpenMP/Fortran
'has_device_addr' works.

[Bug fortran/108558] OpenMP/Fortran 'has_device_addr' clause getting lost?

[cvs-commit at gcc dot gnu.org]

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

--- Comment #3 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Tobias Burnus <burnus@gcc.gnu.org>:

https://gcc.gnu.org/g:2325c8920bbc99edcc9fffaa79577c528df41eb8

commit r13-5438-g2325c8920bbc99edcc9fffaa79577c528df41eb8
Author: Tobias Burnus <tobias@codesourcery.com>
Date:   Fri Jan 27 11:32:19 2023 +0100

    OpenMP/Fortran: Fix has_device_addr clause splitting [PR108558]

    gcc/fortran/ChangeLog:

            PR fortran/108558
            * trans-openmp.cc (gfc_split_omp_clauses): Handle has_device_addr.

    libgomp/ChangeLog:

            PR fortran/108558
            * testsuite/libgomp.fortran/has_device_addr.f90: New test.

[Bug fortran/108558] OpenMP/Fortran 'has_device_addr' clause getting lost?

[burnus at gcc dot gnu.org]

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

--- Comment #4 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Fixed on mainline (GCC 13); still need to backport to GCC 12 as the original
feature was added in r12-7161-gbbb7f8604e1dfc08f44354cfd93d2287f2fdd489

[Bug fortran/108558] OpenMP/Fortran 'has_device_addr' clause getting lost?

[cvs-commit at gcc dot gnu.org]

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

--- Comment #5 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The releases/gcc-12 branch has been updated by Tobias Burnus
<burnus@gcc.gnu.org>:

https://gcc.gnu.org/g:591ec4820aa4e6d757ddc76cae1d92d445daf72c

commit r12-9090-g591ec4820aa4e6d757ddc76cae1d92d445daf72c
Author: Tobias Burnus <tobias@codesourcery.com>
Date:   Fri Jan 27 11:32:19 2023 +0100

    OpenMP/Fortran: Fix has_device_addr clause splitting [PR108558]

    gcc/fortran/ChangeLog:

            PR fortran/108558
            * trans-openmp.cc (gfc_split_omp_clauses): Handle has_device_addr.

    libgomp/ChangeLog:

            PR fortran/108558
            * testsuite/libgomp.fortran/has_device_addr.f90: New test.

    (cherry picked from commit 2325c8920bbc99edcc9fffaa79577c528df41eb8)

[Bug fortran/108558] OpenMP/Fortran 'has_device_addr' clause getting lost?

[burnus at gcc dot gnu.org]

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

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

           What    |Removed                     |Added
----------------------------------------------------------------------------
             Status|NEW                         |RESOLVED
         Resolution|---                         |FIXED

--- Comment #6 from Tobias Burnus <burnus at gcc dot gnu.org> ---
FIXED for GCC 12 + GCC 13 = all affected. → Close.

Thanks for the report