OpenACC (C/C++): Fix 'acc atomic' parsing

Message ID 404f6b06-5f90-a1f6-7720-b185b0e918b4@mentor.com
State New
Headers show
Series
  • OpenACC (C/C++): Fix 'acc atomic' parsing
Related show

Commit Message

Tobias Burnus Nov. 5, 2020, 12:03 p.m.
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

Comments

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
Tobias Burnus Nov. 6, 2020, 10:43 a.m. | #2
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
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;
+}

Patch

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;
+}