[PATCH][OpenACC] Update deviceptr handling during gimplification

[PATCH][OpenACC] Update deviceptr handling during gimplification

[Cesar Philippidis]

[-- Attachment #1: Type: text/plain, Size: 431 bytes --]

I had previously posted this patch as part of a monster deviceptr patch
here <https://gcc.gnu.org/ml/gcc-patches/2018-06/msg01911.html>. This
patch breaks out the generic gimplifier changes. Essentially, with this
patch, the gimplifier will now transfer deviceptr data clauses using
GOMP_MAP_FORCE_DEVICEPTR.

Is this patch OK for trunk? It bootstrapped / regression tested cleanly
for x86_64 with nvptx offloading.

Thanks,
Cesar

[-- Attachment #2: 0001-OpenACC-Update-deviceptr-handling.patch --]
[-- Type: text/x-patch, Size: 4906 bytes --]

From b5cf37b795ce78c78f3f434ac6999f7094bd86aa Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <cesar@codesourcery.com>
Date: Mon, 7 May 2018 08:23:48 -0700
Subject: [PATCH] [OpenACC] Update deviceptr handling

2018-XX-YY  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/fortran/
	* trans-openmp.c (gfc_omp_finish_clause): Don't create pointer data
	mappings for deviceptr clauses.
	(gfc_trans_omp_clauses): Likewise.
	gcc/
	* gimplify.c (enum gimplify_omp_var_data): Add GOVD_DEVICETPR.
	(omp_notice_variable): Add GOVD_DEVICEPTR attribute when appropriate.
	(gimplify_scan_omp_clauses): Likewise.
	(gimplify_adjust_omp_clauses_1): Set GOMP_MAP_FORCE_DEVICEPTR for
	implicit deviceptr mappings.
	gcc/testsuite/
	* c-c++-common/goacc/deviceptr-4.c: Update expected data mapping.

(cherry picked from openacc-gcc-7-branch commit
d3de16b461545aac1925f0d7c2851c8c49a07d06 and commit
f0514fe1899666bb5b8ee52601f5d4263d4c4646)
---
 gcc/fortran/trans-openmp.c                     |  9 +++++++++
 gcc/gimplify.c                                 | 12 +++++++++++-
 gcc/testsuite/c-c++-common/goacc/deviceptr-4.c |  2 +-
 3 files changed, 21 insertions(+), 2 deletions(-)

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index f038f4c..ca31c88 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1060,6 +1060,8 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p)
     }
 
   tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE;
+  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
+    return;
   if (POINTER_TYPE_P (TREE_TYPE (decl)))
     {
       if (!gfc_omp_privatize_by_reference (decl)
@@ -2111,6 +2113,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	      if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
 		{
 		  if (POINTER_TYPE_P (TREE_TYPE (decl))
+		      && n->u.map_op == OMP_MAP_FORCE_DEVICEPTR)
+		    {
+		      OMP_CLAUSE_DECL (node) = decl;
+		      goto finalize_map_clause;
+		    }
+		  else if (POINTER_TYPE_P (TREE_TYPE (decl))
 		      && (gfc_omp_privatize_by_reference (decl)
 			  || GFC_DECL_GET_SCALAR_POINTER (decl)
 			  || GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
@@ -2282,6 +2290,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		  ptr2 = fold_convert (sizetype, ptr2);
 		  OMP_CLAUSE_SIZE (node3)
 		    = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
+		finalize_map_clause:;
 		}
 	      switch (n->u.map_op)
 		{
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 4a109ae..bcf862f 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -105,6 +105,9 @@ enum gimplify_omp_var_data
   /* Flag for GOVD_MAP: must be present already.  */
   GOVD_MAP_FORCE_PRESENT = 524288,
 
+  /* Flag for OpenACC deviceptrs.  */
+  GOVD_DEVICEPTR = (1<<21),
+
   GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
 			   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
 			   | GOVD_LOCAL)
@@ -7232,6 +7235,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
 		        error ("variable %qE declared in enclosing "
 			       "%<host_data%> region", DECL_NAME (decl));
 		      nflags |= GOVD_MAP;
+		      nflags |= (n2->value & GOVD_DEVICEPTR);
 		      if (octx->region_type == ORT_ACC_DATA
 			  && (n2->value & GOVD_MAP_0LEN_ARRAY))
 			nflags |= GOVD_MAP_0LEN_ARRAY;
@@ -8213,6 +8217,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
 	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
 	    flags |= GOVD_MAP_ALWAYS_TO;
+	  else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
+	    flags |= GOVD_DEVICEPTR;
 	  goto do_add;
 
 	case OMP_CLAUSE_DEPEND:
@@ -8828,7 +8834,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
       /* Not all combinations of these GOVD_MAP flags are actually valid.  */
       switch (flags & (GOVD_MAP_TO_ONLY
 		       | GOVD_MAP_FORCE
-		       | GOVD_MAP_FORCE_PRESENT))
+		       | GOVD_MAP_FORCE_PRESENT
+		       | GOVD_DEVICEPTR))
 	{
 	case 0:
 	  kind = GOMP_MAP_TOFROM;
@@ -8845,6 +8852,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
 	case GOVD_MAP_FORCE_PRESENT:
 	  kind = GOMP_MAP_FORCE_PRESENT;
 	  break;
+	case GOVD_DEVICEPTR:
+	  kind = GOMP_MAP_FORCE_DEVICEPTR;
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
index db1b916..79a5162 100644
--- a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
@@ -8,4 +8,4 @@ subr (int *a)
   a[0] += 1.0;
 }
 
-/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(tofrom:a" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(force_deviceptr:a" 1 "gimple" } } */
-- 
2.7.4

Re: [PATCH][OpenACC] Update deviceptr handling during gimplification

[Julian Brown]

[-- Attachment #1: Type: text/plain, Size: 956 bytes --]

On Tue, 7 Aug 2018 15:09:38 -0700
Cesar Philippidis <cesar_philippidis@mentor.com> wrote:

> I had previously posted this patch as part of a monster deviceptr
> patch here
> <https://gcc.gnu.org/ml/gcc-patches/2018-06/msg01911.html>. This
> patch breaks out the generic gimplifier changes. Essentially, with
> this patch, the gimplifier will now transfer deviceptr data clauses
> using GOMP_MAP_FORCE_DEVICEPTR.
> 
> Is this patch OK for trunk? It bootstrapped / regression tested
> cleanly for x86_64 with nvptx offloading.

This patch also appears to fix the attached test case, which had been
associated with a different deviceptr-related patch on the og8 branch
(the other parts of which are upstream already). Perhaps you'd like to
incorporate this test into your patch? It was by James Norris
originally, IIUC.

Thanks,

Julian

ChangeLog

2018-xx-xx  James Norris  <???>

    libgomp/
    * testsuite/libgomp.oacc-fortran/deviceptr-1.f90: New test.

[-- Attachment #2: deviceptr-1.f90.diff --]
[-- Type: text/x-patch, Size: 3563 bytes --]

--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90
@@ -0,0 +1,197 @@
+! { dg-do run }
+
+! Test the deviceptr clause with various directives
+! and in combination with other directives where
+! the deviceptr variable is implied.
+
+subroutine subr1 (a, b)
+  implicit none
+  integer, parameter :: N = 8
+  integer :: a(N)
+  integer :: b(N)
+  integer :: i = 0
+
+  !$acc data deviceptr (a)
+
+  !$acc parallel copy (b)
+    do i = 1, N
+      a(i) = i * 2
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc end data
+
+end subroutine
+
+subroutine subr2 (a, b)
+  implicit none
+  integer, parameter :: N = 8
+  integer :: a(N)
+  !$acc declare deviceptr (a)
+  integer :: b(N)
+  integer :: i = 0
+
+  !$acc parallel copy (b)
+    do i = 1, N
+      a(i) = i * 4
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+end subroutine
+
+subroutine subr3 (a, b)
+  implicit none
+  integer, parameter :: N = 8
+  integer :: a(N)
+  !$acc declare deviceptr (a)
+  integer :: b(N)
+  integer :: i = 0
+
+  !$acc kernels copy (b)
+    do i = 1, N
+      a(i) = i * 8
+      b(i) = a(i)
+    end do
+  !$acc end kernels
+
+end subroutine
+
+subroutine subr4 (a, b)
+  implicit none
+  integer, parameter :: N = 8
+  integer :: a(N)
+  integer :: b(N)
+  integer :: i = 0
+
+  !$acc parallel deviceptr (a) copy (b)
+    do i = 1, N
+      a(i) = i * 16
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+end subroutine
+
+subroutine subr5 (a, b)
+  implicit none
+  integer, parameter :: N = 8
+  integer :: a(N)
+  integer :: b(N)
+  integer :: i = 0
+
+  !$acc kernels deviceptr (a) copy (b)
+    do i = 1, N
+      a(i) = i * 32
+      b(i) = a(i)
+    end do
+  !$acc end kernels
+
+end subroutine
+
+subroutine subr6 (a, b)
+  implicit none
+  integer, parameter :: N = 8
+  integer :: a(N)
+  integer :: b(N)
+  integer :: i = 0
+
+  !$acc parallel deviceptr (a) copy (b)
+    do i = 1, N
+      b(i) = i
+    end do
+  !$acc end parallel
+
+end subroutine
+
+subroutine subr7 (a, b)
+  implicit none
+  integer, parameter :: N = 8
+  integer :: a(N)
+  integer :: b(N)
+  integer :: i = 0
+
+  !$acc data deviceptr (a)
+
+  !$acc parallel copy (b)
+    do i = 1, N
+      a(i) = i * 2
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc parallel copy (b)
+    do i = 1, N
+      a(i) = b(i) * 2
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc end data
+
+end subroutine
+
+program main
+  use iso_c_binding, only: c_ptr, c_f_pointer
+  implicit none
+  type (c_ptr) :: cp
+  integer, parameter :: N = 8
+  integer, pointer :: fp(:)
+  integer :: i = 0
+  integer :: b(N)
+
+  interface
+    function acc_malloc (s) bind (C)
+      use iso_c_binding, only: c_ptr, c_size_t
+      integer (c_size_t), value :: s
+      type (c_ptr) :: acc_malloc
+    end function
+  end interface
+
+  cp = acc_malloc (N * sizeof (fp(N)))
+  call c_f_pointer (cp, fp, [N])
+
+  call subr1 (fp, b)
+
+  do i = 1, N
+    if (b(i) .ne. i * 2) call abort
+  end do
+
+  call subr2 (fp, b)
+
+  do i = 1, N
+    if (b(i) .ne. i * 4) call abort
+  end do
+
+  call subr3 (fp, b)
+
+  do i = 1, N
+    if (b(i) .ne. i * 8) call abort
+  end do
+
+  call subr4 (fp, b)
+
+  do i = 1, N
+    if (b(i) .ne. i * 16) call abort
+  end do
+
+  call subr5 (fp, b)
+
+  do i = 1, N
+    if (b(i) .ne. i * 32) call abort
+  end do
+
+  call subr6 (fp, b)
+
+  do i = 1, N
+    if (b(i) .ne. i) call abort
+  end do
+
+  call subr7 (fp, b)
+
+  do i = 1, N
+    if (b(i) .ne. i * 4) call abort
+  end do
+
+end program main

Re: [PATCH][OpenACC] Update deviceptr handling during gimplification

[Cesar Philippidis]

On 09/25/2018 05:55 PM, Julian Brown wrote:
> On Tue, 7 Aug 2018 15:09:38 -0700
> Cesar Philippidis <cesar_philippidis@mentor.com> wrote:
> 
>> I had previously posted this patch as part of a monster deviceptr
>> patch here
>> <https://gcc.gnu.org/ml/gcc-patches/2018-06/msg01911.html>. This
>> patch breaks out the generic gimplifier changes. Essentially, with
>> this patch, the gimplifier will now transfer deviceptr data clauses
>> using GOMP_MAP_FORCE_DEVICEPTR.
>>
>> Is this patch OK for trunk? It bootstrapped / regression tested
>> cleanly for x86_64 with nvptx offloading.
> 
> This patch also appears to fix the attached test case, which had been
> associated with a different deviceptr-related patch on the og8 branch
> (the other parts of which are upstream already). Perhaps you'd like to
> incorporate this test into your patch? It was by James Norris
> originally, IIUC.

Ok, I'll do that. Thanks for updating those tests.

Cesar

Re: [PATCH][OpenACC] Update deviceptr handling during gimplification

[Jakub Jelinek]

On Tue, Aug 07, 2018 at 03:09:38PM -0700, Cesar Philippidis wrote:
> I had previously posted this patch as part of a monster deviceptr patch
> here <https://gcc.gnu.org/ml/gcc-patches/2018-06/msg01911.html>. This
> patch breaks out the generic gimplifier changes. Essentially, with this
> patch, the gimplifier will now transfer deviceptr data clauses using
> GOMP_MAP_FORCE_DEVICEPTR.
> 
> Is this patch OK for trunk? It bootstrapped / regression tested cleanly
> for x86_64 with nvptx offloading.

See my comments on the monster deviceptr patch.
> 
> 2018-XX-YY  Cesar Philippidis  <cesar@codesourcery.com>
> 
> 	gcc/fortran/
> 	* trans-openmp.c (gfc_omp_finish_clause): Don't create pointer data
> 	mappings for deviceptr clauses.
> 	(gfc_trans_omp_clauses): Likewise.
> 	gcc/
> 	* gimplify.c (enum gimplify_omp_var_data): Add GOVD_DEVICETPR.
> 	(omp_notice_variable): Add GOVD_DEVICEPTR attribute when appropriate.
> 	(gimplify_scan_omp_clauses): Likewise.
> 	(gimplify_adjust_omp_clauses_1): Set GOMP_MAP_FORCE_DEVICEPTR for
> 	implicit deviceptr mappings.
> 	gcc/testsuite/
> 	* c-c++-common/goacc/deviceptr-4.c: Update expected data mapping.

	Jakub