diff options
author | Julian Brown <julian@codesourcery.com> | 2019-09-20 13:53:10 -0700 |
---|---|---|
committer | Thomas Schwinge <thomas@codesourcery.com> | 2022-01-17 08:57:20 +0100 |
commit | fbb438808e9b53a6e6b179a5787d609443acaad6 (patch) | |
tree | b051bbb82e3de22737ca797ecb4fa4223d30f8b4 | |
parent | 025a707d285832cec91c28811db02b74cd08a59a (diff) | |
download | gcc-fbb438808e9b53a6e6b179a5787d609443acaad6.zip gcc-fbb438808e9b53a6e6b179a5787d609443acaad6.tar.gz gcc-fbb438808e9b53a6e6b179a5787d609443acaad6.tar.bz2 |
Test cases for references in OpenACC 'private' clauses
libgomp/
* testsuite/libgomp.oacc-fortran/privatized-ref-1.f95: New test.
* testsuite/libgomp.oacc-c++/privatized-ref-2.C: New test.
* testsuite/libgomp.oacc-c++/privatized-ref-3.C: New test.
Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
3 files changed, 211 insertions, 0 deletions
diff --git a/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C new file mode 100644 index 0000000..7091091 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C @@ -0,0 +1,66 @@ +/* { dg-do run } */ + +#include <stdlib.h> + +void workers (void) +{ + double res[65536]; + int i; + +#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) + /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */ + { + int i, j; +#pragma acc loop gang + for (i = 0; i < 256; i++) + { +#pragma acc loop worker + for (j = 0; j < 256; j++) + { + int tmpvar; + int &tmpref = tmpvar; + tmpref = (i * 256 + j) * 99; + res[i * 256 + j] = tmpref; + } + } + } + + for (i = 0; i < 65536; i++) + if (res[i] != i * 99) + abort (); +} + +void vectors (void) +{ + double res[65536]; + int i; + +#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) + /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */ + { + int i, j; +#pragma acc loop gang worker + for (i = 0; i < 256; i++) + { +#pragma acc loop vector + for (j = 0; j < 256; j++) + { + int tmpvar; + int &tmpref = tmpvar; + tmpref = (i * 256 + j) * 101; + res[i * 256 + j] = tmpref; + } + } + } + + for (i = 0; i < 65536; i++) + if (res[i] != i * 101) + abort (); +} + +int main (int argc, char *argv[]) +{ + workers (); + vectors (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C new file mode 100644 index 0000000..478876e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C @@ -0,0 +1,69 @@ +/* { dg-do run } */ + +/*TODO + { dg-xfail-run-if TODO { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } } */ + +#include <stdlib.h> + +void workers (void) +{ + double res[65536]; + int i; + +#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) + /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */ + { + int i, j; + int tmpvar; + int &tmpref = tmpvar; +#pragma acc loop gang + for (i = 0; i < 256; i++) + { +#pragma acc loop worker private(tmpref) + for (j = 0; j < 256; j++) + { + tmpref = (i * 256 + j) * 99; + res[i * 256 + j] = tmpref; + } + } + } + + for (i = 0; i < 65536; i++) + if (res[i] != i * 99) + abort (); +} + +void vectors (void) +{ + double res[65536]; + int i; + +#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) + /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */ + { + int i, j; + int tmpvar; + int &tmpref = tmpvar; +#pragma acc loop gang worker + for (i = 0; i < 256; i++) + { +#pragma acc loop vector private(tmpref) + for (j = 0; j < 256; j++) + { + tmpref = (i * 256 + j) * 101; + res[i * 256 + j] = tmpref; + } + } + } + + for (i = 0; i < 65536; i++) + if (res[i] != i * 101) + abort (); +} + +int main (int argc, char *argv[]) +{ + workers (); + vectors (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 new file mode 100644 index 0000000..bb0910b --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 @@ -0,0 +1,76 @@ +! { dg-do run } + +!TODO +! { dg-xfail-run-if TODO { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } } + +program main + implicit none + integer :: myint + integer :: i + real :: res(65536), tmp + + res(:) = 0.0 + + myint = 5 + call workers(myint, res) + + do i=1,65536 + tmp = i * 99 + if (res(i) .ne. tmp) stop 1 + end do + + res(:) = 0.0 + + myint = 7 + call vectors(myint, res) + + do i=1,65536 + tmp = i * 101 + if (res(i) .ne. tmp) stop 2 + end do + +contains + + subroutine workers(t1, res) + implicit none + integer :: t1 + integer :: i, j + real, intent(out) :: res(:) + + !$acc parallel copyout(res) num_gangs(64) num_workers(64) + ! { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } + + !$acc loop gang + do i=0,255 + !$acc loop worker private(t1) + do j=1,256 + t1 = (i * 256 + j) * 99 + res(i * 256 + j) = t1 + end do + end do + + !$acc end parallel + end subroutine workers + + subroutine vectors(t1, res) + implicit none + integer :: t1 + integer :: i, j + real, intent(out) :: res(:) + + !$acc parallel copyout(res) num_gangs(64) num_workers(64) + ! { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } + + !$acc loop gang worker + do i=0,255 + !$acc loop vector private(t1) + do j=1,256 + t1 = (i * 256 + j) * 101 + res(i * 256 + j) = t1 + end do + end do + + !$acc end parallel + end subroutine vectors + +end program main |