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 6F2CE389A131; Fri, 18 Mar 2022 16:28:19 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 6F2CE389A131 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.90,192,1643702400"; d="scan'208";a="73169615" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 18 Mar 2022 08:28:20 -0800 IronPort-SDR: PAsE8CNyX2f4ceRkaKsCCIY2qiSxws/5mI80BmtqTIKL8qQ8jA2gUcghagjJAUuBqhhe/IQuMG NP3HDkirVmma3yyJEw6x3v+9ZvxoHVMZ5Y9xrun/UwSLjSVcj7fC6AubCFi1tfxwMmn/jOl4lM QVwXICxqiiFFSnzd2csXctpWwsQl0l6qszUn8a9SpxMhC/ySre6+VLSF55UvQdUfSthneObVRF enbWLqRccWQBr73CBAMqLwFYHIHzfKbGVxOrM1otZr5rM8o5N3rqgnGXYqc74bg/xviLiU+nUq Gpo= From: Julian Brown To: CC: Jakub Jelinek , Thomas Schwinge , Tobias Burnus , Fortran List Subject: [PATCH v2 11/11] OpenMP: Support OpenMP 5.0 "declare mapper" directives for C Date: Fri, 18 Mar 2022 09:28:06 -0700 Message-ID: X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 Content-Transfer-Encoding: 8bit Content-Type: text/plain X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-06.mgc.mentorg.com (139.181.222.6) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-11.1 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_STOCKGEN, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: fortran@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Fortran mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Fri, 18 Mar 2022 16:28:25 -0000 This patch adds support for "declare mapper" directives (and the "mapper" modifier on "map" clauses) for C. As for C++, arrays of custom-mapped objects are not supported yet. I've taken hints from the existing C support for "declare reduction" directives: this works a little differently from C++ for things such as looking up user-defined reductions (or user-defined mappers, in our case). Some support functions have been pulled out of the C++ FE and shared with the C implementation: several language hooks have been added to facilitate that, given the above differences. (Fortran FE support is TBD.) 2022-03-17 Julian Brown gcc/c-family/ * c-common.h (omp_mapper_list, c_omp_find_nested_mappers, c_omp_instantiate_mappers): Add forward declarations/prototypes. * c-omp.cc (c_omp_find_nested_mappers): New function. (remap_mapper_decl_info): New struct. (remap_mapper_decl_1, omp_instantiate_mapper, c_omp_instantiate_mappers): Add functions. * c-decl.cc (c_omp_mapper_id, c_omp_mapper_decl, c_omp_mapper_lookup, c_omp_extract_mapper_directive, c_omp_map_array_section, c_omp_scan_mapper_bindings_r, c_omp_scan_mapper_bindings): New functions. gcc/c/ * c-objc-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 for C. * c-parser.cc (c_parser_omp_clause_map): Add KIND parameter. Handle mapper modifier. (c_parser_omp_all_clauses): Update call to c_parser_omp_clause_map with new kind argument. (c_parser_omp_target): Instantiate explicit mappers and record bindings for implicit mappers. (c_parser_omp_declare_mapper): Parse "declare mapper" directives. (c_parser_omp_declare): Support "declare mapper". * c-tree.h (c_omp_finish_mapper_clauses, c_omp_mapper_lookup, c_omp_extract_mapper_directive, c_omp_map_array_section, c_omp_mapper_id, c_omp_mapper_decl, c_omp_scan_mapper_bindings, c_omp_instantiate_mappers): Add prototypes. * c-typeck.cc (c_finish_omp_clauses): Handle GOMP_MAP_PUSH_MAPPER_NAME and GOMP_MAP_POP_MAPPER_NAME. (c_omp_finish_mapper_clauses): New function (langhook). gcc/cp/ * cp-objcp-common.h (LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define langhooks for C++. * cp-tree.h (cxx_omp_mapper_lookup, cxx_omp_extract_mapper_directive, cxx_omp_map_array_section): Add prototypes. * parser.cc (cp_parser_omp_target): Use new name for c_omp_instantiate_mappers. * pt.cc (tsubst_omp_clauses): Use new name for c_omp_instantiate_mappers. (omp_mapper_lookup): Rename to... (cxx_omp_mapper_lookup): This. (omp_extract_mapper_directive): Rename to... (cxx_omp_extract_mapper_directive): This. (cxx_omp_map_array_section): New function. (remap_mapper_decl_info, remap_mapper_decl_1, omp_instantiate_mapper, omp_instantiate_mappers, mapper_list, find_nested_mappers): Remove. (omp_target_walk_data): Rename mapper_list to omp_mapper_list. (finish_omp_target_clauses_r): Likewise. Use renamed cxx_omp_mapper_lookup, cxx_omp_extract_mapper_directive and c_omp_find_nested_mappers. (finish_omp_target_clauses): Likewise. gcc/ * gimplify.cc (omp_instantiate_mapper): Use omp_map_array_section langhook to handle (singleton only, for now) array sections. Diagnose attempts to use length >1 array sections with custom mappers. (gimplify_scan_omp_clauses): Use omp_extract_mapper_directive langhook. * langhooks-def.h (lhd_omp_mapper_lookup, lhd_omp_extract_mapper_directive, lhd_omp_map_array_section): Add prototypes. (LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define lang hooks. (LANG_HOOKS_DECLS): Add LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, LANG_HOOKS_OMP_MAP_ARRAY_SECTION. * langhooks.cc (lhd_omp_mapper_lookup, lhd_omp_extract_mapper_directive, lhd_omp_map_array_section): New default definitions of langhooks. * langhooks.h (lang_hooks_for_decls): Add omp_mapper_lookup, omp_extract_mapper_directive, omp_map_array_section. * omp-general.h (omp_mapper_list): New. gcc/testsuite/ * g++.dg/gomp/declare-mapper-3.C: Remove from here. * c-c++-common/gomp/declare-mapper-3.c: Move test here, make C-compatible. * g++.dg/gomp/declare-mapper-4.C: Remove from here. * c-c++-common/gomp/declare-mapper-4.c: Move test here, make C-compatible. * c-c++-common/gomp/declare-mapper-5.c: New test. * c-c++-common/gomp/declare-mapper-6.c: New test. * c-c++-common/gomp/declare-mapper-7.c: New test. * c-c++-common/gomp/declare-mapper-8.c: New test. * c-c++-common/gomp/declare-mapper-9.c: New test. * gcc.dg/gomp/declare-mapper-10.c: New test. * gcc.dg/gomp/declare-mapper-11.c: New test. * c-c++-common/gomp/declare-mapper-12.c: New test. libgomp/ * testsuite/libgomp.c-c++-common/declare-mapper-9.c: New test. * testsuite/libgomp.c-c++-common/declare-mapper-10.c: New test. * testsuite/libgomp.c-c++-common/declare-mapper-11.c: New test. * testsuite/libgomp.c-c++-common/declare-mapper-12.c: New test. * testsuite/libgomp.c-c++-common/declare-mapper-13.c: New test. * testsuite/libgomp.c-c++-common/declare-mapper-14.c: New test. --- gcc/c-family/c-common.h | 3 + gcc/c-family/c-omp.cc | 300 ++++++++++++++ gcc/c/c-decl.cc | 169 ++++++++ gcc/c/c-objc-common.h | 12 + gcc/c/c-parser.cc | 298 +++++++++++++- gcc/c/c-tree.h | 8 + gcc/c/c-typeck.cc | 16 + gcc/cp/cp-objcp-common.h | 7 + gcc/cp/cp-tree.h | 3 + gcc/cp/parser.cc | 2 +- gcc/cp/pt.cc | 2 +- gcc/cp/semantics.cc | 383 ++---------------- gcc/gimplify.cc | 57 +-- gcc/langhooks-def.h | 10 + gcc/langhooks.cc | 26 ++ gcc/langhooks.h | 12 + gcc/omp-general.h | 32 ++ .../c-c++-common/gomp/declare-mapper-12.c | 22 + .../gomp/declare-mapper-3.c} | 9 +- .../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 | 24 ++ .../c-c++-common/gomp/declare-mapper-7.c | 30 ++ .../c-c++-common/gomp/declare-mapper-8.c | 43 ++ .../c-c++-common/gomp/declare-mapper-9.c | 34 ++ gcc/testsuite/g++.dg/gomp/declare-mapper-4.C | 74 ---- gcc/testsuite/gcc.dg/gomp/declare-mapper-10.c | 61 +++ gcc/testsuite/gcc.dg/gomp/declare-mapper-11.c | 33 ++ .../libgomp.c-c++-common/declare-mapper-10.c | 58 +++ .../libgomp.c-c++-common/declare-mapper-11.c | 57 +++ .../libgomp.c-c++-common/declare-mapper-12.c | 85 ++++ .../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 | 60 +++ 34 files changed, 1679 insertions(+), 467 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c rename gcc/testsuite/{g++.dg/gomp/declare-mapper-3.C => c-c++-common/gomp/declare-mapper-3.c} (75%) 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 delete mode 100644 gcc/testsuite/g++.dg/gomp/declare-mapper-4.C create mode 100644 gcc/testsuite/gcc.dg/gomp/declare-mapper-10.c create mode 100644 gcc/testsuite/gcc.dg/gomp/declare-mapper-11.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 e592e7fd368..adebd0a2605 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1252,6 +1252,9 @@ extern tree c_omp_check_context_selector (location_t, tree); extern void c_omp_mark_declare_variant (location_t, tree, tree); extern const char *c_omp_map_clause_name (tree, bool); extern void c_omp_adjust_map_clauses (tree, bool); +struct omp_mapper_list; +extern void c_omp_find_nested_mappers (struct omp_mapper_list *, tree); +extern tree c_omp_instantiate_mappers (tree); class c_omp_address_inspector { diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index 77255dd587a..789da097bb0 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -3396,6 +3396,306 @@ c_omp_address_inspector::get_attachment_point (tree expr) return get_origin (baseptr); } +/* Given a mapper function MAPPER_FN, recursively scan through the map clauses + for that mapper, and if any of those should use a (named or unnamed) mapper + themselves, add it to MLIST. */ + +void +c_omp_find_nested_mappers (omp_mapper_list *mlist, tree mapper_fn) +{ + tree mapper = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn); + tree mapper_name = NULL_TREE; + + if (mapper == error_mark_node) + return; + + gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER); + + for (tree clause = OMP_DECLARE_MAPPER_CLAUSES (mapper); + clause; + clause = OMP_CLAUSE_CHAIN (clause)) + { + tree expr = OMP_CLAUSE_DECL (clause); + enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (clause); + tree elem_type; + + if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME) + { + mapper_name = expr; + continue; + } + else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME) + { + mapper_name = NULL_TREE; + continue; + } + + gcc_assert (TREE_CODE (expr) != TREE_LIST); + if (TREE_CODE (expr) == OMP_ARRAY_SECTION) + { + while (TREE_CODE (expr) == OMP_ARRAY_SECTION) + expr = TREE_OPERAND (expr, 0); + + elem_type = TREE_TYPE (expr); + } + else + elem_type = TREE_TYPE (expr); + + /* This might be too much... or not enough? */ + while (TREE_CODE (elem_type) == ARRAY_TYPE + || TREE_CODE (elem_type) == POINTER_TYPE + || TREE_CODE (elem_type) == REFERENCE_TYPE) + elem_type = TREE_TYPE (elem_type); + + elem_type = TYPE_MAIN_VARIANT (elem_type); + + if (AGGREGATE_TYPE_P (elem_type) + && !mlist->contains (mapper_name, elem_type)) + { + tree nested_mapper_fn + = lang_hooks.decls.omp_mapper_lookup (mapper_name, elem_type); + + if (nested_mapper_fn) + { + mlist->add_mapper (mapper_name, elem_type, nested_mapper_fn); + c_omp_find_nested_mappers (mlist, nested_mapper_fn); + } + else if (mapper_name) + { + error ("mapper %qE not found for type %qT", mapper_name, + elem_type); + continue; + } + } + } +} + +struct remap_mapper_decl_info +{ + tree dummy_var; + tree expr; +}; + +/* Helper for rewriting DUMMY_VAR into EXPR in a map clause decl. */ + +static tree +remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data) +{ + remap_mapper_decl_info *map_info = (remap_mapper_decl_info *) data; + + if (operand_equal_p (*tp, map_info->dummy_var)) + { + *tp = map_info->expr; + *walk_subtrees = 0; + } + + return NULL_TREE; +} + +/* Instantiate a mapper MAPPER for expression EXPR, adding new clauses to + OUTLIST. OUTER_KIND is the mapping kind to use if not already specified in + the mapper declaration. */ + +static tree * +omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, + enum gomp_map_kind outer_kind) +{ + tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper); + tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper); + tree mapper_name = NULL_TREE; + + remap_mapper_decl_info map_info; + map_info.dummy_var = dummy_var; + map_info.expr = expr; + + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + { + tree unshared = unshare_expr (c); + enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (c); + tree t = OMP_CLAUSE_DECL (unshared); + tree type = NULL_TREE; + bool nonunit_array_with_mapper = false; + + if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME) + { + mapper_name = t; + continue; + } + else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME) + { + mapper_name = NULL_TREE; + continue; + } + + if (TREE_CODE (t) == OMP_ARRAY_SECTION) + { + location_t loc = OMP_CLAUSE_LOCATION (c); + tree t2 = lang_hooks.decls.omp_map_array_section (loc, t); + + if (t2 == t) + { + nonunit_array_with_mapper = true; + /* We'd want use the mapper for the element type if this worked: + look that one up. */ + type = TREE_TYPE (TREE_TYPE (t)); + } + else + { + t = t2; + type = TREE_TYPE (t); + } + } + else + type = TREE_TYPE (t); + + gcc_assert (type); + + if (type == error_mark_node) + continue; + + walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL); + + if (OMP_CLAUSE_MAP_KIND (unshared) == GOMP_MAP_UNSET) + OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind); + + type = TYPE_MAIN_VARIANT (type); + + tree mapper_fn = lang_hooks.decls.omp_mapper_lookup (mapper_name, type); + + if (mapper_fn && nonunit_array_with_mapper) + { + sorry ("user-defined mapper with non-unit length array section"); + continue; + } + else if (mapper_fn) + { + tree nested_mapper + = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn); + if (nested_mapper != mapper) + { + if (clause_kind == GOMP_MAP_UNSET) + clause_kind = outer_kind; + + outlist = omp_instantiate_mapper (outlist, nested_mapper, + t, clause_kind); + continue; + } + } + else if (mapper_name) + { + error ("mapper %qE not found for type %qT", mapper_name, type); + continue; + } + + *outlist = unshared; + outlist = &OMP_CLAUSE_CHAIN (unshared); + } + + return outlist; +} + +/* Given a list of CLAUSES, scan each clause and invoke a user-defined mapper + appropriate to the type of the data in that clause, if such a mapper is + visible in the current parsing context. */ + +tree +c_omp_instantiate_mappers (tree clauses) +{ + tree c, *pc, mapper_name = NULL_TREE; + + for (pc = &clauses, c = clauses; c; c = *pc) + { + bool using_mapper = false; + + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_MAP: + { + tree t = OMP_CLAUSE_DECL (c); + tree type = NULL_TREE; + bool nonunit_array_with_mapper = false; + + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME) + { + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME) + mapper_name = OMP_CLAUSE_DECL (c); + else + mapper_name = NULL_TREE; + pc = &OMP_CLAUSE_CHAIN (c); + continue; + } + + if (TREE_CODE (t) == OMP_ARRAY_SECTION) + { + location_t loc = OMP_CLAUSE_LOCATION (c); + tree t2 = lang_hooks.decls.omp_map_array_section (loc, t); + + if (t2 == t) + { + /* !!! Array sections of size >1 with mappers for elements + are hard to support. Do something here. */ + nonunit_array_with_mapper = true; + type = TREE_TYPE (TREE_TYPE (t)); + } + else + { + t = t2; + type = TREE_TYPE (t); + } + } + else + type = TREE_TYPE (t); + + if (type == NULL_TREE || type == error_mark_node) + { + pc = &OMP_CLAUSE_CHAIN (c); + continue; + } + + enum gomp_map_kind kind = OMP_CLAUSE_MAP_KIND (c); + if (kind == GOMP_MAP_UNSET) + kind = GOMP_MAP_TOFROM; + + type = TYPE_MAIN_VARIANT (type); + + tree mapper_fn + = lang_hooks.decls.omp_mapper_lookup (mapper_name, type); + + if (mapper_fn && nonunit_array_with_mapper) + { + sorry ("user-defined mapper with non-unit length " + "array section"); + using_mapper = true; + } + else if (mapper_fn) + { + tree mapper + = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn); + pc = omp_instantiate_mapper (pc, mapper, t, kind); + using_mapper = true; + } + else if (mapper_name) + { + error ("mapper %qE not found for type %qT", mapper_name, type); + using_mapper = true; + } + } + break; + + default: + ; + } + + if (using_mapper) + *pc = OMP_CLAUSE_CHAIN (c); + else + pc = &OMP_CLAUSE_CHAIN (c); + } + + return clauses; +} + static const struct c_omp_directive omp_directives[] = { /* Keep this alphabetically sorted by the first word. Non-null second/third if any should precede null ones. */ diff --git a/gcc/c/c-decl.cc b/gcc/c/c-decl.cc index c701f07befe..64e5faf7137 100644 --- a/gcc/c/c-decl.cc +++ b/gcc/c/c-decl.cc @@ -12458,6 +12458,175 @@ c_check_omp_declare_reduction_r (tree *tp, int *, void *data) return NULL_TREE; } +/* Return identifier to look up for omp declare reduction. */ + +tree +c_omp_mapper_id (tree mapper_id) +{ + const char *p = NULL; + + const char prefix[] = "omp declare mapper "; + + if (mapper_id == NULL_TREE) + p = ""; + else if (TREE_CODE (mapper_id) == IDENTIFIER_NODE) + p = IDENTIFIER_POINTER (mapper_id); + else + return error_mark_node; + + size_t lenp = sizeof (prefix); + size_t len = strlen (p); + char *name = XALLOCAVEC (char, lenp + len); + memcpy (name, prefix, lenp - 1); + memcpy (name + lenp - 1, p, len + 1); + return get_identifier (name); +} + +/* Lookup MAPPER_ID in the current scope, or create an artificial + VAR_DECL, bind it into the current scope and return it. */ + +tree +c_omp_mapper_decl (tree mapper_id) +{ + struct c_binding *b = I_SYMBOL_BINDING (mapper_id); + if (b != NULL && B_IN_CURRENT_SCOPE (b)) + return b->decl; + + tree decl = build_decl (BUILTINS_LOCATION, VAR_DECL, + mapper_id, integer_type_node); + DECL_ARTIFICIAL (decl) = 1; + DECL_EXTERNAL (decl) = 1; + TREE_STATIC (decl) = 1; + TREE_PUBLIC (decl) = 0; + bind (mapper_id, decl, current_scope, true, false, BUILTINS_LOCATION); + return decl; +} + +/* Lookup MAPPER_ID in the first scope where it has entry for TYPE. */ + +tree +c_omp_mapper_lookup (tree mapper_id, tree type) +{ + if (TREE_CODE (type) != RECORD_TYPE + && TREE_CODE (type) != UNION_TYPE) + return NULL_TREE; + + mapper_id = c_omp_mapper_id (mapper_id); + + struct c_binding *b = I_SYMBOL_BINDING (mapper_id); + while (b) + { + tree t; + for (t = DECL_INITIAL (b->decl); t; t = TREE_CHAIN (t)) + if (comptypes (TREE_PURPOSE (t), type)) + return TREE_VALUE (t); + b = b->shadowed; + } + return NULL_TREE; +} + +/* For C, we record a pointer to the mapper itself without wrapping it in an + artificial function or similar. So, just return it. */ + +tree +c_omp_extract_mapper_directive (tree mapper) +{ + return mapper; +} + +/* For now we can handle singleton OMP_ARRAY_SECTIONs with custom mappers, but + nothing more complicated. */ + +tree +c_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; + + t = build_array_ref (loc, t, low); + } + + return t; +} + +/* Helper function for below function. */ + +static tree +c_omp_scan_mapper_bindings_r (tree *tp, int *walk_subtrees, void *ptr) +{ + tree t = *tp; + omp_mapper_list *mlist = (omp_mapper_list *) ptr; + tree aggr_type = NULL_TREE; + + if (TREE_CODE (t) == SIZEOF_EXPR + || TREE_CODE (t) == ALIGNOF_EXPR) + { + *walk_subtrees = 0; + return NULL_TREE; + } + + if (TREE_CODE (t) == OMP_CLAUSE) + return 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 = c_omp_mapper_lookup (NULL_TREE, aggr_type); + if (mapper_fn) + mlist->add_mapper (NULL_TREE, aggr_type, mapper_fn); + } + + return NULL_TREE; +} + +/* Scan an offload region's body, and record uses of struct- or union-typed + variables. Add _mapper_binding_ fake clauses to *CLAUSES_PTR. */ + +void +c_omp_scan_mapper_bindings (location_t loc, tree *clauses_ptr, tree body) +{ + hash_set seen_types; + auto_vec mappers; + omp_mapper_list mlist (&seen_types, &mappers); + + walk_tree_without_duplicates (&body, c_omp_scan_mapper_bindings_r, &mlist); + + unsigned int i; + tree mapper; + FOR_EACH_VEC_ELT (mappers, i, mapper) + c_omp_find_nested_mappers (&mlist, mapper); + + FOR_EACH_VEC_ELT (mappers, i, mapper) + { + if (mapper == error_mark_node) + continue; + tree mapper_name = OMP_DECLARE_MAPPER_ID (mapper); + tree decl = OMP_DECLARE_MAPPER_DECL (mapper); + + 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; + + OMP_CLAUSE_CHAIN (c) = *clauses_ptr; + *clauses_ptr = c; + } +} bool c_check_in_current_scope (tree decl) diff --git a/gcc/c/c-objc-common.h b/gcc/c/c-objc-common.h index 0b60df9750f..a1fdc52054f 100644 --- a/gcc/c/c-objc-common.h +++ b/gcc/c/c-objc-common.h @@ -122,6 +122,18 @@ along with GCC; see the file COPYING3. If not see #undef LANG_HOOKS_OMP_CLAUSE_ASSIGN_OP #define LANG_HOOKS_OMP_CLAUSE_ASSIGN_OP c_omp_clause_copy_ctor +#undef LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES +#define LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES c_omp_finish_mapper_clauses + +#undef LANG_HOOKS_OMP_MAPPER_LOOKUP +#define LANG_HOOKS_OMP_MAPPER_LOOKUP c_omp_mapper_lookup + +#undef LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE +#define LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE c_omp_extract_mapper_directive + +#undef LANG_HOOKS_OMP_MAP_ARRAY_SECTION +#define LANG_HOOKS_OMP_MAP_ARRAY_SECTION c_omp_map_array_section + #undef LANG_HOOKS_TREE_INLINING_VAR_MOD_TYPE_P #define LANG_HOOKS_TREE_INLINING_VAR_MOD_TYPE_P c_vla_unspec_p #endif /* GCC_C_OBJC_COMMON */ diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 1ca03b6a632..c774e9cc567 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -16263,10 +16263,9 @@ c_parser_omp_clause_depend (c_parser *parser, tree list) always | close */ static tree -c_parser_omp_clause_map (c_parser *parser, tree list) +c_parser_omp_clause_map (c_parser *parser, tree list, enum gomp_map_kind kind) { location_t clause_loc = c_parser_peek_token (parser)->location; - enum gomp_map_kind kind = GOMP_MAP_TOFROM; tree nl, c; matching_parens parens; @@ -16285,11 +16284,27 @@ c_parser_omp_clause_map (c_parser *parser, tree list) if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA) pos++; + else if ((c_parser_peek_nth_token_raw (parser, pos + 1)->type + == CPP_OPEN_PAREN) + && ((c_parser_peek_nth_token_raw (parser, pos + 2)->type + == CPP_NAME) + || ((c_parser_peek_nth_token_raw (parser, pos + 2)->type + == CPP_KEYWORD) + && (c_parser_peek_nth_token_raw (parser, + pos + 2)->keyword + == RID_DEFAULT))) + && (c_parser_peek_nth_token_raw (parser, pos + 3)->type + == CPP_CLOSE_PAREN) + && (c_parser_peek_nth_token_raw (parser, pos + 4)->type + == CPP_COMMA)) + pos += 4; pos++; } int always_modifier = 0; int close_modifier = 0; + int mapper_modifier = 0; + tree mapper_name = NULL_TREE; for (int pos = 1; pos < map_kind_pos; ++pos) { c_token *tok = c_parser_peek_token (parser); @@ -16310,6 +16325,7 @@ c_parser_omp_clause_map (c_parser *parser, tree list) return list; } always_modifier++; + c_parser_consume_token (parser); } else if (strcmp ("close", p) == 0) { @@ -16320,6 +16336,60 @@ c_parser_omp_clause_map (c_parser *parser, tree list) return list; } close_modifier++; + c_parser_consume_token (parser); + } + else if (strcmp ("mapper", p) == 0) + { + c_parser_consume_token (parser); + + matching_parens mparens; + if (mparens.require_open (parser)) + { + if (mapper_modifier) + { + c_parser_error (parser, "too many % modifiers"); + /* Assume it's a well-formed mapper modifier, even if it + seems to be in the wrong place. */ + c_parser_consume_token (parser); + mparens.require_close (parser); + parens.skip_until_found_close (parser); + return list; + } + + tok = c_parser_peek_token (parser); + + switch (tok->type) + { + case CPP_NAME: + { + mapper_name = tok->value; + c_parser_consume_token (parser); + } + break; + + case CPP_KEYWORD: + if (tok->keyword == RID_DEFAULT) + { + c_parser_consume_token (parser); + break; + } + /* Fallthrough. */ + + default: + error_at (tok->location, + "expected identifier or %"); + return list; + } + + if (!mparens.require_close (parser)) + { + parens.skip_until_found_close (parser); + return list; + } + + mapper_modifier++; + pos += 3; + } } else { @@ -16329,8 +16399,6 @@ c_parser_omp_clause_map (c_parser *parser, tree list) parens.skip_until_found_close (parser); return list; } - - c_parser_consume_token (parser); } if (c_parser_next_token_is (parser, CPP_NAME) @@ -16363,8 +16431,30 @@ c_parser_omp_clause_map (c_parser *parser, tree list) nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list, true); + tree last_new = NULL_TREE; + for (c = nl; 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) = nl; + nl = 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; + } parens.skip_until_found_close (parser); return nl; @@ -17157,7 +17247,7 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, c_name = "depend"; break; case PRAGMA_OMP_CLAUSE_MAP: - clauses = c_parser_omp_clause_map (parser, clauses); + clauses = c_parser_omp_clause_map (parser, clauses, GOMP_MAP_TOFROM); c_name = "map"; break; case PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR: @@ -21157,7 +21247,7 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p) { location_t loc = c_parser_peek_token (parser)->location; c_parser_consume_pragma (parser); - tree *pc = NULL, stmt, block; + tree *pc = NULL, stmt, block, body, clauses; if (context != pragma_stmt && context != pragma_compound) { @@ -21312,10 +21402,9 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p) stmt = make_node (OMP_TARGET); TREE_TYPE (stmt) = void_type_node; - OMP_TARGET_CLAUSES (stmt) - = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, - "#pragma omp target", false); - for (tree c = OMP_TARGET_CLAUSES (stmt); c; c = OMP_CLAUSE_CHAIN (c)) + clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, + "#pragma omp target", false); + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION) { tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); @@ -21324,14 +21413,19 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p) OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = nc; } - OMP_TARGET_CLAUSES (stmt) - = c_finish_omp_clauses (OMP_TARGET_CLAUSES (stmt), C_ORT_OMP_TARGET); - c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); + clauses = c_omp_instantiate_mappers (clauses); + clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_TARGET); + c_omp_adjust_map_clauses (clauses, true); - pc = &OMP_TARGET_CLAUSES (stmt); keep_next_level (); block = c_begin_compound_stmt (true); - add_stmt (c_parser_omp_structured_block (parser, if_p)); + body = c_parser_omp_structured_block (parser, if_p); + + c_omp_scan_mapper_bindings (loc, &clauses, body); + + add_stmt (body); + OMP_TARGET_CLAUSES (stmt) = clauses; + pc = &OMP_TARGET_CLAUSES (stmt); OMP_TARGET_BODY (stmt) = c_end_compound_stmt (loc, block, true); SET_EXPR_LOCATION (stmt, loc); @@ -22545,6 +22639,172 @@ c_parser_omp_declare_reduction (c_parser *parser, enum pragma_context context) } +/* OpenMP 5.0 + #pragma omp declare mapper ([mapper-identifier :] type var) \ + [clause [ [,] clause ] ... ] new-line */ + +static void +c_parser_omp_declare_mapper (c_parser *parser, enum pragma_context context) +{ + tree type, mapper_name = NULL_TREE, var = NULL_TREE, fndecl, stmt, stmtlist; + tree maplist = NULL_TREE, mapper_id, mapper_decl, t; + c_token *token; + bool nested; + + if (context == pragma_struct || context == pragma_param) + { + error ("%<#pragma omp declare reduction%> not at file or block scope"); + goto fail; + } + + if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + goto fail; + + token = c_parser_peek_token (parser); + + if (c_parser_peek_2nd_token (parser)->type == CPP_COLON) + { + switch (token->type) + { + case CPP_NAME: + mapper_name = token->value; + c_parser_consume_token (parser); + break; + case CPP_KEYWORD: + if (token->keyword == RID_DEFAULT) + { + mapper_name = NULL_TREE; + c_parser_consume_token (parser); + break; + } + /* Fallthrough. */ + default: + error_at (token->location, "expected identifier or %"); + c_parser_skip_to_pragma_eol (parser, false); + return; + } + + if (!c_parser_require (parser, CPP_COLON, "expected %<:%>")) + goto fail; + } + + mapper_id = c_omp_mapper_id (mapper_name); + mapper_decl = c_omp_mapper_decl (mapper_id); + + { + location_t loc = c_parser_peek_token (parser)->location; + struct c_type_name *ctype = c_parser_type_name (parser); + type = groktypename (ctype, NULL, NULL); + if (type == error_mark_node) + goto fail; + if (TREE_CODE (type) != RECORD_TYPE + && TREE_CODE (type) != UNION_TYPE) + { + error_at (loc, "%qT is not a struct or union type in " + "%<#pragma omp declare mapper%>", type); + c_parser_skip_to_pragma_eol (parser, false); + return; + } + for (tree t = DECL_INITIAL (mapper_decl); t; t = TREE_CHAIN (t)) + if (comptypes (TREE_PURPOSE (t), type)) + { + error_at (loc, "redeclaration of %qs %<#pragma omp declare " + "mapper%> for type %qT", IDENTIFIER_POINTER (mapper_id) + + sizeof ("omp declare mapper ") - 1, + type); + tree prevmapper = TREE_VALUE (t); + /* Hmm, this location might not be very accurate. */ + location_t ploc + = DECL_SOURCE_LOCATION (OMP_DECLARE_MAPPER_DECL (prevmapper)); + error_at (ploc, "previous %<#pragma omp declare mapper%>"); + c_parser_skip_to_pragma_eol (parser, false); + return; + } + } + + token = c_parser_peek_token (parser); + if (token->type == CPP_NAME) + { + var = build_decl (token->location, VAR_DECL, token->value, type); + c_parser_consume_token (parser); + DECL_ARTIFICIAL (var) = 1; + } + else + { + error_at (token->location, "expected identifier"); + goto fail; + } + + if (!c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>")) + goto fail; + + nested = current_function_decl != NULL_TREE; + if (nested) + c_push_function_context (); + + fndecl = build_decl (BUILTINS_LOCATION, FUNCTION_DECL, mapper_id, + default_function_type); + current_function_decl = fndecl; + allocate_struct_function (fndecl, true); + push_scope (); + stmtlist = push_stmt_list (); + pushdecl (var); + DECL_CONTEXT (var) = fndecl; + + while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL)) + { + location_t here; + pragma_omp_clause c_kind; + here = c_parser_peek_token (parser)->location; + c_kind = c_parser_omp_clause_name (parser); + if (c_kind != PRAGMA_OMP_CLAUSE_MAP) + { + error_at (here, "unexpected clause"); + goto fail; + } + maplist = c_parser_omp_clause_map (parser, maplist, GOMP_MAP_UNSET); + } + + if (maplist == NULL_TREE) + { + error_at (input_location, "missing % clause"); + goto fail; + } + + stmt = make_node (OMP_DECLARE_MAPPER); + TREE_TYPE (stmt) = void_type_node; + OMP_DECLARE_MAPPER_ID (stmt) = mapper_name; + OMP_DECLARE_MAPPER_TYPE (stmt) = type; + OMP_DECLARE_MAPPER_DECL (stmt) = var; + OMP_DECLARE_MAPPER_CLAUSES (stmt) = maplist; + + add_stmt (stmt); + + pop_stmt_list (stmtlist); + pop_scope (); + + if (cfun->language != NULL) + { + ggc_free (cfun->language); + cfun->language = NULL; + } + set_cfun (NULL); + current_function_decl = NULL_TREE; + + if (nested) + c_pop_function_context (); + + c_parser_skip_to_pragma_eol (parser); + + t = tree_cons (type, stmt, DECL_INITIAL (mapper_decl)); + DECL_INITIAL (mapper_decl) = t; + + return; + + fail: + c_parser_skip_to_pragma_eol (parser); +} + /* OpenMP 4.0 #pragma omp declare simd declare-simd-clauses[optseq] new-line #pragma omp declare reduction (reduction-id : typename-list : expression) \ @@ -22574,6 +22834,12 @@ c_parser_omp_declare (c_parser *parser, enum pragma_context context) c_parser_omp_declare_reduction (parser, context); return false; } + if (strcmp (p, "mapper") == 0) + { + c_parser_consume_token (parser); + c_parser_omp_declare_mapper (parser, context); + return false; + } if (!flag_openmp) /* flag_openmp_simd */ { c_parser_skip_to_pragma_eol (parser, false); diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h index 962b9b23ed6..37fb47566e3 100644 --- a/gcc/c/c-tree.h +++ b/gcc/c/c-tree.h @@ -761,6 +761,10 @@ extern tree c_finish_omp_task (location_t, tree, tree); extern void c_finish_omp_cancel (location_t, tree); extern void c_finish_omp_cancellation_point (location_t, tree); extern tree c_finish_omp_clauses (tree, enum c_omp_region_type); +extern tree c_omp_finish_mapper_clauses (tree); +extern tree c_omp_mapper_lookup (tree, tree); +extern tree c_omp_extract_mapper_directive (tree); +extern tree c_omp_map_array_section (location_t, tree); extern tree c_build_va_arg (location_t, tree, location_t, tree); extern tree c_finish_transaction (location_t, tree, int); extern bool c_tree_equal (tree, tree); @@ -812,6 +816,10 @@ extern tree c_omp_reduction_id (enum tree_code, tree); extern tree c_omp_reduction_decl (tree); extern tree c_omp_reduction_lookup (tree, tree); extern tree c_check_omp_declare_reduction_r (tree *, int *, void *); +extern tree c_omp_mapper_id (tree); +extern tree c_omp_mapper_decl (tree); +extern void c_omp_scan_mapper_bindings (location_t, tree *, tree); +extern tree c_omp_instantiate_mappers (tree); extern bool c_check_in_current_scope (tree); extern void c_pushtag (location_t, tree, tree); extern void c_bind (location_t, tree, bool); diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 98212c6b7f5..d909b61f623 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -14877,6 +14877,13 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_FROM: case OMP_CLAUSE__CACHE_: t = OMP_CLAUSE_DECL (c); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)) + { + remove = true; + break; + } if (TREE_CODE (t) == OMP_ARRAY_SECTION) { if (handle_omp_array_sections (c, ort)) @@ -15642,6 +15649,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) return clauses; } +/* Do processing necessary to make CLAUSES well-formed, where CLAUSES result + from implicit instantiation of user-defined mappers (in gimplify.cc). */ + +tree +c_omp_finish_mapper_clauses (tree clauses) +{ + return c_finish_omp_clauses (clauses, C_ORT_OMP); +} + /* Return code to initialize DST with a copy constructor from SRC. C doesn't have copy constructors nor assignment operators, only for _Atomic vars we need to perform __atomic_load from src into a temporary diff --git a/gcc/cp/cp-objcp-common.h b/gcc/cp/cp-objcp-common.h index 6a0df9cc913..cb5ff2a0acb 100644 --- a/gcc/cp/cp-objcp-common.h +++ b/gcc/cp/cp-objcp-common.h @@ -184,6 +184,13 @@ extern tree cxx_simulate_record_decl (location_t, const char *, #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_MAPPABLE_TYPE diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index 8f634197dcc..7344c1ec794 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -8218,6 +8218,9 @@ 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/parser.cc b/gcc/cp/parser.cc index 47e99dddd34..279864d29b1 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -44665,7 +44665,7 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, OMP_CLAUSE_CHAIN (c) = nc; } if (!processing_template_decl) - clauses = omp_instantiate_mappers (clauses); + clauses = c_omp_instantiate_mappers (clauses); clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET); c_omp_adjust_map_clauses (clauses, true); diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index f09248b09f1..fb995e34ab7 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -17865,7 +17865,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort, if (ort != C_ORT_OMP_DECLARE_SIMD) { if (ort == C_ORT_OMP_TARGET) - new_clauses = omp_instantiate_mappers (new_clauses); + new_clauses = c_omp_instantiate_mappers (new_clauses); new_clauses = finish_omp_clauses (new_clauses, ort); if (linear_no_step) for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc)) diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 84ae3e16d72..21234be3c31 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -5992,8 +5992,8 @@ omp_mapper_id (tree mapper_id, tree type) return get_identifier (name); } -static tree -omp_mapper_lookup (tree id, tree type) +tree +cxx_omp_mapper_lookup (tree id, tree type) { if (TREE_CODE (type) != RECORD_TYPE && TREE_CODE (type) != UNION_TYPE) @@ -6002,8 +6002,8 @@ omp_mapper_lookup (tree id, tree type) return lookup_name (id); } -static tree -omp_extract_mapper_directive (tree fndecl) +tree +cxx_omp_extract_mapper_directive (tree fndecl) { if (BASELINK_P (fndecl)) /* See through BASELINK nodes to the underlying function. */ @@ -6027,6 +6027,31 @@ omp_extract_mapper_directive (tree fndecl) return body; } +/* For now we can handle singleton OMP_ARRAY_SECTIONs with custom mappers, but + nothing more complicated. */ + +tree +cxx_omp_map_array_section (location_t loc, tree t) +{ + tree low = TREE_OPERAND (t, 1); + tree len = TREE_OPERAND (t, 2); + + if (len && integer_onep (len)) + { + t = TREE_OPERAND (t, 0); + + if (!low) + low = integer_zero_node; + + if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE) + t = convert_from_reference (t); + + t = build_array_ref (loc, t, low); + } + + return t; +} + /* Helper function for cp_parser_omp_declare_reduction_exprs and tsubst_omp_udr. Remove CLEANUP_STMT for data (omp_priv variable). @@ -6793,242 +6818,6 @@ cp_oacc_check_attachments (tree c) return false; } -struct remap_mapper_decl_info -{ - tree dummy_var; - tree expr; -}; - -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 tree * -omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, - enum gomp_map_kind outer_kind) -{ - tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper); - tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper); - tree mapper_name = NULL_TREE; - - remap_mapper_decl_info map_info; - map_info.dummy_var = dummy_var; - map_info.expr = expr; - - for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) - { - tree unshared = unshare_expr (c); - enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (c); - tree t = OMP_CLAUSE_DECL (unshared); - tree type = NULL_TREE; - bool nonunit_array_with_mapper = false; - - if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME) - { - mapper_name = t; - continue; - } - else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME) - { - mapper_name = NULL_TREE; - continue; - } - - if (TREE_CODE (t) == OMP_ARRAY_SECTION) - { - tree low = TREE_OPERAND (t, 1); - tree len = TREE_OPERAND (t, 2); - - if (len && integer_onep (len)) - { - t = TREE_OPERAND (t, 0); - - if (POINTER_TYPE_P (TREE_TYPE (t)) - || TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE) - type = TREE_TYPE (TREE_TYPE (t)); - - if (!low) - low = integer_zero_node; - - if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE) - t = convert_from_reference (t); - - t = build_array_ref (OMP_CLAUSE_LOCATION (c), t, low); - } - else - { - type = TREE_TYPE (t); - nonunit_array_with_mapper = true; - } - } - else - type = TREE_TYPE (t); - - gcc_assert (type); - - if (type == error_mark_node) - continue; - - walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL); - - if (OMP_CLAUSE_MAP_KIND (unshared) == GOMP_MAP_UNSET) - OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind); - - type = TYPE_MAIN_VARIANT (type); - - tree mapper_fn = 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 = omp_extract_mapper_directive (mapper_fn); - if (nested_mapper != mapper) - { - if (clause_kind == GOMP_MAP_UNSET) - clause_kind = outer_kind; - - outlist = omp_instantiate_mapper (outlist, nested_mapper, - t, clause_kind); - continue; - } - } - else if (mapper_name) - { - error ("mapper %qE not found for type %qT", mapper_name, type); - continue; - } - - *outlist = unshared; - outlist = &OMP_CLAUSE_CHAIN (unshared); - } - - return outlist; -} - -tree -omp_instantiate_mappers (tree clauses) -{ - tree c, *pc, mapper_name = NULL_TREE; - - for (pc = &clauses, c = clauses; c; c = *pc) - { - bool using_mapper = false; - - switch (OMP_CLAUSE_CODE (c)) - { - case OMP_CLAUSE_MAP: - { - tree t = OMP_CLAUSE_DECL (c); - tree type = NULL_TREE; - bool nonunit_array_with_mapper = false; - - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME) - { - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME) - mapper_name = OMP_CLAUSE_DECL (c); - else - mapper_name = NULL_TREE; - pc = &OMP_CLAUSE_CHAIN (c); - continue; - } - - gcc_assert (TREE_CODE (t) != TREE_LIST); - - if (TREE_CODE (t) == OMP_ARRAY_SECTION) - { - tree low = TREE_OPERAND (t, 1); - tree len = TREE_OPERAND (t, 2); - - if (len && integer_onep (len)) - { - t = TREE_OPERAND (t, 0); - - if (!TREE_TYPE (t)) - { - pc = &OMP_CLAUSE_CHAIN (c); - continue; - } - - if (POINTER_TYPE_P (TREE_TYPE (t)) - || TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE) - type = TREE_TYPE (TREE_TYPE (t)); - - if (!low) - low = integer_zero_node; - } - else - { - /* !!! Array sections of size >1 with mappers for elements - are hard to support. Do something here. */ - nonunit_array_with_mapper = true; - 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 = 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 = omp_extract_mapper_directive (mapper_fn); - pc = omp_instantiate_mapper (pc, mapper, t, kind); - using_mapper = true; - } - else if (mapper_name) - { - error ("mapper %qE not found for type %qT", mapper_name, type); - using_mapper = true; - } - } - break; - - default: - ; - } - - if (using_mapper) - *pc = OMP_CLAUSE_CHAIN (c); - else - pc = &OMP_CLAUSE_CHAIN (c); - } - - return clauses; -} - /* For all elements of CLAUSES, validate them vs OpenMP constraints. Remove any elements from the list that are invalid. */ @@ -9640,108 +9429,6 @@ finish_omp_construct (enum tree_code code, tree body, tree clauses) return add_stmt (stmt); } -struct mapper_list -{ - hash_set *seen_types; - vec *mappers; - - mapper_list (hash_set *s, vec *m) - : seen_types (s), mappers (m) { } - - void add_mapper (tree name, tree 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, tree type) - { - if (!name) - name = void_node; - - return seen_types->contains ({ name, type }); - } -}; - -static void -find_nested_mappers (mapper_list *mlist, tree mapper_fn) -{ - tree mapper = 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); //TREE_CHAIN (expr); - - 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 - = omp_mapper_lookup (mapper_name, elem_type); - - if (nested_mapper_fn) - { - mlist->add_mapper (mapper_name, elem_type, nested_mapper_fn); - find_nested_mappers (mlist, nested_mapper_fn); - } - else if (mapper_name) - { - error ("mapper %qE not found for type %qT", mapper_name, - elem_type); - continue; - } - } - } -} - /* Used to walk OpenMP target directive body. */ struct omp_target_walk_data @@ -9768,7 +9455,7 @@ struct omp_target_walk_data variables when recording lambda_objects_accessed. */ hash_set local_decls; - mapper_list *mappers; + omp_mapper_list *mappers; }; /* Helper function of finish_omp_target_clauses, called via @@ -9782,7 +9469,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; - mapper_list *mlist = data->mappers; + omp_mapper_list *mlist = data->mappers; tree aggr_type = NULL_TREE; /* References inside of these expression codes shouldn't incur any @@ -9808,7 +9495,7 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr) if (aggr_type) { - tree mapper_fn = omp_mapper_lookup (NULL_TREE, 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); } @@ -9918,7 +9605,7 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr) hash_set seen_types; auto_vec mapper_fns; - mapper_list mlist (&seen_types, &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); @@ -9926,13 +9613,13 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr) unsigned int i; tree mapper_fn; FOR_EACH_VEC_ELT (mapper_fns, i, mapper_fn) - find_nested_mappers (&mlist, mapper_fn); + c_omp_find_nested_mappers (&mlist, mapper_fn); auto_vec new_clauses; FOR_EACH_VEC_ELT (mapper_fns, i, mapper_fn) { - tree mapper = omp_extract_mapper_directive (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); diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 6155d11170f..861159687a7 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -10464,24 +10464,35 @@ omp_instantiate_mapper (hash_map *implicit_mappers, continue; } - tree decl = OMP_CLAUSE_DECL (clause), unshared; + tree decl = OMP_CLAUSE_DECL (clause), unshared, type; + bool nonunit_array_with_mapper = false; - if (TREE_CODE (decl) == OMP_ARRAY_SECTION - && TREE_OPERAND (decl, 2) - && integer_onep (TREE_OPERAND (decl, 2))) + if (TREE_CODE (decl) == OMP_ARRAY_SECTION) { - unshared = build_omp_clause (OMP_CLAUSE_LOCATION (clause), - OMP_CLAUSE_CODE (clause)); - tree low = TREE_OPERAND (decl, 1); - if (!low || integer_zerop (low)) - OMP_CLAUSE_DECL (unshared) - = build_fold_indirect_ref (TREE_OPERAND (decl, 0)); + 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 - OMP_CLAUSE_DECL (unshared) = decl; - OMP_CLAUSE_SIZE (unshared) = OMP_CLAUSE_SIZE (clause); + { + 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); + { + unshared = unshare_expr (clause); + type = TREE_TYPE (decl); + } walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL); @@ -10489,12 +10500,18 @@ omp_instantiate_mapper (hash_map *implicit_mappers, OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind); decl = OMP_CLAUSE_DECL (unshared); - tree type = TYPE_MAIN_VARIANT (TREE_TYPE (decl)); + type = TYPE_MAIN_VARIANT (type); tree *nested_mapper_p = implicit_mappers->get ({ mapper_name, type }); if (nested_mapper_p && *nested_mapper_p != mapper) { + if (nonunit_array_with_mapper) + { + sorry ("user-defined mapper with non-unit length array section"); + continue; + } + if (clause_kind == GOMP_MAP_UNSET) clause_kind = outer_kind; @@ -11505,16 +11522,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, tree var = OMP_CLAUSE__MAPPER_BINDING__DECL (c); tree type = TYPE_MAIN_VARIANT (TREE_TYPE (var)); tree fndecl = OMP_CLAUSE__MAPPER_BINDING__MAPPER (c); - tree mapper = DECL_SAVED_TREE (fndecl); - if (TREE_CODE (mapper) == BIND_EXPR) - mapper = BIND_EXPR_BODY (mapper); - if (TREE_CODE (mapper) == STATEMENT_LIST) - { - tree_stmt_iterator tsi = tsi_start (mapper); - gcc_assert (TREE_CODE (tsi_stmt (tsi)) == DECL_EXPR); - tsi_next (&tsi); - mapper = tsi_stmt (tsi); - } + tree mapper + = lang_hooks.decls.omp_extract_mapper_directive (fndecl); gcc_assert (mapper != NULL_TREE && TREE_CODE (mapper) == OMP_DECLARE_MAPPER); ctx->implicit_mappers->put ({ name, type }, mapper); diff --git a/gcc/langhooks-def.h b/gcc/langhooks-def.h index fa49092636a..37237666aa9 100644 --- a/gcc/langhooks-def.h +++ b/gcc/langhooks-def.h @@ -85,6 +85,9 @@ 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_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); @@ -272,6 +275,10 @@ extern tree lhd_unit_size_without_reusable_padding (tree); #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 @@ -306,6 +313,9 @@ extern tree lhd_unit_size_without_reusable_padding (tree); 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 fc51dbe720a..fe4a5177584 100644 --- a/gcc/langhooks.cc +++ b/gcc/langhooks.cc @@ -643,6 +643,32 @@ 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 3bdc12badc9..8cce3b958bb 100644 --- a/gcc/langhooks.h +++ b/gcc/langhooks.h @@ -310,6 +310,18 @@ struct lang_hooks_for_decls 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 f676cc7c493..242212b652c 100644 --- a/gcc/omp-general.h +++ b/gcc/omp-general.h @@ -201,4 +201,36 @@ struct default_hash_traits } }; +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, tree 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, tree type) + { + if (!name) + name = void_node; + + return seen_types->contains ({ name, type }); + } +}; + #endif /* GCC_OMP_GENERAL_H */ 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 00000000000..dffb19db03c --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c @@ -0,0 +1,22 @@ +/* { dg-do compile } */ + +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/g++.dg/gomp/declare-mapper-3.C b/gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c similarity index 75% rename from gcc/testsuite/g++.dg/gomp/declare-mapper-3.C rename to gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c index 92212fd0dbd..2c18610b7cc 100644 --- a/gcc/testsuite/g++.dg/gomp/declare-mapper-3.C +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c @@ -1,6 +1,8 @@ // { dg-do compile } // { dg-additional-options "-fdump-tree-gimple" } +#include + // Test named mapper invocation. struct S { @@ -11,10 +13,11 @@ struct S { int main (int argc, char *argv[]) { int N = 1024; -#pragma omp declare mapper (mapN:S s) map(to:s.ptr, s.size) map(s.ptr[:N]) +#pragma omp declare mapper (mapN:struct S s) map(to:s.ptr, s.size) \ + map(s.ptr[:N]) - S s; - s.ptr = new int[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\(to: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" } } 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 00000000000..39e3ab11419 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c @@ -0,0 +1,78 @@ +/* { dg-do compile } */ +/* { 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 00000000000..a4ff3406811 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c @@ -0,0 +1,26 @@ +/* { dg-do compile } */ + +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 "previous 'pragma omp declare mapper' declaration" "" { 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 "redeclaration of 'pragma omp declare mapper'" "" { 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 "previous 'pragma omp declare mapper' declaration" "" { 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 "redeclaration of 'pragma omp declare mapper'" "" { 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 00000000000..4805d9457cb --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c @@ -0,0 +1,24 @@ +/* { dg-do compile } */ + +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 } */ +/* { dg-error "expected '\\)' before '\\\]' token" "" { target c++ } .-3 } */ + +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 00000000000..d7b99eed80e --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c @@ -0,0 +1,30 @@ +/* { dg-do compile } */ + +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 } */ + /* { dg-error "expected '\\)' before '\\\]' token" "" { target c++ } .-3 } */ + 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 00000000000..dadca282711 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c @@ -0,0 +1,43 @@ +/* { dg-do compile } */ + +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 00000000000..502a902d072 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c @@ -0,0 +1,34 @@ +/* { dg-do compile } */ + +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 "previous 'pragma omp declare mapper' declaration" "" { 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 "redeclaration of 'pragma omp declare mapper'" "" { 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 "previous 'pragma omp declare mapper' declaration" "" { 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'" "" { target c++ } .-2 } */ +} diff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-4.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-4.C deleted file mode 100644 index 85bef470332..00000000000 --- a/gcc/testsuite/g++.dg/gomp/declare-mapper-4.C +++ /dev/null @@ -1,74 +0,0 @@ -// { dg-do compile } -// { dg-additional-options "-fdump-tree-original" } - -// Check mapper binding clauses. - -struct Y { - int z; -}; - -struct Z { - int z; -}; - -#pragma omp declare mapper (Y y) map(tofrom: y) -#pragma omp declare mapper (Z z) map(tofrom: z) - -int foo (void) -{ - Y yy; - Z zz; - int dummy; - -#pragma omp target data map(dummy) - { - #pragma omp target - { - yy.z++; - zz.z++; - } - yy.z++; - } - return yy.z; -} - -struct P -{ - Z *zp; -}; - -int bar (void) -{ - Y yy; - Z zz; - P pp; - Z t; - int dummy; - - pp.zp = &t; - -#pragma omp declare mapper (Y y) map(tofrom: y.z) -#pragma omp declare mapper (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(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" } } -// { dg-final { scan-tree-dump {mapper_binding\(struct Z,omp declare mapper ~1Z\) mapper_binding\(struct P,omp declare mapper ~1P\)} "original" } } diff --git a/gcc/testsuite/gcc.dg/gomp/declare-mapper-10.c b/gcc/testsuite/gcc.dg/gomp/declare-mapper-10.c new file mode 100644 index 00000000000..efc9c136915 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/declare-mapper-10.c @@ -0,0 +1,61 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +// "omp declare mapper" support -- check expansion in gimple. + +#include + +struct S { + int *ptr; + int size; +}; + +#define N 64 + +#pragma omp declare mapper (struct S w) map(w.size, w.ptr, w.ptr[:w.size]) +#pragma omp declare mapper (foo:struct S w) map(to:w.size, w.ptr) \ + map(w.ptr[:w.size]) + +int main (int argc, char *argv[]) +{ + struct S s; + s.ptr = (int *) malloc (sizeof (int) * N); + s.size = N; + +#pragma omp declare mapper (bar:struct 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\(tofrom: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" { target c++ } } } */ +/* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(to:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 1 "gimple" { target c++ } } } */ diff --git a/gcc/testsuite/gcc.dg/gomp/declare-mapper-11.c b/gcc/testsuite/gcc.dg/gomp/declare-mapper-11.c new file mode 100644 index 00000000000..927065e5ea6 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/declare-mapper-11.c @@ -0,0 +1,33 @@ +// { dg-do compile } + +// Error-checking tests for "omp declare mapper". + +typedef struct { + int *ptr; + int size; +} S; + +typedef struct { + int z; +} Z; + +int main (int argc, char *argv[]) +{ +#pragma omp declare mapper (S v) map(v.size, v.ptr[:v.size]) +/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-1 } */ + + /* 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' for type 'S'" "" { target c } .-1 } */ + + /* ...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'" "" { target c } .-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 or union type in '#pragma omp declare mapper'" "" { target c } .-1 } */ + + 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 00000000000..d7bfc2f08b5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c @@ -0,0 +1,58 @@ +#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 00000000000..3c501dfd33a --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c @@ -0,0 +1,57 @@ +#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 00000000000..e81ad9ab2f5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c @@ -0,0 +1,85 @@ +#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 00000000000..c4784ebafdd --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c @@ -0,0 +1,55 @@ +/* { dg-do run } */ + +#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 00000000000..3e6027e3050 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c @@ -0,0 +1,57 @@ +/* { dg-do run } */ + +#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 00000000000..d263d7453c7 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c @@ -0,0 +1,60 @@ +#include +#include +#include + +#define N 64 + +struct A { + int *arr1; + float *arr2; + int arr3[N]; +}; + +int +main (int argc, char *argv[]) +{ + struct A var; + + memset (&var, 0, sizeof var); + var.arr1 = (int *) calloc (N, sizeof (int)); + var.arr2 = (float *) calloc (N, sizeof (float)); + + { + #pragma omp declare mapper (struct A x) map(to: x.arr1) \ + map(tofrom: x.arr1[0:N]) + #pragma omp target + { + for (int i = 0; i < N; i++) + var.arr1[i]++; + } + } + + { + #pragma omp declare mapper (struct A x) map(to: x.arr2) \ + map(tofrom: x.arr2[0:N]) + #pragma omp target + { + for (int i = 0; i < N; i++) + var.arr2[i]++; + } + } + + { + #pragma omp declare mapper (struct A x) map(tofrom: x.arr3[0:N]) + #pragma omp target + { + for (int i = 0; i < N; i++) + var.arr3[i]++; + } + } + + for (int i = 0; i < N; i++) + { + assert (var.arr1[i] == 1); + assert (var.arr2[i] == 1); + assert (var.arr3[i] == 1); + } + + free (var.arr1); + free (var.arr2); +} -- 2.29.2