From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 9F1AF3856DC6; Tue, 5 Sep 2023 19:29:05 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 9F1AF3856DC6 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="6.02,229,1688457600"; d="scan'208";a="16179083" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 05 Sep 2023 11:29:03 -0800 IronPort-SDR: 111guDmtmD4rE0IY+UG1dcYHcG3gieUkFqd9HNRLkug2XpJnvC1ut3WvkUNt0lJ4B5dMbGGPAR ef4/WCJBQilD5VWq8XPZplVbkRmzhtXAJN0g1f6bBv0sBHI4iThr6pA4WJvGPnl+8VodYAWyRA 4nvPha73Myt4Xqaf8lDBHHiB5q5g1gBNAQs66yURWGN98JFS6JHGj8d/8daPLCv6HFmluCcAPi +O9vdM9RmsVrw7TqfQ9X4rmwVascSnUGzn6RBU7qrECe6P3bNMDt7eDMDF+rQ8SFmG/NVfzay5 yU4= From: Julian Brown To: CC: , , Subject: [PATCH 3/8] OpenMP: C++ "declare mapper" support Date: Tue, 5 Sep 2023 12:28:23 -0700 Message-ID: X-Mailer: git-send-email 2.41.0 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-12.mgc.mentorg.com (139.181.222.12) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.0 required=5.0 tests=BAYES_00,GIT_PATCH_0,HEADER_FROM_DIFFERENT_DOMAINS,KAM_DMARC_STATUS,KAM_SHORT,SPF_HELO_PASS,SPF_PASS,TXREP,URIBL_BLACK 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 patch adds support for OpenMP 5.0 "declare mapper" functionality for C++. This is based on the version of the patch posted for the og13 branch here: https://gcc.gnu.org/pipermail/gcc-patches/2023-June/623353.html The following follow up patches/fixes have also been incorporated into this version: "OpenMP: Expand "declare mapper" mappers for target {enter,exit,} data directives": https://gcc.gnu.org/pipermail/gcc-patches/2023-July/623780.html "OpenMP: Introduce C_ORT_{,OMP_}DECLARE_MAPPER c_omp_region_type types": https://gcc.gnu.org/pipermail/gcc-patches/2023-August/627005.html 2023-09-05 Julian Brown gcc/c-family/ * c-common.h (c_omp_region_type): Add C_ORT_EXIT_DATA, C_ORT_DECLARE_MAPPER, C_ORT_OMP_EXIT_DATA, C_ORT_OMP_DECLARE_MAPPER values. (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_split_map_kind, omp_join_map_kind, omp_map_decayed_kind, 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, cxx_omp_mapper_lookup, cxx_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): Handle MAPPER_BINDING clauses. (gimplify_adjust_omp_clauses): Instantiate implicit declared mappers. (gimplify_omp_declare_mapper): New function. (gimplify_expr): Call above function. * langhooks-def.h (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. * c-c++-common/gomp/declare-mapper-15.c: Likewise. * c-c++-common/gomp/declare-mapper-16.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 | 7 + gcc/c-family/c-omp.cc | 477 ++++++++++++++++++ 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 | 305 ++++++++++- gcc/cp/pt.cc | 35 +- gcc/cp/semantics.cc | 191 ++++++- gcc/fortran/parse.cc | 3 + gcc/gimplify.cc | 263 ++++++++++ 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-15.c | 59 +++ .../c-c++-common/gomp/declare-mapper-16.c | 39 ++ .../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 | 10 +- 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 +++ 51 files changed, 2907 insertions(+), 34 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-15.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-16.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 4d1a6b1db6b9..b4e49c7d0cbc 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1275,8 +1275,12 @@ enum c_omp_region_type C_ORT_ACC = 1 << 1, C_ORT_DECLARE_SIMD = 1 << 2, C_ORT_TARGET = 1 << 3, + C_ORT_EXIT_DATA = 1 << 4, + C_ORT_DECLARE_MAPPER = 1 << 6, C_ORT_OMP_DECLARE_SIMD = C_ORT_OMP | C_ORT_DECLARE_SIMD, C_ORT_OMP_TARGET = C_ORT_OMP | C_ORT_TARGET, + C_ORT_OMP_EXIT_DATA = C_ORT_OMP | C_ORT_EXIT_DATA, + C_ORT_OMP_DECLARE_MAPPER = C_ORT_OMP | C_ORT_DECLARE_MAPPER, C_ORT_ACC_TARGET = C_ORT_ACC | C_ORT_TARGET }; @@ -1312,6 +1316,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, enum c_omp_region_type); 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 0d668925da11..b73f9682f460 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -3938,6 +3938,483 @@ 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; +} + +static enum gomp_map_kind +omp_split_map_kind (enum gomp_map_kind op, bool *force_p, bool *always_p, + bool *present_p) +{ + *force_p = *always_p = *present_p = false; + + switch (op) + { + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_FORCE_PRESENT: + *force_p = true; + break; + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_ALWAYS_TOFROM: + *always_p = true; + break; + case GOMP_MAP_ALWAYS_PRESENT_TO: + case GOMP_MAP_ALWAYS_PRESENT_FROM: + case GOMP_MAP_ALWAYS_PRESENT_TOFROM: + *always_p = true; + /* Fallthrough. */ + case GOMP_MAP_PRESENT_ALLOC: + case GOMP_MAP_PRESENT_TO: + case GOMP_MAP_PRESENT_FROM: + case GOMP_MAP_PRESENT_TOFROM: + *present_p = true; + break; + default: + ; + } + + switch (op) + { + case GOMP_MAP_ALLOC: + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_PRESENT_ALLOC: + return GOMP_MAP_ALLOC; + case GOMP_MAP_TO: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_PRESENT_TO: + case GOMP_MAP_ALWAYS_PRESENT_TO: + return GOMP_MAP_TO; + case GOMP_MAP_FROM: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_PRESENT_FROM: + case GOMP_MAP_ALWAYS_PRESENT_FROM: + return GOMP_MAP_FROM; + case GOMP_MAP_TOFROM: + case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_PRESENT_TOFROM: + case GOMP_MAP_ALWAYS_PRESENT_TOFROM: + return GOMP_MAP_TOFROM; + default: + ; + } + + return op; +} + +static enum gomp_map_kind +omp_join_map_kind (enum gomp_map_kind op, bool force_p, bool always_p, + bool present_p) +{ + gcc_assert (!force_p || !(always_p || present_p)); + + switch (op) + { + case GOMP_MAP_ALLOC: + if (force_p) + return GOMP_MAP_FORCE_ALLOC; + else if (present_p) + return GOMP_MAP_PRESENT_ALLOC; + break; + + case GOMP_MAP_TO: + if (force_p) + return GOMP_MAP_FORCE_TO; + else if (always_p && present_p) + return GOMP_MAP_ALWAYS_PRESENT_TO; + else if (always_p) + return GOMP_MAP_ALWAYS_TO; + else if (present_p) + return GOMP_MAP_PRESENT_TO; + break; + + case GOMP_MAP_FROM: + if (force_p) + return GOMP_MAP_FORCE_FROM; + else if (always_p && present_p) + return GOMP_MAP_ALWAYS_PRESENT_FROM; + else if (always_p) + return GOMP_MAP_ALWAYS_FROM; + else if (present_p) + return GOMP_MAP_PRESENT_FROM; + break; + + case GOMP_MAP_TOFROM: + if (force_p) + return GOMP_MAP_FORCE_TOFROM; + else if (always_p && present_p) + return GOMP_MAP_ALWAYS_PRESENT_TOFROM; + else if (always_p) + return GOMP_MAP_ALWAYS_TOFROM; + else if (present_p) + return GOMP_MAP_PRESENT_TOFROM; + break; + + default: + ; + } + + return op; +} + +/* Map kind decay (OpenMP 5.2, 5.8.8 "declare mapper Directive"). Return the + map kind to use given MAPPER_KIND specified in the mapper and INVOKED_AS + specified on the clause that invokes the mapper. See also + fortran/trans-openmp.cc:omp_map_decayed_kind. */ + +static enum gomp_map_kind +omp_map_decayed_kind (enum gomp_map_kind mapper_kind, + enum gomp_map_kind invoked_as, bool exit_p) +{ + if (invoked_as == GOMP_MAP_RELEASE || invoked_as == GOMP_MAP_DELETE) + return invoked_as; + + bool force_p, always_p, present_p; + + invoked_as = omp_split_map_kind (invoked_as, &force_p, &always_p, &present_p); + gomp_map_kind decay_to; + + switch (mapper_kind) + { + case GOMP_MAP_ALLOC: + if (exit_p && invoked_as == GOMP_MAP_FROM) + decay_to = GOMP_MAP_RELEASE; + else + decay_to = GOMP_MAP_ALLOC; + break; + + case GOMP_MAP_TO: + if (invoked_as == GOMP_MAP_FROM) + decay_to = exit_p ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC; + else if (invoked_as == GOMP_MAP_ALLOC) + decay_to = GOMP_MAP_ALLOC; + else + decay_to = GOMP_MAP_TO; + break; + + case GOMP_MAP_FROM: + if (invoked_as == GOMP_MAP_ALLOC || invoked_as == GOMP_MAP_TO) + decay_to = GOMP_MAP_ALLOC; + else + decay_to = GOMP_MAP_FROM; + break; + + case GOMP_MAP_TOFROM: + case GOMP_MAP_UNSET: + decay_to = invoked_as; + break; + + default: + gcc_unreachable (); + } + + return omp_join_map_kind (decay_to, force_p, always_p, present_p); +} + +/* 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 (location_t loc, tree *outlist, tree mapper, tree expr, + enum gomp_map_kind outer_kind, + enum c_omp_region_type ort) +{ + 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) + { + 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); + + OMP_CLAUSE_LOCATION (unshared) = loc; + + enum gomp_map_kind decayed_kind + = omp_map_decayed_kind (clause_kind, outer_kind, + (ort & C_ORT_EXIT_DATA) != 0); + OMP_CLAUSE_SET_MAP_KIND (unshared, decayed_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) + { + outlist = omp_instantiate_mapper (loc, outlist, nested_mapper, + t, outer_kind, ort); + 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, enum c_omp_region_type ort) +{ + 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 (OMP_CLAUSE_LOCATION (c), + pc, mapper, t, kind, ort); + 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 6eae8a50207f..455ab83d5162 100644 --- a/gcc/cp/constexpr.cc +++ b/gcc/cp/constexpr.cc @@ -3300,6 +3300,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; @@ -7160,6 +7163,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: @@ -9725,6 +9729,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 206e791fcfd7..8d006d2e33c7 100644 --- a/gcc/cp/cp-gimplify.cc +++ b/gcc/cp/cp-gimplify.cc @@ -2405,6 +2405,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 80893aa17520..3182b6a2c908 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 eaec574efbe8..d2561fee8092 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -2878,7 +2878,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. */ @@ -4369,6 +4372,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 VAR_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) \ @@ -7750,10 +7758,13 @@ extern tree finish_qualified_id_expr (tree, tree, bool, bool, extern void simplify_aggr_init_expr (tree *); extern void finalize_nrv (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 &); @@ -8338,6 +8349,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 792ab330dd05..6779de0fb7a6 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)) { @@ -8604,14 +8610,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 2083d3b7b047..aec54702f06a 100644 --- a/gcc/cp/decl2.cc +++ b/gcc/cp/decl2.cc @@ -5821,10 +5821,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_context) diff --git a/gcc/cp/error.cc b/gcc/cp/error.cc index c3082d46dbfe..68692a0b80f3 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 (VAR_P (t) && 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 2186d4a65f42..d2de6d480f1f 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -40726,13 +40726,12 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind, 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; @@ -40750,12 +40749,17 @@ 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 present_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); @@ -40778,6 +40782,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) { @@ -40791,6 +40796,71 @@ 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 if (strcmp ("present", p) == 0) { @@ -40804,19 +40874,19 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) return list; } present_modifier = true; - } + cp_lexer_consume_token (parser->lexer); + } else { - cp_parser_error (parser, "% clause with map-type modifier other" - " than %, % or %"); + cp_parser_error (parser, "% clause with map-type modifier " + "other than %, %, " + "% or %"); 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) @@ -40872,8 +40942,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; } @@ -41684,7 +41776,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: @@ -45472,7 +45564,10 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, - "#pragma omp target data", pragma_tok); + "#pragma omp target data", pragma_tok, false); + if (!processing_template_decl) + clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP); + clauses = finish_omp_clauses (clauses, C_ORT_OMP); c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) @@ -45586,7 +45681,11 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok, tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK, - "#pragma omp target enter data", pragma_tok); + "#pragma omp target enter data", pragma_tok, + false); + if (!processing_template_decl) + clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP); + clauses = finish_omp_clauses (clauses, C_ORT_OMP); c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) @@ -45701,7 +45800,11 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok, tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK, - "#pragma omp target exit data", pragma_tok); + "#pragma omp target exit data", pragma_tok, + false); + if (!processing_template_decl) + clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_EXIT_DATA); + clauses = finish_omp_clauses (clauses, C_ORT_OMP_EXIT_DATA); c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) @@ -45993,6 +46096,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, C_ORT_OMP_TARGET); clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET); c_omp_adjust_map_clauses (clauses, true); @@ -48251,6 +48356,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) \ @@ -48291,6 +48562,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); @@ -48304,7 +48581,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 68963d12b3d7..a1dadf4dd461 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -15479,6 +15479,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; @@ -18312,8 +18319,10 @@ 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_DECLARE_SIMD && ort != C_ORT_OMP_DECLARE_MAPPER) { + if (ort & C_ORT_OMP) + new_clauses = c_omp_instantiate_mappers (new_clauses, ort); new_clauses = finish_omp_clauses (new_clauses, ort); if (linear_no_step) for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc)) @@ -19732,7 +19741,9 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl) case OMP_TARGET_UPDATE: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: - tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), C_ORT_OMP, args, + tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), + (TREE_CODE (t) == OMP_TARGET_EXIT_DATA + ? C_ORT_OMP_EXIT_DATA : C_ORT_OMP), args, complain, in_decl); t = copy_node (t); OMP_STANDALONE_CLAUSES (t) = tmp; @@ -19880,6 +19891,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_MAPPER, 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; @@ -27248,7 +27275,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 e182914266f3..d6f2da65410f 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 @@ -6077,6 +6078,101 @@ 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); + + if (TYPE_PTR_P (TREE_TYPE (t))) + t = build_array_ref (loc, t, low); + else + t = error_mark_node; + } + + return t; +} + /* Helper function for cp_parser_omp_declare_reduction_exprs and tsubst_omp_udr. Remove CLEANUP_STMT for data (omp_priv variable). @@ -6558,6 +6654,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. */ @@ -8098,6 +8217,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: @@ -9550,6 +9675,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 @@ -9563,6 +9690,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. */ @@ -9576,6 +9704,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); @@ -9678,10 +9827,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 e797402b59fc..b81804755f12 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 ffc487a3a483..51bb64707f0b 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); } @@ -11679,6 +11682,218 @@ 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 we already have clauses pertaining to a struct variable, then we don't + want to implicitly invoke a user-defined mapper. */ + if ((flags & GOVD_EXPLICIT) != 0 && AGGREGATE_TYPE_P (TREE_TYPE (decl))) + 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 (mapper_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. */ @@ -12390,6 +12605,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; @@ -13421,6 +13647,30 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, || 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; @@ -17200,6 +17450,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 @@ -18122,6 +18381,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 c6d18526360c..27a8916eb7fa 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 9a1a9eccca96..6b32bb8e0c92 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 cca75285fc2e..ca613a9d224c 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 5c7fa3a4aafb..a993b1294157 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-15.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c new file mode 100644 index 000000000000..0db83b6fd335 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c @@ -0,0 +1,59 @@ +/* { dg-do compile { target c++ } } */ +/* { dg-options "-fopenmp -fdump-tree-gimple" } */ + +typedef struct { + int a, b, c, d; +} S; + +int main () +{ + S s; + #pragma omp declare mapper (S x) map(alloc: x.a) map(to: x.b) \ + map(from: x.c) map(tofrom: x.d) + + #pragma omp target enter data map(to: s) + + /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(to:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(to:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + #pragma omp target exit data map(from: s) + + /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: 4\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(from:s\.c \[len: [0-9]+\]\) map\(from:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + + #pragma omp target enter data map(alloc: s) + + /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(alloc:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(alloc:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + #pragma omp target exit data map(release: s) + + /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(release:s\.c \[len: [0-9]+\]\) map\(release:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + + #pragma omp target enter data map(present, to: s) + + /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(force_present:s\.a \[len: [0-9]+\]\) map\(force_present:s\.b \[len: [0-9]+\]\) map\(force_present:s\.c \[len: [0-9]+\]\) map\(force_present:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + #pragma omp target exit data map(present, from: s) + + /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(force_present:s\.c \[len: [0-9]+\]\) map\(force_present:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + + #pragma omp target enter data map(always, to: s) + + /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(always,to:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(always,to:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + #pragma omp target exit data map(always, from: s) + + /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(always,from:s\.c \[len: [0-9]+\]\) map\(always,from:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + + #pragma omp target enter data map(always, present, to: s) + + /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(force_present:s\.a \[len: [0-9]+\]\) map\(always,present,to:s\.b \[len: [0-9]+\]\) map\(force_present:s\.c \[len: [0-9]+\]\) map\(always,present,to:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + #pragma omp target exit data map(always, present, from: s) + + /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(always,present,from:s\.c \[len: [0-9]+\]\) map\(always,present,from:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c new file mode 100644 index 000000000000..f8711f60a39a --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c @@ -0,0 +1,39 @@ +/* { dg-do compile { target c++ } } */ +/* { dg-options "-fopenmp -fdump-tree-gimple" } */ + +typedef struct { + int a, b, c, d; +} S; + +int main () +{ + S s = { 0, 0, 0, 0 }; + #pragma omp declare mapper (S x) map(alloc: x.a) map(to: x.b) \ + map(from: x.c) map(tofrom: x.d) + + #pragma omp target data map(s) + /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(to:s\.b \[len: [0-9]+\]\) map\(from:s\.c \[len: [0-9]+\]\) map\(tofrom:s\.d \[len: [0-9]+\]\)} 3 "gimple" } } */ + { + #pragma omp target + { + s.a++; + s.b++; + s.c++; + s.d++; + } + } + + #pragma omp target data map(alloc: s) + /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(alloc:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(alloc:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + { + #pragma omp target + { + s.a++; + s.b++; + s.c++; + s.d++; + } + } + + return 0; +} 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 014ed35ab415..789088396e86 100644 --- a/gcc/testsuite/c-c++-common/gomp/map-6.c +++ b/gcc/testsuite/c-c++-common/gomp/map-6.c @@ -13,19 +13,19 @@ foo (void) #pragma omp target map (to:a) ; - #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" } */ ; - #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" } */ ; - #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" } */ ; - #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" } */ ; - #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" } */ ; 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..8af3bac10718 --- /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\(alloc: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 668808a29d07..41c6560babf5 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -347,6 +347,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 b6c029c346ef..51aca38638af 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -1015,6 +1015,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_ALWAYS_PRESENT_TOFROM: pp_string (pp, "always,present,tofrom"); 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 (); } @@ -1082,6 +1091,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)) @@ -3851,6 +3877,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 420857b110c4..8929d4af9399 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -269,6 +269,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 */ @@ -360,6 +361,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 5b6b1bab9db6..3eaf79193516 100644 --- a/gcc/tree.def +++ b/gcc/tree.def @@ -1291,6 +1291,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 088e39da532a..7b0fd62aab24 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1572,6 +1572,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) @@ -2019,6 +2026,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 20c722665680..4d9ebdbe9f79 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -209,7 +209,13 @@ enum gomp_map_kind GOMP_MAP_PRESENT_ALLOC = (GOMP_MAP_LAST | 4), GOMP_MAP_PRESENT_TO = (GOMP_MAP_LAST | 5), GOMP_MAP_PRESENT_FROM = (GOMP_MAP_LAST | 6), - GOMP_MAP_PRESENT_TOFROM = (GOMP_MAP_LAST | 7) + GOMP_MAP_PRESENT_TOFROM = (GOMP_MAP_LAST | 7), + /* 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 | 8), + /* Used to record the name of a named mapper. */ + GOMP_MAP_PUSH_MAPPER_NAME = (GOMP_MAP_LAST | 9), + GOMP_MAP_POP_MAPPER_NAME = (GOMP_MAP_LAST | 10) }; #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.41.0