[gomp4] Fix handling of subarrays with update directive

[James Norris <jnorris@codesourcery.com>]

[-- 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;
 }