* OpenACC atomic directive
@ 2015-11-02 13:10 Thomas Schwinge
2015-11-02 13:41 ` Jakub Jelinek
0 siblings, 1 reply; 4+ messages in thread
From: Thomas Schwinge @ 2015-11-02 13:10 UTC (permalink / raw)
To: Jakub Jelinek, gcc-patches; +Cc: Chung-Lin Tang
[-- Attachment #1: Type: text/plain, Size: 106353 bytes --]
Hi!
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?
commit 826c7022d0e2b9e225215b168a95487823dce925
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Mon Nov 2 10:35:44 2015 +0100
OpenACC atomic directive
YYYY-MM-DD Thomas Schwinge <thomas@codesourcery.com>
Chung-Lin Tang <cltang@codesourcery.com>
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/
* omp-low.c (check_omp_nesting_restrictions): Allow
GIMPLE_OMP_ATOMIC_LOAD, GIMPLE_OMP_ATOMIC_STORE inside OpenACC
contexts.
YYYY-MM-DD Thomas Schwinge <thomas@codesourcery.com>
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.
YYYY-MM-DD James Norris <jnorris@codesourcery.com>
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.
YYYY-MM-DD Julian Brown <julian@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
libgomp/
* 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.
---
gcc/c-family/c-pragma.c | 1 +
gcc/c-family/c-pragma.h | 1 +
gcc/c/c-parser.c | 3 +
gcc/cp/parser.c | 4 +
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/c-c++-common/goacc-gomp/nesting-1.c | 44 +-
.../c-c++-common/goacc-gomp/nesting-fail-1.c | 26 -
.../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 ++++
27 files changed, 4744 insertions(+), 43 deletions(-)
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/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/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/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/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/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 #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: OpenACC atomic directive
2015-11-02 13:10 OpenACC atomic directive Thomas Schwinge
@ 2015-11-02 13:41 ` Jakub Jelinek
2015-11-03 11:30 ` Thomas Schwinge
0 siblings, 1 reply; 4+ messages in thread
From: Jakub Jelinek @ 2015-11-02 13:41 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches, Chung-Lin Tang
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.
Jakub
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: OpenACC atomic directive
2015-11-02 13:41 ` Jakub Jelinek
@ 2015-11-03 11:30 ` Thomas Schwinge
2021-05-18 9:48 ` Thomas Schwinge
0 siblings, 1 reply; 4+ messages in thread
From: Thomas Schwinge @ 2015-11-03 11:30 UTC (permalink / raw)
To: Jakub Jelinek, gcc-patches; +Cc: Chung-Lin Tang
[-- Attachment #1: Type: text/plain, Size: 114742 bytes --]
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 #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: OpenACC atomic directive
2015-11-03 11:30 ` Thomas Schwinge
@ 2021-05-18 9:48 ` Thomas Schwinge
0 siblings, 0 replies; 4+ messages in thread
From: Thomas Schwinge @ 2021-05-18 9:48 UTC (permalink / raw)
To: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 1917 bytes --]
Hi!
On 2015-11-03T12:30:09+0100, I wrote:
> 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.
> [...], committed in r229703:
> --- gcc/omp-low.c
> +++ gcc/omp-low.c
> @@ -3212,7 +3212,10 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
| /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP builtin)
| inside an OpenACC CTX. */
| if (!(is_gimple_omp (stmt)
| && is_gimple_omp_oacc (stmt)))
> {
> 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");
(The 'gimple_code (stmt) == [...]' expressions later got hoisted out of
the 'ctx_' loop.)
As came up in a different context, it's actually conceptually clearer to
make sure that 'is_gimple_omp_oacc' isn't even used for such "ambiguous"
codes, and here appears to be the only place we do. Pushed "[OMP]
Tighten 'is_gimple_omp_oacc'" to master branch in commit
e1cca88019ab1208f8389606dd18a25cd50c20e1, see attached.
Grüße
Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-OMP-Tighten-is_gimple_omp_oacc.patch --]
[-- Type: text/x-diff, Size: 1953 bytes --]
From e1cca88019ab1208f8389606dd18a25cd50c20e1 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 7 May 2021 10:13:49 +0200
Subject: [PATCH] [OMP] Tighten 'is_gimple_omp_oacc'
No overall change in behavior.
gcc/
* gimple.h (is_gimple_omp_oacc): Tighten.
* omp-low.c (check_omp_nesting_restrictions): Adjust.
---
gcc/gimple.h | 8 ++++++++
gcc/omp-low.c | 11 ++++++-----
2 files changed, 14 insertions(+), 5 deletions(-)
diff --git a/gcc/gimple.h b/gcc/gimple.h
index 3ec86f5f082..91b92b4a4d1 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -6501,6 +6501,14 @@ is_gimple_omp_oacc (const gimple *stmt)
gcc_assert (is_gimple_omp (stmt));
switch (gimple_code (stmt))
{
+ case GIMPLE_OMP_ATOMIC_LOAD:
+ case GIMPLE_OMP_ATOMIC_STORE:
+ case GIMPLE_OMP_CONTINUE:
+ case GIMPLE_OMP_RETURN:
+ /* Codes shared between OpenACC and OpenMP cannot be used to disambiguate
+ the two. */
+ gcc_unreachable ();
+
case GIMPLE_OMP_FOR:
switch (gimple_omp_for_kind (stmt))
{
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index cadca7e201f..d1136d181b3 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -3016,11 +3016,12 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
/* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP builtin)
inside an OpenACC CTX. */
- if (!(is_gimple_omp (stmt)
- && is_gimple_omp_oacc (stmt))
- /* Except for atomic codes that we share with OpenMP. */
- && !(gimple_code (stmt) == GIMPLE_OMP_ATOMIC_LOAD
- || gimple_code (stmt) == GIMPLE_OMP_ATOMIC_STORE))
+ if (gimple_code (stmt) == GIMPLE_OMP_ATOMIC_LOAD
+ || gimple_code (stmt) == GIMPLE_OMP_ATOMIC_STORE)
+ /* ..., except for the atomic codes that OpenACC shares with OpenMP. */
+ ;
+ else if (!(is_gimple_omp (stmt)
+ && is_gimple_omp_oacc (stmt)))
{
if (oacc_get_fn_attrib (cfun->decl) != NULL)
{
--
2.30.2
^ permalink raw reply [flat|nested] 4+ messages in thread
end of thread, other threads:[~2021-05-18 9:49 UTC | newest]
Thread overview: 4+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-11-02 13:10 OpenACC atomic directive Thomas Schwinge
2015-11-02 13:41 ` Jakub Jelinek
2015-11-03 11:30 ` Thomas Schwinge
2021-05-18 9:48 ` Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).