https://gcc.gnu.org/g:e34748375a0124c3f5f4abb3ade24913fe688c99
commit e34748375a0124c3f5f4abb3ade24913fe688c99 Author: Tobias Burnus <tbur...@baylibre.com> Date: Mon Jan 27 12:32:35 2025 +0100 OpenMP: 'interop' construct - add C/C++ parser support, improve Fortran parsing Add middle end support for the 'interop' directive and the 'init', 'use', and 'destroy' clauses - but fail with a sorry, unimplemented in gimplify.cc. For Fortran, generate the tree code, update the internal representation, add some more diagnostic checks and update for newer specification changes ('fr' only takes a single value, but it integer expressions are permitted again [like with the old syntax] not only constant identifiers). For C and C++, this patch adds the full parser support for 'interop'. Still missing is actually handling the directive in the middle end and in libgomp. The GOMP_INTEROP_IFR_* internal values have been changed to have space for vendor specific values that are adjacent to the existing values but negative, if needed. gcc/c-family/ChangeLog: * c-common.h (enum c_omp_region_type): Add C_ORT_INTEROP and C_ORT_OMP_INTEROP. (c_omp_interop_t_p): New prototype. * c-omp.cc (c_omp_interop_t_p): Check whether the type is omp_interop_t. (c_omp_directives): Uncomment 'interop'. * c-pragma.cc (omp_pragmas): Add 'interop'. * c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_INTEROP. (enum pragma_omp_clause): Add init, use, and destroy clauses. gcc/c/ChangeLog: * c-parser.cc (INCLUDE_STRING): Define. (c_parser_pragma): Handle 'interop' directive. (c_parser_omp_clause_name): Handle init, use, and destroy clauses. (c_parser_omp_all_clauses): Likewise; use C_ORT_OMP_INTEROP, if 'use' is permitted, for c_finish_omp_clauses. (c_parser_omp_clause_destroy, c_parser_omp_modifier_prefer_type, c_parser_omp_clause_init, c_parser_omp_clause_use, OMP_INTEROP_CLAUSE_MASK, c_parser_omp_interop): New. * c-typeck.cc (c_finish_omp_clauses): Add missing OPT_Wopenmp to a warning; handle new clauses. gcc/cp/ChangeLog: * parser.cc (INCLUDE_STRING): Define. (cp_parser_omp_clause_name): Handle init, use, and destroy clauses. (cp_parser_omp_all_clauses): Likewise; use C_ORT_OMP_INTEROP, if 'use' is permitted, for c_finish_omp_clauses. (cp_parser_omp_modifier_prefer_type, cp_parser_omp_clause_init, OMP_INTEROP_CLAUSE_MASK, cp_parser_omp_interop): New. (cp_parser_pragma): Handle 'interop' directive. * pt.cc (tsubst_omp_clauses): Handle init, use, and destroy clauses. (tsubst_stmt): Handle OMP_INTEROP. * semantics.cc (cp_omp_init_prefer_type_update): New. (finish_omp_clauses): Handle init, use, and destroy clauses and add clause check for 'depend' on 'interop'. gcc/fortran/ChangeLog: * gfortran.h (gfc_omp_namelist): Cleanup interop internal representation. * dump-parse-tree.cc (show_omp_namelist): Update for changed internal representation. * match.cc (gfc_free_omp_namelist): Likewise. * openmp.cc (gfc_match_omp_prefer_type, gfc_match_omp_init): Likewise; also handle some corner cases better and update for newer 6.0 changes related to 'fr'. (resolve_omp_clauses): Add type-check for interop variables. * trans-openmp.cc (gfc_trans_omp_clauses): Handle init, use and destroy clauses. (gfc_trans_openmp_interop): New. (gfc_trans_omp_directive): Call it. gcc/ChangeLog: * gimplify.cc (gimplify_expr): Handle OMP_INTEROP by printing "sorry, uninplemented". * omp-api.h (omp_get_fr_id_from_name): Change return type to 'char'. * omp-general.cc (omp_get_fr_id_from_name): Likewise; return GOMP_INTEROP_IFR_UNKNOWN not 0 if not found. (omp_get_name_from_fr_id): Return "<unknown>" not NULL if not found (used for dumps). * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_DESTROY, OMP_CLAUSE_USE, and OMP_CLAUSE_INIT. * tree-pretty-print.cc (dump_omp_init_prefer_type): New. (dump_omp_clause): Handle init, use and destroy clauses. (dump_generic_node): Handle interop directive. * tree.cc (omp_clause_num_ops, omp_clause_code_name): Add new init/use/destroy clauses. * tree.def (OACC_LOOP): Fix comment. (OMP_INTEROP): Add. * tree.h (OMP_INTEROP_CLAUSES, OMP_CLAUSE_INIT_TARGET, OMP_CLAUSE_INIT_TARGETSYNC, OMP_CLAUSE_INIT_PREFER_TYPE): New. include/ChangeLog: * gomp-constants.h (GOMP_INTEROP_IFR_NONE): Rename ... (GOMP_INTEROP_IFR_UNKNOWN): ... to this. And change value. (GOMP_INTEROP_IFR_SEPARATOR): Likewise. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/interop-1.f90: Update for parser changes, spec changes and add new tests. * gfortran.dg/gomp/interop-2.f90: Likewise. * gfortran.dg/gomp/interop-3.f90: Likewise. * c-c++-common/gomp/interop-1.c: New test. * c-c++-common/gomp/interop-2.c: New test. * c-c++-common/gomp/interop-3.c: New test. * c-c++-common/gomp/interop-4.c: New test. * g++.dg/gomp/interop-5.C: New test. * gfortran.dg/gomp/interop-4.f90: New test. (cherry picked from commit 8f0c8e577a56891fa104c818834ddafe268722bb) Diff: --- gcc/ChangeLog.omp | 25 ++ gcc/c-family/ChangeLog.omp | 15 + gcc/c-family/c-common.h | 7 +- gcc/c-family/c-omp.cc | 22 +- gcc/c-family/c-pragma.cc | 1 + gcc/c-family/c-pragma.h | 4 + gcc/c/ChangeLog.omp | 16 + gcc/c/c-parser.cc | 436 ++++++++++++++++++++++++ gcc/c/c-typeck.cc | 47 ++- gcc/cp/ChangeLog.omp | 18 + gcc/cp/parser.cc | 485 ++++++++++++++++++++++++++- gcc/cp/pt.cc | 38 +++ gcc/cp/semantics.cc | 124 ++++++- gcc/fortran/ChangeLog.omp | 19 ++ gcc/fortran/dump-parse-tree.cc | 69 ++-- gcc/fortran/gfortran.h | 3 +- gcc/fortran/match.cc | 9 +- gcc/fortran/openmp.cc | 234 ++++++------- gcc/fortran/trans-openmp.cc | 59 +++- gcc/gimplify.cc | 5 + gcc/omp-api.h | 2 +- gcc/omp-general.cc | 6 +- gcc/testsuite/ChangeLog.omp | 16 + gcc/testsuite/c-c++-common/gomp/interop-1.c | 119 +++++++ gcc/testsuite/c-c++-common/gomp/interop-2.c | 127 +++++++ gcc/testsuite/c-c++-common/gomp/interop-3.c | 82 +++++ gcc/testsuite/c-c++-common/gomp/interop-4.c | 75 +++++ gcc/testsuite/g++.dg/gomp/interop-5.C | 90 +++++ gcc/testsuite/gfortran.dg/gomp/interop-1.f90 | 43 ++- gcc/testsuite/gfortran.dg/gomp/interop-2.f90 | 40 ++- gcc/testsuite/gfortran.dg/gomp/interop-3.f90 | 21 +- gcc/testsuite/gfortran.dg/gomp/interop-4.f90 | 56 ++++ gcc/tree-core.h | 13 + gcc/tree-pretty-print.cc | 87 +++++ gcc/tree.cc | 6 + gcc/tree.def | 6 +- gcc/tree.h | 12 + include/ChangeLog.omp | 9 + include/gomp-constants.h | 7 +- 39 files changed, 2241 insertions(+), 212 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 5b5017ee56ac..cf11ac6d3242 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,28 @@ +2025-01-27 Tobias Burnus <tbur...@baylibre.com> + + Backported from master: + 2024-11-22 Tobias Burnus <tbur...@baylibre.com> + + * gimplify.cc (gimplify_expr): Handle OMP_INTEROP by printing + "sorry, uninplemented". + * omp-api.h (omp_get_fr_id_from_name): Change return type to + 'char'. + * omp-general.cc (omp_get_fr_id_from_name): Likewise; return + GOMP_INTEROP_IFR_UNKNOWN not 0 if not found. + (omp_get_name_from_fr_id): Return "<unknown>" not NULL + if not found (used for dumps). + * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_DESTROY, + OMP_CLAUSE_USE, and OMP_CLAUSE_INIT. + * tree-pretty-print.cc (dump_omp_init_prefer_type): New. + (dump_omp_clause): Handle init, use and destroy clauses. + (dump_generic_node): Handle interop directive. + * tree.cc (omp_clause_num_ops, omp_clause_code_name): Add new + init/use/destroy clauses. + * tree.def (OACC_LOOP): Fix comment. + (OMP_INTEROP): Add. + * tree.h (OMP_INTEROP_CLAUSES, OMP_CLAUSE_INIT_TARGET, + OMP_CLAUSE_INIT_TARGETSYNC, OMP_CLAUSE_INIT_PREFER_TYPE): New. + 2025-01-27 Paul-Antoine Arras <par...@baylibre.com> Backported from master: diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp index 75a1c1795410..852acb8a2039 100644 --- a/gcc/c-family/ChangeLog.omp +++ b/gcc/c-family/ChangeLog.omp @@ -1,3 +1,18 @@ +2025-01-27 Tobias Burnus <tbur...@baylibre.com> + + Backported from master: + 2024-11-22 Tobias Burnus <tbur...@baylibre.com> + + * c-common.h (enum c_omp_region_type): Add C_ORT_INTEROP + and C_ORT_OMP_INTEROP. + (c_omp_interop_t_p): New prototype. + * c-omp.cc (c_omp_interop_t_p): Check whether the type is + omp_interop_t. + (c_omp_directives): Uncomment 'interop'. + * c-pragma.cc (omp_pragmas): Add 'interop'. + * c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_INTEROP. + (enum pragma_omp_clause): Add init, use, and destroy clauses. + 2025-01-27 Paul-Antoine Arras <par...@baylibre.com> Backported from master: diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index 7d3b68f0111c..50ad1554ea47 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1280,11 +1280,13 @@ enum c_omp_region_type C_ORT_DECLARE_SIMD = 1 << 2, C_ORT_TARGET = 1 << 3, C_ORT_EXIT_DATA = 1 << 4, - C_ORT_UPDATE = 1 << 5, - C_ORT_DECLARE_MAPPER = 1 << 6, + C_ORT_INTEROP = 1 << 5, + C_ORT_UPDATE = 1 << 6, + C_ORT_DECLARE_MAPPER = 1 << 7, C_ORT_OMP_DECLARE_SIMD = C_ORT_OMP | C_ORT_DECLARE_SIMD, C_ORT_OMP_TARGET = C_ORT_OMP | C_ORT_TARGET, C_ORT_OMP_EXIT_DATA = C_ORT_OMP | C_ORT_EXIT_DATA, + C_ORT_OMP_INTEROP = C_ORT_OMP | C_ORT_INTEROP, C_ORT_OMP_UPDATE = C_ORT_OMP | C_ORT_UPDATE, C_ORT_OMP_DECLARE_MAPPER = C_ORT_OMP | C_ORT_DECLARE_MAPPER, C_ORT_ACC_TARGET = C_ORT_ACC | C_ORT_TARGET @@ -1299,6 +1301,7 @@ extern void c_finish_omp_barrier (location_t); extern tree c_finish_omp_atomic (location_t, enum tree_code, enum tree_code, tree, tree, tree, tree, tree, tree, bool, enum omp_memory_order, bool, bool = false); +extern bool c_omp_interop_t_p (tree); extern bool c_omp_depend_t_p (tree); extern void c_finish_omp_depobj (location_t, tree, enum omp_clause_depend_kind, tree); diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index 20c2609836fb..4fadcfe43b25 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -664,6 +664,24 @@ c_finish_omp_atomic (location_t loc, enum tree_code code, } +/* Return true if TYPE is the implementation's omp_interop_t. */ + +bool +c_omp_interop_t_p (tree type) +{ + type = TYPE_MAIN_VARIANT (type); + return (TREE_CODE (type) == ENUMERAL_TYPE + && TYPE_NAME (type) + && ((TREE_CODE (TYPE_NAME (type)) == TYPE_DECL + ? DECL_NAME (TYPE_NAME (type)) : TYPE_NAME (type)) + == get_identifier ("omp_interop_t")) + && TYPE_FILE_SCOPE_P (type) + && COMPLETE_TYPE_P (type) + && TREE_CODE (TYPE_SIZE (type)) == INTEGER_CST + && !compare_tree_int (TYPE_SIZE (type), + tree_to_uhwi (TYPE_SIZE (ptr_type_node)))); +} + /* Return true if TYPE is the implementation's omp_depend_t. */ bool @@ -5061,8 +5079,8 @@ const struct c_omp_directive c_omp_directives[] = { C_OMP_DIR_CONSTRUCT, true }, /* { "groupprivate", nullptr, nullptr, PRAGMA_OMP_GROUPPRIVATE, C_OMP_DIR_DECLARATIVE, false }, */ - /* { "interop", nullptr, nullptr, PRAGMA_OMP_INTEROP, - C_OMP_DIR_STANDALONE, false }, */ + { "interop", nullptr, nullptr, PRAGMA_OMP_INTEROP, + C_OMP_DIR_STANDALONE, false }, { "loop", nullptr, nullptr, PRAGMA_OMP_LOOP, C_OMP_DIR_CONSTRUCT, true }, { "masked", nullptr, nullptr, PRAGMA_OMP_MASKED, diff --git a/gcc/c-family/c-pragma.cc b/gcc/c-family/c-pragma.cc index 1f4c8aa091bd..5e2632849c9f 100644 --- a/gcc/c-family/c-pragma.cc +++ b/gcc/c-family/c-pragma.cc @@ -1530,6 +1530,7 @@ static const struct omp_pragma_def omp_pragmas[] = { { "error", PRAGMA_OMP_ERROR }, { "end", PRAGMA_OMP_END }, { "flush", PRAGMA_OMP_FLUSH }, + { "interop", PRAGMA_OMP_INTEROP }, { "metadirective", PRAGMA_OMP_METADIRECTIVE }, { "nothing", PRAGMA_OMP_NOTHING }, { "requires", PRAGMA_OMP_REQUIRES }, diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index 8b14fcd8be29..bea0d9121cb2 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -61,6 +61,7 @@ enum pragma_kind { PRAGMA_OMP_END, PRAGMA_OMP_FLUSH, PRAGMA_OMP_FOR, + PRAGMA_OMP_INTEROP, PRAGMA_OMP_LOOP, PRAGMA_OMP_NOTHING, PRAGMA_OMP_MASKED, @@ -112,6 +113,7 @@ enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_DEFAULT, PRAGMA_OMP_CLAUSE_DEFAULTMAP, PRAGMA_OMP_CLAUSE_DEPEND, + PRAGMA_OMP_CLAUSE_DESTROY, PRAGMA_OMP_CLAUSE_DETACH, PRAGMA_OMP_CLAUSE_DEVICE, PRAGMA_OMP_CLAUSE_DEVICE_TYPE, @@ -131,6 +133,7 @@ enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_IN_REDUCTION, PRAGMA_OMP_CLAUSE_INBRANCH, PRAGMA_OMP_CLAUSE_INDIRECT, + PRAGMA_OMP_CLAUSE_INIT, PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR, PRAGMA_OMP_CLAUSE_LASTPRIVATE, PRAGMA_OMP_CLAUSE_LINEAR, @@ -167,6 +170,7 @@ enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_TO, PRAGMA_OMP_CLAUSE_UNIFORM, PRAGMA_OMP_CLAUSE_UNTIED, + PRAGMA_OMP_CLAUSE_USE, PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR, PRAGMA_OMP_CLAUSE_USES_ALLOCATORS, diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp index 56656442267f..4323cd1b239d 100644 --- a/gcc/c/ChangeLog.omp +++ b/gcc/c/ChangeLog.omp @@ -1,3 +1,19 @@ +2025-01-27 Tobias Burnus <tbur...@baylibre.com> + + Backported from master: + 2024-11-22 Tobias Burnus <tbur...@baylibre.com> + + * c-parser.cc (INCLUDE_STRING): Define. + (c_parser_pragma): Handle 'interop' directive. + (c_parser_omp_clause_name): Handle init, use, and destroy clauses. + (c_parser_omp_all_clauses): Likewise; use C_ORT_OMP_INTEROP, if + 'use' is permitted, for c_finish_omp_clauses. + (c_parser_omp_clause_destroy, c_parser_omp_modifier_prefer_type, + c_parser_omp_clause_init, c_parser_omp_clause_use, + OMP_INTEROP_CLAUSE_MASK, c_parser_omp_interop): New. + * c-typeck.cc (c_finish_omp_clauses): Add missing OPT_Wopenmp to + a warning; handle new clauses. + 2025-01-27 Paul-Antoine Arras <par...@baylibre.com> Backported from master: diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 7c657997799a..df369a7090f3 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -37,6 +37,7 @@ along with GCC; see the file COPYING3. If not see #include "config.h" #define INCLUDE_MEMORY +#define INCLUDE_STRING #include "system.h" #include "coretypes.h" #include "target.h" @@ -1760,6 +1761,7 @@ 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_omp_interop (c_parser *); static void c_parser_oacc_routine (c_parser *, enum pragma_context); /* These Objective-C parser functions are only ever called when @@ -14755,6 +14757,15 @@ c_parser_pragma (c_parser *parser, enum pragma_context context, bool *if_p) c_parser_omp_flush (parser); return false; + case PRAGMA_OMP_INTEROP: + if (context != pragma_compound) + { + construct = "omp interop"; + goto in_compound; + } + c_parser_omp_interop (parser); + return false; + case PRAGMA_OMP_TASKWAIT: if (context != pragma_compound) { @@ -15083,6 +15094,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OACC_CLAUSE_DELETE; else if (!strcmp ("depend", p)) result = PRAGMA_OMP_CLAUSE_DEPEND; + else if (!strcmp ("destroy", p)) + result = PRAGMA_OMP_CLAUSE_DESTROY; else if (!strcmp ("detach", p)) result = PRAGMA_OACC_CLAUSE_DETACH; else if (!strcmp ("device", p)) @@ -15141,6 +15154,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OACC_CLAUSE_INDEPENDENT; else if (!strcmp ("indirect", p)) result = PRAGMA_OMP_CLAUSE_INDIRECT; + else if (!strcmp ("init", p)) + result = PRAGMA_OMP_CLAUSE_INIT; else if (!strcmp ("is_device_ptr", p)) result = PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR; break; @@ -15261,6 +15276,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_UNIFORM; else if (!strcmp ("untied", p)) result = PRAGMA_OMP_CLAUSE_UNTIED; + else if (!strcmp ("use", p)) + result = PRAGMA_OMP_CLAUSE_USE; else if (!strcmp ("use_device", p)) result = PRAGMA_OACC_CLAUSE_USE_DEVICE; else if (!strcmp ("use_device_addr", p)) @@ -20130,6 +20147,385 @@ c_parser_omp_clause_detach (c_parser *parser, tree list) return u; } +/* OpenMP 5.0: + destroy ( variable-list ) */ + +static tree +c_parser_omp_clause_destroy (c_parser *parser, tree list) +{ + return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_DESTROY, list); +} + +/* OpenMP 5.1: + prefer_type ( const-int-expr-or-string-literal-list ) + + OpenMP 6.0: + prefer_type ( { preference-selector-list }, { ... } ) + + with preference-selector being: + fr ( identifier-or-string-literal-list ) + attr ( string-list ) + + Data format: + For the foreign runtime identifiers, string values are converted to + their integer value; unknown string or integer values are set to + GOMP_INTEROP_IFR_KNOWN. + + Each item (a) GOMP_INTEROP_IFR_SEPARATOR + (b) for any 'fr', its integer value. + Note: Spec only permits 1 'fr' entry (6.0; changed after TR13) + (c) GOMP_INTEROP_IFR_SEPARATOR + (d) list of \0-terminated non-empty strings for 'attr' + (e) '\0' + Tailing '\0'. */ + +static tree +c_parser_omp_modifier_prefer_type (c_parser *parser) +{ + if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + return error_mark_node; + + std::string str; + + /* Old Format: const-int-expr-or-string-literal-list */ + if (!c_parser_next_token_is (parser, CPP_OPEN_BRACE)) + while (true) + { + str += (char) GOMP_INTEROP_IFR_SEPARATOR; + if (c_parser_next_token_is (parser, CPP_STRING)) + { + c_expr cval = c_parser_string_literal (parser, false, false); + if (cval.value == error_mark_node) + return error_mark_node; + if ((size_t) TREE_STRING_LENGTH (cval.value) + != strlen (TREE_STRING_POINTER (cval.value)) + 1) + { + error_at (cval.get_location (), "string literal must " + "not contain %<\\0%>"); + parser->error = true; + return error_mark_node; + } + + char c = omp_get_fr_id_from_name (TREE_STRING_POINTER (cval.value)); + if (c == GOMP_INTEROP_IFR_UNKNOWN) + warning_at (cval.get_location (), OPT_Wopenmp, + "unknown foreign runtime identifier %qs", + TREE_STRING_POINTER (cval.value)); + str += c; + } + else + { + c_expr cval = c_parser_expr_no_commas (parser, NULL); + tree value = c_fully_fold (cval.value, false, NULL); + if (INTEGRAL_TYPE_P (TREE_TYPE (value)) + && TREE_CODE (value) != INTEGER_CST) + value = convert_lvalue_to_rvalue (cval.get_start (), cval, + false, true).value; + if (TREE_CODE (value) != INTEGER_CST + || !tree_fits_shwi_p (value)) + { + c_parser_error (parser, "expected string literal or constant " + "integer expression"); + return error_mark_node; + } + HOST_WIDE_INT n = tree_to_shwi (value); + if (n < 1 || n > GOMP_INTEROP_IFR_LAST) + { + warning_at (cval.get_location (), OPT_Wopenmp, + "unknown foreign runtime identifier %qwd", n); + n = GOMP_INTEROP_IFR_UNKNOWN; + } + str += (char) n; + } + str += (char) GOMP_INTEROP_IFR_SEPARATOR; + str += '\0'; + if (c_parser_next_token_is (parser, CPP_COMMA)) + { + c_parser_consume_token (parser); + continue; + } + if (!c_parser_require (parser, CPP_CLOSE_PAREN, + "expected %<,%> or %<)%>")) + return error_mark_node; + str += '\0'; + tree res = build_string (str.length (), str.data ()); + TREE_TYPE (res) = build_array_type_nelts (unsigned_char_type_node, + str.length ()); + return res; + } + + /* New format. */ + std::string str2; + while (true) + { + if (!c_parser_require (parser, CPP_OPEN_BRACE, "expected %<{%>")) + return error_mark_node; + str += (char) GOMP_INTEROP_IFR_SEPARATOR; + str2.clear (); + bool has_fr = false; + while (true) + { + c_token *tok = c_parser_peek_token (parser); + if (tok->type != CPP_NAME + || (strcmp("fr", IDENTIFIER_POINTER (tok->value)) != 0 + && strcmp("attr", IDENTIFIER_POINTER (tok->value)) != 0)) + { + c_parser_error (parser, "expected %<fr%> or %<attr%> preference " + "selector"); + return error_mark_node; + } + c_parser_consume_token (parser); + bool is_fr = IDENTIFIER_POINTER (tok->value)[0] == 'f'; + if (is_fr && has_fr) + { + c_parser_error (parser, "duplicated %<fr%> preference selector"); + return error_mark_node; + } + if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + return error_mark_node; + while (true) + { + if (c_parser_next_token_is (parser, CPP_STRING)) + { + c_expr cval = c_parser_string_literal (parser, false, false); + tree value = cval.value; + if (value == error_mark_node) + return error_mark_node; + if ((size_t) TREE_STRING_LENGTH (value) + != strlen (TREE_STRING_POINTER (value)) + 1) + { + error_at (cval.get_location (), "string literal must " + "not contain %<\\0%>"); + parser->error = true; + return error_mark_node; + } + if (!is_fr) + { + if (!startswith (TREE_STRING_POINTER (value), "ompx_")) + { + error_at (cval.get_location (), + "%<attr%> string literal must start with " + "%<ompx_%>"); + parser->error = true; + return error_mark_node; + } + if (strchr (TREE_STRING_POINTER (value), ',')) + { + error_at (cval.get_location (), + "%<attr%> string literal must not contain " + "a comma"); + parser->error = true; + return error_mark_node; + } + str2 += TREE_STRING_POINTER (value); + str2 += '\0'; + } + else + { + if (*TREE_STRING_POINTER (value) == '\0') + { + c_parser_error (parser, "non-empty string literal expected"); + return error_mark_node; + } + char c = omp_get_fr_id_from_name (TREE_STRING_POINTER (value)); + if (c == GOMP_INTEROP_IFR_UNKNOWN) + warning_at (cval.get_location (), OPT_Wopenmp, + "unknown foreign runtime identifier %qs", + TREE_STRING_POINTER (value)); + str += c; + has_fr = true; + } + } + else if (!is_fr) + { + c_parser_error (parser, "expected string literal"); + return error_mark_node; + } + else + { + c_expr cval = c_parser_expr_no_commas (parser, NULL); + tree value = c_fully_fold (cval.value, false, NULL); + if (INTEGRAL_TYPE_P (TREE_TYPE (value)) + && TREE_CODE (value) != INTEGER_CST) + value = convert_lvalue_to_rvalue (cval.get_start (), cval, + false, true).value; + + if (TREE_CODE (value) != INTEGER_CST + || !tree_fits_shwi_p (value)) + { + c_parser_error (parser, "expected string literal or " + "constant integer expression"); + return error_mark_node; + } + HOST_WIDE_INT n = tree_to_shwi (value); + if (n < 1 || n > GOMP_INTEROP_IFR_LAST) + { + warning_at (cval.get_location (), OPT_Wopenmp, + "unknown foreign runtime identifier %qwd", n); + n = GOMP_INTEROP_IFR_UNKNOWN; + } + str += (char) n; + has_fr = true; + } + if (!is_fr + && c_parser_next_token_is (parser, CPP_COMMA)) + { + c_parser_consume_token (parser); + continue; + } + if (!c_parser_require (parser, CPP_CLOSE_PAREN, + is_fr ? G_("expected %<)%>") + : G_("expected %<)%> or %<,%>"))) + return error_mark_node; + break; + } + if (c_parser_next_token_is (parser, CPP_COMMA)) + { + c_parser_consume_token (parser); + continue; + } + if (c_parser_next_token_is (parser, CPP_CLOSE_BRACE)) + break; + c_parser_error (parser, "expected %<,%> or %<}%>"); + return error_mark_node; + } + str += (char) GOMP_INTEROP_IFR_SEPARATOR; + str += str2; + str += '\0'; + c_parser_consume_token (parser); + if (c_parser_next_token_is (parser, CPP_CLOSE_PAREN)) + break; + if (!c_parser_require (parser, CPP_COMMA, "expected %<)%> or %<,%>")) + return error_mark_node; + } + c_parser_consume_token (parser); + str += '\0'; + tree res = build_string (str.length (), str.data ()); + TREE_TYPE (res) = build_array_type_nelts (unsigned_char_type_node, + str.length ()); + return res; +} + +/* OpenMP 5.1: + init ( [init-modifier-list : ] variable-list ) + + Modifiers: + target + targetsync + prefer_type (preference-specification) */ + +static tree +c_parser_omp_clause_init (c_parser *parser, tree list) +{ + location_t loc = c_parser_peek_token (parser)->location; + + matching_parens parens; + if (!parens.require_open (parser)) + return list; + + unsigned pos = 0, raw_pos = 1; + while (c_parser_peek_nth_token_raw (parser, raw_pos)->type == CPP_NAME) + { + pos++; raw_pos++; + if (c_parser_peek_nth_token_raw (parser, raw_pos)->type == CPP_OPEN_PAREN) + { + raw_pos++; + c_parser_check_balanced_raw_token_sequence (parser, &raw_pos); + if (c_parser_peek_nth_token_raw (parser, raw_pos)->type != CPP_CLOSE_PAREN) + { + pos = 0; + break; + } + raw_pos++; + } + if (c_parser_peek_nth_token_raw (parser, raw_pos)->type == CPP_COLON) + break; + if (c_parser_peek_nth_token_raw (parser, raw_pos)->type != CPP_COMMA) + { + pos = 0; + break; + } + pos++; + raw_pos++; + } + + bool target = false; + bool targetsync = false; + tree prefer_type_tree = NULL_TREE; + + for (unsigned pos2 = 0; pos2 < pos; ++pos2) + { + c_token *tok = c_parser_peek_token (parser); + if (tok->type == CPP_COMMA) + { + c_parser_consume_token (parser); + continue; + } + + const char *p = IDENTIFIER_POINTER (tok->value); + if (strcmp ("targetsync", p) == 0) + { + if (targetsync) + error_at (tok->location, "duplicate %<targetsync%> modifier"); + targetsync = true; + c_parser_consume_token (parser); + } + else if (strcmp ("target", p) == 0) + { + if (target) + error_at (tok->location, "duplicate %<target%> modifier"); + target = true; + c_parser_consume_token (parser); + } + else if (strcmp ("prefer_type", p) == 0) + { + if (prefer_type_tree != NULL_TREE) + error_at (tok->location, "duplicate %<prefer_type%> modifier"); + c_parser_consume_token (parser); + prefer_type_tree = c_parser_omp_modifier_prefer_type (parser); + if (prefer_type_tree == error_mark_node) + return list; + } + else + { + c_parser_error (parser, "%<init%> clause with modifier other than " + "%<prefer_type%>, %<target%> or " + "%<targetsync%>"); + parens.skip_until_found_close (parser); + return list; + } + } + if (pos) + { + c_token *tok = c_parser_peek_token (parser); + gcc_checking_assert (tok->type == CPP_COLON); + c_parser_consume_token (parser); + } + + tree nl = c_parser_omp_variable_list (parser, loc, OMP_CLAUSE_INIT, list); + parens.skip_until_found_close (parser); + + for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) + { + if (target) + OMP_CLAUSE_INIT_TARGET (c) = 1; + if (targetsync) + OMP_CLAUSE_INIT_TARGETSYNC (c) = 1; + if (prefer_type_tree) + OMP_CLAUSE_INIT_PREFER_TYPE (c) = prefer_type_tree; + } + return nl; +} + +/* OpenMP 5.0: + use ( variable-list ) */ + +static tree +c_parser_omp_clause_use (c_parser *parser, tree list) +{ + return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_USE, list); +} + /* Parse all OpenACC clauses. The set clauses allowed by the directive is a bitmask in MASK. Return the list of clauses found. */ @@ -20629,6 +21025,18 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_omp_clause_doacross (parser, clauses); c_name = "doacross"; break; + case PRAGMA_OMP_CLAUSE_DESTROY: + clauses = c_parser_omp_clause_destroy (parser, clauses); + c_name = "destroy"; + break; + case PRAGMA_OMP_CLAUSE_INIT: + clauses = c_parser_omp_clause_init (parser, clauses); + c_name = "init"; + break; + case PRAGMA_OMP_CLAUSE_USE: + clauses = c_parser_omp_clause_use (parser, clauses); + c_name = "use"; + break; case PRAGMA_OMP_CLAUSE_MAP: clauses = c_parser_omp_clause_map (parser, clauses, GOMP_MAP_TOFROM); c_name = "map"; @@ -20735,6 +21143,8 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, { if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_UNIFORM)) != 0) return c_finish_omp_clauses (clauses, C_ORT_OMP_DECLARE_SIMD); + if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE)) != 0) + return c_finish_omp_clauses (clauses, C_ORT_OMP_INTEROP); return c_finish_omp_clauses (clauses, C_ORT_OMP); } @@ -24065,6 +24475,32 @@ c_parser_omp_masked (location_t loc, c_parser *parser, clauses); } +/* OpenMP 5.1: + # pragma omp interop clauses[opt] new-line */ + +#define OMP_INTEROP_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DESTROY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INIT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE)) + +static void +c_parser_omp_interop (c_parser *parser) +{ + location_t loc = c_parser_peek_token (parser)->location; + c_parser_consume_pragma (parser); + tree clauses = c_parser_omp_all_clauses (parser, + OMP_INTEROP_CLAUSE_MASK, + "#pragma omp interop"); + tree stmt = make_node (OMP_INTEROP); + TREE_TYPE (stmt) = void_type_node; + OMP_INTEROP_CLAUSES (stmt) = clauses; + SET_EXPR_LOCATION (stmt, loc); + add_stmt (stmt); +} + /* OpenMP 2.5: # pragma omp ordered new-line structured-block diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 91d016aaed5e..018ba344d892 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -14968,6 +14968,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) tree *full_seen = NULL; bool partial_seen = false; bool openacc = (ort & C_ORT_ACC) != 0; + bool init_seen = false; + bool init_use_destroy_seen = false; + tree init_no_targetsync_clause = NULL_TREE; + tree depend_clause = NULL_TREE; bitmap_obstack_initialize (NULL); bitmap_initialize (&generic_head, &bitmap_default_obstack); @@ -15594,7 +15598,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } else if (bitmap_bit_p (&aligned_head, DECL_UID (t))) { - warning_at (OMP_CLAUSE_LOCATION (c), 0, + warning_at (OMP_CLAUSE_LOCATION (c), OPT_Wopenmp, "%qE appears more than once in %<allocate%> clauses", t); remove = true; @@ -15746,6 +15750,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) pc = &OMP_CLAUSE_CHAIN (c); continue; case OMP_CLAUSE_DEPEND: + depend_clause = c; + /* FALLTHRU */ case OMP_CLAUSE_AFFINITY: t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST @@ -16497,6 +16503,32 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } break; + case OMP_CLAUSE_INIT: + init_seen = true; + if (!OMP_CLAUSE_INIT_TARGETSYNC (c)) + init_no_targetsync_clause = c; + /* FALLTHRU */ + case OMP_CLAUSE_DESTROY: + case OMP_CLAUSE_USE: + init_use_destroy_seen = true; + t = OMP_CLAUSE_DECL (c); + if (bitmap_bit_p (&generic_head, DECL_UID (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD appears more than once in action clauses", t); + remove = true; + } + else if (/* ort == C_ORT_OMP_INTEROP [uncomment for depobj init] */ + !c_omp_interop_t_p (TREE_TYPE (OMP_CLAUSE_DECL (c)))) + error_at (OMP_CLAUSE_LOCATION (c), + "%qD must be of %<omp_interop_t%>", OMP_CLAUSE_DECL (c)); + else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE + && TREE_READONLY (OMP_CLAUSE_DECL (c))) + error_at (OMP_CLAUSE_LOCATION (c), + "%qD shall not be const", OMP_CLAUSE_DECL (c)); + bitmap_set_bit (&generic_head, DECL_UID (t)); + pc = &OMP_CLAUSE_CHAIN (c); + break; default: gcc_unreachable (); } @@ -16768,6 +16800,19 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } } + if (ort == C_ORT_OMP_INTEROP + && depend_clause + && (!init_use_destroy_seen + || (init_seen && init_no_targetsync_clause))) + { + error_at (OMP_CLAUSE_LOCATION (depend_clause), + "%<depend%> clause requires action clauses with " + "%<targetsync%> interop-type"); + if (init_no_targetsync_clause) + inform (OMP_CLAUSE_LOCATION (init_no_targetsync_clause), + "%<init%> clause lacks the %<targetsync%> modifier"); + } + bitmap_obstack_release (NULL); return clauses; } diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp index 58896c23731d..41d2eb6f157a 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,3 +1,21 @@ +2025-01-27 Tobias Burnus <tbur...@baylibre.com> + + Backported from master: + 2024-11-22 Tobias Burnus <tbur...@baylibre.com> + + * parser.cc (INCLUDE_STRING): Define. + (cp_parser_omp_clause_name): Handle init, use, and destroy clauses. + (cp_parser_omp_all_clauses): Likewise; use C_ORT_OMP_INTEROP, if + 'use' is permitted, for c_finish_omp_clauses. + (cp_parser_omp_modifier_prefer_type, cp_parser_omp_clause_init, + OMP_INTEROP_CLAUSE_MASK, cp_parser_omp_interop): New. + (cp_parser_pragma): Handle 'interop' directive. + * pt.cc (tsubst_omp_clauses): Handle init, use, and destroy clauses. + (tsubst_stmt): Handle OMP_INTEROP. + * semantics.cc (cp_omp_init_prefer_type_update): New. + (finish_omp_clauses): Handle init, use, and destroy clauses + and add clause check for 'depend' on 'interop'. + 2025-01-27 Paul-Antoine Arras <par...@baylibre.com> Backported from master: diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 94a65bee7d96..e23aaf1c5a12 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -21,6 +21,7 @@ along with GCC; see the file COPYING3. If not see #include "config.h" #include "omp-selectors.h" #define INCLUDE_MEMORY +#define INCLUDE_STRING #include "system.h" #include "coretypes.h" #include "cp-tree.h" @@ -38064,6 +38065,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_DEFAULTMAP; else if (!strcmp ("depend", p)) result = PRAGMA_OMP_CLAUSE_DEPEND; + else if (!strcmp ("destroy", p)) + result = PRAGMA_OMP_CLAUSE_DESTROY; else if (!strcmp ("detach", p)) result = PRAGMA_OACC_CLAUSE_DETACH; else if (!strcmp ("device", p)) @@ -38122,6 +38125,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OACC_CLAUSE_INDEPENDENT; else if (!strcmp ("indirect", p)) result = PRAGMA_OMP_CLAUSE_INDIRECT; + else if (!strcmp ("init", p)) + result = PRAGMA_OMP_CLAUSE_INIT; else if (!strcmp ("is_device_ptr", p)) result = PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR; break; @@ -38238,6 +38243,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_UNIFORM; else if (!strcmp ("untied", p)) result = PRAGMA_OMP_CLAUSE_UNTIED; + else if (!strcmp ("use", p)) + result = PRAGMA_OMP_CLAUSE_USE; else if (!strcmp ("use_device", p)) result = PRAGMA_OACC_CLAUSE_USE_DEVICE; else if (!strcmp ("use_device_addr", p)) @@ -42655,6 +42662,418 @@ cp_parser_omp_clause_device_type (cp_parser *parser, tree list, return list; } +/* OpenMP 5.1: + prefer_type ( const-int-expr-or-string-literal-list ) + + OpenMP 6.0: + prefer_type ( { preference-selector-list }, { ... } ) + + with preference-selector being: + fr ( identifier-or-string-literal-list ) + attr ( string-list ) + + Data format: + For the foreign runtime identifiers, string values are converted to + their integer value; unknown string or integer values are set to + GOMP_INTEROP_IFR_KNOWN. + + Each item (a) GOMP_INTEROP_IFR_SEPARATOR + (b) for any 'fr', its integer value. + Note: Spec only permits 1 'fr' entry (6.0; changed after TR13) + (c) GOMP_INTEROP_IFR_SEPARATOR + (d) list of \0-terminated non-empty strings for 'attr' + (e) '\0' + Tailing '\0'. + + When processing a template: + attr and fr strings are processed normally. + for integer expressions, set fr to UNKNOWN and keep a separate list + of those expressions - and store it as + tree_cons (bytestring, fr_tree list). */ + +static tree +cp_parser_omp_modifier_prefer_type (cp_parser *parser) +{ + if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + return error_mark_node; + + unsigned fr_cnt = 0; + auto_vec<tree> fr_list; + std::string str; + + /* Old Format: const-int-expr-or-string-literal-list */ + if (cp_lexer_peek_token (parser->lexer)->type != CPP_OPEN_BRACE) + while (true) + { + str += (char) GOMP_INTEROP_IFR_SEPARATOR; + if (cp_lexer_peek_token (parser->lexer)->type == CPP_STRING) + { + cp_expr cval = cp_parser_unevaluated_string_literal (parser); + if (*cval == error_mark_node) + return error_mark_node; + if ((size_t) TREE_STRING_LENGTH (*cval) + != strlen (TREE_STRING_POINTER (*cval)) + 1) + { + error_at (cval.get_location (), "string literal must " + "not contain %<\\0%>"); + return error_mark_node; + } + char c = omp_get_fr_id_from_name (TREE_STRING_POINTER (*cval)); + if (c == GOMP_INTEROP_IFR_UNKNOWN) + warning_at (cval.get_location (), OPT_Wopenmp, + "unknown foreign runtime identifier %qs", + TREE_STRING_POINTER (*cval)); + str += c; + } + else + { + location_t loc = cp_lexer_peek_token (parser->lexer)->location; + tree value = cp_parser_assignment_expression (parser); + value = cp_fully_fold (value); + if (!processing_template_decl) + { + if (TREE_CODE (value) != INTEGER_CST + || !tree_fits_shwi_p (value)) + { + cp_parser_error (parser, + "expected string literal or constant " + "integer expression"); + return error_mark_node; + } + HOST_WIDE_INT n = tree_to_shwi (value); + if (n < 1 || n > GOMP_INTEROP_IFR_LAST) + { + warning_at (EXPR_LOC_OR_LOC (value, loc), OPT_Wopenmp, + "unknown foreign runtime identifier %qwd", n); + n = GOMP_INTEROP_IFR_UNKNOWN; + } + str += (char) n; + } + else + { + str += (char) GOMP_INTEROP_IFR_UNKNOWN; + for (unsigned n = fr_list.length (); n < fr_cnt; n++) + fr_list.safe_push (NULL_TREE); + if (EXPR_LOCATION (value) == UNKNOWN_LOCATION) + value = build1_loc (loc, NOP_EXPR, TREE_TYPE (value), value); + fr_list.safe_push (value); + } + } + fr_cnt++; + str += (char) GOMP_INTEROP_IFR_SEPARATOR; + str += '\0'; + if (cp_lexer_peek_token (parser->lexer)->type == CPP_COMMA) + { + cp_lexer_consume_token (parser->lexer); + continue; + } + if (cp_lexer_peek_token (parser->lexer)->type != CPP_CLOSE_PAREN) + { + cp_parser_error (parser, "expected %<,%> or %<)%>"); + return error_mark_node; + } + cp_lexer_consume_token (parser->lexer); + str += '\0'; + tree res = build_string (str.length (), str.data ()); + TREE_TYPE (res) = build_array_type_nelts (unsigned_char_type_node, + str.length ()); + if (!fr_list.is_empty ()) + { + tree t = make_tree_vec (fr_list.length ()); + for (unsigned i = 0; i < fr_list.length (); i++) + TREE_VEC_ELT (t, i) = fr_list[i]; + res = tree_cons (res, t, NULL_TREE); + } + return res; + } + + /* New format. */ + std::string str2; + while (true) + { + if (!cp_parser_require (parser, CPP_OPEN_BRACE, RT_OPEN_BRACE)) + return error_mark_node; + str += (char) GOMP_INTEROP_IFR_SEPARATOR; + str2.clear (); + bool has_fr = false; + while (true) + { + cp_token *tok = cp_lexer_peek_token (parser->lexer); + if (tok->type != CPP_NAME + || (strcmp("fr", IDENTIFIER_POINTER (tok->u.value)) != 0 + && strcmp("attr", IDENTIFIER_POINTER (tok->u.value)) != 0)) + { + cp_parser_error (parser, "expected %<fr%> or %<attr%> preference " + "selector"); + return error_mark_node; + } + cp_lexer_consume_token (parser->lexer); + bool is_fr = IDENTIFIER_POINTER (tok->u.value)[0] == 'f'; + if (is_fr && has_fr) + { + cp_parser_error (parser, "duplicated %<fr%> preference selector"); + return error_mark_node; + } + if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + return error_mark_node; + while (true) + { + if (cp_lexer_peek_token (parser->lexer)->type == CPP_STRING) + { + cp_expr cval = cp_parser_unevaluated_string_literal (parser); + if (*cval == error_mark_node) + return error_mark_node; + if ((size_t) TREE_STRING_LENGTH (*cval) + != strlen (TREE_STRING_POINTER (*cval)) + 1) + { + error_at (cval.get_location (), "string literal must " + "not contain %<\\0%>"); + return error_mark_node; + } + if (!is_fr) + { + if (!startswith (TREE_STRING_POINTER (*cval), "ompx_")) + { + error_at (cval.get_location (), + "%<attr%> string literal must start with " + "%<ompx_%>"); + return error_mark_node; + } + if (strchr (TREE_STRING_POINTER (*cval), ',')) + { + error_at (cval.get_location (), + "%<attr%> string literal must not contain " + "a comma"); + return error_mark_node; + } + str2 += TREE_STRING_POINTER (*cval); + str2 += '\0'; + } + else + { + if (*TREE_STRING_POINTER (*cval) == '\0') + { + cp_parser_error (parser, + "non-empty string literal expected"); + return error_mark_node; + } + char c + = omp_get_fr_id_from_name (TREE_STRING_POINTER (*cval)); + if (c == GOMP_INTEROP_IFR_UNKNOWN) + warning_at (cval.get_location (), OPT_Wopenmp, + "unknown foreign runtime identifier %qs", + TREE_STRING_POINTER (*cval)); + str += c; + has_fr = true; + } + } + else if (!is_fr) + { + cp_parser_error (parser, "expected string literal"); + return error_mark_node; + } + else + { + location_t loc = cp_lexer_peek_token (parser->lexer)->location; + tree value = cp_parser_assignment_expression (parser); + value = cp_fully_fold (value); + if (!processing_template_decl) + { + if (TREE_CODE (value) != INTEGER_CST + || !tree_fits_shwi_p (value)) + { + cp_parser_error (parser, + "expected string literal or " + "constant integer expression"); + return error_mark_node; + } + HOST_WIDE_INT n = tree_to_shwi (value); + if (n < 1 || n > GOMP_INTEROP_IFR_LAST) + { + warning_at (EXPR_LOC_OR_LOC (value, loc), OPT_Wopenmp, + "unknown foreign runtime identifier %qwd", + n); + n = GOMP_INTEROP_IFR_UNKNOWN; + } + str += (char) n; + } + else + { + str += (char) GOMP_INTEROP_IFR_UNKNOWN; + for (unsigned n = fr_list.length (); n < fr_cnt; n++) + fr_list.safe_push (NULL_TREE); + if (EXPR_LOCATION (value) == UNKNOWN_LOCATION) + value = build1_loc (loc, NOP_EXPR, TREE_TYPE (value), + value); + fr_list.safe_push (value); + } + has_fr = true; + } + if (!is_fr + && cp_lexer_peek_token (parser->lexer)->type == CPP_COMMA) + { + cp_lexer_consume_token (parser->lexer); + continue; + } + if (cp_lexer_peek_token (parser->lexer)->type != CPP_CLOSE_PAREN) + { + cp_parser_error (parser, + is_fr ? G_("expected %<)%>") + : G_("expected %<)%> or %<,%>")); + return error_mark_node; + } + cp_lexer_consume_token (parser->lexer); + break; + } + if (cp_lexer_peek_token (parser->lexer)->type == CPP_COMMA) + { + cp_lexer_consume_token (parser->lexer); + continue; + } + if (cp_lexer_peek_token (parser->lexer)->type != CPP_CLOSE_BRACE) + { + cp_parser_error (parser, "expected %<,%> or %<}%>"); + return error_mark_node; + } + cp_lexer_consume_token (parser->lexer); + break; + } + fr_cnt++; + str += (char) GOMP_INTEROP_IFR_SEPARATOR; + str += str2; + str += '\0'; + if (cp_lexer_peek_token (parser->lexer)->type == CPP_CLOSE_PAREN) + break; + if (cp_lexer_peek_token (parser->lexer)->type != CPP_COMMA) + { + cp_parser_error (parser, "expected %<)%> or %<,%>"); + return error_mark_node; + } + cp_lexer_consume_token (parser->lexer); + } + cp_lexer_consume_token (parser->lexer); + str += '\0'; + tree res = build_string (str.length (), str.data ()); + TREE_TYPE (res) = build_array_type_nelts (unsigned_char_type_node, + str.length ()); + if (!fr_list.is_empty ()) + { + tree t = make_tree_vec (fr_list.length ()); + for (unsigned i = 0; i < fr_list.length (); i++) + TREE_VEC_ELT (t, i) = fr_list[i]; + res = tree_cons (res, t, NULL_TREE); + } + return res; +} + +/* OpenMP 5.1: + init ( [init-modifier-list : ] variable-list ) + + Modifiers: + target + targetsync + prefer_type (preference-specification) */ + +static tree +cp_parser_omp_clause_init (cp_parser *parser, tree list) +{ + if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + return list; + + unsigned pos = 0, raw_pos = 1; + while (cp_lexer_peek_nth_token (parser->lexer, raw_pos)->type == CPP_NAME) + { + pos++; raw_pos++; + if (cp_lexer_peek_nth_token (parser->lexer, raw_pos)->type + == CPP_OPEN_PAREN) + { + unsigned n = cp_parser_skip_balanced_tokens (parser, raw_pos); + if (n == raw_pos) + { + pos = 0; + break; + } + raw_pos = n; + } + if (cp_lexer_peek_nth_token (parser->lexer, raw_pos)->type == CPP_COLON) + break; + if (cp_lexer_peek_nth_token (parser->lexer, raw_pos)->type != CPP_COMMA) + { + pos = 0; + break; + } + pos++; + raw_pos++; + } + + bool target = false; + bool targetsync = false; + tree prefer_type_tree = NULL_TREE; + + for (unsigned pos2 = 0; pos2 < pos; ++pos2) + { + cp_token *tok = cp_lexer_peek_token (parser->lexer); + if (tok->type == CPP_COMMA) + { + cp_lexer_consume_token (parser->lexer); + continue; + } + const char *p = IDENTIFIER_POINTER (tok->u.value); + if (strcmp ("targetsync", p) == 0) + { + if (targetsync) + error_at (tok->location, "duplicate %<targetsync%> modifier"); + targetsync = true; + cp_lexer_consume_token (parser->lexer); + } + else if (strcmp ("target", p) == 0) + { + if (target) + error_at (tok->location, "duplicate %<target%> modifier"); + target = true; + cp_lexer_consume_token (parser->lexer); + } + else if (strcmp ("prefer_type", p) == 0) + { + if (prefer_type_tree != NULL_TREE) + error_at (tok->location, "duplicate %<prefer_type%> modifier"); + cp_lexer_consume_token (parser->lexer); + prefer_type_tree = cp_parser_omp_modifier_prefer_type (parser); + if (prefer_type_tree == error_mark_node) + return error_mark_node; + } + else + { + cp_parser_error (parser, "%<init%> clause with modifier other than " + "%<prefer_type%>, %<target%> or " + "%<targetsync%>"); + cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + } + if (pos) + { + gcc_checking_assert (cp_lexer_peek_token (parser->lexer)->type + == CPP_COLON); + cp_lexer_consume_token (parser->lexer); + } + + tree nl = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_INIT, list, + NULL); + for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) + { + if (target) + OMP_CLAUSE_INIT_TARGET (c) = 1; + if (targetsync) + OMP_CLAUSE_INIT_TARGETSYNC (c) = 1; + if (prefer_type_tree) + OMP_CLAUSE_INIT_PREFER_TYPE (c) = prefer_type_tree; + } + return nl; +} + /* OpenACC: async [( int-expr )] */ @@ -43284,6 +43703,27 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, token->location); c_name = "doacross"; break; + case PRAGMA_OMP_CLAUSE_DESTROY: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_DESTROY, + clauses); + c_name = "destroy"; + break; + case PRAGMA_OMP_CLAUSE_INIT: + { + /* prefer_type parsing fails often such that many follow-up errors + are printed and recovery by cp_parser_skip_to_closing_parenthesis + will might skip to EOF, leading to an ICE elsewhere. */ + tree nc = cp_parser_omp_clause_init (parser, clauses); + if (nc == error_mark_node) + goto saw_error; + clauses = nc; + } + c_name = "init"; + break; + case PRAGMA_OMP_CLAUSE_USE: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE, clauses); + c_name = "use"; + break; case PRAGMA_OMP_CLAUSE_DETACH: clauses = cp_parser_omp_clause_detach (parser, clauses); c_name = "detach"; @@ -43385,8 +43825,9 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, { if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_UNIFORM)) != 0) return finish_omp_clauses (clauses, C_ORT_OMP_DECLARE_SIMD); - else - return finish_omp_clauses (clauses, C_ORT_OMP); + if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE)) != 0) + return finish_omp_clauses (clauses, C_ORT_OMP_INTEROP); + return finish_omp_clauses (clauses, C_ORT_OMP); } return clauses; } @@ -51779,6 +52220,30 @@ cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok, return false; } +/* OpenMP 5.1: + # pragma omp interop clauses[opt] new-line */ + +#define OMP_INTEROP_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DESTROY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INIT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE)) + +static void +cp_parser_omp_interop (cp_parser *parser, cp_token *pragma_tok) +{ + location_t loc = cp_lexer_peek_token (parser->lexer)->location; + tree clauses = cp_parser_omp_all_clauses (parser, OMP_INTEROP_CLAUSE_MASK, + "#pragma omp interop", pragma_tok); + tree stmt = make_node (OMP_INTEROP); + TREE_TYPE (stmt) = void_type_node; + OMP_INTEROP_CLAUSES (stmt) = clauses; + SET_EXPR_LOCATION (stmt, loc); + add_stmt (stmt); +} + /* OpenMP 5.0 #pragma omp requires clauses[optseq] new-line */ @@ -53055,6 +53520,22 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context, bool *if_p) } break; + case PRAGMA_OMP_INTEROP: + switch (context) + { + case pragma_compound: + cp_parser_omp_interop (parser, pragma_tok); + return false; + case pragma_stmt: + error_at (pragma_tok->location, "%<#pragma %s%> may only be " + "used in compound statements", "omp interop"); + ret = true; + break; + default: + goto bad_stmt; + } + break; + case PRAGMA_OMP_TASKWAIT: switch (context) { diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index d3e6a76fd81c..5c9aff0f1c1d 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -17836,6 +17836,36 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort, = tsubst_stmt (OMP_CLAUSE_LINEAR_STEP (oc), args, complain, in_decl); break; + case OMP_CLAUSE_INIT: + if (ort == C_ORT_OMP_INTEROP + && OMP_CLAUSE_INIT_PREFER_TYPE (nc) + && TREE_CODE (OMP_CLAUSE_INIT_PREFER_TYPE (nc)) == TREE_LIST + && (OMP_CLAUSE_CHAIN (nc) == NULL_TREE + || OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (nc) ) != OMP_CLAUSE_INIT + || (OMP_CLAUSE_INIT_PREFER_TYPE (nc) + != OMP_CLAUSE_INIT_PREFER_TYPE (OMP_CLAUSE_CHAIN (nc) )))) + { + tree pref_list = OMP_CLAUSE_INIT_PREFER_TYPE (nc); + tree fr_list = TREE_VALUE (pref_list); + int len = TREE_VEC_LENGTH (fr_list); + for (int i = 0; i < len; i++) + { + tree *fr_expr = &TREE_VEC_ELT (fr_list, i); + /* Preserve NOP_EXPR to have a location. */ + if (*fr_expr && TREE_CODE (*fr_expr) == NOP_EXPR) + TREE_OPERAND (*fr_expr, 0) + = tsubst_expr (TREE_OPERAND (*fr_expr, 0), args, complain, + in_decl); + else + *fr_expr = tsubst_expr (*fr_expr, args, complain, in_decl); + } + } + /* FALLTHRU */ + case OMP_CLAUSE_DESTROY: + case OMP_CLAUSE_USE: + OMP_CLAUSE_OPERAND (nc, 0) + = tsubst_stmt (OMP_CLAUSE_OPERAND (oc, 0), args, complain, in_decl); + break; case OMP_CLAUSE_NOWAIT: case OMP_CLAUSE_DEFAULT: case OMP_CLAUSE_UNTIED: @@ -19693,6 +19723,14 @@ tsubst_stmt (tree t, tree args, tsubst_flags_t complain, tree in_decl) } break; + case OMP_INTEROP: + tmp = tsubst_omp_clauses (OMP_INTEROP_CLAUSES (t), C_ORT_OMP_INTEROP, + args, complain, in_decl); + t = copy_node (t); + OMP_INTEROP_CLAUSES (t) = tmp; + add_stmt (t); + break; + case MUST_NOT_THROW_EXPR: { tree op0 = RECUR (TREE_OPERAND (t, 0)); diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index d0c415f4692f..600f90ba4c4b 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -7379,6 +7379,77 @@ cp_oacc_check_attachments (tree c) return false; } +/* Update OMP_CLAUSE_INIT_PREFER_TYPE in case template substitution + happened. */ + +static void +cp_omp_init_prefer_type_update (tree c) +{ + if (processing_template_decl + || OMP_CLAUSE_INIT_PREFER_TYPE (c) == NULL_TREE + || TREE_CODE (OMP_CLAUSE_INIT_PREFER_TYPE (c)) != TREE_LIST) + return; + + tree t = TREE_PURPOSE (OMP_CLAUSE_INIT_PREFER_TYPE (c)); + char *str = const_cast<char *> (TREE_STRING_POINTER (t)); + tree fr_list = TREE_VALUE (OMP_CLAUSE_INIT_PREFER_TYPE (c)); + int len = TREE_VEC_LENGTH (fr_list); + int cnt = 0; + + while (str[0] == (char) GOMP_INTEROP_IFR_SEPARATOR) + { + str++; + if (str[0] == (char) GOMP_INTEROP_IFR_UNKNOWN) + { + /* Assume a no or a single 'fr'. */ + gcc_checking_assert (str[1] == (char) GOMP_INTEROP_IFR_SEPARATOR); + location_t loc = UNKNOWN_LOCATION; + tree value = TREE_VEC_ELT (fr_list, cnt); + if (value != NULL_TREE && value != error_mark_node) + { + loc = EXPR_LOCATION (value); + if (value && TREE_CODE (value) == NOP_EXPR) + value = TREE_OPERAND (value, 0); + value = cp_fully_fold (value); + } + if (value != NULL_TREE && value != error_mark_node) + { + if (TREE_CODE (value) != INTEGER_CST + || !tree_fits_shwi_p (value)) + error_at (loc, + "expected string literal or " + "constant integer expression instead of %qE", value); // FIXME of 'qE' and no 'loc'? + else + { + HOST_WIDE_INT n = tree_to_shwi (value); + if (n < 1 || n > GOMP_INTEROP_IFR_LAST) + { + warning_at (loc, OPT_Wopenmp, + "unknown foreign runtime identifier %qwd", n); + n = GOMP_INTEROP_IFR_UNKNOWN; + } + str[0] = (char) n; + } + } + str++; + } + else if (str[0] != (char) GOMP_INTEROP_IFR_SEPARATOR) + { + /* Assume a no or a single 'fr'. */ + gcc_checking_assert (str[1] == (char) GOMP_INTEROP_IFR_SEPARATOR); + str++; + } + str++; + while (str[0] != '\0') + str += strlen (str) + 1; + str++; + cnt++; + if (cnt >= len) + break; + } + OMP_CLAUSE_INIT_PREFER_TYPE (c) = t; +} + /* For all elements of CLAUSES, validate them vs OpenMP constraints. Remove any elements from the list that are invalid. */ @@ -7408,8 +7479,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bool mergeable_seen = false; bool implicit_moved = false; bool target_in_reduction_seen = false; - bool partial_seen = false; bool num_tasks_seen = false; + bool partial_seen = false; + bool init_seen = false; + bool init_use_destroy_seen = false; + tree init_no_targetsync_clause = NULL_TREE; + tree depend_clause = NULL_TREE; bitmap_obstack_initialize (NULL); bitmap_initialize (&generic_head, &bitmap_default_obstack); @@ -8613,6 +8688,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) pc = &OMP_CLAUSE_CHAIN (c); continue; case OMP_CLAUSE_DEPEND: + depend_clause = c; + /* FALLTHRU */ case OMP_CLAUSE_AFFINITY: t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST @@ -9638,7 +9715,37 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) OMP_CLAUSE_PARTIAL_EXPR (c) = t; break; - + case OMP_CLAUSE_INIT: + init_seen = true; + cp_omp_init_prefer_type_update (c); + if (!OMP_CLAUSE_INIT_TARGETSYNC (c)) + init_no_targetsync_clause = c; + /* FALLTHRU */ + case OMP_CLAUSE_DESTROY: + case OMP_CLAUSE_USE: + init_use_destroy_seen = true; + t = OMP_CLAUSE_DECL (c); + if (bitmap_bit_p (&generic_head, DECL_UID (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD appears more than once in action clauses", t); + remove = true; + } + if (!processing_template_decl) + { + if (/* ort == C_ORT_OMP_INTEROP [uncomment for depobj init] */ + !c_omp_interop_t_p (TREE_TYPE (OMP_CLAUSE_DECL (c)))) + error_at (OMP_CLAUSE_LOCATION (c), + "%qD must be of %<omp_interop_t%>", + OMP_CLAUSE_DECL (c)); + else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE + && TREE_READONLY (OMP_CLAUSE_DECL (c))) + error_at (OMP_CLAUSE_LOCATION (c), + "%qD shall not be const", OMP_CLAUSE_DECL (c)); + } + bitmap_set_bit (&generic_head, DECL_UID (t)); + pc = &OMP_CLAUSE_CHAIN (c); + break; default: gcc_unreachable (); } @@ -10076,6 +10183,19 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) pc = &OMP_CLAUSE_CHAIN (c); } + if (ort == C_ORT_OMP_INTEROP + && depend_clause + && (!init_use_destroy_seen + || (init_seen && init_no_targetsync_clause))) + { + error_at (OMP_CLAUSE_LOCATION (depend_clause), + "%<depend%> clause requires action clauses with " + "%<targetsync%> interop-type"); + if (init_no_targetsync_clause) + inform (OMP_CLAUSE_LOCATION (init_no_targetsync_clause), + "%<init%> clause lacks the %<targetsync%> modifier"); + } + bitmap_obstack_release (NULL); return clauses; } diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp index 65c4c761e331..c6652f0a9734 100644 --- a/gcc/fortran/ChangeLog.omp +++ b/gcc/fortran/ChangeLog.omp @@ -1,3 +1,22 @@ +2025-01-27 Tobias Burnus <tbur...@baylibre.com> + + Backported from master: + 2024-11-22 Tobias Burnus <tbur...@baylibre.com> + + * gfortran.h (gfc_omp_namelist): Cleanup interop internal + representation. + * dump-parse-tree.cc (show_omp_namelist): Update for changed + internal representation. + * match.cc (gfc_free_omp_namelist): Likewise. + * openmp.cc (gfc_match_omp_prefer_type, gfc_match_omp_init): + Likewise; also handle some corner cases better and update for + newer 6.0 changes related to 'fr'. + (resolve_omp_clauses): Add type-check for interop variables. + * trans-openmp.cc (gfc_trans_omp_clauses): Handle init, use + and destroy clauses. + (gfc_trans_openmp_interop): New. + (gfc_trans_omp_directive): Call it. + 2025-01-27 Paul-Antoine Arras <par...@baylibre.com> Backported from master: diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc index 60c03cf0b733..a75dfb42b2fc 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -1539,63 +1539,42 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n) fputs ("target,", dumpfile); if (n->u.init.targetsync) fputs ("targetsync,", dumpfile); - if (n->u2.init_interop_fr) + if (n->u2.init_interop) { - char *attr_str = n->u.init.attr; - int idx = 0; - int fr_id; + char *str = n->u2.init_interop; fputs ("prefer_type(", dumpfile); - do + while (str[0] == (char) GOMP_INTEROP_IFR_SEPARATOR) { - fr_id = n->u2.init_interop_fr[idx]; + bool has_fr = false; fputc ('{', dumpfile); - if (fr_id != GOMP_INTEROP_IFR_NONE) + str++; + while (str[0] != (char) GOMP_INTEROP_IFR_SEPARATOR) { - fputs ("fr(", dumpfile); - do - { - const char *fr_str = omp_get_name_from_fr_id (fr_id); - if (fr_str) - fprintf (dumpfile, "\"%s\"", fr_str); - else - fprintf (dumpfile, "%d", fr_id); - fr_id = n->u2.init_interop_fr[++idx]; - if (fr_id != GOMP_INTEROP_IFR_SEPARATOR) - fputc (',', dumpfile); - } - while (fr_id != GOMP_INTEROP_IFR_SEPARATOR); - fputc (')', dumpfile); - if (attr_str && (attr_str[0] != ' ' || attr_str[1] != '\0')) + if (has_fr) fputc (',', dumpfile); + has_fr = true; + fputs ("fr(\"", dumpfile); + fputs (omp_get_name_from_fr_id (str[0]), dumpfile); + fputs ("\")", dumpfile); + str++; } - else - fr_id = n->u2.init_interop_fr[++idx]; - if (attr_str && attr_str[0] == ' ' && attr_str[1] == '\0') - attr_str += 2; - else if (attr_str) + str++; + if (has_fr && str[0] != '\0') + fputc (',', dumpfile); + while (str[0] != '\0') { fputs ("attr(\"", dumpfile); - do - { - fputs ((char *) attr_str, dumpfile); - fputc ('"', dumpfile); - attr_str += strlen (attr_str) + 1; - if (attr_str[0] == '\0') - break; - fputs (",\"", dumpfile); - } - while (true); - fputc (')', dumpfile); + fputs (str, dumpfile); + fputs ("\")", dumpfile); + str += strlen (str) + 1; + if (str[0] != '\0') + fputc (',', dumpfile); } + str++; fputc ('}', dumpfile); - fr_id = n->u2.init_interop_fr[++idx]; - if (fr_id == GOMP_INTEROP_IFR_SEPARATOR) - break; - fputc (',', dumpfile); - if (attr_str) - ++attr_str; + if (str[0] != '\0') + fputs (", ", dumpfile); } - while (true); fputc (')', dumpfile); } fputc (':', dumpfile); diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 97ed99c8b195..96f884b5b965 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1393,7 +1393,6 @@ typedef struct gfc_omp_namelist bool present_modifier; struct { - char *attr; int len; bool target; bool targetsync; @@ -1408,7 +1407,7 @@ typedef struct gfc_omp_namelist gfc_expr *allocator; struct gfc_symbol *traits_sym; struct gfc_omp_namelist *duplicate_of; - char *init_interop_fr; + char *init_interop; } u2; struct gfc_omp_namelist *next; locus where; diff --git a/gcc/fortran/match.cc b/gcc/fortran/match.cc index 3ae0f72a00ad..0f01e28dc692 100644 --- a/gcc/fortran/match.cc +++ b/gcc/fortran/match.cc @@ -5548,7 +5548,7 @@ gfc_free_omp_namelist (gfc_omp_namelist *name, int list) bool free_init = (list == OMP_LIST_INIT); gfc_omp_namelist *n; gfc_expr *last_allocator = NULL; - char *last_init_attr = NULL; + char *last_init_interop = NULL; for (; name; name = n) { @@ -5572,11 +5572,10 @@ gfc_free_omp_namelist (gfc_omp_namelist *name, int list) { } /* name->u2.traits_sym: shall not call gfc_free_symbol here. */ else if (free_init) { - if (name->u.init.attr != last_init_attr) + if (name->u2.init_interop != last_init_interop) { - last_init_attr = name->u.init.attr; - free (name->u.init.attr); - free (name->u2.init_interop_fr); + last_init_interop = name->u2.init_interop; + free (name->u2.init_interop); } } else if (free_mapper && name->u2.udm) diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index edc88d549800..4fb1c3bb09f3 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -1912,45 +1912,42 @@ error: prefer_type ( <const-int-expr|string literal> [, ...] or prefer_type ( '{' <fr(...) | attr (...)>, ...] '}' [, '{' ... '}' ] ) - where 'fr' takes an integer named constant or a string literal - and 'attr takes a string literal, starting with 'ompx_') + where 'fr' takes a constant expression or a string literal + and 'attr takes a list of string literals, starting with 'ompx_') For the foreign runtime identifiers, string values are converted to - their integer value; unknown string or integer values are set to 0. - - For the simple syntax, pref_int_array contains alternatingly the - fr_id integer value and GOMP_INTEROP_IFR_SEPARATOR followed by a - GOMP_INTEROP_IFR_SEPARATOR as last item. - For the complex syntax, it contains the values associated with a - 'fr(...)' followed by GOMP_INTEROP_IFR_SEPARATOR. If there is no - 'fr' in a curly-brace block, it is GOMP_INTEROP_IFR_NONE followed - by GOMP_INTEROP_IFR_SEPARATOR. An additional GOMP_INTEROP_IFR_SEPARATOR - at the end terminates the array. - - For attributes, if the simply syntax is used, it is NULL - likewise if no - 'attr' appears. For the complex syntax it is: For reach curly-brace block, - it is \0\0 is no attr appears and otherwise a concatenation (including - the \0) of all 'attr' strings followed by a tailing '\0'. At the end, - another '\0' follows. */ + their integer value; unknown string or integer values are set to + GOMP_INTEROP_IFR_KNOWN. + + Data format: + For the foreign runtime identifiers, string values are converted to + their integer value; unknown string or integer values are set to 0. + + Each item (a) GOMP_INTEROP_IFR_SEPARATOR + (b) for any 'fr', its integer value. + Note: Spec only permits 1 'fr' entry (6.0; changed after TR13) + (c) GOMP_INTEROP_IFR_SEPARATOR + (d) list of \0-terminated non-empty strings for 'attr' + (e) '\0' + Tailing '\0'. */ static match -gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_len) +gfc_match_omp_prefer_type (char **type_str, int *type_str_len) { gfc_expr *e; - int cnt_brace_grp = 0; - std::vector<char> int_list; - std::string attr_string; + std::string type_string, attr_string; /* New syntax. */ if (gfc_peek_ascii_char () == '{') do { + attr_string.clear (); + type_string += (char) GOMP_INTEROP_IFR_SEPARATOR; if (gfc_match ("{ ") != MATCH_YES) { gfc_error ("Expected %<{%> at %C"); return MATCH_ERROR; } bool fr_found = false; - bool attr_found = false; do { if (gfc_match ("fr ( ") == MATCH_YES) @@ -1964,24 +1961,27 @@ gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l fr_found = true; do { - if (gfc_match_expr (&e) != MATCH_YES) - return MATCH_ERROR; - if (e->expr_type != EXPR_CONSTANT - || e->ref != NULL + bool found_literal = false; + match m = MATCH_YES; + if (gfc_match_literal_constant (&e, false) == MATCH_YES) + found_literal = true; + else + m = gfc_match_expr (&e); + if (m != MATCH_YES || !gfc_resolve_expr (e) + || e->rank != 0 + || e->expr_type != EXPR_CONSTANT || (e->ts.type != BT_INTEGER - && e->ts.type != BT_CHARACTER) + && (!found_literal || e->ts.type != BT_CHARACTER)) || (e->ts.type == BT_INTEGER - && (!e->symtree - || e->symtree->n.sym->attr.flavor != FL_PARAMETER - || !mpz_fits_sint_p (e->value.integer))) + && !mpz_fits_sint_p (e->value.integer)) || (e->ts.type == BT_CHARACTER && (e->ts.kind != gfc_default_character_kind - || e->value.character.length == 0))) + || e->value.character.length == 0))) { - gfc_error ("Expected scalar integer parameter or " - "non-empty default-kind character literal " - "at %L", &e->where); + gfc_error ("Expected constant scalar integer expression" + " or non-empty default-kind character " + "literal at %L", &e->where); gfc_free_expr (e); return MATCH_ERROR; } @@ -1992,10 +1992,11 @@ gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l val = mpz_get_si (e->value.integer); if (val < 1 || val > GOMP_INTEROP_IFR_LAST) { - gfc_warning (OPT_Wopenmp, - "Unknown foreign runtime identifier " - "%qd at %L", val, &e->where); - val = 0; + gfc_warning_now (OPT_Wopenmp, + "Unknown foreign runtime " + "identifier %qd at %L", + val, &e->where); + val = GOMP_INTEROP_IFR_UNKNOWN; } } else @@ -2011,40 +2012,30 @@ gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l return MATCH_ERROR; } val = omp_get_fr_id_from_name (str); - if (val == 0) - gfc_warning (OPT_Wopenmp, - "Unknown foreign runtime identifier %qs " - "at %L", str, &e->where); + if (val == GOMP_INTEROP_IFR_UNKNOWN) + gfc_warning_now (OPT_Wopenmp, + "Unknown foreign runtime identifier " + "%qs at %L", str, &e->where); } - int_list.push_back (val); - if (gfc_match (", ") == MATCH_YES) - continue; + + type_string += (char) val; if (gfc_match (") ") == MATCH_YES) break; - gfc_error ("Expected %<,%> or %<)%> at %C"); + gfc_error ("Expected %<)%> at %C"); return MATCH_ERROR; } while (true); } else if (gfc_match ("attr ( ") == MATCH_YES) { - attr_found = true; - if (attr_string.empty ()) - for (int i = 0; i < cnt_brace_grp; ++i) - { - /* Add dummy elements for previous curly-brace blocks. */ - attr_string += ' '; - attr_string += '\0'; - attr_string += '\0'; - } do { - if (gfc_match_expr (&e) != MATCH_YES) - return MATCH_ERROR; - if (e->expr_type != EXPR_CONSTANT + if (gfc_match_literal_constant (&e, false) != MATCH_YES + || !gfc_resolve_expr (e) + || e->expr_type != EXPR_CONSTANT || e->rank != 0 || e->ts.type != BT_CHARACTER - || e->ts.kind != gfc_default_character_kind) + || e->ts.kind != gfc_default_character_kind) { gfc_error ("Expected default-kind character literal " "at %L", &e->where); @@ -2093,21 +2084,9 @@ gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l return MATCH_ERROR; } while (true); - ++cnt_brace_grp; - if (!fr_found) - int_list.push_back (GOMP_INTEROP_IFR_NONE); - int_list.push_back (GOMP_INTEROP_IFR_SEPARATOR); - if (!attr_string.empty ()) - { - if (!attr_found) - { - /* Dummy entry. */ - attr_string += ' '; - attr_string += '\0'; - } - attr_string += '\0'; - } - + type_string += (char) GOMP_INTEROP_IFR_SEPARATOR; + type_string += attr_string; + type_string += '\0'; if (gfc_match (", ") == MATCH_YES) continue; if (gfc_match (") ") == MATCH_YES) @@ -2119,12 +2098,19 @@ gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l else do { - if (gfc_match_expr (&e) != MATCH_YES) - return MATCH_ERROR; - if (!gfc_resolve_expr (e) + type_string += (char) GOMP_INTEROP_IFR_SEPARATOR; + bool found_literal = false; + match m = MATCH_YES; + if (gfc_match_literal_constant (&e, false) == MATCH_YES) + found_literal = true; + else + m = gfc_match_expr (&e); + if (m != MATCH_YES + || !gfc_resolve_expr (e) || e->rank != 0 || e->expr_type != EXPR_CONSTANT - || (e->ts.type != BT_INTEGER && e->ts.type != BT_CHARACTER) + || (e->ts.type != BT_INTEGER + && (!found_literal || e->ts.type != BT_CHARACTER)) || (e->ts.type == BT_INTEGER && !mpz_fits_sint_p (e->value.integer)) || (e->ts.type == BT_CHARACTER @@ -2143,9 +2129,9 @@ gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l val = mpz_get_si (e->value.integer); if (val < 1 || val > GOMP_INTEROP_IFR_LAST) { - gfc_warning (OPT_Wopenmp, - "Unknown foreign runtime identifier %qd at %L", - val, &e->where); + gfc_warning_now (OPT_Wopenmp, + "Unknown foreign runtime identifier %qd at %L", + val, &e->where); val = 0; } } @@ -2161,13 +2147,14 @@ gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l return MATCH_ERROR; } val = omp_get_fr_id_from_name (str); - if (val == 0) - gfc_warning (OPT_Wopenmp, - "Unknown foreign runtime identifier %qs at %L", - str, &e->where); + if (val == GOMP_INTEROP_IFR_UNKNOWN) + gfc_warning_now (OPT_Wopenmp, + "Unknown foreign runtime identifier %qs at %L", + str, &e->where); } - int_list.push_back (val); - int_list.push_back (GOMP_INTEROP_IFR_SEPARATOR); + type_string += (char) val; + type_string += (char) GOMP_INTEROP_IFR_SEPARATOR; + type_string += '\0'; gfc_free_expr (e); if (gfc_match (", ") == MATCH_YES) continue; @@ -2177,17 +2164,10 @@ gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l return MATCH_ERROR; } while (true); - int_list.push_back (GOMP_INTEROP_IFR_SEPARATOR); - *fr_int_array = XNEWVEC (char, int_list.size ()); - memcpy (*fr_int_array, int_list.data (), sizeof (char) * int_list.size ()); - - if (!attr_string.empty ()) - { - attr_string += '\0'; - *attr_str_len = attr_string.length(); - *attr_str = XNEWVEC (char, attr_string.length ()); - memcpy (*attr_str, attr_string.data (), attr_string.length ()); - } + type_string += '\0'; + *type_str_len = type_string.length(); + *type_str = XNEWVEC (char, type_string.length ()); + memcpy (*type_str, type_string.data (), type_string.length ()); return MATCH_YES; } @@ -2199,21 +2179,19 @@ static match gfc_match_omp_init (gfc_omp_namelist **list) { bool target = false, targetsync = false; - char *fr_int_array = NULL; - char *attr_str = NULL; - int attr_str_len = 0; + char *type_str = NULL; + int type_str_len = 0; match m; locus old_loc = gfc_current_locus; do { if (gfc_match ("prefer_type ( ") == MATCH_YES) { - if (fr_int_array) + if (type_str) { gfc_error ("Duplicate %<prefer_type%> modifier at %C"); return MATCH_ERROR; } - m = gfc_match_omp_prefer_type (&fr_int_array, &attr_str, - &attr_str_len); + m = gfc_match_omp_prefer_type (&type_str, &type_str_len); if (m != MATCH_YES) return m; if (gfc_match (", ") == MATCH_YES) @@ -2225,16 +2203,21 @@ gfc_match_omp_init (gfc_omp_namelist **list) } if (gfc_match ("targetsync ") == MATCH_YES) { + if (targetsync) + { + /* Avoid the word 'modifier' as it could be also be no clauses and + twice a variable named 'targetsync', which is also invalid. */ + gfc_error ("Duplicate %<targetsync%> at %C"); + return MATCH_ERROR; + } targetsync = true; if (gfc_match (", ") == MATCH_YES) continue; if (gfc_match (": ") == MATCH_YES) break; gfc_char_t c = gfc_peek_char (); - if (!fr_int_array - && (c == ')' - || (gfc_current_form != FORM_FREE - && (c == '_' || ISALPHA (c))))) + if (!type_str && (c == ')' || (gfc_current_form != FORM_FREE + && (c == '_' || ISALPHA (c))))) { gfc_current_locus = old_loc; break; @@ -2244,16 +2227,19 @@ gfc_match_omp_init (gfc_omp_namelist **list) } if (gfc_match ("target ") == MATCH_YES) { + if (target) + { + gfc_error ("Duplicate %<target%> at %C"); + return MATCH_ERROR; + } target = true; if (gfc_match (", ") == MATCH_YES) continue; if (gfc_match (": ") == MATCH_YES) break; gfc_char_t c = gfc_peek_char (); - if (!fr_int_array - && (c == ')' - || (gfc_current_form != FORM_FREE - && (c == '_' || ISALPHA (c))))) + if (!type_str && (c == ')' || (gfc_current_form != FORM_FREE + && (c == '_' || ISALPHA (c))))) { gfc_current_locus = old_loc; break; @@ -2261,7 +2247,7 @@ gfc_match_omp_init (gfc_omp_namelist **list) gfc_error ("Expected %<,%> or %<:%> at %C"); return MATCH_ERROR; } - if (fr_int_array) + if (type_str) { gfc_error ("Expected %<target%> or %<targetsync%> at %C"); return MATCH_ERROR; @@ -2278,9 +2264,8 @@ gfc_match_omp_init (gfc_omp_namelist **list) { n->u.init.target = target; n->u.init.targetsync = targetsync; - n->u.init.attr = attr_str; - n->u.init.len = attr_str_len; - n->u2.init_interop_fr = fr_int_array; + n->u.init.len = type_str_len; + n->u2.init_interop = type_str; } return MATCH_YES; } @@ -9584,6 +9569,21 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, break; } } + if (code && code->op == EXEC_OMP_INTEROP) + for (list = OMP_LIST_INIT; list <= OMP_LIST_DESTROY; list++) + for (n = omp_clauses->lists[list]; n; n = n->next) + { + if (n->sym->ts.type != BT_INTEGER + || n->sym->ts.kind != gfc_index_integer_kind + || n->sym->attr.dimension + || n->sym->attr.flavor != FL_VARIABLE) + gfc_error ("%qs at %L in %qs clause must be a scalar integer " + "variable of %<omp_interop_kind%> kind", n->sym->name, + &n->where, clause_names[list]); + if (list != OMP_LIST_USE && n->sym->attr.intent == INTENT_IN) + gfc_error ("%qs at %L in %qs clause must be definable", + n->sym->name, &n->where, clause_names[list]); + } verify_omp_clauses_symbol_dups (code, omp_clauses, ns, openacc); diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index d4465e40bbe7..a23e56cb856f 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -4687,12 +4687,56 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, case OMP_LIST_SCAN_EX: clause_code = OMP_CLAUSE_EXCLUSIVE; goto add_clause; + case OMP_LIST_USE: + clause_code = OMP_CLAUSE_USE; + goto add_clause; + case OMP_LIST_DESTROY: + clause_code = OMP_CLAUSE_DESTROY; + goto add_clause; add_clause: omp_clauses = gfc_trans_omp_variable_list (clause_code, n, omp_clauses, declare_simd); break; + + case OMP_LIST_INIT: + { + tree pref_type = NULL_TREE; + const char *last = NULL; + for (; n != NULL; n = n->next) + if (n->sym->attr.referenced) + { + tree t = gfc_trans_omp_variable (n->sym, false); + if (t == error_mark_node) + continue; + tree node = build_omp_clause (input_location, + OMP_CLAUSE_INIT); + OMP_CLAUSE_DECL (node) = t; + if (n->u.init.target) + OMP_CLAUSE_INIT_TARGET (node) = 1; + if (n->u.init.targetsync) + OMP_CLAUSE_INIT_TARGETSYNC (node) = 1; + if (last != n->u2.init_interop) + { + last = n->u2.init_interop; + if (n->u2.init_interop == NULL) + pref_type = NULL_TREE; + else + { + pref_type = build_string (n->u.init.len, + n->u2.init_interop); + TREE_TYPE (pref_type) + = build_array_type_nelts (unsigned_char_type_node, + n->u.init.len); + } + } + OMP_CLAUSE_INIT_PREFER_TYPE (node) = pref_type; + omp_clauses = gfc_trans_add_clause (node, omp_clauses); + } + break; + } + case OMP_LIST_ALIGNED: for (; n != NULL; n = n->next) if (n->sym->attr.referenced || declare_simd) @@ -10672,6 +10716,18 @@ gfc_trans_omp_target_update (gfc_code *code) return gfc_finish_block (&block); } +static tree +gfc_trans_openmp_interop (gfc_code *code, gfc_omp_clauses *clauses) +{ + stmtblock_t block; + gfc_start_block (&block); + tree omp_clauses = gfc_trans_omp_clauses (&block, clauses, code->loc); + tree stmt = build1_loc (input_location, OMP_INTEROP, void_type_node, + omp_clauses); + gfc_add_expr_to_block (&block, stmt); + return gfc_finish_block (&block); +} + static tree gfc_trans_omp_workshare (gfc_code *code, gfc_omp_clauses *clauses) { @@ -11050,8 +11106,7 @@ gfc_trans_omp_directive (gfc_code *code) case EXEC_OMP_WORKSHARE: return gfc_trans_omp_workshare (code, code->ext.omp_clauses); case EXEC_OMP_INTEROP: - sorry ("%<!$OMP INTEROP%>"); - return build_empty_stmt (input_location); + return gfc_trans_openmp_interop (code, code->ext.omp_clauses); default: gcc_unreachable (); } diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 977275eb9c99..99b227e41f26 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -20603,6 +20603,11 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, break; } + case OMP_INTEROP: + sorry_at (EXPR_LOCATION (*expr_p), + "%<#pragma omp interop%> not yet supported"); + ret = GS_ERROR; + break; case OMP_ATOMIC: case OMP_ATOMIC_READ: case OMP_ATOMIC_CAPTURE_OLD: diff --git a/gcc/omp-api.h b/gcc/omp-api.h index 1b877f257f09..166a39525c4e 100644 --- a/gcc/omp-api.h +++ b/gcc/omp-api.h @@ -29,7 +29,7 @@ along with GCC; see the file COPYING3. If not see extern bool omp_runtime_api_procname (const char *name); extern bool omp_runtime_api_call (const_tree fndecl); -extern int omp_get_fr_id_from_name (const char *); +extern char omp_get_fr_id_from_name (const char *); extern const char *omp_get_name_from_fr_id (int); #endif diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc index ddd2cacc13b0..ec76efa04315 100644 --- a/gcc/omp-general.cc +++ b/gcc/omp-general.cc @@ -3569,7 +3569,7 @@ static const char* omp_interop_fr_str[] = {"cuda", "cuda_driver", "opencl", /* Returns the foreign-runtime ID if found or 0 otherwise. */ -int +char omp_get_fr_id_from_name (const char *str) { static_assert (GOMP_INTEROP_IFR_LAST == ARRAY_SIZE (omp_interop_fr_str), ""); @@ -3577,7 +3577,7 @@ omp_get_fr_id_from_name (const char *str) for (unsigned i = 0; i < ARRAY_SIZE (omp_interop_fr_str); ++i) if (!strcmp (str, omp_interop_fr_str[i])) return i + 1; - return 0; + return GOMP_INTEROP_IFR_UNKNOWN; } /* Returns the string value to a foreign-runtime integer value or NULL if value @@ -3587,7 +3587,7 @@ const char * omp_get_name_from_fr_id (int fr_id) { if (fr_id < 1 || fr_id > (int) ARRAY_SIZE (omp_interop_fr_str)) - return NULL; + return "<unknown>"; return omp_interop_fr_str[fr_id-1]; } diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index db5c67dc4da6..f5b97f749c7d 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,19 @@ +2025-01-27 Tobias Burnus <tbur...@baylibre.com> + + Backported from master: + 2024-11-22 Tobias Burnus <tbur...@baylibre.com> + + * gfortran.dg/gomp/interop-1.f90: Update for parser changes, + spec changes and add new tests. + * gfortran.dg/gomp/interop-2.f90: Likewise. + * gfortran.dg/gomp/interop-3.f90: Likewise. + * c-c++-common/gomp/interop-1.c: New test. + * c-c++-common/gomp/interop-2.c: New test. + * c-c++-common/gomp/interop-3.c: New test. + * c-c++-common/gomp/interop-4.c: New test. + * g++.dg/gomp/interop-5.C: New test. + * gfortran.dg/gomp/interop-4.f90: New test. + 2025-01-27 Paul-Antoine Arras <par...@baylibre.com> Backported from master: diff --git a/gcc/testsuite/c-c++-common/gomp/interop-1.c b/gcc/testsuite/c-c++-common/gomp/interop-1.c new file mode 100644 index 000000000000..de3a4ba4b6bc --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/interop-1.c @@ -0,0 +1,119 @@ +/* { dg-do compile { target { c || c++11 } } } */ +/* { dg-additional-options "-std=c23" { target c } } */ +/* C++11 and C23 because of 'constexpr'. */ + +/* { dg-prune-output "sorry, unimplemented: '#pragma omp interop' not yet supported" } */ + +/* The following definitions are in omp_lib, which cannot be included + in gcc/testsuite/ */ + +#if __cplusplus >= 201103L +# define __GOMP_UINTPTR_T_ENUM : __UINTPTR_TYPE__ +#else +# define __GOMP_UINTPTR_T_ENUM +#endif + +typedef enum omp_interop_t __GOMP_UINTPTR_T_ENUM +{ + omp_interop_none = 0, + __omp_interop_t_max__ = __UINTPTR_MAX__ +} omp_interop_t; + +typedef enum omp_interop_fr_t +{ + omp_ifr_cuda = 1, + omp_ifr_cuda_driver = 2, + omp_ifr_opencl = 3, + omp_ifr_sycl = 4, + omp_ifr_hip = 5, + omp_ifr_level_zero = 6, + omp_ifr_hsa = 7, + omp_ifr_last = omp_ifr_hsa +} omp_interop_fr_t; + + +// --------------------------------- + +void f() +{ + constexpr omp_interop_fr_t ifr_scalar = omp_ifr_hsa; + constexpr omp_interop_fr_t ifr_array[] = {omp_ifr_cuda, omp_ifr_hip}; + constexpr char my_string[] = "cuda"; + omp_interop_t obj1, obj2, obj3, obj4, obj5; + int x; + + #pragma omp interop init(obj1) init(target,targetsync : obj2, obj3) nowait // OK + #pragma omp interop init(obj1) init (targetsync : obj2, obj3) nowait // OK + #pragma omp interop init(obj1) init (targetsync , target : obj2, obj3) nowait // OK + + #pragma omp interop init(obj1) init(target,targetsync,target: obj2, obj3) nowait // { dg-error "duplicate 'target' modifier" } + #pragma omp interop init(obj1) init(target,targetsync, targetsync : obj2, obj3) nowait // { dg-error "duplicate 'targetsync' modifier" } + + #pragma omp interop init(prefer_type("cuda", omp_ifr_opencl, omp_ifr_level_zero, "hsa"), targetsync : obj1) \ + destroy(obj2, obj3) depend(inout: x) use(obj4, obj5) device(device_num: 0) + + #pragma omp interop init(prefer_type("cu" "da"), targetsync : obj1) // OK + + #pragma omp assume contains(interop) + { + #pragma omp interop init(prefer_type("cuða") : obj3) // { dg-warning "unknown foreign runtime identifier 'cu\[^'\]*a'" } + } + + #pragma omp interop init(prefer_type("cu\0da") : obj3) // { dg-error "string literal must not contain '\\\\0'" } + + #pragma omp interop depend(inout: x) , use(obj2), destroy(obj3) // OK, use or destroy might have 'targetsync' + + #pragma omp interop depend(inout: x) use(obj2), destroy(obj3) // Likewise + + #pragma omp interop depend(inout: x) init(targetsync : obj5) use(obj2), destroy(obj3) init(prefer_type("cuda"), targetsync : obj4) // OK + + #pragma omp interop init ( target , prefer_type( { fr("hsa") }, "hip") : obj1) // { dg-error "expected '\{' before string constant" } + + #pragma omp interop init ( target , prefer_type( { fr("hsa"), attr("ompx_nothing") , fr("hsa" ) }) :obj1) // { dg-error "duplicated 'fr' preference selector before '\\(' token" } + + #pragma omp interop init ( prefer_type( 4, omp_ifr_hip*4) : obj1) // { dg-warning "unknown foreign runtime identifier '20'" } + #pragma omp interop init ( prefer_type( __builtin_sin(3.3) : obj1) + // { dg-error "'prefer_type' undeclared \\(first use in this function\\)" "" { target c } .-1 } + // { dg-error "'prefer_type' has not been declared" "" { target c++ } .-2 } + // { dg-error "expected '\\)' before '\\(' token" "" { target *-*-* } .-3 } + + #pragma omp interop init ( prefer_type( __builtin_sin(3.3) ) : obj1) // { dg-error "expected string literal or constant integer expression before '\\)' token" } + #pragma omp interop init ( prefer_type( {fr(4 ) }) : obj1) // OK + #pragma omp interop init ( prefer_type( {fr("cu\0da" ) }) : obj1) // { dg-error "string literal must not contain '\\\\0'" } + #pragma omp interop init ( prefer_type( {fr("cuda\0") }) : obj1) // { dg-error "string literal must not contain '\\\\0'" } + #pragma omp interop init ( prefer_type( {fr("cuda" ) }) : obj1) // OK + #pragma omp interop init ( prefer_type( {fr(omp_ifr_level_zero ) }, {fr(omp_ifr_hip)}) : obj1) // OK + #pragma omp interop init ( prefer_type( {fr("cuda", "cuda_driver") }) : obj1) // { dg-error "53: expected '\\)' before ',' token" } + #pragma omp interop init ( prefer_type( {fr(my_string) }) : obj1) // { dg-error "56: expected string literal or constant integer expression before '\\)' token" } + #pragma omp interop init ( prefer_type( {fr("hello" }) : obj1) // { dg-error "expected '\\)' before '\\(' token" } + // { dg-error "'prefer_type' has not been declared" "" { target c++ } .-1 } + #pragma omp interop init ( prefer_type( {fr("hello") }) : obj1) + /* { dg-warning "unknown foreign runtime identifier 'hello' \\\[-Wopenmp\\\]" "" { target *-*-* } .-1 } */ + + #pragma omp interop init ( prefer_type( {fr(x) }) : obj1) // { dg-error "expected string literal or constant integer expression before '\\)' token" } + + #pragma omp interop init ( prefer_type( {fr(ifr_scalar ) }) : obj1) // OK + #pragma omp interop init ( prefer_type( {fr(ifr_array ) }) : obj1) // { dg-error "expected string literal or constant integer expression before '\\)' token" } + // OK in C++, for C: constexpr arrays are not part of C23; however, they are/were under consideration for C2y. + #pragma omp interop init ( prefer_type( {fr(ifr_array[0] ) }) : obj1) + // { dg-error "expected string literal or constant integer expression before '\\)' token" "" { target c } .-1 } + + #pragma omp interop init ( prefer_type( omp_ifr_level_zero, omp_ifr_hip ) : obj1) // OK + #pragma omp interop init ( prefer_type( omp_ifr_level_zero +1 ) : obj1) // OK + #pragma omp interop init ( prefer_type( x ) : obj1) // { dg-error "expected string literal or constant integer expression before '\\)' token" } + + #pragma omp interop init ( prefer_type( ifr_scalar ) : obj1) // OK + #pragma omp interop init ( prefer_type( ifr_array ) : obj1) // { dg-error "expected string literal or constant integer expression before '\\)' token" } + // OK in C++, for C: constexpr arrays are not part of C23; however, they are/were under consideration for C2y. + #pragma omp interop init ( prefer_type( ifr_array[1] ) : obj1) + // { dg-error "expected string literal or constant integer expression before '\\)' token" "" { target c } .-1 } + + #pragma omp interop init ( prefer_type( 4, omp_ifr_hip*4) : obj1) // { dg-warning "unknown foreign runtime identifier '20'" } + #pragma omp interop init ( prefer_type( 4, 1, 3) : obj1) + + #pragma omp interop init ( prefer_type( {fr("cuda") }, {fr(omp_ifr_hsa)} , {attr("ompx_a") } , {fr(omp_ifr_hip) }) : obj1) + #pragma omp interop init ( prefer_type( {fr("cuda") }, {fr(omp_ifr_hsa,omp_ifr_level_zero)} , {attr("ompx_a") } , {fr(omp_ifr_hip) }) : obj1) // { dg-error "73: expected '\\)' before ',' token" } + #pragma omp interop init ( prefer_type( {fr("cuda",5) }, {fr(omp_ifr_hsa,omp_ifr_level_zero)} , {attr("ompx_a") } , {fr(omp_ifr_hip) }) : obj1) // { dg-error "53: expected '\\)' before ',' token" } + #pragma omp interop init ( prefer_type( {fr("sycl"), attr("ompx_1", "ompx_2"), attr("ompx_3") }, {attr("ompx_4", "ompx_5"),fr(omp_ifr_level_zero)} ) : obj1) + #pragma omp interop init ( prefer_type( { fr(5), attr("ompx_1") }, {fr(omp_ifr_hsa)} , {attr("ompx_a") } ) : obj1) +} diff --git a/gcc/testsuite/c-c++-common/gomp/interop-2.c b/gcc/testsuite/c-c++-common/gomp/interop-2.c new file mode 100644 index 000000000000..57fd688d55fe --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/interop-2.c @@ -0,0 +1,127 @@ +/* { dg-do compile { target { c || c++11 } } } */ +/* { dg-additional-options "-std=c23" { target c } } */ +/* C++11 and C23 because of 'constexpr'. */ + +/* { dg-prune-output "sorry, unimplemented: '#pragma omp interop' not yet supported" } */ + +/* The following definitions are in omp_lib, which cannot be included + in gcc/testsuite/ */ + +#if __cplusplus >= 201103L +# define __GOMP_UINTPTR_T_ENUM : __UINTPTR_TYPE__ +#else +# define __GOMP_UINTPTR_T_ENUM +#endif + +typedef enum omp_interop_t __GOMP_UINTPTR_T_ENUM +{ + omp_interop_none = 0, + __omp_interop_t_max__ = __UINTPTR_MAX__ +} omp_interop_t; + +typedef enum omp_interop_fr_t +{ + omp_ifr_cuda = 1, + omp_ifr_cuda_driver = 2, + omp_ifr_opencl = 3, + omp_ifr_sycl = 4, + omp_ifr_hip = 5, + omp_ifr_level_zero = 6, + omp_ifr_hsa = 7, + omp_ifr_last = omp_ifr_hsa +} omp_interop_fr_t; + + +// --------------------------------- + +void f(const omp_interop_t ocp) +{ + constexpr omp_interop_t oce = omp_interop_none; + const omp_interop_t occ = omp_interop_none; + omp_interop_t od[5]; + omp_interop_t *op; + short o2; + float of; + + #pragma omp interop init (ocp) // { dg-error "'ocp' shall not be const" } + #pragma omp interop init (oce) // { dg-error "'oce' shall not be const" } + #pragma omp interop init (occ) // { dg-error "'occ' shall not be const" } + #pragma omp interop init (od) // { dg-error "'od' must be of 'omp_interop_t'" } + #pragma omp interop init (od[1])// { dg-error "expected '\\)' before '\\\[' token" } + // { dg-error "'od' must be of 'omp_interop_t'" "" { target *-*-* } .-1 } + #pragma omp interop init (op) // { dg-error "'op' must be of 'omp_interop_t'" } + #pragma omp interop init (*op) + // { dg-error "expected identifier before '\\*' token" "" { target c } .-1 } + // { dg-error "expected unqualified-id before '\\*' token" "" { target c++ } .-2 } + #pragma omp interop init (o2) // { dg-error "'o2' must be of 'omp_interop_t'" } + #pragma omp interop init (of) // { dg-error "'of' must be of 'omp_interop_t'" } + + #pragma omp interop use (ocp) // OK + #pragma omp interop use (oce) // odd but okay + #pragma omp interop use (occ) // okayish + #pragma omp interop use (od) // { dg-error "'od' must be of 'omp_interop_t'" } + #pragma omp interop use (od[1])// { dg-error "expected '\\)' before '\\\[' token" } + // { dg-error "'od' must be of 'omp_interop_t'" "" { target *-*-* } .-1 } + #pragma omp interop use (op) // { dg-error "'op' must be of 'omp_interop_t'" } + #pragma omp interop use (*op) + // { dg-error "expected identifier before '\\*' token" "" { target c } .-1 } + // { dg-error "expected unqualified-id before '\\*' token" "" { target c++ } .-2 } + #pragma omp interop use (o2) // { dg-error "'o2' must be of 'omp_interop_t'" } + #pragma omp interop use (of) // { dg-error "'of' must be of 'omp_interop_t'" } + + #pragma omp interop destroy (ocp) // { dg-error "'ocp' shall not be const" } + #pragma omp interop destroy (oce) // { dg-error "'oce' shall not be const" } + #pragma omp interop destroy (occ) // { dg-error "'occ' shall not be const" } + #pragma omp interop destroy (od) // { dg-error "'od' must be of 'omp_interop_t'" } + #pragma omp interop destroy (od[1])// { dg-error "expected '\\)' before '\\\[' token" } + // { dg-error "'od' must be of 'omp_interop_t'" "" { target *-*-* } .-1 } + #pragma omp interop destroy (op) // { dg-error "'op' must be of 'omp_interop_t'" } + #pragma omp interop destroy (*op) + // { dg-error "expected identifier before '\\*' token" "" { target c } .-1 } + // { dg-error "expected unqualified-id before '\\*' token" "" { target c++ } .-2 } + #pragma omp interop destroy (o2) // { dg-error "'o2' must be of 'omp_interop_t'" } + #pragma omp interop destroy (of) // { dg-error "'of' must be of 'omp_interop_t'" } +} + +void g() +{ + omp_interop_t obj1, obj2, obj3, obj4, obj5; + int x; + + #pragma omp interop init ( prefer_type( {fr("") }) : obj1) // { dg-error "non-empty string literal expected before '\\)' token" } + #pragma omp interop init ( prefer_type( {fr("hip") , attr(omp_ifr_cuda) }) : obj1) ! { dg-error "expected string literal before 'omp_ifr_cuda'" } + + #pragma omp interop init ( prefer_type( {fr("hip") , attr("myooption") }) : obj1) // { dg-error "'attr' string literal must start with 'ompx_'" } + #pragma omp interop init ( prefer_type( {fr("hip") , attr("ompx_option") , attr("ompx_") } ) : obj1) + #pragma omp interop init ( prefer_type( {fr("hip") , attr("ompx_option") }, { attr("ompx_") } ) : obj1) + #pragma omp interop init ( prefer_type( {fr("hip") , attr("ompx_option") } { attr("ompx_") } ) : obj1) // { dg-error "expected '\\)' or ',' before '\{' token" } + #pragma omp interop init ( prefer_type( {fr("hip") , attr("ompx_option") ) : obj1) + // { dg-error "expected ',' or '\}' before '\\)' token" "" { target c } .-1 } + // { dg-error "prefer_type' has not been declared" "" { target c++ } .-2 } + // { dg-error "expected '\\)' before '\\(' token" "" { target c++ } .-3 } + + #pragma omp interop init ( prefer_type( {fr("hip") attr("ompx_option") ) : obj1) + // { dg-error "expected ',' or '\}' before 'attr'" "" { target c } .-1 } + // { dg-error "prefer_type' has not been declared" "" { target c++ } .-2 } + // { dg-error "expected '\\)' before '\\(' token" "" { target c++ } .-3 } + #pragma omp interop init ( prefer_type( {fr("hip")}), prefer_type("cuda") : obj1) // { dg-error "duplicate 'prefer_type' modifier" } + + #pragma omp interop init ( prefer_type( {attr("ompx_option1,ompx_option2") } ) : obj1) // { dg-error "'attr' string literal must not contain a comma" } + + #pragma omp interop init ( prefer_type( {attr("ompx_option1,ompx_option2") ) : obj1) + // { dg-error "'attr' string literal must not contain a comma" "" { target c } .-1 } + // { dg-error "prefer_type' has not been declared" "" { target c++ } .-2 } + // { dg-error "expected '\\)' before '\\(' token" "" { target c++ } .-3 } + + #pragma omp interop init ( targetsync other ) : obj1) + // { dg-error "'targetsync' undeclared \\(first use in this function\\)" "" { target c } .-1 } + // { dg-error "'targetsync' has not been declared" "" { target c++ } .-2 } + // { dg-error "expected '\\)' before 'other'" "" { target *-*-* } .-3 } + // { dg-error "expected an OpenMP clause before ':' token" "" { target *-*-* } .-4 } + + #pragma omp interop init ( prefer_type( {fr("cuda") } ), other : obj1) // { dg-error "'init' clause with modifier other than 'prefer_type', 'target' or 'targetsync' before 'other'" } + #pragma omp interop init ( prefer_type( {fr("cuda") } ), obj1) + // { dg-error "'prefer_type' undeclared \\(first use in this function\\)" "" { target c } .-1 } + // { dg-error "'prefer_type' has not been declared" "" { target c++ } .-2 } + // { dg-error "expected '\\)' before '\\(' token" "" { target *-*-* } .-3 } +} diff --git a/gcc/testsuite/c-c++-common/gomp/interop-3.c b/gcc/testsuite/c-c++-common/gomp/interop-3.c new file mode 100644 index 000000000000..42478bf760d6 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/interop-3.c @@ -0,0 +1,82 @@ +/* { dg-prune-output "sorry, unimplemented: '#pragma omp interop' not yet supported" } */ + +/* The following definitions are in omp_lib, which cannot be included + in gcc/testsuite/ */ + +#if __cplusplus >= 201103L +# define __GOMP_UINTPTR_T_ENUM : __UINTPTR_TYPE__ +#else +# define __GOMP_UINTPTR_T_ENUM +#endif + +typedef enum omp_interop_t __GOMP_UINTPTR_T_ENUM +{ + omp_interop_none = 0, + __omp_interop_t_max__ = __UINTPTR_MAX__ +} omp_interop_t; + +typedef enum omp_interop_fr_t +{ + omp_ifr_cuda = 1, + omp_ifr_cuda_driver = 2, + omp_ifr_opencl = 3, + omp_ifr_sycl = 4, + omp_ifr_hip = 5, + omp_ifr_level_zero = 6, + omp_ifr_hsa = 7, + omp_ifr_last = omp_ifr_hsa +} omp_interop_fr_t; + + +// --------------------------------- + +void f() +{ + omp_interop_t obj1, obj2, obj3, obj4, obj5; + omp_interop_t target, targetsync, prefer_type; + int x; + + #pragma omp interop init(obj1) init(target,targetsync : obj2, obj3) nowait + + #pragma omp interop init(prefer_type("cuda", omp_ifr_opencl, omp_ifr_level_zero, "hsa"), targetsync : obj1) \ + destroy(obj2, obj3) depend(inout: x) use(obj4, obj5) device(device_num: 0) + + #pragma omp assume contains(interop) + { + #pragma omp interop init(prefer_type("cu da") : obj3) // { dg-warning "unknown foreign runtime identifier 'cu da'" } + } + + #pragma omp interop init(obj1, obj2, obj1), use(obj4) destroy(obj4) + // { dg-error "'obj4' appears more than once in action clauses" "" { target *-*-* } .-1 } + // { dg-error "'obj1' appears more than once in action clauses" "" { target *-*-* } .-2 } + + #pragma omp interop depend(inout: x) // { dg-error "'depend' clause requires action clauses with 'targetsync' interop-type" } + + #pragma omp interop depend(inout: x) , use(obj2), destroy(obj3) // OK, use or destroy might have 'targetsync' + + #pragma omp interop depend(inout: x) use(obj2), destroy(obj3) // Likewise + + #pragma omp interop depend(inout: x) use(obj2), destroy(obj3) init(obj4) // { dg-error "'depend' clause requires action clauses with 'targetsync' interop-type" } + // { dg-note "69: 'init' clause lacks the 'targetsync' modifier" "" { target c } .-1 } + // { dg-note "70: 'init' clause lacks the 'targetsync' modifier" "" { target c++ } .-2 } + + #pragma omp interop depend(inout: x) init(targetsync : obj5) use(obj2), destroy(obj3) init(obj4) // { dg-error "'depend' clause requires action clauses with 'targetsync' interop-type" } + // { dg-note "'init' clause lacks the 'targetsync' modifier" "" { target *-*-* } .-1 } + #pragma omp interop depend(inout: x) init(targetsync : obj5) use(obj2), destroy(obj3) init(prefer_type("cuda"), targetsync : obj4) // OK + + #pragma omp interop init(target, targetsync, prefer_type, obj1) + #pragma omp interop init(prefer_type, obj1, target, targetsync) + +// Duplicated variable name or duplicated modifier: + #pragma omp interop init(target, targetsync,target : obj1) // { dg-error "duplicate 'target' modifier" } + #pragma omp interop init(target, targetsync,target) // { dg-error "'target' appears more than once in action clauses" } + #pragma omp interop init(target : target, targetsync,target) // { dg-error "'target' appears more than once in action clauses" } + + #pragma omp interop init(target, targetsync,targetsync : obj1) // { dg-error "duplicate 'targetsync' modifier" } + #pragma omp interop init(target, targetsync,targetsync) // { dg-error "targetsync' appears more than once in action clause" } + #pragma omp interop init(target : target, targetsync,targetsync) // { dg-error "targetsync' appears more than once in action clause" } + + #pragma omp interop init(, targetsync, prefer_type, obj1, target) + // { dg-error "expected identifier before ',' token" "" { target c } .-1 } + // { dg-error "expected unqualified-id before ',' token" "" { target c++ } .-2 } +} diff --git a/gcc/testsuite/c-c++-common/gomp/interop-4.c b/gcc/testsuite/c-c++-common/gomp/interop-4.c new file mode 100644 index 000000000000..1f9c987108b1 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/interop-4.c @@ -0,0 +1,75 @@ +/* { dg-additional-options "-fdump-tree-original" } */ + +/* The following definitions are in omp_lib, which cannot be included + in gcc/testsuite/ */ + +#if __cplusplus >= 201103L +# define __GOMP_UINTPTR_T_ENUM : __UINTPTR_TYPE__ +#else +# define __GOMP_UINTPTR_T_ENUM +#endif + +typedef enum omp_interop_t __GOMP_UINTPTR_T_ENUM +{ + omp_interop_none = 0, + __omp_interop_t_max__ = __UINTPTR_MAX__ +} omp_interop_t; + +typedef enum omp_interop_fr_t +{ + omp_ifr_cuda = 1, + omp_ifr_cuda_driver = 2, + omp_ifr_opencl = 3, + omp_ifr_sycl = 4, + omp_ifr_hip = 5, + omp_ifr_level_zero = 6, + omp_ifr_hsa = 7, + omp_ifr_last = omp_ifr_hsa +} omp_interop_fr_t; + +void +f() +{ + omp_interop_t obj1, obj2, obj3, obj4, obj5, obj6, obj7; + int x[6]; + + #pragma omp interop init ( obj1, obj2) use (obj3) destroy(obj4) init(obj5) destroy(obj6) use(obj7) /* { dg-message "'#pragma omp interop' not yet supported" } */ + /* { dg-final { scan-tree-dump-times "#pragma omp interop use\\(obj7\\) destroy\\(obj6\\) init\\(obj5\\) destroy\\(obj4\\) use\\(obj3\\) init\\(obj2\\) init\\(obj1\\)\[\r\n\]" 1 "original" } } */ + + #pragma omp interop nowait init (targetsync : obj1, obj2) use (obj3) destroy(obj4) init(target, targetsync : obj5) destroy(obj6) use(obj7) depend(inout: x) /* { dg-message "'#pragma omp interop' not yet supported" } */ + /* { dg-final { scan-tree-dump-times "#pragma omp interop depend\\(inout:x\\) use\\(obj7\\) destroy\\(obj6\\) init\\(target, targetsync: obj5\\) destroy\\(obj4\\) use\\(obj3\\) init\\(targetsync: obj2\\) init\\(targetsync: obj1\\) nowait\[\r\n\]" 1 "original" } } */ + + #pragma omp interop init ( obj1, obj2) init (target: obj3) init(targetsync : obj4) init(target,targetsync: obj5) /* { dg-message "'#pragma omp interop' not yet supported" } */ + /* { dg-final { scan-tree-dump-times "#pragma omp interop init\\(target, targetsync: obj5\\) init\\(targetsync: obj4\\) init\\(target: obj3\\) init\\(obj2\\) init\\(obj1\\)\[\r\n\]" 1 "original" } } */ + + /* -------------------------------------------- */ + + #pragma omp interop init (target, prefer_type(omp_ifr_cuda, omp_ifr_cuda+1, "hsa", "myPrivateInterop", omp_ifr_cuda-2) : obj1, obj2) init (target: obj3) init(prefer_type(omp_ifr_hip, "sycl", omp_ifr_opencl), targetsync : obj4, obj7) init(target,prefer_type("level_zero", omp_ifr_level_zero+0),targetsync: obj5) /* { dg-message "'#pragma omp interop' not yet supported" } */ + /* + { dg-warning "unknown foreign runtime identifier 'myPrivateInterop' \\\[-Wopenmp\\\]" "" { target *-*-* } .-2 } + { dg-warning "unknown foreign runtime identifier '-1' \\\[-Wopenmp\\\]" "" { target *-*-* } .-3 } + + { dg!final { scan-tree-dump-times "#pragma omp interop init\\(prefer_type\\({fr\\(\"level_zero\"\\)}, {fr\\(\"level_zero\"\\)}\\), target, targetsync: obj5\\) init\\(prefer_type\\({fr\\(\"hip\"\\)}, {fr\\(\"sycl\"\\)}, {fr\\(\"opencl\"\\)}\\), targetsync: obj7\\) init\\(prefer_type\\({fr\\(\"hip\"\\)}, {fr\\(\"sycl\"\\)}, {fr\\(\"opencl\"\\)}\\), targetsync: obj4\\) init\\(target: obj3\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\)}, {fr\\(\"hsa\"\\)}, {fr\\(\"<unknown>\"\\)}, {fr\\(\"<unknown>\"\\)}\\), target: obj2\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\)}, {fr\\(\"hsa\"\\)}, {fr\\(\"<unknown>\"\\)}, {fr\\(\"<unknown>\"\\)}\\), target: obj1\\)\[\r\n\]" 1 "original" } } + */ + + +/* -------------------------------------------- */ + + #pragma omp interop init ( target, prefer_type( {fr("hip"), attr("ompx_gnu_prio:1", "ompx_gnu_debug")}, {attr("ompx_gnu_nicest"), attr("ompx_something")}) : obj1, obj2) init ( prefer_type( {fr("cuda")}, {fr(omp_ifr_cuda_driver), attr("ompx_nix")}, {fr("best")}), targetsync : obj3, obj4) nowait use(obj5) /* { dg-message "'#pragma omp interop' not yet supported" } */ + /* + { dg-warning "unknown foreign runtime identifier 'best' \\\[-Wopenmp\\\]" "" { target *-*-* } .-2 } + + { dg-final { scan-tree-dump-times "#pragma omp interop use\\(obj5\\) nowait init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\),attr\\(\"ompx_nix\"\\)}, {fr\\(\"<unknown>\"\\)}\\), targetsync: obj4\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\),attr\\(\"ompx_nix\"\\)}, {fr\\(\"<unknown>\"\\)}\\), targetsync: obj3\\) init\\(prefer_type\\({fr\\(\"hip\"\\),attr\\(\"ompx_gnu_prio:1\"\\),attr\\(\"ompx_gnu_debug\"\\)}, {attr\\(\"ompx_gnu_nicest\"\\),attr\\(\"ompx_something\"\\)}\\), target: obj2\\) init\\(prefer_type\\({fr\\(\"hip\"\\),attr\\(\"ompx_gnu_prio:1\"\\),attr\\(\"ompx_gnu_debug\"\\)}, {attr\\(\"ompx_gnu_nicest\"\\),attr\\(\"ompx_something\"\\)}\\), target: obj1\\)\[\r\n\]" 1 "original" } } + */ + +} + +void +g (int *y) +{ + omp_interop_t io1, io2, io3, io4, io5; + + [[omp::directive (interop,init(prefer_type({fr("level_zero")}, {fr(omp_ifr_sycl),attr("ompx_in_order"),attr("ompx_queue:in_order")}), targetsync : io1, io2),use(io3),destroy(io4,io5),depend(inout:y),nowait)]]; /* { dg-message "'#pragma omp interop' not yet supported" } */ + + /* { dg-final { scan-tree-dump-times "#pragma omp interop nowait depend\\(inout:y\\) destroy\\(io5\\) destroy\\(io4\\) use\\(io3\\) init\\(prefer_type\\(\{fr\\(\"level_zero\"\\)\}, \{fr\\(\"sycl\"\\),attr\\(\"ompx_in_order\"\\),attr\\(\"ompx_queue:in_order\"\\)\}\\), targetsync: io2\\) init\\(prefer_type\\(\{fr\\(\"level_zero\"\\)\}, \{fr\\(\"sycl\"\\),attr\\(\"ompx_in_order\"\\),attr\\(\"ompx_queue:in_order\"\\)\}\\), targetsync: io1\\)\[\r\n\]" 1 "original" } } */ +} diff --git a/gcc/testsuite/g++.dg/gomp/interop-5.C b/gcc/testsuite/g++.dg/gomp/interop-5.C new file mode 100644 index 000000000000..5109dc4e4271 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/interop-5.C @@ -0,0 +1,90 @@ +/* { dg-do compile { target c++11 } } */ +/* { dg-additional-options "-fdump-tree-original" } */ + +/* { dg-prune-output "sorry, unimplemented: '#pragma omp interop' not yet supported" } */ + +/* The following definitions are in omp_lib, which cannot be included + in gcc/testsuite/ */ + +#if __cplusplus >= 201103L +# define __GOMP_UINTPTR_T_ENUM : __UINTPTR_TYPE__ +#else +# define __GOMP_UINTPTR_T_ENUM +#endif + +typedef enum omp_interop_t __GOMP_UINTPTR_T_ENUM +{ + omp_interop_none = 0, + __omp_interop_t_max__ = __UINTPTR_MAX__ +} omp_interop_t; + +typedef enum omp_interop_fr_t +{ + omp_ifr_cuda = 1, + omp_ifr_cuda_driver = 2, + omp_ifr_opencl = 3, + omp_ifr_sycl = 4, + omp_ifr_hip = 5, + omp_ifr_level_zero = 6, + omp_ifr_hsa = 7, + omp_ifr_last = omp_ifr_hsa +} omp_interop_fr_t; + +template<typename T, typename T2, typename T3> +void +f () +{ + T obj1, obj2, obj3, obj4, obj5, obj6, obj7; + T2 x[6]; + constexpr T3 ifr_hip = omp_ifr_hip; + constexpr T3 ifr_cuda = omp_ifr_cuda; + constexpr T3 ifr_cuda_driver = omp_ifr_cuda_driver; + constexpr T3 ifr_opencl = omp_ifr_opencl; + constexpr T3 ifr_level_zero = (T3) (omp_ifr_sycl + 2); + constexpr T3 ifr_invalid = (T3) 99; + + #pragma omp interop init ( obj1, obj2) use (obj3) destroy(obj4) init(obj5) destroy(obj6) use(obj7) /* { dg-message "'#pragma omp interop' not yet supported" } */ + /* { dg-final { scan-tree-dump-times "#pragma omp interop use\\(obj7\\) destroy\\(obj6\\) init\\(obj5\\) destroy\\(obj4\\) use\\(obj3\\) init\\(obj2\\) init\\(obj1\\)\[\r\n\]" 2 "original" } } */ + + #pragma omp interop nowait init (targetsync : obj1, obj2) use (obj3) destroy(obj4) init(target, targetsync : obj5) destroy(obj6) use(obj7) depend(inout: x) /* { dg-message "'#pragma omp interop' not yet supported" } */ + /* { dg-final { scan-tree-dump-times "#pragma omp interop depend\\(inout:x\\) use\\(obj7\\) destroy\\(obj6\\) init\\(target, targetsync: obj5\\) destroy\\(obj4\\) use\\(obj3\\) init\\(targetsync: obj2\\) init\\(targetsync: obj1\\) nowait\[\r\n\]" 2 "original" } } */ + + #pragma omp interop init ( obj1, obj2) init (target: obj3) init(targetsync : obj4) init(target,targetsync: obj5) /* { dg-message "'#pragma omp interop' not yet supported" } */ + /* { dg-final { scan-tree-dump-times "#pragma omp interop init\\(target, targetsync: obj5\\) init\\(targetsync: obj4\\) init\\(target: obj3\\) init\\(obj2\\) init\\(obj1\\)\[\r\n\]" 2 "original" } } */ + + /* -------------------------------------------- */ + + #pragma omp interop init (target, prefer_type(ifr_invalid, 123, ifr_cuda, ifr_cuda+1, "hsa", "myPrivateInterop", ifr_cuda-2) : obj1, obj2) + /* + { dg-warning "49: unknown foreign runtime identifier '99' \\\[-Wopenmp\\\]" "" { target *-*-* } .-2 } + { dg-warning "62: unknown foreign runtime identifier '123' \\\[-Wopenmp\\\]" "" { target *-*-* } .-3 } + { dg-warning "96: unknown foreign runtime identifier 'myPrivateInterop' \\\[-Wopenmp\\\]" "" { target *-*-* } .-4 } + { dg-warning "124: unknown foreign runtime identifier '-1' \\\[-Wopenmp\\\]" "" { target *-*-* } .-5 } + + { dg-final { scan-tree-dump-times "#pragma omp interop init\\(prefer_type\\(\{fr\\(\"<unknown>\"\\)\}, \{fr\\(\"<unknown>\"\\)\}, \{fr\\(\"cuda\"\\)\}, \{fr\\(\"cuda_driver\"\\)\}, \{fr\\(\"hsa\"\\)\}, \{fr\\(\"<unknown>\"\\)\}, \{fr\\(\"<unknown>\"\\)\}\\), target: obj2\\) init\\(prefer_type\\(\{fr\\(\"<unknown>\"\\)\}, \{fr\\(\"<unknown>\"\\)\}, \{fr\\(\"cuda\"\\)\}, \{fr\\(\"cuda_driver\"\\)\}, \{fr\\(\"hsa\"\\)\}, \{fr\\(\"<unknown>\"\\)\}, \{fr\\(\"<unknown>\"\\)\}\\), target: obj1\\)\[\r\n\]" 2 "original" } } + */ + + #pragma omp interop init (target, prefer_type(ifr_cuda, ifr_cuda+1, "hsa", "myPrivateInterop", ifr_cuda-2) : obj1, obj2) init (target: obj3) init(prefer_type(ifr_hip, "sycl", ifr_opencl), targetsync : obj4, obj7) init(target,prefer_type("level_zero", ifr_level_zero+0),targetsync: obj5) /* { dg-message "'#pragma omp interop' not yet supported" } */ + /* + { dg-warning "unknown foreign runtime identifier 'myPrivateInterop' \\\[-Wopenmp\\\]" "" { target *-*-* } .-2 } + { dg-warning "unknown foreign runtime identifier '-1' \\\[-Wopenmp\\\]" "" { target *-*-* } .-3 } + + { dg-final { scan-tree-dump-times "#pragma omp interop init\\(prefer_type\\({fr\\(\"level_zero\"\\)}, {fr\\(\"level_zero\"\\)}\\), target, targetsync: obj5\\) init\\(prefer_type\\({fr\\(\"hip\"\\)}, {fr\\(\"sycl\"\\)}, {fr\\(\"opencl\"\\)}\\), targetsync: obj7\\) init\\(prefer_type\\({fr\\(\"hip\"\\)}, {fr\\(\"sycl\"\\)}, {fr\\(\"opencl\"\\)}\\), targetsync: obj4\\) init\\(target: obj3\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\)}, {fr\\(\"hsa\"\\)}, {fr\\(\"<unknown>\"\\)}, {fr\\(\"<unknown>\"\\)}\\), target: obj2\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\)}, {fr\\(\"hsa\"\\)}, {fr\\(\"<unknown>\"\\)}, {fr\\(\"<unknown>\"\\)}\\), target: obj1\\)\[\r\n\]" 2 "original" } } + */ + +/* -------------------------------------------- */ + + #pragma omp interop init ( target, prefer_type( {fr("hip"), attr("ompx_gnu_prio:1", "ompx_gnu_debug")}, {attr("ompx_gnu_nicest"), attr("ompx_something")}) : obj1, obj2) init ( prefer_type( {fr("cuda")}, {fr(ifr_cuda_driver), attr("ompx_nix")}, {fr("best")}), targetsync : obj3, obj4) nowait use(obj5) /* { dg-message "'#pragma omp interop' not yet supported" } */ + /* + { dg-warning "unknown foreign runtime identifier 'best' \\\[-Wopenmp\\\]" "" { target *-*-* } .-2 } + + { dg-final { scan-tree-dump-times "#pragma omp interop use\\(obj5\\) nowait init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\),attr\\(\"ompx_nix\"\\)}, {fr\\(\"<unknown>\"\\)}\\), targetsync: obj4\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\),attr\\(\"ompx_nix\"\\)}, {fr\\(\"<unknown>\"\\)}\\), targetsync: obj3\\) init\\(prefer_type\\({fr\\(\"hip\"\\),attr\\(\"ompx_gnu_prio:1\"\\),attr\\(\"ompx_gnu_debug\"\\)}, {attr\\(\"ompx_gnu_nicest\"\\),attr\\(\"ompx_something\"\\)}\\), target: obj2\\) init\\(prefer_type\\({fr\\(\"hip\"\\),attr\\(\"ompx_gnu_prio:1\"\\),attr\\(\"ompx_gnu_debug\"\\)}, {attr\\(\"ompx_gnu_nicest\"\\),attr\\(\"ompx_something\"\\)}\\), target: obj1\\)\[\r\n\]" 2 "original" } } + */ +} + +void +g (int *y) +{ + f<omp_interop_t, int, omp_interop_fr_t> (); + f<omp_interop_t, int, int> (); +} diff --git a/gcc/testsuite/gfortran.dg/gomp/interop-1.f90 b/gcc/testsuite/gfortran.dg/gomp/interop-1.f90 index b7d2164812cc..a16c3845afdd 100644 --- a/gcc/testsuite/gfortran.dg/gomp/interop-1.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/interop-1.f90 @@ -33,18 +33,27 @@ integer(omp_interop_fr_kind), parameter :: ifr_array(2) = [omp_ifr_cuda, omp_ifr integer(omp_interop_kind) :: obj1, obj2, obj3, obj4, obj5 integer :: x -!$omp interop init(obj1) init(target,targetsync,target,targetsync : obj2, obj3) nowait +!$omp interop init(obj1) init(target,targetsync : obj2, obj3) nowait ! OK +!$omp interop init(obj1) init (targetsync : obj2, obj3) nowait ! OK +!$omp interop init(obj1) init (targetsync , target : obj2, obj3) nowait ! OK -!$omp interop init(prefer_type("cu"//"da", omp_ifr_opencl, omp_ifr_level_zero, "hsa"), targetsync : obj1) & +!$omp interop init(obj1) init(target,targetsync,target: obj2, obj3) nowait ! { dg-error "Duplicate 'target'" } +!$omp interop init(obj1) init(target,targetsync, targetsync : obj2, obj3) nowait ! { dg-error "Duplicate 'targetsync'" } + +!$omp interop init(prefer_type("cuda", omp_ifr_opencl, omp_ifr_level_zero, "hsa"), targetsync : obj1) & !$omp& destroy(obj2, obj3) depend(inout: x) use(obj4, obj5) device(device_num: 0) +!$omp interop init(prefer_type("cu" // "da"), targetsync : obj1) ! { dg-error "37: Expected ',' or '\\)'" } +! { dg-warning "Unknown foreign runtime identifier 'cu' at \\(1\\) \\\[-Wopenmp\\\]" "" { target *-*-* } .-1 } + !$omp assume contains(interop) - !$omp interop init(prefer_type("cu"//char(1)//"da") : obj3) ! { dg-warning "Unknown foreign runtime identifier 'cu\\\\x01da'" } + !$omp interop init(prefer_type("cuða") : obj3) ! { dg-warning "Unknown foreign runtime identifier 'cu\[^'\]*a'" } !$omp end assume -!$omp interop init(prefer_type("cu"//char(0)//"da") : obj3) ! { dg-error "Unexpected null character in character literal" } +!$omp interop init(prefer_type("cu"//char(0)//"da") : obj3) ! { dg-error "36: Expected ',' or '\\)'" } +! { dg-warning "Unknown foreign runtime identifier 'cu' at \\(1\\) \\\[-Wopenmp\\\]" "" { target *-*-* } .-1 } -!$omp interop depend(inout: x) , use(obj2), destroy(obj3) ! OK, use or destory might have 'targetsync' +!$omp interop depend(inout: x) , use(obj2), destroy(obj3) ! OK, use or destroy might have 'targetsync' !$omp interop depend(inout: x) use(obj2), destroy(obj3) ! Likewise @@ -56,15 +65,19 @@ integer :: x !$omp interop init ( prefer_type( 4, omp_ifr_hip*4) : obj1) ! { dg-warning "Unknown foreign runtime identifier '20'" } !$omp interop init ( prefer_type( sin(3.3) : obj1) ! { dg-error "Expected constant scalar integer expression or non-empty default-kind character literal" } -!$omp interop init ( prefer_type( {fr(4 ) }) : obj1) ! { dg-error "Expected scalar integer parameter or non-empty default-kind character literal" } -!$omp interop init ( prefer_type( {fr(4_"cuda" ) }) : obj1) ! { dg-error "Expected scalar integer parameter or non-empty default-kind character literal" } +!$omp interop init ( prefer_type( {fr(4 ) }) : obj1) ! OK +!$omp interop init ( prefer_type( {fr(4_"cuda" ) }) : obj1) ! { dg-error "Expected constant scalar integer expression or non-empty default-kind character literal" } !$omp interop init ( prefer_type( {fr(c_char_"cuda") }) : obj1) ! OK !$omp interop init ( prefer_type( {fr(1_"cuda" ) }) : obj1) ! OK !$omp interop init ( prefer_type( {fr(omp_ifr_level_zero ) }, {fr(omp_ifr_hip)}) : obj1) ! OK -!$omp interop init ( prefer_type( {fr(omp_ifr_level_zero + 1) }) : obj1) ! { dg-error "Expected scalar integer parameter or non-empty default-kind character literal" } -!$omp interop init ( prefer_type( {fr(x) }) : obj1) ! { dg-error "Expected scalar integer parameter or non-empty default-kind character literal" } -!$omp interop init ( prefer_type( {fr(ifr_array ) }) : obj1) ! { dg-error "Expected scalar integer parameter or non-empty default-kind character literal" } -!$omp interop init ( prefer_type( {fr(ifr_array(1) ) }) : obj1) ! { dg-error "Expected scalar integer parameter or non-empty default-kind character literal" } +!$omp interop init ( prefer_type( {fr("cuda" // "_driver") }) : obj1) ! { dg-error "46: Expected '\\)'" } +!$omp interop init ( prefer_type( {fr(trim("cuda" // "_driver")) }) : obj1) ! { dg-error "38: Expected constant scalar integer expression or non-empty default-kind character literal" } +!$omp interop init ( prefer_type( {fr("hello" }) : obj1) ! { dg-error "47: Expected '\\)'" } +! { dg-warning "Unknown foreign runtime identifier 'hello' at \\(1\\) \\\[-Wopenmp\\\]" "" { target *-*-* } .-1 } + +!$omp interop init ( prefer_type( {fr(x) }) : obj1) ! { dg-error "Expected constant scalar integer expression or non-empty default-kind character literal" } +!$omp interop init ( prefer_type( {fr(ifr_array ) }) : obj1) ! { dg-error "Expected constant scalar integer expression or non-empty default-kind character literal" } +!$omp interop init ( prefer_type( {fr(ifr_array(1) ) }) : obj1) !$omp interop init ( prefer_type( omp_ifr_level_zero, omp_ifr_hip ) : obj1) ! OK !$omp interop init ( prefer_type( omp_ifr_level_zero +1 ) : obj1) ! OK @@ -75,8 +88,10 @@ integer :: x !$omp interop init ( prefer_type( 4, omp_ifr_hip*4) : obj1) ! { dg-warning "Unknown foreign runtime identifier '20'" } !$omp interop init ( prefer_type( 4, 1, 3) : obj1) -!$omp interop init ( prefer_type( {fr("cuda","sycl") }, {fr(omp_ifr_hsa,omp_ifr_level_zero)} , {attr("ompx_a") } , {fr(omp_ifr_hip) }) : obj1) -!$omp interop init ( prefer_type( {fr("cuda","sycl"), attr("ompx_1", "ompx_2"), attr("ompx_3") }, {attr("ompx_4", "ompx_5"),fr(omp_ifr_hsa,omp_ifr_level_zero)} ) : obj1) -!$omp interop init ( prefer_type( { fr("cuda","sycl"), attr("ompx_1") }, {fr(omp_ifr_hsa,omp_ifr_level_zero)} , {attr("ompx_a") } ) : obj1) +!$omp interop init ( prefer_type( {fr("cuda") }, {fr(omp_ifr_hsa)} , {attr("ompx_a") } , {fr(omp_ifr_hip) }) : obj1) +!$omp interop init ( prefer_type( {fr("cuda") }, {fr(omp_ifr_hsa,omp_ifr_level_zero)} , {attr("ompx_a") } , {fr(omp_ifr_hip) }) : obj1) ! { dg-error "65: Expected '\\)'" } +!$omp interop init ( prefer_type( {fr("cuda",5) }, {fr(omp_ifr_hsa,omp_ifr_level_zero)} , {attr("ompx_a") } , {fr(omp_ifr_hip) }) : obj1) ! { dg-error "45: Expected '\\)' at" } +!$omp interop init ( prefer_type( {fr("sycl"), attr("ompx_1", "ompx_2"), attr("ompx_3") }, {attr("ompx_4", "ompx_5"),fr(omp_ifr_level_zero)} ) : obj1) +!$omp interop init ( prefer_type( { fr(5), attr("ompx_1") }, {fr(omp_ifr_hsa)} , {attr("ompx_a") } ) : obj1) end diff --git a/gcc/testsuite/gfortran.dg/gomp/interop-2.f90 b/gcc/testsuite/gfortran.dg/gomp/interop-2.f90 index f3391bf88f0b..b3130117fb2b 100644 --- a/gcc/testsuite/gfortran.dg/gomp/interop-2.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/interop-2.f90 @@ -17,16 +17,50 @@ module m integer (omp_interop_fr_kind), parameter :: omp_ifr_hsa = 7 end module m -program main +subroutine s(ointent) use m implicit none +integer(omp_interop_kind), parameter :: op = 0 +integer(omp_interop_kind),intent(in) :: ointent +integer(omp_interop_kind) :: od(5) +integer(1) :: o1 +integer, parameter :: mykind = mod (omp_interop_kind, 100) ! remove saving the 'comes from c_int' info +real(mykind) :: or + +!$omp interop init (op) ! { dg-error "'op' at \\(1\\) in 'INIT' clause must be a scalar integer variable of 'omp_interop_kind' kind" } + ! { dg-error "Object 'op' is not a variable at \\(1\\)" "" { target *-*-* } .-1 } +!$omp interop init (ointent) ! { dg-error "'ointent' at \\(1\\) in 'INIT' clause must be definable" } +!$omp interop init (od) ! { dg-error "'od' at \\(1\\) in 'INIT' clause must be a scalar integer variable of 'omp_interop_kind' kind" } +!$omp interop init (od(1)) ! { dg-error "Syntax error in OpenMP variable list" } +!$omp interop init (o1) ! { dg-error "'o1' at \\(1\\) in 'INIT' clause must be a scalar integer variable of 'omp_interop_kind' kind" } +!$omp interop init (or) ! { dg-error "'or' at \\(1\\) in 'INIT' clause must be a scalar integer variable of 'omp_interop_kind' kind" } + +!$omp interop use (op) ! { dg-error "'op' at \\(1\\) in 'USE' clause must be a scalar integer variable of 'omp_interop_kind' kind" } + ! { dg-error "Object 'op' is not a variable at \\(1\\)" "" { target *-*-* } .-1 } +!$omp interop use (ointent) ! okay +!$omp interop use (od) ! { dg-error "'od' at \\(1\\) in 'USE' clause must be a scalar integer variable of 'omp_interop_kind' kind" } +!$omp interop use (od(1)) ! { dg-error "Syntax error in OpenMP variable list" } +!$omp interop use (o1) ! { dg-error "'o1' at \\(1\\) in 'USE' clause must be a scalar integer variable of 'omp_interop_kind' kind" } +!$omp interop use (or) ! { dg-error "'or' at \\(1\\) in 'USE' clause must be a scalar integer variable of 'omp_interop_kind' kind" } + +!$omp interop destroy (op) ! { dg-error "'op' at \\(1\\) in 'DESTROY' clause must be a scalar integer variable of 'omp_interop_kind' kind" } + ! { dg-error "Object 'op' is not a variable at \\(1\\)" "" { target *-*-* } .-1 } +!$omp interop destroy (ointent) ! { dg-error "'ointent' at \\(1\\) in 'DESTROY' clause must be definable" } +!$omp interop destroy (od) ! { dg-error "'od' at \\(1\\) in 'DESTROY' clause must be a scalar integer variable of 'omp_interop_kind' kind" } +!$omp interop destroy (od(1)) ! { dg-error "Syntax error in OpenMP variable list" } +!$omp interop destroy (o1) ! { dg-error "'o1' at \\(1\\) in 'DESTROY' clause must be a scalar integer variable of 'omp_interop_kind' kind" } +!$omp interop destroy (or) ! { dg-error "'or' at \\(1\\) in 'DESTROY' clause must be a scalar integer variable of 'omp_interop_kind' kind" } -!$omp requires reverse_offload +end subroutine + +program main +use m +implicit none integer(omp_interop_kind) :: obj1, obj2, obj3, obj4, obj5 integer :: x -!$omp interop init ( prefer_type( {fr(1_"") }) : obj1) ! { dg-error "Expected scalar integer parameter or non-empty default-kind character literal" } +!$omp interop init ( prefer_type( {fr(1_"") }) : obj1) ! { dg-error "Expected constant scalar integer expression or non-empty default-kind character literal" } !$omp interop init ( prefer_type( {fr(1_"hip") , attr(omp_ifr_cuda) }) : obj1) ! { dg-error "Expected default-kind character literal" } !$omp interop init ( prefer_type( {fr(1_"hip") , attr("myooption") }) : obj1) ! { dg-error "Character literal at .1. must start with 'ompx_'" } diff --git a/gcc/testsuite/gfortran.dg/gomp/interop-3.f90 b/gcc/testsuite/gfortran.dg/gomp/interop-3.f90 index 462ed4f2e4bc..a3bbfcaf2ca5 100644 --- a/gcc/testsuite/gfortran.dg/gomp/interop-3.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/interop-3.f90 @@ -21,19 +21,17 @@ program main use m implicit none -!$omp requires reverse_offload - integer(omp_interop_kind) :: obj1, obj2, obj3, obj4, obj5 integer(omp_interop_kind) :: target, targetsync,prefer_type integer :: x -!$omp interop init(obj1) init(target,targetsync,target,targetsync : obj2, obj3) nowait +!$omp interop init(obj1) init(target,targetsync : obj2, obj3) nowait -!$omp interop init(prefer_type("cu"//"da", omp_ifr_opencl, omp_ifr_level_zero, "hsa"), targetsync : obj1) & +!$omp interop init(prefer_type(1_"cuda", omp_ifr_opencl, omp_ifr_level_zero, "hsa"), targetsync : obj1) & !$omp& destroy(obj2, obj3) depend(inout: x) use(obj4, obj5) device(device_num: 0) !$omp assume contains(interop) - !$omp interop init(prefer_type("cu"//char(1)//"da") : obj3) ! { dg-warning "Unknown foreign runtime identifier 'cu\\\\x01da'" } + !$omp interop init(prefer_type("cu da") : obj3) ! { dg-warning "Unknown foreign runtime identifier 'cu da'" } !$omp end assume !$omp interop init(obj1, obj2, obj1), use(obj4) destroy(obj4) @@ -42,7 +40,7 @@ integer :: x !$omp interop depend(inout: x) ! { dg-error "DEPEND clause at .1. requires action clause with 'targetsync' interop-type" } -!$omp interop depend(inout: x) , use(obj2), destroy(obj3) ! OK, use or destory might have 'targetsync' +!$omp interop depend(inout: x) , use(obj2), destroy(obj3) ! OK, use or destroy might have 'targetsync' !$omp interop depend(inout: x) use(obj2), destroy(obj3) ! Likewise @@ -53,7 +51,16 @@ integer :: x !$omp interop init(target, targetsync, prefer_type, obj1) !$omp interop init(prefer_type, obj1, target, targetsync) -!$omp interop init(target, targetsync,target) ! { dg-error "Symbol 'target' present on multiple clauses" } + +! Duplicated variable name or duplicated modifier: +!$omp interop init(target, targetsync,target : obj1) ! { dg-error "Duplicate 'target' at \\(1\\)" } +!$omp interop init(target, targetsync,target) ! { dg-error "Duplicate 'target' at \\(1\\)" } +!$omp interop init(target : target, targetsync,target) ! { dg-error "Symbol 'target' present on multiple clauses" } + +!$omp interop init(target, targetsync,targetsync : obj1) ! { dg-error "Duplicate 'targetsync' at \\(1\\)" } +!$omp interop init(target, targetsync,targetsync) ! { dg-error "Duplicate 'targetsync' at \\(1\\)" } +!$omp interop init(target : target, targetsync,targetsync) ! { dg-error "Symbol 'targetsync' present on multiple clauses" } + !$omp interop init(, targetsync, prefer_type, obj1, target) ! { dg-error "Syntax error in OpenMP variable list" } end diff --git a/gcc/testsuite/gfortran.dg/gomp/interop-4.f90 b/gcc/testsuite/gfortran.dg/gomp/interop-4.f90 new file mode 100644 index 000000000000..8783f4cfb5fd --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/interop-4.f90 @@ -0,0 +1,56 @@ +! { dg-additional-options "-fdump-tree-original" } + +module m + use iso_c_binding + implicit none + + ! The following definitions are in omp_lib, which cannot be included + ! in gcc/testsuite/ + integer, parameter :: omp_interop_kind = c_intptr_t + integer, parameter :: omp_interop_fr_kind = c_int + + integer (omp_interop_kind), parameter :: omp_interop_none = 0_omp_interop_kind + integer (omp_interop_fr_kind), parameter :: omp_ifr_cuda = 1 + integer (omp_interop_fr_kind), parameter :: omp_ifr_cuda_driver = 2 + integer (omp_interop_fr_kind), parameter :: omp_ifr_opencl = 3 + integer (omp_interop_fr_kind), parameter :: omp_ifr_sycl = 4 + integer (omp_interop_fr_kind), parameter :: omp_ifr_hip = 5 + integer (omp_interop_fr_kind), parameter :: omp_ifr_level_zero = 6 + integer (omp_interop_fr_kind), parameter :: omp_ifr_hsa = 7 +end module m + +subroutine s +use m +implicit none + +integer(omp_interop_kind) :: obj1, obj2, obj3, obj4, obj5, obj6, obj7 +integer :: x(6) + +!$omp interop init ( obj1, obj2) use (obj3) destroy(obj4) init(obj5) destroy(obj6) use(obj7) ! { dg-message "'#pragma omp interop' not yet supported" } +! { dg-final { scan-tree-dump-times "#pragma omp interop init\\(obj1\\) init\\(obj2\\) init\\(obj5\\) use\\(obj3\\) use\\(obj7\\) destroy\\(obj4\\) destroy\\(obj6\\)\[\r\n\]" 1 "original" } } + +!$omp interop nowait init (targetsync : obj1, obj2) use (obj3) destroy(obj4) init(target, targetsync : obj5) destroy(obj6) use(obj7) depend(inout: x) ! { dg-message "'#pragma omp interop' not yet supported" } +! { dg-final { scan-tree-dump-times "#pragma omp interop depend\\(inout:x\\) init\\(targetsync: obj1\\) init\\(targetsync: obj2\\) init\\(target, targetsync: obj5\\) use\\(obj3\\) use\\(obj7\\) destroy\\(obj4\\) destroy\\(obj6\\) nowait\[\r\n\]" 1 "original" } } + +!$omp interop init ( obj1, obj2) init (target: obj3) init(targetsync : obj4) init(target,targetsync: obj5) ! { dg-message "'#pragma omp interop' not yet supported" } +! { dg-final { scan-tree-dump-times "#pragma omp interop init\\(obj1\\) init\\(obj2\\) init\\(target: obj3\\) init\\(targetsync: obj4\\) init\\(target, targetsync: obj5\\)\[\r\n\]" 1 "original" } } + +! -------------------------------------------- + +!$omp interop init (target, prefer_type(omp_ifr_cuda, omp_ifr_cuda+1, "hsa", "myPrivateInterop", omp_ifr_cuda-2) : obj1, obj2) init (target: obj3) init(prefer_type(omp_ifr_hip, "sycl", omp_ifr_opencl), targetsync : obj4, obj7) init(target,prefer_type("level_zero", omp_ifr_level_zero+0),targetsync: obj5) ! { dg-message "'#pragma omp interop' not yet supported" } +! +! { dg-warning "Unknown foreign runtime identifier 'myPrivateInterop' at \\(1\\) \\\[-Wopenmp\\\]" "" { target *-*-* } .-2 } +! { dg-warning "Unknown foreign runtime identifier '-1' at \\(1\\) \\\[-Wopenmp\\\]" "" { target *-*-* } .-3 } +! +! { dg-final { scan-tree-dump-times "#pragma omp interop init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\)}, {fr\\(\"hsa\"\\)}, {fr\\(\"<unknown>\"\\)}, {fr\\(\"<unknown>\"\\)}\\), target: obj1\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\)}, {fr\\(\"hsa\"\\)}, {fr\\(\"<unknown>\"\\)}, {fr\\(\"<unknown>\"\\)}\\), target: obj2\\) init\\(target: obj3\\) init\\(prefer_type\\({fr\\(\"hip\"\\)}, {fr\\(\"sycl\"\\)}, {fr\\(\"opencl\"\\)}\\), targetsync: obj4\\) init\\(prefer_type\\({fr\\(\"hip\"\\)}, {fr\\(\"sycl\"\\)}, {fr\\(\"opencl\"\\)}\\), targetsync: obj7\\) init\\(prefer_type\\({fr\\(\"level_zero\"\\)}, {fr\\(\"level_zero\"\\)}\\), target, targetsync: obj5\\)\[\r\n\]" 1 "original" } } + + +! -------------------------------------------- + +!$omp interop init ( target, prefer_type( {fr(1_"hip"), attr("ompx_gnu_prio:1", 1_"ompx_gnu_debug")}, {attr("ompx_gnu_nicest"), attr("ompx_something")}) : obj1, obj2) init ( prefer_type( {fr("cuda")}, {fr(omp_ifr_cuda_driver), attr("ompx_nix")}, {fr("best")}), targetsync : obj3, obj4) nowait use(obj5) ! { dg-message "'#pragma omp interop' not yet supported" } +! +! ! { dg-warning "Unknown foreign runtime identifier 'best' at \\(1\\) \\\[-Wopenmp\\\]" "" { target *-*-* } .-2 } +! +! { dg-final { scan-tree-dump-times "#pragma omp interop init\\(prefer_type\\({fr\\(\"hip\"\\),attr\\(\"ompx_gnu_prio:1\"\\),attr\\(\"ompx_gnu_debug\"\\)}, {attr\\(\"ompx_gnu_nicest\"\\),attr\\(\"ompx_something\"\\)}\\), target: obj1\\) init\\(prefer_type\\({fr\\(\"hip\"\\),attr\\(\"ompx_gnu_prio:1\"\\),attr\\(\"ompx_gnu_debug\"\\)}, {attr\\(\"ompx_gnu_nicest\"\\),attr\\(\"ompx_something\"\\)}\\), target: obj2\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\),attr\\(\"ompx_nix\"\\)}, {fr\\(\"<unknown>\"\\)}\\), targetsync: obj3\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\),attr\\(\"ompx_nix\"\\)}, {fr\\(\"<unknown>\"\\)}\\), targetsync: obj4\\) use\\(obj5\\) nowait\[\r\n\]" 1 "original" } } + +end diff --git a/gcc/tree-core.h b/gcc/tree-core.h index d205722cff3d..e94f0d2896bd 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -380,6 +380,19 @@ enum omp_clause_code { /* Range END above for: OMP_CLAUSE_SIZE */ + /* OpenMP clause: destroy (variable-list ). */ + OMP_CLAUSE_DESTROY, + + /* Range START below for: OMP_CLAUSE_INIT_PREFER_TYPE */ + + /* OpenMP clause: init ( [modifier-list : ] variable-list ). */ + OMP_CLAUSE_INIT, + + /* Range END above for: OMP_CLAUSE_INIT_PREFER_TYPE */ + + /* OpenMP clause: use (variable-list ). */ + OMP_CLAUSE_USE, + /* OpenACC clause: gang [(gang-argument-list)]. Where gang-argument-list: [gang-argument-list, ] gang-argument diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index d994faead42c..3bb684609ba7 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -451,6 +451,49 @@ dump_omp_iterators (pretty_printer *pp, tree iter, int spc, dump_flags_t flags) pp_right_paren (pp); } +/* Dump OpenMP's prefer_type of the init clause. */ + +static void +dump_omp_init_prefer_type (pretty_printer *pp, tree t) +{ + if (t == NULL_TREE) + return; + pp_string (pp, "prefer_type("); + const char *str = TREE_STRING_POINTER (t); + while (str[0] == (char) GOMP_INTEROP_IFR_SEPARATOR) + { + bool has_fr = false; + pp_character (pp, '{'); + str++; + while (str[0] != (char) GOMP_INTEROP_IFR_SEPARATOR) + { + if (has_fr) + pp_character (pp, ','); + has_fr = true; + pp_string (pp, "fr(\""); + pp_string (pp, omp_get_name_from_fr_id (str[0])); + pp_string (pp, "\")"); + str++; + } + str++; + if (has_fr && str[0] != '\0') + pp_character (pp, ','); + while (str[0] != '\0') + { + pp_string (pp, "attr(\""); + pp_string (pp, str); + pp_string (pp, "\")"); + str += strlen (str) + 1; + if (str[0] != '\0') + pp_character (pp, ','); + } + str++; + pp_character (pp, '}'); + if (str[0] != '\0') + pp_string (pp, ", "); + } + pp_right_paren (pp); +} /* Dump OMP clause CLAUSE, without following OMP_CLAUSE_CHAIN. @@ -604,6 +647,44 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) pp_right_paren (pp); break; + case OMP_CLAUSE_DESTROY: + pp_string (pp, "destroy("); + dump_generic_node (pp, OMP_CLAUSE_DECL (clause), + spc, flags, false); + pp_right_paren (pp); + break; + + case OMP_CLAUSE_INIT: + pp_string (pp, "init("); + dump_omp_init_prefer_type (pp, OMP_CLAUSE_INIT_PREFER_TYPE (clause)); + if (OMP_CLAUSE_INIT_TARGET (clause)) + { + if (OMP_CLAUSE_INIT_PREFER_TYPE (clause)) + pp_string (pp, ", "); + pp_string (pp, "target"); + } + if (OMP_CLAUSE_INIT_TARGETSYNC (clause)) + { + if (OMP_CLAUSE_INIT_PREFER_TYPE (clause) || OMP_CLAUSE_INIT_TARGET (clause)) + pp_string (pp, ", "); + pp_string (pp, "targetsync"); + } + if (OMP_CLAUSE_INIT_PREFER_TYPE (clause) + || OMP_CLAUSE_INIT_TARGET (clause) + || OMP_CLAUSE_INIT_TARGETSYNC (clause)) + pp_string (pp, ": "); + dump_generic_node (pp, OMP_CLAUSE_DECL (clause), + spc, flags, false); + pp_right_paren (pp); + break; + + case OMP_CLAUSE_USE: + pp_string (pp, "use("); + dump_generic_node (pp, OMP_CLAUSE_DECL (clause), + spc, flags, false); + pp_right_paren (pp); + break; + case OMP_CLAUSE_SELF: pp_string (pp, "self("); dump_generic_node (pp, OMP_CLAUSE_SELF_EXPR (clause), @@ -4078,6 +4159,12 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags, dump_omp_clauses (pp, OMP_DISPATCH_CLAUSES (node), spc, flags); goto dump_omp_body; + case OMP_INTEROP: + pp_string (pp, "#pragma omp interop"); + dump_omp_clauses (pp, OMP_INTEROP_CLAUSES (node), spc, flags); + is_expr = false; + break; + case OMP_SECTION: pp_string (pp, "#pragma omp section"); goto dump_omp_body; diff --git a/gcc/tree.cc b/gcc/tree.cc index 6d3124d1f3ac..fae1e920e9c2 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -271,6 +271,9 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_DOACROSS */ 3, /* OMP_CLAUSE__MAPPER_BINDING_ */ 2, /* OMP_CLAUSE__CACHE_ */ + 1, /* OMP_CLAUSE_DESTROY */ + 2, /* OMP_CLAUSE_INIT */ + 1, /* OMP_CLAUSE_USE */ 2, /* OMP_CLAUSE_GANG */ 1, /* OMP_CLAUSE_ASYNC */ 1, /* OMP_CLAUSE_WAIT */ @@ -372,6 +375,9 @@ const char * const omp_clause_code_name[] = "doacross", "_mapper_binding_", "_cache_", + "destroy", + "init", + "use", "gang", "async", "wait", diff --git a/gcc/tree.def b/gcc/tree.def index 6c2bf6a5479d..0b2e44c2b8ec 100644 --- a/gcc/tree.def +++ b/gcc/tree.def @@ -1238,7 +1238,7 @@ DEFTREECODE (OMP_TILE, "omp_tile", tcc_statement, 7) Operands like for OMP_FOR. */ DEFTREECODE (OMP_UNROLL, "omp_unroll", tcc_statement, 7) -/* OpenMP - #pragma acc loop [clause1 ... clauseN] +/* OpenACC - #pragma acc loop [clause1 ... clauseN] Operands like for OMP_FOR. */ DEFTREECODE (OACC_LOOP, "oacc_loop", tcc_statement, 7) @@ -1303,6 +1303,10 @@ DEFTREECODE (OMP_SCAN, "omp_scan", tcc_statement, 2) Operand 1: OMP_DISPATCH_CLAUSES: List of clauses. */ DEFTREECODE (OMP_DISPATCH, "omp_dispatch", tcc_statement, 2) +/* OpenMP - #pragma omp interop [clause1 ... clauseN] + Operand 0: OMP_INTEROP_CLAUSES: List of clauses. */ +DEFTREECODE (OMP_INTEROP, "omp_inteorp", tcc_statement, 1) + /* 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 812b08a6a113..e6fbad685ef8 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1546,6 +1546,9 @@ class auto_suppress_location_wrappers #define OMP_FOR_PRE_BODY(NODE) TREE_OPERAND (OMP_LOOPING_CHECK (NODE), 5) #define OMP_FOR_ORIG_DECLS(NODE) TREE_OPERAND (OMP_LOOPING_CHECK (NODE), 6) +#define OMP_INTEROP_CLAUSES(NODE)\ + TREE_OPERAND (OMP_INTEROP_CHECK (NODE), 0) + #define OMP_LOOPXFORM_CHECK(NODE) TREE_RANGE_CHECK (NODE, OMP_TILE, OMP_UNROLL) #define OMP_LOOPXFORM_LOWERED(NODE) \ (OMP_LOOPXFORM_CHECK (NODE)->base.public_flag) @@ -1832,6 +1835,15 @@ class auto_suppress_location_wrappers #define OMP_CLAUSE_MOTION_PRESENT(NODE) \ (OMP_CLAUSE_RANGE_CHECK (NODE, OMP_CLAUSE_FROM, OMP_CLAUSE_TO)->base.deprecated_flag) +#define OMP_CLAUSE_INIT_TARGET(NODE) \ + (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_INIT)->base.public_flag) +#define OMP_CLAUSE_INIT_TARGETSYNC(NODE) \ + (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_INIT)->base.deprecated_flag) +#define OMP_CLAUSE_INIT_PREFER_TYPE(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE), \ + OMP_CLAUSE_INIT, \ + OMP_CLAUSE_INIT), 1) + /* Nonzero if this map clause is for array (rather than pointer) based array section with zero bias. Both the non-decl OMP_CLAUSE_MAP and corresponding OMP_CLAUSE_MAP with GOMP_MAP_POINTER are marked with this flag. */ diff --git a/include/ChangeLog.omp b/include/ChangeLog.omp index 86d142dc79bc..a80d1168a334 100644 --- a/include/ChangeLog.omp +++ b/include/ChangeLog.omp @@ -1,3 +1,12 @@ +2025-01-27 Tobias Burnus <tbur...@baylibre.com> + + Backported from master: + 2024-11-22 Tobias Burnus <tbur...@baylibre.com> + + * gomp-constants.h (GOMP_INTEROP_IFR_NONE): Rename ... + (GOMP_INTEROP_IFR_UNKNOWN): ... to this. And change value. + (GOMP_INTEROP_IFR_SEPARATOR): Likewise. + 2025-01-23 Tobias Burnus <tbur...@baylibre.com> Backported from master: diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 3266de04e4f0..ed162f45cd51 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -430,10 +430,11 @@ enum gomp_map_kind #define GOMP_REQUIRES_REVERSE_OFFLOAD 0x80 #define GOMP_REQUIRES_TARGET_USED 0x200 -/* Interop foreign-runtime data. */ +/* Interop foreign-runtime data; + OpenMP defines positive values; reserve 0 and negative for GCC. */ #define GOMP_INTEROP_IFR_LAST 7 -#define GOMP_INTEROP_IFR_SEPARATOR -1 -#define GOMP_INTEROP_IFR_NONE -2 +#define GOMP_INTEROP_IFR_SEPARATOR ((char)(-__INT8_MAX__-1)) +#define GOMP_INTEROP_IFR_UNKNOWN ((char)(-__INT8_MAX__)) /* HSA specific data structures. */