* [gomp4] Fix handling of subarrays with update directive
@ 2016-01-22 18:50 ` James Norris
2016-03-23 8:43 ` Thomas Schwinge
0 siblings, 1 reply; 11+ messages in thread
From: James Norris @ 2016-01-22 18:50 UTC (permalink / raw)
To: GCC Patches, Thomas Schwinge
[-- 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;
}
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gomp4] Fix handling of subarrays with update directive
2016-01-22 18:50 ` [gomp4] Fix handling of subarrays with update directive James Norris
@ 2016-03-23 8:43 ` Thomas Schwinge
2016-03-23 10:33 ` Jakub Jelinek
0 siblings, 1 reply; 11+ messages in thread
From: Thomas Schwinge @ 2016-03-23 8:43 UTC (permalink / raw)
To: James Norris, GCC Patches, Jakub Jelinek; +Cc: Daichi Fukuoka
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
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gomp4] Fix handling of subarrays with update directive
2016-03-23 8:43 ` Thomas Schwinge
@ 2016-03-23 10:33 ` Jakub Jelinek
2016-03-23 10:37 ` Fukuoka Daichi
2016-03-23 13:15 ` James Norris
0 siblings, 2 replies; 11+ messages in thread
From: Jakub Jelinek @ 2016-03-23 10:33 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: James Norris, GCC Patches, Daichi Fukuoka
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
^ permalink raw reply [flat|nested] 11+ messages in thread
* RE: [gomp4] Fix handling of subarrays with update directive
2016-03-23 10:33 ` Jakub Jelinek
@ 2016-03-23 10:37 ` Fukuoka Daichi
2016-03-23 11:10 ` Jakub Jelinek
2016-03-23 13:15 ` James Norris
1 sibling, 1 reply; 11+ messages in thread
From: Fukuoka Daichi @ 2016-03-23 10:37 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: James Norris, Thomas Schwinge, GCC Patches
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
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gomp4] Fix handling of subarrays with update directive
2016-03-23 10:37 ` Fukuoka Daichi
@ 2016-03-23 11:10 ` Jakub Jelinek
2016-03-23 11:31 ` Fukuoka Daichi
0 siblings, 1 reply; 11+ messages in thread
From: Jakub Jelinek @ 2016-03-23 11:10 UTC (permalink / raw)
To: Fukuoka Daichi; +Cc: James Norris, Thomas Schwinge, GCC Patches
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
^ permalink raw reply [flat|nested] 11+ messages in thread
* RE: [gomp4] Fix handling of subarrays with update directive
2016-03-23 11:10 ` Jakub Jelinek
@ 2016-03-23 11:31 ` Fukuoka Daichi
0 siblings, 0 replies; 11+ messages in thread
From: Fukuoka Daichi @ 2016-03-23 11:31 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: James Norris, Thomas Schwinge, GCC Patches
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
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gomp4] Fix handling of subarrays with update directive
2016-03-23 10:33 ` Jakub Jelinek
2016-03-23 10:37 ` Fukuoka Daichi
@ 2016-03-23 13:15 ` James Norris
2016-03-23 13:18 ` Jakub Jelinek
1 sibling, 1 reply; 11+ messages in thread
From: James Norris @ 2016-03-23 13:15 UTC (permalink / raw)
To: Jakub Jelinek, Thomas Schwinge; +Cc: GCC Patches, Daichi Fukuoka
[-- 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
+
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gomp4] Fix handling of subarrays with update directive
2016-03-23 13:15 ` James Norris
@ 2016-03-23 13:18 ` Jakub Jelinek
2016-04-07 11:29 ` [gcc-5] " Thomas Schwinge
0 siblings, 1 reply; 11+ messages in thread
From: Jakub Jelinek @ 2016-03-23 13:18 UTC (permalink / raw)
To: James Norris; +Cc: Thomas Schwinge, GCC Patches, Daichi Fukuoka
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
^ permalink raw reply [flat|nested] 11+ messages in thread
* [gcc-5] Fix handling of subarrays with update directive
2016-03-23 13:18 ` Jakub Jelinek
@ 2016-04-07 11:29 ` Thomas Schwinge
2016-04-07 11:30 ` Jakub Jelinek
0 siblings, 1 reply; 11+ messages in thread
From: Thomas Schwinge @ 2016-04-07 11:29 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek; +Cc: Daichi Fukuoka, James Norris
[-- 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 --]
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gcc-5] Fix handling of subarrays with update directive
2016-04-07 11:29 ` [gcc-5] " Thomas Schwinge
@ 2016-04-07 11:30 ` Jakub Jelinek
2016-04-07 11:48 ` Thomas Schwinge
0 siblings, 1 reply; 11+ messages in thread
From: Jakub Jelinek @ 2016-04-07 11:30 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches, Daichi Fukuoka, James Norris
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
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gcc-5] Fix handling of subarrays with update directive
2016-04-07 11:30 ` Jakub Jelinek
@ 2016-04-07 11:48 ` Thomas Schwinge
0 siblings, 0 replies; 11+ messages in thread
From: Thomas Schwinge @ 2016-04-07 11:48 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek; +Cc: Daichi Fukuoka, James Norris
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
^ permalink raw reply [flat|nested] 11+ messages in thread
end of thread, other threads:[~2016-04-07 11:48 UTC | newest]
Thread overview: 11+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
[not found] <56A0E84A.5030107@mentor.com>
[not found] ` <56F13A4E.20907@mentor.com>
2016-01-22 18:50 ` [gomp4] Fix handling of subarrays with update directive James Norris
2016-03-23 8:43 ` Thomas Schwinge
2016-03-23 10:33 ` Jakub Jelinek
2016-03-23 10:37 ` Fukuoka Daichi
2016-03-23 11:10 ` Jakub Jelinek
2016-03-23 11:31 ` Fukuoka Daichi
2016-03-23 13:15 ` James Norris
2016-03-23 13:18 ` Jakub Jelinek
2016-04-07 11:29 ` [gcc-5] " Thomas Schwinge
2016-04-07 11:30 ` Jakub Jelinek
2016-04-07 11:48 ` Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).