Index: ChangeLog.gomp =================================================================== --- ChangeLog.gomp (revision 232740) +++ ChangeLog.gomp (revision 232741) @@ -1,3 +1,11 @@ +2016-01-22 James Norris + + * 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 * 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; }