OpenACC piggybacks on OpenACC for the atomic parsing; however, there
are two issues:
* Newer OpenMP versions added additional clauses such as 'seq_cst',
which do not exist in OpenACC.
* OpenACC 2.6 added 'acc atomic update capture' (besides 'acc atomic capture',
which was not accepted.
Actually, while OpenACC 2.6/2.7/3.0 has 'acc atomic update capture' in the
syntax, it never explicitly states that this matches 'atomic capture'.
NOTE: I did not check whether the supported expressions by OpenMP 5.0/the
current GCC implementation is the same as in OpenACC 2.6/2.7/3.x.
Any comments?
Tobias
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander
Walter
OpenACC (C/C++): Fix 'acc atomic' parsing
gcc/c/ChangeLog:
* c-parser.c (c_parser_omp_atomic): Add openacc parameter and update
OpenACC matching.
(c_parser_omp_construct): Update call.
gcc/cp/ChangeLog:
* parser.c (cp_parser_omp_atomic): Add openacc parameter and update
OpenACC matching.
(cp_parser_omp_construct): Update call.
gcc/testsuite/ChangeLog:
* c-c++-common/goacc-gomp/atomic.c: New test.
* c-c++-common/goacc/atomic.c: New test.
gcc/c/c-parser.c | 40 +++++++++++++++---------
gcc/cp/parser.c | 39 ++++++++++++++---------
gcc/testsuite/c-c++-common/goacc-gomp/atomic.c | 43 ++++++++++++++++++++++++++
gcc/testsuite/c-c++-common/goacc/atomic.c | 30 ++++++++++++++++++
4 files changed, 124 insertions(+), 28 deletions(-)
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index fc97aa3f95f..79037d98f76 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -17304,7 +17304,7 @@ c_parser_oacc_wait (location_t loc, c_parser *parser, char *p_name)
LOC is the location of the #pragma token. */
static void
-c_parser_omp_atomic (location_t loc, c_parser *parser)
+c_parser_omp_atomic (location_t loc, c_parser *parser, bool openacc)
{
tree lhs = NULL_TREE, rhs = NULL_TREE, v = NULL_TREE;
tree lhs1 = NULL_TREE, rhs1 = NULL_TREE;
@@ -17343,17 +17343,17 @@ c_parser_omp_atomic (location_t loc, c_parser *parser)
new_code = OMP_ATOMIC;
else if (!strcmp (p, "capture"))
new_code = OMP_ATOMIC_CAPTURE_NEW;
- else if (!strcmp (p, "seq_cst"))
+ else if (!openacc && !strcmp (p, "seq_cst"))
new_memory_order = OMP_MEMORY_ORDER_SEQ_CST;
- else if (!strcmp (p, "acq_rel"))
+ else if (!openacc && !strcmp (p, "acq_rel"))
new_memory_order = OMP_MEMORY_ORDER_ACQ_REL;
- else if (!strcmp (p, "release"))
+ else if (!openacc && !strcmp (p, "release"))
new_memory_order = OMP_MEMORY_ORDER_RELEASE;
- else if (!strcmp (p, "acquire"))
+ else if (!openacc && !strcmp (p, "acquire"))
new_memory_order = OMP_MEMORY_ORDER_ACQUIRE;
- else if (!strcmp (p, "relaxed"))
+ else if (!openacc && !strcmp (p, "relaxed"))
new_memory_order = OMP_MEMORY_ORDER_RELAXED;
- else if (!strcmp (p, "hint"))
+ else if (!openacc && !strcmp (p, "hint"))
{
c_parser_consume_token (parser);
clauses = c_parser_omp_clause_hint (parser, clauses);
@@ -17362,15 +17362,24 @@ c_parser_omp_atomic (location_t loc, c_parser *parser)
else
{
p = NULL;
- error_at (cloc, "expected %<read%>, %<write%>, %<update%>, "
- "%<capture%>, %<seq_cst%>, %<acq_rel%>, "
- "%<release%>, %<relaxed%> or %<hint%> clause");
+ if (openacc)
+ error_at (cloc, "expected %<read%>, %<write%>, %<update%>, "
+ "or %<capture%> clause");
+ else
+ error_at (cloc, "expected %<read%>, %<write%>, %<update%>, "
+ "%<capture%>, %<seq_cst%>, %<acq_rel%>, "
+ "%<release%>, %<relaxed%> or %<hint%> clause");
}
if (p)
{
if (new_code != ERROR_MARK)
{
- if (code != ERROR_MARK)
+ /* OpenACC permits 'update capture'. */
+ if (openacc
+ && code == OMP_ATOMIC
+ && new_code == OMP_ATOMIC_CAPTURE_NEW)
+ code = new_code;
+ else if (code != ERROR_MARK)
error_at (cloc, "too many atomic clauses");
else
code = new_code;
@@ -17392,7 +17401,9 @@ c_parser_omp_atomic (location_t loc, c_parser *parser)
if (code == ERROR_MARK)
code = OMP_ATOMIC;
- if (memory_order == OMP_MEMORY_ORDER_UNSPECIFIED)
+ if (openacc)
+ memory_order = OMP_MEMORY_ORDER_RELAXED;
+ else if (memory_order == OMP_MEMORY_ORDER_UNSPECIFIED)
{
omp_requires_mask
= (enum omp_requires) (omp_requires_mask
@@ -17448,6 +17459,7 @@ c_parser_omp_atomic (location_t loc, c_parser *parser)
}
break;
case OMP_ATOMIC:
+ /* case OMP_ATOMIC_CAPTURE_NEW: - or update to OpenMP 5.1 */
if (memory_order == OMP_MEMORY_ORDER_ACQ_REL
|| memory_order == OMP_MEMORY_ORDER_ACQUIRE)
{
@@ -21489,7 +21501,7 @@ c_parser_omp_construct (c_parser *parser, bool *if_p)
switch (p_kind)
{
case PRAGMA_OACC_ATOMIC:
- c_parser_omp_atomic (loc, parser);
+ c_parser_omp_atomic (loc, parser, true);
return;
case PRAGMA_OACC_CACHE:
strcpy (p_name, "#pragma acc");
@@ -21516,7 +21528,7 @@ c_parser_omp_construct (c_parser *parser, bool *if_p)
stmt = c_parser_oacc_wait (loc, parser, p_name);
break;
case PRAGMA_OMP_ATOMIC:
- c_parser_omp_atomic (loc, parser);
+ c_parser_omp_atomic (loc, parser, false);
return;
case PRAGMA_OMP_CRITICAL:
stmt = c_parser_omp_critical (loc, parser, if_p);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index b0d5c69f1d6..1d4f353d947 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -37851,7 +37851,7 @@ cp_parser_omp_structured_block (cp_parser *parser, bool *if_p)
where x and v are lvalue expressions with scalar type. */
static void
-cp_parser_omp_atomic (cp_parser *parser, cp_token *pragma_tok)
+cp_parser_omp_atomic (cp_parser *parser, cp_token *pragma_tok, bool openacc)
{
tree lhs = NULL_TREE, rhs = NULL_TREE, v = NULL_TREE, lhs1 = NULL_TREE;
tree rhs1 = NULL_TREE, orig_lhs;
@@ -37886,17 +37886,17 @@ cp_parser_omp_atomic (cp_parser *parser, cp_token *pragma_tok)
new_code = OMP_ATOMIC;
else if (!strcmp (p, "capture"))
new_code = OMP_ATOMIC_CAPTURE_NEW;
- else if (!strcmp (p, "seq_cst"))
+ else if (!openacc && !strcmp (p, "seq_cst"))
new_memory_order = OMP_MEMORY_ORDER_SEQ_CST;
- else if (!strcmp (p, "acq_rel"))
+ else if (!openacc && !strcmp (p, "acq_rel"))
new_memory_order = OMP_MEMORY_ORDER_ACQ_REL;
- else if (!strcmp (p, "release"))
+ else if (!openacc && !strcmp (p, "release"))
new_memory_order = OMP_MEMORY_ORDER_RELEASE;
- else if (!strcmp (p, "acquire"))
+ else if (!openacc && !strcmp (p, "acquire"))
new_memory_order = OMP_MEMORY_ORDER_ACQUIRE;
- else if (!strcmp (p, "relaxed"))
+ else if (!openacc && !strcmp (p, "relaxed"))
new_memory_order = OMP_MEMORY_ORDER_RELAXED;
- else if (!strcmp (p, "hint"))
+ else if (!openacc && !strcmp (p, "hint"))
{
cp_lexer_consume_token (parser->lexer);
clauses = cp_parser_omp_clause_hint (parser, clauses, cloc);
@@ -37905,15 +37905,24 @@ cp_parser_omp_atomic (cp_parser *parser, cp_token *pragma_tok)
else
{
p = NULL;
- error_at (cloc, "expected %<read%>, %<write%>, %<update%>, "
- "%<capture%>, %<seq_cst%>, %<acq_rel%>, "
- "%<release%>, %<relaxed%> or %<hint%> clause");
+ if (openacc)
+ error_at (cloc, "expected %<read%>, %<write%>, %<update%>, "
+ "or %<capture%> clause");
+ else
+ error_at (cloc, "expected %<read%>, %<write%>, %<update%>, "
+ "%<capture%>, %<seq_cst%>, %<acq_rel%>, "
+ "%<release%>, %<relaxed%> or %<hint%> clause");
}
if (p)
{
if (new_code != ERROR_MARK)
{
- if (code != ERROR_MARK)
+ /* OpenACC permits 'update capture'. */
+ if (openacc
+ && code == OMP_ATOMIC
+ && new_code == OMP_ATOMIC_CAPTURE_NEW)
+ code = new_code;
+ else if (code != ERROR_MARK)
error_at (cloc, "too many atomic clauses");
else
code = new_code;
@@ -37935,7 +37944,9 @@ cp_parser_omp_atomic (cp_parser *parser, cp_token *pragma_tok)
if (code == ERROR_MARK)
code = OMP_ATOMIC;
- if (memory_order == OMP_MEMORY_ORDER_UNSPECIFIED)
+ if (openacc)
+ memory_order = OMP_MEMORY_ORDER_RELAXED;
+ else if (memory_order == OMP_MEMORY_ORDER_UNSPECIFIED)
{
omp_requires_mask
= (enum omp_requires) (omp_requires_mask
@@ -43374,7 +43385,7 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
switch (cp_parser_pragma_kind (pragma_tok))
{
case PRAGMA_OACC_ATOMIC:
- cp_parser_omp_atomic (parser, pragma_tok);
+ cp_parser_omp_atomic (parser, pragma_tok, true);
return;
case PRAGMA_OACC_CACHE:
stmt = cp_parser_oacc_cache (parser, pragma_tok);
@@ -43409,7 +43420,7 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
stmt = cp_parser_oacc_wait (parser, pragma_tok);
break;
case PRAGMA_OMP_ATOMIC:
- cp_parser_omp_atomic (parser, pragma_tok);
+ cp_parser_omp_atomic (parser, pragma_tok, false);
return;
case PRAGMA_OMP_CRITICAL:
stmt = cp_parser_omp_critical (parser, pragma_tok, if_p);
diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/atomic.c b/gcc/testsuite/c-c++-common/goacc-gomp/atomic.c
new file mode 100644
index 00000000000..4d18f238f3b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc-gomp/atomic.c
@@ -0,0 +1,43 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+#pragma omp requires atomic_default_mem_order(acq_rel)
+
+void
+foo ()
+{
+ int i, v;
+
+#pragma omp atomic read
+ i = v;
+
+#pragma acc atomic read
+ i = v;
+
+#pragma omp atomic write
+ i = v;
+
+#pragma acc atomic write
+ i = v;
+
+#pragma omp atomic update
+ i += 1;
+
+#pragma acc atomic update
+ i += 1;
+
+#pragma omp atomic capture
+ v = i += 1;
+
+#pragma acc atomic capture
+ v = i += 1;
+#pragma acc atomic update capture
+ v = i += 1;
+}
+
+/* { dg-final { scan-tree-dump-times "i = #pragma omp atomic read acquire" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "i = #pragma omp atomic read relaxed" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp atomic release" 2 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp atomic relaxed" 2 "original" } } */
+/* { dg-final { scan-tree-dump-times "v = #pragma omp atomic capture acq_rel" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "v = #pragma omp atomic capture relaxed" 2 "original" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/atomic.c b/gcc/testsuite/c-c++-common/goacc/atomic.c
new file mode 100644
index 00000000000..ff3b25e4b37
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/atomic.c
@@ -0,0 +1,30 @@
+/* { dg-do compile } */
+
+void
+foo ()
+{
+ int i, v;
+#pragma acc atomic read bar /* { dg-error "expected 'read', 'write', 'update', or 'capture' clause" } */
+ i = v; /* { dg-error "expected end of line before 'bar'" "" { target *-*-* } .-1 } */
+
+#pragma acc atomic read write /* { dg-error "too many atomic clauses" } */
+ i = v;
+
+#pragma acc atomic read seq_cst /* { dg-error "expected 'read', 'write', 'update', or 'capture' clause" } */
+ i = v; /* { dg-error "expected end of line before 'seq_cst'" "" { target *-*-* } .-1 } */
+
+#pragma acc atomic read relaxed /* { dg-error "expected 'read', 'write', 'update', or 'capture' clause" } */
+ i = v; /* { dg-error "expected end of line before 'relaxed'" "" { target *-*-* } .-1 } */
+
+#pragma acc atomic update hint(1) /* { dg-error "expected 'read', 'write', 'update', or 'capture' clause" } */
+ i += 1; /* { dg-error "expected end of line before 'hint'" "" { target *-*-* } .-1 } */
+
+#pragma acc atomic update update capture /* { dg-error "too many atomic clauses" } */
+ v = i += 1;
+
+#pragma acc atomic update capture capture /* { dg-error "too many atomic clauses" } */
+ v = i += 1;
+
+#pragma acc atomic write capture /* { dg-error "too many atomic clauses" } */
+ i = 1;
+}
