From 9e10bfb7c761a1343ecf51cfff60e03839561962 Mon Sep 17 00:00:00 2001 From: tschwinge Date: Tue, 3 Nov 2015 11:28:22 +0000 Subject: OpenACC atomic directive gcc/c-family/ * c-pragma.c (oacc_pragmas): Add "atomic". * c-pragma.h (pragma_kind): Add PRAGMA_OACC_ATOMIC. gcc/c/ * c-parser.c (c_parser_omp_construct): Handle PRAGMA_OACC_ATOMIC. gcc/cp/ * parser.c (cp_parser_omp_construct, cp_parser_pragma): Handle PRAGMA_OACC_ATOMIC. gcc/fortran/ * gfortran.h (gfc_statement): Add ST_OACC_ATOMIC, ST_OACC_END_ATOMIC. (gfc_exec_op): Add EXEC_OACC_ATOMIC. * match.h (gfc_match_oacc_atomic): New prototype. * openmp.c (gfc_match_omp_atomic, gfc_match_oacc_atomic): New wrapper functions around... (gfc_match_omp_oacc_atomic): ... this new function. (oacc_code_to_statement, gfc_resolve_oacc_directive): Handle EXEC_OACC_ATOMIC. * parse.c (decode_oacc_directive): Handle "atomic", "end atomic". (case_exec_markers): Add ST_OACC_ATOMIC. (gfc_ascii_statement): Handle ST_OACC_ATOMIC, ST_OACC_END_ATOMIC. (parse_omp_atomic): Rename to... (parse_omp_oacc_atomic): ... this new function. Add omp_p formal parameter. Adjust all users. (parse_executable): Handle ST_OACC_ATOMIC. (is_oacc): Handle EXEC_OACC_ATOMIC. * resolve.c (gfc_resolve_blocks, gfc_resolve_code): Handle EXEC_OACC_ATOMIC. * st.c (gfc_free_statement): Handle EXEC_OACC_ATOMIC. * trans-openmp.c (gfc_trans_oacc_directive): Handle EXEC_OACC_ATOMIC. * trans.c (trans_code): Handle EXEC_OACC_ATOMIC. gcc/ * builtins.def (DEF_GOMP_BUILTIN): Enable for flag_openacc. * omp-low.c (check_omp_nesting_restrictions): Allow GIMPLE_OMP_ATOMIC_LOAD, GIMPLE_OMP_ATOMIC_STORE inside OpenACC contexts. gcc/testsuite/ * c-c++-common/goacc-gomp/nesting-fail-1.c: Move "atomic" tests from here to... * c-c++-common/goacc-gomp/nesting-1.c: ... here, and expect them to succeed. libgomp/ * testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c: New file. * testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/atomic_rw-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/atomic_update-1.c: Likewise. * testsuite/libgomp.oacc-fortran/atomic_capture-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/atomic_rw-1.f90: New file. * testsuite/libgomp.oacc-fortran/atomic_update-1.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c: New file. * testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/worker-single-1a.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/worker-single-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/worker-single-6.c: Likewise. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@229703 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog | 22 + .../libgomp.oacc-c-c++-common/atomic_capture-1.c | 866 +++++++++++ .../libgomp.oacc-c-c++-common/atomic_capture-2.c | 1626 ++++++++++++++++++++ .../libgomp.oacc-c-c++-common/atomic_rw-1.c | 34 + .../libgomp.oacc-c-c++-common/atomic_update-1.c | 760 +++++++++ .../libgomp.oacc-c-c++-common/par-reduction-1.c | 44 + .../libgomp.oacc-c-c++-common/par-reduction-2.c | 48 + .../libgomp.oacc-c-c++-common/worker-single-1a.c | 28 + .../libgomp.oacc-c-c++-common/worker-single-4.c | 28 + .../libgomp.oacc-c-c++-common/worker-single-6.c | 46 + .../libgomp.oacc-fortran/atomic_capture-1.f90 | 784 ++++++++++ .../testsuite/libgomp.oacc-fortran/atomic_rw-1.f90 | 29 + .../libgomp.oacc-fortran/atomic_update-1.f90 | 338 ++++ 13 files changed, 4653 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_rw-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_update-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-1a.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-4.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-6.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/atomic_capture-1.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/atomic_rw-1.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/atomic_update-1.f90 (limited to 'libgomp') diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 4d1906613f9..7894a7e2c16 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,25 @@ +2015-11-03 Julian Brown + Thomas Schwinge + + * testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c: New file. + * testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/worker-single-1a.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/worker-single-4.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/worker-single-6.c: Likewise. + +2015-11-03 James Norris + + * testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c: New + file. + * testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/atomic_rw-1.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/atomic_update-1.c: Likewise. + * testsuite/libgomp.oacc-fortran/atomic_capture-1.f90: Likewise. + * testsuite/libgomp.oacc-fortran/atomic_rw-1.f90: New file. + * testsuite/libgomp.oacc-fortran/atomic_update-1.f90: Likewise. + 2015-10-29 Nathan Sidwell * openacc.h (enum acc_device_t): Reformat. Ensure layout diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c new file mode 100644 index 00000000000..ad958cd1f9e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c @@ -0,0 +1,866 @@ +/* { dg-do run } */ + +#include + +int +main(int argc, char **argv) +{ + int iexp, igot; + long long lexp, lgot; + int N = 32; + int idata[N]; + long long ldata[N]; + float fexp, fgot; + float fdata[N]; + int i; + + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { +#pragma acc atomic capture + idata[i] = igot++; + } + } + + if (iexp != igot) + abort (); + + igot = 32; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { +#pragma acc atomic capture + idata[i] = igot--; + } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { +#pragma acc atomic capture + idata[i] = ++igot; + } + } + + if (iexp != igot) + abort (); + + igot = 32; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { +#pragma acc atomic capture + idata[i] = --igot; + } + } + + /* BINOP = + */ + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + idata[i] = igot += expr; + } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + idata[i] = igot = igot + expr; + } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + idata[i] = igot = expr + igot; + } + } + + if (iexp != igot) + abort (); + + /* BINOP = * */ + lgot = 1LL; + lexp = 1LL << N; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + ldata[i] = lgot *= expr; + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 1LL << N; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + ldata[i] = lgot = lgot * expr; + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 1LL << N; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + ldata[i] = lgot = expr * lgot; + } + } + + if (lexp != lgot) + abort (); + + /* BINOP = - */ + igot = 32; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + idata[i] = igot -= expr; + } + } + + if (iexp != igot) + abort (); + + igot = 32; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + idata[i] = igot = igot - expr; + } + } + + if (iexp != igot) + abort (); + + igot = 32; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + idata[i] = igot = expr - igot; + } + } + + if (iexp != igot) + abort (); + + + /* BINOP = / */ + lgot = 1LL << 32; + lexp = 1LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + ldata[i] = lgot /= expr; + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL << 32; + lexp = 1LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + ldata[i] = lgot = lgot / expr; + } + } + + if (lexp != lgot) + abort (); + + lgot = 2LL; + lexp = 2LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL << N; + +#pragma acc atomic capture + ldata[i] = lgot = expr / lgot; + } + } + + if (lexp != lgot) + abort (); + + /* BINOP = & */ + igot = ~0; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1 << i; + +#pragma acc atomic capture + idata[i] = igot &= expr; + } + } + + if (iexp != igot) + abort (); + + igot = ~0; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1 << i; + +#pragma acc atomic capture + idata[i] = igot = igot & expr; + } + } + + if (iexp != igot) + abort (); + + igot = ~0; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1 << i; + +#pragma acc atomic capture + idata[i] = igot = expr & igot; + } + } + + if (iexp != igot) + abort (); + + /* BINOP = ^ */ + igot = ~0; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1 << i; + +#pragma acc atomic capture + idata[i] = igot ^= expr; + } + } + + if (iexp != igot) + abort (); + + igot = ~0; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1 << i; + +#pragma acc atomic capture + idata[i] = igot = igot ^ expr; + } + } + + if (iexp != igot) + abort (); + + igot = ~0; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1 << i; + +#pragma acc atomic capture + idata[i] = igot = expr ^ igot; + } + } + + if (iexp != igot) + abort (); + + /* BINOP = | */ + igot = 0; + iexp = ~0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1 << i; + +#pragma acc atomic capture + idata[i] = igot |= expr; + } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = ~0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1 << i; + +#pragma acc atomic capture + idata[i] = igot = igot | expr; + } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = ~0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1 << i; + +#pragma acc atomic capture + idata[i] = igot = expr | igot; + } + } + + if (iexp != igot) + abort (); + + /* BINOP = << */ + lgot = 1LL; + lexp = 1LL << N; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + ldata[i] = lgot <<= expr; + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 1LL << N; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + idata[i] = lgot = lgot << expr; + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 2LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel + { + long long expr = 1LL; + +#pragma acc atomic capture + ldata[0] = lgot = expr << lgot; + } + } + + if (lexp != lgot) + abort (); + + /* BINOP = >> */ + lgot = 1LL << N; + lexp = 1LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL; + +#pragma acc atomic capture + ldata[i] = lgot >>= expr; + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL << N; + lexp = 1LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + ldata[i] = lgot = lgot >> expr; + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL << 63; + lexp = 1LL << 32; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel + { + long long expr = 1LL << 32; + +#pragma acc atomic capture + ldata[0] = lgot = expr >> lgot; + } + } + + if (lexp != lgot) + abort (); + + fgot = 0.0; + fexp = 32.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { +#pragma acc atomic capture + fdata[i] = fgot++; + } + } + + if (fexp != fgot) + abort (); + + fgot = 32.0; + fexp = 0.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { +#pragma acc atomic capture + fdata[i] = fgot--; + } + } + + if (fexp != fgot) + abort (); + + fgot = 0.0; + fexp = 32.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { +#pragma acc atomic capture + fdata[i] = ++fgot; + } + } + + if (fexp != fgot) + abort (); + + fgot = 32.0; + fexp = 0.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { +#pragma acc atomic capture + fdata[i] = --fgot; + } + } + + if (fexp != fgot) + abort (); + + /* BINOP = + */ + fgot = 0.0; + fexp = 32.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + fdata[i] = fgot += expr; + } + } + + if (fexp != fgot) + abort (); + + fgot = 0.0; + fexp = 32.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + fdata[i] = fgot = fgot + expr; + } + } + + if (fexp != fgot) + abort (); + + fgot = 0.0; + fexp = 32.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + fdata[i] = fgot = expr + fgot; + } + } + + if (fexp != fgot) + abort (); + + /* BINOP = * */ + fgot = 1.0; + fexp = 8192.0*8192.0*64.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic capture + fdata[i] = fgot *= expr; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1.0; + fexp = 8192.0*8192.0*64.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + fdata[i] = fgot = fgot * expr; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1.0; + fexp = 8192.0*8192.0*64.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic capture + fdata[i] = fgot = expr * fgot; + } + } + + if (fexp != fgot) + abort (); + + /* BINOP = - */ + fgot = 32.0; + fexp = 0.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + fdata[i] = fgot -= expr; + } + } + + if (fexp != fgot) + abort (); + + fgot = 32.0; + fexp = 0.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + fdata[i] = fgot = fgot - expr; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1.0; + fexp = 0.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 32.0; + +#pragma acc atomic capture + fdata[i] = fgot = expr - fgot; + } + } + + for (i = 0; i < N; i++) + if (i % 2 == 0) + { + if (fdata[i] != 31.0) + abort (); + } + else + { + if (fdata[i] != 1.0) + abort (); + } + + + /* BINOP = / */ + fexp = 1.0; + fgot = 8192.0*8192.0*64.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic capture + fdata[i] = fgot /= expr; + } + } + + if (fexp != fgot) + abort (); + + fexp = 1.0; + fgot = 8192.0*8192.0*64.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic capture + fdata[i] = fgot = fgot / expr; + } + } + + if (fexp != fgot) + abort (); + + fexp = 1.0; + fgot = 8192.0*8192.0*64.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel + { + float expr = 8192.0*8192.0*64.0; + +#pragma acc atomic capture + fdata[0] = fgot = expr / fgot; + } + } + + if (fexp != fgot) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c new file mode 100644 index 00000000000..842f2de4722 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c @@ -0,0 +1,1626 @@ +/* { dg-do run } */ + +#include + +int +main(int argc, char **argv) +{ + int iexp, igot, imax, imin; + long long lexp, lgot; + int N = 32; + int i; + int idata[N]; + long long ldata[N]; + float fexp, fgot; + float fdata[N]; + + igot = 1234; + iexp = 31; + + for (i = 0; i < N; i++) + idata[i] = i; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) +#pragma acc atomic capture + { idata[i] = igot; igot = i; } + } + + imax = 0; + imin = N; + + for (i = 0; i < N; i++) + { + imax = idata[i] > imax ? idata[i] : imax; + imin = idata[i] < imin ? idata[i] : imin; + } + + if (imax != 1234 || imin != 0) + abort (); + + return 0; + + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) +#pragma acc atomic capture + { idata[i] = igot; igot++; } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) +#pragma acc atomic capture + { idata[i] = igot; ++igot; } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) +#pragma acc atomic capture + { ++igot; idata[i] = igot; } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) +#pragma acc atomic capture + { igot++; idata[i] = igot; } + } + + if (iexp != igot) + abort (); + + igot = 32; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) +#pragma acc atomic capture + { idata[i] = igot; igot--; } + } + + if (iexp != igot) + abort (); + + igot = 32; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) +#pragma acc atomic capture + { idata[i] = igot; --igot; } + } + + if (iexp != igot) + abort (); + + igot = 32; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) +#pragma acc atomic capture + { --igot; idata[i] = igot; } + } + + if (iexp != igot) + abort (); + + igot = 32; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) +#pragma acc atomic capture + { igot--; idata[i] = igot; } + } + + if (iexp != igot) + abort (); + + /* BINOP = + */ + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + { idata[i] = igot; igot += expr; } + } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + { igot += expr; idata[i] = igot; } + } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + { idata[i] = igot; igot = igot + expr; } + } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + { idata[i] = igot; igot = expr + igot; } + } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + { igot = igot + expr; idata[i] = igot; } + } + } + + if (iexp != igot) + abort (); + + + igot = 0; + iexp = 32; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + { igot = expr + igot; idata[i] = igot; } + } + } + + if (iexp != igot) + abort (); + + /* BINOP = * */ + lgot = 1LL; + lexp = 1LL << 32; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + { ldata[i] = lgot; lgot *= expr; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 1LL << 32; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + { lgot *= expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 1LL << 32; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + { ldata[i] = lgot; lgot = lgot * expr; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 1LL << 32; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + { ldata[i] = lgot; lgot = expr * lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 1LL << 32; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + { lgot = lgot * expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 1LL << 32; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2; + +#pragma acc atomic capture + { lgot = expr * lgot; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + /* BINOP = - */ + igot = 32; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + { idata[i] = igot; igot -= expr; } + } + } + + if (iexp != igot) + abort (); + + igot = 32; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + { igot -= expr; idata[i] = igot; } + } + } + + if (iexp != igot) + abort (); + + igot = 32; + iexp = 0; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + { idata[i] = igot; igot = igot - expr; } + } + } + + if (iexp != igot) + abort (); + + igot = 1; + iexp = 1; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + { idata[i] = igot; igot = expr - igot; } + } + } + + for (i = 0; i < N; i++) + if (i % 2 == 0) + { + if (idata[i] != 1) + abort (); + } + else + { + if (idata[i] != 0) + abort (); + } + + if (iexp != igot) + abort (); + + igot = 1; + iexp = -31; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + { igot = igot - expr; idata[i] = igot; } + } + } + + if (iexp != igot) + abort (); + + igot = 1; + iexp = 1; + +#pragma acc data copy (igot, idata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = 1; + +#pragma acc atomic capture + { igot = expr - igot; idata[i] = igot; } + } + } + + for (i = 0; i < N; i++) + if (i % 2 == 0) + { + if (idata[i] != 0) + abort (); + } + else + { + if (idata[i] != 1) + abort (); + } + + if (iexp != igot) + abort (); + + /* BINOP = / */ + lgot = 1LL << 32; + lexp = 1LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + { ldata[i] = lgot; lgot /= expr; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL << 32; + lexp = 1LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + { lgot /= expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL << 32; + lexp = 1LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + { ldata[i] = lgot; lgot = lgot / expr; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 2LL; + lexp = 2LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL << N; + +#pragma acc atomic capture + { ldata[i] = lgot; lgot = expr / lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 2LL; + lexp = 2LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL << N; + +#pragma acc atomic capture + { lgot = lgot / expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 2LL; + lexp = 2LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL << N; + +#pragma acc atomic capture + { lgot = expr / lgot; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + /* BINOP = & */ + lgot = ~0LL; + lexp = 0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { ldata[i] = lgot; lgot &= expr; } + } + } + + if (lexp != lgot) + abort (); + + lgot = ~0LL; + iexp = 0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { lgot &= expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = ~0LL; + lexp = 0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { ldata[i] = lgot; lgot = lgot & expr; } + } + } + + if (lexp != lgot) + abort (); + + lgot = ~0LL; + lexp = 0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { ldata[i] = lgot; lgot = expr & lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = ~0LL; + iexp = 0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { lgot = lgot & expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = ~0LL; + lexp = 0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { lgot = expr & lgot; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + /* BINOP = ^ */ + lgot = ~0LL; + lexp = 0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1 << i; + +#pragma acc atomic capture + { ldata[i] = lgot; lgot ^= expr; } + } + } + + if (lexp != lgot) + abort (); + + lgot = ~0LL; + iexp = 0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { lgot ^= expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = ~0LL; + lexp = 0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { ldata[i] = lgot; lgot = lgot ^ expr; } + } + } + + if (lexp != lgot) + abort (); + + lgot = ~0LL; + lexp = 0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { ldata[i] = lgot; lgot = expr ^ lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = ~0LL; + iexp = 0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { lgot = lgot ^ expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = ~0LL; + lexp = 0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { lgot = expr ^ lgot; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + /* BINOP = | */ + lgot = 0LL; + lexp = ~0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1 << i; + +#pragma acc atomic capture + { ldata[i] = lgot; lgot |= expr; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 0LL; + iexp = ~0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { lgot |= expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 0LL; + lexp = ~0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { ldata[i] = lgot; lgot = lgot | expr; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 0LL; + lexp = ~0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { ldata[i] = lgot; lgot = expr | lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 0LL; + iexp = ~0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { lgot = lgot | expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 0LL; + lexp = ~0LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = ~(1 << i); + +#pragma acc atomic capture + { lgot = expr | lgot; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + /* BINOP = << */ + lgot = 1LL; + lexp = 1LL << N; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL; + +#pragma acc atomic capture + { ldata[i] = lgot; lgot <<= expr; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + iexp = 1LL << N; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL; + +#pragma acc atomic capture + { lgot <<= expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 1LL << N; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL; + +#pragma acc atomic capture + { ldata[i] = lgot; lgot = lgot << expr; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 2LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < 1; i++) + { + long long expr = 1LL; + +#pragma acc atomic capture + { ldata[i] = lgot; lgot = expr << lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 2LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < 1; i++) + { + long long expr = 1LL; + +#pragma acc atomic capture + { lgot = lgot << expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 2LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < 1; i++) + { + long long expr = 1LL; + +#pragma acc atomic capture + { lgot = expr << lgot; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + /* BINOP = >> */ + lgot = 1LL << N; + lexp = 1LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL; + +#pragma acc atomic capture + { ldata[i] = lgot; lgot >>= expr; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL << N; + iexp = 1LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL; + +#pragma acc atomic capture + { lgot >>= expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL << N; + lexp = 1LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL; + +#pragma acc atomic capture + { ldata[i] = lgot; lgot = lgot >> expr; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 1LL << (N - 1); + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < 1; i++) + { + long long expr = 1LL << N; + +#pragma acc atomic capture + { ldata[i] = lgot; lgot = expr >> lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL << N; + lexp = 1LL; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL; + +#pragma acc atomic capture + { lgot = lgot >> expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 1LL << (N - 1); + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < 1; i++) + { + long long expr = 1LL << N; + +#pragma acc atomic capture + { lgot = expr >> lgot; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + // FLOAT FLOAT FLOAT + + /* BINOP = + */ + fgot = 0.0; + fexp = 32.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + { fdata[i] = fgot; fgot += expr; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 0.0; + fexp = 32.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + { fgot += expr; fdata[i] = fgot; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 0.0; + fexp = 32.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + { idata[i] = fgot; fgot = fgot + expr; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 0.0; + fexp = 32.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + { fdata[i] = fgot; fgot = expr + fgot; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 0.0; + fexp = 32.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + { fgot = fgot + expr; fdata[i] = fgot; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 0.0; + fexp = 32.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + { fgot = expr + fgot; fdata[i] = fgot; } + } + } + + if (fexp != fgot) + abort (); + + /* BINOP = * */ + fgot = 1.0; + fexp = 8192.0*8192.0*64.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic capture + { fdata[i] = fgot; fgot *= expr; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 1.0; + fexp = 8192.0*8192.0*64.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic capture + { fgot *= expr; fdata[i] = fgot; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 1.0; + fexp = 8192.0*8192.0*64.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic capture + { fdata[i] = fgot; fgot = fgot * expr; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 1.0; + fexp = 8192.0*8192.0*64.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic capture + { fdata[i] = fgot; fgot = expr * fgot; } + } + } + + if (fexp != fgot) + abort (); + + lgot = 1LL; + lexp = 1LL << 32; + +#pragma acc data copy (lgot, ldata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2LL; + +#pragma acc atomic capture + { lgot = lgot * expr; ldata[i] = lgot; } + } + } + + if (lexp != lgot) + abort (); + + fgot = 1.0; + fexp = 8192.0*8192.0*64.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 2; + +#pragma acc atomic capture + { fgot = expr * fgot; fdata[i] = fgot; } + } + } + + if (fexp != fgot) + abort (); + + /* BINOP = - */ + fgot = 32.0; + fexp = 0.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + { fdata[i] = fgot; fgot -= expr; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 32.0; + fexp = 0.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + { fgot -= expr; fdata[i] = fgot; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 32.0; + fexp = 0.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + { fdata[i] = fgot; fgot = fgot - expr; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 1.0; + fexp = 1.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + { fdata[i] = fgot; fgot = expr - fgot; } + } + } + + for (i = 0; i < N; i++) + if (i % 2 == 0) + { + if (fdata[i] != 1.0) + abort (); + } + else + { + if (fdata[i] != 0.0) + abort (); + } + + if (fexp != fgot) + abort (); + + fgot = 1.0; + fexp = -31.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + { fgot = fgot - expr; fdata[i] = fgot; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 1.0; + fexp = 1.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + { fgot = expr - fgot; fdata[i] = fgot; } + } + } + + for (i = 0; i < N; i++) + if (i % 2 == 0) + { + if (fdata[i] != 0.0) + abort (); + } + else + { + if (fdata[i] != 1.0) + abort (); + } + + if (fexp != fgot) + abort (); + + /* BINOP = / */ + fgot = 8192.0*8192.0*64.0; + fexp = 1.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic capture + { fdata[i] = fgot; fgot /= expr; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 8192.0*8192.0*64.0; + fexp = 1.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic capture + { fgot /= expr; fdata[i] = fgot; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 8192.0*8192.0*64.0; + fexp = 1.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic capture + { fdata[i] = fgot; fgot = fgot / expr; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 8192.0*8192.0*64.0; + fexp = 1.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; + +#pragma acc atomic capture + { fdata[i] = fgot; fgot = expr / fgot; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 4.0; + fexp = 4.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL << N; + +#pragma acc atomic capture + { fgot = fgot / expr; fdata[i] = fgot; } + } + } + + if (fexp != fgot) + abort (); + + fgot = 4.0; + fexp = 4.0; + +#pragma acc data copy (fgot, fdata[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic capture + { fgot = expr / fgot; fdata[i] = fgot; } + } + } + + if (fexp != fgot) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_rw-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_rw-1.c new file mode 100644 index 00000000000..ae4f22e0d29 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_rw-1.c @@ -0,0 +1,34 @@ +/* { dg-do run } */ + +#include + +int +main(int argc, char **argv) +{ + int v1, v2; + int x; + + x = 99; + +#pragma acc parallel copy (v1, v2, x) + { + +#pragma acc atomic read + v1 = x; + +#pragma acc atomic write + x = 32; + +#pragma acc atomic read + v2 = x; + + } + + if (v1 != 99) + abort (); + + if (v2 != 32) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_update-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_update-1.c new file mode 100644 index 00000000000..18ee3aa945b --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_update-1.c @@ -0,0 +1,760 @@ +/* { dg-do run } */ + +#include + +int +main(int argc, char **argv) +{ + float fexp, fgot; + int iexp, igot; + long long lexp, lgot; + int N = 32; + int i; + + fgot = 1234.0; + fexp = 1235.0; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < 1; i++) +#pragma acc atomic update + fgot++; + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + fexp = fgot - N; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { +#pragma acc atomic update + fgot--; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + fexp = fgot + N; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { +#pragma acc atomic update + ++fgot; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + fexp = fgot - N; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { +#pragma acc atomic update + --fgot; + } + } + + if (fexp != fgot) + abort (); + + /* BINOP = + */ + + fgot = 1234.0; + fexp = fgot + N; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; +#pragma acc atomic update + fgot += expr; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + fexp = fgot + N; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; +#pragma acc atomic update + fgot = fgot + expr; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + fexp = fgot + N; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; +#pragma acc atomic update + fgot = expr + fgot; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 0.5; +#pragma acc atomic update + fgot = (expr + expr) + fgot; + } + } + + if (fexp != fgot) + abort (); + + /* BINOP = * */ + + fgot = 1234.0; + fexp = 1234.0; + + for (i = 0; i < N; i++) + fexp *= 2.0; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; +#pragma acc atomic update + fgot *= expr; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + fexp = 1234.0; + + for (i = 0; i < N; i++) + fexp = fexp * 2.0; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; +#pragma acc atomic update + fgot = fgot * expr; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + fexp = 1234.0; + + for (i = 0; i < N; i++) + fexp = 2.0 * fexp; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; +#pragma acc atomic update + fgot = expr * fgot; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; +#pragma acc atomic update + fgot = (expr + expr) * fgot; + } + } + + if (fexp != fgot) + abort (); + + /* BINOP = - */ + + fgot = 1234.0; + fexp = 1234.0; + + for (i = 0; i < N; i++) + fexp -= 2.0; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; +#pragma acc atomic update + fgot -= expr; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + fexp = 1234.0; + + for (i = 0; i < N; i++) + fexp = fexp - 2.0; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; +#pragma acc atomic update + fgot = fgot - expr; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + fexp = 1234.0; + + for (i = 0; i < N; i++) + fexp = 2.0 - fexp; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic update + fgot = expr - fgot; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; +#pragma acc atomic update + fgot = (expr + expr) - fgot; + } + } + + if (fexp != fgot) + abort (); + + /* BINOP = / */ + + fgot = 1234.0; + fexp = 1234.0; + + for (i = 0; i < N; i++) + fexp /= 2.0; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; +#pragma acc atomic update + fgot /= expr; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + fexp = 1234.0; + + for (i = 0; i < N; i++) + fexp = fexp / 2.0; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic update + fgot = fgot / expr; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + fexp = 1234.0; + + for (i = 0; i < N; i++) + fexp = 2.0 / fexp; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 2.0; + +#pragma acc atomic update + fgot = expr / fgot; + } + } + + if (fexp != fgot) + abort (); + + fgot = 1234.0; + fexp = 1234.0; + + for (i = 0; i < N; i++) + fexp = 2.0 / fexp; + +#pragma acc data copy (fgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + float expr = 1.0; +#pragma acc atomic update + fgot = (expr + expr) / fgot; + } + } + + if (fexp != fgot) + abort (); + + /* BINOP = & */ + + igot = ~0; + iexp = 0; + +#pragma acc data copy (igot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = ~(1 << i); + +#pragma acc atomic update + igot &= expr; + } + } + + if (iexp != igot) + abort (); + + igot = ~0; + iexp = 0; + +#pragma acc data copy (igot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = ~(1 << i); +#pragma acc atomic update + igot = igot / expr; + } + } + + if (iexp != igot) + abort (); + + igot = ~0; + iexp = 0; + +#pragma acc data copy (igot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = ~(1 << i); +#pragma acc atomic update + igot = expr & igot; + } + } + + if (iexp != igot) + abort (); + + igot = ~0; + iexp = 0; + +#pragma acc data copy (igot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = ~(1 << i); + int zero = 0; + +#pragma acc atomic update + igot = (expr + zero) & igot; + } + } + + if (iexp != igot) + abort (); + + /* BINOP = ^ */ + + igot = ~0; + iexp = 0; + +#pragma acc data copy (igot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = (1 << i); + +#pragma acc atomic update + igot ^= expr; + } + } + + if (iexp != igot) + abort (); + + igot = ~0; + iexp = 0; + +#pragma acc data copy (igot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = (1 << i); + +#pragma acc atomic update + igot = igot ^ expr; + } + } + + if (iexp != igot) + abort (); + + igot = ~0; + iexp = 0; + +#pragma acc data copy (igot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = (1 << i); + +#pragma acc atomic update + igot = expr ^ igot; + } + } + + if (iexp != igot) + abort (); + + igot = ~0; + iexp = 0; + +#pragma acc data copy (igot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = (1 << i); + int zero = 0; + +#pragma acc atomic update + igot = (expr + zero) ^ igot; + } + } + + if (iexp != igot) + abort (); + + /* BINOP = | */ + + igot = 0; + iexp = ~0; + +#pragma acc data copy (igot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = (1 << i); + +#pragma acc atomic update + igot |= expr; + } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = ~0; + +#pragma acc data copy (igot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = (1 << i); + +#pragma acc atomic update + igot = igot | expr; + } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = ~0; + +#pragma acc data copy (igot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = (1 << i); + +#pragma acc atomic update + igot = expr | igot; + } + } + + if (iexp != igot) + abort (); + + igot = 0; + iexp = ~0; + +#pragma acc data copy (igot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + int expr = (1 << i); + int zero = 0; + +#pragma acc atomic update + igot = (expr + zero) | igot; + } + } + + if (iexp != igot) + abort (); + + /* BINOP = << */ + + lgot = 1LL; + lexp = 1LL << N; + +#pragma acc data copy (lgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL; + +#pragma acc atomic update + lgot <<= expr; + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 1LL << N; + +#pragma acc data copy (lgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL; + +#pragma acc atomic update + lgot = lgot << expr; + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 2LL; + +#pragma acc data copy (lgot) + { +#pragma acc parallel loop + for (i = 0; i < 1; i++) + { + long long expr = 1LL; + +#pragma acc atomic update + lgot = expr << lgot; + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 2LL; + +#pragma acc data copy (lgot) + { +#pragma acc parallel loop + for (i = 0; i < 1; i++) + { + long long expr = 1LL; + long long zero = 0LL; + +#pragma acc atomic update + lgot = (expr + zero) << lgot; + } + } + + if (lexp != lgot) + abort (); + + /* BINOP = >> */ + + lgot = 1LL << N; + lexp = 1LL; + +#pragma acc data copy (lgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL; + +#pragma acc atomic update + lgot >>= expr; + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL << N; + lexp = 1LL; + +#pragma acc data copy (lgot) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + { + long long expr = 1LL; + +#pragma acc atomic update + lgot = lgot >> expr; + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 1LL << (N - 1); + +#pragma acc data copy (lgot) + { +#pragma acc parallel loop + for (i = 0; i < 1; i++) + { + long long expr = 1LL << N; + +#pragma acc atomic update + lgot = expr >> lgot; + } + } + + if (lexp != lgot) + abort (); + + lgot = 1LL; + lexp = 1LL << (N - 1); + +#pragma acc data copy (lgot) + { +#pragma acc parallel loop + for (i = 0; i < 1; i++) + { + long long expr = 1LL << N; + long long zero = 0LL; + +#pragma acc atomic update + lgot = (expr + zero) >> lgot; + } + } + + if (lexp != lgot) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c new file mode 100644 index 00000000000..dbe82fe67e1 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c @@ -0,0 +1,44 @@ +#include + +int +main (int argc, char *argv[]) +{ + int res, res2 = 0; + +#if defined(ACC_DEVICE_TYPE_host) +# define GANGS 1 +#else +# define GANGS 256 +#endif + #pragma acc parallel num_gangs(GANGS) num_workers(1) vector_length(1) \ + copy(res2) + { + #pragma acc atomic + res2 += 5; + } + res = GANGS * 5; + + assert (res == res2); +#undef GANGS + + res = res2 = 1; + +#if defined(ACC_DEVICE_TYPE_host) +# define GANGS 1 +#else +# define GANGS 8 +#endif + #pragma acc parallel num_gangs(GANGS) num_workers(1) vector_length(1) \ + copy(res2) + { + #pragma acc atomic + res2 *= 5; + } + for (int i = 0; i < GANGS; ++i) + res *= 5; + + assert (res == res2); +#undef GANGS + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c new file mode 100644 index 00000000000..12ab552284c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c @@ -0,0 +1,48 @@ +#include +#include + +int +main (int argc, char *argv[]) +{ + int res, res2 = 0; + +#if defined(ACC_DEVICE_TYPE_host) +# define GANGS 1 +#else +# define GANGS 256 +#endif + #pragma acc parallel num_gangs(GANGS) num_workers(1) vector_length(1) \ + copy(res2) async(1) + { + #pragma acc atomic + res2 += 5; + } + res = GANGS * 5; + + acc_wait (1); + + assert (res == res2); +#undef GANGS + + res = res2 = 1; + +#if defined(ACC_DEVICE_TYPE_host) +# define GANGS 1 +#else +# define GANGS 8 +#endif + #pragma acc parallel num_gangs(GANGS) num_workers(1) vector_length(1) \ + copy(res2) async(1) + { + #pragma acc atomic + res2 *= 5; + } + for (int i = 0; i < GANGS; ++i) + res *= 5; + + acc_wait (1); + + assert (res == res2); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-1a.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-1a.c new file mode 100644 index 00000000000..99c6dfbf8a3 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-1a.c @@ -0,0 +1,28 @@ +#include + +/* Test worker-single/vector-single mode. */ + +int +main (int argc, char *argv[]) +{ + int arr[32], i; + + for (i = 0; i < 32; i++) + arr[i] = 0; + + #pragma acc parallel copy(arr) num_gangs(8) num_workers(8) vector_length(32) + { + int j; + #pragma acc loop gang + for (j = 0; j < 32; j++) + { + #pragma acc atomic + arr[j]++; + } + } + + for (i = 0; i < 32; i++) + assert (arr[i] == 1); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-4.c new file mode 100644 index 00000000000..84080d0a8e5 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-4.c @@ -0,0 +1,28 @@ +#include + +/* Test worker-single/vector-partitioned mode. */ + +int +main (int argc, char *argv[]) +{ + int arr[32], i; + + for (i = 0; i < 32; i++) + arr[i] = i; + + #pragma acc parallel copy(arr) num_gangs(1) num_workers(8) vector_length(32) + { + int k; + #pragma acc loop vector + for (k = 0; k < 32; k++) + { + #pragma acc atomic + arr[k]++; + } + } + + for (i = 0; i < 32; i++) + assert (arr[i] == i + 1); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-6.c new file mode 100644 index 00000000000..cbc3e37a3a0 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-6.c @@ -0,0 +1,46 @@ +#include + +#if defined(ACC_DEVICE_TYPE_host) +#define ACTUAL_GANGS 1 +#else +#define ACTUAL_GANGS 8 +#endif + +/* Test worker-single, vector-partitioned, gang-redundant mode. */ + +int +main (int argc, char *argv[]) +{ + int n, arr[32], i; + + for (i = 0; i < 32; i++) + arr[i] = 0; + + n = 0; + + #pragma acc parallel copy(n, arr) num_gangs(ACTUAL_GANGS) num_workers(8) \ + vector_length(32) + { + int j; + + #pragma acc atomic + n++; + + #pragma acc loop vector + for (j = 0; j < 32; j++) + { + #pragma acc atomic + arr[j] += 1; + } + + #pragma acc atomic + n++; + } + + assert (n == ACTUAL_GANGS * 2); + + for (i = 0; i < 32; i++) + assert (arr[i] == ACTUAL_GANGS); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/atomic_capture-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/atomic_capture-1.f90 new file mode 100644 index 00000000000..27c5c9e57bd --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/atomic_capture-1.f90 @@ -0,0 +1,784 @@ +! { dg-do run } + +program main + integer igot, iexp, itmp + real fgot, fexp, ftmp + logical lgot, lexp, ltmp + integer, parameter :: N = 32 + + igot = 0 + iexp = N * 2 + + !$acc parallel copy (igot, itmp) + do i = 1, N + !$acc atomic capture + itmp = igot + igot = i + i + !$acc end atomic + end do + !$acc end parallel + + if (igot /= iexp) call abort + if (itmp /= iexp - 2) call abort + + fgot = 1234.0 + fexp = 1266.0 + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + ftmp = fgot + fgot = fgot + 1.0 + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= fexp - 1.0) call abort + if (fgot /= fexp) call abort + + fgot = 1.0 + fexp = 2.0**32 + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + ftmp = fgot + fgot = fgot * 2.0 + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= fexp / 2.0) call abort + if (fgot /= fexp) call abort + + fgot = 32.0 + fexp = fgot - N + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + ftmp = fgot + fgot = fgot - 1.0 + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= fexp + 1.0) call abort + if (fgot /= fexp) call abort + + fgot = 2**32.0 + fexp = 1.0 + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + ftmp = fgot + fgot = fgot / 2.0 + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= fgot * 2.0) call abort + if (fgot /= fexp) call abort + + lgot = .TRUE. + lexp = .FALSE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + ltmp = lgot + lgot = lgot .and. .FALSE. + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. .not. lexp) call abort + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .FALSE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + ltmp = lgot + lgot = lgot .or. .FALSE. + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. lexp) call abort + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .FALSE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + ltmp = lgot + lgot = lgot .eqv. .TRUE. + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. lexp) call abort + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .TRUE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + ltmp = lgot + lgot = lgot .neqv. .TRUE. + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. .not. lexp) call abort + if (lgot .neqv. lexp) call abort + + fgot = 1234.0 + fexp = 1266.0 + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + ftmp = fgot + fgot = 1.0 + fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= fexp - 1.0) call abort + if (fgot /= fexp) call abort + + fgot = 1.0 + fexp = 2.0**32 + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + ftmp = fgot + fgot = 2.0 * fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= fexp / 2.0) call abort + if (fgot /= fexp) call abort + + fgot = 32.0 + fexp = 32.0 + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + ftmp = fgot + fgot = 2.0 - fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= 2.0 - fexp) call abort + if (fgot /= fexp) call abort + + fgot = 2.0**16 + fexp = 2.0**16 + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + ftmp = fgot + fgot = 2.0 / fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= 2.0 / fexp) call abort + if (fgot /= fexp) call abort + + lgot = .TRUE. + lexp = .FALSE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + ltmp = lgot + lgot = .FALSE. .and. lgot + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. .not. lexp) call abort + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .FALSE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + ltmp = lgot + lgot = .FALSE. .or. lgot + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. lexp) call abort + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .FALSE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + ltmp = lgot + lgot = .TRUE. .eqv. lgot + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. lexp) call abort + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .TRUE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + ltmp = lgot + lgot = .TRUE. .neqv. lgot + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. .not. lexp) call abort + if (lgot .neqv. lexp) call abort + + igot = 1 + iexp = N + + !$acc parallel loop copy (igot, itmp) + do i = 1, N + !$acc atomic capture + itmp = igot + igot = max (igot, i) + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= iexp - 1) call abort + if (igot /= iexp) call abort + + igot = N + iexp = 1 + + !$acc parallel loop copy (igot, itmp) + do i = 1, N + !$acc atomic capture + itmp = igot + igot = min (igot, i) + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= iexp) call abort + if (igot /= iexp) call abort + + igot = -1 + iexp = 0 + + !$acc parallel loop copy (igot, itmp) + do i = 0, N - 1 + iexpr = ibclr (-2, i) + !$acc atomic capture + itmp = igot + igot = iand (igot, iexpr) + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= ibset (iexp, N - 1)) call abort + if (igot /= iexp) call abort + + igot = 0 + iexp = -1 + + !$acc parallel loop copy (igot, itmp) + do i = 0, N - 1 + iexpr = lshift (1, i) + !$acc atomic capture + itmp = igot + igot = ior (igot, iexpr) + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= ieor (iexp, lshift (1, N - 1))) call abort + if (igot /= iexp) call abort + + igot = -1 + iexp = 0 + + !$acc parallel loop copy (igot, itmp) + do i = 0, N - 1 + iexpr = lshift (1, i) + !$acc atomic capture + itmp = igot + igot = ieor (igot, iexpr) + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= ior (iexp, lshift (1, N - 1))) call abort + if (igot /= iexp) call abort + + igot = 1 + iexp = N + + !$acc parallel loop copy (igot, itmp) + do i = 1, N + !$acc atomic capture + itmp = igot + igot = max (i, igot) + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= iexp - 1) call abort + if (igot /= iexp) call abort + + igot = N + iexp = 1 + + !$acc parallel loop copy (igot, itmp) + do i = 1, N + !$acc atomic capture + itmp = igot + igot = min (i, igot) + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= iexp) call abort + if (igot /= iexp) call abort + + igot = -1 + iexp = 0 + + !$acc parallel loop copy (igot, itmp) + do i = 0, N - 1 + iexpr = ibclr (-2, i) + !$acc atomic capture + itmp = igot + igot = iand (iexpr, igot) + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= ibset (iexp, N - 1)) call abort + if (igot /= iexp) call abort + + igot = 0 + iexp = -1 + !! + !$acc parallel loop copy (igot, itmp) + do i = 0, N - 1 + iexpr = lshift (1, i) + !$acc atomic capture + itmp = igot + igot = ior (iexpr, igot) + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= ieor (iexp, lshift (1, N - 1))) call abort + if (igot /= iexp) call abort + + igot = -1 + iexp = 0 + + !$acc parallel loop copy (igot, itmp) + do i = 0, N - 1 + iexpr = lshift (1, i) + !$acc atomic capture + itmp = igot + igot = ieor (iexpr, igot) + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= ior (iexp, lshift (1, N - 1))) call abort + if (igot /= iexp) call abort + + fgot = 1234.0 + fexp = 1266.0 + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + fgot = fgot + 1.0 + ftmp = fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= fexp) call abort + if (fgot /= fexp) call abort + + fgot = 1.0 + fexp = 2.0**32 + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + fgot = fgot * 2.0 + ftmp = fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= fexp) call abort + if (fgot /= fexp) call abort + + fgot = 32.0 + fexp = fgot - N + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + fgot = fgot - 1.0 + ftmp = fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= fexp) call abort + if (fgot /= fexp) call abort + + fgot = 2**32.0 + fexp = 1.0 + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + fgot = fgot / 2.0 + ftmp = fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= fexp) call abort + if (fgot /= fexp) call abort + + lgot = .TRUE. + lexp = .FALSE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + lgot = lgot .and. .FALSE. + ltmp = lgot + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. lexp) call abort + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .FALSE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + lgot = lgot .or. .FALSE. + ltmp = lgot + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. lexp) call abort + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .FALSE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + lgot = lgot .eqv. .TRUE. + ltmp = lgot + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. lexp) call abort + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .TRUE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + lgot = lgot .neqv. .TRUE. + ltmp = lgot + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. lexp) call abort + if (lgot .neqv. lexp) call abort + + fgot = 1234.0 + fexp = 1266.0 + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + fgot = 1.0 + fgot + ftmp = fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= fexp) call abort + if (fgot /= fexp) call abort + + fgot = 1.0 + fexp = 2.0**32 + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + fgot = 2.0 * fgot + ftmp = fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= fexp) call abort + if (fgot /= fexp) call abort + + fgot = 32.0 + fexp = 32.0 + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + fgot = 2.0 - fgot + ftmp = fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= fexp) call abort + if (fgot /= fexp) call abort + + fgot = 2.0**16 + fexp = 2.0**16 + + !$acc parallel loop copy (fgot, ftmp) + do i = 1, N + !$acc atomic capture + fgot = 2.0 / fgot + ftmp = fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (ftmp /= fexp) call abort + if (fgot /= fexp) call abort + + lgot = .TRUE. + lexp = .FALSE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + lgot = .FALSE. .and. lgot + ltmp = lgot + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. lexp) call abort + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .FALSE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + lgot = .FALSE. .or. lgot + ltmp = lgot + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. lexp) call abort + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .FALSE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + lgot = .TRUE. .eqv. lgot + ltmp = lgot + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. lexp) call abort + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .TRUE. + + !$acc parallel copy (lgot, ltmp) + !$acc atomic capture + lgot = .TRUE. .neqv. lgot + ltmp = lgot + !$acc end atomic + !$acc end parallel + + if (ltmp .neqv. lexp) call abort + if (lgot .neqv. lexp) call abort + + igot = 1 + iexp = N + + !$acc parallel loop copy (igot, itmp) + do i = 1, N + !$acc atomic capture + igot = max (igot, i) + itmp = igot + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= iexp) call abort + if (igot /= iexp) call abort + + igot = N + iexp = 1 + + !$acc parallel loop copy (igot, itmp) + do i = 1, N + !$acc atomic capture + igot = min (igot, i) + itmp = igot + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= iexp) call abort + if (igot /= iexp) call abort + + igot = -1 + iexp = 0 + + !$acc parallel loop copy (igot, itmp) + do i = 0, N - 1 + iexpr = ibclr (-2, i) + !$acc atomic capture + igot = iand (igot, iexpr) + itmp = igot + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= iexp) call abort + if (igot /= iexp) call abort + + igot = 0 + iexp = -1 + + !$acc parallel loop copy (igot, itmp) + do i = 0, N - 1 + iexpr = lshift (1, i) + !$acc atomic capture + igot = ior (igot, iexpr) + itmp = igot + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= iexp) call abort + if (igot /= iexp) call abort + + igot = -1 + iexp = 0 + + !$acc parallel loop copy (igot, itmp) + do i = 0, N - 1 + iexpr = lshift (1, i) + !$acc atomic capture + igot = ieor (igot, iexpr) + itmp = igot + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= iexp) call abort + if (igot /= iexp) call abort + + igot = 1 + iexp = N + + !$acc parallel loop copy (igot, itmp) + do i = 1, N + !$acc atomic capture + igot = max (i, igot) + itmp = igot + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= iexp) call abort + if (igot /= iexp) call abort + + igot = N + iexp = 1 + + !$acc parallel loop copy (igot, itmp) + do i = 1, N + !$acc atomic capture + igot = min (i, igot) + itmp = igot + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= iexp) call abort + if (igot /= iexp) call abort + + igot = -1 + iexp = 0 + + !$acc parallel loop copy (igot, itmp) + do i = 0, N - 1 + iexpr = ibclr (-2, i) + !$acc atomic capture + igot = iand (iexpr, igot) + itmp = igot + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= iexp) call abort + if (igot /= iexp) call abort + + igot = 0 + iexp = -1 + + !$acc parallel loop copy (igot, itmp) + do i = 0, N - 1 + iexpr = lshift (1, i) + !$acc atomic capture + igot = ior (iexpr, igot) + itmp = igot + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= iexp) call abort + if (igot /= iexp) call abort + + igot = -1 + iexp = 0 + + !$acc parallel loop copy (igot, itmp) + do i = 0, N - 1 + iexpr = lshift (1, i) + !$acc atomic capture + igot = ieor (iexpr, igot) + itmp = igot + !$acc end atomic + end do + !$acc end parallel loop + + if (itmp /= iexp) call abort + if (igot /= iexp) call abort + +end program diff --git a/libgomp/testsuite/libgomp.oacc-fortran/atomic_rw-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/atomic_rw-1.f90 new file mode 100644 index 00000000000..51ec9aaca87 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/atomic_rw-1.f90 @@ -0,0 +1,29 @@ +! { dg-do run } + +program main + integer v1, v2 + integer x + + x = 99 + + !$acc parallel copy (v1, v2, x) + + !$acc atomic read + v1 = x; + !$acc end atomic + + !$acc atomic write + x = 32; + !$acc end atomic + + !$acc atomic read + v2 = x; + !$acc end atomic + + !$acc end parallel + + if (v1 .ne. 99) call abort + + if (v2 .ne. 32) call abort + +end program main diff --git a/libgomp/testsuite/libgomp.oacc-fortran/atomic_update-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/atomic_update-1.f90 new file mode 100644 index 00000000000..6607c77c935 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/atomic_update-1.f90 @@ -0,0 +1,338 @@ +! { dg-do run } + +program main + integer igot, iexp, iexpr + real fgot, fexp + integer i + integer, parameter :: N = 32 + logical lgot, lexp + + fgot = 1234.0 + fexp = 1266.0 + + !$acc parallel loop copy (fgot) + do i = 1, N + !$acc atomic update + fgot = fgot + 1.0 + !$acc end atomic + end do + !$acc end parallel loop + + if (fgot /= fexp) call abort + + fgot = 1.0 + fexp = 2.0**32 + + !$acc parallel loop copy (fgot) + do i = 1, N + !$acc atomic update + fgot = fgot * 2.0 + !$acc end atomic + end do + !$acc end parallel loop + + if (fgot /= fexp) call abort + + fgot = 32.0 + fexp = fgot - N + + !$acc parallel loop copy (fgot) + do i = 1, N + !$acc atomic update + fgot = fgot - 1.0 + !$acc end atomic + end do + !$acc end parallel loop + + if (fgot /= fexp) call abort + + fgot = 2**32.0 + fexp = 1.0 + + !$acc parallel loop copy (fgot) + do i = 1, N + !$acc atomic update + fgot = fgot / 2.0 + !$acc end atomic + end do + !$acc end parallel loop + + if (fgot /= fexp) call abort + + lgot = .TRUE. + lexp = .FALSE. + + !$acc parallel copy (lgot) + !$acc atomic update + lgot = lgot .and. .FALSE. + !$acc end atomic + !$acc end parallel + + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .FALSE. + + !$acc parallel copy (lgot) + !$acc atomic update + lgot = lgot .or. .FALSE. + !$acc end atomic + !$acc end parallel + + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .FALSE. + + !$acc parallel copy (lgot) + !$acc atomic update + lgot = lgot .eqv. .TRUE. + !$acc end atomic + !$acc end parallel + + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .TRUE. + + !$acc parallel copy (lgot) + !$acc atomic update + lgot = lgot .neqv. .TRUE. + !$acc end atomic + !$acc end parallel + + if (lgot .neqv. lexp) call abort + + fgot = 1234.0 + fexp = 1266.0 + + !$acc parallel loop copy (fgot) + do i = 1, N + !$acc atomic update + fgot = 1.0 + fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (fgot /= fexp) call abort + + fgot = 1.0 + fexp = 2.0**32 + + !$acc parallel loop copy (fgot) + do i = 1, N + !$acc atomic update + fgot = 2.0 * fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (fgot /= fexp) call abort + + fgot = 32.0 + fexp = 32.0 + + !$acc parallel loop copy (fgot) + do i = 1, N + !$acc atomic update + fgot = 2.0 - fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (fgot /= fexp) call abort + + fgot = 2.0**16 + fexp = 2.0**16 + + !$acc parallel loop copy (fgot) + do i = 1, N + !$acc atomic update + fgot = 2.0 / fgot + !$acc end atomic + end do + !$acc end parallel loop + + if (fgot /= fexp) call abort + + lgot = .TRUE. + lexp = .FALSE. + + !$acc parallel copy (lgot) + !$acc atomic update + lgot = .FALSE. .and. lgot + !$acc end atomic + !$acc end parallel + + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .FALSE. + + !$acc parallel copy (lgot) + !$acc atomic update + lgot = .FALSE. .or. lgot + !$acc end atomic + !$acc end parallel + + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .FALSE. + + !$acc parallel copy (lgot) + !$acc atomic update + lgot = .TRUE. .eqv. lgot + !$acc end atomic + !$acc end parallel + + if (lgot .neqv. lexp) call abort + + lgot = .FALSE. + lexp = .TRUE. + + !$acc parallel copy (lgot) + !$acc atomic update + lgot = .TRUE. .neqv. lgot + !$acc end atomic + !$acc end parallel + + if (lgot .neqv. lexp) call abort + + igot = 1 + iexp = N + + !$acc parallel loop copy (igot) + do i = 1, N + !$acc atomic update + igot = max (igot, i) + !$acc end atomic + end do + !$acc end parallel loop + + if (igot /= iexp) call abort + + igot = N + iexp = 1 + + !$acc parallel loop copy (igot) + do i = 1, N + !$acc atomic update + igot = min (igot, i) + !$acc end atomic + end do + !$acc end parallel loop + + if (igot /= iexp) call abort + + igot = -1 + iexp = 0 + + !$acc parallel loop copy (igot) + do i = 0, N - 1 + iexpr = ibclr (-2, i) + !$acc atomic update + igot = iand (igot, iexpr) + !$acc end atomic + end do + !$acc end parallel loop + + if (igot /= iexp) call abort + + igot = 0 + iexp = -1 + + !$acc parallel loop copy (igot) + do i = 0, N - 1 + iexpr = lshift (1, i) + !$acc atomic update + igot = ior (igot, iexpr) + !$acc end atomic + end do + !$acc end parallel loop + + if (igot /= iexp) call abort + + igot = -1 + iexp = 0 + + !$acc parallel loop copy (igot) + do i = 0, N - 1 + iexpr = lshift (1, i) + !$acc atomic update + igot = ieor (igot, iexpr) + !$acc end atomic + end do + !$acc end parallel loop + + if (igot /= iexp) call abort + + igot = 1 + iexp = N + + !$acc parallel loop copy (igot) + do i = 1, N + !$acc atomic update + igot = max (i, igot) + !$acc end atomic + end do + !$acc end parallel loop + + if (igot /= iexp) call abort + + igot = N + iexp = 1 + + !$acc parallel loop copy (igot) + do i = 1, N + !$acc atomic update + igot = min (i, igot) + !$acc end atomic + end do + !$acc end parallel loop + + if (igot /= iexp) call abort + + igot = -1 + iexp = 0 + + !$acc parallel loop copy (igot) + do i = 0, N - 1 + iexpr = ibclr (-2, i) + !$acc atomic update + igot = iand (iexpr, igot) + !$acc end atomic + end do + !$acc end parallel loop + + if (igot /= iexp) call abort + + igot = 0 + iexp = -1 + + !$acc parallel loop copy (igot) + do i = 0, N - 1 + iexpr = lshift (1, i) + !$acc atomic update + igot = ior (iexpr, igot) + !$acc end atomic + end do + !$acc end parallel loop + + if (igot /= iexp) call abort + + igot = -1 + iexp = 0 + + !$acc parallel loop copy (igot) + do i = 0, N - 1 + iexpr = lshift (1, i) + !$acc atomic update + igot = ieor (iexpr, igot) + !$acc end atomic + end do + !$acc end parallel loop + + if (igot /= iexp) call abort + +end program -- cgit v1.2.1