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.

diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 13de2fe48f96d52a937d147288f21da7012baf1f..4352214df3b7cb4b83a9667669b4ee28186da79c 100644 (file)
--- 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 faa50a4fd86b76733ec4e5d04bf96573c975b977..0c3e3e2889c60c08d91569d41b324ffe0e1b4c7c 100644 (file)
--- 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 cfebde8b118149accdaec8e32333e08a1310e3bd..80fd7990bbbb840773cefba28329b6a7c50874a3 100644 (file)
--- 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 df829403c34f138b8e80526485f80e4b15853bb4..fe0a47a6948beb55fa653c6bce4f4d894aba82c0 100644 (file)
--- 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 2a48d4af5276706b03c92dcbd5cab2e199c81147..a2e70fca0b373ab730ba0e43b3cc306c63144b83 100644 (file)
--- 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 (file)
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 (file)
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" } }