This simple patch adds the 'if' clause to 'acc wait'. while OpenACC 3.0
also added it to the init, shutdown and set directive this patch only
handles 'acc wait' as that's used in the ICON code.
Committed as Rev. r16-1634-g2b077252cafa50
Tobias
commit 2b077252cafa5045498a0e0c480ee6d48c136232
Author: Tobias Burnus <tbur...@baylibre.com>
Date: Mon Jun 23 23:24:56 2025 +0200
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.
---
gcc/c-family/c-omp.cc | 9 ++++-
gcc/c/c-parser.cc | 3 +-
gcc/cp/parser.cc | 3 +-
gcc/fortran/openmp.cc | 2 +-
gcc/fortran/trans-openmp.cc | 4 ++
gcc/testsuite/c-c++-common/goacc/acc-wait-1.c | 51 ++++++++++++++++++++++++++
gcc/testsuite/gfortran.dg/goacc/acc-wait-1.f90 | 47 ++++++++++++++++++++++++
7 files changed, 114 insertions(+), 5 deletions(-)
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 13de2fe48f9..4352214df3b 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 faa50a4fd86..0c3e3e2889c 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 cfebde8b118..80fd7990bbb 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 df829403c34..fe0a47a6948 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 2a48d4af527..a2e70fca0b3 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 00000000000..bc7ff022f17
--- /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 00000000000..dd19b415663
--- /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" } }
