https://gcc.gnu.org/g:2b077252cafa5045498a0e0c480ee6d48c136232
commit r16-1634-g2b077252cafa5045498a0e0c480ee6d48c136232
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.
Diff:
---
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 13de2fe48f96..4352214df3b7 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 faa50a4fd86b..0c3e3e2889c6 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 cfebde8b1181..80fd7990bbbb 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 df829403c34f..fe0a47a6948b 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 2a48d4af5276..a2e70fca0b37 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 000000000000..bc7ff022f173
--- /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 000000000000..dd19b4156637
--- /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" } }
