This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

Re: OpenACC atomic directive


Hi!

On Mon, 2 Nov 2015 14:40:54 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> On Mon, Nov 02, 2015 at 02:09:38PM +0100, Thomas Schwinge wrote:
> > The OpenACC atomic directive matches OpenMP's atomic directive (got that
> > clarified by the OpenACC committee), so they can share the same
> > implementation.  OK for trunk?
> 
> Ok.

Thanks for the speedy review!

Testing the x86 -m32 multilib, I noticed one problem:

    $ build-gcc/gcc/xgcc -Bbuild-gcc/gcc/ source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c -fopenacc -foffload=disable -O2 -m32 -march=i486 -c
    source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c: In function 'main':
    source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c:564:9: internal compiler error: Segmentation fault
     #pragma acc atomic capture
             ^
    0xaf041f crash_signal
            [...]/source-gcc/gcc/toplev.c:336
    0xd7c908 contains_struct_check(tree_node*, tree_node_structure_enum, char const*, int, char const*)
            [...]/source-gcc/gcc/tree.h:3032
    0xd7c908 build_call_expr_loc_array(unsigned int, tree_node*, int, tree_node**)
            [...]/source-gcc/gcc/tree.c:10991
    0xd7cb46 build_call_expr(tree_node*, int, ...)
            [...]/source-gcc/gcc/tree.c:11041
    0x9eaf7b expand_omp_atomic_mutex
            [...]/source-gcc/gcc/omp-low.c:11903
    0x9eaf7b expand_omp_atomic
            [...]/source-gcc/gcc/omp-low.c:11990
    [...]

gcc/omp-low.c:

     11892	expand_omp_atomic_mutex (basic_block load_bb, basic_block store_bb,
     11893				 tree addr, tree loaded_val, tree stored_val)
     11894	{
     [...]
     11902	  t = builtin_decl_explicit (BUILT_IN_GOMP_ATOMIC_START);
     11903	  t = build_call_expr (t, 0);

As discussed in
<http://news.gmane.org/find-root.php?message_id=%3C20150709155837.GW10247%40tucnak.redhat.com%3E>
a while ago (and as present on gomp-4_0-branch ever since then), I fixed
this as follows:

--- gcc/builtins.def
+++ gcc/builtins.def
@@ -182,7 +182,8 @@ along with GCC; see the file COPYING3.  If not see
 #define DEF_GOMP_BUILTIN(ENUM, NAME, TYPE, ATTRS) \
   DEF_BUILTIN (ENUM, "__builtin_" NAME, BUILT_IN_NORMAL, TYPE, TYPE,    \
                false, true, true, ATTRS, false, \
-	       (flag_openmp \
+	       (flag_openacc \
+		|| flag_openmp \
 		|| flag_tree_parallelize_loops > 1 \
 		|| flag_cilkplus \
 		|| flag_offload_abi != OFFLOAD_ABI_UNSET))

With that addded, committed in r229703:

commit 9e10bfb7c761a1343ecf51cfff60e03839561962
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Tue Nov 3 11:28:22 2015 +0000

    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
---
 gcc/ChangeLog                                      |    8 +
 gcc/builtins.def                                   |    3 +-
 gcc/c-family/ChangeLog                             |    6 +
 gcc/c-family/c-pragma.c                            |    1 +
 gcc/c-family/c-pragma.h                            |    1 +
 gcc/c/ChangeLog                                    |    5 +
 gcc/c/c-parser.c                                   |    3 +
 gcc/cp/ChangeLog                                   |    6 +
 gcc/cp/parser.c                                    |    4 +
 gcc/fortran/ChangeLog                              |   27 +
 gcc/fortran/gfortran.h                             |    3 +-
 gcc/fortran/match.h                                |    1 +
 gcc/fortran/openmp.c                               |   22 +-
 gcc/fortran/parse.c                                |   40 +-
 gcc/fortran/resolve.c                              |    2 +
 gcc/fortran/st.c                                   |    1 +
 gcc/fortran/trans-openmp.c                         |    2 +
 gcc/fortran/trans.c                                |    1 +
 gcc/omp-low.c                                      |    5 +-
 gcc/testsuite/ChangeLog                            |    7 +
 gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c  |   44 +-
 .../c-c++-common/goacc-gomp/nesting-fail-1.c       |   26 -
 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 ++++
 35 files changed, 4827 insertions(+), 44 deletions(-)

diff --git gcc/ChangeLog gcc/ChangeLog
index 5459551..947b9a7 100644
--- gcc/ChangeLog
+++ gcc/ChangeLog
@@ -1,3 +1,11 @@
+2015-11-03  Thomas Schwinge  <thomas@codesourcery.com>
+	    Chung-Lin Tang  <cltang@codesourcery.com>
+
+	* 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.
+
 2015-11-03  Bilyan Borisov  <bilyan.borisov@arm.com>
 
 	* config/aarch64/aarch64-simd-builtins.def (fmulx): New.
diff --git gcc/builtins.def gcc/builtins.def
index 2cb8251..886b45c 100644
--- gcc/builtins.def
+++ gcc/builtins.def
@@ -182,7 +182,8 @@ along with GCC; see the file COPYING3.  If not see
 #define DEF_GOMP_BUILTIN(ENUM, NAME, TYPE, ATTRS) \
   DEF_BUILTIN (ENUM, "__builtin_" NAME, BUILT_IN_NORMAL, TYPE, TYPE,    \
                false, true, true, ATTRS, false, \
-	       (flag_openmp \
+	       (flag_openacc \
+		|| flag_openmp \
 		|| flag_tree_parallelize_loops > 1 \
 		|| flag_cilkplus \
 		|| flag_offload_abi != OFFLOAD_ABI_UNSET))
diff --git gcc/c-family/ChangeLog gcc/c-family/ChangeLog
index 3ed9421..7f56722 100644
--- gcc/c-family/ChangeLog
+++ gcc/c-family/ChangeLog
@@ -1,3 +1,9 @@
+2015-11-03  Thomas Schwinge  <thomas@codesourcery.com>
+	    Chung-Lin Tang  <cltang@codesourcery.com>
+
+	* c-pragma.c (oacc_pragmas): Add "atomic".
+	* c-pragma.h (pragma_kind): Add PRAGMA_OACC_ATOMIC.
+
 2015-10-30  Evgeny Stupachenko  <evstupac@gmail.com>
 
 	* c-common.c (handle_target_clones_attribute): New.
diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c
index d99c2af..ac11838 100644
--- gcc/c-family/c-pragma.c
+++ gcc/c-family/c-pragma.c
@@ -1204,6 +1204,7 @@ static vec<pragma_ns_name> registered_pp_pragmas;
 
 struct omp_pragma_def { const char *name; unsigned int id; };
 static const struct omp_pragma_def oacc_pragmas[] = {
+  { "atomic", PRAGMA_OACC_ATOMIC },
   { "cache", PRAGMA_OACC_CACHE },
   { "data", PRAGMA_OACC_DATA },
   { "enter", PRAGMA_OACC_ENTER_DATA },
diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h
index cec920f..69e7392 100644
--- gcc/c-family/c-pragma.h
+++ gcc/c-family/c-pragma.h
@@ -27,6 +27,7 @@ along with GCC; see the file COPYING3.  If not see
 enum pragma_kind {
   PRAGMA_NONE = 0,
 
+  PRAGMA_OACC_ATOMIC,
   PRAGMA_OACC_CACHE,
   PRAGMA_OACC_DATA,
   PRAGMA_OACC_ENTER_DATA,
diff --git gcc/c/ChangeLog gcc/c/ChangeLog
index 20df086..7ddafa3 100644
--- gcc/c/ChangeLog
+++ gcc/c/ChangeLog
@@ -1,3 +1,8 @@
+2015-11-03  Thomas Schwinge  <thomas@codesourcery.com>
+	    Chung-Lin Tang  <cltang@codesourcery.com>
+
+	* c-parser.c (c_parser_omp_construct): Handle PRAGMA_OACC_ATOMIC.
+
 2015-10-29  Andrew MacLeod  <amacleod@redhat.com>
 
 	* c-array-notation.c: Reorder #include's and remove duplicates.
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 90038d5..ec88c65 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -16243,6 +16243,9 @@ c_parser_omp_construct (c_parser *parser)
 
   switch (p_kind)
     {
+    case PRAGMA_OACC_ATOMIC:
+      c_parser_omp_atomic (loc, parser);
+      return;
     case PRAGMA_OACC_CACHE:
       strcpy (p_name, "#pragma acc");
       stmt = c_parser_oacc_cache (loc, parser);
diff --git gcc/cp/ChangeLog gcc/cp/ChangeLog
index ad7087e..884da0f 100644
--- gcc/cp/ChangeLog
+++ gcc/cp/ChangeLog
@@ -1,3 +1,9 @@
+2015-11-03  Thomas Schwinge  <thomas@codesourcery.com>
+	    Chung-Lin Tang  <cltang@codesourcery.com>
+
+	* parser.c (cp_parser_omp_construct, cp_parser_pragma): Handle
+	PRAGMA_OACC_ATOMIC.
+
 2015-10-31  Ville Voutilainen  <ville.voutilainen@gmail.com>
 
 	Remove the implementation of N3994, terse range-for loops.
diff --git gcc/cp/parser.c gcc/cp/parser.c
index 24cb47f..a90bf3b 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -35464,6 +35464,9 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok)
 
   switch (pragma_tok->pragma_kind)
     {
+    case PRAGMA_OACC_ATOMIC:
+      cp_parser_omp_atomic (parser, pragma_tok);
+      return;
     case PRAGMA_OACC_CACHE:
       stmt = cp_parser_oacc_cache (parser, pragma_tok);
       break;
@@ -36040,6 +36043,7 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context)
       cp_parser_omp_declare (parser, pragma_tok, context);
       return false;
 
+    case PRAGMA_OACC_ATOMIC:
     case PRAGMA_OACC_CACHE:
     case PRAGMA_OACC_DATA:
     case PRAGMA_OACC_ENTER_DATA:
diff --git gcc/fortran/ChangeLog gcc/fortran/ChangeLog
index a7c2cef..108e2fc 100644
--- gcc/fortran/ChangeLog
+++ gcc/fortran/ChangeLog
@@ -1,3 +1,30 @@
+2015-11-03  Thomas Schwinge  <thomas@codesourcery.com>
+	    Chung-Lin Tang  <cltang@codesourcery.com>
+
+	* 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.
+
 2015-10-31  Cesar Philippidis  <cesar@codesourcery.com>
 
 	PR Bootstrap/68168
diff --git gcc/fortran/gfortran.h gcc/fortran/gfortran.h
index 13e730f..e13b4d4 100644
--- gcc/fortran/gfortran.h
+++ gcc/fortran/gfortran.h
@@ -209,6 +209,7 @@ enum gfc_statement
   ST_OACC_END_LOOP, ST_OACC_DECLARE, ST_OACC_UPDATE, ST_OACC_WAIT,
   ST_OACC_CACHE, ST_OACC_KERNELS_LOOP, ST_OACC_END_KERNELS_LOOP,
   ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA, ST_OACC_ROUTINE,
+  ST_OACC_ATOMIC, ST_OACC_END_ATOMIC,
   ST_OMP_ATOMIC, ST_OMP_BARRIER, ST_OMP_CRITICAL, ST_OMP_END_ATOMIC,
   ST_OMP_END_CRITICAL, ST_OMP_END_DO, ST_OMP_END_MASTER, ST_OMP_END_ORDERED,
   ST_OMP_END_PARALLEL, ST_OMP_END_PARALLEL_DO, ST_OMP_END_PARALLEL_SECTIONS,
@@ -2322,7 +2323,7 @@ enum gfc_exec_op
   EXEC_OACC_KERNELS_LOOP, EXEC_OACC_PARALLEL_LOOP,
   EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS, EXEC_OACC_DATA, EXEC_OACC_HOST_DATA,
   EXEC_OACC_LOOP, EXEC_OACC_UPDATE, EXEC_OACC_WAIT, EXEC_OACC_CACHE,
-  EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA,
+  EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA, EXEC_OACC_ATOMIC,
   EXEC_OMP_CRITICAL, EXEC_OMP_DO, EXEC_OMP_FLUSH, EXEC_OMP_MASTER,
   EXEC_OMP_ORDERED, EXEC_OMP_PARALLEL, EXEC_OMP_PARALLEL_DO,
   EXEC_OMP_PARALLEL_SECTIONS, EXEC_OMP_PARALLEL_WORKSHARE,
diff --git gcc/fortran/match.h gcc/fortran/match.h
index 1b51a88..a52c189 100644
--- gcc/fortran/match.h
+++ gcc/fortran/match.h
@@ -124,6 +124,7 @@ gfc_common_head *gfc_get_common (const char *, int);
 /* openmp.c.  */
 
 /* OpenACC directive matchers.  */
+match gfc_match_oacc_atomic (void);
 match gfc_match_oacc_cache (void);
 match gfc_match_oacc_wait (void);
 match gfc_match_oacc_update (void);
diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c
index e59139c..929a739 100644
--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -2452,8 +2452,8 @@ gfc_match_omp_ordered (void)
 }
 
 
-match
-gfc_match_omp_atomic (void)
+static match
+gfc_match_omp_oacc_atomic (bool omp_p)
 {
   gfc_omp_atomic_op op = GFC_OMP_ATOMIC_UPDATE;
   int seq_cst = 0;
@@ -2491,13 +2491,24 @@ gfc_match_omp_atomic (void)
       gfc_error ("Unexpected junk after $OMP ATOMIC statement at %C");
       return MATCH_ERROR;
     }
-  new_st.op = EXEC_OMP_ATOMIC;
+  new_st.op = (omp_p ? EXEC_OMP_ATOMIC : EXEC_OACC_ATOMIC);
   if (seq_cst)
     op = (gfc_omp_atomic_op) (op | GFC_OMP_ATOMIC_SEQ_CST);
   new_st.ext.omp_atomic = op;
   return MATCH_YES;
 }
 
+match
+gfc_match_oacc_atomic (void)
+{
+  return gfc_match_omp_oacc_atomic (false);
+}
+
+match
+gfc_match_omp_atomic (void)
+{
+  return gfc_match_omp_oacc_atomic (true);
+}
 
 match
 gfc_match_omp_barrier (void)
@@ -4317,6 +4328,8 @@ oacc_code_to_statement (gfc_code *code)
       return ST_OACC_KERNELS_LOOP;
     case EXEC_OACC_LOOP:
       return ST_OACC_LOOP;
+    case EXEC_OACC_ATOMIC:
+      return ST_OACC_ATOMIC;
     default:
       gcc_unreachable ();
     }
@@ -4661,6 +4674,9 @@ gfc_resolve_oacc_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED)
     case EXEC_OACC_LOOP:
       resolve_oacc_loop (code);
       break;
+    case EXEC_OACC_ATOMIC:
+      resolve_omp_atomic (code);
+      break;
     default:
       break;
     }
diff --git gcc/fortran/parse.c gcc/fortran/parse.c
index 650135b..b98dda1 100644
--- gcc/fortran/parse.c
+++ gcc/fortran/parse.c
@@ -637,6 +637,9 @@ decode_oacc_directive (void)
 
   switch (c)
     {
+    case 'a':
+      match ("atomic", gfc_match_oacc_atomic, ST_OACC_ATOMIC);
+      break;
     case 'c':
       match ("cache", gfc_match_oacc_cache, ST_OACC_CACHE);
       break;
@@ -645,6 +648,7 @@ decode_oacc_directive (void)
       match ("declare", gfc_match_oacc_declare, ST_OACC_DECLARE);
       break;
     case 'e':
+      match ("end atomic", gfc_match_omp_eos, ST_OACC_END_ATOMIC);
       match ("end data", gfc_match_omp_eos, ST_OACC_END_DATA);
       match ("end host_data", gfc_match_omp_eos, ST_OACC_END_HOST_DATA);
       match ("end kernels loop", gfc_match_omp_eos, ST_OACC_END_KERNELS_LOOP);
@@ -1373,7 +1377,8 @@ next_statement (void)
   case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: \
   case ST_CRITICAL: \
   case ST_OACC_PARALLEL_LOOP: case ST_OACC_PARALLEL: case ST_OACC_KERNELS: \
-  case ST_OACC_DATA: case ST_OACC_HOST_DATA: case ST_OACC_LOOP: case ST_OACC_KERNELS_LOOP
+  case ST_OACC_DATA: case ST_OACC_HOST_DATA: case ST_OACC_LOOP: \
+  case ST_OACC_KERNELS_LOOP: case ST_OACC_ATOMIC
 
 /* Declaration statements */
 
@@ -1937,6 +1942,12 @@ gfc_ascii_statement (gfc_statement st)
     case ST_OACC_ROUTINE:
       p = "!$ACC ROUTINE";
       break;
+    case ST_OACC_ATOMIC:
+      p = "!ACC ATOMIC";
+      break;
+    case ST_OACC_END_ATOMIC:
+      p = "!ACC END ATOMIC";
+      break;
     case ST_OMP_ATOMIC:
       p = "!$OMP ATOMIC";
       break;
@@ -4316,14 +4327,24 @@ parse_omp_do (gfc_statement omp_st)
 /* Parse the statements of OpenMP atomic directive.  */
 
 static gfc_statement
-parse_omp_atomic (void)
+parse_omp_oacc_atomic (bool omp_p)
 {
-  gfc_statement st;
+  gfc_statement st, st_atomic, st_end_atomic;
   gfc_code *cp, *np;
   gfc_state_data s;
   int count;
 
-  accept_statement (ST_OMP_ATOMIC);
+  if (omp_p)
+    {
+      st_atomic = ST_OMP_ATOMIC;
+      st_end_atomic = ST_OMP_END_ATOMIC;
+    }
+  else
+    {
+      st_atomic = ST_OACC_ATOMIC;
+      st_end_atomic = ST_OACC_END_ATOMIC;
+    }
+  accept_statement (st_atomic);
 
   cp = gfc_state_stack->tail;
   push_state (&s, COMP_OMP_STRUCTURED_BLOCK, NULL);
@@ -4350,7 +4371,7 @@ parse_omp_atomic (void)
   pop_state ();
 
   st = next_statement ();
-  if (st == ST_OMP_END_ATOMIC)
+  if (st == st_end_atomic)
     {
       gfc_clear_new_st ();
       gfc_commit_symbols ();
@@ -4646,7 +4667,7 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only)
 		  continue;
 
 		case ST_OMP_ATOMIC:
-		  st = parse_omp_atomic ();
+		  st = parse_omp_oacc_atomic (true);
 		  continue;
 
 		default:
@@ -4865,8 +4886,12 @@ parse_executable (gfc_statement st)
 	    return st;
 	  continue;
 
+	case ST_OACC_ATOMIC:
+	  st = parse_omp_oacc_atomic (false);
+	  continue;
+
 	case ST_OMP_ATOMIC:
-	  st = parse_omp_atomic ();
+	  st = parse_omp_oacc_atomic (true);
 	  continue;
 
 	default:
@@ -5782,6 +5807,7 @@ is_oacc (gfc_state_data *sd)
     case EXEC_OACC_CACHE:
     case EXEC_OACC_ENTER_DATA:
     case EXEC_OACC_EXIT_DATA:
+    case EXEC_OACC_ATOMIC:
       return true;
 
     default:
diff --git gcc/fortran/resolve.c gcc/fortran/resolve.c
index 1049c0c..bf2837c 100644
--- gcc/fortran/resolve.c
+++ gcc/fortran/resolve.c
@@ -9372,6 +9372,7 @@ gfc_resolve_blocks (gfc_code *b, gfc_namespace *ns)
 	case EXEC_OACC_CACHE:
 	case EXEC_OACC_ENTER_DATA:
 	case EXEC_OACC_EXIT_DATA:
+	case EXEC_OACC_ATOMIC:
 	case EXEC_OMP_ATOMIC:
 	case EXEC_OMP_CRITICAL:
 	case EXEC_OMP_DISTRIBUTE:
@@ -10644,6 +10645,7 @@ start:
 	case EXEC_OACC_CACHE:
 	case EXEC_OACC_ENTER_DATA:
 	case EXEC_OACC_EXIT_DATA:
+	case EXEC_OACC_ATOMIC:
 	  gfc_resolve_oacc_directive (code, ns);
 	  break;
 
diff --git gcc/fortran/st.c gcc/fortran/st.c
index 116af15..629b51d 100644
--- gcc/fortran/st.c
+++ gcc/fortran/st.c
@@ -240,6 +240,7 @@ gfc_free_statement (gfc_code *p)
       gfc_free_omp_namelist (p->ext.omp_namelist);
       break;
 
+    case EXEC_OACC_ATOMIC:
     case EXEC_OMP_ATOMIC:
     case EXEC_OMP_BARRIER:
     case EXEC_OMP_MASTER:
diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c
index 7e01e72..5f4c382 100644
--- gcc/fortran/trans-openmp.c
+++ gcc/fortran/trans-openmp.c
@@ -4409,6 +4409,8 @@ gfc_trans_oacc_directive (gfc_code *code)
       return gfc_trans_oacc_executable_directive (code);
     case EXEC_OACC_WAIT:
       return gfc_trans_oacc_wait_directive (code);
+    case EXEC_OACC_ATOMIC:
+      return gfc_trans_omp_atomic (code);
     default:
       gcc_unreachable ();
     }
diff --git gcc/fortran/trans.c gcc/fortran/trans.c
index 4337fcb..9495450 100644
--- gcc/fortran/trans.c
+++ gcc/fortran/trans.c
@@ -1903,6 +1903,7 @@ trans_code (gfc_code * code, tree cond)
 	case EXEC_OACC_PARALLEL_LOOP:
 	case EXEC_OACC_ENTER_DATA:
 	case EXEC_OACC_EXIT_DATA:
+	case EXEC_OACC_ATOMIC:
 	  res = gfc_trans_oacc_directive (code);
 	  break;
 
diff --git gcc/omp-low.c gcc/omp-low.c
index d0264e9..ccf0b63 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -3212,7 +3212,10 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
     {
       for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
 	if (is_gimple_omp (ctx_->stmt)
-	    && is_gimple_omp_oacc (ctx_->stmt))
+	    && is_gimple_omp_oacc (ctx_->stmt)
+	    /* Except for atomic codes that we share with OpenMP.  */
+	    && ! (gimple_code (stmt) == GIMPLE_OMP_ATOMIC_LOAD
+		  || gimple_code (stmt) == GIMPLE_OMP_ATOMIC_STORE))
 	  {
 	    error_at (gimple_location (stmt),
 		      "non-OpenACC construct inside of OpenACC region");
diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog
index 1de0ea1..c1ac2b0 100644
--- gcc/testsuite/ChangeLog
+++ gcc/testsuite/ChangeLog
@@ -1,3 +1,10 @@
+2015-11-03  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* 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.
+
 2015-11-03  Bilyan Borisov  <bilyan.borisov@arm.com>
 
 	* gcc/testsuite/gcc.target/aarch64/simd/vmulx_f32_1.c: New.
diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
index 1c17818..dabba8c 100644
--- gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
+++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
@@ -1,12 +1,46 @@
 void
-f_omp_parallel (void)
+f_acc_data (void)
 {
-#pragma omp parallel
+#pragma acc data
   {
     int i;
+#pragma omp atomic write
+    i = 0;
+  }
+}
+
+void
+f_acc_kernels (void)
+{
+#pragma acc kernels
+  {
+    int i;
+#pragma omp atomic write
+    i = 0;
+  }
+}
 
-#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
-    for (i = 0; i < 2; ++i)
-      ;
+void
+f_acc_loop (void)
+{
+  int i;
+
+#pragma acc parallel
+#pragma acc loop
+  for (i = 0; i < 2; ++i)
+    {
+#pragma omp atomic write
+      i = 0;
+    }
+}
+
+void
+f_acc_parallel (void)
+{
+#pragma acc parallel
+  {
+    int i;
+#pragma omp atomic write
+    i = 0;
   }
 }
diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
index 0c8ea54..e98258c 100644
--- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
+++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
@@ -216,12 +216,6 @@ f_acc_parallel (void)
 
 #pragma acc parallel
   {
-#pragma omp atomic write
-    i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
-  }
-
-#pragma acc parallel
-  {
 #pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
     ;
   }
@@ -286,12 +280,6 @@ f_acc_kernels (void)
 
 #pragma acc kernels
   {
-#pragma omp atomic write
-    i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
-  }
-
-#pragma acc kernels
-  {
 #pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
     ;
   }
@@ -356,12 +344,6 @@ f_acc_data (void)
 
 #pragma acc data
   {
-#pragma omp atomic write
-    i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
-  }
-
-#pragma acc data
-  {
 #pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
     ;
   }
@@ -442,14 +424,6 @@ f_acc_loop (void)
 #pragma acc loop
   for (i = 0; i < 2; ++i)
     {
-#pragma omp atomic write
-      i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
-    }
-
-#pragma acc parallel
-#pragma acc loop
-  for (i = 0; i < 2; ++i)
-    {
 #pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
       ;
     }
diff --git libgomp/ChangeLog libgomp/ChangeLog
index 4d19066..7894a7e 100644
--- libgomp/ChangeLog
+++ 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 libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c
new file mode 100644
index 0000000..ad958cd
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c
new file mode 100644
index 0000000..842f2de
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_rw-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_rw-1.c
new file mode 100644
index 0000000..ae4f22e
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_update-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_update-1.c
new file mode 100644
index 0000000..18ee3aa
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c
new file mode 100644
index 0000000..dbe82fe
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c
new file mode 100644
index 0000000..12ab552
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-1a.c libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-1a.c
new file mode 100644
index 0000000..99c6dfb
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-4.c libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-4.c
new file mode 100644
index 0000000..84080d0
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-6.c libgomp/testsuite/libgomp.oacc-c-c++-common/worker-single-6.c
new file mode 100644
index 0000000..cbc3e37
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.oacc-fortran/atomic_capture-1.f90 libgomp/testsuite/libgomp.oacc-fortran/atomic_capture-1.f90
new file mode 100644
index 0000000..27c5c9e
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.oacc-fortran/atomic_rw-1.f90 libgomp/testsuite/libgomp.oacc-fortran/atomic_rw-1.f90
new file mode 100644
index 0000000..51ec9aa
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.oacc-fortran/atomic_update-1.f90 libgomp/testsuite/libgomp.oacc-fortran/atomic_update-1.f90
new file mode 100644
index 0000000..6607c77
--- /dev/null
+++ 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


GrÃÃe
 Thomas

Attachment: signature.asc
Description: PGP signature


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]