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

13 files changed, 4653 insertions, 0 deletions

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  <julian@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+
+	* 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  <jnorris@codesourcery.com>
+
+	* 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  <nathan@codesourcery.com>
 
 	* 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 <stdlib.h>
+
+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 <stdlib.h>
+
+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 <stdlib.h>
+
+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 <stdlib.h>
+
+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 <assert.h>
+
+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 <assert.h>
+#include <openacc.h>
+
+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 <assert.h>
+
+/* 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 <assert.h>
+
+/* 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 <assert.h>
+
+#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