From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id C070C382EF36; Fri, 23 Dec 2022 12:14:55 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org C070C382EF36 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.96,268,1665475200"; d="scan'208";a="90736108" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 23 Dec 2022 04:14:55 -0800 IronPort-SDR: QOHVADHPYv263kv6Hnaykm6ATOjNEqlUYxky0VUrtiXsjMPGeaEmYByBF66QTLIR3rwWy+HEjg /cYgAWd4/+uDY35/hdze2K/i1URrJhLU+tHSbaTItmx31O4twxeD8lfKTlP7b0TywlHf6qVayy /KzquY1dcYuRcKts9tvA1G1BOJGJKxL5QBRHavUtl8mo5VVXJd0LCXzP6LdnQmZOswtKtgw8E7 bM0lX56PMM5SBr5FH4obPf11s6XrEAs7SEce/3+isJzcOwf+wCHxC4VMdINbastzr1r+PKXW2B NX0= From: Julian Brown To: CC: , Tobias Burnus , "Jakub Jelinek" , Thomas Schwinge Subject: [PATCH v6 10/11] OpenMP: Support OpenMP 5.0 "declare mapper" directives for C Date: Fri, 23 Dec 2022 04:13:03 -0800 Message-ID: <59d66ab42744d047328f75b9a1282782cfa7ca7f.1671796516.git.julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 Content-Transfer-Encoding: 8bit Content-Type: text/plain X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) 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_STOCKGEN,SPF_HELO_PASS,SPF_PASS,TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: This 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). This version of the patch removes some unnecessary function setup/teardown code from c_parser_omp_declare_mapper, and has been rebased (hence simplified) wrt. refactoring done higher up this patch series. 2022-12-23 Julian Brown gcc/c/ * 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. * 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/testsuite/ * c-c++-common/gomp/declare-mapper-4.c: Enable for C. * 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: Enable for C. * gcc.dg/gomp/declare-mapper-10.c: New test. * gcc.dg/gomp/declare-mapper-11.c: New test. libgomp/ * testsuite/libgomp.c-c++-common/declare-mapper-9.c: Enable for C. * 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/c-decl.cc | 169 +++++++++++ gcc/c/c-objc-common.h | 12 + gcc/c/c-parser.cc | 277 +++++++++++++++++- gcc/c/c-tree.h | 8 + gcc/c/c-typeck.cc | 15 + .../c-c++-common/gomp/declare-mapper-12.c | 2 +- .../c-c++-common/gomp/declare-mapper-4.c | 2 +- .../c-c++-common/gomp/declare-mapper-5.c | 2 +- .../c-c++-common/gomp/declare-mapper-6.c | 2 +- .../c-c++-common/gomp/declare-mapper-7.c | 2 +- .../c-c++-common/gomp/declare-mapper-8.c | 2 +- .../c-c++-common/gomp/declare-mapper-9.c | 2 +- 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 | 2 +- .../libgomp.c-c++-common/declare-mapper-11.c | 2 +- .../libgomp.c-c++-common/declare-mapper-12.c | 2 +- .../libgomp.c-c++-common/declare-mapper-13.c | 2 +- .../libgomp.c-c++-common/declare-mapper-14.c | 2 +- .../libgomp.c-c++-common/declare-mapper-9.c | 2 +- 20 files changed, 572 insertions(+), 29 deletions(-) create mode 100644 gcc/testsuite/gcc.dg/gomp/declare-mapper-10.c create mode 100644 gcc/testsuite/gcc.dg/gomp/declare-mapper-11.c diff --git a/gcc/c/c-decl.cc b/gcc/c/c-decl.cc index e47ca6718b3e..de5a41ee0c02 100644 --- a/gcc/c/c-decl.cc +++ b/gcc/c/c-decl.cc @@ -13084,6 +13084,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 0b60df9750f6..a1fdc52054fd 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 31e5c49cbe83..5dca50850b38 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -17221,10 +17221,9 @@ c_parser_omp_clause_doacross (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; @@ -17243,11 +17242,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); @@ -17268,6 +17283,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) { @@ -17278,6 +17294,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 { @@ -17287,8 +17357,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) @@ -17321,8 +17389,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; @@ -18125,7 +18215,7 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, c_name = "doacross"; 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: @@ -22169,7 +22259,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) { @@ -22324,10 +22414,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); @@ -22336,14 +22425,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); @@ -23644,6 +23738,151 @@ 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, stmt, stmtlist; + tree maplist = NULL_TREE, mapper_id, mapper_decl, t; + c_token *token; + + 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; + + push_scope (); + stmtlist = push_stmt_list (); + pushdecl (var); + DECL_CONTEXT (var) = current_function_decl; + + 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) = type; + OMP_DECLARE_MAPPER_ID (stmt) = mapper_name; + OMP_DECLARE_MAPPER_DECL (stmt) = var; + OMP_DECLARE_MAPPER_CLAUSES (stmt) = maplist; + + add_stmt (stmt); + + pop_stmt_list (stmtlist); + pop_scope (); + + 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) \ @@ -23673,6 +23912,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 075518e59e81..964623feb472 100644 --- a/gcc/c/c-tree.h +++ b/gcc/c/c-tree.h @@ -813,6 +813,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); @@ -864,6 +868,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 b2d4679e2a08..5c1b21392838 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -15336,6 +15336,12 @@ c_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: @@ -16178,6 +16184,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/testsuite/c-c++-common/gomp/declare-mapper-12.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c index c4d017036c5e..dffb19db03cd 100644 --- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c @@ -1,4 +1,4 @@ -/* { dg-do compile { target c++ } } */ +/* { dg-do compile } */ struct XYZ { int a; diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c index 6d933e4bf6f4..39e3ab114199 100644 --- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c @@ -1,4 +1,4 @@ -/* { dg-do compile { target c++ } } */ +/* { dg-do compile } */ /* { dg-additional-options "-fdump-tree-original" } */ /* Check mapper binding clauses. */ diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c index f675a8c68902..86f14e76cbf3 100644 --- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c @@ -1,4 +1,4 @@ -/* { dg-do compile { target c++ } } */ +/* { dg-do compile } */ typedef struct S_ { int *myarr; diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c index a2f6c08cdfdd..c13eb8b5816e 100644 --- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c @@ -1,4 +1,4 @@ -/* { dg-do compile { target c++ } } */ +/* { dg-do compile } */ int x = 5; diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c index 1b1be9dbb666..0f8dd25a18dc 100644 --- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c @@ -1,4 +1,4 @@ -/* { dg-do compile { target c++ } } */ +/* { dg-do compile } */ struct Q { int *arr1; diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c index 86ddb942072c..dadca282711c 100644 --- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c @@ -1,4 +1,4 @@ -/* { dg-do compile { target c++ } } */ +/* { dg-do compile } */ struct Q { int *arr1; diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c index 54e58426910e..b568c5a477f0 100644 --- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c @@ -1,4 +1,4 @@ -/* { dg-do compile { target c++ } } */ +/* { dg-do compile } */ int x = 5; 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 000000000000..efc9c1369158 --- /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 000000000000..927065e5ea63 --- /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 index b0fa40929fbc..ca5aef4d9d38 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c @@ -1,4 +1,4 @@ -/* { dg-do run { target c++ } } */ +/* { dg-do run } */ #include #include diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c index b509ddc412c5..942d6a5a6ad2 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c @@ -1,4 +1,4 @@ -/* { dg-do run { target c++ } } */ +/* { dg-do run } */ #include #include diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c index cf8919c22edf..cbedee51683f 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c @@ -1,4 +1,4 @@ -/* { dg-do run { target c++ } } */ +/* { dg-do run } */ #include #include diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c index 99b7eedad90f..c4784ebafdd5 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c @@ -1,4 +1,4 @@ -/* { dg-do run { target c++ } } */ +/* { dg-do run } */ #include diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c index e7108da25fef..3e6027e30508 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c @@ -1,4 +1,4 @@ -/* { dg-do run { target c++ } } */ +/* { dg-do run } */ #include #include diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c index 9f85df53998a..324d53567787 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c @@ -1,4 +1,4 @@ -/* { dg-do run { target c++ } } */ +/* { dg-do run } */ #include #include -- 2.29.2