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.

4 files changed, 110 insertions, 10 deletions

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index fc97aa3..dedfb84 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 e7bfbf6..f030cad 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 0000000..4d18f23
--- /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 0000000..ff3b25e
--- /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;
+}