git
/
gcc
mirror of git://gcc.gnu.org/git/gcc.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

re PR libgomp/69414 ([OpenACC] "!$acc update self" does not provide expected result) 

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.

Co-Authored-By: Daichi Fukuoka <dc-fukuoka@sgi.com>

From-SVN: r234428
This commit is contained in:
James Norris 2016-03-23 14:38:55 +00:00 committed by James Norris
parent 4d8989d5b0
commit b6d1f2b546
5 changed files with 420 additions and 9 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
libgomp/ChangeLog
View File

@ -1,3 +1,12 @@
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-03-23 Martin Liska <mliska@suse.cz>
PR hsa/70337

6
libgomp/oacc-mem.c
View File

@ -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);

85
libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c
View File

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

87
libgomp/testsuite/libgomp.oacc-c-c++-common/update-1.c
View File

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

242
libgomp/testsuite/libgomp.oacc-fortran/update-1.f90 Normal file
View File

@ -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