* [Patch] OpenACC (C/C++): Fix 'acc atomic' parsing
@ 2020-11-05 12:03 Tobias Burnus
2020-11-05 12:13 ` Jakub Jelinek
0 siblings, 1 reply; 3+ messages in thread
From: Tobias Burnus @ 2020-11-05 12:03 UTC (permalink / raw)
To: gcc-patches, Thomas Schwinge; +Cc: Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 812 bytes --]
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
[-- Attachment #2: acc-atomic.diff --]
[-- Type: text/x-patch, Size: 11386 bytes --]
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;
+}
^ permalink raw reply [flat|nested] 3+ messages in thread
* Re: [Patch] OpenACC (C/C++): Fix 'acc atomic' parsing
2020-11-05 12:03 [Patch] OpenACC (C/C++): Fix 'acc atomic' parsing Tobias Burnus
@ 2020-11-05 12:13 ` Jakub Jelinek
2020-11-06 10:43 ` Tobias Burnus
0 siblings, 1 reply; 3+ messages in thread
From: Jakub Jelinek @ 2020-11-05 12:13 UTC (permalink / raw)
To: Tobias Burnus; +Cc: gcc-patches, Thomas Schwinge
On Thu, Nov 05, 2020 at 01:03:38PM +0100, Tobias Burnus wrote:
> 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?
> 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");
Wouldn't it be much simpler and more readable to do:
else if (!strcmp (p, "capture"))
new_code = OMP_ATOMIC_CAPTURE_NEW;
+ else if (openacc)
+ {
+ p = NULL;
+ error_at (cloc, "expected %<read%>, %<write%>, %<update%>, "
+ "or %<capture%> clause");
+ }
else if (!strcmp (p, "seq_cst"))
new_memory_order = OMP_MEMORY_ORDER_SEQ_CST;
... - handling of other openmp only clauses here
else
{
p = NULL;
error_at (cloc, "expected %<read%>, %<write%>, %<update%>, "
"%<capture%>, %<seq_cst%>, %<acq_rel%>, "
"%<release%>, %<relaxed%> or %<hint%> clause");
}
?
Ditto C++.
Otherwise LGTM, but I have no idea what OpenACC actually says...
Jakub
^ permalink raw reply [flat|nested] 3+ messages in thread
* Re: [Patch] OpenACC (C/C++): Fix 'acc atomic' parsing
2020-11-05 12:13 ` Jakub Jelinek
@ 2020-11-06 10:43 ` Tobias Burnus
0 siblings, 0 replies; 3+ messages in thread
From: Tobias Burnus @ 2020-11-06 10:43 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: gcc-patches, Thomas Schwinge
[-- Attachment #1: Type: text/plain, Size: 1325 bytes --]
On 05.11.20 13:13, Jakub Jelinek wrote:
> Wouldn't it be much simpler and more readable to do:
> else if (!strcmp (p, "capture"))
> new_code = OMP_ATOMIC_CAPTURE_NEW;
> + else if (openacc)
> + {
> + p = NULL;
> + error_at (cloc, "expected %<read%>, %<write%>, %<update%>, "
> + "or %<capture%> clause");
> + }
Thanks for looking through the patch – and the suggestion.
It is simpler – and also avoids issues when the OpenMP adds more clauses.
> Otherwise LGTM, but I have no idea what OpenACC actually says...
I have now installed it as commit r11-4774-ga2c11935b010ee55f7ccd14d27f62c6fbed3745e.
Regarding OpenACC, permitted are (since 2.5):
"|#pragma acc atomic [atomic-clause] new-line ||#pragma acc atomic update capture new-line||"Where atomic-clause is one of read, write, update, or capture." I did
note that I misread the spec – 'update capture' is not permitted for
Fortran, only for C/C++. (This is now OpenACC spec Issue #333, which
also asks to better specify what 'update capture' means.)|
Tobias
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
[-- Attachment #2: committed.diff --]
[-- Type: text/x-patch, Size: 9383 bytes --]
commit a2c11935b010ee55f7ccd14d27f62c6fbed3745e
Author: Tobias Burnus <tobias@codesourcery.com>
Date: Fri Nov 6 11:13:47 2020 +0100
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 | 24 +++++++++++---
gcc/cp/parser.c | 23 +++++++++++---
gcc/testsuite/c-c++-common/goacc-gomp/atomic.c | 43 ++++++++++++++++++++++++++
gcc/testsuite/c-c++-common/goacc/atomic.c | 30 ++++++++++++++++++
4 files changed, 110 insertions(+), 10 deletions(-)
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index fc97aa3f95f..dedfb8472d0 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,6 +17343,12 @@ 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 (openacc)
+ {
+ p = NULL;
+ error_at (cloc, "expected %<read%>, %<write%>, %<update%>, "
+ "or %<capture%> clause");
+ }
else if (!strcmp (p, "seq_cst"))
new_memory_order = OMP_MEMORY_ORDER_SEQ_CST;
else if (!strcmp (p, "acq_rel"))
@@ -17370,7 +17376,12 @@ c_parser_omp_atomic (location_t loc, c_parser *parser)
{
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 +17403,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 +17461,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 +21503,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 +21530,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 e7bfbf649a5..f030cad18b2 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -37994,7 +37994,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;
@@ -38029,6 +38029,12 @@ 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 (openacc)
+ {
+ p = NULL;
+ error_at (cloc, "expected %<read%>, %<write%>, %<update%>, "
+ "or %<capture%> clause");
+ }
else if (!strcmp (p, "seq_cst"))
new_memory_order = OMP_MEMORY_ORDER_SEQ_CST;
else if (!strcmp (p, "acq_rel"))
@@ -38056,7 +38062,12 @@ cp_parser_omp_atomic (cp_parser *parser, cp_token *pragma_tok)
{
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;
@@ -38078,7 +38089,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
@@ -43517,7 +43530,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);
@@ -43552,7 +43565,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;
+}
^ permalink raw reply [flat|nested] 3+ messages in thread
end of thread, other threads:[~2020-11-06 10:43 UTC | newest]
Thread overview: 3+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2020-11-05 12:03 [Patch] OpenACC (C/C++): Fix 'acc atomic' parsing Tobias Burnus
2020-11-05 12:13 ` Jakub Jelinek
2020-11-06 10:43 ` Tobias Burnus
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).