This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [gomp4] Fix handling of subarrays with update directive
- From: James Norris <jnorris at codesourcery dot com>
- To: Jakub Jelinek <jakub at redhat dot com>, Thomas Schwinge <thomas at codesourcery dot com>
- Cc: GCC Patches <gcc-patches at gcc dot gnu dot org>, Daichi Fukuoka <dc-fukuoka at sgi dot com>
- Date: Wed, 23 Mar 2016 08:05:19 -0500
- Subject: Re: [gomp4] Fix handling of subarrays with update directive
- Authentication-results: sourceware.org; auth=none
- References: <56A0E84A dot 5030107 at mentor dot com> <56F13A4E dot 20907 at mentor dot com> <56A279FE dot 3020600 at codesourcery dot com> <87io0dbnhp dot fsf at hertz dot schwinge dot homeip dot net> <20160323102421 dot GU3017 at tucnak dot redhat dot com>
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.
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
+