From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id CA28C384585E; Fri, 23 Dec 2022 12:14:13 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org CA28C384585E Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.96,268,1665475200"; d="scan'208";a="91397254" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 23 Dec 2022 04:14:13 -0800 IronPort-SDR: JwTMrSToyIi6P4Ncg1amNYcG9a0JvMXP+6ICPRfJPdS00txy2jweqYQFhjeMT5Rq2ewssFb6Zv DCpUnbG2yYLZDJcZmnQGidDQbCvDNFoKbpmpeoZAxsmL48SUvx+K2uNFOlI9qn8/Hq/EKolryn Yi/thB43BaNFJ6FwBtw3XwvTR/hoEOww5FT6hRImgdzEEz/GGKrTlht2YRpsDI0gXhe7CjShGl Gtzw4WwANLkvDXNgl8ozCIKKhPuAcknJhsnTSn3KqVTqB+EUQkGmWvlT1WkTOv4gs3T6Yw3cjH q24= From: Julian Brown To: CC: , Tobias Burnus , "Jakub Jelinek" , Thomas Schwinge Subject: [PATCH v6 08/11] OpenMP: C++ "declare mapper" support Date: Fri, 23 Dec 2022 04:13:01 -0800 Message-ID: <7af31eda70f95a2694db119ba644209f1d855011.1671796516.git.julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 Content-Transfer-Encoding: 8bit Content-Type: text/plain X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-13.mgc.mentorg.com (139.181.222.13) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00,GIT_PATCH_0,HEADER_FROM_DIFFERENT_DOMAINS,KAM_DMARC_STATUS,RCVD_IN_MSPIKE_H2,SPF_HELO_PASS,SPF_PASS,TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: This is a new version of the patch to support OpenMP 5.0 "declare mapper" functionality for C++. As with the previously-posted version, arrays of structs whose elements would be mapped via a user-defined mapper remain unsupported. This version of the patch uses a magic VAR_DECL instead of a magic FUNCTION_DECL for representing mappers, which simplifies parsing somewhat, and slightly reduces the number of places that need special-case handling in the FE. We use the DECL_INITIAL of the VAR_DECL to hold the OMP_DECLARE_MAPPER definition. To make types agree, we use the type of the object to be mapped for both the var decl and the OMP_DECLARE_MAPPER node itself. Hence the OMP_DECLARE_MAPPER looks like a magic constant of struct type in this particular case. The magic var decl can go in all the places that the "declare mapper" function decl went previously: at the top level of the program, within a class definition (including template classes), and within a function definition (including template functions). In the class case we conceptually use the C++-17-ism of definining the var decl "inline static", equivalent to e.g.: [template ...] class bla { static inline omp declare mapper ... = #define omp declare mapper ..." }; (though of course we don't restrict the "declare mapper"-in-class syntax to C++-17.) The new representation necessitates some changes to template instantiation -- declare mappers may trigger implicitly, so we must make sure they are instantiated before they are needed (see changes to mark_used, etc.). I've rearranged the processing done by the gimplify_scan_omp_clauses and gimplify_adjust_omp_clauses functions so the order of the phases can remain intact in the presence of declared mappers. To do this, most gimplification of clauses in gimplify_scan_omp_clauses has been moved to gimplify_adjust_omp_clauses. This allows e.g. struct sibling-list handling and topological clause sorting to work with the non-gimplified form of clauses in the latter function -- including those that arise from mapper expansion. This seems to work well now. Relative to the last-posted version, this patch brings forward various refactoring that was previously done by the C and Fortran "declare mapper" support patches -- aiming to reduce churn. E.g. nested mapper finding and mapper instantiation has been moved to c-family/c-omp.cc so it can be shared between C and C++, and omp_name_type in omp-general.h (used as the key to hash mapper definitions) is already templatized ready for Fortran support. This patch does not synthesize default mappers that map each of a struct's elements individually: whole-struct mappings are still done by copying the block of memory containing the struct. That works fine apart from cases where a struct has a member that is a reference (to a pointer). We could fix that by synthesizing a mapper for such cases (only), but that hasn't been attempted yet. (I think that means Jakub's concerns about blow-up of element mappings won't be a problem until that's done.) New tests added in {gcc,libgomp}/c-c++-common have been restricted to C++ for now. 2022-11-30 Julian Brown gcc/c-family/ * c-common.h (omp_mapper_list): Add forward declaration. (c_omp_find_nested_mappers, c_omp_instantiate_mappers): Add prototypes. * c-omp.cc (c_omp_find_nested_mappers): New function. (remap_mapper_decl_info): New struct. (remap_mapper_decl_1, omp_instantiate_mapper, c_omp_instantiate_mappers): New functions. gcc/cp/ * constexpr.cc (reduced_constant_expression_p): Add OMP_DECLARE_MAPPER case. (cxx_eval_constant_expression, potential_constant_expression_1): Likewise. * cp-gimplify.cc (cxx_omp_finish_mapper_clauses): New function. * cp-objcp-common.h (LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES, LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define langhooks. * cp-tree.h (lang_decl_base): Add omp_declare_mapper_p field. Recount spare bits comment. (DECL_OMP_DECLARE_MAPPER_P): New macro. (omp_mapper_id, cp_check_omp_declare_mapper, omp_instantiate_mappers, cxx_omp_finish_mapper_clauses, cxx_omp_mapper_lookup, cxx_omp_extract_mapper_directive, cxx_omp_map_array_section: Add prototypes. * decl.cc (check_initializer): Add OpenMP declare mapper support. (cp_finish_decl): Set DECL_INITIAL for OpenMP declare mapper var decls as appropriate. * decl2.cc (mark_used): Instantiate OpenMP "declare mapper" magic var decls. * error.cc (dump_omp_declare_mapper): New function. (dump_simple_decl): Use above. * parser.cc (cp_parser_omp_clause_map): Add KIND parameter. Support "mapper" modifier. (cp_parser_omp_all_clauses): Add KIND argument to cp_parser_omp_clause_map call. (cp_parser_omp_target): Call omp_instantiate_mappers before finish_omp_clauses. (cp_parser_omp_declare_mapper): New function. (cp_parser_omp_declare): Add "declare mapper" support. * pt.cc (tsubst_decl): Adjust name of "declare mapper" magic var decls once we know their type. (tsubst_omp_clauses): Call omp_instantiate_mappers before finish_omp_clauses, for target regions. (tsubst_expr): Support OMP_DECLARE_MAPPER nodes. (instantiate_decl): Instantiate initialiser (i.e definition) for OpenMP declare mappers. * semantics.cc (gimplify.h): Include. (omp_mapper_id, omp_mapper_lookup, omp_extract_mapper_directive, cxx_omp_map_array_section, cp_check_omp_declare_mapper): New functions. (finish_omp_clauses): Delete GOMP_MAP_PUSH_MAPPER_NAME and GOMP_MAP_POP_MAPPER_NAME artificial clauses. (omp_target_walk_data): Add MAPPERS field. (finish_omp_target_clauses_r): Scan for uses of struct/union/class type variables. (finish_omp_target_clauses): Create artificial mapper binding clauses for used structs/unions/classes in offload region. gcc/fortran/ * parse.cc (tree.h, fold-const.h, tree-hash-traits.h): Add includes (for additions to omp-general.h). gcc/ * gimplify.cc (gimplify_omp_ctx): Add IMPLICIT_MAPPERS field. (new_omp_context): Initialise IMPLICIT_MAPPERS hash map. (delete_omp_context): Delete IMPLICIT_MAPPERS hash map. (instantiate_mapper_info): New structs. (remap_mapper_decl_1, omp_mapper_copy_decl, omp_instantiate_mapper, omp_instantiate_implicit_mappers): New functions. (gimplify_scan_omp_clauses): Remove PREV_LIST_P tracking and use mapping groups instead. Remove most gimplification of OMP_CLAUSE_MAP nodes from here, but still populate ctx->variables splay tree. Handle MAPPER_BINDING clauses. (gimplify_adjust_omp_clauses): Move most gimplification of OMP clauses here. Instantiate implicit declared mappers. (gimplify_omp_declare_mapper): New function. (gimplify_expr): Call above function. * langhooks.def (lhd_omp_finish_mapper_clauses, lhd_omp_mapper_lookup, lhd_omp_extract_mapper_directive, lhd_omp_map_array_section): Add prototypes. (LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES, LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define macros. (LANG_HOOK_DECLS): Add above macros. * langhooks.cc (lhd_omp_finish_mapper_clauses, lhd_omp_mapper_lookup, lhd_omp_extract_mapper_directive, lhd_omp_map_array_section): New dummy functions. * langhooks.h (lang_hooks_for_decls): Add OMP_FINISH_MAPPER_CLAUSES, OMP_MAPPER_LOOKUP, OMP_EXTRACT_MAPPER_DIRECTIVE, OMP_MAP_ARRAY_SECTION hooks. * omp-general.h (omp_name_type): Add templatized struct, hash type traits (for omp_name_type specialization). (omp_mapper_list): Add struct. * tree-core.h (omp_clause_code): Add OMP_CLAUSE__MAPPER_BINDING_. * tree-pretty-print.cc (dump_omp_clause): Support GOMP_MAP_UNSET, GOMP_MAP_PUSH_MAPPER_NAME, GOMP_MAP_POP_MAPPER_NAME artificial mapping clauses. Support OMP_CLAUSE__MAPPER_BINDING_ and OMP_DECLARE_MAPPER. * tree.cc (omp_clause_num_ops, omp_clause_code_name): Add OMP_CLAUSE__MAPPER_BINDING_. * tree.def (OMP_DECLARE_MAPPER): New tree code. * tree.h (OMP_DECLARE_MAPPER_ID, OMP_DECLARE_MAPPER_DECL, OMP_DECLARE_MAPPER_CLAUSES): New defines. (OMP_CLAUSE__MAPPER_BINDING__ID, OMP_CLAUSE__MAPPER_BINDING__DECL, OMP_CLAUSE__MAPPER_BINDING__MAPPER): New defines. include/ * gomp-constants.h (gomp_map_kind): Add GOMP_MAP_UNSET, GOMP_MAP_PUSH_MAPPER_NAME, GOMP_MAP_POP_MAPPER_NAME artificial mapping clause types. gcc/testsuite/ * c-c++-common/gomp/map-6.c: Update error scan output. * c-c++-common/gomp/declare-mapper-3.c: New test (only enabled for C++ for now). * c-c++-common/gomp/declare-mapper-4.c: Likewise. * c-c++-common/gomp/declare-mapper-5.c: Likewise. * c-c++-common/gomp/declare-mapper-6.c: Likewise. * c-c++-common/gomp/declare-mapper-7.c: Likewise. * c-c++-common/gomp/declare-mapper-8.c: Likewise. * c-c++-common/gomp/declare-mapper-9.c: Likewise. * c-c++-common/gomp/declare-mapper-12.c: Likewise. * g++.dg/gomp/declare-mapper-1.C: New test. * g++.dg/gomp/declare-mapper-2.C: New test. libgomp/ * testsuite/libgomp.c++/declare-mapper-1.C: New test. * testsuite/libgomp.c++/declare-mapper-2.C: New test. * testsuite/libgomp.c++/declare-mapper-3.C: New test. * testsuite/libgomp.c++/declare-mapper-4.C: New test. * testsuite/libgomp.c++/declare-mapper-5.C: New test. * testsuite/libgomp.c++/declare-mapper-6.C: New test. * testsuite/libgomp.c++/declare-mapper-7.C: New test. * testsuite/libgomp.c++/declare-mapper-8.C: New test. * testsuite/libgomp.c-c++-common/declare-mapper-9.c: New test (only enabled for C++ for now). * testsuite/libgomp.c-c++-common/declare-mapper-10.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-11.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-12.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-13.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-14.c: Likewise. --- gcc/c-family/c-common.h | 3 + gcc/c-family/c-omp.cc | 300 ++++++ gcc/cp/constexpr.cc | 9 + gcc/cp/cp-gimplify.cc | 6 + gcc/cp/cp-objcp-common.h | 9 + gcc/cp/cp-tree.h | 17 +- gcc/cp/decl.cc | 27 +- gcc/cp/decl2.cc | 9 +- gcc/cp/error.cc | 25 + gcc/cp/parser.cc | 284 +++++- gcc/cp/pt.cc | 29 +- gcc/cp/semantics.cc | 188 +++- gcc/fortran/parse.cc | 3 + gcc/gimplify.cc | 958 ++++++++++++------ gcc/langhooks-def.h | 13 + gcc/langhooks.cc | 35 + gcc/langhooks.h | 16 + gcc/omp-general.h | 86 ++ .../c-c++-common/gomp/declare-mapper-12.c | 22 + .../c-c++-common/gomp/declare-mapper-3.c | 30 + .../c-c++-common/gomp/declare-mapper-4.c | 78 ++ .../c-c++-common/gomp/declare-mapper-5.c | 26 + .../c-c++-common/gomp/declare-mapper-6.c | 23 + .../c-c++-common/gomp/declare-mapper-7.c | 29 + .../c-c++-common/gomp/declare-mapper-8.c | 43 + .../c-c++-common/gomp/declare-mapper-9.c | 34 + gcc/testsuite/c-c++-common/gomp/map-6.c | 6 +- gcc/testsuite/g++.dg/gomp/declare-mapper-1.C | 58 ++ gcc/testsuite/g++.dg/gomp/declare-mapper-2.C | 30 + gcc/tree-core.h | 4 + gcc/tree-pretty-print.cc | 41 + gcc/tree.cc | 2 + gcc/tree.def | 7 + gcc/tree.h | 19 + include/gomp-constants.h | 8 +- .../testsuite/libgomp.c++/declare-mapper-1.C | 87 ++ .../testsuite/libgomp.c++/declare-mapper-2.C | 55 + .../testsuite/libgomp.c++/declare-mapper-3.C | 63 ++ .../testsuite/libgomp.c++/declare-mapper-4.C | 63 ++ .../testsuite/libgomp.c++/declare-mapper-5.C | 52 + .../testsuite/libgomp.c++/declare-mapper-6.C | 37 + .../testsuite/libgomp.c++/declare-mapper-7.C | 48 + .../testsuite/libgomp.c++/declare-mapper-8.C | 61 ++ .../libgomp.c-c++-common/declare-mapper-10.c | 60 ++ .../libgomp.c-c++-common/declare-mapper-11.c | 59 ++ .../libgomp.c-c++-common/declare-mapper-12.c | 87 ++ .../libgomp.c-c++-common/declare-mapper-13.c | 55 + .../libgomp.c-c++-common/declare-mapper-14.c | 57 ++ .../libgomp.c-c++-common/declare-mapper-9.c | 62 ++ 49 files changed, 2998 insertions(+), 325 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c create mode 100644 gcc/testsuite/g++.dg/gomp/declare-mapper-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/declare-mapper-2.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-1.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-2.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-3.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-4.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-5.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-6.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-7.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-8.C create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index 510ee2d24e0c..3472f180543e 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1281,6 +1281,9 @@ extern enum omp_clause_defaultmap_kind c_omp_predetermined_mapping (tree); extern tree c_omp_check_context_selector (location_t, tree); extern void c_omp_mark_declare_variant (location_t, tree, tree); extern void c_omp_adjust_map_clauses (tree, bool); +template struct omp_mapper_list; +extern void c_omp_find_nested_mappers (struct omp_mapper_list *, tree); +extern tree c_omp_instantiate_mappers (tree); namespace omp_addr_tokenizer { struct omp_addr_token; } typedef omp_addr_tokenizer::omp_addr_token omp_addr_token; diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index 2de3fc2be759..32699adc664c 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -3939,6 +3939,306 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr, return error_mark_node; } +/* Given a mapper function MAPPER_FN, recursively scan through the map clauses + for that mapper, and if any of those should use a (named or unnamed) mapper + themselves, add it to MLIST. */ + +void +c_omp_find_nested_mappers (omp_mapper_list *mlist, tree mapper_fn) +{ + tree mapper = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn); + tree mapper_name = NULL_TREE; + + if (mapper == error_mark_node) + return; + + gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER); + + for (tree clause = OMP_DECLARE_MAPPER_CLAUSES (mapper); + clause; + clause = OMP_CLAUSE_CHAIN (clause)) + { + tree expr = OMP_CLAUSE_DECL (clause); + enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (clause); + tree elem_type; + + if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME) + { + mapper_name = expr; + continue; + } + else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME) + { + mapper_name = NULL_TREE; + continue; + } + + gcc_assert (TREE_CODE (expr) != TREE_LIST); + if (TREE_CODE (expr) == OMP_ARRAY_SECTION) + { + while (TREE_CODE (expr) == OMP_ARRAY_SECTION) + expr = TREE_OPERAND (expr, 0); + + elem_type = TREE_TYPE (expr); + } + else + elem_type = TREE_TYPE (expr); + + /* This might be too much... or not enough? */ + while (TREE_CODE (elem_type) == ARRAY_TYPE + || TREE_CODE (elem_type) == POINTER_TYPE + || TREE_CODE (elem_type) == REFERENCE_TYPE) + elem_type = TREE_TYPE (elem_type); + + elem_type = TYPE_MAIN_VARIANT (elem_type); + + if (AGGREGATE_TYPE_P (elem_type) + && !mlist->contains (mapper_name, elem_type)) + { + tree nested_mapper_fn + = lang_hooks.decls.omp_mapper_lookup (mapper_name, elem_type); + + if (nested_mapper_fn) + { + mlist->add_mapper (mapper_name, elem_type, nested_mapper_fn); + c_omp_find_nested_mappers (mlist, nested_mapper_fn); + } + else if (mapper_name) + { + error ("mapper %qE not found for type %qT", mapper_name, + elem_type); + continue; + } + } + } +} + +struct remap_mapper_decl_info +{ + tree dummy_var; + tree expr; +}; + +/* Helper for rewriting DUMMY_VAR into EXPR in a map clause decl. */ + +static tree +remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data) +{ + remap_mapper_decl_info *map_info = (remap_mapper_decl_info *) data; + + if (operand_equal_p (*tp, map_info->dummy_var)) + { + *tp = map_info->expr; + *walk_subtrees = 0; + } + + return NULL_TREE; +} + +/* Instantiate a mapper MAPPER for expression EXPR, adding new clauses to + OUTLIST. OUTER_KIND is the mapping kind to use if not already specified in + the mapper declaration. */ + +static tree * +omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, + enum gomp_map_kind outer_kind) +{ + tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper); + tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper); + tree mapper_name = NULL_TREE; + + remap_mapper_decl_info map_info; + map_info.dummy_var = dummy_var; + map_info.expr = expr; + + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + { + tree unshared = unshare_expr (c); + enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (c); + tree t = OMP_CLAUSE_DECL (unshared); + tree type = NULL_TREE; + bool nonunit_array_with_mapper = false; + + if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME) + { + mapper_name = t; + continue; + } + else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME) + { + mapper_name = NULL_TREE; + continue; + } + + if (TREE_CODE (t) == OMP_ARRAY_SECTION) + { + location_t loc = OMP_CLAUSE_LOCATION (c); + tree t2 = lang_hooks.decls.omp_map_array_section (loc, t); + + if (t2 == t) + { + nonunit_array_with_mapper = true; + /* We'd want use the mapper for the element type if this worked: + look that one up. */ + type = TREE_TYPE (TREE_TYPE (t)); + } + else + { + t = t2; + type = TREE_TYPE (t); + } + } + else + type = TREE_TYPE (t); + + gcc_assert (type); + + if (type == error_mark_node) + continue; + + walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL); + + if (OMP_CLAUSE_MAP_KIND (unshared) == GOMP_MAP_UNSET) + OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind); + + type = TYPE_MAIN_VARIANT (type); + + tree mapper_fn = lang_hooks.decls.omp_mapper_lookup (mapper_name, type); + + if (mapper_fn && nonunit_array_with_mapper) + { + sorry ("user-defined mapper with non-unit length array section"); + continue; + } + else if (mapper_fn) + { + tree nested_mapper + = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn); + if (nested_mapper != mapper) + { + if (clause_kind == GOMP_MAP_UNSET) + clause_kind = outer_kind; + + outlist = omp_instantiate_mapper (outlist, nested_mapper, + t, clause_kind); + continue; + } + } + else if (mapper_name) + { + error ("mapper %qE not found for type %qT", mapper_name, type); + continue; + } + + *outlist = unshared; + outlist = &OMP_CLAUSE_CHAIN (unshared); + } + + return outlist; +} + +/* Given a list of CLAUSES, scan each clause and invoke a user-defined mapper + appropriate to the type of the data in that clause, if such a mapper is + visible in the current parsing context. */ + +tree +c_omp_instantiate_mappers (tree clauses) +{ + tree c, *pc, mapper_name = NULL_TREE; + + for (pc = &clauses, c = clauses; c; c = *pc) + { + bool using_mapper = false; + + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_MAP: + { + tree t = OMP_CLAUSE_DECL (c); + tree type = NULL_TREE; + bool nonunit_array_with_mapper = false; + + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME) + { + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME) + mapper_name = OMP_CLAUSE_DECL (c); + else + mapper_name = NULL_TREE; + pc = &OMP_CLAUSE_CHAIN (c); + continue; + } + + if (TREE_CODE (t) == OMP_ARRAY_SECTION) + { + location_t loc = OMP_CLAUSE_LOCATION (c); + tree t2 = lang_hooks.decls.omp_map_array_section (loc, t); + + if (t2 == t) + { + /* !!! Array sections of size >1 with mappers for elements + are hard to support. Do something here. */ + nonunit_array_with_mapper = true; + type = TREE_TYPE (TREE_TYPE (t)); + } + else + { + t = t2; + type = TREE_TYPE (t); + } + } + else + type = TREE_TYPE (t); + + if (type == NULL_TREE || type == error_mark_node) + { + pc = &OMP_CLAUSE_CHAIN (c); + continue; + } + + enum gomp_map_kind kind = OMP_CLAUSE_MAP_KIND (c); + if (kind == GOMP_MAP_UNSET) + kind = GOMP_MAP_TOFROM; + + type = TYPE_MAIN_VARIANT (type); + + tree mapper_fn + = lang_hooks.decls.omp_mapper_lookup (mapper_name, type); + + if (mapper_fn && nonunit_array_with_mapper) + { + sorry ("user-defined mapper with non-unit length " + "array section"); + using_mapper = true; + } + else if (mapper_fn) + { + tree mapper + = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn); + pc = omp_instantiate_mapper (pc, mapper, t, kind); + using_mapper = true; + } + else if (mapper_name) + { + error ("mapper %qE not found for type %qT", mapper_name, type); + using_mapper = true; + } + } + break; + + default: + ; + } + + if (using_mapper) + *pc = OMP_CLAUSE_CHAIN (c); + else + pc = &OMP_CLAUSE_CHAIN (c); + } + + return clauses; +} + const struct c_omp_directive c_omp_directives[] = { /* Keep this alphabetically sorted by the first word. Non-null second/third if any should precede null ones. */ diff --git a/gcc/cp/constexpr.cc b/gcc/cp/constexpr.cc index 6994d2b25e04..944ef5d70de7 100644 --- a/gcc/cp/constexpr.cc +++ b/gcc/cp/constexpr.cc @@ -3273,6 +3273,9 @@ reduced_constant_expression_p (tree t) /* Even if we can't lower this yet, it's constant. */ return true; + case OMP_DECLARE_MAPPER: + return true; + case CONSTRUCTOR: /* And we need to handle PTRMEM_CST wrapped in a CONSTRUCTOR. */ tree field; @@ -7088,6 +7091,7 @@ cxx_eval_constant_expression (const constexpr_ctx *ctx, tree t, case LABEL_EXPR: case CASE_LABEL_EXPR: case PREDICT_EXPR: + case OMP_DECLARE_MAPPER: return t; case PARM_DECL: @@ -9499,6 +9503,11 @@ potential_constant_expression_1 (tree t, bool want_rval, bool strict, bool now, "expression", t); return false; + case OMP_DECLARE_MAPPER: + /* This can be used to initialize VAR_DECLs: it's treated as a magic + constant. */ + return true; + case ASM_EXPR: if (flags & tf_error) inline_asm_in_constexpr_error (loc, fundef_p); diff --git a/gcc/cp/cp-gimplify.cc b/gcc/cp/cp-gimplify.cc index 1e0b4f8c1cf5..3fa20b5a203e 100644 --- a/gcc/cp/cp-gimplify.cc +++ b/gcc/cp/cp-gimplify.cc @@ -2319,6 +2319,12 @@ cxx_omp_finish_clause (tree c, gimple_seq *, bool /* openacc */) } } +tree +cxx_omp_finish_mapper_clauses (tree clauses) +{ + return finish_omp_clauses (clauses, C_ORT_OMP); +} + /* Return true if DECL's DECL_VALUE_EXPR (if any) should be disregarded in OpenMP construct, because it is going to be remapped during OpenMP lowering. SHARED is true if DECL diff --git a/gcc/cp/cp-objcp-common.h b/gcc/cp/cp-objcp-common.h index f4ba0c9e012b..ad2590fbe856 100644 --- a/gcc/cp/cp-objcp-common.h +++ b/gcc/cp/cp-objcp-common.h @@ -184,6 +184,15 @@ extern tree cxx_simulate_record_decl (location_t, const char *, #define LANG_HOOKS_OMP_CLAUSE_DTOR cxx_omp_clause_dtor #undef LANG_HOOKS_OMP_FINISH_CLAUSE #define LANG_HOOKS_OMP_FINISH_CLAUSE cxx_omp_finish_clause +#undef LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES +#define LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES cxx_omp_finish_mapper_clauses +#undef LANG_HOOKS_OMP_MAPPER_LOOKUP +#define LANG_HOOKS_OMP_MAPPER_LOOKUP cxx_omp_mapper_lookup +#undef LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE +#define LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE \ + cxx_omp_extract_mapper_directive +#undef LANG_HOOKS_OMP_MAP_ARRAY_SECTION +#define LANG_HOOKS_OMP_MAP_ARRAY_SECTION cxx_omp_map_array_section #undef LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE #define LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE cxx_omp_privatize_by_reference #undef LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index 8157d8fd0e5f..b395d7be0220 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -2882,7 +2882,10 @@ struct GTY(()) lang_decl_base { /* VAR_DECL or FUNCTION_DECL has keyed decls. */ unsigned module_keyed_decls_p : 1; - /* 12 spare bits. */ + /* VAR_DECL being used to represent an OpenMP declared mapper. */ + unsigned omp_declare_mapper_p : 1; + + /* 10 spare bits. */ }; /* True for DECL codes which have template info and access. */ @@ -4356,6 +4359,11 @@ get_vec_init_expr (tree t) #define DECL_OMP_DECLARE_REDUCTION_P(NODE) \ (LANG_DECL_FN_CHECK (DECL_COMMON_CHECK (NODE))->omp_declare_reduction_p) +/* Nonzero if NODE is an artificial FUNCTION_DECL for + #pragma omp declare mapper. */ +#define DECL_OMP_DECLARE_MAPPER_P(NODE) \ + (DECL_LANG_SPECIFIC (VAR_DECL_CHECK (NODE))->u.base.omp_declare_mapper_p) + /* Nonzero if DECL has been declared threadprivate by #pragma omp threadprivate. */ #define CP_DECL_THREADPRIVATE_P(DECL) \ @@ -7716,10 +7724,13 @@ extern tree finish_qualified_id_expr (tree, tree, bool, bool, extern void simplify_aggr_init_expr (tree *); extern void finalize_nrv (tree *, tree, tree); extern tree omp_reduction_id (enum tree_code, tree, tree); +extern tree omp_mapper_id (tree, tree); extern tree cp_remove_omp_priv_cleanup_stmt (tree *, int *, void *); extern bool cp_check_omp_declare_reduction (tree); +extern bool cp_check_omp_declare_mapper (tree); extern void finish_omp_declare_simd_methods (tree); extern tree finish_omp_clauses (tree, enum c_omp_region_type); +extern tree omp_instantiate_mappers (tree); extern tree push_omp_privatization_clauses (bool); extern void pop_omp_privatization_clauses (tree); extern void save_omp_privatization_clauses (vec &); @@ -8279,6 +8290,10 @@ extern tree cxx_omp_clause_copy_ctor (tree, tree, tree); extern tree cxx_omp_clause_assign_op (tree, tree, tree); extern tree cxx_omp_clause_dtor (tree, tree); extern void cxx_omp_finish_clause (tree, gimple_seq *, bool); +extern tree cxx_omp_finish_mapper_clauses (tree); +extern tree cxx_omp_mapper_lookup (tree, tree); +extern tree cxx_omp_extract_mapper_directive (tree); +extern tree cxx_omp_map_array_section (location_t, tree); extern bool cxx_omp_privatize_by_reference (const_tree); extern bool cxx_omp_disregard_value_expr (tree, bool); extern void cp_fold_function (tree); diff --git a/gcc/cp/decl.cc b/gcc/cp/decl.cc index 3c0355a1c395..9914ccf7c0fc 100644 --- a/gcc/cp/decl.cc +++ b/gcc/cp/decl.cc @@ -7441,6 +7441,12 @@ check_initializer (tree decl, tree init, int flags, vec **cleanups) } else if (!init && DECL_REALLY_EXTERN (decl)) ; + else if (flag_openmp + && VAR_P (decl) + && DECL_LANG_SPECIFIC (decl) + && DECL_OMP_DECLARE_MAPPER_P (decl) + && TREE_CODE (init) == OMP_DECLARE_MAPPER) + return NULL_TREE; else if (init || type_build_ctor_call (type) || TYPE_REF_P (type)) { @@ -8584,14 +8590,23 @@ cp_finish_decl (tree decl, tree init, bool init_const_expr_p, varpool_node::get_create (decl); } + if (flag_openmp + && VAR_P (decl) + && DECL_LANG_SPECIFIC (decl) + && DECL_OMP_DECLARE_MAPPER_P (decl) + && init) + { + gcc_assert (TREE_CODE (init) == OMP_DECLARE_MAPPER); + DECL_INITIAL (decl) = init; + } /* Convert the initializer to the type of DECL, if we have not already initialized DECL. */ - if (!DECL_INITIALIZED_P (decl) - /* If !DECL_EXTERNAL then DECL is being defined. In the - case of a static data member initialized inside the - class-specifier, there can be an initializer even if DECL - is *not* defined. */ - && (!DECL_EXTERNAL (decl) || init)) + else if (!DECL_INITIALIZED_P (decl) + /* If !DECL_EXTERNAL then DECL is being defined. In the + case of a static data member initialized inside the + class-specifier, there can be an initializer even if DECL + is *not* defined. */ + && (!DECL_EXTERNAL (decl) || init)) { cleanups = make_tree_vector (); init = check_initializer (decl, init, flags, &cleanups); diff --git a/gcc/cp/decl2.cc b/gcc/cp/decl2.cc index dac74198b9f1..2ac769695004 100644 --- a/gcc/cp/decl2.cc +++ b/gcc/cp/decl2.cc @@ -5785,10 +5785,15 @@ mark_used (tree decl, tsubst_flags_t complain /* = tf_warning_or_error */) /* If DECL has a deduced return type, we need to instantiate it now to find out its type. For OpenMP user defined reductions, we need them - instantiated for reduction clauses which inline them by hand directly. */ + instantiated for reduction clauses which inline them by hand directly. + OpenMP declared mappers are used implicitly so must be instantiated + before they can be detected. */ if (undeduced_auto_decl (decl) || (TREE_CODE (decl) == FUNCTION_DECL - && DECL_OMP_DECLARE_REDUCTION_P (decl))) + && DECL_OMP_DECLARE_REDUCTION_P (decl)) + || (TREE_CODE (decl) == VAR_DECL + && DECL_LANG_SPECIFIC (decl) + && DECL_OMP_DECLARE_MAPPER_P (decl))) maybe_instantiate_decl (decl); if (processing_template_decl || in_template_function ()) diff --git a/gcc/cp/error.cc b/gcc/cp/error.cc index 6262028da272..52b66975d238 100644 --- a/gcc/cp/error.cc +++ b/gcc/cp/error.cc @@ -1137,12 +1137,37 @@ dump_global_iord (cxx_pretty_printer *pp, tree t) pp_printf (pp, p, DECL_SOURCE_FILE (t)); } +/* Write a representation of OpenMP "declare mapper" T to PP in a manner + suitable for error messages. */ + +static void +dump_omp_declare_mapper (cxx_pretty_printer *pp, tree t, int flags) +{ + pp_string (pp, "#pragma omp declare mapper"); + if (t == NULL_TREE || t == error_mark_node) + return; + pp_space (pp); + pp_cxx_left_paren (pp); + if (OMP_DECLARE_MAPPER_ID (t)) + { + pp_cxx_tree_identifier (pp, OMP_DECLARE_MAPPER_ID (t)); + pp_colon (pp); + } + dump_type (pp, TREE_TYPE (t), flags); + pp_cxx_right_paren (pp); +} + static void dump_simple_decl (cxx_pretty_printer *pp, tree t, tree type, int flags) { if (TREE_CODE (t) == VAR_DECL && DECL_NTTP_OBJECT_P (t)) return dump_expr (pp, DECL_INITIAL (t), flags); + if (TREE_CODE (t) == VAR_DECL + && DECL_LANG_SPECIFIC (t) + && DECL_OMP_DECLARE_MAPPER_P (t)) + return dump_omp_declare_mapper (pp, DECL_INITIAL (t), flags); + if (flags & TFF_DECL_SPECIFIERS) { if (concept_definition_p (t)) diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 1a45dad6b8e4..4e2a621b53bc 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -40551,13 +40551,12 @@ cp_parser_omp_clause_doacross (cp_parser *parser, tree list, location_t loc) map ( [map-type-modifier[,] ...] map-kind: variable-list ) map-type-modifier: - always | close */ + always | close | mapper ( mapper-name ) */ static tree -cp_parser_omp_clause_map (cp_parser *parser, tree list) +cp_parser_omp_clause_map (cp_parser *parser, tree list, enum gomp_map_kind kind) { tree nlist, c; - enum gomp_map_kind kind = GOMP_MAP_TOFROM; if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) return list; @@ -40575,11 +40574,16 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA) pos++; + else if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type + == CPP_OPEN_PAREN) + pos = cp_parser_skip_balanced_tokens (parser, pos + 1); pos++; } bool always_modifier = false; bool close_modifier = false; + bool mapper_modifier = false; + tree mapper_name = NULL_TREE; for (int pos = 1; pos < map_kind_pos; ++pos) { cp_token *tok = cp_lexer_peek_token (parser->lexer); @@ -40602,6 +40606,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) return list; } always_modifier = true; + cp_lexer_consume_token (parser->lexer); } else if (strcmp ("close", p) == 0) { @@ -40615,20 +40620,83 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) return list; } close_modifier = true; + cp_lexer_consume_token (parser->lexer); + } + else if (strcmp ("mapper", p) == 0) + { + cp_lexer_consume_token (parser->lexer); + + matching_parens parens; + if (parens.require_open (parser)) + { + if (mapper_modifier) + { + cp_parser_error (parser, "too many % modifiers"); + /* Assume it's a well-formed mapper modifier, even if it + seems to be in the wrong place. */ + cp_lexer_consume_token (parser->lexer); + parens.require_close (parser); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/ + true); + return list; + } + + tok = cp_lexer_peek_token (parser->lexer); + switch (tok->type) + { + case CPP_NAME: + { + cp_expr e = cp_parser_identifier (parser); + if (e != error_mark_node) + mapper_name = e; + else + goto err; + } + break; + + case CPP_KEYWORD: + if (tok->keyword == RID_DEFAULT) + { + cp_lexer_consume_token (parser->lexer); + break; + } + /* Fallthrough. */ + + default: + err: + cp_parser_error (parser, + "expected identifier or %"); + return list; + } + + if (!parens.require_close (parser)) + { + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/ + true); + return list; + } + + mapper_modifier = true; + pos += 3; + } } else { cp_parser_error (parser, "%<#pragma omp target%> with " - "modifier other than % or " - "% on % clause"); + "modifier other than %, % " + "or % on % clause"); cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, /*or_comma=*/false, /*consume_paren=*/true); return list; } - - cp_lexer_consume_token (parser->lexer); } if (cp_lexer_next_token_is (parser->lexer, CPP_NAME) @@ -40674,8 +40742,30 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) NULL, true); finish_scope (); + tree last_new = NULL_TREE; + for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c)) - OMP_CLAUSE_SET_MAP_KIND (c, kind); + { + OMP_CLAUSE_SET_MAP_KIND (c, kind); + last_new = c; + } + + if (mapper_name) + { + tree name = build_omp_clause (input_location, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_PUSH_MAPPER_NAME); + OMP_CLAUSE_DECL (name) = mapper_name; + OMP_CLAUSE_CHAIN (name) = nlist; + nlist = name; + + gcc_assert (last_new); + + name = build_omp_clause (input_location, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_POP_MAPPER_NAME); + OMP_CLAUSE_DECL (name) = null_pointer_node; + OMP_CLAUSE_CHAIN (name) = OMP_CLAUSE_CHAIN (last_new); + OMP_CLAUSE_CHAIN (last_new) = name; + } return nlist; } @@ -41486,7 +41576,7 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, c_name = "detach"; break; case PRAGMA_OMP_CLAUSE_MAP: - clauses = cp_parser_omp_clause_map (parser, clauses); + clauses = cp_parser_omp_clause_map (parser, clauses, GOMP_MAP_TOFROM); c_name = "map"; break; case PRAGMA_OMP_CLAUSE_DEVICE: @@ -45654,6 +45744,8 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = nc; } + if (!processing_template_decl) + clauses = c_omp_instantiate_mappers (clauses); clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET); c_omp_adjust_map_clauses (clauses, true); @@ -47902,6 +47994,172 @@ cp_parser_omp_declare_reduction (cp_parser *parser, cp_token *pragma_tok, obstack_free (&declarator_obstack, p); } +/* OpenMP 5.0 + #pragma omp declare mapper([mapper-identifier:]type var) \ + [clause[[,] clause] ... ] new-line */ + +static void +cp_parser_omp_declare_mapper (cp_parser *parser, cp_token *pragma_tok, + enum pragma_context) +{ + cp_token *token = NULL; + tree type = NULL_TREE, vardecl = NULL_TREE, block = NULL_TREE; + bool block_scope = false; + /* Don't create location wrapper nodes within "declare mapper" + directives. */ + auto_suppress_location_wrappers sentinel; + tree mapper_name = NULL_TREE; + tree mapper_id, id, placeholder, mapper, maplist = NULL_TREE; + + if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + goto fail; + + if (current_function_decl) + block_scope = true; + + token = cp_lexer_peek_token (parser->lexer); + + if (cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON)) + { + switch (token->type) + { + case CPP_NAME: + { + cp_expr e = cp_parser_identifier (parser); + if (e != error_mark_node) + mapper_name = e; + else + goto fail; + } + break; + + case CPP_KEYWORD: + if (token->keyword == RID_DEFAULT) + { + mapper_name = NULL_TREE; + cp_lexer_consume_token (parser->lexer); + break; + } + /* Fallthrough. */ + + default: + cp_parser_error (parser, "expected identifier or %"); + } + + if (!cp_parser_require (parser, CPP_COLON, RT_COLON)) + goto fail; + } + + { + const char *saved_message = parser->type_definition_forbidden_message; + parser->type_definition_forbidden_message + = G_("types may not be defined within %"); + type_id_in_expr_sentinel s (parser); + type = cp_parser_type_id (parser); + parser->type_definition_forbidden_message = saved_message; + } + + if (dependent_type_p (type)) + mapper_id = omp_mapper_id (mapper_name, NULL_TREE); + else + mapper_id = omp_mapper_id (mapper_name, type); + + vardecl = build_lang_decl (VAR_DECL, mapper_id, type); + DECL_ARTIFICIAL (vardecl) = 1; + TREE_STATIC (vardecl) = 1; + TREE_PUBLIC (vardecl) = 0; + DECL_EXTERNAL (vardecl) = 0; + DECL_DECLARED_CONSTEXPR_P (vardecl) = 1; + DECL_INITIALIZED_BY_CONSTANT_EXPRESSION_P (vardecl) = 1; + DECL_OMP_DECLARE_MAPPER_P (vardecl) = 1; + + keep_next_level (true); + block = begin_omp_structured_block (); + + if (block_scope) + DECL_CONTEXT (vardecl) = current_function_decl; + else if (current_class_type) + DECL_CONTEXT (vardecl) = current_class_type; + else + DECL_CONTEXT (vardecl) = current_namespace; + + if (processing_template_decl) + vardecl = push_template_decl (vardecl); + + id = cp_parser_declarator_id (parser, false); + + if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN)) + { + finish_omp_structured_block (block); + goto fail; + } + + placeholder = build_lang_decl (VAR_DECL, id, type); + DECL_CONTEXT (placeholder) = DECL_CONTEXT (vardecl); + if (processing_template_decl) + placeholder = push_template_decl (placeholder); + pushdecl (placeholder); + cp_finish_decl (placeholder, NULL_TREE, 0, NULL_TREE, 0); + DECL_ARTIFICIAL (placeholder) = 1; + TREE_USED (placeholder) = 1; + + while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL)) + { + pragma_omp_clause c_kind = cp_parser_omp_clause_name (parser); + if (c_kind != PRAGMA_OMP_CLAUSE_MAP) + { + if (c_kind != PRAGMA_OMP_CLAUSE_NONE) + cp_parser_error (parser, "unexpected clause"); + finish_omp_structured_block (block); + goto fail; + } + maplist = cp_parser_omp_clause_map (parser, maplist, GOMP_MAP_UNSET); + if (maplist == NULL_TREE) + break; + } + + if (maplist == NULL_TREE) + { + cp_parser_error (parser, "missing % clause"); + finish_omp_structured_block (block); + goto fail; + } + + mapper = make_node (OMP_DECLARE_MAPPER); + TREE_TYPE (mapper) = type; + OMP_DECLARE_MAPPER_ID (mapper) = mapper_name; + OMP_DECLARE_MAPPER_DECL (mapper) = placeholder; + OMP_DECLARE_MAPPER_CLAUSES (mapper) = maplist; + + finish_omp_structured_block (block); + + DECL_INITIAL (vardecl) = mapper; + + if (current_class_type) + { + if (processing_template_decl) + { + retrofit_lang_decl (vardecl); + SET_DECL_VAR_DECLARED_INLINE_P (vardecl); + } + finish_static_data_member_decl (vardecl, mapper, + /*init_const_expr_p=*/true, NULL_TREE, 0); + finish_member_declaration (vardecl); + } + else if (processing_template_decl && block_scope) + add_decl_expr (vardecl); + else + pushdecl (vardecl); + + cp_check_omp_declare_mapper (vardecl); + + cp_parser_require_pragma_eol (parser, pragma_tok); + return; + +fail: + cp_parser_skip_to_pragma_eol (parser, pragma_tok); +} + /* OpenMP 4.0 #pragma omp declare simd declare-simd-clauses[optseq] new-line #pragma omp declare reduction (reduction-id : typename-list : expression) \ @@ -47942,6 +48200,12 @@ cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok, context); return false; } + if (strcmp (p, "mapper") == 0) + { + cp_lexer_consume_token (parser->lexer); + cp_parser_omp_declare_mapper (parser, pragma_tok, context); + return false; + } if (!flag_openmp) /* flag_openmp_simd */ { cp_parser_skip_to_pragma_eol (parser, pragma_tok); @@ -47955,7 +48219,7 @@ cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok, } } cp_parser_error (parser, "expected %, %, " - "% or %"); + "%, % or %"); cp_parser_require_pragma_eol (parser, pragma_tok); return false; } diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index 6cf3ed8d63ad..fadc4c6fa369 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -15320,6 +15320,13 @@ tsubst_decl (tree t, tree args, tsubst_flags_t complain) TYPE_ALIGN (TREE_TYPE (t))); } + if (flag_openmp + && VAR_P (t) + && DECL_LANG_SPECIFIC (t) + && DECL_OMP_DECLARE_MAPPER_P (t) + && strchr (IDENTIFIER_POINTER (DECL_NAME (t)), '~') == NULL) + DECL_NAME (r) = omp_mapper_id (DECL_NAME (t), TREE_TYPE (r)); + layout_decl (r, 0); } break; @@ -18146,6 +18153,8 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort, new_clauses = nreverse (new_clauses); if (ort != C_ORT_OMP_DECLARE_SIMD) { + if (ort == C_ORT_OMP_TARGET) + new_clauses = c_omp_instantiate_mappers (new_clauses); new_clauses = finish_omp_clauses (new_clauses, ort); if (linear_no_step) for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc)) @@ -19698,6 +19707,22 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl) } break; + case OMP_DECLARE_MAPPER: + { + t = copy_node (t); + + tree decl = OMP_DECLARE_MAPPER_DECL (t); + decl = tsubst (decl, args, complain, in_decl); + tree type = tsubst (TREE_TYPE (t), args, complain, in_decl); + tree clauses = OMP_DECLARE_MAPPER_CLAUSES (t); + clauses = tsubst_omp_clauses (clauses, C_ORT_OMP_DECLARE_SIMD, args, + complain, in_decl); + TREE_TYPE (t) = type; + OMP_DECLARE_MAPPER_DECL (t) = decl; + OMP_DECLARE_MAPPER_CLAUSES (t) = clauses; + RETURN (t); + } + case TRANSACTION_EXPR: { int flags = 0; @@ -26949,7 +26974,9 @@ instantiate_decl (tree d, bool defer_ok, bool expl_inst_class_mem_p) || (external_p && VAR_P (d)) /* Handle here a deleted function too, avoid generating its body (c++/61080). */ - || deleted_p) + || deleted_p + /* We need the initializer for an OpenMP declare mapper. */ + || (VAR_P (d) && DECL_LANG_SPECIFIC (d) && DECL_OMP_DECLARE_MAPPER_P (d))) { /* The definition of the static data member is now required so we must substitute the initializer. */ diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index f43816dafff0..bb5ee13aba19 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -45,6 +45,7 @@ along with GCC; see the file COPYING3. If not see #include "gomp-constants.h" #include "predict.h" #include "memmodel.h" +#include "gimplify.h" /* There routines provide a modular interface to perform many parsing operations. They may therefore be used during actual parsing, or @@ -5972,6 +5973,98 @@ omp_reduction_lookup (location_t loc, tree id, tree type, tree *baselinkp, return id; } +/* Return identifier to look up for omp declare mapper. */ + +tree +omp_mapper_id (tree mapper_id, tree type) +{ + const char *p = NULL; + const char *m = NULL; + + if (mapper_id == NULL_TREE) + p = ""; + else if (TREE_CODE (mapper_id) == IDENTIFIER_NODE) + p = IDENTIFIER_POINTER (mapper_id); + else + return error_mark_node; + + if (type != NULL_TREE) + m = mangle_type_string (TYPE_MAIN_VARIANT (type)); + + const char prefix[] = "omp declare mapper "; + size_t lenp = sizeof (prefix); + if (strncmp (p, prefix, lenp - 1) == 0) + lenp = 1; + size_t len = strlen (p); + size_t lenm = m ? strlen (m) + 1 : 0; + char *name = XALLOCAVEC (char, lenp + len + lenm); + memcpy (name, prefix, lenp - 1); + memcpy (name + lenp - 1, p, len + 1); + if (m) + { + name[lenp + len - 1] = '~'; + memcpy (name + lenp + len, m, lenm); + } + return get_identifier (name); +} + +tree +cxx_omp_mapper_lookup (tree id, tree type) +{ + if (TREE_CODE (type) != RECORD_TYPE + && TREE_CODE (type) != UNION_TYPE) + return NULL_TREE; + id = omp_mapper_id (id, type); + return lookup_name (id); +} + +tree +cxx_omp_extract_mapper_directive (tree vardecl) +{ + gcc_assert (TREE_CODE (vardecl) == VAR_DECL); + + /* Instantiate the decl if we haven't already. */ + mark_used (vardecl); + tree body = DECL_INITIAL (vardecl); + + if (TREE_CODE (body) == STATEMENT_LIST) + { + tree_stmt_iterator tsi = tsi_start (body); + gcc_assert (TREE_CODE (tsi_stmt (tsi)) == DECL_EXPR); + tsi_next (&tsi); + body = tsi_stmt (tsi); + } + + gcc_assert (TREE_CODE (body) == OMP_DECLARE_MAPPER); + + return body; +} + +/* For now we can handle singleton OMP_ARRAY_SECTIONs with custom mappers, but + nothing more complicated. */ + +tree +cxx_omp_map_array_section (location_t loc, tree t) +{ + tree low = TREE_OPERAND (t, 1); + tree len = TREE_OPERAND (t, 2); + + if (len && integer_onep (len)) + { + t = TREE_OPERAND (t, 0); + + if (!low) + low = integer_zero_node; + + if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE) + t = convert_from_reference (t); + + t = build_array_ref (loc, t, low); + } + + return t; +} + /* Helper function for cp_parser_omp_declare_reduction_exprs and tsubst_omp_udr. Remove CLEANUP_STMT for data (omp_priv variable). @@ -6453,6 +6546,29 @@ finish_omp_reduction_clause (tree c, bool *need_default_ctor, bool *need_dtor) return false; } +/* Check an instance of an "omp declare mapper" function. */ + +bool +cp_check_omp_declare_mapper (tree udm) +{ + tree type = TREE_TYPE (udm); + location_t loc = DECL_SOURCE_LOCATION (udm); + + if (type == error_mark_node) + return false; + + if (!processing_template_decl + && TREE_CODE (type) != RECORD_TYPE + && TREE_CODE (type) != UNION_TYPE) + { + error_at (loc, "%qT is not a struct, union or class type in " + "%<#pragma omp declare mapper%>", type); + return false; + } + + return true; +} + /* Called from finish_struct_1. linear(this) or linear(this:step) clauses might not be finalized yet because the class has been incomplete when parsing #pragma omp declare simd methods. Fix those up now. */ @@ -7993,6 +8109,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_MAP: if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved) goto move_implicit; + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME) + { + remove = true; + break; + } /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: @@ -9442,6 +9564,8 @@ struct omp_target_walk_data /* Local variables declared inside a BIND_EXPR, used to filter out such variables when recording lambda_objects_accessed. */ hash_set local_decls; + + omp_mapper_list *mappers; }; /* Helper function of finish_omp_target_clauses, called via @@ -9455,6 +9579,7 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr) struct omp_target_walk_data *data = (struct omp_target_walk_data *) ptr; tree current_object = data->current_object; tree current_closure = data->current_closure; + omp_mapper_list *mlist = data->mappers; /* References inside of these expression codes shouldn't incur any form of mapping, so return early. */ @@ -9468,6 +9593,27 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr) if (TREE_CODE (t) == OMP_CLAUSE) return NULL_TREE; + if (!processing_template_decl) + { + tree aggr_type = NULL_TREE; + + if (TREE_CODE (t) == COMPONENT_REF + && AGGREGATE_TYPE_P (TREE_TYPE (TREE_OPERAND (t, 0)))) + aggr_type = TREE_TYPE (TREE_OPERAND (t, 0)); + else if ((TREE_CODE (t) == VAR_DECL + || TREE_CODE (t) == PARM_DECL + || TREE_CODE (t) == RESULT_DECL) + && AGGREGATE_TYPE_P (TREE_TYPE (t))) + aggr_type = TREE_TYPE (t); + + if (aggr_type) + { + tree mapper_fn = cxx_omp_mapper_lookup (NULL_TREE, aggr_type); + if (mapper_fn) + mlist->add_mapper (NULL_TREE, aggr_type, mapper_fn); + } + } + if (current_object) { tree this_expr = TREE_OPERAND (current_object, 0); @@ -9570,10 +9716,48 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr) else data.current_closure = NULL_TREE; - cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r, &data); - auto_vec new_clauses; + if (!processing_template_decl) + { + hash_set > seen_types; + auto_vec mapper_fns; + omp_mapper_list mlist (&seen_types, &mapper_fns); + data.mappers = &mlist; + + cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r, + &data); + + unsigned int i; + tree mapper_fn; + FOR_EACH_VEC_ELT (mapper_fns, i, mapper_fn) + c_omp_find_nested_mappers (&mlist, mapper_fn); + + FOR_EACH_VEC_ELT (mapper_fns, i, mapper_fn) + { + tree mapper = cxx_omp_extract_mapper_directive (mapper_fn); + if (mapper == error_mark_node) + continue; + tree mapper_name = OMP_DECLARE_MAPPER_ID (mapper); + tree decl = OMP_DECLARE_MAPPER_DECL (mapper); + if (BASELINK_P (mapper_fn)) + mapper_fn = BASELINK_FUNCTIONS (mapper_fn); + + tree c = build_omp_clause (loc, OMP_CLAUSE__MAPPER_BINDING_); + OMP_CLAUSE__MAPPER_BINDING__ID (c) = mapper_name; + OMP_CLAUSE__MAPPER_BINDING__DECL (c) = decl; + OMP_CLAUSE__MAPPER_BINDING__MAPPER (c) = mapper_fn; + + new_clauses.safe_push (c); + } + } + else + { + data.mappers = NULL; + cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r, + &data); + } + tree omp_target_this_expr = NULL_TREE; tree *explicit_this_deref_map = NULL; if (data.this_expr_accessed) diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc index bc2b2188eea7..c1015078d17f 100644 --- a/gcc/fortran/parse.cc +++ b/gcc/fortran/parse.cc @@ -27,6 +27,9 @@ along with GCC; see the file COPYING3. If not see #include "match.h" #include "parse.h" #include "tree-core.h" +#include "tree.h" +#include "fold-const.h" +#include "tree-hash-traits.h" #include "omp-general.h" /* Current statement label. Zero means no statement label. Because new_st diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 97f69a9fa935..050282ce0587 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -219,6 +219,7 @@ struct gimplify_omp_ctx { struct gimplify_omp_ctx *outer_context; splay_tree variables; + hash_map, tree> *implicit_mappers; hash_set *privatized_types; tree clauses; /* Iteration variables in an OMP_FOR. */ @@ -452,6 +453,7 @@ new_omp_context (enum omp_region_type region_type) c = XCNEW (struct gimplify_omp_ctx); c->outer_context = gimplify_omp_ctxp; c->variables = splay_tree_new (splay_tree_compare_decl_uid, 0, 0); + c->implicit_mappers = new hash_map, tree>; c->privatized_types = new hash_set; c->location = input_location; c->region_type = region_type; @@ -475,6 +477,7 @@ delete_omp_context (struct gimplify_omp_ctx *c) { splay_tree_delete (c->variables); delete c->privatized_types; + delete c->implicit_mappers; c->loop_iter_var.release (); XDELETE (c); } @@ -11539,6 +11542,227 @@ error_out: return success; } +struct instantiate_mapper_info +{ + tree *mapper_clauses_p; + struct gimplify_omp_ctx *omp_ctx; + gimple_seq *pre_p; +}; + +/* Helper function for omp_instantiate_mapper. */ + +static tree +remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data) +{ + copy_body_data *id = (copy_body_data *) data; + + if (DECL_P (*tp)) + { + tree replacement = remap_decl (*tp, id); + if (*tp != replacement) + { + *tp = unshare_expr (replacement); + *walk_subtrees = 0; + } + } + + return NULL_TREE; +} + +/* A copy_decl implementation (for use with tree-inline.cc functions) that + only transform decls or SSA names that are part of a map we already + prepared. */ + +static tree +omp_mapper_copy_decl (tree var, copy_body_data *cb) +{ + tree *repl = cb->decl_map->get (var); + + if (repl) + return *repl; + + return var; +} + +static tree * +omp_instantiate_mapper (gimple_seq *pre_p, + hash_map, tree> *implicit_mappers, + tree mapperfn, tree expr, enum gomp_map_kind outer_kind, + tree *mapper_clauses_p) +{ + tree mapper_name = NULL_TREE; + tree mapper = lang_hooks.decls.omp_extract_mapper_directive (mapperfn); + gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER); + + tree clause = OMP_DECLARE_MAPPER_CLAUSES (mapper); + tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper); + + /* The "extraction map" is used to map the mapper variable in the "declare + mapper" directive, and also any temporary variables that have been created + as part of expanding the mapper function's body (which are expanded as a + "bind" expression in the pre_p sequence). */ + hash_map extraction_map; + + extraction_map.put (dummy_var, expr); + extraction_map.put (expr, expr); + + /* This copy_body_data is only used to remap the decls in the + OMP_DECLARE_MAPPER tree node expansion itself. All relevant decls should + already be in the current function. */ + copy_body_data id; + memset (&id, 0, sizeof (id)); + id.src_fn = current_function_decl; + id.dst_fn = current_function_decl; + id.src_cfun = cfun; + id.decl_map = &extraction_map; + id.copy_decl = omp_mapper_copy_decl; + id.transform_call_graph_edges = CB_CGE_DUPLICATE; // ??? + id.transform_new_cfg = true; // ??? + + for (; clause; clause = OMP_CLAUSE_CHAIN (clause)) + { + enum gomp_map_kind map_kind = OMP_CLAUSE_MAP_KIND (clause); + tree *nested_mapper_p = NULL; + + if (map_kind == GOMP_MAP_PUSH_MAPPER_NAME) + { + mapper_name = OMP_CLAUSE_DECL (clause); + continue; + } + else if (map_kind == GOMP_MAP_POP_MAPPER_NAME) + { + mapper_name = NULL_TREE; + continue; + } + + tree decl = OMP_CLAUSE_DECL (clause); + tree unshared, type; + bool nonunit_array_with_mapper = false; + + if (TREE_CODE (decl) == OMP_ARRAY_SECTION) + { + location_t loc = OMP_CLAUSE_LOCATION (clause); + tree tmp = lang_hooks.decls.omp_map_array_section (loc, decl); + if (tmp == decl) + { + unshared = unshare_expr (clause); + nonunit_array_with_mapper = true; + type = TREE_TYPE (TREE_TYPE (decl)); + } + else + { + unshared = build_omp_clause (OMP_CLAUSE_LOCATION (clause), + OMP_CLAUSE_CODE (clause)); + OMP_CLAUSE_DECL (unshared) = tmp; + OMP_CLAUSE_SIZE (unshared) + = DECL_P (tmp) ? DECL_SIZE_UNIT (tmp) + : TYPE_SIZE_UNIT (TREE_TYPE (tmp)); + type = TREE_TYPE (tmp); + } + } + else + { + unshared = unshare_expr (clause); + type = TREE_TYPE (decl); + } + + walk_tree (&unshared, remap_mapper_decl_1, &id, NULL); + + if (OMP_CLAUSE_MAP_KIND (unshared) == GOMP_MAP_UNSET) + OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind); + + decl = OMP_CLAUSE_DECL (unshared); + type = TYPE_MAIN_VARIANT (type); + + nested_mapper_p = implicit_mappers->get ({ mapper_name, type }); + + if (nested_mapper_p && *nested_mapper_p != mapperfn) + { + if (nonunit_array_with_mapper) + { + sorry ("user-defined mapper with non-unit length array section"); + continue; + } + + if (map_kind == GOMP_MAP_UNSET) + map_kind = outer_kind; + + mapper_clauses_p + = omp_instantiate_mapper (pre_p, implicit_mappers, + *nested_mapper_p, decl, map_kind, + mapper_clauses_p); + continue; + } + + *mapper_clauses_p = unshared; + mapper_clauses_p = &OMP_CLAUSE_CHAIN (unshared); + } + + return mapper_clauses_p; +} + +static int +omp_instantiate_implicit_mappers (splay_tree_node n, void *data) +{ + tree decl = (tree) n->key; + instantiate_mapper_info *im_info = (instantiate_mapper_info *) data; + gimplify_omp_ctx *ctx = im_info->omp_ctx; + tree *mapper_p = NULL; + tree type = TREE_TYPE (decl); + bool ref_p = false; + unsigned flags = n->value; + + if (flags & (GOVD_EXPLICIT | GOVD_LOCAL)) + return 0; + if ((flags & GOVD_SEEN) == 0) + return 0; + + if (TREE_CODE (type) == REFERENCE_TYPE) + { + ref_p = true; + type = TREE_TYPE (type); + } + + type = TYPE_MAIN_VARIANT (type); + + if (DECL_P (decl) && type && AGGREGATE_TYPE_P (type)) + { + gcc_assert (ctx); + mapper_p = ctx->implicit_mappers->get ({ NULL_TREE, type }); + } + + /* If we already have clauses pertaining to a struct variable, then we don't + want to implicitly invoke a user-defined mapper. */ + bool handled_p = false; + splay_tree_node an + = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); + if (an) + { + unsigned flags = an->value; + if ((flags & GOVD_EXPLICIT) != 0 + && AGGREGATE_TYPE_P (TREE_TYPE (decl))) + handled_p = true; + } + + if (mapper_p && !handled_p) + { + /* If we have a reference, map the pointed-to object rather than the + reference itself. */ + if (ref_p) + decl = build_fold_indirect_ref (decl); + + im_info->mapper_clauses_p + = omp_instantiate_mapper (im_info->pre_p, ctx->implicit_mappers, + *mapper_p, decl, GOMP_MAP_TOFROM, + im_info->mapper_clauses_p); + /* Make sure we don't map the same variable implicitly in + gimplify_adjust_omp_clauses_1 also. */ + n->value |= GOVD_EXPLICIT; + } + + return 0; +} + /* Scan the OMP clauses in *LIST_P, installing mappings into a new and previous omp contexts. */ @@ -11550,7 +11774,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, using namespace omp_addr_tokenizer; struct gimplify_omp_ctx *ctx, *outer_ctx; tree c; - tree *prev_list_p = NULL, *orig_list_p = list_p; + tree *orig_list_p = list_p; int handled_depend_iterators = -1; int nowait = -1; @@ -11581,75 +11805,30 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } - if (code == OMP_TARGET - || code == OMP_TARGET_DATA - || code == OMP_TARGET_ENTER_DATA - || code == OMP_TARGET_EXIT_DATA) - { - vec *groups; - groups = omp_gather_mapping_groups (list_p); - if (groups) - { - hash_map *grpmap; - grpmap = omp_index_mapping_groups (groups); + vec *groups = omp_gather_mapping_groups (list_p); + hash_map *grpmap = NULL; + unsigned grpnum = 0; + tree *grp_start_p = NULL, grp_end = NULL_TREE; - omp_resolve_clause_dependencies (code, groups, grpmap); - omp_build_struct_sibling_lists (code, region_type, groups, &grpmap, - list_p); - - omp_mapping_group *outlist = NULL; - bool enter_exit = (code == OMP_TARGET_ENTER_DATA - || code == OMP_TARGET_EXIT_DATA); - - /* Topological sorting may fail if we have duplicate nodes, which - we should have detected and shown an error for already. Skip - sorting in that case. */ - if (seen_error ()) - goto failure; - - delete grpmap; - delete groups; - - /* Rebuild now we have struct sibling lists. */ - groups = omp_gather_mapping_groups (list_p); - grpmap = omp_index_mapping_groups (groups); - - outlist = omp_tsort_mapping_groups (groups, grpmap, enter_exit); - outlist = omp_segregate_mapping_groups (outlist); - list_p = omp_reorder_mapping_groups (groups, outlist, list_p); - - failure: - delete grpmap; - delete groups; - } - } - else if (region_type & ORT_ACC) - { - vec *groups; - groups = omp_gather_mapping_groups (list_p); - if (groups) - { - hash_map *grpmap; - grpmap = omp_index_mapping_groups (groups); - - oacc_resolve_clause_dependencies (groups, grpmap); - omp_build_struct_sibling_lists (code, region_type, groups, &grpmap, - list_p); - - delete groups; - delete grpmap; - } - } + if (groups) + grpmap = omp_index_mapping_groups (groups); while ((c = *list_p) != NULL) { bool remove = false; bool notice_outer = true; + bool map_descriptor; const char *check_non_private = NULL; unsigned int flags; tree decl; auto_vec addr_tokens; + if (grp_end && c == OMP_CLAUSE_CHAIN (grp_end)) + { + grp_start_p = NULL; + grp_end = NULL_TREE; + } + switch (OMP_CLAUSE_CODE (c)) { case OMP_CLAUSE_PRIVATE: @@ -11954,99 +12133,37 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, goto do_add; case OMP_CLAUSE_MAP: + if (!grp_start_p) + { + grp_start_p = list_p; + grp_end = (*groups)[grpnum].grp_end; + grpnum++; + } decl = OMP_CLAUSE_DECL (c); + if (error_operand_p (decl)) + { + remove = true; + break; + } + if (!omp_parse_expr (addr_tokens, decl)) { remove = true; break; } - if (error_operand_p (decl)) - remove = true; - switch (code) - { - case OMP_TARGET: - break; - case OACC_DATA: - if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE) - break; - goto check_firstprivate; - case OACC_ENTER_DATA: - case OACC_EXIT_DATA: - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH - && addr_tokens[0]->type == ARRAY_BASE) - remove = true; - /* FALLTHRU */ - case OMP_TARGET_DATA: - case OMP_TARGET_ENTER_DATA: - case OMP_TARGET_EXIT_DATA: - case OACC_HOST_DATA: - check_firstprivate: - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER - || (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) - /* For target {,enter ,exit }data only the array slice is - mapped, but not the pointer to it. */ - remove = true; - break; - default: - break; - } - /* For Fortran, not only the pointer to the data is mapped but also - the address of the pointer, the array descriptor etc.; for - 'exit data' - and in particular for 'delete:' - having an 'alloc:' - does not make sense. Likewise, for 'update' only transferring the - data itself is needed as the rest has been handled in previous - directives. However, for 'exit data', the array descriptor needs - to be delete; hence, we turn the MAP_TO_PSET into a MAP_DELETE. - - NOTE: Generally, it is not safe to perform "enter data" operations - on arrays where the data *or the descriptor* may go out of scope - before a corresponding "exit data" operation -- and such a - descriptor may be synthesized temporarily, e.g. to pass an - explicit-shape array to a function expecting an assumed-shape - argument. Performing "enter data" inside the called function - would thus be problematic. */ - - tree desc; - if (DECL_P (decl) - || (prev_list_p - && (desc = OMP_CLAUSE_CHAIN (*prev_list_p)) - && OMP_CLAUSE_CODE (desc) == OMP_CLAUSE_MAP - && (OMP_CLAUSE_MAP_KIND (desc) == GOMP_MAP_RELEASE - || OMP_CLAUSE_MAP_KIND (desc) == GOMP_MAP_DELETE) - && DECL_P (OMP_CLAUSE_DECL (desc)))) - { - /* We only do the transformations here for non-component - mappings. Array descriptors for component mappings are - handled in omp_accumulate_sibling_list. */ - if (code == OMP_TARGET_EXIT_DATA - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET) - { - enum gomp_map_kind k - = (OMP_CLAUSE_MAP_KIND (*prev_list_p) == GOMP_MAP_DELETE - ? GOMP_MAP_DELETE : GOMP_MAP_RELEASE); - OMP_CLAUSE_SET_MAP_KIND (c, k); - } - else if ((code == OMP_TARGET_EXIT_DATA - || code == OMP_TARGET_UPDATE) - && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)) - remove = true; - else if (code == OMP_TARGET_EXIT_DATA - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC - && OMP_CLAUSE_CHAIN (c) - && (OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) - == OMP_CLAUSE_MAP) - && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) - == GOMP_MAP_ATTACH_DETACH) - || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) - == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)) - && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL - (OMP_CLAUSE_CHAIN (c)))) == REFERENCE_TYPE) - OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_RELEASE); - } + if (code == OMP_TARGET_EXIT_DATA + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC + && OMP_CLAUSE_CHAIN (c) + && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP + && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) + == GOMP_MAP_ATTACH_DETACH) + || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)) + && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL + (OMP_CLAUSE_CHAIN (c)))) == REFERENCE_TYPE) + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_RELEASE); if (remove) break; @@ -12066,41 +12183,51 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, DECL_NAME (decl)); } } - if (OMP_CLAUSE_SIZE (c) == NULL_TREE) - OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl) - : TYPE_SIZE_UNIT (TREE_TYPE (decl)); - if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, - NULL, is_gimple_val, fb_rvalue) == GS_ERROR) + + map_descriptor = false; + + /* This condition checks if we're mapping an array descriptor that + isn't inside a derived type -- these have special handling, and + are not handled as structs in omp_build_struct_sibling_lists. + See that function for further details. */ + if (*grp_start_p != grp_end + && OMP_CLAUSE_CHAIN (*grp_start_p) + && OMP_CLAUSE_CHAIN (*grp_start_p) != grp_end) { - remove = true; - break; - } - else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER - || (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_REFERENCE) - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) - && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST) - { - OMP_CLAUSE_SIZE (c) - = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL, - false); - if ((region_type & ORT_TARGET) != 0) - omp_add_variable (ctx, OMP_CLAUSE_SIZE (c), - GOVD_FIRSTPRIVATE | GOVD_SEEN); + tree grp_mid = OMP_CLAUSE_CHAIN (*grp_start_p); + if (OMP_CLAUSE_CODE (grp_mid) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_TO_PSET + && DECL_P (OMP_CLAUSE_DECL (grp_mid))) + map_descriptor = true; } - if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT_UNORD) - && (addr_tokens[0]->type == STRUCTURE_BASE - || addr_tokens[0]->type == ARRAY_BASE) - && addr_tokens[0]->u.structure_base_kind == BASE_DECL) + /* Adding the decl for a struct access: we haven't created + GOMP_MAP_STRUCT nodes yet, so this statement needs to predict + whether they will be created in gimplify_adjust_omp_clauses. */ + if (c == grp_end + && addr_tokens[0]->type == STRUCTURE_BASE + && addr_tokens[0]->u.structure_base_kind == BASE_DECL + && !map_descriptor) { gcc_assert (addr_tokens[1]->type == ACCESS_METHOD); /* If we got to this struct via a chain of pointers, maybe we want to map it implicitly instead. */ if (omp_access_chain_p (addr_tokens, 1)) break; + omp_mapping_group *wholestruct; + if (!(region_type & ORT_ACC) + && omp_mapped_by_containing_struct (grpmap, + OMP_CLAUSE_DECL (c), + &wholestruct)) + break; decl = addr_tokens[1]->expr; + if (splay_tree_lookup (ctx->variables, (splay_tree_key) decl)) + break; + /* Standalone attach or detach clauses for a struct element + should not inhibit implicit mapping of the whole struct. */ + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + break; flags = GOVD_MAP | GOVD_EXPLICIT; gcc_assert (addr_tokens[1]->u.access_kind != ACCESS_DIRECT @@ -12108,14 +12235,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, goto do_add_decl; } - if (TREE_CODE (decl) == TARGET_EXPR) - { - if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL, - is_gimple_lvalue, fb_lvalue) - == GS_ERROR) - remove = true; - } - else if (!DECL_P (decl)) + if (!DECL_P (decl)) { tree d = decl, *pd; if (TREE_CODE (d) == ARRAY_REF) @@ -12138,56 +12258,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, pd = &TREE_OPERAND (decl, 0); decl = TREE_OPERAND (decl, 0); } - /* An "attach/detach" operation on an update directive should - behave as a GOMP_MAP_ALWAYS_POINTER. Beware that - unlike attach or detach map kinds, GOMP_MAP_ALWAYS_POINTER - depends on the previous mapping. */ - if (code == OACC_UPDATE - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) - OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + if (addr_tokens[0]->type == STRUCTURE_BASE + && addr_tokens[0]->u.structure_base_kind == BASE_DECL + && addr_tokens[1]->type == ACCESS_METHOD + && (addr_tokens[1]->u.access_kind == ACCESS_POINTER + || (addr_tokens[1]->u.access_kind + == ACCESS_POINTER_OFFSET)) + && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))) { - if (TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) - == ARRAY_TYPE) - remove = true; - else - { - gomp_map_kind k = ((code == OACC_EXIT_DATA - || code == OMP_TARGET_EXIT_DATA) - ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH); - OMP_CLAUSE_SET_MAP_KIND (c, k); - } - } - - tree cref = decl; - - while (TREE_CODE (cref) == ARRAY_REF) - cref = TREE_OPERAND (cref, 0); - - if (TREE_CODE (cref) == INDIRECT_REF) - cref = TREE_OPERAND (cref, 0); - - if (TREE_CODE (cref) == COMPONENT_REF) - { - tree base = cref; - while (base && !DECL_P (base)) - { - tree innerbase = omp_get_base_pointer (base); - if (!innerbase) - break; - base = innerbase; - } - if (base - && DECL_P (base) - && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) - && POINTER_TYPE_P (TREE_TYPE (base))) - { - splay_tree_node n - = splay_tree_lookup (ctx->variables, - (splay_tree_key) base); - n->value |= GOVD_SEEN; - } + tree base = addr_tokens[1]->expr; + splay_tree_node n + = splay_tree_lookup (ctx->variables, + (splay_tree_key) base); + n->value |= GOVD_SEEN; } if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c)) @@ -12298,27 +12382,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } } } - else if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, - fb_lvalue) == GS_ERROR) - { - remove = true; - break; - } - - if (!remove - && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER - && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH - && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET - && OMP_CLAUSE_CHAIN (c) - && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP - && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) - == GOMP_MAP_ALWAYS_POINTER) - || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) - == GOMP_MAP_ATTACH_DETACH) - || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) - == GOMP_MAP_TO_PSET))) - prev_list_p = list_p; - break; } flags = GOVD_MAP | GOVD_EXPLICIT; @@ -12326,43 +12389,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM) flags |= GOVD_MAP_ALWAYS_TO; - if ((code == OMP_TARGET - || code == OMP_TARGET_DATA - || code == OMP_TARGET_ENTER_DATA - || code == OMP_TARGET_EXIT_DATA) - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) - { - for (struct gimplify_omp_ctx *octx = outer_ctx; octx; - octx = octx->outer_context) - { - splay_tree_node n - = splay_tree_lookup (octx->variables, - (splay_tree_key) OMP_CLAUSE_DECL (c)); - /* If this is contained in an outer OpenMP region as a - firstprivate value, remove the attach/detach. */ - if (n && (n->value & GOVD_FIRSTPRIVATE)) - { - OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FIRSTPRIVATE_POINTER); - goto do_add; - } - } - - enum gomp_map_kind map_kind = (code == OMP_TARGET_EXIT_DATA - ? GOMP_MAP_DETACH - : GOMP_MAP_ATTACH); - OMP_CLAUSE_SET_MAP_KIND (c, map_kind); - } - else if ((code == OACC_ENTER_DATA - || code == OACC_EXIT_DATA - || code == OACC_PARALLEL) - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) - { - enum gomp_map_kind map_kind = (code == OACC_EXIT_DATA - ? GOMP_MAP_DETACH - : GOMP_MAP_ATTACH); - OMP_CLAUSE_SET_MAP_KIND (c, map_kind); - } - goto do_add; case OMP_CLAUSE_AFFINITY: @@ -12451,6 +12477,17 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } goto do_notice; + case OMP_CLAUSE__MAPPER_BINDING_: + { + tree name = OMP_CLAUSE__MAPPER_BINDING__ID (c); + tree var = OMP_CLAUSE__MAPPER_BINDING__DECL (c); + tree type = TYPE_MAIN_VARIANT (TREE_TYPE (var)); + tree fndecl = OMP_CLAUSE__MAPPER_BINDING__MAPPER (c); + ctx->implicit_mappers->put ({ name, type }, fndecl); + remove = true; + break; + } + case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_ADDR: flags = GOVD_EXPLICIT; @@ -12998,6 +13035,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, list_p = &OMP_CLAUSE_CHAIN (c); } + if (groups) + { + delete grpmap; + delete groups; + } + ctx->clauses = *orig_list_p; gimplify_omp_ctxp = ctx; } @@ -13465,15 +13508,99 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, } } + if (code == OMP_TARGET + || code == OMP_TARGET_DATA + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA) + { + tree mapper_clauses = NULL_TREE; + instantiate_mapper_info im_info; + + im_info.mapper_clauses_p = &mapper_clauses; + im_info.omp_ctx = ctx; + im_info.pre_p = pre_p; + + splay_tree_foreach (ctx->variables, + omp_instantiate_implicit_mappers, + (void *) &im_info); + + if (mapper_clauses) + { + mapper_clauses + = lang_hooks.decls.omp_finish_mapper_clauses (mapper_clauses); + + /* Stick the implicitly-expanded mapper clauses at the end of the + clause list. */ + tree *tail = list_p; + while (*tail) + tail = &OMP_CLAUSE_CHAIN (*tail); + *tail = mapper_clauses; + } + + vec *groups; + groups = omp_gather_mapping_groups (list_p); + hash_map *grpmap = NULL; + + if (groups) + { + grpmap = omp_index_mapping_groups (groups); + + omp_resolve_clause_dependencies (code, groups, grpmap); + omp_build_struct_sibling_lists (code, ctx->region_type, groups, + &grpmap, list_p); + + omp_mapping_group *outlist = NULL; + + delete grpmap; + delete groups; + + /* Rebuild now we have struct sibling lists. */ + groups = omp_gather_mapping_groups (list_p); + grpmap = omp_index_mapping_groups (groups); + + bool enter_exit = (code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA); + + outlist = omp_tsort_mapping_groups (groups, grpmap, enter_exit); + outlist = omp_segregate_mapping_groups (outlist); + list_p = omp_reorder_mapping_groups (groups, outlist, list_p); + + delete grpmap; + delete groups; + } + } + else if (ctx->region_type & ORT_ACC) + { + vec *groups; + groups = omp_gather_mapping_groups (list_p); + if (groups) + { + hash_map *grpmap; + grpmap = omp_index_mapping_groups (groups); + + oacc_resolve_clause_dependencies (groups, grpmap); + omp_build_struct_sibling_lists (code, ctx->region_type, groups, + &grpmap, list_p); + + delete groups; + delete grpmap; + } + } + tree attach_list = NULL_TREE; tree *attach_tail = &attach_list; + tree *grp_start_p = NULL, grp_end = NULL_TREE; + while ((c = *list_p) != NULL) { splay_tree_node n; bool remove = false; bool move_attach = false; + if (grp_end && c == OMP_CLAUSE_CHAIN (grp_end)) + grp_end = NULL_TREE; + switch (OMP_CLAUSE_CODE (c)) { case OMP_CLAUSE_FIRSTPRIVATE: @@ -13628,26 +13755,67 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, break; case OMP_CLAUSE_MAP: - if (code == OMP_TARGET_EXIT_DATA - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER) + decl = OMP_CLAUSE_DECL (c); + if (!grp_end) { + grp_start_p = list_p; + grp_end = *omp_group_last (grp_start_p); + } + switch (code) + { + case OMP_TARGET: + break; + case OACC_DATA: + if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE) + break; + goto check_firstprivate; + case OACC_ENTER_DATA: + case OACC_EXIT_DATA: + case OMP_TARGET_DATA: + case OMP_TARGET_ENTER_DATA: + case OMP_TARGET_EXIT_DATA: + case OACC_HOST_DATA: + check_firstprivate: + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) + /* For target {,enter ,exit }data only the array slice is + mapped, but not the pointer to it. */ + remove = true; + if (code == OMP_TARGET_EXIT_DATA + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER) + remove = true; + break; + default: + break; + } + if (remove) + break; + if (OMP_CLAUSE_SIZE (c) == NULL_TREE) + OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl) + : TYPE_SIZE_UNIT (TREE_TYPE (decl)); + gimplify_omp_ctxp = ctx->outer_context; + if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL, is_gimple_val, + fb_rvalue) == GS_ERROR) + { + gimplify_omp_ctxp = ctx; remove = true; break; } - /* If we have a target region, we can push all the attaches to the - end of the list (we may have standalone "attach" operations - synthesized for GOMP_MAP_STRUCT nodes that must be processed after - the attachment point AND the pointed-to block have been mapped). - If we have something else, e.g. "enter data", we need to keep - "attach" nodes together with the previous node they attach to so - that separate "exit data" operations work properly (see - libgomp/target.c). */ - if ((ctx->region_type & ORT_TARGET) != 0 - && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH - || (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))) - move_attach = true; - decl = OMP_CLAUSE_DECL (c); + else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST) + { + OMP_CLAUSE_SIZE (c) + = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL, + false); + if ((ctx->region_type & ORT_TARGET) != 0) + omp_add_variable (ctx, OMP_CLAUSE_SIZE (c), + GOVD_FIRSTPRIVATE | GOVD_SEEN); + } + gimplify_omp_ctxp = ctx; /* Data clauses associated with reductions must be compatible with present_or_copy. Warn and adjust the clause if that is not the case. */ @@ -13684,7 +13852,54 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, remove = true; break; } - if (!DECL_P (decl)) + /* For Fortran, not only the pointer to the data is mapped but also + the address of the pointer, the array descriptor etc.; for + 'exit data' - and in particular for 'delete:' - having an 'alloc:' + does not make sense. Likewise, for 'update' only transferring the + data itself is needed as the rest has been handled in previous + directives. However, for 'exit data', the array descriptor needs + to be deleted; hence, we turn the MAP_TO_PSET into a MAP_DELETE. + + NOTE: Generally, it is not safe to perform "enter data" operations + on arrays where the data *or the descriptor* may go out of scope + before a corresponding "exit data" operation -- and such a + descriptor may be synthesized temporarily, e.g. to pass an + explicit-shape array to a function expecting an assumed-shape + argument. Performing "enter data" inside the called function + would thus be problematic. */ + if (code == OMP_TARGET_EXIT_DATA + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET + && grp_start_p + && OMP_CLAUSE_CHAIN (*grp_start_p) == c) + OMP_CLAUSE_SET_MAP_KIND (c, OMP_CLAUSE_MAP_KIND (*grp_start_p) + == GOMP_MAP_DELETE + ? GOMP_MAP_DELETE : GOMP_MAP_RELEASE); + else if ((code == OMP_TARGET_EXIT_DATA || code == OMP_TARGET_UPDATE) + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)) + { + remove = true; + break; + } + else if (code == OMP_TARGET_EXIT_DATA + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC + && OMP_CLAUSE_CHAIN (c) + && (OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) + == OMP_CLAUSE_MAP) + && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) + == GOMP_MAP_ATTACH_DETACH) + || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)) + && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL + (OMP_CLAUSE_CHAIN (c)))) == REFERENCE_TYPE) + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_RELEASE); + if (TREE_CODE (decl) == TARGET_EXPR) + { + if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL, + is_gimple_lvalue, fb_lvalue) == GS_ERROR) + remove = true; + } + else if (!DECL_P (decl)) { if ((ctx->region_type & ORT_TARGET) != 0 && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) @@ -13707,8 +13922,123 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, } } } + + tree d = decl, *pd; + if (TREE_CODE (d) == ARRAY_REF) + { + while (TREE_CODE (d) == ARRAY_REF) + d = TREE_OPERAND (d, 0); + if (TREE_CODE (d) == COMPONENT_REF + && TREE_CODE (TREE_TYPE (d)) == ARRAY_TYPE) + decl = d; + } + pd = &OMP_CLAUSE_DECL (c); + if (d == decl + && TREE_CODE (decl) == INDIRECT_REF + && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF + && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) + == REFERENCE_TYPE) + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)) + { + pd = &TREE_OPERAND (decl, 0); + decl = TREE_OPERAND (decl, 0); + } + + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + switch (code) + { + case OACC_ENTER_DATA: + case OACC_EXIT_DATA: + if (TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) + == ARRAY_TYPE) + remove = true; + else if (code == OACC_ENTER_DATA) + goto change_to_attach; + /* Fallthrough. */ + case OMP_TARGET_EXIT_DATA: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DETACH); + break; + case OACC_UPDATE: + /* An "attach/detach" operation on an update directive + should behave as a GOMP_MAP_ALWAYS_POINTER. Note that + both GOMP_MAP_ATTACH_DETACH and GOMP_MAP_ALWAYS_POINTER + kinds depend on the previous mapping (for non-TARGET + regions). */ + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); + break; + default: + change_to_attach: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH); + if ((ctx->region_type & ORT_TARGET) != 0) + move_attach = true; + } + else if ((ctx->region_type & ORT_TARGET) != 0 + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))) + move_attach = true; + + /* If we have e.g. map(struct: *var), don't gimplify the + argument since omp-low.cc wants to see the decl itself. */ + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT) + break; + + /* We've already partly gimplified this in + gimplify_scan_omp_clauses. Don't do any more. */ + if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c)) + break; + + gimplify_omp_ctxp = ctx->outer_context; + if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, + fb_lvalue) == GS_ERROR) + remove = true; + gimplify_omp_ctxp = ctx; + break; } + + if ((code == OMP_TARGET + || code == OMP_TARGET_DATA + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA) + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + { + bool firstprivatize = false; + + for (struct gimplify_omp_ctx *octx = ctx->outer_context; octx; + octx = octx->outer_context) + { + splay_tree_node n + = splay_tree_lookup (octx->variables, + (splay_tree_key) OMP_CLAUSE_DECL (c)); + /* If this is contained in an outer OpenMP region as a + firstprivate value, remove the attach/detach. */ + if (n && (n->value & GOVD_FIRSTPRIVATE)) + { + firstprivatize = true; + break; + } + } + + enum gomp_map_kind map_kind; + if (firstprivatize) + map_kind = GOMP_MAP_FIRSTPRIVATE_POINTER; + else if (code == OMP_TARGET_EXIT_DATA) + map_kind = GOMP_MAP_DETACH; + else + map_kind = GOMP_MAP_ATTACH; + OMP_CLAUSE_SET_MAP_KIND (c, map_kind); + } + else if ((ctx->region_type & ORT_ACC) != 0 + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + { + enum gomp_map_kind map_kind = (code == OACC_EXIT_DATA + ? GOMP_MAP_DETACH + : GOMP_MAP_ATTACH); + OMP_CLAUSE_SET_MAP_KIND (c, map_kind); + } + n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); if ((ctx->region_type & ORT_TARGET) != 0 && !(n->value & GOVD_SEEN) @@ -13783,6 +14113,21 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, || ((n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)) == 0)); } + + /* If we have a target region, we can push all the attaches to the + end of the list (we may have standalone "attach" operations + synthesized for GOMP_MAP_STRUCT nodes that must be processed after + the attachment point AND the pointed-to block have been mapped). + If we have something else, e.g. "enter data", we need to keep + "attach" nodes together with the previous node they attach to so + that separate "exit data" operations work properly (see + libgomp/target.c). */ + if ((ctx->region_type & ORT_TARGET) != 0 + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))) + move_attach = true; + break; case OMP_CLAUSE_TO: @@ -16990,6 +17335,15 @@ gimplify_omp_ordered (tree expr, gimple_seq body) return gimple_build_omp_ordered (body, OMP_ORDERED_CLAUSES (expr)); } +/* Gimplify an OMP_DECLARE_MAPPER node (by just removing it). */ + +static enum gimplify_status +gimplify_omp_declare_mapper (tree *expr_p) +{ + *expr_p = NULL_TREE; + return GS_ALL_DONE; +} + /* Convert the GENERIC expression tree *EXPR_P to GIMPLE. If the expression produces a value to be used as an operand inside a GIMPLE statement, the value will be stored back in *EXPR_P. This value will @@ -17912,6 +18266,10 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, ret = gimplify_omp_atomic (expr_p, pre_p); break; + case OMP_DECLARE_MAPPER: + ret = gimplify_omp_declare_mapper (expr_p); + break; + case TRANSACTION_EXPR: ret = gimplify_transaction (expr_p, pre_p); break; diff --git a/gcc/langhooks-def.h b/gcc/langhooks-def.h index 0f1bfbf255c9..fa117884fc90 100644 --- a/gcc/langhooks-def.h +++ b/gcc/langhooks-def.h @@ -85,6 +85,10 @@ extern enum omp_clause_defaultmap_kind lhd_omp_predetermined_mapping (tree); extern tree lhd_omp_assignment (tree, tree, tree); extern void lhd_omp_finish_clause (tree, gimple_seq *, bool); extern tree lhd_omp_array_size (tree, gimple_seq *); +extern tree lhd_omp_finish_mapper_clauses (tree); +extern tree lhd_omp_mapper_lookup (tree, tree); +extern tree lhd_omp_extract_mapper_directive (tree); +extern tree lhd_omp_map_array_section (location_t, tree); struct gimplify_omp_ctx; extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *, tree); @@ -273,6 +277,11 @@ extern tree lhd_unit_size_without_reusable_padding (tree); #define LANG_HOOKS_OMP_CLAUSE_LINEAR_CTOR NULL #define LANG_HOOKS_OMP_CLAUSE_DTOR hook_tree_tree_tree_null #define LANG_HOOKS_OMP_FINISH_CLAUSE lhd_omp_finish_clause +#define LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES lhd_omp_finish_mapper_clauses +#define LANG_HOOKS_OMP_MAPPER_LOOKUP lhd_omp_mapper_lookup +#define LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE \ + lhd_omp_extract_mapper_directive +#define LANG_HOOKS_OMP_MAP_ARRAY_SECTION lhd_omp_map_array_section #define LANG_HOOKS_OMP_ALLOCATABLE_P hook_bool_tree_false #define LANG_HOOKS_OMP_SCALAR_P lhd_omp_scalar_p #define LANG_HOOKS_OMP_SCALAR_TARGET_P hook_bool_tree_false @@ -307,6 +316,10 @@ extern tree lhd_unit_size_without_reusable_padding (tree); LANG_HOOKS_OMP_CLAUSE_LINEAR_CTOR, \ LANG_HOOKS_OMP_CLAUSE_DTOR, \ LANG_HOOKS_OMP_FINISH_CLAUSE, \ + LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES, \ + LANG_HOOKS_OMP_MAPPER_LOOKUP, \ + LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, \ + LANG_HOOKS_OMP_MAP_ARRAY_SECTION, \ LANG_HOOKS_OMP_ALLOCATABLE_P, \ LANG_HOOKS_OMP_SCALAR_P, \ LANG_HOOKS_OMP_SCALAR_TARGET_P, \ diff --git a/gcc/langhooks.cc b/gcc/langhooks.cc index 7d9756759900..649ee4be9429 100644 --- a/gcc/langhooks.cc +++ b/gcc/langhooks.cc @@ -642,6 +642,41 @@ lhd_omp_array_size (tree, gimple_seq *) return NULL_TREE; } +/* Finalize clause list C after expanding custom mappers for implicitly-mapped + variables. */ + +tree +lhd_omp_finish_mapper_clauses (tree c) +{ + return c; +} + +/* Look up an OpenMP "declare mapper" mapper. */ + +tree +lhd_omp_mapper_lookup (tree, tree) +{ + return NULL_TREE; +} + +/* Given the representation used by the front-end to contain a mapper + directive, return the statement for the directive itself. */ + +tree +lhd_omp_extract_mapper_directive (tree) +{ + return error_mark_node; +} + +/* Return a simplified form for OMP_ARRAY_SECTION argument, or + error_mark_node if impossible. */ + +tree +lhd_omp_map_array_section (location_t, tree) +{ + return error_mark_node; +} + /* Return true if DECL is a scalar variable (for the purpose of implicit firstprivatization & mapping). Only if alloc_ptr_ok are allocatables and pointers accepted. */ diff --git a/gcc/langhooks.h b/gcc/langhooks.h index b1b2b0e10f09..1d5d6aef95a8 100644 --- a/gcc/langhooks.h +++ b/gcc/langhooks.h @@ -313,6 +313,22 @@ struct lang_hooks_for_decls /* Do language specific checking on an implicitly determined clause. */ void (*omp_finish_clause) (tree clause, gimple_seq *pre_p, bool); + /* Finish language-specific processing on mapping nodes after expanding + user-defined mappers. */ + tree (*omp_finish_mapper_clauses) (tree clauses); + + /* Find a mapper in the current parsing context, given a NAME (or + NULL_TREE) and TYPE. */ + tree (*omp_mapper_lookup) (tree name, tree type); + + /* Return the statement for the mapper directive definition, from the + representation used to contain it (e.g. an inline function + declaration). */ + tree (*omp_extract_mapper_directive) (tree fndecl); + + /* Return a simplified form for OMP_ARRAY_SECTION argument. */ + tree (*omp_map_array_section) (location_t, tree t); + /* Return true if DECL is an allocatable variable (for the purpose of implicit mapping). */ bool (*omp_allocatable_p) (tree decl); diff --git a/gcc/omp-general.h b/gcc/omp-general.h index f9b7ef3426c4..83401f6d9d00 100644 --- a/gcc/omp-general.h +++ b/gcc/omp-general.h @@ -152,6 +152,92 @@ get_openacc_privatization_dump_flags () extern tree omp_build_component_ref (tree obj, tree field); +template +struct omp_name_type +{ + tree name; + T type; +}; + +template <> +struct default_hash_traits > + : typed_noop_remove > +{ + GTY((skip)) typedef omp_name_type value_type; + GTY((skip)) typedef omp_name_type compare_type; + + static hashval_t + hash (omp_name_type p) + { + return p.name ? iterative_hash_expr (p.name, TYPE_UID (p.type)) + : TYPE_UID (p.type); + } + + static const bool empty_zero_p = true; + + static bool + is_empty (omp_name_type p) + { + return p.type == NULL; + } + + static bool + is_deleted (omp_name_type) + { + return false; + } + + static bool + equal (const omp_name_type &a, const omp_name_type &b) + { + if (a.name == NULL_TREE && b.name == NULL_TREE) + return a.type == b.type; + else if (a.name == NULL_TREE || b.name == NULL_TREE) + return false; + else + return a.name == b.name && a.type == b.type; + } + + static void + mark_empty (omp_name_type &e) + { + e.type = NULL; + } +}; + +template +struct omp_mapper_list +{ + hash_set> *seen_types; + vec *mappers; + + omp_mapper_list (hash_set> *s, vec *m) + : seen_types (s), mappers (m) { } + + void add_mapper (tree name, T type, tree mapperfn) + { + /* We can't hash a NULL_TREE... */ + if (!name) + name = void_node; + + omp_name_type n_t = { name, type }; + + if (seen_types->contains (n_t)) + return; + + seen_types->add (n_t); + mappers->safe_push (mapperfn); + } + + bool contains (tree name, T type) + { + if (!name) + name = void_node; + + return seen_types->contains ({ name, type }); + } +}; + namespace omp_addr_tokenizer { /* These are the ways of accessing a variable that have special-case handling diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c new file mode 100644 index 000000000000..c4d017036c5e --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c @@ -0,0 +1,22 @@ +/* { dg-do compile { target c++ } } */ + +struct XYZ { + int a; + int *b; + int c; +}; + +#pragma omp declare mapper(struct XYZ t) +/* { dg-error "missing 'map' clause" "" { target c } .-1 } */ +/* { dg-error "missing 'map' clause before end of line" "" { target c++ } .-2 } */ + +struct ABC { + int *a; + int b; + int c; +}; + +#pragma omp declare mapper(struct ABC d) firstprivate(d.b) +/* { dg-error "unexpected clause" "" { target c } .-1 } */ +/* { dg-error "expected end of line before '\\(' token" "" { target c } .-2 } */ +/* { dg-error "unexpected clause before '\\(' token" "" { target c++ } .-3 } */ diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c new file mode 100644 index 000000000000..983d979d68c5 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c @@ -0,0 +1,30 @@ +// { dg-do compile { target c++ } } +// { dg-additional-options "-fdump-tree-gimple" } + +#include + +// Test named mapper invocation. + +struct S { + int *ptr; + int size; +}; + +int main (int argc, char *argv[]) +{ + int N = 1024; +#pragma omp declare mapper (mapN:struct S s) map(to:s.ptr, s.size) \ + map(s.ptr[:N]) + + struct S s; + s.ptr = (int *) malloc (sizeof (int) * N); + +#pragma omp target map(mapper(mapN), tofrom: s) +// { dg-final { scan-tree-dump {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} "gimple" } } + { + for (int i = 0; i < N; i++) + s.ptr[i]++; + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c new file mode 100644 index 000000000000..6d933e4bf6f4 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c @@ -0,0 +1,78 @@ +/* { dg-do compile { target c++ } } */ +/* { dg-additional-options "-fdump-tree-original" } */ + +/* Check mapper binding clauses. */ + +struct Y { + int z; +}; + +struct Z { + int z; +}; + +#pragma omp declare mapper (struct Y y) map(tofrom: y) +#pragma omp declare mapper (struct Z z) map(tofrom: z) + +int foo (void) +{ + struct Y yy; + struct Z zz; + int dummy; + +#pragma omp target data map(dummy) + { + #pragma omp target + { + yy.z++; + zz.z++; + } + yy.z++; + } + return yy.z; +} + +struct P +{ + struct Z *zp; +}; + +int bar (void) +{ + struct Y yy; + struct Z zz; + struct P pp; + struct Z t; + int dummy; + + pp.zp = &t; + +#pragma omp declare mapper (struct Y y) map(tofrom: y.z) +#pragma omp declare mapper (struct Z z) map(tofrom: z.z) + +#pragma omp target data map(dummy) + { + #pragma omp target + { + yy.z++; + zz.z++; + } + yy.z++; + } + + #pragma omp declare mapper(struct P x) map(to:x.zp) map(tofrom:*x.zp) + + #pragma omp target + { + zz = *pp.zp; + } + + return zz.z; +} + +/* { dg-final { scan-tree-dump-times {mapper_binding\(struct Y,omp declare mapper ~1Y\) mapper_binding\(struct Z,omp declare mapper ~1Z\)} 2 "original" { target c++ } } } */ +/* { dg-final { scan-tree-dump {mapper_binding\(struct Z,omp declare mapper ~1Z\) mapper_binding\(struct P,omp declare mapper ~1P\)} "original" { target c++ } } } */ + +/* { dg-final { scan-tree-dump {mapper_binding\(struct Z,#pragma omp declare mapper \(struct Z z\) map\(tofrom:z\)\) mapper_binding\(struct Y,#pragma omp declare mapper \(struct Y y\) map\(tofrom:y\)\)} "original" { target c } } } */ +/* { dg-final { scan-tree-dump {mapper_binding\(struct Z,#pragma omp declare mapper \(struct Z z\) map\(tofrom:z\.z\)\) mapper_binding\(struct Y,#pragma omp declare mapper \(struct Y y\) map\(tofrom:y\.z\)\)} "original" { target c } } } */ +/* { dg-final { scan-tree-dump {mapper_binding\(struct P,#pragma omp declare mapper \(struct P x\) map\(tofrom:\(x\.zp\)\[0:1\]\) map\(to:x.zp\)\) mapper_binding\(struct Z,#pragma omp declare mapper \(struct Z z\) map\(tofrom:z\.z\)\)} "original" { target c } } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c new file mode 100644 index 000000000000..f675a8c68902 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c @@ -0,0 +1,26 @@ +/* { dg-do compile { target c++ } } */ + +typedef struct S_ { + int *myarr; + int size; +} S; + +#pragma omp declare mapper (named: struct S_ v) map(to:v.size, v.myarr) \ + map(tofrom: v.myarr[0:v.size]) +/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-2 } */ +/* { dg-note "'#pragma omp declare mapper \\(named: S_\\)' previously defined here" "" { target c++ } .-3 } */ + +#pragma omp declare mapper (named: S v) map(to:v.size, v.myarr) \ + map(tofrom: v.myarr[0:v.size]) +/* { dg-error "redeclaration of 'named' '#pragma omp declare mapper' for type 'S' \\\{aka 'struct S_'\\\}" "" { target c } .-2 } */ +/* { dg-error "redefinition of '#pragma omp declare mapper \\(named: S\\)'" "" { target c++ } .-3 } */ + +#pragma omp declare mapper (struct S_ v) map(to:v.size, v.myarr) \ + map(tofrom: v.myarr[0:v.size]) +/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-2 } */ +/* { dg-note "'#pragma omp declare mapper \\(S_\\)' previously defined here" "" { target c++ } .-3 } */ + +#pragma omp declare mapper (S v) map(to:v.size, v.myarr) \ + map(tofrom: v.myarr[0:v.size]) +/* { dg-error "redeclaration of '' '#pragma omp declare mapper' for type 'S' \\\{aka 'struct S_'\\\}" "" { target c } .-2 } */ +/* { dg-error "redefinition of '#pragma omp declare mapper \\(S\\)'" "" { target c++ } .-3 } */ diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c new file mode 100644 index 000000000000..a2f6c08cdfdd --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c @@ -0,0 +1,23 @@ +/* { dg-do compile { target c++ } } */ + +int x = 5; + +struct Q { + int *arr1; + int *arr2; + int *arr3; +}; + +#pragma omp declare mapper (struct Q myq) map(myq.arr2[0:x]) + +struct R { + int *arr1; + int *arr2; + int *arr3; +}; + +#pragma omp declare mapper (struct R myr) map(myr.arr3[0:y]) +/* { dg-error "'y' undeclared" "" { target c } .-1 } */ +/* { dg-error "'y' was not declared in this scope" "" { target c++ } .-2 } */ + +int y = 7; diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c new file mode 100644 index 000000000000..1b1be9dbb666 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c @@ -0,0 +1,29 @@ +/* { dg-do compile { target c++ } } */ + +struct Q { + int *arr1; + int *arr2; + int *arr3; +}; + +int foo (void) +{ + int x = 5; + #pragma omp declare mapper (struct Q myq) map(myq.arr2[0:x]) + return x; +} + +struct R { + int *arr1; + int *arr2; + int *arr3; +}; + +int bar (void) +{ + #pragma omp declare mapper (struct R myr) map(myr.arr3[0:y]) + /* { dg-error "'y' undeclared" "" { target c } .-1 } */ + /* { dg-error "'y' was not declared in this scope" "" { target c++ } .-2 } */ + int y = 7; + return y; +} diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c new file mode 100644 index 000000000000..86ddb942072c --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c @@ -0,0 +1,43 @@ +/* { dg-do compile { target c++ } } */ + +struct Q { + int *arr1; + int *arr2; + int *arr3; + int len; +}; + +struct R { + struct Q qarr[5]; +}; + +struct R2 { + struct Q *qptr; +}; + +#pragma omp declare mapper (struct Q myq) map(myq.arr1[0:myq.len]) \ + map(myq.arr2[0:myq.len]) \ + map(myq.arr3[0:myq.len]) + +#pragma omp declare mapper (struct R myr) map(myr.qarr[2:3]) + +#pragma omp declare mapper (struct R2 myr2) map(myr2.qptr[2:3]) + +int main (int argc, char *argv[]) +{ + struct R r; + struct R2 r2; + int N = 256; + +#pragma omp target +/* { dg-message "sorry, unimplemented: user-defined mapper with non-unit length array section" "" { target *-*-* } .-1 } */ + { + for (int i = 2; i < 5; i++) + for (int j = 0; j < N; j++) + { + r.qarr[i].arr1[j]++; + r2.qptr[i].arr2[j]++; + } + } +} + diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c new file mode 100644 index 000000000000..54e58426910e --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c @@ -0,0 +1,34 @@ +/* { dg-do compile { target c++ } } */ + +int x = 5; + +struct Q { + int *arr1; + int *arr2; + int *arr3; +}; + +int y = 5; + +#pragma omp declare mapper (struct Q myq) map(myq.arr2[0:x]) +/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-1 } */ +/* { dg-note "'#pragma omp declare mapper \\(Q\\)' previously defined here" "" { target c++ } .-2 } */ + +#pragma omp declare mapper (struct Q myq) map(myq.arr2[0:y]) +/* { dg-error "redeclaration of '' '#pragma omp declare mapper' for type 'struct Q'" "" { target c } .-1 } */ +/* { dg-error "redefinition of '#pragma omp declare mapper \\(Q\\)'" "" { target c++ } .-2 } */ + +struct R { + int *arr1; +}; + +void foo (void) +{ +#pragma omp declare mapper (struct R myr) map(myr.arr1[0:x]) +/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-1 } */ +/* { dg-note "'#pragma omp declare mapper \\(R\\)' previously declared here" "" { target c++ } .-2 } */ + +#pragma omp declare mapper (struct R myr) map(myr.arr1[0:y]) +/* { dg-error "redeclaration of '' '#pragma omp declare mapper' for type 'struct R'" "" { target c } .-1 } */ +/* { dg-error "redeclaration of '#pragma omp declare mapper \\(R\\)'" "" { target c++ } .-2 } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c index c749db845b0a..77054b6b56aa 100644 --- a/gcc/testsuite/c-c++-common/gomp/map-6.c +++ b/gcc/testsuite/c-c++-common/gomp/map-6.c @@ -13,10 +13,12 @@ foo (void) #pragma omp target map (to:a) ; - #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" } */ + #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" "" { target c } } */ +/* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'mapper'" "" { target c++ } .-1 } */ ; - #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" } */ + #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" "" { target c } } */ +/* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'mapper'" "" { target c++ } .-1 } */ ; #pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */ diff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C new file mode 100644 index 000000000000..3177d20adbc2 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C @@ -0,0 +1,58 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } + +// "omp declare mapper" support -- check expansion in gimple. + +struct S { + int *ptr; + int size; +}; + +#define N 64 + +#pragma omp declare mapper (S w) map(w.size, w.ptr, w.ptr[:w.size]) +#pragma omp declare mapper (foo:S w) map(to:w.size, w.ptr) map(w.ptr[:w.size]) + +int main (int argc, char *argv[]) +{ + S s; + s.ptr = new int[N]; + s.size = N; + +#pragma omp declare mapper (bar:S w) map(w.size, w.ptr, w.ptr[:w.size]) + +#pragma omp target + { + for (int i = 0; i < N; i++) + s.ptr[i]++; + } + +#pragma omp target map(tofrom: s) + { + for (int i = 0; i < N; i++) + s.ptr[i]++; + } + +#pragma omp target map(mapper(default), tofrom: s) + { + for (int i = 0; i < N; i++) + s.ptr[i]++; + } + +#pragma omp target map(mapper(foo), alloc: s) + { + for (int i = 0; i < N; i++) + s.ptr[i]++; + } + +#pragma omp target map(mapper(bar), tofrom: s) + { + for (int i = 0; i < N; i++) + s.ptr[i]++; + } + + return 0; +} + +// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(tofrom:s\.size \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 4 "gimple" } } +// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 1 "gimple" } } diff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-2.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-2.C new file mode 100644 index 000000000000..7df72c76e2af --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/declare-mapper-2.C @@ -0,0 +1,30 @@ +// { dg-do compile } + +// Error-checking tests for "omp declare mapper". + +struct S { + int *ptr; + int size; +}; + +struct Z { + int z; +}; + +int main (int argc, char *argv[]) +{ +#pragma omp declare mapper (S v) map(v.size, v.ptr[:v.size]) // { dg-note "'#pragma omp declare mapper \\(S\\)' previously declared here" } + + /* This one's a duplicate. */ +#pragma omp declare mapper (default: S v) map (to: v.size) map (v) // { dg-error "redeclaration of '#pragma omp declare mapper \\(S\\)'" } + + /* ...and this one doesn't use a "base language identifier" for the mapper + name. */ +#pragma omp declare mapper (case: S v) map (to: v.size) // { dg-error "expected identifier or 'default' before 'case'" } + // { dg-error "expected ':' before 'case'" "" { target *-*-* } .-1 } + + /* A non-struct/class/union type isn't supposed to work. */ +#pragma omp declare mapper (name:Z [5]foo) map (foo[0].z) // { dg-error "'Z \\\[5\\\]' is not a struct, union or class type in '#pragma omp declare mapper'" } + + return 0; +} diff --git a/gcc/tree-core.h b/gcc/tree-core.h index e146b133dbd6..c23ad0e7a6f6 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -348,6 +348,10 @@ enum omp_clause_code { /* OpenMP clause: doacross ({source,sink}:vec). */ OMP_CLAUSE_DOACROSS, + /* OpenMP mapper binding: record implicit mappers in scope for aggregate + types used within an offload region. */ + OMP_CLAUSE__MAPPER_BINDING_, + /* Internal structure to hold OpenACC cache directive's variable-list. #pragma acc cache (variable-list). */ OMP_CLAUSE__CACHE_, diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 1399667cbdb6..2818cecaa8f7 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -994,6 +994,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: pp_string (pp, "attach_zero_length_array_section"); break; + case GOMP_MAP_UNSET: + pp_string (pp, "unset"); + break; + case GOMP_MAP_PUSH_MAPPER_NAME: + pp_string (pp, "push_mapper"); + break; + case GOMP_MAP_POP_MAPPER_NAME: + pp_string (pp, "pop_mapper"); + break; default: gcc_unreachable (); } @@ -1057,6 +1066,23 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) spc, flags, false); goto print_clause_size; + case OMP_CLAUSE__MAPPER_BINDING_: + pp_string (pp, "mapper_binding("); + if (OMP_CLAUSE__MAPPER_BINDING__ID (clause)) + { + dump_generic_node (pp, OMP_CLAUSE__MAPPER_BINDING__ID (clause), spc, + flags, false); + pp_comma (pp); + } + dump_generic_node (pp, + TREE_TYPE (OMP_CLAUSE__MAPPER_BINDING__DECL (clause)), + spc, flags, false); + pp_comma (pp); + dump_generic_node (pp, OMP_CLAUSE__MAPPER_BINDING__MAPPER (clause), spc, + flags, false); + pp_right_paren (pp); + break; + case OMP_CLAUSE_NUM_TEAMS: pp_string (pp, "num_teams("); if (OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (clause)) @@ -3825,6 +3851,21 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags, is_expr = false; break; + case OMP_DECLARE_MAPPER: + pp_string (pp, "#pragma omp declare mapper ("); + if (OMP_DECLARE_MAPPER_ID (node)) + { + dump_generic_node (pp, OMP_DECLARE_MAPPER_ID (node), spc, flags, + false); + pp_colon (pp); + } + dump_generic_node (pp, TREE_TYPE (node), spc, flags, false); + pp_space (pp); + dump_generic_node (pp, OMP_DECLARE_MAPPER_DECL (node), spc, flags, false); + pp_right_paren (pp); + dump_omp_clauses (pp, OMP_DECLARE_MAPPER_CLAUSES (node), spc, flags); + break; + case TRANSACTION_EXPR: if (TRANSACTION_EXPR_OUTER (node)) pp_string (pp, "__transaction_atomic [[outer]]"); diff --git a/gcc/tree.cc b/gcc/tree.cc index 581d4489438e..14823a0ffc85 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -271,6 +271,7 @@ unsigned const char omp_clause_num_ops[] = 2, /* OMP_CLAUSE_MAP */ 1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */ 1, /* OMP_CLAUSE_DOACROSS */ + 3, /* OMP_CLAUSE__MAPPER_BINDING_ */ 2, /* OMP_CLAUSE__CACHE_ */ 2, /* OMP_CLAUSE_GANG */ 1, /* OMP_CLAUSE_ASYNC */ @@ -362,6 +363,7 @@ const char * const omp_clause_code_name[] = "map", "has_device_addr", "doacross", + "_mapper_binding_", "_cache_", "gang", "async", diff --git a/gcc/tree.def b/gcc/tree.def index f82bf257aedc..c876f06fd354 100644 --- a/gcc/tree.def +++ b/gcc/tree.def @@ -1250,6 +1250,13 @@ DEFTREECODE (OMP_SECTION, "omp_section", tcc_statement, 1) Operand 0: OMP_MASTER_BODY: Master section body. */ DEFTREECODE (OMP_MASTER, "omp_master", tcc_statement, 1) +/* OpenMP - #pragma omp declare mapper ([id:] type var) [clause1 ... clauseN] + Operand 0: Identifier. + Operand 1: Variable decl. + Operand 2: List of clauses. + The type of the construct is used for the type to be mapped. */ +DEFTREECODE (OMP_DECLARE_MAPPER, "omp_declare_mapper", tcc_statement, 3) + /* OpenACC - #pragma acc cache (variable1 ... variableN) Operand 0: OACC_CACHE_CLAUSES: List of variables (transformed into OMP_CLAUSE__CACHE_ clauses). */ diff --git a/gcc/tree.h b/gcc/tree.h index ac3dc6181e2a..2fdb9462f0ce 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1536,6 +1536,13 @@ class auto_suppress_location_wrappers #define OMP_TARGET_EXIT_DATA_CLAUSES(NODE)\ TREE_OPERAND (OMP_TARGET_EXIT_DATA_CHECK (NODE), 0) +#define OMP_DECLARE_MAPPER_ID(NODE) \ + TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 0) +#define OMP_DECLARE_MAPPER_DECL(NODE) \ + TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 1) +#define OMP_DECLARE_MAPPER_CLAUSES(NODE) \ + TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 2) + #define OMP_SCAN_BODY(NODE) TREE_OPERAND (OMP_SCAN_CHECK (NODE), 0) #define OMP_SCAN_CLAUSES(NODE) TREE_OPERAND (OMP_SCAN_CHECK (NODE), 1) @@ -1976,6 +1983,18 @@ class auto_suppress_location_wrappers #define OMP_CLAUSE__SCANTEMP__CONTROL(NODE) \ TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__SCANTEMP_)) +#define OMP_CLAUSE__MAPPER_BINDING__ID(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \ + OMP_CLAUSE__MAPPER_BINDING_), 0) + +#define OMP_CLAUSE__MAPPER_BINDING__DECL(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \ + OMP_CLAUSE__MAPPER_BINDING_), 1) + +#define OMP_CLAUSE__MAPPER_BINDING__MAPPER(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \ + OMP_CLAUSE__MAPPER_BINDING_), 2) + /* SSA_NAME accessors. */ /* Whether SSA_NAME NODE is a virtual operand. This simply caches the diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 696f5d4060e5..4e913cb2ebba 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -188,7 +188,13 @@ enum gomp_map_kind /* An attach or detach operation. Rewritten to the appropriate type during gimplification, depending on directive (i.e. "enter data" or parallel/kernels region vs. "exit data"). */ - GOMP_MAP_ATTACH_DETACH = (GOMP_MAP_LAST | 3) + GOMP_MAP_ATTACH_DETACH = (GOMP_MAP_LAST | 3), + /* Unset, used for "declare mapper" maps with no explicit data movement + specified. These use the movement specified at the invocation site. */ + GOMP_MAP_UNSET = (GOMP_MAP_LAST | 4), + /* Used to record the name of a named mapper. */ + GOMP_MAP_PUSH_MAPPER_NAME = (GOMP_MAP_LAST | 5), + GOMP_MAP_POP_MAPPER_NAME = (GOMP_MAP_LAST | 6) }; #define GOMP_MAP_COPY_TO_P(X) \ diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-1.C b/libgomp/testsuite/libgomp.c++/declare-mapper-1.C new file mode 100644 index 000000000000..aba4f4265392 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-1.C @@ -0,0 +1,87 @@ +// { dg-do run } + +#include +#include + +#define N 64 + +struct points +{ + double *x; + double *y; + double *z; + size_t len; +}; + +#pragma omp declare mapper(points p) map(to:p.x, p.y, p.z) \ + map(p.x[0:p.len]) \ + map(p.y[0:p.len]) \ + map(p.z[0:p.len]) + +struct shape +{ + points tmp; + points *pts; + int metadata[128]; +}; + +#pragma omp declare mapper(shape s) map(tofrom:s.pts, *s.pts) map(alloc:s.tmp) + +void +alloc_points (points *pts, size_t sz) +{ + pts->x = new double[sz]; + pts->y = new double[sz]; + pts->z = new double[sz]; + pts->len = sz; + for (int i = 0; i < sz; i++) + pts->x[i] = pts->y[i] = pts->z[i] = 0; +} + +int main (int argc, char *argv[]) +{ + shape myshape; + points mypts; + + myshape.pts = &mypts; + + alloc_points (&myshape.tmp, N); + myshape.pts = new points; + alloc_points (myshape.pts, N); + + #pragma omp target map(myshape) + { + for (int i = 0; i < N; i++) + { + myshape.pts->x[i]++; + myshape.pts->y[i]++; + myshape.pts->z[i]++; + } + } + + for (int i = 0; i < N; i++) + { + assert (myshape.pts->x[i] == 1); + assert (myshape.pts->y[i] == 1); + assert (myshape.pts->z[i] == 1); + } + + #pragma omp target + { + for (int i = 0; i < N; i++) + { + myshape.pts->x[i]++; + myshape.pts->y[i]++; + myshape.pts->z[i]++; + } + } + + for (int i = 0; i < N; i++) + { + assert (myshape.pts->x[i] == 2); + assert (myshape.pts->y[i] == 2); + assert (myshape.pts->z[i] == 2); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-2.C b/libgomp/testsuite/libgomp.c++/declare-mapper-2.C new file mode 100644 index 000000000000..d848fdb73692 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-2.C @@ -0,0 +1,55 @@ +// { dg-do run } + +#include + +#define N 256 + +struct doublebuf +{ + int buf_a[N][N]; + int buf_b[N][N]; +}; + +#pragma omp declare mapper(lo:doublebuf b) map(b.buf_a[0:N/2][0:N]) \ + map(b.buf_b[0:N/2][0:N]) + +#pragma omp declare mapper(hi:doublebuf b) map(b.buf_a[N/2:N/2][0:N]) \ + map(b.buf_b[N/2:N/2][0:N]) + +int main (int argc, char *argv[]) +{ + doublebuf db; + + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + db.buf_a[i][j] = db.buf_b[i][j] = 0; + + #pragma omp target map(mapper(lo), tofrom:db) + { + for (int i = 0; i < N / 2; i++) + for (int j = 0; j < N; j++) + { + db.buf_a[i][j]++; + db.buf_b[i][j]++; + } + } + + #pragma omp target map(mapper(hi), tofrom:db) + { + for (int i = N / 2; i < N; i++) + for (int j = 0; j < N; j++) + { + db.buf_a[i][j]++; + db.buf_b[i][j]++; + } + } + + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + { + assert (db.buf_a[i][j] == 1); + assert (db.buf_b[i][j] == 1); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-3.C b/libgomp/testsuite/libgomp.c++/declare-mapper-3.C new file mode 100644 index 000000000000..ea9b7ded75b6 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-3.C @@ -0,0 +1,63 @@ +// { dg-do run } + +#include +#include + +struct S { + int *myarr; +}; + +#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:20]) + +namespace A { +#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:100]) +} + +namespace B { +#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[100:100]) +} + +namespace A +{ + void incr_a (S my_s) + { +#pragma omp target + { + for (int i = 0; i < 100; i++) + my_s.myarr[i]++; + } + } +} + +namespace B +{ + void incr_b (S my_s) + { +#pragma omp target + { + for (int i = 100; i < 200; i++) + my_s.myarr[i]++; + } + } +} + +int main (int argc, char *argv[]) +{ + S my_s; + + my_s.myarr = (int *) calloc (200, sizeof (int)); + +#pragma omp target + { + for (int i = 0; i < 20; i++) + my_s.myarr[i]++; + } + + A::incr_a (my_s); + B::incr_b (my_s); + + for (int i = 0; i < 200; i++) + assert (my_s.myarr[i] == (i < 20) ? 2 : 1); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-4.C b/libgomp/testsuite/libgomp.c++/declare-mapper-4.C new file mode 100644 index 000000000000..f194e63b5b7f --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-4.C @@ -0,0 +1,63 @@ +// { dg-do run } + +#include +#include + +struct S { + int *myarr; +}; + +#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:20]) + +namespace A { +#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:100]) +} + +namespace B { +#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[100:100]) +} + +namespace A +{ + void incr_a (S &my_s) + { +#pragma omp target + { + for (int i = 0; i < 100; i++) + my_s.myarr[i]++; + } + } +} + +namespace B +{ + void incr_b (S &my_s) + { +#pragma omp target + { + for (int i = 100; i < 200; i++) + my_s.myarr[i]++; + } + } +} + +int main (int argc, char *argv[]) +{ + S my_s; + + my_s.myarr = (int *) calloc (200, sizeof (int)); + +#pragma omp target + { + for (int i = 0; i < 20; i++) + my_s.myarr[i]++; + } + + A::incr_a (my_s); + B::incr_b (my_s); + + for (int i = 0; i < 200; i++) + assert (my_s.myarr[i] == (i < 20) ? 2 : 1); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-5.C b/libgomp/testsuite/libgomp.c++/declare-mapper-5.C new file mode 100644 index 000000000000..0030de8791a0 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-5.C @@ -0,0 +1,52 @@ +// { dg-do run } + +#include + +struct S +{ + int *myarr; + int len; +}; + +class C +{ + S smemb; +#pragma omp declare mapper (custom:S s) map(to:s.myarr) \ + map(tofrom:s.myarr[0:s.len]) + +public: + C(int l) + { + smemb.myarr = new int[l]; + smemb.len = l; + for (int i = 0; i < l; i++) + smemb.myarr[i] = 0; + } + void bump(); + void check(); +}; + +void +C::bump () +{ +#pragma omp target map(mapper(custom), tofrom: smemb) + { + for (int i = 0; i < smemb.len; i++) + smemb.myarr[i]++; + } +} + +void +C::check () +{ + for (int i = 0; i < smemb.len; i++) + assert (smemb.myarr[i] == 1); +} + +int main (int argc, char *argv[]) +{ + C test (100); + test.bump (); + test.check (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-6.C b/libgomp/testsuite/libgomp.c++/declare-mapper-6.C new file mode 100644 index 000000000000..14ed10df7025 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-6.C @@ -0,0 +1,37 @@ +// { dg-do run } + +#include + +template +void adjust (T param) +{ +#pragma omp declare mapper (T x) map(to:x.len, x.base) \ + map(tofrom:x.base[0:x.len]) + +#pragma omp target + for (int i = 0; i < param.len; i++) + param.base[i]++; +} + +struct S { + int len; + int *base; +}; + +int main (int argc, char *argv[]) +{ + S a; + + a.len = 100; + a.base = new int[a.len]; + + for (int i = 0; i < a.len; i++) + a.base[i] = 0; + + adjust (a); + + for (int i = 0; i < a.len; i++) + assert (a.base[i] == 1); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-7.C b/libgomp/testsuite/libgomp.c++/declare-mapper-7.C new file mode 100644 index 000000000000..ab6320997148 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-7.C @@ -0,0 +1,48 @@ +// { dg-do run } + +#include + +struct S +{ + int *myarr; +}; + +struct T +{ + S *s; +}; + +#pragma omp declare mapper (s100: S x) map(to: x.myarr) \ + map(tofrom: x.myarr[0:100]) + +void +bump (T t) +{ + /* Here we have an implicit/default mapper invoking a named mapper. We + need to make sure that can be located properly at gimplification + time. */ +#pragma omp declare mapper (T t) map(to:t.s) map(mapper(s100), tofrom: t.s[0]) + +#pragma omp target + for (int i = 0; i < 100; i++) + t.s->myarr[i]++; +} + +int main (int argc, char *argv[]) +{ + S my_s; + T my_t; + + my_s.myarr = new int[100]; + my_t.s = &my_s; + + for (int i = 0; i < 100; i++) + my_s.myarr[i] = 0; + + bump (my_t); + + for (int i = 0; i < 100; i++) + assert (my_s.myarr[i] == 1); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-8.C b/libgomp/testsuite/libgomp.c++/declare-mapper-8.C new file mode 100644 index 000000000000..3818e5264d35 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-8.C @@ -0,0 +1,61 @@ +// { dg-do run } + +#include + +struct S +{ + int *myarr; + int len; +}; + +template +class C +{ + T memb; +#pragma omp declare mapper (T t) map(to:t.len, t.myarr) \ + map(tofrom:t.myarr[0:t.len]) + +public: + C(int sz); + ~C(); + void bump(); + void check(); +}; + +template +C::C(int sz) +{ + memb.myarr = new int[sz]; + for (int i = 0; i < sz; i++) + memb.myarr[i] = 0; + memb.len = sz; +} + +template +C::~C() +{ + delete[] memb.myarr; +} + +template +void C::bump() +{ +#pragma omp target map(memb) + for (int i = 0; i < memb.len; i++) + memb.myarr[i]++; +} + +template +void C::check() +{ + for (int i = 0; i < memb.len; i++) + assert (memb.myarr[i] == 1); +} + +int main(int argc, char *argv[]) +{ + C c_int(100); + c_int.bump(); + c_int.check(); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c new file mode 100644 index 000000000000..b0fa40929fbc --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c @@ -0,0 +1,60 @@ +/* { dg-do run { target c++ } } */ + +#include +#include +#include + +#define N 64 + +typedef struct { + int *arr; + int size; +} B; + +#pragma omp declare mapper (mapB : B myb) map(to: myb.size, myb.arr) \ + map(tofrom: myb.arr[0:myb.size]) + +struct A { + int *arr1; + B *arr2; + int arr3[N]; +}; + +int +main (int argc, char *argv[]) +{ + struct A var; + + memset (&var, 0, sizeof var); + var.arr1 = (int *) calloc (N, sizeof (int)); + var.arr2 = (B *) malloc (sizeof (B)); + var.arr2->arr = (int *) calloc (N, sizeof (float)); + var.arr2->size = N; + + { + #pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr2) \ + map(tofrom: x.arr1[0:N]) \ + map(mapper(mapB), tofrom: x.arr2[0:1]) + #pragma omp target + { + for (int i = 0; i < N; i++) + { + var.arr1[i]++; + var.arr2->arr[i]++; + } + } + } + + for (int i = 0; i < N; i++) + { + assert (var.arr1[i] == 1); + assert (var.arr2->arr[i] == 1); + assert (var.arr3[i] == 0); + } + + free (var.arr1); + free (var.arr2->arr); + free (var.arr2); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c new file mode 100644 index 000000000000..b509ddc412c5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c @@ -0,0 +1,59 @@ +/* { dg-do run { target c++ } } */ + +#include +#include +#include + +#define N 64 + +typedef struct B_tag { + int *arr; + int size; +} B; + +#pragma omp declare mapper (B myb) map(to: myb.size, myb.arr) \ + map(tofrom: myb.arr[0:myb.size]) + +struct A { + int *arr1; + B *arr2; + int arr3[N]; +}; + +int +main (int argc, char *argv[]) +{ + struct A var; + + memset (&var, 0, sizeof var); + var.arr1 = (int *) calloc (N, sizeof (int)); + var.arr2 = (B *) malloc (sizeof (B)); + var.arr2->arr = (int *) calloc (N, sizeof (int)); + var.arr2->size = N; + + { + #pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr2) \ + map(tofrom: x.arr1[0:N]) map(tofrom: x.arr2[0:1]) + #pragma omp target + { + for (int i = 0; i < N; i++) + { + var.arr1[i]++; + var.arr2->arr[i]++; + } + } + } + + for (int i = 0; i < N; i++) + { + assert (var.arr1[i] == 1); + assert (var.arr2->arr[i] == 1); + assert (var.arr3[i] == 0); + } + + free (var.arr1); + free (var.arr2->arr); + free (var.arr2); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c new file mode 100644 index 000000000000..cf8919c22edf --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c @@ -0,0 +1,87 @@ +/* { dg-do run { target c++ } } */ + +#include +#include +#include + +#define N 64 + +typedef struct { + int *arr; + int size; +} B; + +#pragma omp declare mapper (samename : B myb) map(to: myb.size, myb.arr) \ + map(tofrom: myb.arr[0:myb.size]) + +typedef struct { + int *arr; + int size; +} C; + + +struct A { + int *arr1; + B *arr2; + C *arr3; +}; + +int +main (int argc, char *argv[]) +{ + struct A var; + + memset (&var, 0, sizeof var); + var.arr1 = (int *) calloc (N, sizeof (int)); + var.arr2 = (B *) malloc (sizeof (B)); + var.arr2->arr = (int *) calloc (N, sizeof (int)); + var.arr2->size = N; + var.arr3 = (C *) malloc (sizeof (C)); + var.arr3->arr = (int *) calloc (N, sizeof (int)); + var.arr3->size = N; + + { + #pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr2) \ + map(tofrom: x.arr1[0:N]) \ + map(mapper(samename), tofrom: x.arr2[0:1]) + #pragma omp target + { + for (int i = 0; i < N; i++) + { + var.arr1[i]++; + var.arr2->arr[i]++; + } + } + } + + { + #pragma omp declare mapper (samename : C myc) map(to: myc.size, myc.arr) \ + map(tofrom: myc.arr[0:myc.size]) + #pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr3) \ + map(tofrom: x.arr1[0:N]) \ + map(mapper(samename), tofrom: *x.arr3) + #pragma omp target + { + for (int i = 0; i < N; i++) + { + var.arr1[i]++; + var.arr3->arr[i]++; + } + } + } + + for (int i = 0; i < N; i++) + { + assert (var.arr1[i] == 2); + assert (var.arr2->arr[i] == 1); + assert (var.arr3->arr[i] == 1); + } + + free (var.arr1); + free (var.arr2->arr); + free (var.arr2); + free (var.arr3->arr); + free (var.arr3); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c new file mode 100644 index 000000000000..99b7eedad90f --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c @@ -0,0 +1,55 @@ +/* { dg-do run { target c++ } } */ + +#include + +struct T { + int a; + int b; + int c; +}; + +void foo (void) +{ + struct T x; + x.a = x.b = x.c = 0; + +#pragma omp target + { + x.a++; + x.c++; + } + + assert (x.a == 1); + assert (x.b == 0); + assert (x.c == 1); +} + +// An identity mapper. This should do the same thing as the default! +#pragma omp declare mapper (struct T v) map(v) + +void bar (void) +{ + struct T x; + x.a = x.b = x.c = 0; + +#pragma omp target + { + x.b++; + } + +#pragma omp target map(x) + { + x.a++; + } + + assert (x.a == 1); + assert (x.b == 1); + assert (x.c == 0); +} + +int main (int argc, char *argv[]) +{ + foo (); + bar (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c new file mode 100644 index 000000000000..e7108da25fef --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c @@ -0,0 +1,57 @@ +/* { dg-do run { target c++ } } */ + +#include +#include + +struct Z { + int *arr; +}; + +void baz (struct Z *zarr, int len) +{ +#pragma omp declare mapper (struct Z myvar) map(to: myvar.arr) \ + map(tofrom: myvar.arr[0:len]) + zarr[0].arr = (int *) calloc (len, sizeof (int)); + zarr[5].arr = (int *) calloc (len, sizeof (int)); + +#pragma omp target map(zarr, *zarr) + { + for (int i = 0; i < len; i++) + zarr[0].arr[i]++; + } + +#pragma omp target map(zarr, zarr[5]) + { + for (int i = 0; i < len; i++) + zarr[5].arr[i]++; + } + +#pragma omp target map(zarr[5]) + { + for (int i = 0; i < len; i++) + zarr[5].arr[i]++; + } + +#pragma omp target map(zarr, zarr[5:1]) + { + for (int i = 0; i < len; i++) + zarr[5].arr[i]++; + } + + for (int i = 0; i < len; i++) + assert (zarr[0].arr[i] == 1); + + for (int i = 0; i < len; i++) + assert (zarr[5].arr[i] == 3); + + free (zarr[5].arr); + free (zarr[0].arr); +} + +int +main (int argc, char *argv[]) +{ + struct Z myzarr[10]; + baz (myzarr, 256); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c new file mode 100644 index 000000000000..9f85df53998a --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c @@ -0,0 +1,62 @@ +/* { dg-do run { target c++ } } */ + +#include +#include +#include + +#define N 64 + +struct A { + int *arr1; + float *arr2; + int arr3[N]; +}; + +int +main (int argc, char *argv[]) +{ + struct A var; + + memset (&var, 0, sizeof var); + var.arr1 = (int *) calloc (N, sizeof (int)); + var.arr2 = (float *) calloc (N, sizeof (float)); + + { + #pragma omp declare mapper (struct A x) map(to: x.arr1) \ + map(tofrom: x.arr1[0:N]) + #pragma omp target + { + for (int i = 0; i < N; i++) + var.arr1[i]++; + } + } + + { + #pragma omp declare mapper (struct A x) map(to: x.arr2) \ + map(tofrom: x.arr2[0:N]) + #pragma omp target + { + for (int i = 0; i < N; i++) + var.arr2[i]++; + } + } + + { + #pragma omp declare mapper (struct A x) map(tofrom: x.arr3[0:N]) + #pragma omp target + { + for (int i = 0; i < N; i++) + var.arr3[i]++; + } + } + + for (int i = 0; i < N; i++) + { + assert (var.arr1[i] == 1); + assert (var.arr2[i] == 1); + assert (var.arr3[i] == 1); + } + + free (var.arr1); + free (var.arr2); +} -- 2.29.2