[gomp4] Fix handling of subarrays with update directive

[gomp4] Fix handling of subarrays with update directive

[James Norris]

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

Hi,

The attached patch fixes a defect reported with gcc 5.2
(https://gcc.gnu.org/bugzilla/show_bug.cgi?id=69414).
It is also the case, the issue is present in the gomp4
branch. The patch also adds additional testing.

Committed to gomp4 after bootstrap and regtesting.

Thanks,
Jim


[-- Attachment #2: update.patch --]
[-- Type: text/x-patch, Size: 15265 bytes --]

Index: ChangeLog.gomp
===================================================================
--- ChangeLog.gomp	(revision 232740)
+++ ChangeLog.gomp	(revision 232741)
@@ -1,3 +1,11 @@
+2016-01-22  James Norris  <jnorris@codesourcery.com>
+
+	* oacc-mem.c (delete_copyout, update_dev_host): Fix device address.
+	* testsuite/libgomp.oacc-c-c++-common/update-1.c: Additional tests.
+	* testsuite/libgomp.oacc-c-c++-common/update-1-2.c: Likewise.
+	* testsuite/libgomp.oacc-fortran/update-1.f90: New file.
+	* testsuite/libgomp.oacc-fortran/update-1-2.f90: Likewise.
+
 2016-01-22  Nathan Sidwell  <nathan@codesourcery.com>
 
 	* testsuite/libgomp.oacc-c-c++-common/routine-1.c: Specify vector.
Index: oacc-mem.c
===================================================================
--- oacc-mem.c	(revision 232740)
+++ oacc-mem.c	(revision 232741)
@@ -509,7 +509,7 @@
       gomp_fatal ("[%p,%d] is not mapped", (void *)h, (int)s);
     }
 
-  d = (void *) (n->tgt->tgt_start + n->tgt_offset);
+  d = (void *) (n->tgt->tgt_start + n->tgt_offset + h - n->host_start);
 
   host_size = n->host_end - n->host_start;
 
@@ -562,7 +562,7 @@
       gomp_fatal ("[%p,%d] is not mapped", h, (int)s);
     }
 
-  d = (void *) (n->tgt->tgt_start + n->tgt_offset);
+  d = (void *) (n->tgt->tgt_start + n->tgt_offset + h - n->host_start);
 
   gomp_mutex_unlock (&acc_dev->lock);
 
Index: testsuite/libgomp.oacc-fortran/update-1-2.f90
===================================================================
--- testsuite/libgomp.oacc-fortran/update-1-2.f90	(revision 0)
+++ testsuite/libgomp.oacc-fortran/update-1-2.f90	(revision 232741)
@@ -0,0 +1,239 @@
+! Copy of update-1.f90 with self exchanged with host for !$acc update
+
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+program update
+  use openacc
+  implicit none 
+  integer, parameter :: N = 8
+  real :: a(N), b(N)
+  integer i
+
+  do i = 1, N
+    a(i) = 3.0
+    b(i) = 0.0
+  end do
+
+  !$acc enter data copyin (a, b)
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update self (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 3.0) call abort
+    if (b(i) .ne. 3.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 5.0
+    b(i) = 1.0
+  end do
+
+  !$acc update device (a, b)
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do 
+  !$acc end parallel
+
+  !$acc update self (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 5.0) call abort
+    if (b(i) .ne. 5.0) call abort
+ end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  !$acc parallel present (a, b)
+  do i = 1, N
+    b(i) = a(i)
+  end do
+  !$acc end parallel
+
+  !$acc update self (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 5.0) call abort
+    if (b(i) .ne. 5.0) call abort
+  end do
+
+  do i = 1, N
+    a(i) = 6.0
+    b(i) = 0.0
+  end do
+
+  !$acc update device (a, b)
+
+  do i = 1, N
+    a(i) = 9.0
+  end do
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update self (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 6.0) call abort
+    if (b(i) .ne. 6.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 7.0
+    b(i) = 2.0
+  end do
+
+  !$acc update device (a, b)
+
+  do i = 1, N
+    a(i) = 9.0
+  end do
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update self (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 7.0) call abort
+    if (b(i) .ne. 7.0) call abort
+  end do
+
+  do i = 1, N
+    a(i) = 9.0
+  end do
+
+  !$acc update device (a)
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update self (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 9.0) call abort
+    if (b(i) .ne. 9.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 5.0
+  end do
+
+  !$acc update device (a)
+
+  do i = 1, N
+    a(i) = 6.0
+  end do
+
+  !$acc update device (a(1:rshift (N, 1)))
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update self (a, b)
+
+  do i = 1, rshift (N, 1)
+    if (a(i) .ne. 6.0) call abort
+    if (b(i) .ne. 6.0) call abort
+  end do
+
+  do i = rshift (N, 1) + 1, N
+    if (a(i) .ne. 5.0) call abort
+    if (b(i) .ne. 5.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 0.0
+  end do
+
+  !$acc update device (a(1:4))
+
+  !$acc parallel present (a)
+    do i = 1, N
+      a(i) = a(i) + 1.0
+    end do
+  !$acc end parallel
+
+  !$acc update self (a(5:N))
+
+  do i = 1, rshift (N, 1)
+    if (a(i) .ne. 0.0) call abort
+  end do
+
+  do i = rshift (N, 1) + 1, N
+    if (a(i) .ne. 6.0) call abort
+  end do
+
+  !$acc update self (a(1:4))
+
+  do i = 1, rshift (N, 1)
+    if (a(i) .ne. 1.0) call abort
+  end do
+
+  do i = rshift (N, 1) + 1, N
+    if (a(i) .ne. 6.0) call abort
+  end do
+
+  a(3) = 9
+  a(4) = 9
+  a(5) = 9
+  a(6) = 9
+
+  !$acc update device (a(3:6))
+
+  !$acc parallel present (a(1:N))
+    do i = 1, N
+      a(i) = a(i) + 1.0
+    end do
+  !$acc end parallel
+
+  !$acc update self (a(3:6))
+
+  do i = 1, 2
+    if (a(i) .ne. 1.0) call abort
+  end do
+
+  do i = 3, 6
+    if (a(i) .ne. 10.0) call abort
+  end do
+
+  do i = 7, N
+    if (a(i) .ne. 6.0) call abort
+  end do
+  !$acc exit data delete (a, b)
+
+end program
+
Index: testsuite/libgomp.oacc-fortran/update-1.f90
===================================================================
--- testsuite/libgomp.oacc-fortran/update-1.f90	(revision 0)
+++ testsuite/libgomp.oacc-fortran/update-1.f90	(revision 232741)
@@ -0,0 +1,242 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+program update
+  use openacc
+  implicit none 
+  integer, parameter :: N = 8
+  integer, parameter :: NDIV2 = N / 2
+  real :: a(N), b(N)
+  integer i
+
+  do i = 1, N
+    a(i) = 3.0
+    b(i) = 0.0
+  end do
+
+  !$acc enter data copyin (a, b)
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 3.0) call abort
+    if (b(i) .ne. 3.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 5.0
+    b(i) = 1.0
+  end do
+
+  !$acc update device (a, b)
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do 
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 5.0) call abort
+    if (b(i) .ne. 5.0) call abort
+ end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  !$acc parallel present (a, b)
+  do i = 1, N
+    b(i) = a(i)
+  end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 5.0) call abort
+    if (b(i) .ne. 5.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 6.0
+    b(i) = 0.0
+  end do
+
+  !$acc update device (a, b)
+
+  do i = 1, N
+    a(i) = 9.0
+  end do
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 6.0) call abort
+    if (b(i) .ne. 6.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 7.0
+    b(i) = 2.0
+  end do
+
+  !$acc update device (a, b)
+
+  do i = 1, N
+    a(i) = 9.0
+  end do
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 7.0) call abort
+    if (b(i) .ne. 7.0) call abort
+  end do
+
+  do i = 1, N
+    a(i) = 9.0
+  end do
+
+  !$acc update device (a)
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 9.0) call abort
+    if (b(i) .ne. 9.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 5.0
+  end do
+
+  !$acc update device (a)
+
+  do i = 1, N
+    a(i) = 6.0
+  end do
+
+  !$acc update device (a(1:NDIV2))
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, NDIV2
+    if (a(i) .ne. 6.0) call abort
+    if (b(i) .ne. 6.0) call abort
+  end do
+
+  do i = NDIV2 + 1, N
+    if (a(i) .ne. 5.0) call abort
+    if (b(i) .ne. 5.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 0.0
+  end do
+
+  !$acc update device (a(1:4))
+
+  !$acc parallel present (a)
+    do i = 1, N
+      a(i) = a(i) + 1.0
+    end do
+  !$acc end parallel
+
+  !$acc update host (a(5:N))
+
+  do i = 1, NDIV2
+    if (a(i) .ne. 0.0) call abort
+  end do
+
+  do i = NDIV2 + 1, N
+    if (a(i) .ne. 6.0) call abort
+  end do
+
+  !$acc update host (a(1:4))
+
+  do i = 1, NDIV2
+    if (a(i) .ne. 1.0) call abort
+  end do
+
+  do i = NDIV2 + 1, N
+    if (a(i) .ne. 6.0) call abort
+  end do
+
+  a(3) = 9
+  a(4) = 9
+  a(5) = 9
+  a(6) = 9
+
+  !$acc update device (a(3:6))
+
+  !$acc parallel present (a(1:N))
+    do i = 1, N
+      a(i) = a(i) + 1.0
+    end do
+  !$acc end parallel
+
+  !$acc update host (a(3:6))
+
+  do i = 1, 2
+    if (a(i) .ne. 1.0) call abort
+  end do
+
+  do i = 3, 6
+    if (a(i) .ne. 10.0) call abort
+  end do
+
+  do i = 7, N
+    if (a(i) .ne. 6.0) call abort
+  end do
+
+  !$acc exit data delete (a, b)
+
+end program
+
Index: testsuite/libgomp.oacc-c-c++-common/update-1-2.c
===================================================================
--- testsuite/libgomp.oacc-c-c++-common/update-1-2.c	(revision 232740)
+++ testsuite/libgomp.oacc-c-c++-common/update-1-2.c	(revision 232741)
@@ -13,6 +13,7 @@
 main (int argc, char **argv)
 {
     int N = 8;
+    int NDIV2 = N / 2;
     float *a, *b, *c;
     float *d_a, *d_b, *d_c;
     int i;
@@ -242,7 +243,7 @@
         a[i] = 6.0;
     }
 
-#pragma acc update device (a[0:N >> 1])
+#pragma acc update device (a[0:NDIV2])
 
 #pragma acc parallel present (a[0:N], b[0:N])
     {
@@ -254,7 +255,7 @@
 
 #pragma acc update self (a[0:N], b[0:N])
 
-    for (i = 0; i < (N >> 1); i++)
+    for (i = 0; i < NDIV2; i++)
     {
         if (a[i] != 6.0)
             abort ();
@@ -263,7 +264,7 @@
             abort ();
     }
 
-    for (i = (N >> 1); i < N; i++)
+    for (i = NDIV2; i < N; i++)
     {
         if (a[i] != 5.0)
             abort ();
@@ -278,5 +279,83 @@
     if (!acc_is_present (&b[0], (N * sizeof (float))))
       abort ();
 
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 0.0;
+    }
+
+#pragma acc update device (a[0:4])
+
+#pragma acc parallel present (a[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            a[ii] = a[ii] + 1.0;
+    }
+
+#pragma acc update self (a[4:4])
+
+    for (i = 0; i < NDIV2; i++)
+    {
+        if (a[i] != 0.0)
+            abort ();
+    }
+
+    for (i = NDIV2; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
+#pragma acc update self (a[0:4])
+
+    for (i = 0; i < NDIV2; i++)
+    {
+        if (a[i] != 1.0)
+            abort ();
+    }
+
+    for (i = NDIV2; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
+    a[2] = 9;
+    a[3] = 9;
+    a[4] = 9;
+    a[5] = 9;
+
+#pragma acc update device (a[2:4])
+
+#pragma acc parallel present (a[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            a[ii] = a[ii] + 1.0;
+    }
+
+#pragma acc update self (a[2:4])
+
+    for (i = 0; i < 2; i++)
+    {
+      if (a[i] != 1.0)
+	abort ();
+    }
+
+    for (i = 2; i < 6; i++)
+    {
+      if (a[i] != 10.0)
+	abort ();
+    }
+
+    for (i = 6; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
     return 0;
 }
Index: testsuite/libgomp.oacc-c-c++-common/update-1.c
===================================================================
--- testsuite/libgomp.oacc-c-c++-common/update-1.c	(revision 232740)
+++ testsuite/libgomp.oacc-c-c++-common/update-1.c	(revision 232741)
@@ -11,6 +11,7 @@
 main (int argc, char **argv)
 {
     int N = 8;
+    int NDIV2 = N / 2;
     float *a, *b, *c;
     float *d_a, *d_b, *d_c;
     int i;
@@ -109,7 +110,7 @@
             b[ii] = a[ii];
     }
 
-#pragma acc update self (a[0:N], b[0:N])
+#pragma acc update host (a[0:N], b[0:N])
 
     for (i = 0; i < N; i++)
     {
@@ -240,7 +241,7 @@
         a[i] = 6.0;
     }
 
-#pragma acc update device (a[0:N >> 1])
+#pragma acc update device (a[0:NDIV2])
 
 #pragma acc parallel present (a[0:N], b[0:N])
     {
@@ -252,7 +253,7 @@
 
 #pragma acc update host (a[0:N], b[0:N])
 
-    for (i = 0; i < (N >> 1); i++)
+    for (i = 0; i < NDIV2; i++)
     {
         if (a[i] != 6.0)
             abort ();
@@ -261,7 +262,7 @@
             abort ();
     }
 
-    for (i = (N >> 1); i < N; i++)
+    for (i = NDIV2; i < N; i++)
     {
         if (a[i] != 5.0)
             abort ();
@@ -276,5 +277,83 @@
     if (!acc_is_present (&b[0], (N * sizeof (float))))
       abort ();
 
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 0.0;
+    }
+
+#pragma acc update device (a[0:4])
+
+#pragma acc parallel present (a[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            a[ii] = a[ii] + 1.0;
+    }
+
+#pragma acc update host (a[4:4])
+
+    for (i = 0; i < NDIV2; i++)
+    {
+        if (a[i] != 0.0)
+            abort ();
+    }
+
+    for (i = NDIV2; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
+#pragma acc update host (a[0:4])
+
+    for (i = 0; i < NDIV2; i++)
+    {
+        if (a[i] != 1.0)
+            abort ();
+    }
+
+    for (i = NDIV2; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
+    a[2] = 9;
+    a[3] = 9;
+    a[4] = 9;
+    a[5] = 9;
+
+#pragma acc update device (a[2:4])
+
+#pragma acc parallel present (a[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            a[ii] = a[ii] + 1.0;
+    }
+
+#pragma acc update host (a[2:4])
+
+    for (i = 0; i < 2; i++)
+    {
+      if (a[i] != 1.0)
+	abort ();
+    }
+
+    for (i = 2; i < 6; i++)
+    {
+      if (a[i] != 10.0)
+	abort ();
+    }
+
+    for (i = 6; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
     return 0;
 }

Re: [gomp4] Fix handling of subarrays with update directive

[Thomas Schwinge]

Hi!

On Fri, 22 Jan 2016 12:50:38 -0600, James Norris <jnorris@codesourcery.com> wrote:
> The attached patch fixes a defect reported with gcc 5.2
> (https://gcc.gnu.org/bugzilla/show_bug.cgi?id=69414).
> It is also the case, the issue is present in the gomp4
> branch. The patch also adds additional testing.
> 
> Committed to gomp4 after bootstrap and regtesting.

Jakub, is that following two-line patch (libgomp/oacc-mem.c, plus test
cases) OK for trunk?

> --- ChangeLog.gomp	(revision 232740)
> +++ ChangeLog.gomp	(revision 232741)
> @@ -1,3 +1,11 @@
> +2016-01-22  James Norris  <jnorris@codesourcery.com>
> +
> +	* oacc-mem.c (delete_copyout, update_dev_host): Fix device address.

Given that Daichi Fukuoka submitted the patch in
<https://gcc.gnu.org/PR69414>, please acknowledge that in the ChangeLog.
Please also include the test case submitted there; rename their file from
gcc1.f90 to pr69414.f90 (or similar).

> --- oacc-mem.c	(revision 232740)
> +++ oacc-mem.c	(revision 232741)
> @@ -509,7 +509,7 @@
>        gomp_fatal ("[%p,%d] is not mapped", (void *)h, (int)s);
>      }
>  
> -  d = (void *) (n->tgt->tgt_start + n->tgt_offset);
> +  d = (void *) (n->tgt->tgt_start + n->tgt_offset + h - n->host_start);
>  
>    host_size = n->host_end - n->host_start;
>  
> @@ -562,7 +562,7 @@
>        gomp_fatal ("[%p,%d] is not mapped", h, (int)s);
>      }
>  
> -  d = (void *) (n->tgt->tgt_start + n->tgt_offset);
> +  d = (void *) (n->tgt->tgt_start + n->tgt_offset + h - n->host_start);
>  
>    gomp_mutex_unlock (&acc_dev->lock);
>  

> --- testsuite/libgomp.oacc-fortran/update-1-2.f90	(revision 0)
> +++ testsuite/libgomp.oacc-fortran/update-1-2.f90	(revision 232741)
> @@ -0,0 +1,239 @@
> +! Copy of update-1.f90 with self exchanged with host for !$acc update

Actually, that file contains additional changes.

Anyway, I suggest we don't add new test cases with just self vs. host
clauses exchanged -- their equivalence is something that should (already)
be tested in GCC front end test cases.  I know that we currently have a
distinct libgomp.oacc-c-c++-common/update-1-2.c test case as a "copy of
update-1.c with self exchanged with host for #pragma acc update", once
added by me; I suggest we remove that.  Let's just in the same file use a
mixture of host and self clauses.  (I'll take care of that.)


Grüße
 Thomas

Re: [gomp4] Fix handling of subarrays with update directive

[Jakub Jelinek]

On Wed, Mar 23, 2016 at 09:35:30AM +0100, Thomas Schwinge wrote:
> Hi!
> 
> On Fri, 22 Jan 2016 12:50:38 -0600, James Norris <jnorris@codesourcery.com> wrote:
> > The attached patch fixes a defect reported with gcc 5.2
> > (https://gcc.gnu.org/bugzilla/show_bug.cgi?id=69414).
> > It is also the case, the issue is present in the gomp4
> > branch. The patch also adds additional testing.
> > 
> > Committed to gomp4 after bootstrap and regtesting.
> 
> Jakub, is that following two-line patch (libgomp/oacc-mem.c, plus test
> cases) OK for trunk?
> 
> > --- ChangeLog.gomp	(revision 232740)
> > +++ ChangeLog.gomp	(revision 232741)
> > @@ -1,3 +1,11 @@
> > +2016-01-22  James Norris  <jnorris@codesourcery.com>
> > +
> > +	* oacc-mem.c (delete_copyout, update_dev_host): Fix device address.
> 
> Given that Daichi Fukuoka submitted the patch in
> <https://gcc.gnu.org/PR69414>, please acknowledge that in the ChangeLog.
> Please also include the test case submitted there; rename their file from
> gcc1.f90 to pr69414.f90 (or similar).

Yeah, the oacc-mem.c is a two liner, we don't need copyright assignment for
that, so please mention Daichi-san's name + mail address, and PR libgpmp/69414
in the ChangeLog entry.

> > --- oacc-mem.c	(revision 232740)
> > +++ oacc-mem.c	(revision 232741)
> > @@ -509,7 +509,7 @@
> >        gomp_fatal ("[%p,%d] is not mapped", (void *)h, (int)s);
> >      }
> >  
> > -  d = (void *) (n->tgt->tgt_start + n->tgt_offset);
> > +  d = (void *) (n->tgt->tgt_start + n->tgt_offset + h - n->host_start);

I'd prefer (uintptr_t) h instead of, and you probably need to wrap it then.

> >    host_size = n->host_end - n->host_start;
> >  
> > @@ -562,7 +562,7 @@
> >        gomp_fatal ("[%p,%d] is not mapped", h, (int)s);
> >      }
> >  
> > -  d = (void *) (n->tgt->tgt_start + n->tgt_offset);
> > +  d = (void *) (n->tgt->tgt_start + n->tgt_offset + h - n->host_start);

Ditto.

> > --- testsuite/libgomp.oacc-fortran/update-1-2.f90	(revision 0)
> > +++ testsuite/libgomp.oacc-fortran/update-1-2.f90	(revision 232741)
> > @@ -0,0 +1,239 @@
> > +! Copy of update-1.f90 with self exchanged with host for !$acc update
> 
> Actually, that file contains additional changes.
> 
> Anyway, I suggest we don't add new test cases with just self vs. host
> clauses exchanged -- their equivalence is something that should (already)
> be tested in GCC front end test cases.  I know that we currently have a
> distinct libgomp.oacc-c-c++-common/update-1-2.c test case as a "copy of
> update-1.c with self exchanged with host for #pragma acc update", once
> added by me; I suggest we remove that.  Let's just in the same file use a
> mixture of host and self clauses.  (I'll take care of that.)

Otherwise LGTM, but please repost it with all the testcase changes you want
to make.

	Jakub

RE: [gomp4] Fix handling of subarrays with update directive

[Fukuoka Daichi]

Hi,

> Yeah, the oacc-mem.c is a two liner, we don't need copyright assignment for that, so please mention Daichi-san's name + mail address, and PR libgpmp/69414 in the ChangeLog entry.

If possible, could you add " SGI Japan, Ltd." before/after my name?
Because I am belonging to the company.
Thank you.

Regards,
Daichi Fukuoka

-----Original Message-----
From: Jakub Jelinek [mailto:jakub@redhat.com] 
Sent: Wednesday, March 23, 2016 7:24 PM
To: Thomas Schwinge
Cc: James Norris; GCC Patches; Fukuoka Daichi
Subject: Re: [gomp4] Fix handling of subarrays with update directive

On Wed, Mar 23, 2016 at 09:35:30AM +0100, Thomas Schwinge wrote:
> Hi!
> 
> On Fri, 22 Jan 2016 12:50:38 -0600, James Norris <jnorris@codesourcery.com> wrote:
> > The attached patch fixes a defect reported with gcc 5.2 
> > (https://gcc.gnu.org/bugzilla/show_bug.cgi?id=69414).
> > It is also the case, the issue is present in the gomp4 branch. The 
> > patch also adds additional testing.
> > 
> > Committed to gomp4 after bootstrap and regtesting.
> 
> Jakub, is that following two-line patch (libgomp/oacc-mem.c, plus test
> cases) OK for trunk?
> 
> > --- ChangeLog.gomp	(revision 232740)
> > +++ ChangeLog.gomp	(revision 232741)
> > @@ -1,3 +1,11 @@
> > +2016-01-22  James Norris  <jnorris@codesourcery.com>
> > +
> > +	* oacc-mem.c (delete_copyout, update_dev_host): Fix device address.
> 
> Given that Daichi Fukuoka submitted the patch in 
> <https://gcc.gnu.org/PR69414>, please acknowledge that in the ChangeLog.
> Please also include the test case submitted there; rename their file 
> from
> gcc1.f90 to pr69414.f90 (or similar).

Yeah, the oacc-mem.c is a two liner, we don't need copyright assignment for that, so please mention Daichi-san's name + mail address, and PR libgpmp/69414 in the ChangeLog entry.

> > --- oacc-mem.c	(revision 232740)
> > +++ oacc-mem.c	(revision 232741)
> > @@ -509,7 +509,7 @@
> >        gomp_fatal ("[%p,%d] is not mapped", (void *)h, (int)s);
> >      }
> >  
> > -  d = (void *) (n->tgt->tgt_start + n->tgt_offset);
> > +  d = (void *) (n->tgt->tgt_start + n->tgt_offset + h - 
> > + n->host_start);

I'd prefer (uintptr_t) h instead of, and you probably need to wrap it then.

> >    host_size = n->host_end - n->host_start;
> >  
> > @@ -562,7 +562,7 @@
> >        gomp_fatal ("[%p,%d] is not mapped", h, (int)s);
> >      }
> >  
> > -  d = (void *) (n->tgt->tgt_start + n->tgt_offset);
> > +  d = (void *) (n->tgt->tgt_start + n->tgt_offset + h - 
> > + n->host_start);

Ditto.

> > --- testsuite/libgomp.oacc-fortran/update-1-2.f90	(revision 0)
> > +++ testsuite/libgomp.oacc-fortran/update-1-2.f90	(revision 232741)
> > @@ -0,0 +1,239 @@
> > +! Copy of update-1.f90 with self exchanged with host for !$acc 
> > +update
> 
> Actually, that file contains additional changes.
> 
> Anyway, I suggest we don't add new test cases with just self vs. host 
> clauses exchanged -- their equivalence is something that should 
> (already) be tested in GCC front end test cases.  I know that we 
> currently have a distinct libgomp.oacc-c-c++-common/update-1-2.c test 
> case as a "copy of update-1.c with self exchanged with host for 
> #pragma acc update", once added by me; I suggest we remove that.  
> Let's just in the same file use a mixture of host and self clauses.  
> (I'll take care of that.)

Otherwise LGTM, but please repost it with all the testcase changes you want to make.

	Jakub

Re: [gomp4] Fix handling of subarrays with update directive

[Jakub Jelinek]

On Wed, Mar 23, 2016 at 10:31:33AM +0000, Fukuoka Daichi wrote:
> > Yeah, the oacc-mem.c is a two liner, we don't need copyright assignment for that, so please mention Daichi-san's name + mail address, and PR libgpmp/69414 in the ChangeLog entry.
> 
> If possible, could you add " SGI Japan, Ltd." before/after my name?
> Because I am belonging to the company.

The ChangeLog format is <date>  <name> <surname>  <email>
thus company names aren't listed explicitly, but your email is @sgi.com, so
it will be there.

	Jakub

RE: [gomp4] Fix handling of subarrays with update directive

[Fukuoka Daichi]

Jakub,

> The ChangeLog format is <date>  <name> <surname>  <email> thus company names aren't listed explicitly, but your email is @sgi.com, so it will be there.

I understand the format, thanks.

Regards,
Daichi Fukuoka

-----Original Message-----
From: Jakub Jelinek [mailto:jakub@redhat.com] 
Sent: Wednesday, March 23, 2016 7:37 PM
To: Fukuoka Daichi
Cc: James Norris; Thomas Schwinge; GCC Patches
Subject: Re: [gomp4] Fix handling of subarrays with update directive

On Wed, Mar 23, 2016 at 10:31:33AM +0000, Fukuoka Daichi wrote:
> > Yeah, the oacc-mem.c is a two liner, we don't need copyright assignment for that, so please mention Daichi-san's name + mail address, and PR libgpmp/69414 in the ChangeLog entry.
> 
> If possible, could you add " SGI Japan, Ltd." before/after my name?
> Because I am belonging to the company.

The ChangeLog format is <date>  <name> <surname>  <email> thus company names aren't listed explicitly, but your email is @sgi.com, so it will be there.

	Jakub

Re: [gomp4] Fix handling of subarrays with update directive

[James Norris]

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

Jakub,

On 03/23/2016 05:24 AM, Jakub Jelinek wrote:
> Otherwise LGTM, but please repost it with all the testcase changes you want
> to make.

Attached is the updated patch with the castings added, as well as
the updated tests.

Thanks!
Jim

======

2016-03-23  James Norris  <jnorris@codesourcery.com>
             Daichi Fukuoka <dc-fukuoka@sgi.com>

         * oacc-mem.c (delete_copyout, update_dev_host): Fix device address.
         * testsuite/libgomp.oacc-c-c++-common/update-1.c: Additional tests.
         * testsuite/libgomp.oacc-c-c++-common/update-1-2.c: Likewise.
         * testsuite/libgomp.oacc-fortran/update-1.f90: New file.


[-- Attachment #2: bugfix.patch --]
[-- Type: text/x-patch, Size: 11095 bytes --]

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index fca65e6..66dc6a0 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,11 @@
+2016-03-22  James Norris  <jnorris@codesourcery.com>
+	    Daichi Fukuoka <dc-fukuoka@sgi.com>
+
+	* oacc-mem.c (delete_copyout, update_dev_host): Fix device address.
+	* testsuite/libgomp.oacc-c-c++-common/update-1.c: Additional tests.
+	* testsuite/libgomp.oacc-c-c++-common/update-1-2.c: Likewise.
+	* testsuite/libgomp.oacc-fortran/update-1.f90: New file.
+
 2016-03-16  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* testsuite/libgomp.oacc-fortran/kernels-loop-2.f95: Adjust to
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index f6cc373..487691e 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -509,7 +509,8 @@ delete_copyout (unsigned f, void *h, size_t s)
       gomp_fatal ("[%p,%d] is not mapped", (void *)h, (int)s);
     }
 
-  d = (void *) (n->tgt->tgt_start + n->tgt_offset);
+  d = (void *) (n->tgt->tgt_start + n->tgt_offset
+	        + (uintptr_t) h - n->host_start);
 
   host_size = n->host_end - n->host_start;
 
@@ -562,7 +563,8 @@ update_dev_host (int is_dev, void *h, size_t s)
       gomp_fatal ("[%p,%d] is not mapped", h, (int)s);
     }
 
-  d = (void *) (n->tgt->tgt_start + n->tgt_offset);
+  d = (void *) (n->tgt->tgt_start + n->tgt_offset
+		+ (uintptr_t) h - n->host_start);
 
   gomp_mutex_unlock (&acc_dev->lock);
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c
index c7e7257..82c3192 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c
@@ -13,6 +13,7 @@ int
 main (int argc, char **argv)
 {
     int N = 8;
+    int NDIV2 = N / 2;
     float *a, *b, *c;
     float *d_a, *d_b, *d_c;
     int i;
@@ -242,7 +243,7 @@ main (int argc, char **argv)
         a[i] = 6.0;
     }
 
-#pragma acc update device (a[0:N >> 1])
+#pragma acc update device (a[0:NDIV2])
 
 #pragma acc parallel present (a[0:N], b[0:N])
     {
@@ -254,7 +255,7 @@ main (int argc, char **argv)
 
 #pragma acc update self (a[0:N], b[0:N])
 
-    for (i = 0; i < (N >> 1); i++)
+    for (i = 0; i < NDIV2; i++)
     {
         if (a[i] != 6.0)
             abort ();
@@ -263,7 +264,7 @@ main (int argc, char **argv)
             abort ();
     }
 
-    for (i = (N >> 1); i < N; i++)
+    for (i = NDIV2; i < N; i++)
     {
         if (a[i] != 5.0)
             abort ();
@@ -278,5 +279,83 @@ main (int argc, char **argv)
     if (!acc_is_present (&b[0], (N * sizeof (float))))
       abort ();
 
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 0.0;
+    }
+
+#pragma acc update device (a[0:4])
+
+#pragma acc parallel present (a[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            a[ii] = a[ii] + 1.0;
+    }
+
+#pragma acc update self (a[4:4])
+
+    for (i = 0; i < NDIV2; i++)
+    {
+        if (a[i] != 0.0)
+            abort ();
+    }
+
+    for (i = NDIV2; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
+#pragma acc update self (a[0:4])
+
+    for (i = 0; i < NDIV2; i++)
+    {
+        if (a[i] != 1.0)
+            abort ();
+    }
+
+    for (i = NDIV2; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
+    a[2] = 9;
+    a[3] = 9;
+    a[4] = 9;
+    a[5] = 9;
+
+#pragma acc update device (a[2:4])
+
+#pragma acc parallel present (a[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            a[ii] = a[ii] + 1.0;
+    }
+
+#pragma acc update self (a[2:4])
+
+    for (i = 0; i < 2; i++)
+    {
+      if (a[i] != 1.0)
+	abort ();
+    }
+
+    for (i = 2; i < 6; i++)
+    {
+      if (a[i] != 10.0)
+	abort ();
+    }
+
+    for (i = 6; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
     return 0;
 }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/update-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/update-1.c
index dff139f..1b2a460 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/update-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/update-1.c
@@ -11,6 +11,7 @@ int
 main (int argc, char **argv)
 {
     int N = 8;
+    int NDIV2 = N / 2;
     float *a, *b, *c;
     float *d_a, *d_b, *d_c;
     int i;
@@ -109,7 +110,7 @@ main (int argc, char **argv)
             b[ii] = a[ii];
     }
 
-#pragma acc update self (a[0:N], b[0:N])
+#pragma acc update host (a[0:N], b[0:N])
 
     for (i = 0; i < N; i++)
     {
@@ -240,7 +241,7 @@ main (int argc, char **argv)
         a[i] = 6.0;
     }
 
-#pragma acc update device (a[0:N >> 1])
+#pragma acc update device (a[0:NDIV2])
 
 #pragma acc parallel present (a[0:N], b[0:N])
     {
@@ -252,7 +253,7 @@ main (int argc, char **argv)
 
 #pragma acc update host (a[0:N], b[0:N])
 
-    for (i = 0; i < (N >> 1); i++)
+    for (i = 0; i < NDIV2; i++)
     {
         if (a[i] != 6.0)
             abort ();
@@ -261,7 +262,7 @@ main (int argc, char **argv)
             abort ();
     }
 
-    for (i = (N >> 1); i < N; i++)
+    for (i = NDIV2; i < N; i++)
     {
         if (a[i] != 5.0)
             abort ();
@@ -276,5 +277,83 @@ main (int argc, char **argv)
     if (!acc_is_present (&b[0], (N * sizeof (float))))
       abort ();
 
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 0.0;
+    }
+
+#pragma acc update device (a[0:4])
+
+#pragma acc parallel present (a[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            a[ii] = a[ii] + 1.0;
+    }
+
+#pragma acc update host (a[4:4])
+
+    for (i = 0; i < NDIV2; i++)
+    {
+        if (a[i] != 0.0)
+            abort ();
+    }
+
+    for (i = NDIV2; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
+#pragma acc update host (a[0:4])
+
+    for (i = 0; i < NDIV2; i++)
+    {
+        if (a[i] != 1.0)
+            abort ();
+    }
+
+    for (i = NDIV2; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
+    a[2] = 9;
+    a[3] = 9;
+    a[4] = 9;
+    a[5] = 9;
+
+#pragma acc update device (a[2:4])
+
+#pragma acc parallel present (a[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            a[ii] = a[ii] + 1.0;
+    }
+
+#pragma acc update host (a[2:4])
+
+    for (i = 0; i < 2; i++)
+    {
+      if (a[i] != 1.0)
+	abort ();
+    }
+
+    for (i = 2; i < 6; i++)
+    {
+      if (a[i] != 10.0)
+	abort ();
+    }
+
+    for (i = 6; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
     return 0;
 }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/update-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/update-1.f90
new file mode 100644
index 0000000..4e1d2c0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/update-1.f90
@@ -0,0 +1,242 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+program update
+  use openacc
+  implicit none 
+  integer, parameter :: N = 8
+  integer, parameter :: NDIV2 = N / 2
+  real :: a(N), b(N)
+  integer i
+
+  do i = 1, N
+    a(i) = 3.0
+    b(i) = 0.0
+  end do
+
+  !$acc enter data copyin (a, b)
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 3.0) call abort
+    if (b(i) .ne. 3.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 5.0
+    b(i) = 1.0
+  end do
+
+  !$acc update device (a, b)
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do 
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 5.0) call abort
+    if (b(i) .ne. 5.0) call abort
+ end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  !$acc parallel present (a, b)
+  do i = 1, N
+    b(i) = a(i)
+  end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 5.0) call abort
+    if (b(i) .ne. 5.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 6.0
+    b(i) = 0.0
+  end do
+
+  !$acc update device (a, b)
+
+  do i = 1, N
+    a(i) = 9.0
+  end do
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 6.0) call abort
+    if (b(i) .ne. 6.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 7.0
+    b(i) = 2.0
+  end do
+
+  !$acc update device (a, b)
+
+  do i = 1, N
+    a(i) = 9.0
+  end do
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 7.0) call abort
+    if (b(i) .ne. 7.0) call abort
+  end do
+
+  do i = 1, N
+    a(i) = 9.0
+  end do
+
+  !$acc update device (a)
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 9.0) call abort
+    if (b(i) .ne. 9.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 5.0
+  end do
+
+  !$acc update device (a)
+
+  do i = 1, N
+    a(i) = 6.0
+  end do
+
+  !$acc update device (a(1:NDIV2))
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, NDIV2
+    if (a(i) .ne. 6.0) call abort
+    if (b(i) .ne. 6.0) call abort
+  end do
+
+  do i = NDIV2 + 1, N
+    if (a(i) .ne. 5.0) call abort
+    if (b(i) .ne. 5.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 0.0
+  end do
+
+  !$acc update device (a(1:4))
+
+  !$acc parallel present (a)
+    do i = 1, N
+      a(i) = a(i) + 1.0
+    end do
+  !$acc end parallel
+
+  !$acc update host (a(5:N))
+
+  do i = 1, NDIV2
+    if (a(i) .ne. 0.0) call abort
+  end do
+
+  do i = NDIV2 + 1, N
+    if (a(i) .ne. 6.0) call abort
+  end do
+
+  !$acc update host (a(1:4))
+
+  do i = 1, NDIV2
+    if (a(i) .ne. 1.0) call abort
+  end do
+
+  do i = NDIV2 + 1, N
+    if (a(i) .ne. 6.0) call abort
+  end do
+
+  a(3) = 9
+  a(4) = 9
+  a(5) = 9
+  a(6) = 9
+
+  !$acc update device (a(3:6))
+
+  !$acc parallel present (a(1:N))
+    do i = 1, N
+      a(i) = a(i) + 1.0
+    end do
+  !$acc end parallel
+
+  !$acc update host (a(3:6))
+
+  do i = 1, 2
+    if (a(i) .ne. 1.0) call abort
+  end do
+
+  do i = 3, 6
+    if (a(i) .ne. 10.0) call abort
+  end do
+
+  do i = 7, N
+    if (a(i) .ne. 6.0) call abort
+  end do
+
+  !$acc exit data delete (a, b)
+
+end program
+

Re: [gomp4] Fix handling of subarrays with update directive

[Jakub Jelinek]

On Wed, Mar 23, 2016 at 08:05:19AM -0500, James Norris wrote:
> Jakub,
> 
> On 03/23/2016 05:24 AM, Jakub Jelinek wrote:
> >Otherwise LGTM, but please repost it with all the testcase changes you want
> >to make.
> 
> Attached is the updated patch with the castings added, as well as
> the updated tests.
> 
> Thanks!
> Jim
> 
> ======
> 
> 2016-03-23  James Norris  <jnorris@codesourcery.com>
>             Daichi Fukuoka <dc-fukuoka@sgi.com>
> 

Missing
	PR libgomp/69414
>         * oacc-mem.c (delete_copyout, update_dev_host): Fix device address.
>         * testsuite/libgomp.oacc-c-c++-common/update-1.c: Additional tests.
>         * testsuite/libgomp.oacc-c-c++-common/update-1-2.c: Likewise.
>         * testsuite/libgomp.oacc-fortran/update-1.f90: New file.

Ok with that change.


	Jakub

[gcc-5] Fix handling of subarrays with update directive

[Thomas Schwinge]

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

Hi Jakub!

On Wed, 23 Mar 2016 14:10:31 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> On Wed, Mar 23, 2016 at 08:05:19AM -0500, James Norris wrote:
> > On 03/23/2016 05:24 AM, Jakub Jelinek wrote:
> > 2016-03-23  James Norris  <jnorris@codesourcery.com>
> >             Daichi Fukuoka <dc-fukuoka@sgi.com>

> 	PR libgomp/69414
> >         * oacc-mem.c (delete_copyout, update_dev_host): Fix device address.
> >         * testsuite/libgomp.oacc-c-c++-common/update-1.c: Additional tests.
> >         * testsuite/libgomp.oacc-c-c++-common/update-1-2.c: Likewise.
> >         * testsuite/libgomp.oacc-fortran/update-1.f90: New file.
> 
> Ok with that change.

OK to backport that commit to gcc-5-branch (which it has been reported
against)?


Grüße
 Thomas

[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]

Re: [gcc-5] Fix handling of subarrays with update directive

[Jakub Jelinek]

On Thu, Apr 07, 2016 at 01:28:48PM +0200, Thomas Schwinge wrote:
> Hi Jakub!
> 
> On Wed, 23 Mar 2016 14:10:31 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> > On Wed, Mar 23, 2016 at 08:05:19AM -0500, James Norris wrote:
> > > On 03/23/2016 05:24 AM, Jakub Jelinek wrote:
> > > 2016-03-23  James Norris  <jnorris@codesourcery.com>
> > >             Daichi Fukuoka <dc-fukuoka@sgi.com>
> 
> > 	PR libgomp/69414
> > >         * oacc-mem.c (delete_copyout, update_dev_host): Fix device address.
> > >         * testsuite/libgomp.oacc-c-c++-common/update-1.c: Additional tests.
> > >         * testsuite/libgomp.oacc-c-c++-common/update-1-2.c: Likewise.
> > >         * testsuite/libgomp.oacc-fortran/update-1.f90: New file.
> > 
> > Ok with that change.
> 
> OK to backport that commit to gcc-5-branch (which it has been reported
> against)?

Ok.

	Jakub

Re: [gcc-5] Fix handling of subarrays with update directive

[Thomas Schwinge]

Hi!

On Thu, 7 Apr 2016 13:29:58 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Thu, Apr 07, 2016 at 01:28:48PM +0200, Thomas Schwinge wrote:
> > On Wed, 23 Mar 2016 14:10:31 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> > > On Wed, Mar 23, 2016 at 08:05:19AM -0500, James Norris wrote:
> > > > On 03/23/2016 05:24 AM, Jakub Jelinek wrote:
> > > > 2016-03-23  James Norris  <jnorris@codesourcery.com>
> > > >             Daichi Fukuoka <dc-fukuoka@sgi.com>
> > 
> > > 	PR libgomp/69414
> > > >         * oacc-mem.c (delete_copyout, update_dev_host): Fix device address.
> > > >         * testsuite/libgomp.oacc-c-c++-common/update-1.c: Additional tests.
> > > >         * testsuite/libgomp.oacc-c-c++-common/update-1-2.c: Likewise.
> > > >         * testsuite/libgomp.oacc-fortran/update-1.f90: New file.
> > > 
> > > Ok with that change.
> > 
> > OK to backport that commit to gcc-5-branch (which it has been reported
> > against)?
> 
> Ok.

In r234806 committed to gcc-5-branch (without changes):

commit 09222d2f8af5e1d4b07d56a56b6806b674af2952
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Thu Apr 7 11:43:30 2016 +0000

    [PR libgomp/69414] Fix handling of subarrays with update directive
    
    	libgomp/
    	Backport trunk r234428:
    
    	2016-03-23  James Norris  <jnorris@codesourcery.com>
    		    Daichi Fukuoka <dc-fukuoka@sgi.com>
    
    	PR libgomp/69414
    	* oacc-mem.c (delete_copyout, update_dev_host): Fix device address.
    	* testsuite/libgomp.oacc-c-c++-common/update-1.c: Additional tests.
    	* testsuite/libgomp.oacc-c-c++-common/update-1-2.c: Likewise.
    	* testsuite/libgomp.oacc-fortran/update-1.f90: New file.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gcc-5-branch@234806 138bc75d-0d04-0410-961f-82ee72b054a4
---
 libgomp/ChangeLog                                  |   13 ++
 libgomp/oacc-mem.c                                 |    6 +-
 .../libgomp.oacc-c-c++-common/update-1-2.c         |   85 ++++++-
 .../testsuite/libgomp.oacc-c-c++-common/update-1.c |   87 ++++++-
 .../testsuite/libgomp.oacc-fortran/update-1.f90    |  242 ++++++++++++++++++++
 5 files changed, 424 insertions(+), 9 deletions(-)

diff --git libgomp/ChangeLog libgomp/ChangeLog
index 3da9fa1..ed890e0 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,3 +1,16 @@
+2016-04-07  Thomas Schwinge  <thomas@codesourcery.com>
+
+	Backport trunk r234428:
+
+	2016-03-23  James Norris  <jnorris@codesourcery.com>
+		    Daichi Fukuoka <dc-fukuoka@sgi.com>
+
+	PR libgomp/69414
+	* oacc-mem.c (delete_copyout, update_dev_host): Fix device address.
+	* testsuite/libgomp.oacc-c-c++-common/update-1.c: Additional tests.
+	* testsuite/libgomp.oacc-c-c++-common/update-1-2.c: Likewise.
+	* testsuite/libgomp.oacc-fortran/update-1.f90: New file.
+
 2016-02-16  Tom de Vries  <tom@codesourcery.com>
 
 	backport from trunk:
diff --git libgomp/oacc-mem.c libgomp/oacc-mem.c
index 89ef5fc..c3e12fa 100644
--- libgomp/oacc-mem.c
+++ libgomp/oacc-mem.c
@@ -447,7 +447,8 @@ delete_copyout (unsigned f, void *h, size_t s)
   if (!n)
     gomp_fatal ("[%p,%d] is not mapped", (void *)h, (int)s);
 
-  d = (void *) (n->tgt->tgt_start + n->tgt_offset);
+  d = (void *) (n->tgt->tgt_start + n->tgt_offset
+		+ (uintptr_t) h - n->host_start);
 
   host_size = n->host_end - n->host_start;
 
@@ -490,7 +491,8 @@ update_dev_host (int is_dev, void *h, size_t s)
   if (!n)
     gomp_fatal ("[%p,%d] is not mapped", h, (int)s);
 
-  d = (void *) (n->tgt->tgt_start + n->tgt_offset);
+  d = (void *) (n->tgt->tgt_start + n->tgt_offset
+		+ (uintptr_t) h - n->host_start);
 
   if (is_dev)
     acc_dev->host2dev_func (acc_dev->target_id, d, h, s);
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c
index c7e7257..82c3192 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c
@@ -13,6 +13,7 @@ int
 main (int argc, char **argv)
 {
     int N = 8;
+    int NDIV2 = N / 2;
     float *a, *b, *c;
     float *d_a, *d_b, *d_c;
     int i;
@@ -242,7 +243,7 @@ main (int argc, char **argv)
         a[i] = 6.0;
     }
 
-#pragma acc update device (a[0:N >> 1])
+#pragma acc update device (a[0:NDIV2])
 
 #pragma acc parallel present (a[0:N], b[0:N])
     {
@@ -254,7 +255,7 @@ main (int argc, char **argv)
 
 #pragma acc update self (a[0:N], b[0:N])
 
-    for (i = 0; i < (N >> 1); i++)
+    for (i = 0; i < NDIV2; i++)
     {
         if (a[i] != 6.0)
             abort ();
@@ -263,7 +264,7 @@ main (int argc, char **argv)
             abort ();
     }
 
-    for (i = (N >> 1); i < N; i++)
+    for (i = NDIV2; i < N; i++)
     {
         if (a[i] != 5.0)
             abort ();
@@ -278,5 +279,83 @@ main (int argc, char **argv)
     if (!acc_is_present (&b[0], (N * sizeof (float))))
       abort ();
 
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 0.0;
+    }
+
+#pragma acc update device (a[0:4])
+
+#pragma acc parallel present (a[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            a[ii] = a[ii] + 1.0;
+    }
+
+#pragma acc update self (a[4:4])
+
+    for (i = 0; i < NDIV2; i++)
+    {
+        if (a[i] != 0.0)
+            abort ();
+    }
+
+    for (i = NDIV2; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
+#pragma acc update self (a[0:4])
+
+    for (i = 0; i < NDIV2; i++)
+    {
+        if (a[i] != 1.0)
+            abort ();
+    }
+
+    for (i = NDIV2; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
+    a[2] = 9;
+    a[3] = 9;
+    a[4] = 9;
+    a[5] = 9;
+
+#pragma acc update device (a[2:4])
+
+#pragma acc parallel present (a[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            a[ii] = a[ii] + 1.0;
+    }
+
+#pragma acc update self (a[2:4])
+
+    for (i = 0; i < 2; i++)
+    {
+      if (a[i] != 1.0)
+	abort ();
+    }
+
+    for (i = 2; i < 6; i++)
+    {
+      if (a[i] != 10.0)
+	abort ();
+    }
+
+    for (i = 6; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
     return 0;
 }
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/update-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/update-1.c
index dff139f..1b2a460 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/update-1.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/update-1.c
@@ -11,6 +11,7 @@ int
 main (int argc, char **argv)
 {
     int N = 8;
+    int NDIV2 = N / 2;
     float *a, *b, *c;
     float *d_a, *d_b, *d_c;
     int i;
@@ -109,7 +110,7 @@ main (int argc, char **argv)
             b[ii] = a[ii];
     }
 
-#pragma acc update self (a[0:N], b[0:N])
+#pragma acc update host (a[0:N], b[0:N])
 
     for (i = 0; i < N; i++)
     {
@@ -240,7 +241,7 @@ main (int argc, char **argv)
         a[i] = 6.0;
     }
 
-#pragma acc update device (a[0:N >> 1])
+#pragma acc update device (a[0:NDIV2])
 
 #pragma acc parallel present (a[0:N], b[0:N])
     {
@@ -252,7 +253,7 @@ main (int argc, char **argv)
 
 #pragma acc update host (a[0:N], b[0:N])
 
-    for (i = 0; i < (N >> 1); i++)
+    for (i = 0; i < NDIV2; i++)
     {
         if (a[i] != 6.0)
             abort ();
@@ -261,7 +262,7 @@ main (int argc, char **argv)
             abort ();
     }
 
-    for (i = (N >> 1); i < N; i++)
+    for (i = NDIV2; i < N; i++)
     {
         if (a[i] != 5.0)
             abort ();
@@ -276,5 +277,83 @@ main (int argc, char **argv)
     if (!acc_is_present (&b[0], (N * sizeof (float))))
       abort ();
 
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 0.0;
+    }
+
+#pragma acc update device (a[0:4])
+
+#pragma acc parallel present (a[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            a[ii] = a[ii] + 1.0;
+    }
+
+#pragma acc update host (a[4:4])
+
+    for (i = 0; i < NDIV2; i++)
+    {
+        if (a[i] != 0.0)
+            abort ();
+    }
+
+    for (i = NDIV2; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
+#pragma acc update host (a[0:4])
+
+    for (i = 0; i < NDIV2; i++)
+    {
+        if (a[i] != 1.0)
+            abort ();
+    }
+
+    for (i = NDIV2; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
+    a[2] = 9;
+    a[3] = 9;
+    a[4] = 9;
+    a[5] = 9;
+
+#pragma acc update device (a[2:4])
+
+#pragma acc parallel present (a[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+            a[ii] = a[ii] + 1.0;
+    }
+
+#pragma acc update host (a[2:4])
+
+    for (i = 0; i < 2; i++)
+    {
+      if (a[i] != 1.0)
+	abort ();
+    }
+
+    for (i = 2; i < 6; i++)
+    {
+      if (a[i] != 10.0)
+	abort ();
+    }
+
+    for (i = 6; i < N; i++)
+    {
+        if (a[i] != 6.0)
+            abort ();
+    }
+
     return 0;
 }
diff --git libgomp/testsuite/libgomp.oacc-fortran/update-1.f90 libgomp/testsuite/libgomp.oacc-fortran/update-1.f90
new file mode 100644
index 0000000..4e1d2c0
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/update-1.f90
@@ -0,0 +1,242 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+program update
+  use openacc
+  implicit none 
+  integer, parameter :: N = 8
+  integer, parameter :: NDIV2 = N / 2
+  real :: a(N), b(N)
+  integer i
+
+  do i = 1, N
+    a(i) = 3.0
+    b(i) = 0.0
+  end do
+
+  !$acc enter data copyin (a, b)
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 3.0) call abort
+    if (b(i) .ne. 3.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 5.0
+    b(i) = 1.0
+  end do
+
+  !$acc update device (a, b)
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do 
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 5.0) call abort
+    if (b(i) .ne. 5.0) call abort
+ end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  !$acc parallel present (a, b)
+  do i = 1, N
+    b(i) = a(i)
+  end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 5.0) call abort
+    if (b(i) .ne. 5.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 6.0
+    b(i) = 0.0
+  end do
+
+  !$acc update device (a, b)
+
+  do i = 1, N
+    a(i) = 9.0
+  end do
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 6.0) call abort
+    if (b(i) .ne. 6.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 7.0
+    b(i) = 2.0
+  end do
+
+  !$acc update device (a, b)
+
+  do i = 1, N
+    a(i) = 9.0
+  end do
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 7.0) call abort
+    if (b(i) .ne. 7.0) call abort
+  end do
+
+  do i = 1, N
+    a(i) = 9.0
+  end do
+
+  !$acc update device (a)
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, N
+    if (a(i) .ne. 9.0) call abort
+    if (b(i) .ne. 9.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 5.0
+  end do
+
+  !$acc update device (a)
+
+  do i = 1, N
+    a(i) = 6.0
+  end do
+
+  !$acc update device (a(1:NDIV2))
+
+  !$acc parallel present (a, b)
+    do i = 1, N
+      b(i) = a(i)
+    end do
+  !$acc end parallel
+
+  !$acc update host (a, b)
+
+  do i = 1, NDIV2
+    if (a(i) .ne. 6.0) call abort
+    if (b(i) .ne. 6.0) call abort
+  end do
+
+  do i = NDIV2 + 1, N
+    if (a(i) .ne. 5.0) call abort
+    if (b(i) .ne. 5.0) call abort
+  end do
+
+  if (acc_is_present (a) .neqv. .TRUE.) call abort
+  if (acc_is_present (b) .neqv. .TRUE.) call abort
+
+  do i = 1, N
+    a(i) = 0.0
+  end do
+
+  !$acc update device (a(1:4))
+
+  !$acc parallel present (a)
+    do i = 1, N
+      a(i) = a(i) + 1.0
+    end do
+  !$acc end parallel
+
+  !$acc update host (a(5:N))
+
+  do i = 1, NDIV2
+    if (a(i) .ne. 0.0) call abort
+  end do
+
+  do i = NDIV2 + 1, N
+    if (a(i) .ne. 6.0) call abort
+  end do
+
+  !$acc update host (a(1:4))
+
+  do i = 1, NDIV2
+    if (a(i) .ne. 1.0) call abort
+  end do
+
+  do i = NDIV2 + 1, N
+    if (a(i) .ne. 6.0) call abort
+  end do
+
+  a(3) = 9
+  a(4) = 9
+  a(5) = 9
+  a(6) = 9
+
+  !$acc update device (a(3:6))
+
+  !$acc parallel present (a(1:N))
+    do i = 1, N
+      a(i) = a(i) + 1.0
+    end do
+  !$acc end parallel
+
+  !$acc update host (a(3:6))
+
+  do i = 1, 2
+    if (a(i) .ne. 1.0) call abort
+  end do
+
+  do i = 3, 6
+    if (a(i) .ne. 10.0) call abort
+  end do
+
+  do i = 7, N
+    if (a(i) .ne. 6.0) call abort
+  end do
+
+  !$acc exit data delete (a, b)
+
+end program
+


Grüße
 Thomas