aboutsummaryrefslogtreecommitdiff
path: root/gcc
diff options
context:
space:
mode:
Diffstat (limited to 'gcc')
-rw-r--r--gcc/ChangeLog8
-rw-r--r--gcc/builtins.def3
-rw-r--r--gcc/c-family/ChangeLog6
-rw-r--r--gcc/c-family/c-pragma.c1
-rw-r--r--gcc/c-family/c-pragma.h1
-rw-r--r--gcc/c/ChangeLog5
-rw-r--r--gcc/c/c-parser.c3
-rw-r--r--gcc/cp/ChangeLog6
-rw-r--r--gcc/cp/parser.c4
-rw-r--r--gcc/fortran/ChangeLog27
-rw-r--r--gcc/fortran/gfortran.h3
-rw-r--r--gcc/fortran/match.h1
-rw-r--r--gcc/fortran/openmp.c22
-rw-r--r--gcc/fortran/parse.c40
-rw-r--r--gcc/fortran/resolve.c2
-rw-r--r--gcc/fortran/st.c1
-rw-r--r--gcc/fortran/trans-openmp.c2
-rw-r--r--gcc/fortran/trans.c1
-rw-r--r--gcc/omp-low.c5
-rw-r--r--gcc/testsuite/ChangeLog7
-rw-r--r--gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c44
-rw-r--r--gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c26
22 files changed, 174 insertions, 44 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 5459551..947b9a7 100644
--- a/gcc/ChangeLog
+++ b/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 a/gcc/builtins.def b/gcc/builtins.def
index 2cb8251..886b45c 100644
--- a/gcc/builtins.def
+++ b/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 a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog
index 3ed9421..7f56722 100644
--- a/gcc/c-family/ChangeLog
+++ b/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 a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c
index d99c2af..ac11838 100644
--- a/gcc/c-family/c-pragma.c
+++ b/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 a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index cec920f..69e7392 100644
--- a/gcc/c-family/c-pragma.h
+++ b/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 a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index 20df086..7ddafa3 100644
--- a/gcc/c/ChangeLog
+++ b/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 a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 90038d5..ec88c65 100644
--- a/gcc/c/c-parser.c
+++ b/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 a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index ad7087e..884da0f 100644
--- a/gcc/cp/ChangeLog
+++ b/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 a/gcc/cp/parser.c b/gcc/cp/parser.c
index 24cb47f..a90bf3b 100644
--- a/gcc/cp/parser.c
+++ b/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 a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog
index a7c2cef..108e2fc 100644
--- a/gcc/fortran/ChangeLog
+++ b/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 a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 13e730f..e13b4d4 100644
--- a/gcc/fortran/gfortran.h
+++ b/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 a/gcc/fortran/match.h b/gcc/fortran/match.h
index 1b51a88..a52c189 100644
--- a/gcc/fortran/match.h
+++ b/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 a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index e59139c..929a739 100644
--- a/gcc/fortran/openmp.c
+++ b/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 a/gcc/fortran/parse.c b/gcc/fortran/parse.c
index 650135b..b98dda1 100644
--- a/gcc/fortran/parse.c
+++ b/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 a/gcc/fortran/resolve.c b/gcc/fortran/resolve.c
index 1049c0c..bf2837c 100644
--- a/gcc/fortran/resolve.c
+++ b/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 a/gcc/fortran/st.c b/gcc/fortran/st.c
index 116af15..629b51d 100644
--- a/gcc/fortran/st.c
+++ b/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 a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 7e01e72..5f4c382 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/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 a/gcc/fortran/trans.c b/gcc/fortran/trans.c
index 4337fcb..9495450 100644
--- a/gcc/fortran/trans.c
+++ b/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 a/gcc/omp-low.c b/gcc/omp-low.c
index d0264e9..ccf0b63 100644
--- a/gcc/omp-low.c
+++ b/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 a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 1de0ea1..c1ac2b0 100644
--- a/gcc/testsuite/ChangeLog
+++ b/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 a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
index 1c17818..dabba8c 100644
--- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
+++ b/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 a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
index 0c8ea54..e98258c 100644
--- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
+++ b/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" } */
;
}
generated by cgit v1.1 at 2024-10-03 03:16:11 +0000