OpenACC: Add 'if' clause to 'acc wait' directive

OpenACC 3.0 added the 'if' clause to four directives; this patch only adds
it to 'acc wait'.

gcc/c-family/ChangeLog:

	* c-omp.cc (c_finish_oacc_wait): Handle if clause.

gcc/c/ChangeLog:

	* c-parser.cc (OACC_WAIT_CLAUSE_MASK): Add if clause.

gcc/cp/ChangeLog:

	* parser.cc (OACC_WAIT_CLAUSE_MASK): Ass if clause.

gcc/fortran/ChangeLog:

	* openmp.cc (OACC_WAIT_CLAUSES): Add if clause.
	* trans-openmp.cc (gfc_trans_oacc_wait_directive): Handle it.

gcc/testsuite/ChangeLog:

	* c-c++-common/goacc/acc-wait-1.c: New test.
	* gfortran.dg/goacc/acc-wait-1.f90: New test.

7 files changed, 114 insertions, 5 deletions

diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 13de2fe..4352214 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -52,8 +52,8 @@ c_finish_oacc_wait (location_t loc, tree parms, tree clauses)
   vec_alloc (args, nparms + 2);
   stmt = builtin_decl_explicit (BUILT_IN_GOACC_WAIT);
 
-  if (omp_find_clause (clauses, OMP_CLAUSE_ASYNC))
-    t = OMP_CLAUSE_ASYNC_EXPR (clauses);
+  if ((t = omp_find_clause (clauses, OMP_CLAUSE_ASYNC)))
+    t = OMP_CLAUSE_ASYNC_EXPR (t);
   else
     t = build_int_cst (integer_type_node, GOMP_ASYNC_SYNC);
 
@@ -71,6 +71,11 @@ c_finish_oacc_wait (location_t loc, tree parms, tree clauses)
 
   stmt = build_call_expr_loc_vec (loc, stmt, args);
 
+  t = omp_find_clause (clauses, OMP_CLAUSE_IF);
+  if (t)
+    stmt = build3_loc (input_location, COND_EXPR, void_type_node,
+		       OMP_CLAUSE_IF_EXPR (t), stmt, NULL_TREE);
+
   vec_free (args);
 
   return stmt;
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index faa50a4..0c3e3e2 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -22501,7 +22501,8 @@ c_parser_oacc_update (c_parser *parser)
 */
 
 #define OACC_WAIT_CLAUSE_MASK						\
-	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) )
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) )
 
 static tree
 c_parser_oacc_wait (location_t loc, c_parser *parser, char *p_name)
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index cfebde8..80fd799 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -49704,7 +49704,8 @@ cp_parser_oacc_update (cp_parser *parser, cp_token *pragma_tok)
 */
 
 #define OACC_WAIT_CLAUSE_MASK					\
-	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC))
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF))
 
 static tree
 cp_parser_oacc_wait (cp_parser *parser, cp_token *pragma_tok)
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index df82940..fe0a47a 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -4474,7 +4474,7 @@ error:
    | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE | OMP_CLAUSE_FINALIZE	      \
    | OMP_CLAUSE_DETACH)
 #define OACC_WAIT_CLAUSES \
-  omp_mask (OMP_CLAUSE_ASYNC)
+  omp_mask (OMP_CLAUSE_ASYNC) | OMP_CLAUSE_IF
 #define OACC_ROUTINE_CLAUSES \
   (omp_mask (OMP_CLAUSE_GANG) | OMP_CLAUSE_WORKER | OMP_CLAUSE_VECTOR	      \
    | OMP_CLAUSE_SEQ							      \
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 2a48d4a..a2e70fc 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -6048,6 +6048,10 @@ gfc_trans_oacc_wait_directive (gfc_code *code)
     args->quick_push (gfc_convert_expr_to_tree (&block, el->expr));
 
   stmt = build_call_expr_loc_vec (loc, stmt, args);
+  if (clauses->if_expr)
+    stmt = build3_loc (input_location, COND_EXPR, void_type_node,
+		       gfc_convert_expr_to_tree (&block, clauses->if_expr),
+		       stmt, NULL_TREE);
   gfc_add_expr_to_block (&block, stmt);
 
   vec_free (args);
diff --git a/gcc/testsuite/c-c++-common/goacc/acc-wait-1.c b/gcc/testsuite/c-c++-common/goacc/acc-wait-1.c
new file mode 100644
index 0000000..bc7ff02
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/acc-wait-1.c
@@ -0,0 +1,51 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-original" }
+
+void f0 (int x)
+{
+  #pragma acc wait(x) if(false) async
+}
+// { dg-final { scan-tree-dump-times "if \\(0\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(-1, 1, x\\);" 1 "original" { target c } } }
+// { dg-final { scan-tree-dump-not "__builtin_GOACC_wait \\(-1, 1, x\\);" "original" { target c++ } } }
+
+
+void f1 (int y, int ia)
+{
+  #pragma acc wait(y) if(true) async(ia)
+}
+// { dg-final { scan-tree-dump-times "if \\(1\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(ia, 1, y\\);" 1 "original" { target c } } }
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_wait \\(ia, 1, y\\)" 1 "original" { target c++ } } }
+// { dg-final { scan-tree-dump-not "if \\(\[^\\n\\r\]+\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(ia, 1, y\\);" "original" { target c++ } } }
+
+
+void fl (int z, bool ll)
+{
+  #pragma acc wait(z) if(ll) async(3)
+}
+// { dg-final { scan-tree-dump-times "if \\(ll != 0\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(3, 1, z\\);" 1 "original" { target c } } }
+// { dg-final { scan-tree-dump-times "if \\(ll\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(3, 1, z\\);" 1 "original" { target c++ } } }
+
+
+void a0 (int a)
+{
+  #pragma acc wait(a) if(false)
+}
+// { dg-final { scan-tree-dump-times "if \\(0\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(-2, 1, a\\);" 1 "original" { target c } } }
+// { dg-final { scan-tree-dump-not "__builtin_GOACC_wait \\(-2, 1, a\\);" "original" { target c++ } } }
+
+
+void a1 (int b)
+{
+  #pragma acc wait(b) if(true)
+}
+// { dg-final { scan-tree-dump-times "if \\(1\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(-2, 1, b\\);" 1 "original" { target c } } }
+// { dg-final { scan-tree-dump-times "__builtin_GOACC_wait \\(-2, 1, b\\)" 1 "original" { target c++ } } }
+// { dg-final { scan-tree-dump-not "if \\(\[^\\n\\r\]+\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(-2, 1, b\\);" "original" { target c++ } } }
+
+
+void al (int c, bool qq)
+{
+  #pragma acc wait(c) if(qq)
+}
+// { dg-final { scan-tree-dump-times "if \\(qq != 0\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(-2, 1, c\\);" 1 "original" { target c } } }
+// { dg-final { scan-tree-dump-times "if \\(qq\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(-2, 1, c\\);" 1 "original" { target c++ } } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/acc-wait-1.f90 b/gcc/testsuite/gfortran.dg/goacc/acc-wait-1.f90
new file mode 100644
index 0000000..dd19b41
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/acc-wait-1.f90
@@ -0,0 +1,47 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+
+subroutine f0(x)
+  implicit none
+  integer, value :: x
+  !$acc wait(x) if(.false.) async
+end
+
+subroutine f1(y, ia)
+  implicit none
+  integer, value :: y, ia
+  !$acc wait(y) if(.true.) async(ia)
+end
+
+subroutine fl(z, ll)
+  implicit none
+  integer, value :: z
+  logical, value :: ll
+  !$acc wait(z) if(ll) async(3)
+end
+
+subroutine a0(a)
+  implicit none
+  integer, value :: a
+  !$acc wait(a) if(.false.)
+end
+
+subroutine a1(b)
+  implicit none
+  integer, value :: b
+  !$acc wait(b) if(.true.)
+end
+
+subroutine al(c, qq)
+  implicit none
+  integer, value :: c
+  logical, value :: qq
+  !$acc wait(c) if(qq)
+end
+
+! { dg-final { scan-tree-dump-times "D\.\[0-9\]+ = x;\[\\n\\r\]+ *if \\(0\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(-1, 1, D\.\[0-9\]+\\);" 1 "original" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9\]+ = ia;\[\\n\\r\]+ *D\.\[0-9\]+ = y;\[\\n\\r\]+ *if \\(1\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(D\.\[0-9\]+, 1, D\.\[0-9\]+\\);" 1 "original" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9\]+ = z;\[\\n\\r\]+ *D\.\[0-9\]+ = ll;\[\\n\\r\]+ *if \\(D\.\[0-9\]+\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(3, 1, D\.\[0-9\]+\\);" 1 "original" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9\]+ = a;\[\\n\\r\]+ *if \\(0\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(-2, 1, D\.\[0-9\]+\\);" 1 "original" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9\]+ = b;\[\\n\\r\]+ *if \\(1\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(-2, 1, D\.\[0-9\]+\\);" 1 "original" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9\]+ = c;\[\\n\\r\]+ *D\.\[0-9\]+ = qq;\[\\n\\r\]+ *if \\(D\.\[0-9\]+\\)\[\\n\\r\]+ *\{\[\\n\\r\]+ *__builtin_GOACC_wait \\(-2, 1, D\.\[0-9\]+\\);" 1 "original" } }