* [PATCH 1/7] OpenMP: dispatch + adjust_args tree data structures and front-end interfaces
2024-05-27 11:54 [PATCH 0/7] OpenMP: dispatch + adjust_args support Paul-Antoine Arras
@ 2024-05-27 11:54 ` Paul-Antoine Arras
2024-05-27 11:54 ` [PATCH 2/7] OpenMP: middle-end support for dispatch + adjust_args Paul-Antoine Arras
` (5 subsequent siblings)
6 siblings, 0 replies; 10+ messages in thread
From: Paul-Antoine Arras @ 2024-05-27 11:54 UTC (permalink / raw)
To: gcc-patches; +Cc: Paul-Antoine Arras
This patch introduces the OMP_DISPATCH tree node, as well as two new clauses
`nocontext` and `novariants`. It defines/exposes interfaces that will be
used in subsequent patches that add front-end and middle-end support, but
nothing generates these nodes yet.
It also adds support for new OpenMP context selectors: `dispatch` as trait
selector and `need_device_ptr` as pseudo-trait set selector. The purpose of the
latter is for the C++ front-end to store the list of arguments (that need to be
converted to device pointers) until the declaration of the variant function
becomes available.
gcc/ChangeLog:
* builtin-types.def (BT_FN_PTR_CONST_PTR_INT): New.
* omp-selectors.h (enum omp_tss_code): Add
OMP_TRAIT_SET_NEED_DEVICE_PTR.
(enum omp_ts_code): Add OMP_TRAIT_CONSTRUCT_DISPATCH.
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_NOVARIANTS and
OMP_CLAUSE_NOCONTEXT.
* tree-pretty-print.cc (dump_omp_clause): Handle OMP_CLAUSE_NOVARIANTS
and OMP_CLAUSE_NOCONTEXT.
(dump_generic_node): Handle OMP_DISPATCH.
* tree.cc (omp_clause_num_ops): Add OMP_CLAUSE_NOVARIANTS and
OMP_CLAUSE_NOCONTEXT.
(omp_clause_code_name): Add "novariants" and "nocontext".
* tree.def (OMP_DISPATCH): New.
* tree.h (OMP_DISPATCH_BODY): New macro.
(OMP_DISPATCH_CLAUSES): New macro.
(OMP_CLAUSE_NOVARIANTS_EXPR): New macro.
(OMP_CLAUSE_NOCONTEXT_EXPR): New macro.
---
gcc/builtin-types.def | 1 +
gcc/omp-selectors.h | 3 +++
gcc/tree-core.h | 7 +++++++
gcc/tree-pretty-print.cc | 21 +++++++++++++++++++++
gcc/tree.cc | 4 ++++
gcc/tree.def | 5 +++++
gcc/tree.h | 7 +++++++
7 files changed, 48 insertions(+)
diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index c97d6bad1de..ef7aaf67d13 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -677,6 +677,7 @@ DEF_FUNCTION_TYPE_2 (BT_FN_INT_FEXCEPT_T_PTR_INT, BT_INT, BT_FEXCEPT_T_PTR,
DEF_FUNCTION_TYPE_2 (BT_FN_INT_CONST_FEXCEPT_T_PTR_INT, BT_INT,
BT_CONST_FEXCEPT_T_PTR, BT_INT)
DEF_FUNCTION_TYPE_2 (BT_FN_PTR_CONST_PTR_UINT8, BT_PTR, BT_CONST_PTR, BT_UINT8)
+DEF_FUNCTION_TYPE_2 (BT_FN_PTR_CONST_PTR_INT, BT_PTR, BT_CONST_PTR, BT_INT)
DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR_PTR, BT_FN_VOID_PTR_PTR)
diff --git a/gcc/omp-selectors.h b/gcc/omp-selectors.h
index c61808ec0ad..12bc9e9afa0 100644
--- a/gcc/omp-selectors.h
+++ b/gcc/omp-selectors.h
@@ -31,6 +31,8 @@ enum omp_tss_code {
OMP_TRAIT_SET_TARGET_DEVICE,
OMP_TRAIT_SET_IMPLEMENTATION,
OMP_TRAIT_SET_USER,
+ OMP_TRAIT_SET_NEED_DEVICE_PTR, // pseudo-set selector used to convey argument
+ // list until variant has a decl
OMP_TRAIT_SET_LAST,
OMP_TRAIT_SET_INVALID = -1
};
@@ -55,6 +57,7 @@ enum omp_ts_code {
OMP_TRAIT_CONSTRUCT_PARALLEL,
OMP_TRAIT_CONSTRUCT_FOR,
OMP_TRAIT_CONSTRUCT_SIMD,
+ OMP_TRAIT_CONSTRUCT_DISPATCH,
OMP_TRAIT_LAST,
OMP_TRAIT_INVALID = -1
};
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 9fa74342919..ed6ffdab87f 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -533,6 +533,13 @@ enum omp_clause_code {
/* OpenACC clause: nohost. */
OMP_CLAUSE_NOHOST,
+
+ /* OpenMP clause: novariants (scalar-expression). */
+ OMP_CLAUSE_NOVARIANTS,
+
+ /* OpenMP clause: nocontext (scalar-expression). */
+ OMP_CLAUSE_NOCONTEXT,
+
};
#undef DEFTREESTRUCT
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index f9ad8562078..bbae3a98e9a 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -506,6 +506,22 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case OMP_CLAUSE_EXCLUSIVE:
name = "exclusive";
goto print_remap;
+ case OMP_CLAUSE_NOVARIANTS:
+ pp_string (pp, "novariants");
+ pp_left_paren (pp);
+ gcc_assert (OMP_CLAUSE_NOVARIANTS_EXPR (clause));
+ dump_generic_node (pp, OMP_CLAUSE_NOVARIANTS_EXPR (clause), spc, flags,
+ false);
+ pp_right_paren (pp);
+ break;
+ case OMP_CLAUSE_NOCONTEXT:
+ pp_string (pp, "nocontext");
+ pp_left_paren (pp);
+ gcc_assert (OMP_CLAUSE_NOCONTEXT_EXPR (clause));
+ dump_generic_node (pp, OMP_CLAUSE_NOCONTEXT_EXPR (clause), spc, flags,
+ false);
+ pp_right_paren (pp);
+ break;
case OMP_CLAUSE__LOOPTEMP_:
name = "_looptemp_";
goto print_remap;
@@ -3918,6 +3934,11 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags,
dump_omp_clauses (pp, OMP_SECTIONS_CLAUSES (node), spc, flags);
goto dump_omp_body;
+ case OMP_DISPATCH:
+ pp_string (pp, "#pragma omp dispatch");
+ dump_omp_clauses (pp, OMP_DISPATCH_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
+
case OMP_SECTION:
pp_string (pp, "#pragma omp section");
goto dump_omp_body;
diff --git a/gcc/tree.cc b/gcc/tree.cc
index 6564b002dc1..f0007e24fc4 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -328,6 +328,8 @@ unsigned const char omp_clause_num_ops[] =
0, /* OMP_CLAUSE_IF_PRESENT */
0, /* OMP_CLAUSE_FINALIZE */
0, /* OMP_CLAUSE_NOHOST */
+ 1, /* OMP_CLAUSE_NOVARIANTS */
+ 1, /* OMP_CLAUSE_NOCONTEXT */
};
const char * const omp_clause_code_name[] =
@@ -421,6 +423,8 @@ const char * const omp_clause_code_name[] =
"if_present",
"finalize",
"nohost",
+ "novariants",
+ "nocontext",
};
/* Unless specific to OpenACC, we tend to internally maintain OpenMP-centric
diff --git a/gcc/tree.def b/gcc/tree.def
index 24128e1e039..ff7ca96d29c 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1290,6 +1290,11 @@ DEFTREECODE (OMP_MASKED, "omp_masked", tcc_statement, 2)
Operand 1: OMP_SCAN_CLAUSES: List of clauses. */
DEFTREECODE (OMP_SCAN, "omp_scan", tcc_statement, 2)
+/* OpenMP - #pragma omp dispatch [clause1 ... clauseN]
+ Operand 0: OMP_DISPATCH_BODY: Expression statement including a target call.
+ Operand 1: OMP_DISPATCH_CLAUSES: List of clauses. */
+DEFTREECODE (OMP_DISPATCH, "omp_dispatch", tcc_statement, 2)
+
/* OpenMP - #pragma omp section
Operand 0: OMP_SECTION_BODY: Section body. */
DEFTREECODE (OMP_SECTION, "omp_section", tcc_statement, 1)
diff --git a/gcc/tree.h b/gcc/tree.h
index ee2aae332a4..f23a7366de0 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1599,6 +1599,9 @@ class auto_suppress_location_wrappers
#define OMP_SCAN_BODY(NODE) TREE_OPERAND (OMP_SCAN_CHECK (NODE), 0)
#define OMP_SCAN_CLAUSES(NODE) TREE_OPERAND (OMP_SCAN_CHECK (NODE), 1)
+#define OMP_DISPATCH_BODY(NODE) TREE_OPERAND (OMP_DISPATCH_CHECK (NODE), 0)
+#define OMP_DISPATCH_CLAUSES(NODE) TREE_OPERAND (OMP_DISPATCH_CHECK (NODE), 1)
+
#define OMP_CLAUSE_SIZE(NODE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE), \
OMP_CLAUSE_FROM, \
@@ -1742,6 +1745,10 @@ class auto_suppress_location_wrappers
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_HINT), 0)
#define OMP_CLAUSE_FILTER_EXPR(NODE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_FILTER), 0)
+#define OMP_CLAUSE_NOVARIANTS_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NOVARIANTS), 0)
+#define OMP_CLAUSE_NOCONTEXT_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NOCONTEXT), 0)
#define OMP_CLAUSE_GRAINSIZE_EXPR(NODE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_GRAINSIZE),0)
--
2.45.1
^ permalink raw reply [flat|nested] 10+ messages in thread* [PATCH 2/7] OpenMP: middle-end support for dispatch + adjust_args
2024-05-27 11:54 [PATCH 0/7] OpenMP: dispatch + adjust_args support Paul-Antoine Arras
2024-05-27 11:54 ` [PATCH 1/7] OpenMP: dispatch + adjust_args tree data structures and front-end interfaces Paul-Antoine Arras
@ 2024-05-27 11:54 ` Paul-Antoine Arras
2024-05-27 11:54 ` [PATCH 3/7] OpenMP: C front-end " Paul-Antoine Arras
` (4 subsequent siblings)
6 siblings, 0 replies; 10+ messages in thread
From: Paul-Antoine Arras @ 2024-05-27 11:54 UTC (permalink / raw)
To: gcc-patches; +Cc: Paul-Antoine Arras
This patch adds middle-end support for the `dispatch` construct and the
`adjust_args` clause. The heavy lifting is done in `gimplify_omp_dispatch` and
`gimplify_call_expr` respectively. For `adjust_args`, this mostly consists in
emitting a call to `gomp_get_mapped_ptr` for the adequate device.
For dispatch, the following steps are performed:
* Handle the device clause, if any. This may affect `need_device_ptr` arguments.
* Handle novariants and nocontext clauses, if any. Evaluate compile-time
constants and select a variant, if possible. Otherwise, emit code to handle all
possible cases at run time.
* Create an explicit task, as if the `task` construct was used, that wraps the
body of the `dispatch` statement. Move relevant clauses to the task.
gcc/ChangeLog:
* gimple-low.cc (lower_stmt): Handle GIMPLE_OMP_DISPATCH.
* gimple-pretty-print.cc (dump_gimple_omp_dispatch): New function.
(pp_gimple_stmt_1): Handle GIMPLE_OMP_DISPATCH.
* gimple-walk.cc (walk_gimple_stmt): Likewise.
* gimple.cc (gimple_build_omp_dispatch): New function.
(gimple_copy): Handle GIMPLE_OMP_DISPATCH.
* gimple.def (GIMPLE_OMP_DISPATCH): Define.
* gimple.h (gimple_build_omp_dispatch): Declare.
(gimple_has_substatements): Handle GIMPLE_OMP_DISPATCH.
(gimple_omp_dispatch_clauses): New function.
(gimple_omp_dispatch_clauses_ptr): Likewise.
(gimple_omp_dispatch_set_clauses): Likewise.
(gimple_return_set_retval): Handle GIMPLE_OMP_DISPATCH.
* gimplify.cc (enum omp_region_type): Add ORT_DISPATCH.
(gimplify_call_expr): Handle need_device_ptr arguments.
(is_gimple_stmt): Handle OMP_DISPATCH.
(gimplify_scan_omp_clauses): Handle OMP_CLAUSE_DEVICE in a dispatch
construct. Handle OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT.
(gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_NOVARIANTS and
OMP_CLAUSE_NOCONTEXT.
(omp_construct_selector_matches): Handle OMP_DISPATCH with nocontext
clause.
(omp_has_novariants): New function.
(omp_has_nocontext): Likewise.
(gimplify_omp_dispatch): Likewise.
(gimplify_expr): Handle OMP_DISPATCH.
* gimplify.h (omp_has_novariants): Declare.
(omp_has_nocontext): Declare.
* omp-builtins.def (BUILT_IN_OMP_GET_MAPPED_PTR): Define.
(BUILT_IN_OMP_GET_DEFAULT_DEVICE): Define.
(BUILT_IN_OMP_SET_DEFAULT_DEVICE): Define.
* omp-expand.cc (expand_omp_dispatch): New function.
(expand_omp): Handle GIMPLE_OMP_DISPATCH.
(omp_make_gimple_edges): Likewise.
* omp-general.cc (omp_construct_traits_to_codes): Add OMP_DISPATCH.
(struct omp_ts_info): Add dispatch.
(omp_context_selector_matches): Handle OMP_TRAIT_SET_NEED_DEVICE_PTR.
(omp_resolve_declare_variant): Handle novariants. Adjust
DECL_ASSEMBLER_NAME.
---
gcc/gimple-low.cc | 1 +
gcc/gimple-pretty-print.cc | 33 +++
gcc/gimple-walk.cc | 1 +
gcc/gimple.cc | 20 ++
gcc/gimple.def | 5 +
gcc/gimple.h | 33 ++-
gcc/gimplify.cc | 417 ++++++++++++++++++++++++++++++++++++-
gcc/gimplify.h | 2 +
gcc/omp-builtins.def | 6 +
gcc/omp-expand.cc | 18 ++
gcc/omp-general.cc | 16 +-
gcc/omp-low.cc | 35 ++++
gcc/tree-inline.cc | 7 +
13 files changed, 583 insertions(+), 11 deletions(-)
diff --git a/gcc/gimple-low.cc b/gcc/gimple-low.cc
index e0371988705..712a1ebf776 100644
--- a/gcc/gimple-low.cc
+++ b/gcc/gimple-low.cc
@@ -746,6 +746,7 @@ lower_stmt (gimple_stmt_iterator *gsi, struct lower_data *data)
case GIMPLE_EH_MUST_NOT_THROW:
case GIMPLE_OMP_FOR:
case GIMPLE_OMP_SCOPE:
+ case GIMPLE_OMP_DISPATCH:
case GIMPLE_OMP_SECTIONS:
case GIMPLE_OMP_SECTIONS_SWITCH:
case GIMPLE_OMP_SECTION:
diff --git a/gcc/gimple-pretty-print.cc b/gcc/gimple-pretty-print.cc
index a71e1e0efc7..d9a24ad2169 100644
--- a/gcc/gimple-pretty-print.cc
+++ b/gcc/gimple-pretty-print.cc
@@ -1726,6 +1726,35 @@ dump_gimple_omp_scope (pretty_printer *buffer, const gimple *gs,
}
}
+/* Dump a GIMPLE_OMP_DISPATCH tuple on the pretty_printer BUFFER. */
+
+static void
+dump_gimple_omp_dispatch (pretty_printer *buffer, const gimple *gs, int spc,
+ dump_flags_t flags)
+{
+ if (flags & TDF_RAW)
+ {
+ dump_gimple_fmt (buffer, spc, flags, "%G <%+BODY <%S>%nCLAUSES <", gs,
+ gimple_omp_body (gs));
+ dump_omp_clauses (buffer, gimple_omp_dispatch_clauses (gs), spc, flags);
+ dump_gimple_fmt (buffer, spc, flags, " >");
+ }
+ else
+ {
+ pp_string (buffer, "#pragma omp dispatch");
+ dump_omp_clauses (buffer, gimple_omp_dispatch_clauses (gs), spc, flags);
+ if (!gimple_seq_empty_p (gimple_omp_body (gs)))
+ {
+ newline_and_indent (buffer, spc + 2);
+ pp_left_brace (buffer);
+ pp_newline (buffer);
+ dump_gimple_seq (buffer, gimple_omp_body (gs), spc + 4, flags);
+ newline_and_indent (buffer, spc + 2);
+ pp_right_brace (buffer);
+ }
+ }
+}
+
/* Dump a GIMPLE_OMP_TARGET tuple on the pretty_printer BUFFER. */
static void
@@ -2805,6 +2834,10 @@ pp_gimple_stmt_1 (pretty_printer *buffer, const gimple *gs, int spc,
dump_gimple_omp_scope (buffer, gs, spc, flags);
break;
+ case GIMPLE_OMP_DISPATCH:
+ dump_gimple_omp_dispatch (buffer, gs, spc, flags);
+ break;
+
case GIMPLE_OMP_MASTER:
case GIMPLE_OMP_SECTION:
case GIMPLE_OMP_STRUCTURED_BLOCK:
diff --git a/gcc/gimple-walk.cc b/gcc/gimple-walk.cc
index 9f768ca20fd..1122713a98b 100644
--- a/gcc/gimple-walk.cc
+++ b/gcc/gimple-walk.cc
@@ -707,6 +707,7 @@ walk_gimple_stmt (gimple_stmt_iterator *gsi, walk_stmt_fn callback_stmt,
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
case GIMPLE_OMP_SCOPE:
+ case GIMPLE_OMP_DISPATCH:
case GIMPLE_OMP_SECTIONS:
case GIMPLE_OMP_SINGLE:
case GIMPLE_OMP_TARGET:
diff --git a/gcc/gimple.cc b/gcc/gimple.cc
index a9f968cb038..3a26c74a105 100644
--- a/gcc/gimple.cc
+++ b/gcc/gimple.cc
@@ -1235,6 +1235,21 @@ gimple_build_omp_scope (gimple_seq body, tree clauses)
return p;
}
+/* Build a GIMPLE_OMP_DISPATCH statement.
+
+ BODY is the target function call to be dispatched.
+ CLAUSES are any of the OMP dispatch construct's clauses: ... */
+
+gimple *
+gimple_build_omp_dispatch (gimple_seq body, tree clauses)
+{
+ gimple *p = gimple_alloc (GIMPLE_OMP_DISPATCH, 0);
+ gimple_omp_dispatch_set_clauses (p, clauses);
+ if (body)
+ gimple_omp_set_body (p, body);
+
+ return p;
+}
/* Build a GIMPLE_OMP_TARGET statement.
@@ -2148,6 +2163,11 @@ gimple_copy (gimple *stmt)
gimple_omp_scope_set_clauses (copy, t);
goto copy_omp_body;
+ case GIMPLE_OMP_DISPATCH:
+ t = unshare_expr (gimple_omp_dispatch_clauses (stmt));
+ gimple_omp_dispatch_set_clauses (copy, t);
+ goto copy_omp_body;
+
case GIMPLE_OMP_TARGET:
{
gomp_target *omp_target_stmt = as_a <gomp_target *> (stmt);
diff --git a/gcc/gimple.def b/gcc/gimple.def
index fbcd727f945..21c7405875d 100644
--- a/gcc/gimple.def
+++ b/gcc/gimple.def
@@ -350,6 +350,11 @@ DEFGSCODE(GIMPLE_OMP_SCAN, "gimple_omp_scan", GSS_OMP_SINGLE_LAYOUT)
CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */
DEFGSCODE(GIMPLE_OMP_SCOPE, "gimple_omp_scope", GSS_OMP_SINGLE_LAYOUT)
+/* GIMPLE_OMP_DISPATCH <BODY, CLAUSES> represents #pragma omp dispatch
+ BODY is the target function call to be dispatched.
+ CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */
+DEFGSCODE(GIMPLE_OMP_DISPATCH, "gimple_omp_dispatch", GSS_OMP_SINGLE_LAYOUT)
+
/* OMP_SECTION <BODY> represents #pragma omp section.
BODY is the sequence of statements in the section body. */
DEFGSCODE(GIMPLE_OMP_SECTION, "gimple_omp_section", GSS_OMP)
diff --git a/gcc/gimple.h b/gcc/gimple.h
index bd315ffc2dd..25590a22ffb 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -746,7 +746,7 @@ struct GTY((tag("GSS_OMP_CONTINUE")))
};
/* GIMPLE_OMP_SINGLE, GIMPLE_OMP_ORDERED, GIMPLE_OMP_TASKGROUP,
- GIMPLE_OMP_SCAN, GIMPLE_OMP_MASKED, GIMPLE_OMP_SCOPE. */
+ GIMPLE_OMP_SCAN, GIMPLE_OMP_MASKED, GIMPLE_OMP_SCOPE, GIMPLE_OMP_DISPATCH. */
struct GTY((tag("GSS_OMP_SINGLE_LAYOUT")))
gimple_statement_omp_single_layout : public gimple_statement_omp
@@ -1595,6 +1595,7 @@ gomp_task *gimple_build_omp_task (gimple_seq, tree, tree, tree, tree,
gimple *gimple_build_omp_section (gimple_seq);
gimple *gimple_build_omp_structured_block (gimple_seq);
gimple *gimple_build_omp_scope (gimple_seq, tree);
+gimple *gimple_build_omp_dispatch (gimple_seq, tree);
gimple *gimple_build_omp_master (gimple_seq);
gimple *gimple_build_omp_masked (gimple_seq, tree);
gimple *gimple_build_omp_taskgroup (gimple_seq, tree);
@@ -1886,6 +1887,7 @@ gimple_has_substatements (gimple *g)
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
case GIMPLE_OMP_SCOPE:
+ case GIMPLE_OMP_DISPATCH:
case GIMPLE_OMP_SECTIONS:
case GIMPLE_OMP_SINGLE:
case GIMPLE_OMP_TARGET:
@@ -5437,6 +5439,34 @@ gimple_omp_scope_set_clauses (gimple *gs, tree clauses)
= clauses;
}
+/* Return the clauses associated with OMP_DISPATCH statement GS. */
+
+inline tree
+gimple_omp_dispatch_clauses (const gimple *gs)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_DISPATCH);
+ return static_cast<const gimple_statement_omp_single_layout *> (gs)->clauses;
+}
+
+/* Return a pointer to the clauses associated with OMP dispatch statement
+ GS. */
+
+inline tree *
+gimple_omp_dispatch_clauses_ptr (gimple *gs)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_DISPATCH);
+ return &static_cast<gimple_statement_omp_single_layout *> (gs)->clauses;
+}
+
+/* Set CLAUSES to be the clauses associated with OMP dispatch statement
+ GS. */
+
+inline void
+gimple_omp_dispatch_set_clauses (gimple *gs, tree clauses)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_DISPATCH);
+ static_cast<gimple_statement_omp_single_layout *> (gs)->clauses = clauses;
+}
/* Return the kind of the OMP_FOR statemement G. */
@@ -6771,6 +6801,7 @@ gimple_return_set_retval (greturn *gs, tree retval)
case GIMPLE_OMP_TARGET: \
case GIMPLE_OMP_TEAMS: \
case GIMPLE_OMP_SCOPE: \
+ case GIMPLE_OMP_DISPATCH: \
case GIMPLE_OMP_SECTION: \
case GIMPLE_OMP_STRUCTURED_BLOCK: \
case GIMPLE_OMP_MASTER: \
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index b0ed58ed0f9..1dd69dbb1de 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -161,7 +161,8 @@ enum omp_region_type
{
ORT_WORKSHARE = 0x00,
ORT_TASKGROUP = 0x01,
- ORT_SIMD = 0x04,
+ ORT_DISPATCH = 0x02,
+ ORT_SIMD = 0x04,
ORT_PARALLEL = 0x08,
ORT_COMBINED_PARALLEL = ORT_PARALLEL | 1,
@@ -4051,6 +4052,7 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value)
/* Gimplify the function arguments. */
if (nargs > 0)
{
+ tree device_num = NULL_TREE;
for (i = (PUSH_ARGS_REVERSED ? nargs - 1 : 0);
PUSH_ARGS_REVERSED ? i >= 0 : i < nargs;
PUSH_ARGS_REVERSED ? i-- : i++)
@@ -4061,8 +4063,99 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value)
be the plain PARM_DECL. */
if ((i != 1) || !builtin_va_start_p)
{
- t = gimplify_arg (&CALL_EXPR_ARG (*expr_p, i), pre_p,
- EXPR_LOCATION (*expr_p), ! returns_twice);
+ tree *arg_p = &CALL_EXPR_ARG (*expr_p, i);
+ if (flag_openmp && EXPR_P (CALL_EXPR_FN (*expr_p))
+ && lookup_attribute ("omp declare variant variant",
+ DECL_ATTRIBUTES (TREE_OPERAND (
+ CALL_EXPR_FN (*expr_p), 0)))
+ != NULL_TREE)
+ {
+ tree param
+ = DECL_ARGUMENTS (TREE_OPERAND (CALL_EXPR_FN (*expr_p), 0));
+
+ if (param != NULL_TREE)
+ {
+ for (int param_idx = 0; param_idx < i; param_idx++)
+ param = TREE_CHAIN (param);
+
+ bool is_device_ptr = false;
+ if (gimplify_omp_ctxp != NULL
+ && gimplify_omp_ctxp->code == OMP_DISPATCH)
+ {
+ for (tree c = gimplify_omp_ctxp->clauses; c;
+ c = TREE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c)
+ == OMP_CLAUSE_IS_DEVICE_PTR)
+ {
+ tree decl1 = DECL_NAME (OMP_CLAUSE_DECL (c));
+ tree decl2
+ = tree_strip_nop_conversions (*arg_p);
+ if (TREE_CODE (decl2) == ADDR_EXPR)
+ decl2 = TREE_OPERAND (decl2, 0);
+ gcc_assert (TREE_CODE (decl2) == VAR_DECL
+ || TREE_CODE (decl2)
+ == PARM_DECL);
+ decl2 = DECL_NAME (decl2);
+ if (decl1 == decl2)
+ {
+ is_device_ptr = true;
+ break;
+ }
+ }
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE)
+ device_num = OMP_CLAUSE_OPERAND (c, 0);
+ }
+ }
+
+ if (!is_device_ptr
+ && lookup_attribute ("omp declare variant "
+ "adjust_args need_device_ptr",
+ DECL_ATTRIBUTES (param))
+ != NULL_TREE)
+ {
+ if (device_num == NULL_TREE)
+ {
+ // device_num = omp_get_default_device();
+ tree fn = builtin_decl_explicit (
+ BUILT_IN_OMP_GET_DEFAULT_DEVICE);
+ gcall *call = gimple_build_call (fn, 0);
+ device_num = create_tmp_var (
+ gimple_call_return_type (call));
+ gimple_call_set_lhs (call, device_num);
+ gimplify_seq_add_stmt (pre_p, call);
+ }
+
+ // mapped_arg = omp_get_mapped_ptr(arg, device_num);
+ tree fn = builtin_decl_explicit (
+ BUILT_IN_OMP_GET_MAPPED_PTR);
+ *arg_p = (TREE_CODE (*arg_p) == NOP_EXPR)
+ ? TREE_OPERAND (*arg_p, 0)
+ : *arg_p;
+ gimplify_arg (arg_p, pre_p, loc);
+ gimplify_arg (&device_num, pre_p, loc);
+ call = gimple_build_call (fn, 2, *arg_p, device_num);
+ tree mapped_arg
+ = create_tmp_var (gimple_call_return_type (call));
+ gimple_call_set_lhs (call, mapped_arg);
+ gimplify_seq_add_stmt (pre_p, call);
+
+ *arg_p = mapped_arg;
+
+ // Mark mapped argument as device pointer to ensure
+ // idempotency in gimplification
+ gcc_assert (gimplify_omp_ctxp->code == OMP_DISPATCH);
+ tree c = build_omp_clause (input_location,
+ OMP_CLAUSE_IS_DEVICE_PTR);
+ OMP_CLAUSE_DECL (c) = *arg_p;
+ OMP_CLAUSE_CHAIN (c) = gimplify_omp_ctxp->clauses;
+ gimplify_omp_ctxp->clauses = c;
+ }
+ }
+ }
+
+ t = gimplify_arg (arg_p, pre_p, EXPR_LOCATION (*expr_p),
+ !returns_twice);
if (t == GS_ERROR)
ret = GS_ERROR;
@@ -6307,6 +6400,7 @@ is_gimple_stmt (tree t)
case OACC_LOOP:
case OMP_SCAN:
case OMP_SCOPE:
+ case OMP_DISPATCH:
case OMP_SECTIONS:
case OMP_SECTION:
case OMP_STRUCTURED_BLOCK:
@@ -13080,6 +13174,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
}
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE
+ && code == OMP_DISPATCH)
+ {
+ bool saved_into_ssa = gimplify_ctxp->into_ssa;
+ gimplify_ctxp->into_ssa = false;
+ if (gimplify_expr (&OMP_CLAUSE_DEVICE_ID (c), pre_p, NULL,
+ is_gimple_val, fb_rvalue)
+ == GS_ERROR)
+ remove = true;
+ else if (DECL_P (OMP_CLAUSE_DEVICE_ID (c)))
+ omp_add_variable (ctx, OMP_CLAUSE_DEVICE_ID (c),
+ GOVD_SHARED | GOVD_SEEN);
+ gimplify_ctxp->into_ssa = saved_into_ssa;
+ break;
+ }
/* Fall through. */
case OMP_CLAUSE_PRIORITY:
@@ -13309,6 +13418,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
break;
+ case OMP_CLAUSE_NOVARIANTS:
+ {
+ OMP_CLAUSE_NOVARIANTS_EXPR (c);
+ tree t = gimple_boolify (OMP_CLAUSE_NOVARIANTS_EXPR (c));
+ OMP_CLAUSE_NOVARIANTS_EXPR (c) = t;
+ }
+ break;
+ case OMP_CLAUSE_NOCONTEXT:
+ {
+ OMP_CLAUSE_NOCONTEXT_EXPR (c);
+ tree t = gimple_boolify (OMP_CLAUSE_NOCONTEXT_EXPR (c));
+ OMP_CLAUSE_NOCONTEXT_EXPR (c) = t;
+ }
+ break;
case OMP_CLAUSE_NOHOST:
default:
gcc_unreachable ();
@@ -13763,7 +13886,9 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
{
struct gimplify_omp_ctx *octx;
for (octx = ctx; octx; octx = octx->outer_context)
- if ((octx->region_type & (ORT_PARALLEL | ORT_TASK | ORT_TEAMS)) != 0)
+ if ((octx->region_type
+ & (ORT_DISPATCH | ORT_PARALLEL | ORT_TASK | ORT_TEAMS))
+ != 0)
break;
if (octx)
{
@@ -14574,6 +14699,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
case OMP_CLAUSE_FINALIZE:
case OMP_CLAUSE_INCLUSIVE:
case OMP_CLAUSE_EXCLUSIVE:
+ case OMP_CLAUSE_NOVARIANTS:
+ case OMP_CLAUSE_NOCONTEXT:
break;
case OMP_CLAUSE_NOHOST:
@@ -14663,9 +14790,9 @@ omp_construct_selector_matches (enum tree_code *constructs, int nconstructs,
== ORT_TARGET && ctx->code == OMP_TARGET)
|| ((ctx->region_type & ORT_TEAMS) && ctx->code == OMP_TEAMS)
|| (ctx->region_type == ORT_WORKSHARE && ctx->code == OMP_FOR)
- || (ctx->region_type == ORT_SIMD
- && ctx->code == OMP_SIMD
- && !omp_find_clause (ctx->clauses, OMP_CLAUSE_BIND)))
+ || (ctx->region_type == ORT_SIMD && ctx->code == OMP_SIMD
+ && !omp_find_clause (ctx->clauses, OMP_CLAUSE_BIND))
+ || (ctx->code == OMP_DISPATCH && omp_has_nocontext () != 1))
{
++cnt;
if (scores)
@@ -14783,6 +14910,60 @@ omp_construct_selector_matches (enum tree_code *constructs, int nconstructs,
return 0;
}
+/* Try to evaluate a novariants clause. Return 1 if true, 0 if false or absent,
+ * -1 if run-time evaluation is needed. */
+
+int
+omp_has_novariants (void)
+{
+ for (struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; ctx;
+ ctx = ctx->outer_context)
+ {
+ if (ctx->code == OMP_DISPATCH)
+ {
+ tree c = omp_find_clause (ctx->clauses, OMP_CLAUSE_NOVARIANTS);
+ if (c != NULL_TREE)
+ {
+ if (integer_nonzerop (OMP_CLAUSE_NOVARIANTS_EXPR (c)))
+ return 1;
+ else if (integer_zerop (OMP_CLAUSE_NOVARIANTS_EXPR (c)))
+ return 0;
+ else
+ return -1;
+ }
+ return 0;
+ }
+ }
+ return 0;
+}
+
+/* Try to evaluate a nocontext clause. Return 1 if true, 0 if false or absent,
+ * -1 if run-time evaluation is needed. */
+
+int
+omp_has_nocontext (void)
+{
+ for (struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; ctx;
+ ctx = ctx->outer_context)
+ {
+ if (ctx->code == OMP_DISPATCH)
+ {
+ tree c = omp_find_clause (ctx->clauses, OMP_CLAUSE_NOCONTEXT);
+ if (c != NULL_TREE)
+ {
+ if (integer_nonzerop (OMP_CLAUSE_NOCONTEXT_EXPR (c)))
+ return 1;
+ else if (integer_zerop (OMP_CLAUSE_NOCONTEXT_EXPR (c)))
+ return 0;
+ else
+ return -1;
+ }
+ return 0;
+ }
+ }
+ return 0;
+}
+
/* Gimplify OACC_CACHE. */
static void
@@ -17614,6 +17795,221 @@ gimplify_omp_ordered (tree expr, gimple_seq body)
return gimple_build_omp_ordered (body, OMP_ORDERED_CLAUSES (expr));
}
+/* Gimplify an OMP_DISPATCH construct. */
+
+static enum gimplify_status
+gimplify_omp_dispatch (tree *expr_p, gimple_seq *pre_p)
+{
+ tree expr = *expr_p;
+ gimple_seq body = NULL;
+
+ gimplify_scan_omp_clauses (&OMP_DISPATCH_CLAUSES (expr), pre_p, ORT_DISPATCH,
+ OMP_DISPATCH);
+ push_gimplify_context ();
+
+ // If device clause, adjust ICV
+ tree device
+ = omp_find_clause (OMP_DISPATCH_CLAUSES (expr), OMP_CLAUSE_DEVICE);
+ if (device)
+ {
+ tree t = builtin_decl_explicit (BUILT_IN_OMP_SET_DEFAULT_DEVICE);
+ t = build_call_expr_loc (input_location, t, 1,
+ OMP_CLAUSE_DEVICE_ID (device));
+ gimplify_and_add (t, &body);
+ if (DECL_P (OMP_CLAUSE_DEVICE_ID (device)))
+ omp_notice_variable (gimplify_omp_ctxp, OMP_CLAUSE_DEVICE_ID (device),
+ true);
+ }
+
+ // If the novariants and nocontext clauses are not compile-time constants,
+ // we need to generate code for all possible cases:
+ // if (novariants) // implies nocontext
+ // base()
+ // else if (nocontext)
+ // variant1()
+ // else
+ // variant2()
+ tree dispatch_body = OMP_DISPATCH_BODY (expr);
+ if (TREE_CODE (dispatch_body) == BIND_EXPR)
+ dispatch_body = BIND_EXPR_BODY (dispatch_body);
+ if (TREE_CODE (dispatch_body) == STATEMENT_LIST)
+ {
+ // Fortran FE may insert some pre-call code, for instance when an
+ // array is passed as argument. Skip to the actual call.
+ dispatch_body = expr_last (dispatch_body);
+ }
+ gcc_assert (TREE_CODE (dispatch_body) == CALL_EXPR
+ || TREE_CODE (dispatch_body) == MODIFY_EXPR);
+ tree base_call_expr = dispatch_body;
+ tree dst;
+ if (TREE_CODE (base_call_expr) == MODIFY_EXPR)
+ {
+ dst = TREE_OPERAND (base_call_expr, 0);
+ base_call_expr = TREE_OPERAND (base_call_expr, 1);
+ while (TREE_CODE (base_call_expr) == FLOAT_EXPR
+ || TREE_CODE (base_call_expr) == CONVERT_EXPR
+ || TREE_CODE (base_call_expr) == COMPLEX_EXPR)
+ base_call_expr = TREE_OPERAND (base_call_expr, 0);
+ }
+
+ tree base_fndecl = get_callee_fndecl (STRIP_NOPS (base_call_expr));
+ if (base_fndecl != NULL_TREE)
+ {
+ if (DECL_VIRTUAL_P (base_fndecl))
+ {
+ error_at (
+ EXPR_LOCATION (base_call_expr),
+ "%qD is a virtual function but only a direct call is allowed "
+ "in a dispatch construct",
+ DECL_NAME (base_fndecl));
+ }
+
+ tree variant_fndecl = omp_resolve_declare_variant (base_fndecl);
+ if (base_fndecl != variant_fndecl
+ && (omp_has_novariants () == -1 || omp_has_nocontext () == -1))
+ {
+ tree novariants_clause = NULL_TREE, nocontext_clause = NULL_TREE,
+ novariants_cond = NULL_TREE, nocontext_cond = NULL_TREE;
+ for (tree c = OMP_DISPATCH_CLAUSES (expr); c; c = TREE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NOVARIANTS)
+ {
+ gcc_assert (novariants_cond == NULL_TREE);
+ novariants_clause = c;
+ novariants_cond = OMP_CLAUSE_NOVARIANTS_EXPR (c);
+ }
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NOCONTEXT)
+ {
+ gcc_assert (nocontext_cond == NULL_TREE);
+ nocontext_clause = c;
+ nocontext_cond = OMP_CLAUSE_NOCONTEXT_EXPR (c);
+ }
+ }
+ gcc_assert (novariants_cond != NULL_TREE
+ || nocontext_cond != NULL_TREE);
+
+ enum gimplify_status ret
+ = gimplify_expr (&novariants_cond, &body, NULL, is_gimple_val,
+ fb_rvalue);
+ if (ret == GS_ERROR || ret == GS_UNHANDLED)
+ return ret;
+ ret = gimplify_expr (&nocontext_cond, &body, NULL, is_gimple_val,
+ fb_rvalue);
+ if (ret == GS_ERROR || ret == GS_UNHANDLED)
+ return ret;
+
+ tree base_label = create_artificial_label (UNKNOWN_LOCATION);
+ tree variant1_label = create_artificial_label (UNKNOWN_LOCATION);
+ tree cond_label = create_artificial_label (UNKNOWN_LOCATION);
+ tree variant2_label = create_artificial_label (UNKNOWN_LOCATION);
+ tree end_label = create_artificial_label (UNKNOWN_LOCATION);
+
+ if (novariants_cond != NULL_TREE)
+ {
+ gcond *novariants_cond_stmt
+ = gimple_build_cond_from_tree (novariants_cond, base_label,
+ cond_label);
+ gimplify_seq_add_stmt (&body, novariants_cond_stmt);
+
+ gimplify_seq_add_stmt (&body, gimple_build_label (base_label));
+ tree base_call_expr2 = copy_node (base_call_expr);
+ if (TREE_CODE (dispatch_body) == MODIFY_EXPR)
+ {
+ base_call_expr2 = build2 (MODIFY_EXPR, TREE_TYPE (dst), dst,
+ base_call_expr2);
+ }
+ OMP_CLAUSE_NOVARIANTS_EXPR (novariants_clause)
+ = boolean_true_node;
+ gimplify_and_add (base_call_expr2, &body);
+ gimplify_seq_add_stmt (&body, gimple_build_goto (end_label));
+
+ OMP_CLAUSE_NOVARIANTS_EXPR (novariants_clause)
+ = boolean_false_node;
+ }
+
+ gimplify_seq_add_stmt (&body, gimple_build_label (cond_label));
+ if (nocontext_cond != NULL_TREE)
+ {
+ gcond *nocontext_cond_stmt
+ = gimple_build_cond_from_tree (nocontext_cond, variant1_label,
+ variant2_label);
+ gimplify_seq_add_stmt (&body, nocontext_cond_stmt);
+
+ gimplify_seq_add_stmt (&body,
+ gimple_build_label (variant1_label));
+ tree variant_call_expr = copy_node (base_call_expr);
+ if (TREE_CODE (dispatch_body) == MODIFY_EXPR)
+ {
+ variant_call_expr = build2 (MODIFY_EXPR, TREE_TYPE (dst), dst,
+ variant_call_expr);
+ }
+ OMP_CLAUSE_NOCONTEXT_EXPR (nocontext_clause) = boolean_true_node;
+ gimplify_and_add (variant_call_expr, &body);
+ gimplify_seq_add_stmt (&body, gimple_build_goto (end_label));
+ OMP_CLAUSE_NOCONTEXT_EXPR (nocontext_clause) = boolean_false_node;
+ }
+
+ gimplify_seq_add_stmt (&body, gimple_build_label (variant2_label));
+ tree variant_call_expr = copy_node (base_call_expr);
+ if (TREE_CODE (dispatch_body) == MODIFY_EXPR)
+ {
+ variant_call_expr
+ = build2 (MODIFY_EXPR, TREE_TYPE (dst), dst, variant_call_expr);
+ }
+ gimplify_and_add (variant_call_expr, &body);
+ gimplify_seq_add_stmt (&body, gimple_build_goto (end_label));
+ gimplify_seq_add_stmt (&body, gimple_build_label (end_label));
+ }
+ else
+ gimplify_and_add (OMP_DISPATCH_BODY (expr), &body);
+ }
+ else
+ gimplify_and_add (OMP_DISPATCH_BODY (expr), &body);
+
+ // Wrap dispatch body into a bind
+ gimple *bind = gimple_build_bind (NULL_TREE, body, NULL_TREE);
+ pop_gimplify_context (bind);
+
+ gimplify_adjust_omp_clauses (pre_p, bind, &OMP_DISPATCH_CLAUSES (expr),
+ OMP_DISPATCH);
+
+ // Move relevant clauses to the task construct
+ tree task_clauses = NULL_TREE;
+ tree *task_clauses_ptr = &task_clauses;
+ bool has_nowait = false;
+ for (tree c = OMP_DISPATCH_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SHARED
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND)
+ {
+ *task_clauses_ptr = c;
+ task_clauses_ptr = &OMP_CLAUSE_CHAIN (c);
+ }
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
+ {
+ *task_clauses_ptr
+ = build_omp_clause (input_location, OMP_CLAUSE_SHARED);
+ OMP_CLAUSE_DECL (*task_clauses_ptr) = OMP_CLAUSE_DECL (c);
+ task_clauses_ptr = &OMP_CLAUSE_CHAIN (*task_clauses_ptr);
+ }
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NOWAIT)
+ has_nowait = true;
+ }
+ *task_clauses_ptr = build_omp_clause (input_location, OMP_CLAUSE_IF);
+ OMP_CLAUSE_IF_EXPR (*task_clauses_ptr)
+ = has_nowait ? boolean_true_node : boolean_false_node;
+
+ // Wrap bind into a task
+ gimple *task
+ = gimple_build_omp_task (bind, task_clauses, NULL_TREE, NULL_TREE,
+ NULL_TREE, NULL_TREE, NULL_TREE);
+
+ gimple *stmt = gimple_build_omp_dispatch (task, OMP_DISPATCH_CLAUSES (expr));
+ gimplify_seq_add_stmt (pre_p, stmt);
+ *expr_p = NULL_TREE;
+ return GS_ALL_DONE;
+}
+
/* Convert the GENERIC expression tree *EXPR_P to GIMPLE. If the
expression produces a value to be used as an operand inside a GIMPLE
statement, the value will be stored back in *EXPR_P. This value will
@@ -18540,6 +18936,10 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = gimplify_omp_atomic (expr_p, pre_p);
break;
+ case OMP_DISPATCH:
+ ret = gimplify_omp_dispatch (expr_p, pre_p);
+ break;
+
case TRANSACTION_EXPR:
ret = gimplify_transaction (expr_p, pre_p);
break;
@@ -18865,7 +19265,8 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
&& code != OMP_SECTION
&& code != OMP_STRUCTURED_BLOCK
&& code != OMP_SINGLE
- && code != OMP_SCOPE);
+ && code != OMP_SCOPE
+ && code != OMP_DISPATCH);
}
#endif
diff --git a/gcc/gimplify.h b/gcc/gimplify.h
index ac3cc8eb552..55aece2b65b 100644
--- a/gcc/gimplify.h
+++ b/gcc/gimplify.h
@@ -77,6 +77,8 @@ extern enum gimplify_status gimplify_expr (tree *, gimple_seq *, gimple_seq *,
bool (*) (tree), fallback_t);
int omp_construct_selector_matches (enum tree_code *, int, int *);
+int omp_has_novariants (void);
+int omp_has_nocontext (void);
extern void gimplify_type_sizes (tree, gimple_seq *);
extern void gimplify_one_sizepos (tree *, gimple_seq *);
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index 044d5d087b6..c83edabbcc3 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -76,6 +76,12 @@ DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_TEAM_NUM, "omp_get_team_num",
BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_NUM_TEAMS, "omp_get_num_teams",
BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_MAPPED_PTR, "omp_get_mapped_ptr",
+ BT_FN_PTR_CONST_PTR_INT, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_DEFAULT_DEVICE, "omp_get_default_device",
+ BT_FN_INT, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_OMP_SET_DEFAULT_DEVICE, "omp_set_default_device",
+ BT_FN_INT, ATTR_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_ATOMIC_START, "GOMP_atomic_start",
BT_FN_VOID, ATTR_NOTHROW_LEAF_LIST)
diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index 24287826444..6fa372a550b 100644
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -8636,6 +8636,19 @@ expand_omp_single (struct omp_region *region)
single_succ_edge (exit_bb)->flags = EDGE_FALLTHRU;
}
+/* Expand code for an OpenMP dispatch directive... */
+
+static void
+expand_omp_dispatch (struct omp_region *region)
+{
+ basic_block entry_bb = region->entry;
+ gimple_stmt_iterator si = gsi_last_nondebug_bb (entry_bb);
+ enum gimple_code code = gimple_code (gsi_stmt (si));
+ gcc_assert (code == GIMPLE_OMP_DISPATCH);
+ gsi_remove (&si, true);
+ single_succ_edge (entry_bb)->flags = EDGE_FALLTHRU;
+}
+
/* Generic expansion for OpenMP synchronization directives: master,
ordered and critical. All we need to do here is remove the entry
and exit markers for REGION. */
@@ -10654,6 +10667,10 @@ expand_omp (struct omp_region *region)
expand_omp_single (region);
break;
+ case GIMPLE_OMP_DISPATCH:
+ expand_omp_dispatch (region);
+ break;
+
case GIMPLE_OMP_ORDERED:
{
gomp_ordered *ord_stmt
@@ -11001,6 +11018,7 @@ omp_make_gimple_edges (basic_block bb, struct omp_region **region,
case GIMPLE_OMP_MASTER:
case GIMPLE_OMP_MASKED:
case GIMPLE_OMP_SCOPE:
+ case GIMPLE_OMP_DISPATCH:
case GIMPLE_OMP_CRITICAL:
case GIMPLE_OMP_SECTION:
cur_region = new_omp_region (bb, code, cur_region);
diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc
index 2c095200d5b..d585df400d1 100644
--- a/gcc/omp-general.cc
+++ b/gcc/omp-general.cc
@@ -1037,7 +1037,7 @@ omp_construct_traits_to_codes (tree ctx, int nconstructs,
/* Order must match the OMP_TRAIT_CONSTRUCT_* enumerators in
enum omp_ts_code. */
static enum tree_code code_map[]
- = { OMP_TARGET, OMP_TEAMS, OMP_PARALLEL, OMP_FOR, OMP_SIMD };
+ = { OMP_TARGET, OMP_TEAMS, OMP_PARALLEL, OMP_FOR, OMP_SIMD, OMP_DISPATCH };
for (tree ts = ctx; ts; ts = TREE_CHAIN (ts), i--)
{
@@ -1130,6 +1130,7 @@ const char *omp_tss_map[] =
"target_device",
"implementation",
"user",
+ "need_device_ptr",
NULL
};
@@ -1236,10 +1237,14 @@ struct omp_ts_info omp_ts_map[] =
OMP_TRAIT_PROPERTY_CLAUSE_LIST, false,
NULL
},
+ { "dispatch",
+ (1 << OMP_TRAIT_SET_CONSTRUCT),
+ OMP_TRAIT_PROPERTY_NONE, false,
+ NULL
+ },
{ NULL, 0, OMP_TRAIT_PROPERTY_NONE, false, NULL } /* OMP_TRAIT_LAST */
};
-
/* Return a name from PROP, a property in selectors accepting
name lists. */
@@ -1445,6 +1450,8 @@ omp_context_selector_matches (tree ctx)
for (tree tss = ctx; tss; tss = TREE_CHAIN (tss))
{
enum omp_tss_code set = OMP_TSS_CODE (tss);
+ if (set == OMP_TRAIT_SET_NEED_DEVICE_PTR)
+ continue;
tree selectors = OMP_TSS_TRAIT_SELECTORS (tss);
/* Immediately reject the match if there are any ignored
@@ -2484,6 +2491,9 @@ omp_resolve_declare_variant (tree base)
if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
return omp_resolve_late_declare_variant (base);
+ if (omp_has_novariants () == 1)
+ return base;
+
auto_vec <tree, 16> variants;
auto_vec <bool, 16> defer;
bool any_deferred = false;
@@ -2630,6 +2640,8 @@ omp_resolve_declare_variant (tree base)
(*slot)->variants = entry.variants;
tree alt = build_decl (DECL_SOURCE_LOCATION (base), FUNCTION_DECL,
DECL_NAME (base), TREE_TYPE (base));
+ if (DECL_ASSEMBLER_NAME_SET_P (base))
+ SET_DECL_ASSEMBLER_NAME (alt, DECL_ASSEMBLER_NAME (base));
DECL_ARTIFICIAL (alt) = 1;
DECL_IGNORED_P (alt) = 1;
TREE_STATIC (alt) = 1;
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 4d003f42098..693d8ca7d8d 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -4185,6 +4185,11 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
scan_omp (gimple_omp_body_ptr (stmt), ctx);
break;
+ case GIMPLE_OMP_DISPATCH:
+ ctx = new_omp_context (stmt, ctx);
+ scan_omp (gimple_omp_body_ptr (stmt), ctx);
+ break;
+
case GIMPLE_OMP_SECTIONS:
scan_omp_sections (as_a <gomp_sections *> (stmt), ctx);
break;
@@ -8926,6 +8931,31 @@ lower_omp_scope (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (BLOCK_VARS (block))
TREE_USED (block) = 1;
}
+
+/* Lower code for an OMP dispatch directive. */
+
+static void
+lower_omp_dispatch (gimple_stmt_iterator *gsi_p, omp_context *ctx)
+{
+ tree block;
+ gimple *stmt = gsi_stmt (*gsi_p);
+ gbind *bind;
+
+ push_gimplify_context ();
+
+ block = make_node (BLOCK);
+ bind = gimple_build_bind (NULL, NULL, block);
+ gsi_replace (gsi_p, bind, true);
+
+ lower_omp (gimple_omp_body_ptr (stmt), ctx);
+ gimple_bind_set_body (bind, maybe_catch_exception (gimple_omp_body (stmt)));
+
+ pop_gimplify_context (bind);
+
+ gimple_bind_append_vars (bind, ctx->block_vars);
+ BLOCK_VARS (block) = ctx->block_vars;
+}
+
/* Expand code for an OpenMP master or masked directive. */
static void
@@ -14399,6 +14429,11 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gcc_assert (ctx);
lower_omp_scope (gsi_p, ctx);
break;
+ case GIMPLE_OMP_DISPATCH:
+ ctx = maybe_lookup_ctx (stmt);
+ gcc_assert (ctx);
+ lower_omp_dispatch (gsi_p, ctx);
+ break;
case GIMPLE_OMP_SINGLE:
ctx = maybe_lookup_ctx (stmt);
gcc_assert (ctx);
diff --git a/gcc/tree-inline.cc b/gcc/tree-inline.cc
index f31a34ac410..2e06b706025 100644
--- a/gcc/tree-inline.cc
+++ b/gcc/tree-inline.cc
@@ -1679,6 +1679,12 @@ remap_gimple_stmt (gimple *stmt, copy_body_data *id)
(s1, gimple_omp_scope_clauses (stmt));
break;
+ case GIMPLE_OMP_DISPATCH:
+ s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
+ copy = gimple_build_omp_dispatch (s1,
+ gimple_omp_dispatch_clauses (stmt));
+ break;
+
case GIMPLE_OMP_TASKGROUP:
s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
copy = gimple_build_omp_taskgroup
@@ -4609,6 +4615,7 @@ estimate_num_insns (gimple *stmt, eni_weights *weights)
case GIMPLE_OMP_MASTER:
case GIMPLE_OMP_MASKED:
case GIMPLE_OMP_SCOPE:
+ case GIMPLE_OMP_DISPATCH:
case GIMPLE_OMP_TASKGROUP:
case GIMPLE_OMP_ORDERED:
case GIMPLE_OMP_SCAN:
--
2.45.1
^ permalink raw reply [flat|nested] 10+ messages in thread* [PATCH 3/7] OpenMP: C front-end support for dispatch + adjust_args
2024-05-27 11:54 [PATCH 0/7] OpenMP: dispatch + adjust_args support Paul-Antoine Arras
2024-05-27 11:54 ` [PATCH 1/7] OpenMP: dispatch + adjust_args tree data structures and front-end interfaces Paul-Antoine Arras
2024-05-27 11:54 ` [PATCH 2/7] OpenMP: middle-end support for dispatch + adjust_args Paul-Antoine Arras
@ 2024-05-27 11:54 ` Paul-Antoine Arras
2024-05-27 11:54 ` [PATCH 4/7] OpenMP: C++ " Paul-Antoine Arras
` (3 subsequent siblings)
6 siblings, 0 replies; 10+ messages in thread
From: Paul-Antoine Arras @ 2024-05-27 11:54 UTC (permalink / raw)
To: gcc-patches; +Cc: Paul-Antoine Arras
This patch adds support to the C front-end to parse the `dispatch` construct and
the `adjust_args` clause. It also includes some common C/C++ bits for pragmas
and attributes.
Additional common C/C++ testcases are in a later patch in the series.
gcc/c-family/ChangeLog:
* c-attribs.cc (c_common_gnu_attributes): Add attribute for adjust_args
need_device_ptr.
* c-omp.cc (c_omp_directives): Uncomment dispatch.
* c-pragma.cc (omp_pragmas): Add dispatch.
* c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_DISPATCH.
(enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_NOCONTEXT and
PRAGMA_OMP_CLAUSE_NOVARIANTS.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_dispatch): New function.
(c_parser_omp_clause_name): Handle nocontext and novariants clauses.
(c_parser_omp_clause_novariants): New function.
(c_parser_omp_clause_nocontext): Likewise.
(c_parser_omp_all_clauses): Handle nocontext and novariants clauses.
(c_parser_omp_dispatch_body): New function adapted from
c_parser_expr_no_commas.
(OMP_DISPATCH_CLAUSE_MASK): Define.
(c_parser_omp_dispatch): New function.
(c_finish_omp_declare_variant): Parse adjust_args.
(c_parser_omp_construct): Handle PRAGMA_OMP_DISPATCH.
* c-typeck.cc (c_finish_omp_clauses): Handle OMP_CLAUSE_NOVARIANTS and
OMP_CLAUSE_NOCONTEXT.
gcc/testsuite/ChangeLog:
* gcc.dg/gomp/adjust-args-1.c: New test.
* gcc.dg/gomp/dispatch-1.c: New test.
---
gcc/c-family/c-attribs.cc | 2 +
gcc/c-family/c-omp.cc | 4 +-
gcc/c-family/c-pragma.cc | 1 +
gcc/c-family/c-pragma.h | 3 +
gcc/c/c-parser.cc | 483 +++++++++++++++++++---
gcc/c/c-typeck.cc | 2 +
gcc/testsuite/gcc.dg/gomp/adjust-args-1.c | 32 ++
gcc/testsuite/gcc.dg/gomp/dispatch-1.c | 53 +++
libgomp/testsuite/libgomp.c/dispatch-1.c | 76 ++++
9 files changed, 601 insertions(+), 55 deletions(-)
create mode 100644 gcc/testsuite/gcc.dg/gomp/adjust-args-1.c
create mode 100644 gcc/testsuite/gcc.dg/gomp/dispatch-1.c
create mode 100644 libgomp/testsuite/libgomp.c/dispatch-1.c
diff --git a/gcc/c-family/c-attribs.cc b/gcc/c-family/c-attribs.cc
index 04e39b41bdf..860a068d527 100644
--- a/gcc/c-family/c-attribs.cc
+++ b/gcc/c-family/c-attribs.cc
@@ -556,6 +556,8 @@ const struct attribute_spec c_common_gnu_attributes[] =
handle_omp_declare_variant_attribute, NULL },
{ "omp declare variant variant", 0, -1, true, false, false, false,
handle_omp_declare_variant_attribute, NULL },
+ { "omp declare variant adjust_args need_device_ptr", 0, -1, true, false, false, false,
+ handle_omp_declare_variant_attribute, NULL },
{ "simd", 0, 1, true, false, false, false,
handle_simd_attribute, NULL },
{ "omp declare target", 0, -1, true, false, false, false,
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index c0e02aa422f..e6b42dbd01c 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -4196,8 +4196,8 @@ const struct c_omp_directive c_omp_directives[] = {
C_OMP_DIR_DECLARATIVE, false },
{ "depobj", nullptr, nullptr, PRAGMA_OMP_DEPOBJ,
C_OMP_DIR_STANDALONE, false },
- /* { "dispatch", nullptr, nullptr, PRAGMA_OMP_DISPATCH,
- C_OMP_DIR_CONSTRUCT, false }, */
+ { "dispatch", nullptr, nullptr, PRAGMA_OMP_DISPATCH,
+ C_OMP_DIR_DECLARATIVE, false },
{ "distribute", nullptr, nullptr, PRAGMA_OMP_DISTRIBUTE,
C_OMP_DIR_CONSTRUCT, true },
{ "end", "assumes", nullptr, PRAGMA_OMP_END,
diff --git a/gcc/c-family/c-pragma.cc b/gcc/c-family/c-pragma.cc
index 1237ee6e62b..60fadeee286 100644
--- a/gcc/c-family/c-pragma.cc
+++ b/gcc/c-family/c-pragma.cc
@@ -1526,6 +1526,7 @@ static const struct omp_pragma_def omp_pragmas[] = {
{ "cancellation", PRAGMA_OMP_CANCELLATION_POINT },
{ "critical", PRAGMA_OMP_CRITICAL },
{ "depobj", PRAGMA_OMP_DEPOBJ },
+ { "dispatch", PRAGMA_OMP_DISPATCH },
{ "error", PRAGMA_OMP_ERROR },
{ "end", PRAGMA_OMP_END },
{ "flush", PRAGMA_OMP_FLUSH },
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index ce93a52fa57..061a83d1716 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -55,6 +55,7 @@ enum pragma_kind {
PRAGMA_OMP_CRITICAL,
PRAGMA_OMP_DECLARE,
PRAGMA_OMP_DEPOBJ,
+ PRAGMA_OMP_DISPATCH,
PRAGMA_OMP_DISTRIBUTE,
PRAGMA_OMP_ERROR,
PRAGMA_OMP_END,
@@ -132,9 +133,11 @@ enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_LINK,
PRAGMA_OMP_CLAUSE_MAP,
PRAGMA_OMP_CLAUSE_MERGEABLE,
+ PRAGMA_OMP_CLAUSE_NOCONTEXT,
PRAGMA_OMP_CLAUSE_NOGROUP,
PRAGMA_OMP_CLAUSE_NONTEMPORAL,
PRAGMA_OMP_CLAUSE_NOTINBRANCH,
+ PRAGMA_OMP_CLAUSE_NOVARIANTS,
PRAGMA_OMP_CLAUSE_NOWAIT,
PRAGMA_OMP_CLAUSE_NUM_TASKS,
PRAGMA_OMP_CLAUSE_NUM_TEAMS,
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 00f8bf4376e..c9cd36fb429 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -1733,6 +1733,8 @@ static void c_parser_omp_assumption_clauses (c_parser *, bool);
static void c_parser_omp_allocate (c_parser *);
static void c_parser_omp_assumes (c_parser *);
static bool c_parser_omp_ordered (c_parser *, enum pragma_context, bool *);
+static tree
+c_parser_omp_dispatch (location_t, c_parser *);
static void c_parser_oacc_routine (c_parser *, enum pragma_context);
/* These Objective-C parser functions are only ever called when
@@ -14972,6 +14974,8 @@ c_parser_omp_clause_name (c_parser *parser)
case 'n':
if (!strcmp ("no_create", p))
result = PRAGMA_OACC_CLAUSE_NO_CREATE;
+ else if (!strcmp ("nocontext", p))
+ result = PRAGMA_OMP_CLAUSE_NOCONTEXT;
else if (!strcmp ("nogroup", p))
result = PRAGMA_OMP_CLAUSE_NOGROUP;
else if (!strcmp ("nohost", p))
@@ -14980,6 +14984,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
else if (!strcmp ("notinbranch", p))
result = PRAGMA_OMP_CLAUSE_NOTINBRANCH;
+ else if (!strcmp ("novariants", p))
+ result = PRAGMA_OMP_CLAUSE_NOVARIANTS;
else if (!strcmp ("nowait", p))
result = PRAGMA_OMP_CLAUSE_NOWAIT;
else if (!strcmp ("num_gangs", p))
@@ -19190,6 +19196,60 @@ c_parser_omp_clause_uniform (c_parser *parser, tree list)
return list;
}
+/* OpenMP 5.1
+ novariants ( scalar-expression ) */
+
+static tree
+c_parser_omp_clause_novariants (c_parser *parser, tree list)
+{
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ location_t loc = c_parser_peek_token (parser)->location;
+ c_expr expr = c_parser_expr_no_commas (parser, NULL);
+ tree t = convert_lvalue_to_rvalue (loc, expr, true, true).value;
+ t = c_objc_common_truthvalue_conversion (loc, t);
+ t = c_fully_fold (t, false, NULL);
+ parens.skip_until_found_close (parser);
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_NOVARIANTS, "novariants");
+
+ tree c = build_omp_clause (loc, OMP_CLAUSE_NOVARIANTS);
+ OMP_CLAUSE_NOVARIANTS_EXPR (c) = t;
+ OMP_CLAUSE_CHAIN (c) = list;
+ list = c;
+
+ return list;
+}
+
+/* OpenMP 5.1
+ nocontext ( scalar-expression ) */
+
+static tree
+c_parser_omp_clause_nocontext (c_parser *parser, tree list)
+{
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ location_t loc = c_parser_peek_token (parser)->location;
+ c_expr expr = c_parser_expr_no_commas (parser, NULL);
+ tree t = convert_lvalue_to_rvalue (loc, expr, true, true).value;
+ t = c_objc_common_truthvalue_conversion (loc, t);
+ t = c_fully_fold (t, false, NULL);
+ parens.skip_until_found_close (parser);
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_NOCONTEXT, "nocontext");
+
+ tree c = build_omp_clause (loc, OMP_CLAUSE_NOCONTEXT);
+ OMP_CLAUSE_NOCONTEXT_EXPR (c) = t;
+ OMP_CLAUSE_CHAIN (c) = list;
+ list = c;
+
+ return list;
+}
+
/* OpenMP 5.0:
detach ( event-handle ) */
@@ -19801,6 +19861,14 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses);
c_name = "enter";
break;
+ case PRAGMA_OMP_CLAUSE_NOVARIANTS:
+ c_name = "novariants";
+ clauses = c_parser_omp_clause_novariants (parser, clauses);
+ break;
+ case PRAGMA_OMP_CLAUSE_NOCONTEXT:
+ c_name = "nocontext";
+ clauses = c_parser_omp_clause_nocontext (parser, clauses);
+ break;
default:
c_parser_error (parser, "expected an OpenMP clause");
goto saw_error;
@@ -23458,6 +23526,168 @@ c_parser_omp_scope (location_t loc, c_parser *parser, bool *if_p)
return add_stmt (stmt);
}
+// Adapted from c_parser_expr_no_commas
+static tree
+c_parser_omp_dispatch_body (c_parser *parser)
+{
+ struct c_expr lhs, rhs, ret;
+ struct c_expr orig_expr;
+ location_t expr_loc = c_parser_peek_token (parser)->location;
+ source_range tok_range = c_parser_peek_token (parser)->get_range ();
+ location_t sizeof_arg_loc[3];
+ tree sizeof_arg[3];
+ vec<tree, va_gc> *exprlist;
+ vec<location_t> arg_loc = vNULL;
+ vec<tree, va_gc> *origtypes = NULL;
+ unsigned int literal_zero_mask;
+ location_t start;
+ location_t finish;
+
+ lhs = c_parser_conditional_expression (parser, NULL, NULL);
+ if (TREE_CODE (lhs.value) == CALL_EXPR)
+ return lhs.value;
+ else
+ {
+ location_t op_location = c_parser_peek_token (parser)->location;
+ if (!c_parser_require (parser, CPP_EQ, "expected %<=%>"))
+ return error_mark_node;
+
+ /* Parse function name*/
+ if (!c_parser_next_token_is (parser, CPP_NAME))
+ {
+ c_parser_error (parser, "expected a function name");
+ rhs.set_error ();
+ return rhs.value;
+ }
+ expr_loc = c_parser_peek_token (parser)->location;
+ tree id = c_parser_peek_token (parser)->value;
+ c_parser_consume_token (parser);
+ if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+ return error_mark_node;
+
+ rhs.value = build_external_ref (expr_loc, id, true, &rhs.original_type);
+ set_c_expr_source_range (&rhs, tok_range);
+ /* Parse argument list */
+ {
+ for (int i = 0; i < 3; i++)
+ {
+ sizeof_arg[i] = NULL_TREE;
+ sizeof_arg_loc[i] = UNKNOWN_LOCATION;
+ }
+ literal_zero_mask = 0;
+ if (c_parser_next_token_is (parser, CPP_CLOSE_PAREN))
+ exprlist = NULL;
+ else
+ exprlist = c_parser_expr_list (parser, true, false, &origtypes,
+ sizeof_arg_loc, sizeof_arg, &arg_loc,
+ &literal_zero_mask);
+ c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
+ }
+ orig_expr = rhs;
+ mark_exp_read (rhs.value);
+ if (warn_sizeof_pointer_memaccess)
+ sizeof_pointer_memaccess_warning (sizeof_arg_loc, rhs.value, exprlist,
+ sizeof_arg,
+ sizeof_ptr_memacc_comptypes);
+ if (TREE_CODE (rhs.value) == FUNCTION_DECL)
+ {
+ if (fndecl_built_in_p (rhs.value, BUILT_IN_MEMSET)
+ && vec_safe_length (exprlist) == 3)
+ {
+ tree arg0 = (*exprlist)[0];
+ tree arg2 = (*exprlist)[2];
+ warn_for_memset (expr_loc, arg0, arg2, literal_zero_mask);
+ }
+ if (warn_absolute_value
+ && fndecl_built_in_p (rhs.value, BUILT_IN_NORMAL)
+ && vec_safe_length (exprlist) == 1)
+ warn_for_abs (expr_loc, rhs.value, (*exprlist)[0]);
+ if (parser->omp_for_parse_state
+ && parser->omp_for_parse_state->in_intervening_code
+ && omp_runtime_api_call (rhs.value))
+ {
+ error_at (expr_loc, "calls to the OpenMP runtime API are "
+ "not permitted in intervening code");
+ parser->omp_for_parse_state->fail = true;
+ }
+ }
+
+ start = rhs.get_start ();
+ finish = parser->tokens_buf[0].get_finish ();
+ rhs.value = c_build_function_call_vec (expr_loc, arg_loc, rhs.value,
+ exprlist, origtypes);
+ set_c_expr_source_range (&rhs, start, finish);
+ rhs.m_decimal = 0;
+
+ rhs.original_code = ERROR_MARK;
+ if (TREE_CODE (rhs.value) == INTEGER_CST
+ && TREE_CODE (orig_expr.value) == FUNCTION_DECL
+ && fndecl_built_in_p (orig_expr.value, BUILT_IN_CONSTANT_P))
+ rhs.original_code = C_MAYBE_CONST_EXPR;
+ rhs.original_type = NULL;
+ if (exprlist)
+ {
+ release_tree_vector (exprlist);
+ release_tree_vector (origtypes);
+ }
+ arg_loc.release ();
+
+ /* Build assignment */
+ rhs = convert_lvalue_to_rvalue (expr_loc, rhs, true, true);
+ ret.value
+ = build_modify_expr (op_location, lhs.value, lhs.original_type,
+ NOP_EXPR, expr_loc, rhs.value, rhs.original_type);
+ ret.m_decimal = 0;
+ set_c_expr_source_range (&ret, lhs.get_start (), rhs.get_finish ());
+ ret.original_code = MODIFY_EXPR;
+ ret.original_type = NULL;
+ return ret.value;
+ }
+}
+
+/* OpenMP 5.1:
+ # pragma omp dispatch dispatch-clause[optseq] new-line
+ expression-stmt
+
+ LOC is the location of the #pragma.
+*/
+
+#define OMP_DISPATCH_CLAUSE_MASK \
+ ((OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT))
+
+static tree
+c_parser_omp_dispatch (location_t loc, c_parser *parser)
+{
+ tree stmt = make_node (OMP_DISPATCH);
+ SET_EXPR_LOCATION (stmt, loc);
+ TREE_TYPE (stmt) = void_type_node;
+
+ OMP_DISPATCH_CLAUSES (stmt)
+ = c_parser_omp_all_clauses (parser, OMP_DISPATCH_CLAUSE_MASK,
+ "#pragma omp dispatch");
+
+ // Parse body as expression statement
+ loc = c_parser_peek_token (parser)->location;
+ tree dispatch_body = c_parser_omp_dispatch_body (parser);
+ if (dispatch_body == error_mark_node)
+ {
+ inform (loc, "%<#pragma omp dispatch%> must be followed by a function "
+ "call with optional assignment");
+ c_parser_skip_to_end_of_block_or_statement (parser);
+ return NULL_TREE;
+ }
+
+ c_parser_skip_until_found (parser, CPP_SEMICOLON, "expected %<;%>");
+ OMP_DISPATCH_BODY (stmt) = dispatch_body;
+
+ return add_stmt (stmt);
+}
+
/* OpenMP 3.0:
# pragma omp task task-clause[optseq] new-line
@@ -24438,6 +24668,10 @@ check_clauses:
OpenMP 5.0:
# pragma omp declare variant (identifier) match(context-selector) new-line
+
+ OpenMP 5.1
+ # pragma omp declare variant (identifier) match(context-selector) \
+ adjust_args(adjust-op:argument-list) new-line
*/
#define OMP_DECLARE_SIMD_CLAUSE_MASK \
@@ -24901,19 +25135,39 @@ c_finish_omp_declare_variant (c_parser *parser, tree fndecl, tree parms)
parens.require_close (parser);
- if (c_parser_next_token_is (parser, CPP_COMMA)
- && c_parser_peek_2nd_token (parser)->type == CPP_NAME)
- c_parser_consume_token (parser);
+ vec<tree> adjust_args_list = vNULL;
+ bool has_match = false, has_adjust_args = false;
+ location_t adjust_args_loc;
- const char *clause = "";
- location_t match_loc = c_parser_peek_token (parser)->location;
- if (c_parser_next_token_is (parser, CPP_NAME))
- clause = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
- if (strcmp (clause, "match"))
+ do
{
- c_parser_error (parser, "expected %<match%>");
- goto fail;
- }
+ if (c_parser_next_token_is (parser, CPP_COMMA)
+ && c_parser_peek_2nd_token (parser)->type == CPP_NAME)
+ c_parser_consume_token (parser);
+
+ const char *clause = "";
+ location_t match_loc = c_parser_peek_token (parser)->location;
+ if (c_parser_next_token_is (parser, CPP_NAME))
+ clause = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+
+ enum clause
+ {
+ match,
+ adjust_args
+ } ccode;
+
+ if (strcmp (clause, "match") == 0)
+ ccode = match;
+ else if (strcmp (clause, "adjust_args") == 0)
+ {
+ ccode = adjust_args;
+ adjust_args_loc = match_loc;
+ }
+ else
+ {
+ c_parser_error (parser, "expected %<match%> or %<adjust_args%>");
+ goto fail;
+ }
c_parser_consume_token (parser);
@@ -24923,55 +25177,172 @@ c_finish_omp_declare_variant (c_parser *parser, tree fndecl, tree parms)
if (parms == NULL_TREE)
parms = error_mark_node;
- tree ctx = c_parser_omp_context_selector_specification (parser, parms);
- if (ctx == error_mark_node)
- goto fail;
- ctx = omp_check_context_selector (match_loc, ctx);
- if (ctx != error_mark_node && variant != error_mark_node)
+ if (ccode == match)
{
- if (TREE_CODE (variant) != FUNCTION_DECL)
+ has_match = true;
+ tree ctx = c_parser_omp_context_selector_specification (parser, parms);
+ if (ctx == error_mark_node)
+ goto fail;
+ ctx = omp_check_context_selector (match_loc, ctx);
+ if (ctx != error_mark_node && variant != error_mark_node)
{
- error_at (token->location, "variant %qD is not a function", variant);
- variant = error_mark_node;
- }
- else if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT,
- OMP_TRAIT_CONSTRUCT_SIMD)
- && !comptypes (TREE_TYPE (fndecl), TREE_TYPE (variant)))
- {
- error_at (token->location, "variant %qD and base %qD have "
- "incompatible types", variant, fndecl);
- variant = error_mark_node;
- }
- else if (fndecl_built_in_p (variant)
- && (strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
- "__builtin_", strlen ("__builtin_")) == 0
- || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
- "__sync_", strlen ("__sync_")) == 0
- || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
- "__atomic_", strlen ("__atomic_")) == 0))
- {
- error_at (token->location, "variant %qD is a built-in", variant);
- variant = error_mark_node;
- }
- if (variant != error_mark_node)
- {
- C_DECL_USED (variant) = 1;
- tree construct
- = omp_get_context_selector_list (ctx, OMP_TRAIT_SET_CONSTRUCT);
- omp_mark_declare_variant (match_loc, variant, construct);
- if (omp_context_selector_matches (ctx))
+ if (TREE_CODE (variant) != FUNCTION_DECL)
{
- tree attr
- = tree_cons (get_identifier ("omp declare variant base"),
- build_tree_list (variant, ctx),
- DECL_ATTRIBUTES (fndecl));
- DECL_ATTRIBUTES (fndecl) = attr;
+ error_at (token->location, "variant %qD is not a function",
+ variant);
+ variant = error_mark_node;
+ }
+ else if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT,
+ OMP_TRAIT_CONSTRUCT_SIMD)
+ && !comptypes (TREE_TYPE (fndecl), TREE_TYPE (variant)))
+ {
+ error_at (token->location,
+ "variant %qD and base %qD have "
+ "incompatible types",
+ variant, fndecl);
+ variant = error_mark_node;
+ }
+ else if (fndecl_built_in_p (variant)
+ && (strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
+ "__builtin_", strlen ("__builtin_"))
+ == 0
+ || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
+ "__sync_", strlen ("__sync_"))
+ == 0
+ || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
+ "__atomic_", strlen ("__atomic_"))
+ == 0))
+ {
+ error_at (token->location, "variant %qD is a built-in", variant);
+ variant = error_mark_node;
+ }
+ if (variant != error_mark_node)
+ {
+ C_DECL_USED (variant) = 1;
+ tree construct
+ = omp_get_context_selector_list (ctx, OMP_TRAIT_SET_CONSTRUCT);
+ omp_mark_declare_variant (match_loc, variant, construct);
+ if (omp_context_selector_matches (ctx))
+ {
+ tree attr
+ = tree_cons (get_identifier ("omp declare variant base"),
+ build_tree_list (variant, ctx),
+ DECL_ATTRIBUTES (fndecl));
+ DECL_ATTRIBUTES (fndecl) = attr;
+ }
}
}
}
+ else if (ccode == adjust_args)
+ {
+ has_adjust_args = true;
+ if (c_parser_next_token_is (parser, CPP_NAME)
+ && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+ {
+ const char *p
+ = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+ if (strcmp (p, "need_device_ptr") == 0
+ || strcmp (p, "nothing") == 0)
+ {
+ c_parser_consume_token (parser); // need_device_ptr
+ c_parser_consume_token (parser); // :
+
+ location_t loc = c_parser_peek_token (parser)->location;
+ tree list
+ = c_parser_omp_variable_list (parser, loc, OMP_CLAUSE_ERROR,
+ NULL_TREE);
+
+ tree arg;
+ for (tree c = list; c != NULL_TREE; c = TREE_CHAIN (c))
+ {
+ tree decl = TREE_PURPOSE (c);
+ int idx;
+ for (arg = parms, idx = 0; arg != NULL;
+ arg = TREE_CHAIN (arg), idx++)
+ if (arg == decl)
+ break;
+ if (arg == NULL_TREE)
+ {
+ error_at (loc, "%qD is not a function argument",
+ decl);
+ goto fail;
+ }
+ if (adjust_args_list.contains (arg))
+ {
+ error_at (loc, "%qD is specified more than once",
+ decl);
+ goto fail;
+ }
+ if (strcmp (p, "need_device_ptr") == 0
+ && TREE_CODE (TREE_TYPE (arg)) != POINTER_TYPE)
+ {
+ error_at (loc, "%qD is not a C pointer", decl);
+ goto fail;
+ }
+ adjust_args_list.safe_push (arg);
+ if (strcmp (p, "need_device_ptr") == 0)
+ {
+ tree variant_decl = (TREE_CODE (variant) == NOP_EXPR)
+ ? TREE_OPERAND (variant, 0)
+ : variant;
+ tree variant_parm = DECL_ARGUMENTS (variant_decl);
+ for (int i = 0; i < idx; i++)
+ {
+ variant_parm = TREE_CHAIN (variant_parm);
+ gcc_assert (variant_parm != NULL);
+ }
+ tree attr = tree_cons (
+ get_identifier ("omp declare variant adjust_args "
+ "need_device_ptr"),
+ NULL_TREE, DECL_ATTRIBUTES (variant_parm));
+ DECL_ATTRIBUTES (variant_parm) = attr;
+ }
+ }
+ }
+ else
+ {
+ error_at (c_parser_peek_token (parser)->location,
+ "expected %<nothing%> or %<need_device_ptr%>");
+ goto fail;
+ }
+ }
+ else
+ {
+ error_at (c_parser_peek_token (parser)->location,
+ "expected %<nothing%> or %<need_device_ptr%> "
+ "followed by %<:%>");
+ goto fail;
+ }
+ }
parens.require_close (parser);
+ } while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL));
c_parser_skip_to_pragma_eol (parser);
+
+ if (has_adjust_args)
+ {
+ if (!has_match)
+ {
+ error_at (
+ adjust_args_loc,
+ "an %<adjust_args%> clause can only be specified if the "
+ "%<dispatch%> selector of the construct selector set appears "
+ "in the %<match%> clause");
+ }
+ else
+ {
+ tree attr = lookup_attribute ("omp declare variant base",
+ DECL_ATTRIBUTES (fndecl));
+ tree ctx = TREE_VALUE (TREE_VALUE (attr));
+ if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT,
+ OMP_TRAIT_CONSTRUCT_DISPATCH))
+ error_at (
+ adjust_args_loc,
+ "an %<adjust_args%> clause can only be specified if the "
+ "%<dispatch%> selector of the construct selector set appears "
+ "in the %<match%> clause");
+ }
+ }
}
/* Finalize #pragma omp declare simd or #pragma omp declare variant
@@ -25788,7 +26159,6 @@ c_parser_omp_declare_reduction (c_parser *parser, enum pragma_context context)
types.release ();
}
-
/* OpenMP 4.0
#pragma omp declare simd declare-simd-clauses[optseq] new-line
#pragma omp declare reduction (reduction-id : typename-list : expression) \
@@ -25796,7 +26166,11 @@ c_parser_omp_declare_reduction (c_parser *parser, enum pragma_context context)
#pragma omp declare target new-line
OpenMP 5.0
- #pragma omp declare variant (identifier) match (context-selector) */
+ #pragma omp declare variant (identifier) match (context-selector)
+
+ OpenMP 5.1
+ #pragma omp declare variant (identifier) match (context-selector) \
+ adjust_args(adjust-op:argument-list) */
static bool
c_parser_omp_declare (c_parser *parser, enum pragma_context context)
@@ -26599,6 +26973,9 @@ c_parser_omp_construct (c_parser *parser, bool *if_p)
case PRAGMA_OMP_ASSUME:
c_parser_omp_assume (parser, if_p);
return;
+ case PRAGMA_OMP_DISPATCH:
+ stmt = c_parser_omp_dispatch (loc, parser);
+ break;
default:
gcc_unreachable ();
}
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 2d092357e0f..48b4bd906e8 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -16054,6 +16054,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_FINALIZE:
case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE_INDIRECT:
+ case OMP_CLAUSE_NOVARIANTS:
+ case OMP_CLAUSE_NOCONTEXT:
pc = &OMP_CLAUSE_CHAIN (c);
continue;
diff --git a/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c b/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c
new file mode 100644
index 00000000000..393a44de8e0
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c
@@ -0,0 +1,32 @@
+/* Test parsing of OMP clause adjust_args */
+/* { dg-do compile } */
+
+int b;
+
+int f0 (void *a);
+int g (void *a);
+int f1 (int);
+
+#pragma omp declare variant (f0) match (construct={target}) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */
+int f2 (void *a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (other: a) /* { dg-error "expected 'nothing' or 'need_device_ptr'" } */
+int f3 (int a);
+#pragma omp declare variant (f0) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */
+int f4 (void *a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args () /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */
+int f5 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing) /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */
+int f6 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing:) /* { dg-error "expected identifier before '\\)' token" } */
+int f7 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: z) /* { dg-error "'z' undeclared here \\(not in a function\\)" } */
+int f8 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: a) /* { dg-error "'a' is not a C pointer" } */
+int f9 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (nothing: a) /* { dg-error "'a' is specified more than once" } */
+int f10 (int a);
+#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: a) /* { dg-error "'a' is specified more than once" } */
+int f11 (void *a);
+#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: b) /* { dg-error "'b' is not a function argument" } */
+int f12 (void *a);
+
diff --git a/gcc/testsuite/gcc.dg/gomp/dispatch-1.c b/gcc/testsuite/gcc.dg/gomp/dispatch-1.c
new file mode 100644
index 00000000000..c8f45c12be6
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/dispatch-1.c
@@ -0,0 +1,53 @@
+/* Test parsing of #pragma omp dispatch */
+/* { dg-do compile } */
+
+int f0 (int);
+
+void f1 (void)
+{
+ int a, b;
+ double x;
+ struct {int a; int b;} s;
+ int arr[1];
+
+#pragma omp dispatch
+ int c = f0 (a); /* { dg-error "expected expression before 'int'" } */
+#pragma omp dispatch
+ int f2 (int d); /* { dg-error "expected expression before 'int'" } */
+#pragma omp dispatch
+ a = b; /* { dg-error "expected '\\(' before ';' token" } */
+#pragma omp dispatch
+ s.a = f0(a) + b; /* { dg-error "expected ';' before '\\+' token" } */
+#pragma omp dispatch
+ b = !f0(a); /* { dg-error "expected a function name before '!' token" } */
+#pragma omp dispatch
+ s.b += f0(s.a); /* { dg-error "expected '=' before '\\+=' token" } */
+#pragma omp dispatch
+#pragma omp threadprivate(a) /* { dg-error "expected expression before '#pragma'" } */
+ a = f0(b);
+
+#pragma omp dispatch nocontext(s) /* { dg-error "used struct type value where scalar is required" } */
+ f0(a);
+#pragma omp dispatch nocontext(a, b) /* { dg-error "expected '\\)' before ','" } */
+ f0(a);
+#pragma omp dispatch nocontext(a) nocontext(b) /* { dg-error "too many 'nocontext' clauses" } */
+ f0(a);
+#pragma omp dispatch novariants(s) /* { dg-error "used struct type value where scalar is required" } */
+ f0(a);
+#pragma omp dispatch novariants(a, b) /* { dg-error "expected '\\)' before ','" } */
+ f0(a);
+#pragma omp dispatch novariants(a) novariants(b) /* { dg-error "too many 'novariants' clauses" } */
+ f0(a);
+#pragma omp dispatch nowait nowait /* { dg-error "too many 'nowait' clauses" } */
+ f0(a);
+#pragma omp dispatch device(x) /* { dg-error "expected integer expression before end of line" } */
+ f0(a);
+#pragma omp dispatch device(arr) /* { dg-error "expected integer expression before end of line" } */
+ f0(a);
+#pragma omp dispatch is_device_ptr(x) /* { dg-error "'is_device_ptr' variable is neither a pointer nor an array" } */
+ f0(a);
+#pragma omp dispatch is_device_ptr(&x) /* { dg-error "expected identifier before '&' token" } */
+ f0(a);
+#pragma omp dispatch depend(inout: f0) /* { dg-error "'f0' is not lvalue expression nor array section in 'depend' clause" } */
+ f0(a);
+}
diff --git a/libgomp/testsuite/libgomp.c/dispatch-1.c b/libgomp/testsuite/libgomp.c/dispatch-1.c
new file mode 100644
index 00000000000..0efc075a859
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/dispatch-1.c
@@ -0,0 +1,76 @@
+// Adapted from OpenMP examples
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int baz (double *d_bv, const double *d_av, int n)
+{
+#pragma omp distribute parallel for
+ for (int i = 0; i < n; i++)
+ d_bv[i] = d_av[i] * i;
+ return -3;
+}
+
+int bar (double *d_bv, const double *d_av, int n)
+{
+#pragma omp target is_device_ptr(d_bv, d_av)
+ for (int i = 0; i < n; i++)
+ d_bv[i] = d_av[i] * i;
+ return -2;
+}
+
+#pragma omp declare variant(bar) match(construct={dispatch}) adjust_args(need_device_ptr: bv, av)
+#pragma omp declare variant(baz) match(implementation={vendor(gnu)})
+int foo (double *bv, const double *av, int n)
+{
+ for (int i = 0; i < n; i++)
+ bv[i] = av[i] * i;
+ return -1;
+}
+
+int test (int n)
+{
+ const double e = 2.71828;
+
+ double *av = (double *) malloc (n * sizeof (*av));
+ double *bv = (double *) malloc (n * sizeof (*bv));
+ double *d_bv = (double *) malloc (n * sizeof (*d_bv));
+
+ for (int i = 0; i < n; i++)
+ {
+ av[i] = e * i;
+ bv[i] = 0.0;
+ d_bv[i] = 0.0;
+ }
+
+ int f, last_dev = omp_get_num_devices () - 1;
+#pragma omp target data map(to: av[:n]) map(from: d_bv[:n]) device(last_dev) if (n == 1024)
+ {
+ #pragma omp dispatch nocontext(n > 1024) novariants(n < 1024) device(last_dev)
+ f = foo (d_bv, av, n);
+ }
+
+ foo (bv, av, n);
+ for (int i = 0; i < n; i++)
+ {
+ if (d_bv[i] != bv[i])
+ {
+ fprintf (stderr, "ERROR at %d: %lf (act) != %lf (exp)\n", i, d_bv[i], bv[i]);
+ return 1;
+ }
+ }
+ return f;
+}
+
+int
+main (void)
+{
+ int ret = test(1023);
+ if (ret != -1) return 1;
+ ret = test(1024);
+ if (ret != -2) return 1;
+ ret = test(1025);
+ if (ret != -3) return 1;
+ return 0;
+}
--
2.45.1
^ permalink raw reply [flat|nested] 10+ messages in thread* [PATCH 4/7] OpenMP: C++ front-end support for dispatch + adjust_args
2024-05-27 11:54 [PATCH 0/7] OpenMP: dispatch + adjust_args support Paul-Antoine Arras
` (2 preceding siblings ...)
2024-05-27 11:54 ` [PATCH 3/7] OpenMP: C front-end " Paul-Antoine Arras
@ 2024-05-27 11:54 ` Paul-Antoine Arras
2024-05-27 11:54 ` [PATCH 5/7] OpenMP: common C/C++ testcases " Paul-Antoine Arras
` (2 subsequent siblings)
6 siblings, 0 replies; 10+ messages in thread
From: Paul-Antoine Arras @ 2024-05-27 11:54 UTC (permalink / raw)
To: gcc-patches; +Cc: Paul-Antoine Arras
This patch adds C++ support for the `dispatch` construct and the `adjust_args`
clause. It relies on the c-family bits comprised in the corresponding C front
end patch for pragmas and attributes.
Additional C/C++ common testcases are provided in a subsequent patch in the
series.
gcc/cp/ChangeLog:
* decl.cc (omp_declare_variant_finalize_one): Set adjust_args
need_device_ptr attribute.
* parser.cc (cp_parser_direct_declarator): Update call to
cp_parser_late_return_type_opt.
(cp_parser_late_return_type_opt): Add parameter. Update call to
cp_parser_late_parsing_omp_declare_simd.
(cp_parser_omp_clause_name): Handle nocontext and novariants clauses.
(cp_parser_omp_clause_novariants): New function.
(cp_parser_omp_clause_nocontext): Likewise.
(cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_NOVARIANTS and
PRAGMA_OMP_CLAUSE_NOCONTEXT.
(cp_parser_omp_dispatch_body): New function, inspired from
cp_parser_assignment_expression and cp_parser_postfix_expression.
(OMP_DISPATCH_CLAUSE_MASK): Define.
(cp_parser_omp_dispatch): New function.
(cp_finish_omp_declare_variant): Add parameter. Handle adjust_args
clause.
(cp_parser_late_parsing_omp_declare_simd): Add parameter. Update calls
to cp_finish_omp_declare_variant and cp_finish_omp_declare_variant.
(cp_parser_omp_construct): Handle PRAGMA_OMP_DISPATCH.
(cp_parser_pragma): Likewise.
* pt.cc (tsubst_attribute): Skip pseudo-TSS need_device_ptr.
* semantics.cc (finish_omp_clauses): Handle OMP_CLAUSE_NOCONTEXT and
OMP_CLAUSE_NOVARIANTS.
gcc/testsuite/ChangeLog:
* g++.dg/gomp/adjust-args-1.C: New test.
* g++.dg/gomp/adjust-args-2.C: New test.
* g++.dg/gomp/dispatch-1.C: New test.
* g++.dg/gomp/dispatch-2.C: New test.
---
gcc/cp/decl.cc | 27 +
gcc/cp/parser.cc | 613 ++++++++++++++++++++--
gcc/cp/pt.cc | 3 +
gcc/cp/semantics.cc | 20 +
gcc/testsuite/g++.dg/gomp/adjust-args-1.C | 39 ++
gcc/testsuite/g++.dg/gomp/adjust-args-2.C | 51 ++
gcc/testsuite/g++.dg/gomp/dispatch-1.C | 53 ++
gcc/testsuite/g++.dg/gomp/dispatch-2.C | 62 +++
8 files changed, 822 insertions(+), 46 deletions(-)
create mode 100644 gcc/testsuite/g++.dg/gomp/adjust-args-1.C
create mode 100644 gcc/testsuite/g++.dg/gomp/adjust-args-2.C
create mode 100644 gcc/testsuite/g++.dg/gomp/dispatch-1.C
create mode 100644 gcc/testsuite/g++.dg/gomp/dispatch-2.C
diff --git a/gcc/cp/decl.cc b/gcc/cp/decl.cc
index a992d54dc8f..fe53e59add1 100644
--- a/gcc/cp/decl.cc
+++ b/gcc/cp/decl.cc
@@ -8360,6 +8360,33 @@ omp_declare_variant_finalize_one (tree decl, tree attr)
if (!omp_context_selector_matches (ctx))
return true;
TREE_PURPOSE (TREE_VALUE (attr)) = variant;
+
+ for (tree a = ctx; a != NULL_TREE; a = TREE_CHAIN (a))
+ {
+ if (OMP_TSS_CODE (a) == OMP_TRAIT_SET_NEED_DEVICE_PTR)
+ {
+ tree parm_decl = TREE_VALUE (TREE_VALUE (a));
+ bool found_arg = false;
+ for (tree arg = DECL_ARGUMENTS (variant); arg != NULL;
+ arg = TREE_CHAIN (arg))
+ if (DECL_NAME (arg) == DECL_NAME (parm_decl))
+ {
+ DECL_ATTRIBUTES (arg)
+ = tree_cons (get_identifier (
+ "omp declare variant adjust_args "
+ "need_device_ptr"),
+ NULL_TREE, DECL_ATTRIBUTES (arg));
+ found_arg = true;
+ }
+ if (!found_arg)
+ {
+ error_at (varid_loc,
+ "variant %qD does not have a parameter %qD",
+ variant, parm_decl);
+ return true;
+ }
+ }
+ }
}
}
else if (!processing_template_decl)
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 779625144db..d61c37729a7 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -19,6 +19,7 @@ along with GCC; see the file COPYING3. If not see
<http://www.gnu.org/licenses/>. */
#include "config.h"
+#include "omp-selectors.h"
#define INCLUDE_MEMORY
#include "system.h"
#include "coretypes.h"
@@ -2587,7 +2588,7 @@ static cp_ref_qualifier cp_parser_ref_qualifier_opt
static tree cp_parser_tx_qualifier_opt
(cp_parser *);
static tree cp_parser_late_return_type_opt
- (cp_parser *, cp_declarator *, tree &);
+ (cp_parser *, cp_declarator *, tree &, tree);
static tree cp_parser_declarator_id
(cp_parser *, bool);
static tree cp_parser_type_id
@@ -2622,7 +2623,7 @@ static void cp_parser_ctor_initializer_opt_and_function_body
(cp_parser *, bool);
static tree cp_parser_late_parsing_omp_declare_simd
- (cp_parser *, tree);
+ (cp_parser *, tree, tree);
static tree cp_parser_late_parsing_oacc_routine
(cp_parser *, tree);
@@ -24150,7 +24151,7 @@ cp_parser_direct_declarator (cp_parser* parser,
tree requires_clause = NULL_TREE;
late_return
= cp_parser_late_return_type_opt (parser, declarator,
- requires_clause);
+ requires_clause, params);
cp_finalize_omp_declare_simd (parser, &odsd);
@@ -25014,8 +25015,8 @@ parsing_function_declarator ()
function. */
static tree
-cp_parser_late_return_type_opt (cp_parser* parser, cp_declarator *declarator,
- tree& requires_clause)
+cp_parser_late_return_type_opt (cp_parser *parser, cp_declarator *declarator,
+ tree &requires_clause, tree parms)
{
cp_token *token;
tree type = NULL_TREE;
@@ -25051,8 +25052,8 @@ cp_parser_late_return_type_opt (cp_parser* parser, cp_declarator *declarator,
if (declare_simd_p)
declarator->attributes
- = cp_parser_late_parsing_omp_declare_simd (parser,
- declarator->attributes);
+ = cp_parser_late_parsing_omp_declare_simd (parser, declarator->attributes,
+ parms);
if (oacc_routine_p)
declarator->attributes
= cp_parser_late_parsing_oacc_routine (parser,
@@ -38237,6 +38238,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
case 'n':
if (!strcmp ("no_create", p))
result = PRAGMA_OACC_CLAUSE_NO_CREATE;
+ else if (!strcmp ("nocontext", p))
+ result = PRAGMA_OMP_CLAUSE_NOCONTEXT;
else if (!strcmp ("nogroup", p))
result = PRAGMA_OMP_CLAUSE_NOGROUP;
else if (!strcmp ("nohost", p))
@@ -38245,6 +38248,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
else if (!strcmp ("notinbranch", p))
result = PRAGMA_OMP_CLAUSE_NOTINBRANCH;
+ else if (!strcmp ("novariants", p))
+ result = PRAGMA_OMP_CLAUSE_NOVARIANTS;
else if (!strcmp ("nowait", p))
result = PRAGMA_OMP_CLAUSE_NOWAIT;
else if (!strcmp ("num_gangs", p))
@@ -40647,6 +40652,56 @@ cp_parser_omp_clause_thread_limit (cp_parser *parser, tree list,
return c;
}
+/* OpenMP 5.1
+ novariants ( scalar-expression ) */
+
+static tree
+cp_parser_omp_clause_novariants (cp_parser *parser, tree list, location_t loc)
+{
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ tree t = cp_parser_assignment_expression (parser);
+ if (t == error_mark_node || !parens.require_close (parser))
+ cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_NOVARIANTS, "novariants", loc);
+
+ tree c = build_omp_clause (loc, OMP_CLAUSE_NOVARIANTS);
+ OMP_CLAUSE_NOVARIANTS_EXPR (c) = t;
+ OMP_CLAUSE_CHAIN (c) = list;
+
+ return c;
+}
+
+/* OpenMP 5.1
+ nocontext ( scalar-expression ) */
+
+static tree
+cp_parser_omp_clause_nocontext (cp_parser *parser, tree list, location_t loc)
+{
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ tree t = cp_parser_assignment_expression (parser);
+ if (t == error_mark_node || !parens.require_close (parser))
+ cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_NOCONTEXT, "nocontext", loc);
+
+ tree c = build_omp_clause (loc, OMP_CLAUSE_NOCONTEXT);
+ OMP_CLAUSE_NOCONTEXT_EXPR (c) = t;
+ OMP_CLAUSE_CHAIN (c) = list;
+
+ return c;
+}
+
/* OpenMP 4.0:
aligned ( variable-list )
aligned ( variable-list : constant-expression ) */
@@ -42745,6 +42800,16 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses);
c_name = "enter";
break;
+ case PRAGMA_OMP_CLAUSE_NOVARIANTS:
+ clauses = cp_parser_omp_clause_novariants (parser, clauses,
+ token->location);
+ c_name = "novariants";
+ break;
+ case PRAGMA_OMP_CLAUSE_NOCONTEXT:
+ clauses
+ = cp_parser_omp_clause_nocontext (parser, clauses, token->location);
+ c_name = "nocontext";
+ break;
default:
cp_parser_error (parser, "expected an OpenMP clause");
goto saw_error;
@@ -48653,12 +48718,305 @@ cp_parser_omp_assumes (cp_parser *parser, cp_token *pragma_tok)
return false;
}
+/* Parse a function dispatch structured block:
+
+ lvalue-expression = target-call ( [expression-list] );
+ or
+ target-call ( [expression-list] );
+
+ Inspired from cp_parser_assignment_expression and
+ cp_parser_postfix_expression.
+*/
+
+static tree
+cp_parser_omp_dispatch_body (cp_parser *parser)
+{
+ cp_expr expr;
+ cp_id_kind idk = CP_ID_KIND_NONE;
+
+ /* Parse the binary expressions (lvalue-expression or target-call). */
+ expr = cp_parser_binary_expression (parser, false, false, false,
+ PREC_NOT_OPERATOR, NULL);
+ if (TREE_CODE (expr) == CALL_EXPR || TREE_CODE (expr) == ERROR_MARK)
+ return expr;
+
+ /* We have the lvalue, now deal with the assignment. */
+
+ if (!cp_parser_require (parser, CPP_EQ, RT_EQ))
+ return error_mark_node;
+
+ /* Peek at the next token. */
+ cp_token *token = cp_lexer_peek_token (parser->lexer);
+ location_t loc = token->location;
+ location_t start_loc = get_range_from_loc (line_table, loc).m_start;
+
+ /* Parse function name as primary expression. */
+ cp_expr rhs
+ = cp_parser_primary_expression (parser, false, false, false, false, &idk);
+ if (TREE_CODE (rhs) == ERROR_MARK)
+ return rhs;
+
+ /* Keep looping until the postfix-expression is complete. */
+ bool parens_found = false;
+ while (true)
+ {
+ if (idk == CP_ID_KIND_UNQUALIFIED && identifier_p (rhs)
+ && cp_lexer_next_token_is_not (parser->lexer, CPP_OPEN_PAREN))
+ /* It is not a Koenig lookup function call. */
+ rhs = unqualified_name_lookup_error (rhs);
+
+ /* Peek at the next token. */
+ token = cp_lexer_peek_token (parser->lexer);
+
+ switch (token->type)
+ {
+ case CPP_OPEN_PAREN:
+ /* postfix-expression ( expression-list [opt] ) */
+ {
+ if (parens_found)
+ {
+ cp_parser_error (
+ parser,
+ "only one function call is allowed in a dispatch construct");
+ return error_mark_node;
+ }
+ parens_found = true;
+
+ bool koenig_p;
+ tsubst_flags_t complain = complain_flags (false);
+ vec<tree, va_gc> *args;
+ location_t close_paren_loc = UNKNOWN_LOCATION;
+ location_t combined_loc = UNKNOWN_LOCATION;
+
+ args = (cp_parser_parenthesized_expression_list (
+ parser, non_attr,
+ /*cast_p=*/false, /*allow_expansion_p=*/true,
+ /*non_constant_p=*/NULL,
+ /*close_paren_loc=*/&close_paren_loc,
+ /*wrap_locations_p=*/true));
+
+ if (args == NULL)
+ {
+ rhs = error_mark_node;
+ break;
+ }
+
+ koenig_p = false;
+ if (idk == CP_ID_KIND_UNQUALIFIED || idk == CP_ID_KIND_TEMPLATE_ID)
+ {
+ if (identifier_p (rhs)
+ /* In C++20, we may need to perform ADL for a template
+ name. */
+ || (TREE_CODE (rhs) == TEMPLATE_ID_EXPR
+ && identifier_p (TREE_OPERAND (rhs, 0))))
+ {
+ if (!args->is_empty ())
+ {
+ koenig_p = true;
+ if (!any_type_dependent_arguments_p (args))
+ rhs = perform_koenig_lookup (rhs, args, complain);
+ }
+ else
+ rhs = unqualified_fn_lookup_error (rhs);
+ }
+ /* We do not perform argument-dependent lookup if
+ normal lookup finds a non-function, in accordance
+ with the expected resolution of DR 218. */
+ else if (!args->is_empty () && is_overloaded_fn (rhs))
+ {
+ /* Do not do argument dependent lookup if regular
+ lookup finds a member function or a block-scope
+ function declaration. [basic.lookup.argdep]/3 */
+ bool do_adl_p = true;
+ tree fns = get_fns (rhs);
+ for (lkp_iterator iter (fns); iter; ++iter)
+ {
+ tree fn = STRIP_TEMPLATE (*iter);
+ if ((TREE_CODE (fn) == USING_DECL
+ && DECL_DEPENDENT_P (fn))
+ || DECL_FUNCTION_MEMBER_P (fn)
+ || DECL_LOCAL_DECL_P (fn))
+ {
+ do_adl_p = false;
+ break;
+ }
+ }
+
+ if (do_adl_p)
+ {
+ koenig_p = true;
+ if (!any_type_dependent_arguments_p (args))
+ rhs = perform_koenig_lookup (rhs, args, complain);
+ }
+ }
+ }
+
+ /* Temporarily set input_location to the combined location
+ with call expression range, as e.g. build_out_target_exprs
+ called from convert_default_arg relies on input_location,
+ so updating it only when the call is fully built results
+ in inconsistencies between location handling in templates
+ and outside of templates. */
+ if (close_paren_loc != UNKNOWN_LOCATION)
+ combined_loc
+ = make_location (token->location, start_loc, close_paren_loc);
+ iloc_sentinel ils (combined_loc);
+
+ if (TREE_CODE (rhs) == COMPONENT_REF)
+ {
+ tree instance = TREE_OPERAND (rhs, 0);
+ tree fn = TREE_OPERAND (rhs, 1);
+
+ if (processing_template_decl
+ && (type_dependent_object_expression_p (instance)
+ || (!BASELINK_P (fn) && TREE_CODE (fn) != FIELD_DECL)
+ || type_dependent_expression_p (fn)
+ || any_type_dependent_arguments_p (args)))
+ {
+ maybe_generic_this_capture (instance, fn);
+ rhs = build_min_nt_call_vec (rhs, args);
+ }
+ else if (BASELINK_P (fn))
+ {
+ rhs
+ = (build_new_method_call (instance, fn, &args, NULL_TREE,
+ (idk == CP_ID_KIND_QUALIFIED
+ ? LOOKUP_NORMAL
+ | LOOKUP_NONVIRTUAL
+ : LOOKUP_NORMAL),
+ /*fn_p=*/NULL, complain));
+ }
+ else
+ rhs = finish_call_expr (rhs, &args,
+ /*disallow_virtual=*/false,
+ /*koenig_p=*/false, complain);
+ }
+ else if (TREE_CODE (rhs) == OFFSET_REF
+ || TREE_CODE (rhs) == MEMBER_REF
+ || TREE_CODE (rhs) == DOTSTAR_EXPR)
+ rhs = (build_offset_ref_call_from_tree (rhs, &args, complain));
+ else if (idk == CP_ID_KIND_QUALIFIED)
+ /* A call to a static class member, or a namespace-scope
+ function. */
+ rhs = finish_call_expr (rhs, &args,
+ /*disallow_virtual=*/true, koenig_p,
+ complain);
+ else
+ /* All other function calls. */
+ {
+ if (DECL_P (rhs) && parser->omp_for_parse_state
+ && parser->omp_for_parse_state->in_intervening_code
+ && omp_runtime_api_call (rhs))
+ {
+ error_at (loc, "calls to the OpenMP runtime API are "
+ "not permitted in intervening code");
+ parser->omp_for_parse_state->fail = true;
+ }
+ rhs = finish_call_expr (rhs, &args,
+ /*disallow_virtual=*/false, koenig_p,
+ complain);
+ }
+ if (close_paren_loc != UNKNOWN_LOCATION)
+ rhs.set_location (combined_loc);
+
+ /* The expr is certainly no longer an id. */
+ idk = CP_ID_KIND_NONE;
+
+ release_tree_vector (args);
+ }
+ break;
+
+ case CPP_DOT:
+ case CPP_DEREF:
+ /* postfix-expression . template [opt] id-expression
+ postfix-expression . pseudo-destructor-name
+ postfix-expression -> template [opt] id-expression
+ postfix-expression -> pseudo-destructor-name */
+
+ /* Consume the `.' or `->' operator. */
+ cp_lexer_consume_token (parser->lexer);
+
+ rhs = cp_parser_postfix_dot_deref_expression (parser, token->type,
+ rhs, false, &idk, loc);
+
+ break;
+
+ default:
+ goto finish;
+ }
+ }
+finish:
+ if (!parens_found)
+ {
+ cp_parser_error (parser, "expected %<(%>");
+ return error_mark_node;
+ }
+
+ /* Build the assignment expression. Its default
+ location:
+ LHS = RHS
+ ~~~~^~~~~
+ is the location of the '=' token as the
+ caret, ranging from the start of the lhs to the
+ end of the rhs. */
+ loc = make_location (loc, expr.get_start (), rhs.get_finish ());
+ expr
+ = cp_build_modify_expr (loc, expr, NOP_EXPR, rhs, complain_flags (false));
+
+ return expr;
+}
+
+/* OpenMP 5.1:
+ # pragma omp dispatch dispatch-clause[optseq] new-line
+ expression-stmt
+
+ LOC is the location of the #pragma.
+*/
+
+#define OMP_DISPATCH_CLAUSE_MASK \
+ ((OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT))
+
+static tree
+cp_parser_omp_dispatch (cp_parser *parser, cp_token *pragma_tok)
+{
+ location_t loc = cp_lexer_peek_token (parser->lexer)->location;
+ tree stmt = make_node (OMP_DISPATCH);
+ SET_EXPR_LOCATION (stmt, loc);
+ TREE_TYPE (stmt) = void_type_node;
+
+ OMP_DISPATCH_CLAUSES (stmt)
+ = cp_parser_omp_all_clauses (parser, OMP_DISPATCH_CLAUSE_MASK,
+ "#pragma omp dispatch", pragma_tok);
+
+ // Parse expression statement
+ loc = cp_lexer_peek_token (parser->lexer)->location;
+ tree dispatch_body = cp_parser_omp_dispatch_body (parser);
+ if (dispatch_body == error_mark_node)
+ {
+ inform (loc,
+ "%<#pragma omp dispatch%> must be followed by a direct function "
+ "call with optional assignment");
+ cp_parser_skip_to_end_of_block_or_statement (parser);
+ return NULL_TREE;
+ }
+
+ cp_parser_consume_semicolon_at_end_of_statement (parser);
+ OMP_DISPATCH_BODY (stmt) = dispatch_body;
+
+ return add_stmt (stmt);
+}
+
/* Finalize #pragma omp declare variant after a fndecl has been parsed, and put
that into "omp declare variant base" attribute. */
static tree
cp_finish_omp_declare_variant (cp_parser *parser, cp_token *pragma_tok,
- tree attrs)
+ tree attrs, tree parms)
{
matching_parens parens;
if (!parens.require_open (parser))
@@ -48716,45 +49074,197 @@ cp_finish_omp_declare_variant (cp_parser *parser, cp_token *pragma_tok,
location_t finish_loc = get_finish (varid.get_location ());
location_t varid_loc = make_location (caret_loc, start_loc, finish_loc);
- if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)
- && cp_lexer_nth_token_is (parser->lexer, 2, CPP_NAME))
- cp_lexer_consume_token (parser->lexer);
+ vec<tree> adjust_args_list = vNULL;
+ bool has_match = false, has_adjust_args = false;
+ location_t adjust_args_loc;
+ tree need_device_ptr_list = NULL_TREE, *need_device_ptr_chain_p = NULL;
- const char *clause = "";
- location_t match_loc = cp_lexer_peek_token (parser->lexer)->location;
- if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
- clause = IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value);
- if (strcmp (clause, "match"))
+ do
{
- cp_parser_error (parser, "expected %<match%>");
- goto fail;
+ if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)
+ && cp_lexer_nth_token_is (parser->lexer, 2, CPP_NAME))
+ cp_lexer_consume_token (parser->lexer);
+
+ const char *clause = "";
+ location_t match_loc = cp_lexer_peek_token (parser->lexer)->location;
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ clause
+ = IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value);
+
+ enum clause
+ {
+ match,
+ adjust_args
+ } ccode;
+
+ if (strcmp (clause, "match") == 0)
+ ccode = match;
+ else if (strcmp (clause, "adjust_args") == 0)
+ {
+ ccode = adjust_args;
+ adjust_args_loc = match_loc;
+ }
+ else
+ {
+ cp_parser_error (parser, "expected %<match%> or %<adjust_args%>");
+ goto fail;
+ }
+
+ cp_lexer_consume_token (parser->lexer);
+
+ if (!parens.require_open (parser))
+ goto fail;
+
+ if (ccode == match)
+ {
+ has_match = true;
+ tree ctx
+ = cp_parser_omp_context_selector_specification (parser, true);
+ if (ctx == error_mark_node)
+ goto fail;
+ ctx = omp_check_context_selector (match_loc, ctx);
+ if (ctx != error_mark_node && variant != error_mark_node)
+ {
+ tree match_loc_node
+ = maybe_wrap_with_location (integer_zero_node, match_loc);
+ tree loc_node
+ = maybe_wrap_with_location (integer_zero_node, varid_loc);
+ loc_node
+ = tree_cons (match_loc_node,
+ build_int_cst (integer_type_node, idk),
+ build_tree_list (loc_node, integer_zero_node));
+ attrs = tree_cons (get_identifier ("omp declare variant base"),
+ tree_cons (variant, ctx, loc_node), attrs);
+ if (processing_template_decl)
+ ATTR_IS_DEPENDENT (attrs) = 1;
+ }
+ if (!parens.require_close (parser))
+ goto fail;
+ }
+ else if (ccode == adjust_args)
+ {
+ has_adjust_args = true;
+ cp_token *adjust_op_tok = cp_lexer_peek_token (parser->lexer);
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
+ && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+ {
+ const char *p = IDENTIFIER_POINTER (adjust_op_tok->u.value);
+ if (strcmp (p, "need_device_ptr") == 0
+ || strcmp (p, "nothing") == 0)
+ {
+ cp_lexer_consume_token (parser->lexer); // need_device_ptr
+ cp_lexer_consume_token (parser->lexer); // :
+ location_t arg_loc
+ = cp_lexer_peek_token (parser->lexer)->location;
+
+ tree arg;
+ tree list
+ = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_ERROR,
+ NULL_TREE, NULL);
+
+ for (tree c = list; c != NULL_TREE; c = TREE_CHAIN (c))
+ {
+ tree decl = TREE_PURPOSE (c);
+ int idx;
+ for (arg = parms, idx = 0; arg != NULL;
+ arg = TREE_CHAIN (arg), idx++)
+ if (TREE_VALUE (arg) == decl)
+ break;
+ if (arg == NULL_TREE)
+ {
+ error_at (arg_loc, "%qD is not a function argument",
+ decl);
+ continue;
+ }
+ arg = TREE_VALUE (arg);
+ if (adjust_args_list.contains (arg))
+ {
+ error_at (arg_loc, "%qD is specified more than once",
+ decl);
+ continue;
+ }
+ if (strcmp (p, "need_device_ptr") == 0)
+ {
+ bool is_ptr_or_template
+ = TEMPLATE_PARM_P (TREE_TYPE (arg))
+ || POINTER_TYPE_P (TREE_TYPE (arg));
+ if (!is_ptr_or_template)
+ {
+ error_at (arg_loc, "%qD is not a C pointer",
+ decl);
+ continue;
+ }
+ }
+ adjust_args_list.safe_push (arg);
+ if (strcmp (p, "need_device_ptr") == 0)
+ {
+ tree attr = tree_cons (NULL_TREE, TREE_PURPOSE (c),
+ NULL_TREE);
+ if (need_device_ptr_list == NULL_TREE)
+ {
+ gcc_assert (need_device_ptr_chain_p == NULL);
+ need_device_ptr_list = attr;
+ }
+ else
+ *need_device_ptr_chain_p = attr;
+ need_device_ptr_chain_p = &TREE_CHAIN (attr);
+ }
+ }
+ }
+ else
+ {
+ error_at (adjust_op_tok->location,
+ "expected %<nothing%> or %<need_device_ptr%>");
+ goto fail;
+ }
+ }
+ else
+ {
+ error_at (adjust_op_tok->location,
+ "expected %<nothing%> or %<need_device_ptr%> followed "
+ "by %<:%>");
+ goto fail;
+ }
+ }
+ } while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL));
+
+ if (has_adjust_args)
+ {
+ if (!has_match)
+ {
+ error_at (
+ adjust_args_loc,
+ "an %<adjust_args%> clause can only be specified if the "
+ "%<dispatch%> selector of the construct selector set appears "
+ "in the %<match%> clause");
+ }
+ else
+ {
+ tree ctx = TREE_VALUE (TREE_VALUE (attrs));
+ if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT,
+ OMP_TRAIT_CONSTRUCT_DISPATCH))
+ error_at (
+ adjust_args_loc,
+ "an %<adjust_args%> clause can only be specified if the "
+ "%<dispatch%> selector of the construct selector set appears "
+ "in the %<match%> clause");
+ }
}
- cp_lexer_consume_token (parser->lexer);
-
- if (!parens.require_open (parser))
- goto fail;
-
- tree ctx = cp_parser_omp_context_selector_specification (parser, true);
- if (ctx == error_mark_node)
- goto fail;
- ctx = omp_check_context_selector (match_loc, ctx);
- if (ctx != error_mark_node && variant != error_mark_node)
+ if (need_device_ptr_list)
{
- tree match_loc_node = maybe_wrap_with_location (integer_zero_node,
- match_loc);
- tree loc_node = maybe_wrap_with_location (integer_zero_node, varid_loc);
- loc_node = tree_cons (match_loc_node,
- build_int_cst (integer_type_node, idk),
- build_tree_list (loc_node, integer_zero_node));
- attrs = tree_cons (get_identifier ("omp declare variant base"),
- tree_cons (variant, ctx, loc_node), attrs);
- if (processing_template_decl)
- ATTR_IS_DEPENDENT (attrs) = 1;
+ // We might not have DECL_ARGUMENTS for the variant yet. So we store the
+ // need_device_ptr list in the base function attribute beside the context
+ // selector.
+ gcc_assert (TREE_PURPOSE (attrs)
+ == get_identifier ("omp declare variant base"));
+ gcc_assert (TREE_PURPOSE (TREE_VALUE (attrs)) == variant);
+ TREE_VALUE (TREE_VALUE (attrs))
+ = make_trait_set_selector (OMP_TRAIT_SET_NEED_DEVICE_PTR,
+ need_device_ptr_list,
+ TREE_VALUE (TREE_VALUE (attrs)));
}
- parens.require_close (parser);
- cp_parser_skip_to_pragma_eol (parser, pragma_tok);
return attrs;
}
@@ -48763,7 +49273,8 @@ cp_finish_omp_declare_variant (cp_parser *parser, cp_token *pragma_tok,
been parsed, and put that into "omp declare simd" attribute. */
static tree
-cp_parser_late_parsing_omp_declare_simd (cp_parser *parser, tree attrs)
+cp_parser_late_parsing_omp_declare_simd (cp_parser *parser, tree attrs,
+ tree parms)
{
struct cp_token_cache *ce;
cp_omp_declare_simd_data *data = parser->omp_declare_simd;
@@ -48807,7 +49318,7 @@ cp_parser_late_parsing_omp_declare_simd (cp_parser *parser, tree attrs)
{
gcc_assert (strcmp (kind, "variant") == 0);
attrs
- = cp_finish_omp_declare_variant (parser, pragma_tok, attrs);
+ = cp_finish_omp_declare_variant (parser, pragma_tok, attrs, parms);
}
cp_parser_pop_lexer (parser);
}
@@ -48938,9 +49449,8 @@ cp_parser_late_parsing_omp_declare_simd (cp_parser *parser, tree attrs)
else
{
gcc_assert (strcmp (kind, "variant") == 0);
- attrs
- = cp_finish_omp_declare_variant (parser, pragma_tok,
- attrs);
+ attrs = cp_finish_omp_declare_variant (parser, pragma_tok,
+ attrs, parms);
}
gcc_assert (parser->lexer != lexer);
vec_safe_truncate (lexer->buffer, 0);
@@ -49790,7 +50300,11 @@ cp_parser_omp_declare_reduction (cp_parser *parser, cp_token *pragma_tok,
#pragma omp declare target new-line
OpenMP 5.0
- #pragma omp declare variant (identifier) match (context-selector) */
+ #pragma omp declare variant (identifier) match (context-selector)
+
+ OpenMP 5.1
+ #pragma omp declare variant (identifier) match (context-selector) \
+ adjust_args (adjust-op:argument-list) */
static bool
cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok,
@@ -50645,6 +51159,9 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
case PRAGMA_OMP_ASSUME:
cp_parser_omp_assume (parser, pragma_tok, if_p);
return;
+ case PRAGMA_OMP_DISPATCH:
+ stmt = cp_parser_omp_dispatch (parser, pragma_tok);
+ break;
default:
gcc_unreachable ();
}
@@ -51339,6 +51856,10 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context, bool *if_p)
"%<#pragma omp sections%> construct");
break;
+ case PRAGMA_OMP_DISPATCH:
+ cp_parser_omp_dispatch (parser, pragma_tok);
+ return true;
+
case PRAGMA_IVDEP:
case PRAGMA_UNROLL:
case PRAGMA_NOVECTOR:
diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc
index dfce1b3c359..b8ae4707258 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -12014,6 +12014,9 @@ tsubst_attribute (tree t, tree *decl_p, tree args,
for (tree tss = ctx; tss; tss = TREE_CHAIN (tss))
{
enum omp_tss_code set = OMP_TSS_CODE (tss);
+ if (set == OMP_TRAIT_SET_NEED_DEVICE_PTR)
+ continue;
+
tree selectors = NULL_TREE;
for (tree ts = OMP_TSS_TRAIT_SELECTORS (tss); ts;
ts = TREE_CHAIN (ts))
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index f90c304a65b..248aee4bd5b 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -7613,6 +7613,26 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
OMP_CLAUSE_FINAL_EXPR (c) = t;
break;
+ case OMP_CLAUSE_NOCONTEXT:
+ t = OMP_CLAUSE_NOCONTEXT_EXPR (c);
+ t = maybe_convert_cond (t);
+ if (t == error_mark_node)
+ remove = true;
+ else if (!processing_template_decl)
+ t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
+ OMP_CLAUSE_NOCONTEXT_EXPR (c) = t;
+ break;
+
+ case OMP_CLAUSE_NOVARIANTS:
+ t = OMP_CLAUSE_NOVARIANTS_EXPR (c);
+ t = maybe_convert_cond (t);
+ if (t == error_mark_node)
+ remove = true;
+ else if (!processing_template_decl)
+ t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
+ OMP_CLAUSE_NOVARIANTS_EXPR (c) = t;
+ break;
+
case OMP_CLAUSE_GANG:
/* Operand 1 is the gang static: argument. */
t = OMP_CLAUSE_OPERAND (c, 1);
diff --git a/gcc/testsuite/g++.dg/gomp/adjust-args-1.C b/gcc/testsuite/g++.dg/gomp/adjust-args-1.C
new file mode 100644
index 00000000000..1c6dd8ac97b
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/adjust-args-1.C
@@ -0,0 +1,39 @@
+/* Test parsing of OMP clause adjust_args */
+/* { dg-do compile } */
+
+int b;
+
+int f0 (void *a);
+int g (void *a);
+int f1 (int);
+
+#pragma omp declare variant (f0) match (construct={target}) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */
+int f2 (void *a);
+#pragma omp declare variant (f0) match (construct={dispatch,target}) adjust_args (need_device_ptr: a) /* { dg-error "'int f0.void..' used as a variant with incompatible 'construct' selector sets" } */
+int f2a (void *a);
+#pragma omp declare variant (f0) match (construct={target,dispatch}) adjust_args (need_device_ptr: a) /* { dg-error "'int f0.void..' used as a variant with incompatible 'construct' selector sets" } */
+int f2b (void *a);
+#pragma omp declare variant (f0) match (construct={dispatch},device={arch(gcn)}) adjust_args (need_device_ptr: a) /* { dg-error "'int f0.void..' used as a variant with incompatible 'construct' selector sets" } */
+int f2c (void *a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (other: a) /* { dg-error "expected 'nothing' or 'need_device_ptr'" } */
+int f3 (int a);
+#pragma omp declare variant (f0) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */
+int f4 (void *a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args () /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */
+int f5 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing) /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */
+int f6 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing:) /* { dg-error "expected unqualified-id before '\\)' token" } */
+int f7 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: z) /* { dg-error "'z' has not been declared" } */
+int f8 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: a) /* { dg-error "'a' is not a C pointer" } */
+int f9 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (nothing: a) /* { dg-error "'a' is specified more than once" } */
+int f10 (int a);
+#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: a) /* { dg-error "'a' is specified more than once" } */
+int f11 (void *a);
+#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: b) /* { dg-error "'b' is not a function argument" } */
+int f12 (void *a);
+#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: this) /* { dg-error "expected unqualified-id before 'this'" } */
+int f13 (void *a);
diff --git a/gcc/testsuite/g++.dg/gomp/adjust-args-2.C b/gcc/testsuite/g++.dg/gomp/adjust-args-2.C
new file mode 100644
index 00000000000..a78f06ec193
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/adjust-args-2.C
@@ -0,0 +1,51 @@
+struct S {
+ int a;
+ int g (const void *b);
+ #pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: b)
+ int f0(const void *b);
+ int operator()() { return a; }
+ bool operator!() { return !a; }
+};
+
+template <typename T>
+T f0(T a, T *b);
+
+#pragma omp declare variant (f0) match (construct={dispatch}) adjust_args (need_device_ptr: a, b)
+template <typename T>
+T f1(T a, T *b);
+
+namespace N {
+ class C{
+ public:
+ void g(C *c);
+ #pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: c)
+ void f0(C *c);
+ };
+ void g(C *c);
+ #pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: c)
+ void f0(C *c);
+}
+
+#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: c)
+void f3(N::C *c);
+void f4(S *&s);
+#pragma omp declare variant (f4) match (construct={dispatch}) adjust_args (need_device_ptr: s)
+void f5(S *&s);
+
+void test() {
+ S s, *sp;
+ N::C c;
+ int *a, b;
+ #pragma omp dispatch
+ s.f0(a);
+ #pragma omp dispatch
+ f1(b, a);
+ #pragma omp dispatch
+ c.f0(&c);
+ #pragma omp dispatch
+ N::f0(&c);
+ #pragma omp dispatch
+ f3(&c);
+ #pragma omp dispatch
+ f5(sp);
+}
diff --git a/gcc/testsuite/g++.dg/gomp/dispatch-1.C b/gcc/testsuite/g++.dg/gomp/dispatch-1.C
new file mode 100644
index 00000000000..fb467afcd85
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/dispatch-1.C
@@ -0,0 +1,53 @@
+struct S {
+ int a;
+ void f0(double);
+ int operator()() { return a; }
+ bool operator!() { return !a; }
+};
+
+int f0(int);
+template <typename T>
+T f1(T a, T b);
+void (*f2)(void);
+
+namespace N {
+ class C{};
+ void f0(C);
+ int a;
+}
+
+int test() {
+ int result;
+ double d = 5.0;
+ N::C c;
+ S s;
+ S* sp = &s;
+ int &r = result;
+ #pragma omp dispatch
+ result = f0(5);
+ #pragma omp dispatch
+ r = f0(5);
+ #pragma omp dispatch
+ N::a = ::f0(5);
+ #pragma omp dispatch
+ sp->a = f1<int>(5, 10);
+ #pragma omp dispatch
+ s.a = f1(5, 10);
+ #pragma omp dispatch
+ f2();
+ #pragma omp dispatch
+ N::f0(c);
+ #pragma omp dispatch
+ f0(c);
+ #pragma omp dispatch
+ s.f0(d);
+ #pragma omp dispatch
+ sp->f0(d);
+ #pragma omp dispatch
+ sp->f0(d);
+ #pragma omp dispatch
+ s();
+ #pragma omp dispatch
+ !s;
+ return result;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/dispatch-2.C b/gcc/testsuite/g++.dg/gomp/dispatch-2.C
new file mode 100644
index 00000000000..1bc304e005e
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/dispatch-2.C
@@ -0,0 +1,62 @@
+/* Test parsing of #pragma omp dispatch */
+/* { dg-do compile } */
+
+struct S {
+ int a;
+ int b;
+ virtual int f (double);
+};
+
+int f0 (int);
+
+void f1 (void)
+{
+ int a, b;
+ double x;
+ int arr[1];
+ S s;
+
+#pragma omp dispatch
+ int c = f0 (a); /* { dg-error "expected primary-expression before 'int'" } */
+#pragma omp dispatch
+ int f2 (int d); /* { dg-error "expected primary-expression before 'int'" } */
+#pragma omp dispatch
+ a = b; /* { dg-error "expected '\\(' before ';' token" } */
+#pragma omp dispatch
+ s.a = f0(a) + b; /* { dg-error "expected ';' before '\\+' token" } */
+#pragma omp dispatch
+ b = !f0(a); /* { dg-error "expected primary-expression before '!' token" } */
+#pragma omp dispatch
+ s.b += f0(s.a); /* { dg-error "expected '=' before '\\+=' token" } */
+#pragma omp dispatch
+#pragma omp threadprivate(a) /* { dg-error "'#pragma' is not allowed here" } */
+ a = f0(b);
+#pragma omp dispatch
+ a = s.f(x); /* { dg-error "'f' is a virtual function but only a direct call is allowed in a dispatch construct" } */
+
+#pragma omp dispatch nocontext(s) /* { dg-error "could not convert 's' from 'S' to 'bool'" } */
+ f0(a);
+#pragma omp dispatch nocontext(a, b) /* { dg-error "expected '\\)' before ','" } */
+ f0(a);
+#pragma omp dispatch nocontext(a) nocontext(b) /* { dg-error "too many 'nocontext' clauses" } */
+ f0(a);
+#pragma omp dispatch novariants(s) /* { dg-error "could not convert 's' from 'S' to 'bool'" } */
+ f0(a);
+#pragma omp dispatch novariants(a, b) /* { dg-error "expected '\\)' before ','" } */
+ f0(a);
+#pragma omp dispatch novariants(a) novariants(b) /* { dg-error "too many 'novariants' clauses" } */
+ f0(a);
+#pragma omp dispatch nowait nowait /* { dg-error "too many 'nowait' clauses" } */
+ f0(a);
+#pragma omp dispatch device(x) /* { dg-error "'device' id must be integral" } */
+ f0(a);
+#pragma omp dispatch device(arr) /* { dg-error "'device' id must be integral" } */
+ f0(a);
+#pragma omp dispatch is_device_ptr(x) /* { dg-error "'is_device_ptr' variable is neither a pointer, nor an array nor reference to pointer" } */
+ f0(a);
+#pragma omp dispatch is_device_ptr(&x) /* { dg-error "expected unqualified-id before '&' token" } */
+ f0(a);
+#pragma omp dispatch depend(inout: s.f) /* { dg-error "'s.S::f' is not lvalue expression nor array section in 'depend' clause" } */
+ f0(a);
+
+}
--
2.45.1
^ permalink raw reply [flat|nested] 10+ messages in thread* [PATCH 5/7] OpenMP: common C/C++ testcases for dispatch + adjust_args
2024-05-27 11:54 [PATCH 0/7] OpenMP: dispatch + adjust_args support Paul-Antoine Arras
` (3 preceding siblings ...)
2024-05-27 11:54 ` [PATCH 4/7] OpenMP: C++ " Paul-Antoine Arras
@ 2024-05-27 11:54 ` Paul-Antoine Arras
2024-05-27 11:54 ` [PATCH 6/7] OpenMP: Fortran front-end support " Paul-Antoine Arras
2024-05-27 11:54 ` [PATCH 7/7] OpenMP: update documentation for dispatch and adjust_args Paul-Antoine Arras
6 siblings, 0 replies; 10+ messages in thread
From: Paul-Antoine Arras @ 2024-05-27 11:54 UTC (permalink / raw)
To: gcc-patches; +Cc: Paul-Antoine Arras
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/declare-variant-2.c: Adjust dg-error directives.
* c-c++-common/gomp/adjust-args-1.c: New test.
* c-c++-common/gomp/adjust-args-2.c: New test.
* c-c++-common/gomp/dispatch-1.c: New test.
* c-c++-common/gomp/dispatch-2.c: New test.
* c-c++-common/gomp/dispatch-3.c: New test.
* c-c++-common/gomp/dispatch-4.c: New test.
* c-c++-common/gomp/dispatch-5.c: New test.
* c-c++-common/gomp/dispatch-6.c: New test.
* c-c++-common/gomp/dispatch-7.c: New test.
---
.../c-c++-common/gomp/adjust-args-1.c | 30 +++++++++
.../c-c++-common/gomp/adjust-args-2.c | 31 +++++++++
.../c-c++-common/gomp/declare-variant-2.c | 4 +-
gcc/testsuite/c-c++-common/gomp/dispatch-1.c | 65 +++++++++++++++++++
gcc/testsuite/c-c++-common/gomp/dispatch-2.c | 28 ++++++++
gcc/testsuite/c-c++-common/gomp/dispatch-3.c | 15 +++++
gcc/testsuite/c-c++-common/gomp/dispatch-4.c | 18 +++++
gcc/testsuite/c-c++-common/gomp/dispatch-5.c | 26 ++++++++
gcc/testsuite/c-c++-common/gomp/dispatch-6.c | 19 ++++++
gcc/testsuite/c-c++-common/gomp/dispatch-7.c | 28 ++++++++
10 files changed, 262 insertions(+), 2 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/adjust-args-1.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/adjust-args-2.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-1.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-2.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-3.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-4.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-5.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-6.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/dispatch-7.c
diff --git a/gcc/testsuite/c-c++-common/gomp/adjust-args-1.c b/gcc/testsuite/c-c++-common/gomp/adjust-args-1.c
new file mode 100644
index 00000000000..728abe62092
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/adjust-args-1.c
@@ -0,0 +1,30 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+int f (int a, void *b, float c[2]);
+
+#pragma omp declare variant (f) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: b, c)
+int f0 (int a, void *b, float c[2]);
+#pragma omp declare variant (f) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: b) adjust_args (need_device_ptr: c)
+int f1 (int a, void *b, float c[2]);
+
+int test () {
+ int a;
+ void *b;
+ float c[2];
+ struct {int a;} s;
+
+ s.a = f0 (a, b, c);
+ #pragma omp dispatch
+ s.a = f0 (a, b, c);
+
+ f1 (a, b, c);
+ #pragma omp dispatch
+ s.a = f1 (a, b, c);
+
+ return s.a;
+}
+
+/* { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&c, D\.\[0-9]+\\);" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(b, D\.\[0-9]+\\);" 2 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/adjust-args-2.c b/gcc/testsuite/c-c++-common/gomp/adjust-args-2.c
new file mode 100644
index 00000000000..e36d93a01d9
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/adjust-args-2.c
@@ -0,0 +1,31 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+int f (int a, void *b, float c[2]);
+
+#pragma omp declare variant (f) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: b, c)
+int f0 (int a, void *b, float c[2]);
+#pragma omp declare variant (f) adjust_args (need_device_ptr: b, c) match (construct={dispatch}) adjust_args (nothing: a)
+int f1 (int a, void *b, float c[2]);
+
+void test () {
+ int a;
+ void *b;
+ float c[2];
+
+ #pragma omp dispatch
+ f0 (a, b, c);
+
+ #pragma omp dispatch device (-4852)
+ f0 (a, b, c);
+
+ #pragma omp dispatch device (a + a)
+ f0 (a, b, c);
+}
+
+/* { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&c, D\.\[0-9]+\\);" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(b, D\.\[0-9]+\\);" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&c, -4852\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(b, -4852\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp dispatch device\\(-4852\\)" 1 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c b/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c
index 05e485ef6a8..50d9b2dcf4b 100644
--- a/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/declare-variant-2.c
@@ -8,9 +8,9 @@ void f3 (void);
void f4 (void);
#pragma omp declare variant match(user={condition(0)}) /* { dg-error "expected '\\(' before 'match'" } */
void f5 (void);
-#pragma omp declare variant (f1) /* { dg-error "expected 'match' before end of line" } */
+#pragma omp declare variant (f1) /* { dg-error "expected 'match' or 'adjust_args' before end of line" } */
void f6 (void);
-#pragma omp declare variant (f1) simd /* { dg-error "expected 'match' before 'simd'" } */
+#pragma omp declare variant (f1) simd /* { dg-error "expected 'match' or 'adjust_args' before 'simd'" } */
void f7 (void);
#pragma omp declare variant (f1) match /* { dg-error "expected '\\(' before end of line" } */
void f8 (void);
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-1.c b/gcc/testsuite/c-c++-common/gomp/dispatch-1.c
new file mode 100644
index 00000000000..e77b2f3ecf6
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-1.c
@@ -0,0 +1,65 @@
+#include <stdint.h>
+
+int f0 (int, long, double);
+void f2 (void);
+int f3 (void);
+void (*f4) (void);
+
+void f1 (void)
+{
+ int a, c;
+ long b;
+ double x;
+ struct {int a; float b; short c;} s, *sp;
+ int arr[3];
+
+#pragma omp dispatch
+ c = f0 (a, b, x);
+#pragma omp dispatch
+ x = f0 (a * 4, 2 - b, x * x);
+#pragma omp dispatch
+ s.a = f0 (a, sp->c, x);
+#pragma omp dispatch
+ sp->c = f0 (s.a - 2, b / 3, x * 5);
+#pragma omp dispatch
+ arr[0] = f0 (arr[1], !b, arr[2]);
+#pragma omp dispatch
+ (*sp).c = f0 (s.a, b, x);
+#pragma omp dispatch
+ sp->b = f0 (s.a++, b % 4, --x);
+#pragma omp dispatch
+ f0 (f3(), b, s.b);
+#pragma omp dispatch
+ f2 ();
+#pragma omp dispatch
+ f4 ();
+
+#pragma omp dispatch nocontext(sp->a * x + arr[2])
+ f2 ();
+#pragma omp dispatch nocontext(arr - (intptr_t)(x / s.b))
+ f2 ();
+#pragma omp dispatch nocontext(x == s.c || b != c)
+ f2 ();
+#pragma omp dispatch novariants(b << sp->c)
+ f2 ();
+#pragma omp dispatch novariants(!arr | s.a)
+ f2 ();
+#pragma omp dispatch novariants(s.c ? f3() : a & c)
+ f2 ();
+#pragma omp dispatch nowait
+ f2 ();
+#pragma omp dispatch device(-25373654)
+ f2 ();
+#pragma omp dispatch device(b * (int)(x - sp->b))
+ f2 ();
+#pragma omp dispatch is_device_ptr(arr)
+ f2 ();
+#pragma omp dispatch is_device_ptr(sp)
+ f2 ();
+#pragma omp dispatch depend(inout: sp)
+ f2 ();
+#pragma omp dispatch depend(inoutset: arr[:2])
+ f2 ();
+#pragma omp dispatch depend(mutexinoutset: arr)
+ f2 ();
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-2.c b/gcc/testsuite/c-c++-common/gomp/dispatch-2.c
new file mode 100644
index 00000000000..24ab9545b73
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-2.c
@@ -0,0 +1,28 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+int f0 (void);
+int f1 (void);
+#pragma omp declare variant (f0) match (construct={dispatch})
+#pragma omp declare variant (f1) match (implementation={vendor(gnu)})
+int f2 (void);
+
+int test (void)
+{
+ int a;
+#pragma omp dispatch
+ a = f2 ();
+#pragma omp dispatch novariants(1)
+ a = f2 ();
+#pragma omp dispatch novariants(0)
+ a = f2 ();
+#pragma omp dispatch nocontext(1)
+ a = f2 ();
+#pragma omp dispatch nocontext(0)
+ a = f2 ();
+ return a;
+}
+
+/* { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 3 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-3.c b/gcc/testsuite/c-c++-common/gomp/dispatch-3.c
new file mode 100644
index 00000000000..319f73a84d8
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-3.c
@@ -0,0 +1,15 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void f2 (void);
+
+void test (void)
+{
+#pragma omp dispatch /* { dg-final { scan-tree-dump-times "#pragma omp task if\\(0\\)" 1 "gimple" } } */
+ f2 ();
+#pragma omp dispatch nowait /* { dg-final { scan-tree-dump-times "#pragma omp task if\\(1\\)" 1 "gimple" } } */
+ f2 ();
+}
+
+
+
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-4.c b/gcc/testsuite/c-c++-common/gomp/dispatch-4.c
new file mode 100644
index 00000000000..aeb5c00507f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-4.c
@@ -0,0 +1,18 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void f2 (int a);
+
+void test (void)
+{
+ int a;
+
+#pragma omp dispatch device(-25373654)
+/* { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device \\(-25373654\\);" 1 "gimple" } } */
+ f2 (a);
+#pragma omp dispatch device(a + a)
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = a \\* 2;.*#pragma omp dispatch device\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\).*#pragma omp task shared\\(D\.\[0-9]+\\).*__builtin_omp_set_default_device \\(D\.\[0-9]+\\);" 1 "gimple" } } */
+ f2 (a);
+}
+
+
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-5.c b/gcc/testsuite/c-c++-common/gomp/dispatch-5.c
new file mode 100644
index 00000000000..78a37bac59b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-5.c
@@ -0,0 +1,26 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void f1 (void* p, int arr[]);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: p, arr)
+void f2 (void* p, int arr[]);
+
+void test (void)
+{
+ void *p;
+ int arr[2];
+
+#pragma omp dispatch
+ f2 (p, arr);
+#pragma omp dispatch is_device_ptr(p)
+/* { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p\\) shared\\(arr\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*int D\.\[0-9]+;\[ \t\n\r]*void \\* D\.\[0-9]+;\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&arr, D\.\[0-9]+\\);\[ \t\n\r]*f1 \\(p, D\.\[0-9]+\\);" 1 "gimple" } } */
+ f2 (p, arr);
+#pragma omp dispatch is_device_ptr(arr)
+/* { dg-final { scan-tree-dump-times "#pragma omp task shared\\(arr\\) shared\\(p\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*int D\.\[0-9]+;\[ \t\n\r]*void \\* D\.\[0-9]+;\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(p, D\.\[0-9]+\\);\[ \t\n\r]*f1 \\(D\.\[0-9]+, &arr\\);" 1 "gimple" } } */
+ f2 (p, arr);
+#pragma omp dispatch is_device_ptr(p, arr)
+/* { dg-final { scan-tree-dump-times "#pragma omp task shared\\(arr\\) shared\\(p\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*f1 \\(p, &arr\\);" 1 "gimple" } } */
+ f2 (p, arr);
+}
+
+
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-6.c b/gcc/testsuite/c-c++-common/gomp/dispatch-6.c
new file mode 100644
index 00000000000..7c495179a13
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-6.c
@@ -0,0 +1,19 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-ompexp" } */
+
+void f2 (void* p);
+
+void test (void)
+{
+ void *p;
+
+#pragma omp dispatch
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_task \\(.*, .*, .*, .*, .*, .*, 0B, .*, .*\\);" 1 "ompexp" } } */
+ f2 (p);
+#pragma omp dispatch depend(inout: p)
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+\\\[2] = &p;" 1 "ompexp" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_task \\(.*, .*, .*, .*, .*, .*, &D\.\[0-9]+, .*, .*\\);" 1 "ompexp" } } */
+ f2 (p);
+}
+
+
diff --git a/gcc/testsuite/c-c++-common/gomp/dispatch-7.c b/gcc/testsuite/c-c++-common/gomp/dispatch-7.c
new file mode 100644
index 00000000000..8cc4526fea0
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/dispatch-7.c
@@ -0,0 +1,28 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-omplower" } */
+
+int f0 (void);
+int f1 (void);
+#pragma omp declare variant (f0) match (construct={dispatch})
+#pragma omp declare variant (f1) match (implementation={vendor(gnu)})
+int f2 (void);
+
+int test (void)
+{
+ int a, n;
+#pragma omp dispatch novariants(n < 1024) nocontext(n > 1024)
+ a = f2 ();
+ return a;
+}
+
+/* { dg-final { scan-tree-dump-times "#pragma omp dispatch nocontext\\(0\\) novariants\\(0\\) shared\\(n\\) shared\\(a\\)" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp task shared\\(n\\) shared\\(a\\)" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 1 "gimple" } } */
+
+/* { dg-final { scan-tree-dump-times ".omp_data_o.1.n = n;" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times ".omp_data_o.1.a = &a;" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "n = .omp_data_i->n;" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = .omp_data_i->a;" 3 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "\\*D\.\[0-9]+ = D\.\[0-9]+;" 3 "omplower" } } */
--
2.45.1
^ permalink raw reply [flat|nested] 10+ messages in thread* [PATCH 6/7] OpenMP: Fortran front-end support for dispatch + adjust_args
2024-05-27 11:54 [PATCH 0/7] OpenMP: dispatch + adjust_args support Paul-Antoine Arras
` (4 preceding siblings ...)
2024-05-27 11:54 ` [PATCH 5/7] OpenMP: common C/C++ testcases " Paul-Antoine Arras
@ 2024-05-27 11:54 ` Paul-Antoine Arras
2024-05-28 7:14 ` Tobias Burnus
2024-05-27 11:54 ` [PATCH 7/7] OpenMP: update documentation for dispatch and adjust_args Paul-Antoine Arras
6 siblings, 1 reply; 10+ messages in thread
From: Paul-Antoine Arras @ 2024-05-27 11:54 UTC (permalink / raw)
To: gcc-patches; +Cc: Paul-Antoine Arras
This patch adds support for the `dispatch` construct and the `adjust_args`
clause to the Fortran front-end.
gcc/fortran/ChangeLog:
* dump-parse-tree.cc (show_omp_clauses): Handle novariants and nocontext
clauses.
(show_omp_node): Handle EXEC_OMP_DISPATCH.
(show_code_node): Likewise.
* frontend-passes.cc (gfc_code_walker): Handle novariants and nocontext.
* gfortran.h (enum gfc_statement): Add ST_OMP_DISPATCH.
(symbol_attribute): Add omp_declare_variant_need_device_ptr.
(gfc_omp_clauses): Add novariants and nocontext.
(gfc_omp_declare_variant): Add need_device_ptr_arg_list.
(enum gfc_exec_op): Add EXEC_OMP_DISPATCH.
* match.h (gfc_match_omp_dispatch): Declare.
* openmp.cc (gfc_free_omp_clauses): Free novariants and nocontext
clauses.
(gfc_free_omp_declare_variant_list): Free need_device_ptr_arg_list
namelist.
(enum omp_mask2): Add OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT.
(gfc_match_omp_clauses): Handle OMP_CLAUSE_NOVARIANTS and
OMP_CLAUSE_NOCONTEXT.
(OMP_DISPATCH_CLAUSES): Define.
(gfc_match_omp_dispatch): New function.
(gfc_match_omp_declare_variant): Parse adjust_args.
(resolve_omp_clauses): Handle adjust_args, novariants and nocontext.
Adjust handling of OMP_LIST_IS_DEVICE_PTR.
(icode_code_error_callback): Handle EXEC_OMP_DISPATCH.
(omp_code_to_statement): Likewise.
(resolve_omp_dispatch): New function.
(gfc_resolve_omp_directive): Handle EXEC_OMP_DISPATCH.
* parse.cc (decode_omp_directive): Match dispatch.
(next_statement): Handle ST_OMP_DISPATCH.
(gfc_ascii_statement): Likewise.
(parse_omp_dispatch): New function.
(parse_executable): Handle ST_OMP_DISPATCH.
* resolve.cc (gfc_resolve_blocks): Handle EXEC_OMP_DISPATCH.
* st.cc (gfc_free_statement): Likewise.
* trans-decl.cc (create_function_arglist): Declare.
(gfc_get_extern_function_decl): Call it.
* trans-openmp.cc (gfc_trans_omp_clauses): Handle novariants and
nocontext.
(gfc_trans_omp_dispatch): New function.
(gfc_trans_omp_directive): Handle EXEC_OMP_DISPATCH.
(gfc_trans_omp_declare_variant): Handle adjust_args.
* trans.cc (trans_code): Handle EXEC_OMP_DISPATCH:.
* types.def (BT_FN_PTR_CONST_PTR_INT): Declare.
gcc/testsuite/ChangeLog:
* gfortran.dg/gomp/declare-variant-2.f90: Update dg-error.
* gfortran.dg/gomp/adjust-args-1.f90: New test.
* gfortran.dg/gomp/adjust-args-2.f90: New test.
* gfortran.dg/gomp/adjust-args-3.f90: New test.
* gfortran.dg/gomp/adjust-args-4.f90: New test.
* gfortran.dg/gomp/adjust-args-5.f90: New test.
* gfortran.dg/gomp/dispatch-1.f90: New test.
* gfortran.dg/gomp/dispatch-2.f90: New test.
* gfortran.dg/gomp/dispatch-3.f90: New test.
* gfortran.dg/gomp/dispatch-4.f90: New test.
* gfortran.dg/gomp/dispatch-5.f90: New test.
* gfortran.dg/gomp/dispatch-6.f90: New test.
* gfortran.dg/gomp/dispatch-7.f90: New test.
* gfortran.dg/gomp/dispatch-8.f90: New test.
---
gcc/fortran/dump-parse-tree.cc | 17 ++
gcc/fortran/frontend-passes.cc | 2 +
gcc/fortran/gfortran.h | 11 +-
gcc/fortran/match.h | 1 +
gcc/fortran/openmp.cc | 193 ++++++++++++++++--
gcc/fortran/parse.cc | 38 ++++
gcc/fortran/resolve.cc | 2 +
gcc/fortran/st.cc | 1 +
gcc/fortran/trans-decl.cc | 9 +-
gcc/fortran/trans-openmp.cc | 161 +++++++++++++++
gcc/fortran/trans.cc | 1 +
gcc/fortran/types.def | 1 +
.../gfortran.dg/gomp/adjust-args-1.f90 | 54 +++++
.../gfortran.dg/gomp/adjust-args-2.f90 | 18 ++
.../gfortran.dg/gomp/adjust-args-3.f90 | 26 +++
.../gfortran.dg/gomp/adjust-args-4.f90 | 58 ++++++
.../gfortran.dg/gomp/adjust-args-5.f90 | 58 ++++++
.../gfortran.dg/gomp/declare-variant-2.f90 | 6 +-
gcc/testsuite/gfortran.dg/gomp/dispatch-1.f90 | 77 +++++++
gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90 | 75 +++++++
gcc/testsuite/gfortran.dg/gomp/dispatch-3.f90 | 39 ++++
gcc/testsuite/gfortran.dg/gomp/dispatch-4.f90 | 19 ++
gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90 | 24 +++
gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90 | 38 ++++
gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90 | 27 +++
gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90 | 39 ++++
26 files changed, 976 insertions(+), 19 deletions(-)
create mode 100644 gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/adjust-args-2.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/adjust-args-3.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/adjust-args-4.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/adjust-args-5.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-1.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-3.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-4.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90
diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
index 87a65036a3d..f64dec63655 100644
--- a/gcc/fortran/dump-parse-tree.cc
+++ b/gcc/fortran/dump-parse-tree.cc
@@ -2119,6 +2119,18 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses)
}
if (omp_clauses->assume)
show_omp_assumes (omp_clauses->assume);
+ if (omp_clauses->novariants)
+ {
+ fputs (" NOVARIANTS(", dumpfile);
+ show_expr (omp_clauses->novariants);
+ fputc (')', dumpfile);
+ }
+ if (omp_clauses->nocontext)
+ {
+ fputs (" NOCONTEXT(", dumpfile);
+ show_expr (omp_clauses->nocontext);
+ fputc (')', dumpfile);
+ }
}
/* Show a single OpenMP or OpenACC directive node and everything underneath it
@@ -2156,6 +2168,9 @@ show_omp_node (int level, gfc_code *c)
case EXEC_OMP_CANCEL: name = "CANCEL"; break;
case EXEC_OMP_CANCELLATION_POINT: name = "CANCELLATION POINT"; break;
case EXEC_OMP_CRITICAL: name = "CRITICAL"; break;
+ case EXEC_OMP_DISPATCH:
+ name = "DISPATCH";
+ break;
case EXEC_OMP_DISTRIBUTE: name = "DISTRIBUTE"; break;
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
name = "DISTRIBUTE PARALLEL DO"; break;
@@ -2257,6 +2272,7 @@ show_omp_node (int level, gfc_code *c)
case EXEC_OMP_ASSUME:
case EXEC_OMP_CANCEL:
case EXEC_OMP_CANCELLATION_POINT:
+ case EXEC_OMP_DISPATCH:
case EXEC_OMP_DISTRIBUTE:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
@@ -3498,6 +3514,7 @@ show_code_node (int level, gfc_code *c)
case EXEC_OMP_BARRIER:
case EXEC_OMP_CRITICAL:
case EXEC_OMP_DEPOBJ:
+ case EXEC_OMP_DISPATCH:
case EXEC_OMP_DISTRIBUTE:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
diff --git a/gcc/fortran/frontend-passes.cc b/gcc/fortran/frontend-passes.cc
index 3c06018fdbb..1a0ef50b91d 100644
--- a/gcc/fortran/frontend-passes.cc
+++ b/gcc/fortran/frontend-passes.cc
@@ -5669,6 +5669,8 @@ gfc_code_walker (gfc_code **c, walk_code_fn_t codefn, walk_expr_fn_t exprfn,
WALK_SUBEXPR (co->ext.omp_clauses->num_tasks);
WALK_SUBEXPR (co->ext.omp_clauses->priority);
WALK_SUBEXPR (co->ext.omp_clauses->detach);
+ WALK_SUBEXPR (co->ext.omp_clauses->novariants);
+ WALK_SUBEXPR (co->ext.omp_clauses->nocontext);
for (idx = 0; idx < ARRAY_SIZE (list_types); idx++)
for (n = co->ext.omp_clauses->lists[list_types[idx]];
n; n = n->next)
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index de1a7cd0935..361b4bece15 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -321,7 +321,7 @@ enum gfc_statement
ST_OMP_ALLOCATE, ST_OMP_ALLOCATE_EXEC,
ST_OMP_ALLOCATORS, ST_OMP_END_ALLOCATORS,
/* Note: gfc_match_omp_nothing returns ST_NONE. */
- ST_OMP_NOTHING, ST_NONE
+ ST_OMP_NOTHING, ST_NONE, ST_OMP_DISPATCH
};
/* Types of interfaces that we can have. Assignment interfaces are
@@ -1004,6 +1004,9 @@ typedef struct
ENUM_BITFIELD (gfc_omp_device_type) omp_device_type:2;
unsigned omp_allocate:1;
+ /* Mentioned in OMP DECLARE VARIANT. */
+ unsigned omp_declare_variant_need_device_ptr : 1;
+
/* Mentioned in OACC DECLARE. */
unsigned oacc_declare_create:1;
unsigned oacc_declare_copyin:1;
@@ -1431,6 +1434,7 @@ enum
OMP_LIST_HAS_DEVICE_ADDR,
OMP_LIST_ENTER,
OMP_LIST_USES_ALLOCATORS,
+ OMP_LIST_ADJUST_ARGS,
OMP_LIST_NUM /* Must be the last. */
};
@@ -1576,6 +1580,8 @@ typedef struct gfc_omp_clauses
struct gfc_expr *depobj;
struct gfc_expr *dist_chunk_size;
struct gfc_expr *message;
+ struct gfc_expr *novariants;
+ struct gfc_expr *nocontext;
struct gfc_omp_assumptions *assume;
const char *critical_name;
enum gfc_omp_default_sharing default_sharing;
@@ -1702,6 +1708,7 @@ typedef struct gfc_omp_declare_variant
struct gfc_symtree *variant_proc_symtree;
gfc_omp_set_selector *set_selectors;
+ gfc_omp_namelist *need_device_ptr_arg_list;
bool checked_p : 1; /* Set if previously checked for errors. */
bool error_p : 1; /* Set if error found in directive. */
@@ -3033,7 +3040,7 @@ enum gfc_exec_op
EXEC_OMP_TARGET_TEAMS_LOOP, EXEC_OMP_MASKED, EXEC_OMP_PARALLEL_MASKED,
EXEC_OMP_PARALLEL_MASKED_TASKLOOP, EXEC_OMP_PARALLEL_MASKED_TASKLOOP_SIMD,
EXEC_OMP_MASKED_TASKLOOP, EXEC_OMP_MASKED_TASKLOOP_SIMD, EXEC_OMP_SCOPE,
- EXEC_OMP_ERROR, EXEC_OMP_ALLOCATE, EXEC_OMP_ALLOCATORS
+ EXEC_OMP_ERROR, EXEC_OMP_ALLOCATE, EXEC_OMP_ALLOCATORS, EXEC_OMP_DISPATCH
};
typedef struct gfc_code
diff --git a/gcc/fortran/match.h b/gcc/fortran/match.h
index b09921357fd..448f631275e 100644
--- a/gcc/fortran/match.h
+++ b/gcc/fortran/match.h
@@ -163,6 +163,7 @@ match gfc_match_omp_declare_simd (void);
match gfc_match_omp_declare_target (void);
match gfc_match_omp_declare_variant (void);
match gfc_match_omp_depobj (void);
+match gfc_match_omp_dispatch (void);
match gfc_match_omp_distribute (void);
match gfc_match_omp_distribute_parallel_do (void);
match gfc_match_omp_distribute_parallel_do_simd (void);
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 5246647e6f8..b29f39febe6 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -71,7 +71,7 @@ static const struct gfc_omp_directive gfc_omp_directives[] = {
{"declare target", GFC_OMP_DIR_DECLARATIVE, ST_OMP_DECLARE_TARGET},
{"declare variant", GFC_OMP_DIR_DECLARATIVE, ST_OMP_DECLARE_VARIANT},
{"depobj", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DEPOBJ},
- /* {"dispatch", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DISPATCH}, */
+ {"dispatch", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DISPATCH},
{"distribute", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DISTRIBUTE},
{"do", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DO},
/* "error" becomes GFC_OMP_DIR_EXECUTABLE with at(execution) */
@@ -180,6 +180,8 @@ gfc_free_omp_clauses (gfc_omp_clauses *c)
gfc_free_expr (c->num_tasks);
gfc_free_expr (c->priority);
gfc_free_expr (c->detach);
+ gfc_free_expr (c->novariants);
+ gfc_free_expr (c->nocontext);
gfc_free_expr (c->async_expr);
gfc_free_expr (c->gang_num_expr);
gfc_free_expr (c->gang_static_expr);
@@ -321,6 +323,8 @@ gfc_free_omp_declare_variant_list (gfc_omp_declare_variant *list)
gfc_omp_declare_variant *current = list;
list = list->next;
gfc_free_omp_set_selector_list (current->set_selectors);
+ gfc_free_omp_namelist (current->need_device_ptr_arg_list, false, false,
+ false);
free (current);
}
}
@@ -1098,6 +1102,8 @@ enum omp_mask2
OMP_CLAUSE_ASSUMPTIONS, /* OpenMP 5.1. */
OMP_CLAUSE_USES_ALLOCATORS, /* OpenMP 5.0 */
OMP_CLAUSE_INDIRECT, /* OpenMP 5.1 */
+ OMP_CLAUSE_NOVARIANTS, /* OpenMP 5.1 */
+ OMP_CLAUSE_NOCONTEXT, /* OpenMP 5.1 */
/* This must come last. */
OMP_MASK2_LAST
};
@@ -3215,6 +3221,25 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
c->assume->no_parallelism = needs_space = true;
continue;
}
+
+ if ((mask & OMP_CLAUSE_NOVARIANTS)
+ && (m = gfc_match_dupl_check (!c->novariants, "novariants", true,
+ &c->novariants))
+ != MATCH_NO)
+ {
+ if (m == MATCH_ERROR)
+ goto error;
+ continue;
+ }
+ if ((mask & OMP_CLAUSE_NOCONTEXT)
+ && (m = gfc_match_dupl_check (!c->nocontext, "nocontext", true,
+ &c->nocontext))
+ != MATCH_NO)
+ {
+ if (m == MATCH_ERROR)
+ goto error;
+ continue;
+ }
if ((mask & OMP_CLAUSE_NOGROUP)
&& (m = gfc_match_dupl_check (!c->nogroup, "nogroup"))
!= MATCH_NO)
@@ -4530,6 +4555,9 @@ cleanup:
omp_mask (OMP_CLAUSE_NOWAIT)
#define OMP_ALLOCATORS_CLAUSES \
omp_mask (OMP_CLAUSE_ALLOCATE)
+#define OMP_DISPATCH_CLAUSES \
+ (omp_mask (OMP_CLAUSE_DEVICE) | OMP_CLAUSE_DEPEND | OMP_CLAUSE_NOVARIANTS \
+ | OMP_CLAUSE_NOCONTEXT | OMP_CLAUSE_IS_DEVICE_PTR | OMP_CLAUSE_NOWAIT)
static match
@@ -4843,6 +4871,12 @@ error:
return MATCH_ERROR;
}
+match
+gfc_match_omp_dispatch (void)
+{
+ return match_omp (EXEC_OMP_DISPATCH, OMP_DISPATCH_CLAUSES);
+}
+
match
gfc_match_omp_distribute (void)
{
@@ -6069,6 +6103,7 @@ gfc_match_omp_declare_variant (void)
odv = gfc_get_omp_declare_variant ();
odv->where = gfc_current_locus;
odv->variant_proc_symtree = variant_proc_st;
+ odv->need_device_ptr_arg_list = NULL;
odv->base_proc_symtree = base_proc_st;
odv->next = NULL;
odv->error_p = false;
@@ -6085,13 +6120,29 @@ gfc_match_omp_declare_variant (void)
return MATCH_ERROR;
}
+ bool has_match = false, has_adjust_args = false;
+ locus adjust_args_loc;
+
for (;;)
{
- if (gfc_match (" match") != MATCH_YES)
+ enum clause
+ {
+ match,
+ adjust_args
+ } ccode;
+
+ if (gfc_match (" match") == MATCH_YES)
+ ccode = match;
+ else if (gfc_match (" adjust_args") == MATCH_YES)
+ {
+ ccode = adjust_args;
+ adjust_args_loc = gfc_current_locus;
+ }
+ else
{
if (first_p)
{
- gfc_error ("expected %<match%> at %C");
+ gfc_error ("expected %<match%> or %<adjust_args%> at %C");
return MATCH_ERROR;
}
else
@@ -6104,18 +6155,86 @@ gfc_match_omp_declare_variant (void)
return MATCH_ERROR;
}
- if (gfc_match_omp_context_selector_specification (odv) != MATCH_YES)
- return MATCH_ERROR;
-
- if (gfc_match (" )") != MATCH_YES)
+ if (ccode == match)
{
- gfc_error ("expected %<)%> at %C");
- return MATCH_ERROR;
+ has_match = true;
+ if (gfc_match_omp_context_selector_specification (odv)
+ != MATCH_YES)
+ return MATCH_ERROR;
+ if (gfc_match (" )") != MATCH_YES)
+ {
+ gfc_error ("expected %<)%> at %C");
+ return MATCH_ERROR;
+ }
+ }
+ else if (ccode == adjust_args)
+ {
+ has_adjust_args = true;
+ bool need_device_ptr_p;
+ if (gfc_match (" nothing") == MATCH_YES)
+ need_device_ptr_p = false;
+ else if (gfc_match (" need_device_ptr") == MATCH_YES)
+ need_device_ptr_p = true;
+ else
+ {
+ gfc_error ("expected %<nothing%> or %<need_device_ptr%> at %C");
+ return MATCH_ERROR;
+ }
+ if (need_device_ptr_p)
+ {
+ if (gfc_match_omp_variable_list (" :",
+ &odv->need_device_ptr_arg_list,
+ false)
+ != MATCH_YES)
+ {
+ gfc_error ("expected argument list at %C");
+ return MATCH_ERROR;
+ }
+ for (gfc_omp_namelist *n = odv->need_device_ptr_arg_list;
+ n != NULL; n = n->next)
+ {
+ if (!n->sym->attr.dummy)
+ {
+ gfc_error ("list item %qs at %L is not a dummy argument",
+ n->sym->name, &n->where);
+ return MATCH_ERROR;
+ }
+ if (n->sym->ts.type != BT_DERIVED
+ || !n->sym->ts.u.derived->ts.is_iso_c)
+ {
+ gfc_error ("argument list item %qs in "
+ "%<need_device_ptr%> at %L must be of "
+ "TYPE(C_PTR)",
+ n->sym->name, &n->where);
+ return MATCH_ERROR;
+ }
+ }
+ }
+ else
+ {
+ gfc_omp_namelist *nothing_arg_list = NULL;
+ if (gfc_match_omp_variable_list (" :", ¬hing_arg_list, false)
+ != MATCH_YES)
+ {
+ gfc_error ("expected argument list at %C");
+ return MATCH_ERROR;
+ }
+ gfc_free_omp_namelist (nothing_arg_list, false, false, false);
+ }
}
first_p = false;
}
+ if (has_adjust_args && !has_match)
+ {
+ gfc_error ("an %<adjust_args%> clause at %C can only be specified if the "
+ "%<dispatch%> selector of the construct selector set appears "
+ "in the %<match%> clause",
+ &adjust_args_loc);
+ return MATCH_ERROR;
+ }
+
return MATCH_YES;
}
@@ -7544,7 +7663,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
"DEVICE_RESIDENT", "LINK", "USE_DEVICE",
"CACHE", "IS_DEVICE_PTR", "USE_DEVICE_PTR", "USE_DEVICE_ADDR",
"NONTEMPORAL", "ALLOCATE", "HAS_DEVICE_ADDR", "ENTER",
- "USES_ALLOCATORS" };
+ "USES_ALLOCATORS", "ADJUST_ARGS" };
STATIC_ASSERT (ARRAY_SIZE (clause_names) == OMP_LIST_NUM);
if (omp_clauses == NULL)
@@ -7726,6 +7845,26 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
gfc_error ("FINAL clause at %L requires a scalar LOGICAL expression",
&expr->where);
}
+ if (omp_clauses->novariants)
+ {
+ gfc_expr *expr = omp_clauses->novariants;
+ if (!gfc_resolve_expr (expr) || expr->ts.type != BT_LOGICAL
+ || expr->rank != 0)
+ gfc_error (
+ "NOVARIANTS clause at %L requires a scalar LOGICAL expression",
+ &expr->where);
+ if_without_mod = true;
+ }
+ if (omp_clauses->nocontext)
+ {
+ gfc_expr *expr = omp_clauses->nocontext;
+ if (!gfc_resolve_expr (expr) || expr->ts.type != BT_LOGICAL
+ || expr->rank != 0)
+ gfc_error (
+ "NOCONTEXT clause at %L requires a scalar LOGICAL expression",
+ &expr->where);
+ if_without_mod = true;
+ }
if (omp_clauses->num_threads)
resolve_positive_int_expr (omp_clauses->num_threads, "NUM_THREADS");
if (omp_clauses->chunk_size)
@@ -8675,9 +8814,9 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
last = NULL;
for (n = omp_clauses->lists[list]; n != NULL; )
{
- if (n->sym->ts.type == BT_DERIVED
- && n->sym->ts.u.derived->ts.is_iso_c
- && code->op != EXEC_OMP_TARGET)
+ if ((n->sym->ts.type != BT_DERIVED
+ || !n->sym->ts.u.derived->ts.is_iso_c)
+ && code->op == EXEC_OMP_DISPATCH)
/* Non-TARGET (i.e. DISPATCH) requires a C_PTR. */
gfc_error ("List item %qs in %s clause at %L must be of "
"TYPE(C_PTR)", n->sym->name, name, &n->where);
@@ -10290,6 +10429,7 @@ icode_code_error_callback (gfc_code **codep,
case EXEC_OMP_MASKED_TASKLOOP_SIMD:
case EXEC_OMP_SCOPE:
case EXEC_OMP_ERROR:
+ case EXEC_OMP_DISPATCH:
gfc_error ("%s cannot contain OpenMP directive in intervening code "
"at %L",
state->name, &code->loc);
@@ -11168,6 +11308,8 @@ omp_code_to_statement (gfc_code *code)
return ST_OMP_PARALLEL_LOOP;
case EXEC_OMP_DEPOBJ:
return ST_OMP_DEPOBJ;
+ case EXEC_OMP_DISPATCH:
+ return ST_OMP_DISPATCH;
default:
gcc_unreachable ();
}
@@ -11583,6 +11725,26 @@ resolve_omp_target (gfc_code *code)
#undef GFC_IS_TEAMS_CONSTRUCT
}
+static void
+resolve_omp_dispatch (gfc_code *code)
+{
+ gfc_code *next = code->block->next;
+ if (next == NULL)
+ return;
+ gfc_exec_op op = next->op;
+ if (op != EXEC_CALL
+ && (op != EXEC_ASSIGN || next->expr2->expr_type != EXPR_FUNCTION))
+ gfc_error (
+ "%<OMP DISPATCH%> directive at %L must be followed by a procedure "
+ "call with optional assignment",
+ &code->loc);
+
+ if ((op == EXEC_CALL && next->resolved_sym->attr.proc_pointer)
+ || (op == EXEC_ASSIGN && gfc_expr_attr (next->expr2).proc_pointer))
+ gfc_error ("%<OMP DISPATCH%> directive at %L cannot be followed by a "
+ "procedure pointer",
+ &code->loc);
+}
/* Resolve OpenMP directive clauses and check various requirements
of each directive. */
@@ -11696,6 +11858,11 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns)
code->ext.omp_clauses->if_present = false;
resolve_omp_clauses (code, code->ext.omp_clauses, ns);
break;
+ case EXEC_OMP_DISPATCH:
+ if (code->ext.omp_clauses)
+ resolve_omp_clauses (code, code->ext.omp_clauses, ns);
+ resolve_omp_dispatch (code);
+ break;
default:
break;
}
diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc
index 79c810c86ba..74fc249269d 100644
--- a/gcc/fortran/parse.cc
+++ b/gcc/fortran/parse.cc
@@ -1050,6 +1050,7 @@ decode_omp_directive (void)
break;
case 'd':
matcho ("depobj", gfc_match_omp_depobj, ST_OMP_DEPOBJ);
+ matcho ("dispatch", gfc_match_omp_dispatch, ST_OMP_DISPATCH);
matchs ("distribute parallel do simd",
gfc_match_omp_distribute_parallel_do_simd,
ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD);
@@ -1916,6 +1917,7 @@ next_statement (void)
case ST_OMP_LOOP: case ST_OMP_PARALLEL_LOOP: case ST_OMP_TEAMS_LOOP: \
case ST_OMP_TARGET_PARALLEL_LOOP: case ST_OMP_TARGET_TEAMS_LOOP: \
case ST_OMP_ALLOCATE_EXEC: case ST_OMP_ALLOCATORS: case ST_OMP_ASSUME: \
+ case ST_OMP_DISPATCH: \
case ST_CRITICAL: \
case ST_OACC_PARALLEL_LOOP: case ST_OACC_PARALLEL: case ST_OACC_KERNELS: \
case ST_OACC_DATA: case ST_OACC_HOST_DATA: case ST_OACC_LOOP: \
@@ -2597,6 +2599,9 @@ gfc_ascii_statement (gfc_statement st, bool strip_sentinel)
case ST_OMP_DEPOBJ:
p = "!$OMP DEPOBJ";
break;
+ case ST_OMP_DISPATCH:
+ p = "!$OMP DISPATCH";
+ break;
case ST_OMP_DISTRIBUTE:
p = "!$OMP DISTRIBUTE";
break;
@@ -6183,6 +6188,35 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only)
}
+static gfc_statement
+parse_omp_dispatch (void)
+{
+ gfc_statement st;
+ gfc_code *cp, *np;
+ gfc_state_data s;
+
+ accept_statement (ST_OMP_DISPATCH);
+
+ cp = gfc_state_stack->tail;
+ push_state (&s, COMP_OMP_STRUCTURED_BLOCK, NULL);
+ np = new_level (cp);
+ np->op = cp->op;
+ np->block = NULL;
+
+ st = next_statement ();
+ if (st == ST_CALL || st == ST_ASSIGNMENT)
+ accept_statement (st);
+ else
+ {
+ gfc_error ("%<OMP DISPATCH%> directive must be followed by a procedure "
+ "call with optional assignment at %C");
+ reject_statement ();
+ }
+ pop_state ();
+ st = next_statement ();
+ return st;
+}
+
/* Accept a series of executable statements. We return the first
statement that doesn't fit to the caller. Any block statements are
passed on to the correct handler, which usually passes the buck
@@ -6383,6 +6417,10 @@ parse_executable (gfc_statement st)
st = parse_omp_oacc_atomic (true);
continue;
+ case ST_OMP_DISPATCH:
+ st = parse_omp_dispatch ();
+ continue;
+
default:
return st;
}
diff --git a/gcc/fortran/resolve.cc b/gcc/fortran/resolve.cc
index d7a0856fcca..755d1302ce9 100644
--- a/gcc/fortran/resolve.cc
+++ b/gcc/fortran/resolve.cc
@@ -11378,6 +11378,7 @@ gfc_resolve_blocks (gfc_code *b, gfc_namespace *ns)
case EXEC_OMP_ALLOCATORS:
case EXEC_OMP_ASSUME:
case EXEC_OMP_CRITICAL:
+ case EXEC_OMP_DISPATCH:
case EXEC_OMP_DISTRIBUTE:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
@@ -13054,6 +13055,7 @@ start:
case EXEC_OMP_CRITICAL:
case EXEC_OMP_FLUSH:
case EXEC_OMP_DEPOBJ:
+ case EXEC_OMP_DISPATCH:
case EXEC_OMP_DISTRIBUTE:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
diff --git a/gcc/fortran/st.cc b/gcc/fortran/st.cc
index 6a605ad91d4..90ee1352ba4 100644
--- a/gcc/fortran/st.cc
+++ b/gcc/fortran/st.cc
@@ -222,6 +222,7 @@ gfc_free_statement (gfc_code *p)
case EXEC_OMP_CANCELLATION_POINT:
case EXEC_OMP_CRITICAL:
case EXEC_OMP_DEPOBJ:
+ case EXEC_OMP_DISPATCH:
case EXEC_OMP_DISTRIBUTE:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc
index dca7779528b..4390769146a 100644
--- a/gcc/fortran/trans-decl.cc
+++ b/gcc/fortran/trans-decl.cc
@@ -2124,6 +2124,8 @@ get_proc_pointer_decl (gfc_symbol *sym)
return decl;
}
+static void
+create_function_arglist (gfc_symbol *sym);
/* Get a basic decl for an external function. */
@@ -2377,7 +2379,12 @@ module_sym:
if (sym->formal_ns->omp_declare_simd)
gfc_trans_omp_declare_simd (sym->formal_ns);
if (flag_openmp)
- gfc_trans_omp_declare_variant (sym->formal_ns);
+ {
+ // We need DECL_ARGUMENTS to put attributes on, in case some arguments
+ // need adjustment
+ create_function_arglist (sym->formal_ns->proc_name);
+ gfc_trans_omp_declare_variant (sym->formal_ns);
+ }
}
return fndecl;
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index f867e2240bf..5e4450184d1 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -4233,6 +4233,36 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
}
+ if (clauses->novariants)
+ {
+ tree novariants_var;
+
+ gfc_init_se (&se, NULL);
+ gfc_conv_expr (&se, clauses->novariants);
+ gfc_add_block_to_block (block, &se.pre);
+ novariants_var = gfc_evaluate_now (se.expr, block);
+ gfc_add_block_to_block (block, &se.post);
+
+ c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_NOVARIANTS);
+ OMP_CLAUSE_NOVARIANTS_EXPR (c) = novariants_var;
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
+
+ if (clauses->nocontext)
+ {
+ tree nocontext_var;
+
+ gfc_init_se (&se, NULL);
+ gfc_conv_expr (&se, clauses->nocontext);
+ gfc_add_block_to_block (block, &se.pre);
+ nocontext_var = gfc_evaluate_now (se.expr, block);
+ gfc_add_block_to_block (block, &se.post);
+
+ c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_NOCONTEXT);
+ OMP_CLAUSE_NOCONTEXT_EXPR (c) = nocontext_var;
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
+
if (clauses->num_threads)
{
tree num_threads;
@@ -6311,6 +6341,30 @@ gfc_trans_omp_depobj (gfc_code *code)
return gfc_finish_block (&block);
}
+static tree
+gfc_trans_omp_dispatch (gfc_code *code)
+{
+ stmtblock_t block;
+ gfc_code *next = code->block->next;
+ // assume ill-formed "function dispatch structured
+ // block" have already been rejected by resolve_omp_dispatch
+ gcc_assert (next->op == EXEC_CALL || next->op == EXEC_ASSIGN);
+
+ tree body = gfc_trans_code (next);
+ gfc_start_block (&block);
+ tree omp_clauses
+ = gfc_trans_omp_clauses (&block, code->ext.omp_clauses, code->loc);
+
+ tree stmt = make_node (OMP_DISPATCH);
+ SET_EXPR_LOCATION (stmt, gfc_get_location (&code->loc));
+ TREE_TYPE (stmt) = void_type_node;
+ OMP_DISPATCH_BODY (stmt) = body;
+ OMP_DISPATCH_CLAUSES (stmt) = omp_clauses;
+
+ gfc_add_expr_to_block (&block, stmt);
+ return gfc_finish_block (&block);
+}
+
static tree
gfc_trans_omp_error (gfc_code *code)
{
@@ -8221,6 +8275,8 @@ gfc_trans_omp_directive (gfc_code *code)
case EXEC_OMP_TASKLOOP:
return gfc_trans_omp_do (code, code->op, NULL, code->ext.omp_clauses,
NULL);
+ case EXEC_OMP_DISPATCH:
+ return gfc_trans_omp_dispatch (code);
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
case EXEC_OMP_DISTRIBUTE_SIMD:
@@ -8337,6 +8393,7 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns)
tree base_fn_decl = ns->proc_name->backend_decl;
gfc_namespace *search_ns = ns;
gfc_omp_declare_variant *next;
+ vec<tree> adjust_args_list = vNULL;
for (gfc_omp_declare_variant *odv = search_ns->omp_declare_variant;
search_ns; odv = next)
@@ -8532,6 +8589,19 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns)
variant_proc_sym = NULL;
}
}
+ if (odv->need_device_ptr_arg_list != NULL
+ && omp_get_context_selector (set_selectors, OMP_TRAIT_SET_CONSTRUCT,
+ OMP_TRAIT_CONSTRUCT_DISPATCH)
+ == NULL_TREE)
+ {
+ gfc_error ("an %<adjust_args%> clause can only be "
+ "specified if the "
+ "%<dispatch%> selector of the construct "
+ "selector set appears "
+ "in the %<match%> clause at %L",
+ &odv->where);
+ variant_proc_sym = NULL;
+ }
if (variant_proc_sym != NULL)
{
gfc_set_sym_referenced (variant_proc_sym);
@@ -8548,6 +8618,97 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns)
DECL_ATTRIBUTES (base_fn_decl)
= tree_cons (id, build_tree_list (variant, set_selectors),
DECL_ATTRIBUTES (base_fn_decl));
+
+ // Handle adjust_args
+ for (gfc_omp_namelist *arg_list
+ = odv->need_device_ptr_arg_list;
+ arg_list != NULL; arg_list = arg_list->next)
+ {
+ if (arg_list->sym->backend_decl == NULL_TREE)
+ {
+ gfc_error (
+ "%s at %L is not a base function argument",
+ arg_list->sym->name, &arg_list->where);
+ continue;
+ }
+
+ tree base_fn_arg_decl = arg_list->sym->backend_decl;
+ if (base_fn_arg_decl != error_mark_node)
+ {
+ // Is t specified more than once?
+ if (adjust_args_list.contains (base_fn_arg_decl))
+ {
+ gfc_error (
+ "%qD at %L is specified more than once",
+ base_fn_arg_decl, &arg_list->where);
+ continue;
+ }
+ adjust_args_list.safe_push (base_fn_arg_decl);
+
+ // Handle variant argument
+ tree variant
+ = gfc_get_symbol_decl (variant_proc_sym);
+ tree variant_parm = DECL_ARGUMENTS (variant);
+ int idx;
+ tree arg;
+ for (arg = DECL_ARGUMENTS (base_fn_decl), idx = 0;
+ arg != NULL; arg = TREE_CHAIN (arg), idx++)
+ if (arg == base_fn_arg_decl)
+ break;
+ gcc_assert (arg != NULL_TREE);
+ if (variant_parm == NULL_TREE)
+ {
+ gfc_formal_arglist *arg
+ = variant_proc_sym->formal;
+ for (int i = 0; i < idx; i++)
+ {
+ arg = arg->next;
+ gcc_assert (arg != NULL);
+ }
+
+ // Check we got the right parameter name
+ if (strcmp (arg_list->sym->name, arg->sym->name)
+ != 0)
+ {
+ gfc_error ("%s at %L is not a variant "
+ "function argument",
+ arg_list->sym->name,
+ &arg_list->where);
+ continue;
+ }
+ arg->sym->attr
+ .omp_declare_variant_need_device_ptr
+ = 1;
+ }
+ else
+ {
+ for (int i = 0; i < idx; i++)
+ {
+ variant_parm = TREE_CHAIN (variant_parm);
+ gcc_assert (variant_parm != NULL_TREE);
+ }
+ // Check we got the right parameter name
+ if (strcmp (arg_list->sym->name,
+ IDENTIFIER_POINTER (
+ DECL_NAME (variant_parm)))
+ != 0)
+ {
+ gfc_error ("%s at %L is not a variant "
+ "function argument",
+ arg_list->sym->name,
+ &arg_list->where);
+ continue;
+ }
+
+ tree attr = tree_cons (
+ get_identifier (
+ "omp declare variant adjust_args "
+ "need_device_ptr"),
+ NULL_TREE, DECL_ATTRIBUTES (variant_parm));
+ DECL_ATTRIBUTES (variant_parm) = attr;
+ }
+ }
+ }
}
}
}
diff --git a/gcc/fortran/trans.cc b/gcc/fortran/trans.cc
index badad6ae892..2795cdf7464 100644
--- a/gcc/fortran/trans.cc
+++ b/gcc/fortran/trans.cc
@@ -2596,6 +2596,7 @@ trans_code (gfc_code * code, tree cond)
case EXEC_OMP_CANCELLATION_POINT:
case EXEC_OMP_CRITICAL:
case EXEC_OMP_DEPOBJ:
+ case EXEC_OMP_DISPATCH:
case EXEC_OMP_DISTRIBUTE:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index 390cc9542f7..5047c8f816a 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -120,6 +120,7 @@ DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_INT_BOOL, BT_BOOL, BT_INT, BT_BOOL)
DEF_FUNCTION_TYPE_2 (BT_FN_VOID_PTR_PTRMODE,
BT_VOID, BT_PTR, BT_PTRMODE)
DEF_FUNCTION_TYPE_2 (BT_FN_VOID_CONST_PTR_SIZE, BT_VOID, BT_CONST_PTR, BT_SIZE)
+DEF_FUNCTION_TYPE_2 (BT_FN_PTR_CONST_PTR_INT, BT_PTR, BT_CONST_PTR, BT_INT)
DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR_PTR, BT_FN_VOID_PTR_PTR)
diff --git a/gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90
new file mode 100644
index 00000000000..982c2c2cb9d
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90
@@ -0,0 +1,54 @@
+! Test parsing of OMP clause adjust_args
+! { dg-do compile }
+
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+ integer :: b
+interface
+ integer function f0 (a)
+ import c_ptr
+ type(c_ptr), intent(inout) :: a
+ end function
+ integer function g (a)
+ import c_ptr
+ type(c_ptr), intent(inout) :: a
+ end function
+ integer function f1 (i)
+ integer, intent(in) :: i
+ end function
+
+ integer function f3 (a)
+ import c_ptr
+ type(c_ptr), intent(inout) :: a
+ !$omp declare variant (f1) match (construct={dispatch}) adjust_args (other: a) ! { dg-error "expected 'nothing' or 'need_device_ptr' at .1." }
+end function
+integer function f4 (a)
+import c_ptr
+type(c_ptr), intent(inout) :: a
+ !$omp declare variant (f0) adjust_args (nothing: a) ! { dg-error "an 'adjust_args' clause at .1. can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" }
+end function
+integer function f5 (i)
+integer, intent(inout) :: i
+ !$omp declare variant (f1) match (construct={dispatch}) adjust_args () ! { dg-error "expected 'nothing' or 'need_device_ptr' at .1." }
+end function
+integer function f6 (i)
+integer, intent(inout) :: i
+ !$omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing) ! { dg-error "expected argument list at .1." }
+end function
+integer function f7 (i)
+integer, intent(inout) :: i
+ !$omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing:) ! { dg-error "expected argument list at .1." }
+ end function
+ integer function f9 (i)
+ integer, intent(inout) :: i
+ !$omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: i) ! { dg-error "argument list item 'i' in 'need_device_ptr' at .1. must be of TYPE.C_PTR." }
+end function
+ integer function f12 (a)
+ import c_ptr
+ type(c_ptr), intent(inout) :: a
+ !$omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: b) ! { dg-error "list item 'b' at .1. is not a dummy argument" }
+end function
+
+ end interface
+end module
diff --git a/gcc/testsuite/gfortran.dg/gomp/adjust-args-2.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-2.f90
new file mode 100644
index 00000000000..c65a4839ca5
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-2.f90
@@ -0,0 +1,18 @@
+! Test resolution of OMP clause adjust_args
+! { dg-do compile }
+
+module main
+ implicit none
+interface
+subroutine f1 (i)
+ integer, intent(inout) :: i
+end subroutine
+end interface
+contains
+
+ subroutine f3 (i)
+ integer, intent(inout) :: i
+ !$omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: z) ! { dg-error "Symbol 'z' at .1. has no IMPLICIT type" }
+ end subroutine
+
+end module
diff --git a/gcc/testsuite/gfortran.dg/gomp/adjust-args-3.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-3.f90
new file mode 100644
index 00000000000..b731cb340c1
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-3.f90
@@ -0,0 +1,26 @@
+! Test translation of OMP clause adjust_args
+! { dg-do compile }
+
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+ !type(c_ptr) :: a
+
+contains
+ subroutine base2 (a)
+ type(c_ptr), intent(inout) :: a
+ !$omp declare variant (variant2) match (construct={parallel}) adjust_args (need_device_ptr: a) ! { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause at .1." }
+ end subroutine
+ subroutine base3 (a)
+ type(c_ptr), intent(inout) :: a
+ !$omp declare variant (variant2) match (construct={dispatch}) adjust_args (need_device_ptr: a) adjust_args (need_device_ptr: a) ! { dg-error "'a' at .1. is specified more than once" }
+ end subroutine
+
+ subroutine variant2 (a)
+ type(c_ptr), intent(inout) :: a
+ end subroutine
+ subroutine variant3 (i)
+ integer :: i
+ end subroutine
+
+end module
diff --git a/gcc/testsuite/gfortran.dg/gomp/adjust-args-4.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-4.f90
new file mode 100644
index 00000000000..75e884044b2
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-4.f90
@@ -0,0 +1,58 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+
+ type :: struct
+ integer :: a
+ real :: b
+ end type
+
+ interface
+ integer function f(a, b, c)
+ import c_ptr
+ integer, intent(in) :: a
+ type(c_ptr), intent(inout) :: b
+ type(c_ptr), intent(out) :: c(:)
+ end function
+ integer function f0(a, b, c)
+ import c_ptr
+ integer, intent(in) :: a
+ type(c_ptr), intent(inout) :: b
+ type(c_ptr), intent(out) :: c(:)
+ !$omp declare variant (f) match (construct={dispatch}) &
+ !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b, c)
+ end function
+ integer function f1(a, b, c)
+ import c_ptr
+ integer, intent(in) :: a
+ type(c_ptr), intent(inout) :: b
+ type(c_ptr), intent(out) :: c(:)
+ !$omp declare variant (f) match (construct={dispatch}) &
+ !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b) adjust_args (need_device_ptr: c)
+ end function
+ end interface
+
+contains
+subroutine test
+ integer :: a
+ type(c_ptr) :: b
+ type(c_ptr) :: c(2)
+ type(struct) :: s
+
+ s%a = f0 (a, b, c)
+ !$omp dispatch
+ s%a = f0 (a, b, c)
+
+ s%b = f1 (a, b, c)
+ !$omp dispatch
+ s%b = f1 (a, b, c)
+
+end subroutine
+end module
+
+! { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&parm\.\[0-9]+, D\.\[0-9]+\\);" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&b, D\.\[0-9]+\\);" 2 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/adjust-args-5.f90 b/gcc/testsuite/gfortran.dg/gomp/adjust-args-5.f90
new file mode 100644
index 00000000000..75e884044b2
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/adjust-args-5.f90
@@ -0,0 +1,58 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+
+ type :: struct
+ integer :: a
+ real :: b
+ end type
+
+ interface
+ integer function f(a, b, c)
+ import c_ptr
+ integer, intent(in) :: a
+ type(c_ptr), intent(inout) :: b
+ type(c_ptr), intent(out) :: c(:)
+ end function
+ integer function f0(a, b, c)
+ import c_ptr
+ integer, intent(in) :: a
+ type(c_ptr), intent(inout) :: b
+ type(c_ptr), intent(out) :: c(:)
+ !$omp declare variant (f) match (construct={dispatch}) &
+ !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b, c)
+ end function
+ integer function f1(a, b, c)
+ import c_ptr
+ integer, intent(in) :: a
+ type(c_ptr), intent(inout) :: b
+ type(c_ptr), intent(out) :: c(:)
+ !$omp declare variant (f) match (construct={dispatch}) &
+ !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b) adjust_args (need_device_ptr: c)
+ end function
+ end interface
+
+contains
+subroutine test
+ integer :: a
+ type(c_ptr) :: b
+ type(c_ptr) :: c(2)
+ type(struct) :: s
+
+ s%a = f0 (a, b, c)
+ !$omp dispatch
+ s%a = f0 (a, b, c)
+
+ s%b = f1 (a, b, c)
+ !$omp dispatch
+ s%b = f1 (a, b, c)
+
+end subroutine
+end module
+
+! { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&parm\.\[0-9]+, D\.\[0-9]+\\);" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&b, D\.\[0-9]+\\);" 2 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90
index 7fc5071feff..62d2cb96fac 100644
--- a/gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90
@@ -18,10 +18,10 @@ contains
!$omp declare variant match(user={condition(.false.)}) ! { dg-error "expected '\\(' at .1." }
end subroutine
subroutine f6 ()
- !$omp declare variant (f1) ! { dg-error "expected 'match' at .1." }
+ !$omp declare variant (f1) ! { dg-error "expected 'match' or 'adjust_args' at .1." }
end subroutine
subroutine f7 ()
- !$omp declare variant (f1) simd ! { dg-error "expected 'match' at .1." }
+ !$omp declare variant (f1) simd ! { dg-error "expected 'match' or 'adjust_args' at .1." }
end subroutine
subroutine f8 ()
!$omp declare variant (f1) match ! { dg-error "expected '\\(' at .1." }
@@ -183,7 +183,7 @@ contains
!$omp declare variant (f1) match(construct={requires}) ! { dg-warning "unknown selector 'requires' for context selector set 'construct' at .1." }
end subroutine
subroutine f75 ()
- !$omp declare variant (f1),match(construct={parallel}) ! { dg-error "expected 'match' at .1." }
+ !$omp declare variant (f1),match(construct={parallel}) ! { dg-error "expected 'match' or 'adjust_args' at .1." }
end subroutine
subroutine f76 ()
!$omp declare variant (f1) match(implementation={atomic_default_mem_order("relaxed")}) ! { dg-error "expected identifier at .1." }
diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-1.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-1.f90
new file mode 100644
index 00000000000..12c30904131
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-1.f90
@@ -0,0 +1,77 @@
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+ contains
+
+ subroutine f1 ()
+ integer :: a, b, arr(10)
+ real :: x
+ complex :: c
+ character :: ch
+ logical :: bool
+ type :: struct
+ integer :: a
+ real :: b
+ end type
+ type(struct) :: s
+ type(c_ptr) :: p
+
+ interface
+ subroutine f0 (a, c, bool, s)
+ import :: struct
+ integer, intent(in) :: a
+ complex, intent(out) :: c
+ logical, intent(inout) :: bool
+ type(struct) :: s
+ end subroutine
+ integer function f2 (arr, x, ch, b)
+ integer, intent(inout) :: arr(:)
+ real, intent(in) :: x
+ character, intent(out) :: ch
+ real :: b
+ end function
+ subroutine f3 (p)
+ import :: c_ptr
+ type(c_ptr) :: p
+ end subroutine
+ integer function f4 ()
+ end function
+ end interface
+
+ !$omp dispatch
+ b = f2(arr, x, ch, s%b)
+ !$omp dispatch
+ c = f2(arr(:5), x * 2.4, ch, s%b)
+ !$omp dispatch
+ arr(1) = f2(arr, x, ch, s%b)
+ !$omp dispatch
+ s%a = f2(arr, x, ch, s%b)
+ !$omp dispatch
+ x = f2(arr, x, ch, s%b)
+ !$omp dispatch
+ call f0(a, c, bool, s)
+ !$omp dispatch
+ call f0(f4(), c, bool, s)
+
+ !$omp dispatch nocontext(.TRUE.)
+ call f0(a, c, bool, s)
+ !$omp dispatch nocontext(arr(2) < 10)
+ call f0(a, c, bool, s)
+ !$omp dispatch novariants(.FALSE.)
+ call f0(a, c, bool, s)
+ !$omp dispatch novariants(bool)
+ call f0(a, c, bool, s)
+ !$omp dispatch nowait
+ call f0(a, c, bool, s)
+ !$omp dispatch device(arr(9))
+ call f0(a, c, bool, s)
+ !$omp dispatch device(a + a)
+ call f0(a, c, bool, s)
+ !$omp dispatch device(-25373654)
+ call f0(a, c, bool, s)
+ !$omp dispatch is_device_ptr(p)
+ call f3(p)
+ !$omp dispatch depend(in: a, c, bool) depend(inout: s, arr(:3))
+ call f0(a, c, bool, s)
+ end subroutine
+end module
diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90
new file mode 100644
index 00000000000..f52df4446c4
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90
@@ -0,0 +1,75 @@
+module main
+ implicit none
+ contains
+
+ subroutine f1 ()
+ integer :: a, b, arr(10)
+ real :: x
+ complex :: c
+ character :: ch
+ logical :: bool
+ type :: struct
+ integer :: a
+ real :: b
+ end type
+ type(struct) :: s
+
+ interface
+ subroutine f0 (a, c, bool, s)
+ import :: struct
+ integer, intent(in) :: a
+ complex, intent(out) :: c
+ logical, intent(inout) :: bool
+ type(struct) :: s
+ end subroutine
+ integer function f2 (arr, x, ch, b)
+ integer, intent(inout) :: arr(:)
+ real, intent(in) :: x
+ character, intent(out) :: ch
+ real :: b
+ end function
+ end interface
+ procedure(f0), pointer:: fp => NULL()
+
+ !$omp dispatch !{ dg-error "'OMP DISPATCH' directive at .1. must be followed by a procedure call with optional assignment" }
+50 b = f2(arr, x, ch, s%b) + a
+ !$omp dispatch !{ dg-error "'OMP DISPATCH' directive at .1. must be followed by a procedure call with optional assignment" }
+ a = b
+ !$omp dispatch !{ dg-error "'OMP DISPATCH' directive at .1. must be followed by a procedure call with optional assignment" }
+ b = Not (2)
+ !$omp dispatch
+ !$omp threadprivate(a) !{ dg-error "'OMP DISPATCH' directive must be followed by a procedure call with optional assignment at .1." }
+ a = f2(arr, x, ch, s%b)
+ !$omp dispatch
+ print *, 'This is not allowed here.' !{ dg-error "'OMP DISPATCH' directive must be followed by a procedure call with optional assignment at .1." }
+ !$omp dispatch
+ goto 50 !{ dg-error "'OMP DISPATCH' directive must be followed by a procedure call with optional assignment at .1." }
+ !$omp dispatch !{ dg-error "'OMP DISPATCH' directive at .1. cannot be followed by a procedure pointer" }
+ call fp(a, c, bool, s)
+
+ !$omp dispatch nocontext(s) !{ dg-error "NOCONTEXT clause at .1. requires a scalar LOGICAL expression" }
+ call f0(a, c, bool, s)
+ !$omp dispatch nocontext(a, b) !{ dg-error "Invalid expression after 'nocontext.' at .1." }
+ call f0(a, c, bool, s)
+ !$omp dispatch nocontext(a) nocontext(b) !{ dg-error "Duplicated 'nocontext' clause at .1." }
+ call f0(a, c, bool, s)
+ !$omp dispatch novariants(s) !{ dg-error "NOVARIANTS clause at .1. requires a scalar LOGICAL expression" }
+ call f0(a, c, bool, s)
+ !$omp dispatch novariants(a, b) !{ dg-error "Invalid expression after 'novariants.' at .1." }
+ call f0(a, c, bool, s)
+ !$omp dispatch novariants(a) novariants(b) !{ dg-error "Duplicated 'novariants' clause at .1." }
+ call f0(a, c, bool, s)
+ !$omp dispatch nowait nowait !{ dg-error "Duplicated 'nowait' clause at .1." }
+ call f0(a, c, bool, s)
+ !$omp dispatch device(x) !{ dg-error "DEVICE clause at .1. requires a scalar INTEGER expression" }
+ call f0(a, c, bool, s)
+ !$omp dispatch device(arr) !{ dg-error "DEVICE clause at .1. requires a scalar INTEGER expression" }
+ call f0(a, c, bool, s)
+ !$omp dispatch is_device_ptr(x) !{ dg-error "List item 'x' in IS_DEVICE_PTR clause at .1. must be of TYPE.C_PTR." }
+ call f0(a, c, bool, s)
+ !$omp dispatch is_device_ptr(arr) !{ dg-error "List item 'arr' in IS_DEVICE_PTR clause at .1. must be of TYPE.C_PTR." }
+ call f0(a, c, bool, s)
+ !$omp dispatch depend(inout: f0) !{ dg-error "Object 'f0' is not a variable at .1." }
+ call f0(a, c, bool, s)
+ end subroutine
+end module
diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-3.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-3.f90
new file mode 100644
index 00000000000..84590fd883a
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-3.f90
@@ -0,0 +1,39 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module main
+ implicit none
+ interface
+ integer function f0 ()
+ end function
+
+ integer function f1 ()
+ end function
+
+ integer function f2 ()
+ !$omp declare variant (f0) match (construct={dispatch})
+ !$omp declare variant (f1) match (implementation={vendor(gnu)})
+ end function
+ end interface
+ contains
+
+ integer function test ()
+ integer :: a
+
+ !$omp dispatch
+ a = f2 ()
+ !$omp dispatch novariants(.TRUE.)
+ a = f2 ()
+ !$omp dispatch novariants(.FALSE.)
+ a = f2 ()
+ !$omp dispatch nocontext(.TRUE.)
+ a = f2 ()
+ !$omp dispatch nocontext(.FALSE.)
+ a = f2 ()
+ end function
+end module
+
+
+! { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 3 "gimple" } }
+! { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-4.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-4.f90
new file mode 100644
index 00000000000..149d0613b97
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-4.f90
@@ -0,0 +1,19 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module main
+ implicit none
+ interface
+ subroutine f2 ()
+ end subroutine
+ end interface
+ contains
+
+ subroutine test ()
+ !$omp dispatch ! { dg-final { scan-tree-dump-times "#pragma omp task if\\(0\\)" 1 "gimple" } }
+ call f2 ()
+ !$omp dispatch nowait ! { dg-final { scan-tree-dump-times "#pragma omp task if\\(1\\)" 1 "gimple" } }
+ call f2 ()
+ end subroutine
+end module
+
diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90
new file mode 100644
index 00000000000..e45397f3f96
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90
@@ -0,0 +1,24 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module main
+ implicit none
+ interface
+ subroutine f2 (a)
+ integer, intent(in) :: a
+ end subroutine
+ end interface
+ contains
+
+ subroutine test ()
+ integer :: a
+
+ !$omp dispatch device(-25373654)
+ ! { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device \\(-25373654\\);" 1 "gimple" } }
+ call f2 (a)
+ !$omp dispatch device(a + a)
+ ! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = a.0_1 \\* 2;.*#pragma omp dispatch device\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\).*#pragma omp task shared\\(D\.\[0-9]+\\).*__builtin_omp_set_default_device \\(D\.\[0-9]+\\);" 1 "gimple" } }
+ call f2 (a)
+ end subroutine
+end module
+
diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90
new file mode 100644
index 00000000000..9f4fa2970ca
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90
@@ -0,0 +1,38 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+ interface
+ subroutine f1 (p, p2)
+ import :: c_ptr
+ type(c_ptr), intent(out) :: p
+ type(c_ptr), intent(in) :: p2
+ end subroutine
+ subroutine f2 (p, p2)
+ import :: c_ptr
+ type(c_ptr), intent(out) :: p
+ type(c_ptr), intent(in) :: p2
+ !$omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: p, p2)
+ end subroutine
+ end interface
+ contains
+
+ subroutine test ()
+ type(c_ptr) :: p, p2
+
+ !$omp dispatch
+ call f2 (p, p2)
+ !$omp dispatch is_device_ptr(p)
+ ! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p\\) shared\\(p2\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*integer\\(kind=4\\) D\.\[0-9]+;\[ \t\n\r]*void \\* D\.\[0-9]+;\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&p2, D\.\[0-9]+\\);\[ \t\n\r]*f1 \\(&p, D\.\[0-9]+\\);" 1 "gimple" } }
+ call f2 (p, p2)
+ !$omp dispatch is_device_ptr(p2)
+ ! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p2\\) shared\\(p\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*integer\\(kind=4\\) D\.\[0-9]+;\[ \t\n\r]*void \\* D\.\[0-9]+;\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&p, D\.\[0-9]+\\);\[ \t\n\r]*f1 \\(D\.\[0-9]+, &p2\\);" 1 "gimple" } }
+ call f2 (p, p2)
+ !$omp dispatch is_device_ptr(p, p2)
+ ! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p\\) shared\\(p2\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*f1 \\(&p, &p2\\);" 1 "gimple" } }
+ call f2 (p, p2)
+ end subroutine
+end module
+
diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90
new file mode 100644
index 00000000000..32b6347be67
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90
@@ -0,0 +1,27 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-ompexp" }
+
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+ interface
+ subroutine f2 (p)
+ import :: c_ptr
+ type(c_ptr), intent(out) :: p
+ end subroutine
+ end interface
+ contains
+
+ subroutine test ()
+ type(c_ptr) :: p
+
+ !$omp dispatch
+ ! { dg-final { scan-tree-dump-times "__builtin_GOMP_task \\(.*, .*, .*, .*, .*, .*, 0B, .*, .*\\);" 1 "ompexp" } }
+ call f2 (p)
+ !$omp dispatch depend(inout: p)
+ ! { dg-final { scan-tree-dump-times "D\.\[0-9]+\\\[2] = &p;" 1 "ompexp" } }
+ ! { dg-final { scan-tree-dump-times "__builtin_GOMP_task \\(.*, .*, .*, .*, .*, .*, &D\.\[0-9]+, .*, .*\\);" 1 "ompexp" } }
+ call f2 (p)
+ end subroutine
+end module
+
diff --git a/gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90 b/gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90
new file mode 100644
index 00000000000..6771336aa33
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90
@@ -0,0 +1,39 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple -fdump-tree-omplower" }
+
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+ interface
+ integer function f0 ()
+ end function
+ integer function f1 ()
+ end function
+ integer function f2 ()
+ !$omp declare variant (f0) match (construct={dispatch})
+ !$omp declare variant (f1) match (implementation={vendor(gnu)})
+ end function
+ end interface
+ contains
+
+ subroutine test ()
+ integer :: a, n
+
+ !$omp dispatch novariants(n < 1024) nocontext(n > 1024)
+ a = f2 ()
+ end subroutine
+end module
+
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = n <= 1023;" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = n > 1024;" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp dispatch novariants\\(0\\) nocontext\\(0\\) shared\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 1 "gimple" } }
+
+! { dg-final { scan-tree-dump-times ".omp_data_o.1.D\.\[0-9]+ = D\.\[0-9]+;" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times ".omp_data_o.1.a = &a;" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = .omp_data_i->D\.\[0-9]+;" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = .omp_data_i->a;" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "\\*D\.\[0-9]+ = D\.\[0-9]+;" 3 "omplower" } }
--
2.45.1
^ permalink raw reply [flat|nested] 10+ messages in thread* Re: [PATCH 6/7] OpenMP: Fortran front-end support for dispatch + adjust_args
2024-05-27 11:54 ` [PATCH 6/7] OpenMP: Fortran front-end support " Paul-Antoine Arras
@ 2024-05-28 7:14 ` Tobias Burnus
2024-05-31 10:07 ` Paul-Antoine Arras
0 siblings, 1 reply; 10+ messages in thread
From: Tobias Burnus @ 2024-05-28 7:14 UTC (permalink / raw)
To: Paul-Antoine Arras, gcc-patches
Hi PA, hi all,
two remarks while quickly browsing the code:
Paul-Antoine Arras:
> + if (n->sym->ts.type != BT_DERIVED
> + || !n->sym->ts.u.derived->ts.is_iso_c)
> + {
> + gfc_error ("argument list item %qs in "
> + "%<need_device_ptr%> at %L must be of "
> + "TYPE(C_PTR)",
> + n->sym->name, &n->where);
I think you need to rule out 'c_funptr' as well, e.g. via:
|| (n->sym->ts.u.derived->intmod_sym_id
!= ISOCBINDING_PTR)))
I do note that in openmp.cc, we have one check which checks explicitly
for c_ptr and one existing one which only checks for (c_ptr or
c_funptr); can you fix that one as well?
* * *
But I mainly miss an update to 'module.cc' for the 'declare variant'
change; the 'adjust_args' (for 'need_device_ptr', only) list items have
to be saved in the .mod file - otherwise the following will not work:
<some-file>-aux.f90
! { dg-do compile { target skip-all-targets } }
module my_mod
...
!$omp declare variant ... adjust_args(need_device_ptr: ...)
...
end module
<some-file>.f90
{ dg-do ...
! { dg-additional-sources <some-file>-aux.f90 }
...
call <base-function>
...
!$omp displatch
call <base-function>
end
For C++ modules, it should be fine as those for those, the tree is dumped.
Tobias
^ permalink raw reply [flat|nested] 10+ messages in thread* Re: [PATCH 6/7] OpenMP: Fortran front-end support for dispatch + adjust_args
2024-05-28 7:14 ` Tobias Burnus
@ 2024-05-31 10:07 ` Paul-Antoine Arras
0 siblings, 0 replies; 10+ messages in thread
From: Paul-Antoine Arras @ 2024-05-31 10:07 UTC (permalink / raw)
To: Tobias Burnus, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 1860 bytes --]
Hi Tobias,
Thanks for your comments. Here is an updated patch.
On 28/05/2024 09:14, Tobias Burnus wrote:
> Paul-Antoine Arras:
>> + if (n->sym->ts.type != BT_DERIVED
>> + || !n->sym->ts.u.derived->ts.is_iso_c)
>> + {
>> + gfc_error ("argument list item %qs in "
>> + "%<need_device_ptr%> at %L must be of "
>> + "TYPE(C_PTR)",
>> + n->sym->name, &n->where);
>
> I think you need to rule out 'c_funptr' as well, e.g. via:
>
> || (n->sym->ts.u.derived->intmod_sym_id
> != ISOCBINDING_PTR)))
>
> I do note that in openmp.cc, we have one check which checks explicitly
> for c_ptr and one existing one which only checks for (c_ptr or
> c_funptr); can you fix that one as well?
This is now handled in the new patch.
> But I mainly miss an update to 'module.cc' for the 'declare variant'
> change; the 'adjust_args' (for 'need_device_ptr', only) list items have
> to be saved in the .mod file - otherwise the following will not work:
>
> <some-file>-aux.f90
> ! { dg-do compile { target skip-all-targets } }
> module my_mod
> ...
> !$omp declare variant ... adjust_args(need_device_ptr: ...)
> ...
> end module
>
> <some-file>.f90
> { dg-do ...
> ! { dg-additional-sources <some-file>-aux.f90 }
> ...
> call <base-function>
> ...
> !$omp displatch
> call <base-function>
> end
I added a new testcase along those lines. However, I had to xfail it due
to completely missing support for declare variant (even without
adjust_args) in module.cc. For reference, Tobias created this PR:
https://gcc.gnu.org/PR115271.
--
PA
[-- Attachment #2: fortran.patch --]
[-- Type: text/x-patch, Size: 57911 bytes --]
commit ab1b93e3e6e7cb9b5a7419b7106ea11110324699
Author: Paul-Antoine Arras <parras@baylibre.com>
Date: Fri May 24 19:13:50 2024 +0200
OpenMP: Fortran front-end support for dispatch + adjust_args
This patch adds support for the `dispatch` construct and the `adjust_args`
clause to the Fortran front-end.
Handling of `adjust_args` across translation units is missing due to PR115271.
gcc/fortran/ChangeLog:
* dump-parse-tree.cc (show_omp_clauses): Handle novariants and nocontext
clauses.
(show_omp_node): Handle EXEC_OMP_DISPATCH.
(show_code_node): Likewise.
* frontend-passes.cc (gfc_code_walker): Handle novariants and nocontext.
* gfortran.h (enum gfc_statement): Add ST_OMP_DISPATCH.
(symbol_attribute): Add omp_declare_variant_need_device_ptr.
(gfc_omp_clauses): Add novariants and nocontext.
(gfc_omp_declare_variant): Add need_device_ptr_arg_list.
(enum gfc_exec_op): Add EXEC_OMP_DISPATCH.
* match.h (gfc_match_omp_dispatch): Declare.
* openmp.cc (gfc_free_omp_clauses): Free novariants and nocontext
clauses.
(gfc_free_omp_declare_variant_list): Free need_device_ptr_arg_list
namelist.
(enum omp_mask2): Add OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT.
(gfc_match_omp_clauses): Handle OMP_CLAUSE_NOVARIANTS and
OMP_CLAUSE_NOCONTEXT.
(OMP_DISPATCH_CLAUSES): Define.
(gfc_match_omp_dispatch): New function.
(gfc_match_omp_declare_variant): Parse adjust_args.
(resolve_omp_clauses): Handle adjust_args, novariants and nocontext.
Adjust handling of OMP_LIST_IS_DEVICE_PTR.
(icode_code_error_callback): Handle EXEC_OMP_DISPATCH.
(omp_code_to_statement): Likewise.
(resolve_omp_dispatch): New function.
(gfc_resolve_omp_directive): Handle EXEC_OMP_DISPATCH.
* parse.cc (decode_omp_directive): Match dispatch.
(next_statement): Handle ST_OMP_DISPATCH.
(gfc_ascii_statement): Likewise.
(parse_omp_dispatch): New function.
(parse_executable): Handle ST_OMP_DISPATCH.
* resolve.cc (gfc_resolve_blocks): Handle EXEC_OMP_DISPATCH.
* st.cc (gfc_free_statement): Likewise.
* trans-decl.cc (create_function_arglist): Declare.
(gfc_get_extern_function_decl): Call it.
* trans-openmp.cc (gfc_trans_omp_clauses): Handle novariants and
nocontext.
(gfc_trans_omp_dispatch): New function.
(gfc_trans_omp_directive): Handle EXEC_OMP_DISPATCH.
(gfc_trans_omp_declare_variant): Handle adjust_args.
* trans.cc (trans_code): Handle EXEC_OMP_DISPATCH:.
* types.def (BT_FN_PTR_CONST_PTR_INT): Declare.
gcc/testsuite/ChangeLog:
* gfortran.dg/gomp/declare-variant-2.f90: Update dg-error.
* gfortran.dg/gomp/declare-variant-21.f90: New test (xfail).
* gfortran.dg/gomp/declare-variant-21-aux.f90: New test.
* gfortran.dg/gomp/adjust-args-1.f90: New test.
* gfortran.dg/gomp/adjust-args-2.f90: New test.
* gfortran.dg/gomp/adjust-args-3.f90: New test.
* gfortran.dg/gomp/adjust-args-4.f90: New test.
* gfortran.dg/gomp/adjust-args-5.f90: New test.
* gfortran.dg/gomp/dispatch-1.f90: New test.
* gfortran.dg/gomp/dispatch-2.f90: New test.
* gfortran.dg/gomp/dispatch-3.f90: New test.
* gfortran.dg/gomp/dispatch-4.f90: New test.
* gfortran.dg/gomp/dispatch-5.f90: New test.
* gfortran.dg/gomp/dispatch-6.f90: New test.
* gfortran.dg/gomp/dispatch-7.f90: New test.
* gfortran.dg/gomp/dispatch-8.f90: New test.
diff --git gcc/fortran/dump-parse-tree.cc gcc/fortran/dump-parse-tree.cc
index 87a65036a3d..f64dec63655 100644
--- gcc/fortran/dump-parse-tree.cc
+++ gcc/fortran/dump-parse-tree.cc
@@ -2119,6 +2119,18 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses)
}
if (omp_clauses->assume)
show_omp_assumes (omp_clauses->assume);
+ if (omp_clauses->novariants)
+ {
+ fputs (" NOVARIANTS(", dumpfile);
+ show_expr (omp_clauses->novariants);
+ fputc (')', dumpfile);
+ }
+ if (omp_clauses->nocontext)
+ {
+ fputs (" NOCONTEXT(", dumpfile);
+ show_expr (omp_clauses->nocontext);
+ fputc (')', dumpfile);
+ }
}
/* Show a single OpenMP or OpenACC directive node and everything underneath it
@@ -2156,6 +2168,9 @@ show_omp_node (int level, gfc_code *c)
case EXEC_OMP_CANCEL: name = "CANCEL"; break;
case EXEC_OMP_CANCELLATION_POINT: name = "CANCELLATION POINT"; break;
case EXEC_OMP_CRITICAL: name = "CRITICAL"; break;
+ case EXEC_OMP_DISPATCH:
+ name = "DISPATCH";
+ break;
case EXEC_OMP_DISTRIBUTE: name = "DISTRIBUTE"; break;
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
name = "DISTRIBUTE PARALLEL DO"; break;
@@ -2257,6 +2272,7 @@ show_omp_node (int level, gfc_code *c)
case EXEC_OMP_ASSUME:
case EXEC_OMP_CANCEL:
case EXEC_OMP_CANCELLATION_POINT:
+ case EXEC_OMP_DISPATCH:
case EXEC_OMP_DISTRIBUTE:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
@@ -3498,6 +3514,7 @@ show_code_node (int level, gfc_code *c)
case EXEC_OMP_BARRIER:
case EXEC_OMP_CRITICAL:
case EXEC_OMP_DEPOBJ:
+ case EXEC_OMP_DISPATCH:
case EXEC_OMP_DISTRIBUTE:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
diff --git gcc/fortran/frontend-passes.cc gcc/fortran/frontend-passes.cc
index 3c06018fdbb..1a0ef50b91d 100644
--- gcc/fortran/frontend-passes.cc
+++ gcc/fortran/frontend-passes.cc
@@ -5669,6 +5669,8 @@ gfc_code_walker (gfc_code **c, walk_code_fn_t codefn, walk_expr_fn_t exprfn,
WALK_SUBEXPR (co->ext.omp_clauses->num_tasks);
WALK_SUBEXPR (co->ext.omp_clauses->priority);
WALK_SUBEXPR (co->ext.omp_clauses->detach);
+ WALK_SUBEXPR (co->ext.omp_clauses->novariants);
+ WALK_SUBEXPR (co->ext.omp_clauses->nocontext);
for (idx = 0; idx < ARRAY_SIZE (list_types); idx++)
for (n = co->ext.omp_clauses->lists[list_types[idx]];
n; n = n->next)
diff --git gcc/fortran/gfortran.h gcc/fortran/gfortran.h
index de1a7cd0935..361b4bece15 100644
--- gcc/fortran/gfortran.h
+++ gcc/fortran/gfortran.h
@@ -321,7 +321,7 @@ enum gfc_statement
ST_OMP_ALLOCATE, ST_OMP_ALLOCATE_EXEC,
ST_OMP_ALLOCATORS, ST_OMP_END_ALLOCATORS,
/* Note: gfc_match_omp_nothing returns ST_NONE. */
- ST_OMP_NOTHING, ST_NONE
+ ST_OMP_NOTHING, ST_NONE, ST_OMP_DISPATCH
};
/* Types of interfaces that we can have. Assignment interfaces are
@@ -1004,6 +1004,9 @@ typedef struct
ENUM_BITFIELD (gfc_omp_device_type) omp_device_type:2;
unsigned omp_allocate:1;
+ /* Mentioned in OMP DECLARE VARIANT. */
+ unsigned omp_declare_variant_need_device_ptr : 1;
+
/* Mentioned in OACC DECLARE. */
unsigned oacc_declare_create:1;
unsigned oacc_declare_copyin:1;
@@ -1431,6 +1434,7 @@ enum
OMP_LIST_HAS_DEVICE_ADDR,
OMP_LIST_ENTER,
OMP_LIST_USES_ALLOCATORS,
+ OMP_LIST_ADJUST_ARGS,
OMP_LIST_NUM /* Must be the last. */
};
@@ -1576,6 +1580,8 @@ typedef struct gfc_omp_clauses
struct gfc_expr *depobj;
struct gfc_expr *dist_chunk_size;
struct gfc_expr *message;
+ struct gfc_expr *novariants;
+ struct gfc_expr *nocontext;
struct gfc_omp_assumptions *assume;
const char *critical_name;
enum gfc_omp_default_sharing default_sharing;
@@ -1702,6 +1708,7 @@ typedef struct gfc_omp_declare_variant
struct gfc_symtree *variant_proc_symtree;
gfc_omp_set_selector *set_selectors;
+ gfc_omp_namelist *need_device_ptr_arg_list;
bool checked_p : 1; /* Set if previously checked for errors. */
bool error_p : 1; /* Set if error found in directive. */
@@ -3033,7 +3040,7 @@ enum gfc_exec_op
EXEC_OMP_TARGET_TEAMS_LOOP, EXEC_OMP_MASKED, EXEC_OMP_PARALLEL_MASKED,
EXEC_OMP_PARALLEL_MASKED_TASKLOOP, EXEC_OMP_PARALLEL_MASKED_TASKLOOP_SIMD,
EXEC_OMP_MASKED_TASKLOOP, EXEC_OMP_MASKED_TASKLOOP_SIMD, EXEC_OMP_SCOPE,
- EXEC_OMP_ERROR, EXEC_OMP_ALLOCATE, EXEC_OMP_ALLOCATORS
+ EXEC_OMP_ERROR, EXEC_OMP_ALLOCATE, EXEC_OMP_ALLOCATORS, EXEC_OMP_DISPATCH
};
typedef struct gfc_code
diff --git gcc/fortran/match.h gcc/fortran/match.h
index b09921357fd..448f631275e 100644
--- gcc/fortran/match.h
+++ gcc/fortran/match.h
@@ -163,6 +163,7 @@ match gfc_match_omp_declare_simd (void);
match gfc_match_omp_declare_target (void);
match gfc_match_omp_declare_variant (void);
match gfc_match_omp_depobj (void);
+match gfc_match_omp_dispatch (void);
match gfc_match_omp_distribute (void);
match gfc_match_omp_distribute_parallel_do (void);
match gfc_match_omp_distribute_parallel_do_simd (void);
diff --git gcc/fortran/openmp.cc gcc/fortran/openmp.cc
index 5246647e6f8..b70a3ed138b 100644
--- gcc/fortran/openmp.cc
+++ gcc/fortran/openmp.cc
@@ -71,7 +71,7 @@ static const struct gfc_omp_directive gfc_omp_directives[] = {
{"declare target", GFC_OMP_DIR_DECLARATIVE, ST_OMP_DECLARE_TARGET},
{"declare variant", GFC_OMP_DIR_DECLARATIVE, ST_OMP_DECLARE_VARIANT},
{"depobj", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DEPOBJ},
- /* {"dispatch", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DISPATCH}, */
+ {"dispatch", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DISPATCH},
{"distribute", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DISTRIBUTE},
{"do", GFC_OMP_DIR_EXECUTABLE, ST_OMP_DO},
/* "error" becomes GFC_OMP_DIR_EXECUTABLE with at(execution) */
@@ -180,6 +180,8 @@ gfc_free_omp_clauses (gfc_omp_clauses *c)
gfc_free_expr (c->num_tasks);
gfc_free_expr (c->priority);
gfc_free_expr (c->detach);
+ gfc_free_expr (c->novariants);
+ gfc_free_expr (c->nocontext);
gfc_free_expr (c->async_expr);
gfc_free_expr (c->gang_num_expr);
gfc_free_expr (c->gang_static_expr);
@@ -321,6 +323,8 @@ gfc_free_omp_declare_variant_list (gfc_omp_declare_variant *list)
gfc_omp_declare_variant *current = list;
list = list->next;
gfc_free_omp_set_selector_list (current->set_selectors);
+ gfc_free_omp_namelist (current->need_device_ptr_arg_list, false, false,
+ false);
free (current);
}
}
@@ -1098,6 +1102,8 @@ enum omp_mask2
OMP_CLAUSE_ASSUMPTIONS, /* OpenMP 5.1. */
OMP_CLAUSE_USES_ALLOCATORS, /* OpenMP 5.0 */
OMP_CLAUSE_INDIRECT, /* OpenMP 5.1 */
+ OMP_CLAUSE_NOVARIANTS, /* OpenMP 5.1 */
+ OMP_CLAUSE_NOCONTEXT, /* OpenMP 5.1 */
/* This must come last. */
OMP_MASK2_LAST
};
@@ -3215,6 +3221,25 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
c->assume->no_parallelism = needs_space = true;
continue;
}
+
+ if ((mask & OMP_CLAUSE_NOVARIANTS)
+ && (m = gfc_match_dupl_check (!c->novariants, "novariants", true,
+ &c->novariants))
+ != MATCH_NO)
+ {
+ if (m == MATCH_ERROR)
+ goto error;
+ continue;
+ }
+ if ((mask & OMP_CLAUSE_NOCONTEXT)
+ && (m = gfc_match_dupl_check (!c->nocontext, "nocontext", true,
+ &c->nocontext))
+ != MATCH_NO)
+ {
+ if (m == MATCH_ERROR)
+ goto error;
+ continue;
+ }
if ((mask & OMP_CLAUSE_NOGROUP)
&& (m = gfc_match_dupl_check (!c->nogroup, "nogroup"))
!= MATCH_NO)
@@ -4530,6 +4555,9 @@ cleanup:
omp_mask (OMP_CLAUSE_NOWAIT)
#define OMP_ALLOCATORS_CLAUSES \
omp_mask (OMP_CLAUSE_ALLOCATE)
+#define OMP_DISPATCH_CLAUSES \
+ (omp_mask (OMP_CLAUSE_DEVICE) | OMP_CLAUSE_DEPEND | OMP_CLAUSE_NOVARIANTS \
+ | OMP_CLAUSE_NOCONTEXT | OMP_CLAUSE_IS_DEVICE_PTR | OMP_CLAUSE_NOWAIT)
static match
@@ -4843,6 +4871,12 @@ error:
return MATCH_ERROR;
}
+match
+gfc_match_omp_dispatch (void)
+{
+ return match_omp (EXEC_OMP_DISPATCH, OMP_DISPATCH_CLAUSES);
+}
+
match
gfc_match_omp_distribute (void)
{
@@ -6069,6 +6103,7 @@ gfc_match_omp_declare_variant (void)
odv = gfc_get_omp_declare_variant ();
odv->where = gfc_current_locus;
odv->variant_proc_symtree = variant_proc_st;
+ odv->need_device_ptr_arg_list = NULL;
odv->base_proc_symtree = base_proc_st;
odv->next = NULL;
odv->error_p = false;
@@ -6085,13 +6120,29 @@ gfc_match_omp_declare_variant (void)
return MATCH_ERROR;
}
+ bool has_match = false, has_adjust_args = false;
+ locus adjust_args_loc;
+
for (;;)
{
- if (gfc_match (" match") != MATCH_YES)
+ enum clause
+ {
+ match,
+ adjust_args
+ } ccode;
+
+ if (gfc_match (" match") == MATCH_YES)
+ ccode = match;
+ else if (gfc_match (" adjust_args") == MATCH_YES)
+ {
+ ccode = adjust_args;
+ adjust_args_loc = gfc_current_locus;
+ }
+ else
{
if (first_p)
{
- gfc_error ("expected %<match%> at %C");
+ gfc_error ("expected %<match%> or %<adjust_args%> at %C");
return MATCH_ERROR;
}
else
@@ -6104,18 +6155,88 @@ gfc_match_omp_declare_variant (void)
return MATCH_ERROR;
}
- if (gfc_match_omp_context_selector_specification (odv) != MATCH_YES)
- return MATCH_ERROR;
-
- if (gfc_match (" )") != MATCH_YES)
+ if (ccode == match)
{
- gfc_error ("expected %<)%> at %C");
- return MATCH_ERROR;
+ has_match = true;
+ if (gfc_match_omp_context_selector_specification (odv)
+ != MATCH_YES)
+ return MATCH_ERROR;
+ if (gfc_match (" )") != MATCH_YES)
+ {
+ gfc_error ("expected %<)%> at %C");
+ return MATCH_ERROR;
+ }
+ }
+ else if (ccode == adjust_args)
+ {
+ has_adjust_args = true;
+ bool need_device_ptr_p;
+ if (gfc_match (" nothing") == MATCH_YES)
+ need_device_ptr_p = false;
+ else if (gfc_match (" need_device_ptr") == MATCH_YES)
+ need_device_ptr_p = true;
+ else
+ {
+ gfc_error ("expected %<nothing%> or %<need_device_ptr%> at %C");
+ return MATCH_ERROR;
+ }
+ if (need_device_ptr_p)
+ {
+ if (gfc_match_omp_variable_list (" :",
+ &odv->need_device_ptr_arg_list,
+ false)
+ != MATCH_YES)
+ {
+ gfc_error ("expected argument list at %C");
+ return MATCH_ERROR;
+ }
+ for (gfc_omp_namelist *n = odv->need_device_ptr_arg_list;
+ n != NULL; n = n->next)
+ {
+ if (!n->sym->attr.dummy)
+ {
+ gfc_error ("list item %qs at %L is not a dummy argument",
+ n->sym->name, &n->where);
+ return MATCH_ERROR;
+ }
+ if (n->sym->ts.type != BT_DERIVED
+ || !n->sym->ts.u.derived->ts.is_iso_c
+ || (n->sym->ts.u.derived->intmod_sym_id
+ != ISOCBINDING_PTR))
+ {
+ gfc_error ("argument list item %qs in "
+ "%<need_device_ptr%> at %L must be of "
+ "TYPE(C_PTR)",
+ n->sym->name, &n->where);
+ return MATCH_ERROR;
+ }
+ }
+ }
+ else
+ {
+ gfc_omp_namelist *nothing_arg_list = NULL;
+ if (gfc_match_omp_variable_list (" :", ¬hing_arg_list, false)
+ != MATCH_YES)
+ {
+ gfc_error ("expected argument list at %C");
+ return MATCH_ERROR;
+ }
+ gfc_free_omp_namelist (nothing_arg_list, false, false, false);
+ }
}
first_p = false;
}
+ if (has_adjust_args && !has_match)
+ {
+ gfc_error ("an %<adjust_args%> clause at %C can only be specified if the "
+ "%<dispatch%> selector of the construct selector set appears "
+ "in the %<match%> clause",
+ &adjust_args_loc);
+ return MATCH_ERROR;
+ }
+
return MATCH_YES;
}
@@ -7544,7 +7665,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
"DEVICE_RESIDENT", "LINK", "USE_DEVICE",
"CACHE", "IS_DEVICE_PTR", "USE_DEVICE_PTR", "USE_DEVICE_ADDR",
"NONTEMPORAL", "ALLOCATE", "HAS_DEVICE_ADDR", "ENTER",
- "USES_ALLOCATORS" };
+ "USES_ALLOCATORS", "ADJUST_ARGS" };
STATIC_ASSERT (ARRAY_SIZE (clause_names) == OMP_LIST_NUM);
if (omp_clauses == NULL)
@@ -7726,6 +7847,26 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
gfc_error ("FINAL clause at %L requires a scalar LOGICAL expression",
&expr->where);
}
+ if (omp_clauses->novariants)
+ {
+ gfc_expr *expr = omp_clauses->novariants;
+ if (!gfc_resolve_expr (expr) || expr->ts.type != BT_LOGICAL
+ || expr->rank != 0)
+ gfc_error (
+ "NOVARIANTS clause at %L requires a scalar LOGICAL expression",
+ &expr->where);
+ if_without_mod = true;
+ }
+ if (omp_clauses->nocontext)
+ {
+ gfc_expr *expr = omp_clauses->nocontext;
+ if (!gfc_resolve_expr (expr) || expr->ts.type != BT_LOGICAL
+ || expr->rank != 0)
+ gfc_error (
+ "NOCONTEXT clause at %L requires a scalar LOGICAL expression",
+ &expr->where);
+ if_without_mod = true;
+ }
if (omp_clauses->num_threads)
resolve_positive_int_expr (omp_clauses->num_threads, "NUM_THREADS");
if (omp_clauses->chunk_size)
@@ -8675,14 +8816,18 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
last = NULL;
for (n = omp_clauses->lists[list]; n != NULL; )
{
- if (n->sym->ts.type == BT_DERIVED
- && n->sym->ts.u.derived->ts.is_iso_c
- && code->op != EXEC_OMP_TARGET)
+ if ((n->sym->ts.type != BT_DERIVED
+ || !n->sym->ts.u.derived->ts.is_iso_c
+ || (n->sym->ts.u.derived->intmod_sym_id
+ != ISOCBINDING_PTR))
+ && code->op == EXEC_OMP_DISPATCH)
/* Non-TARGET (i.e. DISPATCH) requires a C_PTR. */
gfc_error ("List item %qs in %s clause at %L must be of "
"TYPE(C_PTR)", n->sym->name, name, &n->where);
else if (n->sym->ts.type != BT_DERIVED
- || !n->sym->ts.u.derived->ts.is_iso_c)
+ || !n->sym->ts.u.derived->ts.is_iso_c
+ || (n->sym->ts.u.derived->intmod_sym_id
+ != ISOCBINDING_PTR))
{
/* For TARGET, non-C_PTR are deprecated and handled as
has_device_addr. */
@@ -10290,6 +10435,7 @@ icode_code_error_callback (gfc_code **codep,
case EXEC_OMP_MASKED_TASKLOOP_SIMD:
case EXEC_OMP_SCOPE:
case EXEC_OMP_ERROR:
+ case EXEC_OMP_DISPATCH:
gfc_error ("%s cannot contain OpenMP directive in intervening code "
"at %L",
state->name, &code->loc);
@@ -11168,6 +11314,8 @@ omp_code_to_statement (gfc_code *code)
return ST_OMP_PARALLEL_LOOP;
case EXEC_OMP_DEPOBJ:
return ST_OMP_DEPOBJ;
+ case EXEC_OMP_DISPATCH:
+ return ST_OMP_DISPATCH;
default:
gcc_unreachable ();
}
@@ -11583,6 +11731,26 @@ resolve_omp_target (gfc_code *code)
#undef GFC_IS_TEAMS_CONSTRUCT
}
+static void
+resolve_omp_dispatch (gfc_code *code)
+{
+ gfc_code *next = code->block->next;
+ if (next == NULL)
+ return;
+ gfc_exec_op op = next->op;
+ if (op != EXEC_CALL
+ && (op != EXEC_ASSIGN || next->expr2->expr_type != EXPR_FUNCTION))
+ gfc_error (
+ "%<OMP DISPATCH%> directive at %L must be followed by a procedure "
+ "call with optional assignment",
+ &code->loc);
+
+ if ((op == EXEC_CALL && next->resolved_sym->attr.proc_pointer)
+ || (op == EXEC_ASSIGN && gfc_expr_attr (next->expr2).proc_pointer))
+ gfc_error ("%<OMP DISPATCH%> directive at %L cannot be followed by a "
+ "procedure pointer",
+ &code->loc);
+}
/* Resolve OpenMP directive clauses and check various requirements
of each directive. */
@@ -11696,6 +11864,11 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns)
code->ext.omp_clauses->if_present = false;
resolve_omp_clauses (code, code->ext.omp_clauses, ns);
break;
+ case EXEC_OMP_DISPATCH:
+ if (code->ext.omp_clauses)
+ resolve_omp_clauses (code, code->ext.omp_clauses, ns);
+ resolve_omp_dispatch (code);
+ break;
default:
break;
}
diff --git gcc/fortran/parse.cc gcc/fortran/parse.cc
index 79c810c86ba..74fc249269d 100644
--- gcc/fortran/parse.cc
+++ gcc/fortran/parse.cc
@@ -1050,6 +1050,7 @@ decode_omp_directive (void)
break;
case 'd':
matcho ("depobj", gfc_match_omp_depobj, ST_OMP_DEPOBJ);
+ matcho ("dispatch", gfc_match_omp_dispatch, ST_OMP_DISPATCH);
matchs ("distribute parallel do simd",
gfc_match_omp_distribute_parallel_do_simd,
ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD);
@@ -1916,6 +1917,7 @@ next_statement (void)
case ST_OMP_LOOP: case ST_OMP_PARALLEL_LOOP: case ST_OMP_TEAMS_LOOP: \
case ST_OMP_TARGET_PARALLEL_LOOP: case ST_OMP_TARGET_TEAMS_LOOP: \
case ST_OMP_ALLOCATE_EXEC: case ST_OMP_ALLOCATORS: case ST_OMP_ASSUME: \
+ case ST_OMP_DISPATCH: \
case ST_CRITICAL: \
case ST_OACC_PARALLEL_LOOP: case ST_OACC_PARALLEL: case ST_OACC_KERNELS: \
case ST_OACC_DATA: case ST_OACC_HOST_DATA: case ST_OACC_LOOP: \
@@ -2597,6 +2599,9 @@ gfc_ascii_statement (gfc_statement st, bool strip_sentinel)
case ST_OMP_DEPOBJ:
p = "!$OMP DEPOBJ";
break;
+ case ST_OMP_DISPATCH:
+ p = "!$OMP DISPATCH";
+ break;
case ST_OMP_DISTRIBUTE:
p = "!$OMP DISTRIBUTE";
break;
@@ -6183,6 +6188,35 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only)
}
+static gfc_statement
+parse_omp_dispatch (void)
+{
+ gfc_statement st;
+ gfc_code *cp, *np;
+ gfc_state_data s;
+
+ accept_statement (ST_OMP_DISPATCH);
+
+ cp = gfc_state_stack->tail;
+ push_state (&s, COMP_OMP_STRUCTURED_BLOCK, NULL);
+ np = new_level (cp);
+ np->op = cp->op;
+ np->block = NULL;
+
+ st = next_statement ();
+ if (st == ST_CALL || st == ST_ASSIGNMENT)
+ accept_statement (st);
+ else
+ {
+ gfc_error ("%<OMP DISPATCH%> directive must be followed by a procedure "
+ "call with optional assignment at %C");
+ reject_statement ();
+ }
+ pop_state ();
+ st = next_statement ();
+ return st;
+}
+
/* Accept a series of executable statements. We return the first
statement that doesn't fit to the caller. Any block statements are
passed on to the correct handler, which usually passes the buck
@@ -6383,6 +6417,10 @@ parse_executable (gfc_statement st)
st = parse_omp_oacc_atomic (true);
continue;
+ case ST_OMP_DISPATCH:
+ st = parse_omp_dispatch ();
+ continue;
+
default:
return st;
}
diff --git gcc/fortran/resolve.cc gcc/fortran/resolve.cc
index d7a0856fcca..755d1302ce9 100644
--- gcc/fortran/resolve.cc
+++ gcc/fortran/resolve.cc
@@ -11378,6 +11378,7 @@ gfc_resolve_blocks (gfc_code *b, gfc_namespace *ns)
case EXEC_OMP_ALLOCATORS:
case EXEC_OMP_ASSUME:
case EXEC_OMP_CRITICAL:
+ case EXEC_OMP_DISPATCH:
case EXEC_OMP_DISTRIBUTE:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
@@ -13054,6 +13055,7 @@ start:
case EXEC_OMP_CRITICAL:
case EXEC_OMP_FLUSH:
case EXEC_OMP_DEPOBJ:
+ case EXEC_OMP_DISPATCH:
case EXEC_OMP_DISTRIBUTE:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
diff --git gcc/fortran/st.cc gcc/fortran/st.cc
index 6a605ad91d4..90ee1352ba4 100644
--- gcc/fortran/st.cc
+++ gcc/fortran/st.cc
@@ -222,6 +222,7 @@ gfc_free_statement (gfc_code *p)
case EXEC_OMP_CANCELLATION_POINT:
case EXEC_OMP_CRITICAL:
case EXEC_OMP_DEPOBJ:
+ case EXEC_OMP_DISPATCH:
case EXEC_OMP_DISTRIBUTE:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
diff --git gcc/fortran/trans-decl.cc gcc/fortran/trans-decl.cc
index dca7779528b..4390769146a 100644
--- gcc/fortran/trans-decl.cc
+++ gcc/fortran/trans-decl.cc
@@ -2124,6 +2124,8 @@ get_proc_pointer_decl (gfc_symbol *sym)
return decl;
}
+static void
+create_function_arglist (gfc_symbol *sym);
/* Get a basic decl for an external function. */
@@ -2377,7 +2379,12 @@ module_sym:
if (sym->formal_ns->omp_declare_simd)
gfc_trans_omp_declare_simd (sym->formal_ns);
if (flag_openmp)
- gfc_trans_omp_declare_variant (sym->formal_ns);
+ {
+ // We need DECL_ARGUMENTS to put attributes on, in case some arguments
+ // need adjustment
+ create_function_arglist (sym->formal_ns->proc_name);
+ gfc_trans_omp_declare_variant (sym->formal_ns);
+ }
}
return fndecl;
diff --git gcc/fortran/trans-openmp.cc gcc/fortran/trans-openmp.cc
index f867e2240bf..5e4450184d1 100644
--- gcc/fortran/trans-openmp.cc
+++ gcc/fortran/trans-openmp.cc
@@ -4233,6 +4233,36 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
}
+ if (clauses->novariants)
+ {
+ tree novariants_var;
+
+ gfc_init_se (&se, NULL);
+ gfc_conv_expr (&se, clauses->novariants);
+ gfc_add_block_to_block (block, &se.pre);
+ novariants_var = gfc_evaluate_now (se.expr, block);
+ gfc_add_block_to_block (block, &se.post);
+
+ c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_NOVARIANTS);
+ OMP_CLAUSE_NOVARIANTS_EXPR (c) = novariants_var;
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
+
+ if (clauses->nocontext)
+ {
+ tree nocontext_var;
+
+ gfc_init_se (&se, NULL);
+ gfc_conv_expr (&se, clauses->nocontext);
+ gfc_add_block_to_block (block, &se.pre);
+ nocontext_var = gfc_evaluate_now (se.expr, block);
+ gfc_add_block_to_block (block, &se.post);
+
+ c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_NOCONTEXT);
+ OMP_CLAUSE_NOCONTEXT_EXPR (c) = nocontext_var;
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
+
if (clauses->num_threads)
{
tree num_threads;
@@ -6311,6 +6341,30 @@ gfc_trans_omp_depobj (gfc_code *code)
return gfc_finish_block (&block);
}
+static tree
+gfc_trans_omp_dispatch (gfc_code *code)
+{
+ stmtblock_t block;
+ gfc_code *next = code->block->next;
+ // assume ill-formed "function dispatch structured
+ // block" have already been rejected by resolve_omp_dispatch
+ gcc_assert (next->op == EXEC_CALL || next->op == EXEC_ASSIGN);
+
+ tree body = gfc_trans_code (next);
+ gfc_start_block (&block);
+ tree omp_clauses
+ = gfc_trans_omp_clauses (&block, code->ext.omp_clauses, code->loc);
+
+ tree stmt = make_node (OMP_DISPATCH);
+ SET_EXPR_LOCATION (stmt, gfc_get_location (&code->loc));
+ TREE_TYPE (stmt) = void_type_node;
+ OMP_DISPATCH_BODY (stmt) = body;
+ OMP_DISPATCH_CLAUSES (stmt) = omp_clauses;
+
+ gfc_add_expr_to_block (&block, stmt);
+ return gfc_finish_block (&block);
+}
+
static tree
gfc_trans_omp_error (gfc_code *code)
{
@@ -8221,6 +8275,8 @@ gfc_trans_omp_directive (gfc_code *code)
case EXEC_OMP_TASKLOOP:
return gfc_trans_omp_do (code, code->op, NULL, code->ext.omp_clauses,
NULL);
+ case EXEC_OMP_DISPATCH:
+ return gfc_trans_omp_dispatch (code);
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
case EXEC_OMP_DISTRIBUTE_SIMD:
@@ -8337,6 +8393,7 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns)
tree base_fn_decl = ns->proc_name->backend_decl;
gfc_namespace *search_ns = ns;
gfc_omp_declare_variant *next;
+ vec<tree> adjust_args_list = vNULL;
for (gfc_omp_declare_variant *odv = search_ns->omp_declare_variant;
search_ns; odv = next)
@@ -8532,6 +8589,19 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns)
variant_proc_sym = NULL;
}
}
+ if (odv->need_device_ptr_arg_list != NULL
+ && omp_get_context_selector (set_selectors, OMP_TRAIT_SET_CONSTRUCT,
+ OMP_TRAIT_CONSTRUCT_DISPATCH)
+ == NULL_TREE)
+ {
+ gfc_error ("an %<adjust_args%> clause can only be "
+ "specified if the "
+ "%<dispatch%> selector of the construct "
+ "selector set appears "
+ "in the %<match%> clause at %L",
+ &odv->where);
+ variant_proc_sym = NULL;
+ }
if (variant_proc_sym != NULL)
{
gfc_set_sym_referenced (variant_proc_sym);
@@ -8548,6 +8618,97 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns)
DECL_ATTRIBUTES (base_fn_decl)
= tree_cons (id, build_tree_list (variant, set_selectors),
DECL_ATTRIBUTES (base_fn_decl));
+
+ // Handle adjust_args
+ for (gfc_omp_namelist *arg_list
+ = odv->need_device_ptr_arg_list;
+ arg_list != NULL; arg_list = arg_list->next)
+ {
+ if (arg_list->sym->backend_decl == NULL_TREE)
+ {
+ gfc_error (
+ "%s at %L is not a base function argument",
+ arg_list->sym->name, &arg_list->where);
+ continue;
+ }
+
+ tree base_fn_arg_decl = arg_list->sym->backend_decl;
+ if (base_fn_arg_decl != error_mark_node)
+ {
+ // Is t specified more than once?
+ if (adjust_args_list.contains (base_fn_arg_decl))
+ {
+ gfc_error (
+ "%qD at %L is specified more than once",
+ base_fn_arg_decl, &arg_list->where);
+ continue;
+ }
+ adjust_args_list.safe_push (base_fn_arg_decl);
+
+ // Handle variant argument
+ tree variant
+ = gfc_get_symbol_decl (variant_proc_sym);
+ tree variant_parm = DECL_ARGUMENTS (variant);
+ int idx;
+ tree arg;
+ for (arg = DECL_ARGUMENTS (base_fn_decl), idx = 0;
+ arg != NULL; arg = TREE_CHAIN (arg), idx++)
+ if (arg == base_fn_arg_decl)
+ break;
+ gcc_assert (arg != NULL_TREE);
+ if (variant_parm == NULL_TREE)
+ {
+ gfc_formal_arglist *arg
+ = variant_proc_sym->formal;
+ for (int i = 0; i < idx; i++)
+ {
+ arg = arg->next;
+ gcc_assert (arg != NULL);
+ }
+
+ // Check we got the right parameter name
+ if (strcmp (arg_list->sym->name, arg->sym->name)
+ != 0)
+ {
+ gfc_error ("%s at %L is not a variant "
+ "function argument",
+ arg_list->sym->name,
+ &arg_list->where);
+ continue;
+ }
+ arg->sym->attr
+ .omp_declare_variant_need_device_ptr
+ = 1;
+ }
+ else
+ {
+ for (int i = 0; i < idx; i++)
+ {
+ variant_parm = TREE_CHAIN (variant_parm);
+ gcc_assert (variant_parm != NULL_TREE);
+ }
+ // Check we got the right parameter name
+ if (strcmp (arg_list->sym->name,
+ IDENTIFIER_POINTER (
+ DECL_NAME (variant_parm)))
+ != 0)
+ {
+ gfc_error ("%s at %L is not a variant "
+ "function argument",
+ arg_list->sym->name,
+ &arg_list->where);
+ continue;
+ }
+
+ tree attr = tree_cons (
+ get_identifier (
+ "omp declare variant adjust_args "
+ "need_device_ptr"),
+ NULL_TREE, DECL_ATTRIBUTES (variant_parm));
+ DECL_ATTRIBUTES (variant_parm) = attr;
+ }
+ }
+ }
}
}
}
diff --git gcc/fortran/trans.cc gcc/fortran/trans.cc
index badad6ae892..2795cdf7464 100644
--- gcc/fortran/trans.cc
+++ gcc/fortran/trans.cc
@@ -2596,6 +2596,7 @@ trans_code (gfc_code * code, tree cond)
case EXEC_OMP_CANCELLATION_POINT:
case EXEC_OMP_CRITICAL:
case EXEC_OMP_DEPOBJ:
+ case EXEC_OMP_DISPATCH:
case EXEC_OMP_DISTRIBUTE:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
diff --git gcc/fortran/types.def gcc/fortran/types.def
index 390cc9542f7..5047c8f816a 100644
--- gcc/fortran/types.def
+++ gcc/fortran/types.def
@@ -120,6 +120,7 @@ DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_INT_BOOL, BT_BOOL, BT_INT, BT_BOOL)
DEF_FUNCTION_TYPE_2 (BT_FN_VOID_PTR_PTRMODE,
BT_VOID, BT_PTR, BT_PTRMODE)
DEF_FUNCTION_TYPE_2 (BT_FN_VOID_CONST_PTR_SIZE, BT_VOID, BT_CONST_PTR, BT_SIZE)
+DEF_FUNCTION_TYPE_2 (BT_FN_PTR_CONST_PTR_INT, BT_PTR, BT_CONST_PTR, BT_INT)
DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR_PTR, BT_FN_VOID_PTR_PTR)
diff --git gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90 gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90
new file mode 100644
index 00000000000..68adb60a397
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/adjust-args-1.f90
@@ -0,0 +1,63 @@
+! Test parsing of OMP clause adjust_args
+! { dg-do compile }
+
+module main
+ use iso_c_binding, only: c_ptr, c_funptr
+ implicit none
+ integer :: b
+ interface
+ integer function f0 (a)
+ import c_ptr
+ type(c_ptr), intent(inout) :: a
+ end function
+ integer function g (a)
+ import c_ptr
+ type(c_ptr), intent(inout) :: a
+ end function
+ integer function h (a)
+ import c_funptr
+ type(c_funptr), intent(inout) :: a
+ end function
+ integer function f1 (i)
+ integer, intent(in) :: i
+ end function
+
+ integer function f3 (a)
+ import c_ptr
+ type(c_ptr), intent(inout) :: a
+ !$omp declare variant (f1) match (construct={dispatch}) adjust_args (other: a) ! { dg-error "expected 'nothing' or 'need_device_ptr' at .1." }
+ end function
+ integer function f4 (a)
+ import c_ptr
+ type(c_ptr), intent(inout) :: a
+ !$omp declare variant (f0) adjust_args (nothing: a) ! { dg-error "an 'adjust_args' clause at .1. can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" }
+ end function
+ integer function f5 (i)
+ integer, intent(inout) :: i
+ !$omp declare variant (f1) match (construct={dispatch}) adjust_args () ! { dg-error "expected 'nothing' or 'need_device_ptr' at .1." }
+ end function
+ integer function f6 (i)
+ integer, intent(inout) :: i
+ !$omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing) ! { dg-error "expected argument list at .1." }
+ end function
+ integer function f7 (i)
+ integer, intent(inout) :: i
+ !$omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing:) ! { dg-error "expected argument list at .1." }
+ end function
+ integer function f9 (i)
+ integer, intent(inout) :: i
+ !$omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: i) ! { dg-error "argument list item 'i' in 'need_device_ptr' at .1. must be of TYPE.C_PTR." }
+ end function
+ integer function f12 (a)
+ import c_ptr
+ type(c_ptr), intent(inout) :: a
+ !$omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: b) ! { dg-error "list item 'b' at .1. is not a dummy argument" }
+ end function
+ integer function f13 (a)
+ import c_funptr
+ type(c_funptr), intent(inout) :: a
+ !$omp declare variant (h) match (construct={dispatch}) adjust_args (need_device_ptr: a) ! { dg-error "argument list item 'a' in 'need_device_ptr' at .1. must be of TYPE.C_PTR." }
+ end function
+
+ end interface
+end module
diff --git gcc/testsuite/gfortran.dg/gomp/adjust-args-2.f90 gcc/testsuite/gfortran.dg/gomp/adjust-args-2.f90
new file mode 100644
index 00000000000..c65a4839ca5
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/adjust-args-2.f90
@@ -0,0 +1,18 @@
+! Test resolution of OMP clause adjust_args
+! { dg-do compile }
+
+module main
+ implicit none
+interface
+subroutine f1 (i)
+ integer, intent(inout) :: i
+end subroutine
+end interface
+contains
+
+ subroutine f3 (i)
+ integer, intent(inout) :: i
+ !$omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: z) ! { dg-error "Symbol 'z' at .1. has no IMPLICIT type" }
+ end subroutine
+
+end module
diff --git gcc/testsuite/gfortran.dg/gomp/adjust-args-3.f90 gcc/testsuite/gfortran.dg/gomp/adjust-args-3.f90
new file mode 100644
index 00000000000..b731cb340c1
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/adjust-args-3.f90
@@ -0,0 +1,26 @@
+! Test translation of OMP clause adjust_args
+! { dg-do compile }
+
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+ !type(c_ptr) :: a
+
+contains
+ subroutine base2 (a)
+ type(c_ptr), intent(inout) :: a
+ !$omp declare variant (variant2) match (construct={parallel}) adjust_args (need_device_ptr: a) ! { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause at .1." }
+ end subroutine
+ subroutine base3 (a)
+ type(c_ptr), intent(inout) :: a
+ !$omp declare variant (variant2) match (construct={dispatch}) adjust_args (need_device_ptr: a) adjust_args (need_device_ptr: a) ! { dg-error "'a' at .1. is specified more than once" }
+ end subroutine
+
+ subroutine variant2 (a)
+ type(c_ptr), intent(inout) :: a
+ end subroutine
+ subroutine variant3 (i)
+ integer :: i
+ end subroutine
+
+end module
diff --git gcc/testsuite/gfortran.dg/gomp/adjust-args-4.f90 gcc/testsuite/gfortran.dg/gomp/adjust-args-4.f90
new file mode 100644
index 00000000000..75e884044b2
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/adjust-args-4.f90
@@ -0,0 +1,58 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+
+ type :: struct
+ integer :: a
+ real :: b
+ end type
+
+ interface
+ integer function f(a, b, c)
+ import c_ptr
+ integer, intent(in) :: a
+ type(c_ptr), intent(inout) :: b
+ type(c_ptr), intent(out) :: c(:)
+ end function
+ integer function f0(a, b, c)
+ import c_ptr
+ integer, intent(in) :: a
+ type(c_ptr), intent(inout) :: b
+ type(c_ptr), intent(out) :: c(:)
+ !$omp declare variant (f) match (construct={dispatch}) &
+ !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b, c)
+ end function
+ integer function f1(a, b, c)
+ import c_ptr
+ integer, intent(in) :: a
+ type(c_ptr), intent(inout) :: b
+ type(c_ptr), intent(out) :: c(:)
+ !$omp declare variant (f) match (construct={dispatch}) &
+ !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b) adjust_args (need_device_ptr: c)
+ end function
+ end interface
+
+contains
+subroutine test
+ integer :: a
+ type(c_ptr) :: b
+ type(c_ptr) :: c(2)
+ type(struct) :: s
+
+ s%a = f0 (a, b, c)
+ !$omp dispatch
+ s%a = f0 (a, b, c)
+
+ s%b = f1 (a, b, c)
+ !$omp dispatch
+ s%b = f1 (a, b, c)
+
+end subroutine
+end module
+
+! { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&parm\.\[0-9]+, D\.\[0-9]+\\);" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&b, D\.\[0-9]+\\);" 2 "gimple" } }
diff --git gcc/testsuite/gfortran.dg/gomp/adjust-args-5.f90 gcc/testsuite/gfortran.dg/gomp/adjust-args-5.f90
new file mode 100644
index 00000000000..75e884044b2
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/adjust-args-5.f90
@@ -0,0 +1,58 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+
+ type :: struct
+ integer :: a
+ real :: b
+ end type
+
+ interface
+ integer function f(a, b, c)
+ import c_ptr
+ integer, intent(in) :: a
+ type(c_ptr), intent(inout) :: b
+ type(c_ptr), intent(out) :: c(:)
+ end function
+ integer function f0(a, b, c)
+ import c_ptr
+ integer, intent(in) :: a
+ type(c_ptr), intent(inout) :: b
+ type(c_ptr), intent(out) :: c(:)
+ !$omp declare variant (f) match (construct={dispatch}) &
+ !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b, c)
+ end function
+ integer function f1(a, b, c)
+ import c_ptr
+ integer, intent(in) :: a
+ type(c_ptr), intent(inout) :: b
+ type(c_ptr), intent(out) :: c(:)
+ !$omp declare variant (f) match (construct={dispatch}) &
+ !$omp& adjust_args (nothing: a) adjust_args (need_device_ptr: b) adjust_args (need_device_ptr: c)
+ end function
+ end interface
+
+contains
+subroutine test
+ integer :: a
+ type(c_ptr) :: b
+ type(c_ptr) :: c(2)
+ type(struct) :: s
+
+ s%a = f0 (a, b, c)
+ !$omp dispatch
+ s%a = f0 (a, b, c)
+
+ s%b = f1 (a, b, c)
+ !$omp dispatch
+ s%b = f1 (a, b, c)
+
+end subroutine
+end module
+
+! { dg-final { scan-tree-dump-times "__builtin_omp_get_default_device \\(\\);" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&parm\.\[0-9]+, D\.\[0-9]+\\);" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&b, D\.\[0-9]+\\);" 2 "gimple" } }
diff --git gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90 gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90
index 7fc5071feff..62d2cb96fac 100644
--- gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90
+++ gcc/testsuite/gfortran.dg/gomp/declare-variant-2.f90
@@ -18,10 +18,10 @@ contains
!$omp declare variant match(user={condition(.false.)}) ! { dg-error "expected '\\(' at .1." }
end subroutine
subroutine f6 ()
- !$omp declare variant (f1) ! { dg-error "expected 'match' at .1." }
+ !$omp declare variant (f1) ! { dg-error "expected 'match' or 'adjust_args' at .1." }
end subroutine
subroutine f7 ()
- !$omp declare variant (f1) simd ! { dg-error "expected 'match' at .1." }
+ !$omp declare variant (f1) simd ! { dg-error "expected 'match' or 'adjust_args' at .1." }
end subroutine
subroutine f8 ()
!$omp declare variant (f1) match ! { dg-error "expected '\\(' at .1." }
@@ -183,7 +183,7 @@ contains
!$omp declare variant (f1) match(construct={requires}) ! { dg-warning "unknown selector 'requires' for context selector set 'construct' at .1." }
end subroutine
subroutine f75 ()
- !$omp declare variant (f1),match(construct={parallel}) ! { dg-error "expected 'match' at .1." }
+ !$omp declare variant (f1),match(construct={parallel}) ! { dg-error "expected 'match' or 'adjust_args' at .1." }
end subroutine
subroutine f76 ()
!$omp declare variant (f1) match(implementation={atomic_default_mem_order("relaxed")}) ! { dg-error "expected identifier at .1." }
diff --git gcc/testsuite/gfortran.dg/gomp/declare-variant-21-aux.f90 gcc/testsuite/gfortran.dg/gomp/declare-variant-21-aux.f90
new file mode 100644
index 00000000000..4e8bb129d40
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/declare-variant-21-aux.f90
@@ -0,0 +1,18 @@
+! { dg-do compile { target skip-all-targets } }
+
+module my_mod
+ use iso_c_binding, only: c_ptr
+ implicit none
+ interface
+ subroutine base_proc (a)
+ use iso_c_binding, only: c_ptr
+ type(c_ptr), intent(inout) :: a
+ end subroutine
+ end interface
+
+contains
+ subroutine variant_proc (a)
+ type(c_ptr), intent(inout) :: a
+ !$omp declare variant (base_proc) match (construct={dispatch}) adjust_args(need_device_ptr: a)
+ end subroutine
+end module
diff --git gcc/testsuite/gfortran.dg/gomp/declare-variant-21.f90 gcc/testsuite/gfortran.dg/gomp/declare-variant-21.f90
new file mode 100644
index 00000000000..022ae04dac0
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/declare-variant-21.f90
@@ -0,0 +1,28 @@
+! { dg-do run }
+! { dg-additional-sources declare-variant-21-aux.f90 }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+! Test XFAILed due to https://gcc.gnu.org/PR115271
+
+
+subroutine base_proc (a)
+ use iso_c_binding, only: c_ptr
+ type(c_ptr), intent(inout) :: a
+end subroutine
+
+program main
+ use iso_c_binding, only: c_ptr
+ use my_mod
+ implicit none
+
+ type(c_ptr) :: a
+
+
+ call base_proc(a)
+ !call variant_proc(a)
+
+ !$omp dispatch
+ call base_proc(a)
+! { dg-final { scan-tree-dump "variant_proc \\(&a\\)" "gimple" { xfail *-*-* } } }
+
+end program main
diff --git gcc/testsuite/gfortran.dg/gomp/dispatch-1.f90 gcc/testsuite/gfortran.dg/gomp/dispatch-1.f90
new file mode 100644
index 00000000000..12c30904131
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/dispatch-1.f90
@@ -0,0 +1,77 @@
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+ contains
+
+ subroutine f1 ()
+ integer :: a, b, arr(10)
+ real :: x
+ complex :: c
+ character :: ch
+ logical :: bool
+ type :: struct
+ integer :: a
+ real :: b
+ end type
+ type(struct) :: s
+ type(c_ptr) :: p
+
+ interface
+ subroutine f0 (a, c, bool, s)
+ import :: struct
+ integer, intent(in) :: a
+ complex, intent(out) :: c
+ logical, intent(inout) :: bool
+ type(struct) :: s
+ end subroutine
+ integer function f2 (arr, x, ch, b)
+ integer, intent(inout) :: arr(:)
+ real, intent(in) :: x
+ character, intent(out) :: ch
+ real :: b
+ end function
+ subroutine f3 (p)
+ import :: c_ptr
+ type(c_ptr) :: p
+ end subroutine
+ integer function f4 ()
+ end function
+ end interface
+
+ !$omp dispatch
+ b = f2(arr, x, ch, s%b)
+ !$omp dispatch
+ c = f2(arr(:5), x * 2.4, ch, s%b)
+ !$omp dispatch
+ arr(1) = f2(arr, x, ch, s%b)
+ !$omp dispatch
+ s%a = f2(arr, x, ch, s%b)
+ !$omp dispatch
+ x = f2(arr, x, ch, s%b)
+ !$omp dispatch
+ call f0(a, c, bool, s)
+ !$omp dispatch
+ call f0(f4(), c, bool, s)
+
+ !$omp dispatch nocontext(.TRUE.)
+ call f0(a, c, bool, s)
+ !$omp dispatch nocontext(arr(2) < 10)
+ call f0(a, c, bool, s)
+ !$omp dispatch novariants(.FALSE.)
+ call f0(a, c, bool, s)
+ !$omp dispatch novariants(bool)
+ call f0(a, c, bool, s)
+ !$omp dispatch nowait
+ call f0(a, c, bool, s)
+ !$omp dispatch device(arr(9))
+ call f0(a, c, bool, s)
+ !$omp dispatch device(a + a)
+ call f0(a, c, bool, s)
+ !$omp dispatch device(-25373654)
+ call f0(a, c, bool, s)
+ !$omp dispatch is_device_ptr(p)
+ call f3(p)
+ !$omp dispatch depend(in: a, c, bool) depend(inout: s, arr(:3))
+ call f0(a, c, bool, s)
+ end subroutine
+end module
diff --git gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90 gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90
new file mode 100644
index 00000000000..d2d555b5932
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/dispatch-2.f90
@@ -0,0 +1,79 @@
+module main
+ use iso_c_binding, only: c_funptr
+ implicit none
+ contains
+
+ subroutine f1 ()
+ integer :: a, b, arr(10)
+ real :: x
+ complex :: c
+ character :: ch
+ logical :: bool
+ type :: struct
+ integer :: a
+ real :: b
+ end type
+ type(struct) :: s
+ type(c_funptr) :: p
+
+ interface
+ subroutine f0 (a, c, bool, s)
+ import :: struct
+ integer, intent(in) :: a
+ complex, intent(out) :: c
+ logical, intent(inout) :: bool
+ type(struct) :: s
+ end subroutine
+ integer function f2 (arr, x, ch, b)
+ integer, intent(inout) :: arr(:)
+ real, intent(in) :: x
+ character, intent(out) :: ch
+ real :: b
+ end function
+ end interface
+ procedure(f0), pointer:: fp => NULL()
+
+ !$omp dispatch !{ dg-error "'OMP DISPATCH' directive at .1. must be followed by a procedure call with optional assignment" }
+50 b = f2(arr, x, ch, s%b) + a
+ !$omp dispatch !{ dg-error "'OMP DISPATCH' directive at .1. must be followed by a procedure call with optional assignment" }
+ a = b
+ !$omp dispatch !{ dg-error "'OMP DISPATCH' directive at .1. must be followed by a procedure call with optional assignment" }
+ b = Not (2)
+ !$omp dispatch
+ !$omp threadprivate(a) !{ dg-error "'OMP DISPATCH' directive must be followed by a procedure call with optional assignment at .1." }
+ a = f2(arr, x, ch, s%b)
+ !$omp dispatch
+ print *, 'This is not allowed here.' !{ dg-error "'OMP DISPATCH' directive must be followed by a procedure call with optional assignment at .1." }
+ !$omp dispatch
+ goto 50 !{ dg-error "'OMP DISPATCH' directive must be followed by a procedure call with optional assignment at .1." }
+ !$omp dispatch !{ dg-error "'OMP DISPATCH' directive at .1. cannot be followed by a procedure pointer" }
+ call fp(a, c, bool, s)
+
+ !$omp dispatch nocontext(s) !{ dg-error "NOCONTEXT clause at .1. requires a scalar LOGICAL expression" }
+ call f0(a, c, bool, s)
+ !$omp dispatch nocontext(a, b) !{ dg-error "Invalid expression after 'nocontext.' at .1." }
+ call f0(a, c, bool, s)
+ !$omp dispatch nocontext(a) nocontext(b) !{ dg-error "Duplicated 'nocontext' clause at .1." }
+ call f0(a, c, bool, s)
+ !$omp dispatch novariants(s) !{ dg-error "NOVARIANTS clause at .1. requires a scalar LOGICAL expression" }
+ call f0(a, c, bool, s)
+ !$omp dispatch novariants(a, b) !{ dg-error "Invalid expression after 'novariants.' at .1." }
+ call f0(a, c, bool, s)
+ !$omp dispatch novariants(a) novariants(b) !{ dg-error "Duplicated 'novariants' clause at .1." }
+ call f0(a, c, bool, s)
+ !$omp dispatch nowait nowait !{ dg-error "Duplicated 'nowait' clause at .1." }
+ call f0(a, c, bool, s)
+ !$omp dispatch device(x) !{ dg-error "DEVICE clause at .1. requires a scalar INTEGER expression" }
+ call f0(a, c, bool, s)
+ !$omp dispatch device(arr) !{ dg-error "DEVICE clause at .1. requires a scalar INTEGER expression" }
+ call f0(a, c, bool, s)
+ !$omp dispatch is_device_ptr(x) !{ dg-error "List item 'x' in IS_DEVICE_PTR clause at .1. must be of TYPE.C_PTR." }
+ call f0(a, c, bool, s)
+ !$omp dispatch is_device_ptr(arr) !{ dg-error "List item 'arr' in IS_DEVICE_PTR clause at .1. must be of TYPE.C_PTR." }
+ call f0(a, c, bool, s)
+ !$omp dispatch is_device_ptr(p) !{ dg-error "List item 'p' in IS_DEVICE_PTR clause at .1. must be of TYPE.C_PTR." }
+ call f0(a, c, bool, s)
+ !$omp dispatch depend(inout: f0) !{ dg-error "Object 'f0' is not a variable at .1." }
+ call f0(a, c, bool, s)
+ end subroutine
+end module
diff --git gcc/testsuite/gfortran.dg/gomp/dispatch-3.f90 gcc/testsuite/gfortran.dg/gomp/dispatch-3.f90
new file mode 100644
index 00000000000..84590fd883a
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/dispatch-3.f90
@@ -0,0 +1,39 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module main
+ implicit none
+ interface
+ integer function f0 ()
+ end function
+
+ integer function f1 ()
+ end function
+
+ integer function f2 ()
+ !$omp declare variant (f0) match (construct={dispatch})
+ !$omp declare variant (f1) match (implementation={vendor(gnu)})
+ end function
+ end interface
+ contains
+
+ integer function test ()
+ integer :: a
+
+ !$omp dispatch
+ a = f2 ()
+ !$omp dispatch novariants(.TRUE.)
+ a = f2 ()
+ !$omp dispatch novariants(.FALSE.)
+ a = f2 ()
+ !$omp dispatch nocontext(.TRUE.)
+ a = f2 ()
+ !$omp dispatch nocontext(.FALSE.)
+ a = f2 ()
+ end function
+end module
+
+
+! { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 3 "gimple" } }
+! { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } }
diff --git gcc/testsuite/gfortran.dg/gomp/dispatch-4.f90 gcc/testsuite/gfortran.dg/gomp/dispatch-4.f90
new file mode 100644
index 00000000000..149d0613b97
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/dispatch-4.f90
@@ -0,0 +1,19 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module main
+ implicit none
+ interface
+ subroutine f2 ()
+ end subroutine
+ end interface
+ contains
+
+ subroutine test ()
+ !$omp dispatch ! { dg-final { scan-tree-dump-times "#pragma omp task if\\(0\\)" 1 "gimple" } }
+ call f2 ()
+ !$omp dispatch nowait ! { dg-final { scan-tree-dump-times "#pragma omp task if\\(1\\)" 1 "gimple" } }
+ call f2 ()
+ end subroutine
+end module
+
diff --git gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90 gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90
new file mode 100644
index 00000000000..e45397f3f96
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/dispatch-5.f90
@@ -0,0 +1,24 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module main
+ implicit none
+ interface
+ subroutine f2 (a)
+ integer, intent(in) :: a
+ end subroutine
+ end interface
+ contains
+
+ subroutine test ()
+ integer :: a
+
+ !$omp dispatch device(-25373654)
+ ! { dg-final { scan-tree-dump-times "__builtin_omp_set_default_device \\(-25373654\\);" 1 "gimple" } }
+ call f2 (a)
+ !$omp dispatch device(a + a)
+ ! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = a.0_1 \\* 2;.*#pragma omp dispatch device\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\).*#pragma omp task shared\\(D\.\[0-9]+\\).*__builtin_omp_set_default_device \\(D\.\[0-9]+\\);" 1 "gimple" } }
+ call f2 (a)
+ end subroutine
+end module
+
diff --git gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90 gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90
new file mode 100644
index 00000000000..9f4fa2970ca
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/dispatch-6.f90
@@ -0,0 +1,38 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+ interface
+ subroutine f1 (p, p2)
+ import :: c_ptr
+ type(c_ptr), intent(out) :: p
+ type(c_ptr), intent(in) :: p2
+ end subroutine
+ subroutine f2 (p, p2)
+ import :: c_ptr
+ type(c_ptr), intent(out) :: p
+ type(c_ptr), intent(in) :: p2
+ !$omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: p, p2)
+ end subroutine
+ end interface
+ contains
+
+ subroutine test ()
+ type(c_ptr) :: p, p2
+
+ !$omp dispatch
+ call f2 (p, p2)
+ !$omp dispatch is_device_ptr(p)
+ ! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p\\) shared\\(p2\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*integer\\(kind=4\\) D\.\[0-9]+;\[ \t\n\r]*void \\* D\.\[0-9]+;\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&p2, D\.\[0-9]+\\);\[ \t\n\r]*f1 \\(&p, D\.\[0-9]+\\);" 1 "gimple" } }
+ call f2 (p, p2)
+ !$omp dispatch is_device_ptr(p2)
+ ! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p2\\) shared\\(p\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*integer\\(kind=4\\) D\.\[0-9]+;\[ \t\n\r]*void \\* D\.\[0-9]+;\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_default_device \\(\\);\[ \t\n\r]*D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(&p, D\.\[0-9]+\\);\[ \t\n\r]*f1 \\(D\.\[0-9]+, &p2\\);" 1 "gimple" } }
+ call f2 (p, p2)
+ !$omp dispatch is_device_ptr(p, p2)
+ ! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(p\\) shared\\(p2\\)\[^\n\r]*\[ \t\n\r]*\{\[ \t\n\r]*p = {CLOBBER};\[ \t\n\r]*f1 \\(&p, &p2\\);" 1 "gimple" } }
+ call f2 (p, p2)
+ end subroutine
+end module
+
diff --git gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90 gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90
new file mode 100644
index 00000000000..32b6347be67
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/dispatch-7.f90
@@ -0,0 +1,27 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-ompexp" }
+
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+ interface
+ subroutine f2 (p)
+ import :: c_ptr
+ type(c_ptr), intent(out) :: p
+ end subroutine
+ end interface
+ contains
+
+ subroutine test ()
+ type(c_ptr) :: p
+
+ !$omp dispatch
+ ! { dg-final { scan-tree-dump-times "__builtin_GOMP_task \\(.*, .*, .*, .*, .*, .*, 0B, .*, .*\\);" 1 "ompexp" } }
+ call f2 (p)
+ !$omp dispatch depend(inout: p)
+ ! { dg-final { scan-tree-dump-times "D\.\[0-9]+\\\[2] = &p;" 1 "ompexp" } }
+ ! { dg-final { scan-tree-dump-times "__builtin_GOMP_task \\(.*, .*, .*, .*, .*, .*, &D\.\[0-9]+, .*, .*\\);" 1 "ompexp" } }
+ call f2 (p)
+ end subroutine
+end module
+
diff --git gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90 gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90
new file mode 100644
index 00000000000..6771336aa33
--- /dev/null
+++ gcc/testsuite/gfortran.dg/gomp/dispatch-8.f90
@@ -0,0 +1,39 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple -fdump-tree-omplower" }
+
+module main
+ use iso_c_binding, only: c_ptr
+ implicit none
+ interface
+ integer function f0 ()
+ end function
+ integer function f1 ()
+ end function
+ integer function f2 ()
+ !$omp declare variant (f0) match (construct={dispatch})
+ !$omp declare variant (f1) match (implementation={vendor(gnu)})
+ end function
+ end interface
+ contains
+
+ subroutine test ()
+ integer :: a, n
+
+ !$omp dispatch novariants(n < 1024) nocontext(n > 1024)
+ a = f2 ()
+ end subroutine
+end module
+
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = n <= 1023;" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = n > 1024;" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp dispatch novariants\\(0\\) nocontext\\(0\\) shared\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp task shared\\(D\.\[0-9]+\\) shared\\(D\.\[0-9]+\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "a = f2 \\\(\\\);" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "a = f1 \\\(\\\);" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "a = f0 \\\(\\\);" 1 "gimple" } }
+
+! { dg-final { scan-tree-dump-times ".omp_data_o.1.D\.\[0-9]+ = D\.\[0-9]+;" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times ".omp_data_o.1.a = &a;" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = .omp_data_i->D\.\[0-9]+;" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "D\.\[0-9]+ = .omp_data_i->a;" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "\\*D\.\[0-9]+ = D\.\[0-9]+;" 3 "omplower" } }
^ permalink raw reply [flat|nested] 10+ messages in thread
* [PATCH 7/7] OpenMP: update documentation for dispatch and adjust_args
2024-05-27 11:54 [PATCH 0/7] OpenMP: dispatch + adjust_args support Paul-Antoine Arras
` (5 preceding siblings ...)
2024-05-27 11:54 ` [PATCH 6/7] OpenMP: Fortran front-end support " Paul-Antoine Arras
@ 2024-05-27 11:54 ` Paul-Antoine Arras
6 siblings, 0 replies; 10+ messages in thread
From: Paul-Antoine Arras @ 2024-05-27 11:54 UTC (permalink / raw)
To: gcc-patches; +Cc: Paul-Antoine Arras
libgomp/ChangeLog:
* libgomp.texi:
---
libgomp/libgomp.texi | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 71d62105a20..b72accd0d26 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -294,8 +294,8 @@ The OpenMP 4.5 specification is fully supported.
@item C/C++'s @code{declare variant} directive: elision support of
preprocessed code @tab N @tab
@item @code{declare variant}: new clauses @code{adjust_args} and
- @code{append_args} @tab N @tab
-@item @code{dispatch} construct @tab N @tab
+ @code{append_args} @tab P @tab Only @code{adjust_args}
+@item @code{dispatch} construct @tab Y @tab
@item device-specific ICV settings with environment variables @tab Y @tab
@item @code{assume} and @code{assumes} directives @tab Y @tab
@item @code{nothing} directive @tab Y @tab
--
2.45.1
^ permalink raw reply [flat|nested] 10+ messages in thread