From patchwork Tue Sep 13 21:04:30 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1198 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a5d:5044:0:0:0:0:0 with SMTP id h4csp2531479wrt; Tue, 13 Sep 2022 14:11:29 -0700 (PDT) X-Google-Smtp-Source: AA6agR5vB0QCJwe355OM5GvzdJlisyep8Rx2l9KHPpg70+NC0YCSShycMdx5LTYZcOO4omCsCXys X-Received: by 2002:a17:907:1c87:b0:741:8199:a59d with SMTP id nb7-20020a1709071c8700b007418199a59dmr23179785ejc.736.1663103489488; Tue, 13 Sep 2022 14:11:29 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1663103489; cv=none; d=google.com; s=arc-20160816; b=OsADy4fOdMv9vqgQbjzmG1zjGjswL+cgi1wdChYizWjKUTprkXSaSCRaHrJR2SjffZ J0HEfKMh09XA5d4o61R9+ag4cP2YEjmeK//w/aC3HQ5YPbgvKTSyib7myG47RYmdPFFF oShRo49i8mdWrrCQCkjmwCbLhaHvyMUNYffjvM8wpLJOPKQBy9uwG8kYTcXzaKBSfvQl pDBJqBf0FAwqEs1yF6pMMLjg3xEBQB8mTJbI1T5dZ/6FnjR+bA9szx9258c4FylVZkn8 TsGSZROoEB3Kd9wsx4nQoarA1bQFH4qrY0SaCAz8KSnWpf6hQiC5W/KjsVDUOtvKmTP8 OAbA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:cc:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:content-transfer-encoding :mime-version:references:in-reply-to:message-id:date:subject:to:from :ironport-sdr:dmarc-filter:delivered-to; bh=c86bq0Lfj/SXcOX9gb5kIwjUMsCoUyEIMu7xu0Sxkp0=; b=NudZF08DJeChJDnKu9facjByAJ35sECypURWRjdzXArE/Xp4jd5AZCqKan39ODYKZ/ MFfUDrsd8v+vypmCMt+ZUeEBwkP+6z9eLq0+oaKJk+6B7kWVx8k9ypkomG8dt5+CFhW1 2tSf90qM+JjCnVFEBPLAJXc3xPNUE3oeN63NvqjPfbVXI6qTQMyoze/gH7s7Bxg+GKGr rjLDRz6jLa9nihPH+Bsz2UQyhgM6KPdCZwhE8Moef2v0yh372TAZVc8BUPVyPw6IsJ+l 3Q6NrCuwaVII1t/0BYfN7Pha9DL9680We+iTFNtelKtytkFiQAEwOHnu+vFEFrDFg4Nw ooBA== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from sourceware.org (server2.sourceware.org. [2620:52:3:1:0:246e:9693:128c]) by mx.google.com with ESMTPS id y22-20020a056402359600b00445f3dd3861si9992775edc.313.2022.09.13.14.11.29 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Tue, 13 Sep 2022 14:11:29 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) client-ip=2620:52:3:1:0:246e:9693:128c; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id B83A93898382 for ; Tue, 13 Sep 2022 21:05:21 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id D35E6388550C; Tue, 13 Sep 2022 21:04:43 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org D35E6388550C 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.93,313,1654588800"; d="scan'208";a="82933104" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 13 Sep 2022 13:04:43 -0800 IronPort-SDR: 8KjrupN6+uioEoP7EShhrlvYoqhAZhihGsZ6YTyK1l7uUxDSRGAowtgPCfW+GtEQK7R+WcBXR0 SsmHER+x91UCUsawAby70Js7GJFIylbZDy0vi5wwCwU2bCMQJSpIOU4eI/vp4yiXd6sEl2UTF4 4wh3qGCw/RkP4sMzU5Plyb9NqC4wM+5VJaqbta7qm7VzFBbaZL+r0Pw1RzEqgQ3zP7OcywLrex nf29oKMoDG09y1yLrE6W5R3LOCcfV40C8MDe92wL1CoQqf8chOG/VwTRd2N7e/IKkLiOOWQuKO b/A= From: Julian Brown To: Subject: [PATCH v3 11/11] FYI/unfinished: OpenMP 5.0 "declare mapper" support for C++ Date: Tue, 13 Sep 2022 14:04:30 -0700 Message-ID: <2d52a6cf5ba904abd98d028a163c1012becf95a6.1663101299.git.julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: Jakub Jelinek , tobias@codesourcery.com, fortran@gcc.gnu.org Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: =?utf-8?q?INBOX?= X-GMAIL-THRID: =?utf-8?q?1743890404437968213?= X-GMAIL-MSGID: =?utf-8?q?1743890404437968213?= This patch implements OpenMP 5.0 "declare mapper" support for C++. This hasn't been fully revised yet following previous review comments, but I am including it in this series to demonstrate the new approach to gimplifying map clauses after "declare mapper" instantiation. The "gimplify_scan_omp_clauses" function is still called twice: firstly (before scanning the offload region's body) with SUPPRESS_GIMPLIFICATION set to true. This sets up variables in the splay tree, etc. but does not gimplify anything. Then, implicit "declare mappers" are instantiated after scanning the region's body, then "gimplify_scan_omp_clauses" is called again, and does the rest of its previous tasks -- builds struct sibling lists, and gimplifies clauses. Then gimplify_adjust_omp_clauses is called, and compilation continues. Does this approach seem OK? Thanks, Julian --- gcc/cp/cp-gimplify.cc | 6 + gcc/cp/cp-objcp-common.h | 2 + gcc/cp/cp-tree.h | 10 + gcc/cp/decl.cc | 18 +- gcc/cp/mangle.cc | 5 +- gcc/cp/name-lookup.cc | 3 +- gcc/cp/parser.cc | 393 +++++++++++++- gcc/cp/pt.cc | 92 +++- gcc/cp/semantics.cc | 495 +++++++++++++++++- gcc/fortran/parse.cc | 3 + gcc/gimplify.cc | 379 ++++++++++++-- gcc/langhooks-def.h | 3 + gcc/langhooks.cc | 9 + gcc/langhooks.h | 4 + gcc/omp-general.h | 52 ++ gcc/testsuite/c-c++-common/gomp/map-6.c | 6 +- gcc/testsuite/g++.dg/gomp/declare-mapper-1.C | 58 ++ gcc/testsuite/g++.dg/gomp/declare-mapper-2.C | 30 ++ gcc/testsuite/g++.dg/gomp/declare-mapper-3.C | 27 + gcc/testsuite/g++.dg/gomp/declare-mapper-4.C | 74 +++ gcc/testsuite/g++.dg/gomp/ind-base-3.C | 1 - gcc/testsuite/g++.dg/gomp/member-array-2.C | 1 - gcc/tree-core.h | 4 + gcc/tree-pretty-print.cc | 42 ++ gcc/tree.cc | 2 + gcc/tree.def | 7 + gcc/tree.h | 21 + include/gomp-constants.h | 8 +- .../testsuite/libgomp.c++/declare-mapper-1.C | 87 +++ .../testsuite/libgomp.c++/declare-mapper-2.C | 55 ++ .../testsuite/libgomp.c++/declare-mapper-3.C | 63 +++ .../testsuite/libgomp.c++/declare-mapper-4.C | 63 +++ .../testsuite/libgomp.c++/declare-mapper-5.C | 52 ++ .../testsuite/libgomp.c++/declare-mapper-6.C | 37 ++ .../testsuite/libgomp.c++/declare-mapper-7.C | 48 ++ .../testsuite/libgomp.c++/declare-mapper-8.C | 61 +++ 36 files changed, 2161 insertions(+), 60 deletions(-) create mode 100644 gcc/testsuite/g++.dg/gomp/declare-mapper-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/declare-mapper-2.C create mode 100644 gcc/testsuite/g++.dg/gomp/declare-mapper-3.C create mode 100644 gcc/testsuite/g++.dg/gomp/declare-mapper-4.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-1.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-2.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-3.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-4.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-5.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-6.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-7.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-8.C diff --git a/gcc/cp/cp-gimplify.cc b/gcc/cp/cp-gimplify.cc index c05be833357..050049f3c1a 100644 --- a/gcc/cp/cp-gimplify.cc +++ b/gcc/cp/cp-gimplify.cc @@ -2282,6 +2282,12 @@ cxx_omp_finish_clause (tree c, gimple_seq *, bool /* openacc */) } } +tree +cxx_omp_finish_mapper_clauses (tree clauses) +{ + return finish_omp_clauses (clauses, C_ORT_OMP); +} + /* Return true if DECL's DECL_VALUE_EXPR (if any) should be disregarded in OpenMP construct, because it is going to be remapped during OpenMP lowering. SHARED is true if DECL diff --git a/gcc/cp/cp-objcp-common.h b/gcc/cp/cp-objcp-common.h index 1a67f14d9b3..b6f72e004d3 100644 --- a/gcc/cp/cp-objcp-common.h +++ b/gcc/cp/cp-objcp-common.h @@ -185,6 +185,8 @@ extern tree cxx_simulate_record_decl (location_t, const char *, #define LANG_HOOKS_OMP_CLAUSE_DTOR cxx_omp_clause_dtor #undef LANG_HOOKS_OMP_FINISH_CLAUSE #define LANG_HOOKS_OMP_FINISH_CLAUSE cxx_omp_finish_clause +#undef LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES +#define LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES cxx_omp_finish_mapper_clauses #undef LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE #define LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE cxx_omp_privatize_by_reference #undef LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index c897da204fe..be41019529d 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -2893,6 +2893,7 @@ struct GTY(()) lang_decl_fn { unsigned this_thunk_p : 1; unsigned omp_declare_reduction_p : 1; + unsigned omp_declare_mapper_p : 1; unsigned has_dependent_explicit_spec_p : 1; unsigned immediate_fn_p : 1; unsigned maybe_deleted : 1; @@ -4289,6 +4290,11 @@ get_vec_init_expr (tree t) #define DECL_OMP_DECLARE_REDUCTION_P(NODE) \ (LANG_DECL_FN_CHECK (DECL_COMMON_CHECK (NODE))->omp_declare_reduction_p) +/* Nonzero if NODE is an artificial FUNCTION_DECL for + #pragma omp declare mapper. */ +#define DECL_OMP_DECLARE_MAPPER_P(NODE) \ + (LANG_DECL_FN_CHECK (DECL_COMMON_CHECK (NODE))->omp_declare_mapper_p) + /* Nonzero if DECL has been declared threadprivate by #pragma omp threadprivate. */ #define CP_DECL_THREADPRIVATE_P(DECL) \ @@ -7657,10 +7663,13 @@ extern tree finish_qualified_id_expr (tree, tree, bool, bool, extern void simplify_aggr_init_expr (tree *); extern void finalize_nrv (tree *, tree, tree); extern tree omp_reduction_id (enum tree_code, tree, tree); +extern tree omp_mapper_id (tree, tree); extern tree cp_remove_omp_priv_cleanup_stmt (tree *, int *, void *); extern bool cp_check_omp_declare_reduction (tree); +extern bool cp_check_omp_declare_mapper (tree); extern void finish_omp_declare_simd_methods (tree); extern tree finish_omp_clauses (tree, enum c_omp_region_type); +extern tree omp_instantiate_mappers (tree); extern tree push_omp_privatization_clauses (bool); extern void pop_omp_privatization_clauses (tree); extern void save_omp_privatization_clauses (vec &); @@ -8212,6 +8221,7 @@ extern tree cxx_omp_clause_copy_ctor (tree, tree, tree); extern tree cxx_omp_clause_assign_op (tree, tree, tree); extern tree cxx_omp_clause_dtor (tree, tree); extern void cxx_omp_finish_clause (tree, gimple_seq *, bool); +extern tree cxx_omp_finish_mapper_clauses (tree); extern bool cxx_omp_privatize_by_reference (const_tree); extern bool cxx_omp_disregard_value_expr (tree, bool); extern void cp_fold_function (tree); diff --git a/gcc/cp/decl.cc b/gcc/cp/decl.cc index b72b2a8456b..48ca679a91a 100644 --- a/gcc/cp/decl.cc +++ b/gcc/cp/decl.cc @@ -1897,6 +1897,18 @@ duplicate_decls (tree newdecl, tree olddecl, bool hiding, bool was_hidden) "previous % declaration"); return error_mark_node; } + else if (TREE_CODE (newdecl) == FUNCTION_DECL + && DECL_OMP_DECLARE_MAPPER_P (newdecl)) + { + /* OMP UDMs are never duplicates either. */ + gcc_assert (DECL_OMP_DECLARE_MAPPER_P (olddecl)); + error_at (newdecl_loc, + "redeclaration of %"); + inform (olddecl_loc, + "previous % declaration"); + return error_mark_node; + + } else if (TREE_CODE (newdecl) == FUNCTION_DECL && ((DECL_TEMPLATE_SPECIALIZATION (olddecl) && (!DECL_TEMPLATE_INFO (newdecl) @@ -17977,7 +17989,8 @@ finish_function (bool inline_p) /* Perform delayed folding before NRV transformation. */ if (!processing_template_decl && !DECL_IMMEDIATE_FUNCTION_P (fndecl) - && !DECL_OMP_DECLARE_REDUCTION_P (fndecl)) + && !DECL_OMP_DECLARE_REDUCTION_P (fndecl) + && !DECL_OMP_DECLARE_MAPPER_P (fndecl)) cp_fold_function (fndecl); /* Set up the named return value optimization, if we can. Candidate @@ -18054,7 +18067,8 @@ finish_function (bool inline_p) /* Genericize before inlining. */ if (!processing_template_decl && !DECL_IMMEDIATE_FUNCTION_P (fndecl) - && !DECL_OMP_DECLARE_REDUCTION_P (fndecl)) + && !DECL_OMP_DECLARE_REDUCTION_P (fndecl) + && !DECL_OMP_DECLARE_MAPPER_P (fndecl)) cp_genericize (fndecl); /* Emit the resumer and destroyer functions now, providing that we have diff --git a/gcc/cp/mangle.cc b/gcc/cp/mangle.cc index 75388e99bfd..d06bff22f23 100644 --- a/gcc/cp/mangle.cc +++ b/gcc/cp/mangle.cc @@ -955,10 +955,11 @@ decl_mangling_context (tree decl) tcontext = CP_DECL_CONTEXT (decl); - /* Ignore the artificial declare reduction functions. */ + /* Ignore the artificial declare reduction and declare mapper functions. */ if (tcontext && TREE_CODE (tcontext) == FUNCTION_DECL - && DECL_OMP_DECLARE_REDUCTION_P (tcontext)) + && (DECL_OMP_DECLARE_REDUCTION_P (tcontext) + || DECL_OMP_DECLARE_MAPPER_P (tcontext))) return decl_mangling_context (tcontext); return tcontext; diff --git a/gcc/cp/name-lookup.cc b/gcc/cp/name-lookup.cc index ce622761a1a..926bf10cc2f 100644 --- a/gcc/cp/name-lookup.cc +++ b/gcc/cp/name-lookup.cc @@ -3343,7 +3343,8 @@ set_decl_context_in_fn (tree ctx, tree decl) gcc_checking_assert (DECL_LOCAL_DECL_P (decl) && (DECL_NAMESPACE_SCOPE_P (decl) || (TREE_CODE (decl) == FUNCTION_DECL - && DECL_OMP_DECLARE_REDUCTION_P (decl)))); + && (DECL_OMP_DECLARE_REDUCTION_P (decl) + || DECL_OMP_DECLARE_MAPPER_P (decl))))); if (!DECL_CONTEXT (decl) /* When parsing the parameter list of a function declarator, diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 80d33c3b674..e87de0a38a6 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -26407,10 +26407,12 @@ cp_parser_class_specifier (cp_parser* parser) { /* OpenMP UDRs need to be parsed before all other functions. */ FOR_EACH_VEC_SAFE_ELT (unparsed_funs_with_definitions, ix, decl) - if (DECL_OMP_DECLARE_REDUCTION_P (decl)) + if (DECL_OMP_DECLARE_REDUCTION_P (decl) + || DECL_OMP_DECLARE_MAPPER_P (decl)) cp_parser_late_parsing_for_member (parser, decl); FOR_EACH_VEC_SAFE_ELT (unparsed_funs_with_definitions, ix, decl) - if (!DECL_OMP_DECLARE_REDUCTION_P (decl)) + if (!DECL_OMP_DECLARE_REDUCTION_P (decl) + && !DECL_OMP_DECLARE_MAPPER_P (decl)) cp_parser_late_parsing_for_member (parser, decl); } else @@ -32392,6 +32394,8 @@ cp_parser_enclosed_template_argument_list (cp_parser* parser) return arguments; } +static bool cp_parser_omp_declare_mapper_maplist (tree, cp_parser *); + /* MEMBER_FUNCTION is a member function, or a friend. If default arguments, or the body of the function have not yet been parsed, parse them now. */ @@ -32454,6 +32458,16 @@ cp_parser_late_parsing_for_member (cp_parser* parser, tree member_function) finish_function (/*inline_p=*/true); cp_check_omp_declare_reduction (member_function); } + else if (DECL_OMP_DECLARE_MAPPER_P (member_function)) + { + parser->lexer->in_pragma = true; + cp_parser_omp_declare_mapper_maplist (member_function, parser); + finish_function (/*inline_p=*/true); + cp_check_omp_declare_mapper (member_function); + /* If this is a template class, this forces the body of the mapper + to be instantiated. */ + DECL_PRESERVE_P (member_function) = 1; + } else /* Now, parse the body of the function. */ cp_parser_function_definition_after_declarator (parser, @@ -39887,13 +39901,12 @@ cp_parser_omp_clause_depend (cp_parser *parser, tree list, location_t loc) map ( [map-type-modifier[,] ...] map-kind: variable-list ) map-type-modifier: - always | close */ + always | close | mapper ( mapper-name ) */ static tree -cp_parser_omp_clause_map (cp_parser *parser, tree list) +cp_parser_omp_clause_map (cp_parser *parser, tree list, enum gomp_map_kind kind) { tree nlist, c; - enum gomp_map_kind kind = GOMP_MAP_TOFROM; if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) return list; @@ -39911,11 +39924,27 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA) pos++; + else if ((cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type + == CPP_OPEN_PAREN) + && ((cp_lexer_peek_nth_token (parser->lexer, pos + 2)->type + == CPP_NAME) + || ((cp_lexer_peek_nth_token (parser->lexer, pos + 2)->type + == CPP_KEYWORD) + && (cp_lexer_peek_nth_token (parser->lexer, + pos + 2)->keyword + == RID_DEFAULT))) + && (cp_lexer_peek_nth_token (parser->lexer, pos + 3)->type + == CPP_CLOSE_PAREN) + && (cp_lexer_peek_nth_token (parser->lexer, pos + 4)->type + == CPP_COMMA)) + pos += 4; pos++; } bool always_modifier = false; bool close_modifier = false; + bool mapper_modifier = false; + tree mapper_name = NULL_TREE; for (int pos = 1; pos < map_kind_pos; ++pos) { cp_token *tok = cp_lexer_peek_token (parser->lexer); @@ -39938,6 +39967,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) return list; } always_modifier = true; + cp_lexer_consume_token (parser->lexer); } else if (strcmp ("close", p) == 0) { @@ -39951,20 +39981,83 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) return list; } close_modifier = true; + cp_lexer_consume_token (parser->lexer); + } + else if (strcmp ("mapper", p) == 0) + { + cp_lexer_consume_token (parser->lexer); + + matching_parens parens; + if (parens.require_open (parser)) + { + if (mapper_modifier) + { + cp_parser_error (parser, "too many % modifiers"); + /* Assume it's a well-formed mapper modifier, even if it + seems to be in the wrong place. */ + cp_lexer_consume_token (parser->lexer); + parens.require_close (parser); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/ + true); + return list; + } + + tok = cp_lexer_peek_token (parser->lexer); + switch (tok->type) + { + case CPP_NAME: + { + cp_expr e = cp_parser_identifier (parser); + if (e != error_mark_node) + mapper_name = e; + else + goto err; + } + break; + + case CPP_KEYWORD: + if (tok->keyword == RID_DEFAULT) + { + cp_lexer_consume_token (parser->lexer); + break; + } + /* Fallthrough. */ + + default: + err: + cp_parser_error (parser, + "expected identifier or %"); + return list; + } + + if (!parens.require_close (parser)) + { + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/ + true); + return list; + } + + mapper_modifier = true; + pos += 3; + } } else { cp_parser_error (parser, "%<#pragma omp target%> with " - "modifier other than % or " - "% on % clause"); + "modifier other than %, % " + "or % on % clause"); cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, /*or_comma=*/false, /*consume_paren=*/true); return list; } - - cp_lexer_consume_token (parser->lexer); } if (cp_lexer_next_token_is (parser->lexer, CPP_NAME) @@ -40005,8 +40098,30 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list, NULL, true); + tree last_new = NULL_TREE; + for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c)) - OMP_CLAUSE_SET_MAP_KIND (c, kind); + { + OMP_CLAUSE_SET_MAP_KIND (c, kind); + last_new = c; + } + + if (mapper_name) + { + tree name = build_omp_clause (input_location, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_PUSH_MAPPER_NAME); + OMP_CLAUSE_DECL (name) = mapper_name; + OMP_CLAUSE_CHAIN (name) = nlist; + nlist = name; + + gcc_assert (last_new); + + name = build_omp_clause (input_location, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_POP_MAPPER_NAME); + OMP_CLAUSE_DECL (name) = null_pointer_node; + OMP_CLAUSE_CHAIN (name) = OMP_CLAUSE_CHAIN (last_new); + OMP_CLAUSE_CHAIN (last_new) = name; + } return nlist; } @@ -40817,7 +40932,7 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, c_name = "detach"; break; case PRAGMA_OMP_CLAUSE_MAP: - clauses = cp_parser_omp_clause_map (parser, clauses); + clauses = cp_parser_omp_clause_map (parser, clauses, GOMP_MAP_TOFROM); c_name = "map"; break; case PRAGMA_OMP_CLAUSE_DEVICE: @@ -45006,6 +45121,8 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = nc; } + if (!processing_template_decl) + clauses = omp_instantiate_mappers (clauses); clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET); c_omp_adjust_map_clauses (clauses, true); @@ -46954,6 +47071,252 @@ cp_parser_omp_declare_reduction (cp_parser *parser, cp_token *pragma_tok, obstack_free (&declarator_obstack, p); } +/* OpenMP 5.0: + Parse a variable name and a list of map clauses for "omp declare mapper" + directives: + + ... var) [clause[[,] clause] ... ] new-line */ + +static bool +cp_parser_omp_declare_mapper_maplist (tree fndecl, cp_parser *parser) +{ + pragma_omp_clause c_kind; + tree maplist = NULL_TREE, stmt = NULL_TREE; + tree mapper_name = NULL_TREE; + tree type = TREE_VALUE (TYPE_ARG_TYPES (TREE_TYPE (fndecl))); + tree id = cp_parser_declarator_id (parser, false); + + if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN)) + return false; + + gcc_assert (TYPE_REF_P (type)); + type = TREE_TYPE (type); + + keep_next_level (true); + tree block = begin_omp_structured_block (); + + tree var = build_lang_decl (VAR_DECL, id, type); + pushdecl (var); + cp_finish_decl (var, NULL_TREE, 0, NULL_TREE, 0); + DECL_ARTIFICIAL (var) = 1; + TREE_USED (var) = 1; + + const char *fnname = IDENTIFIER_POINTER (DECL_NAME (fndecl)); + if (startswith (fnname, "omp declare mapper ")) + fnname += sizeof "omp declare mapper " - 1; + const char *mapname_end = strchr (fnname, '~'); + if (mapname_end && mapname_end != fnname) + { + char *tmp = XALLOCAVEC (char, mapname_end - fnname + 1); + strncpy (tmp, fnname, mapname_end - fnname); + tmp[mapname_end - fnname] = '\0'; + mapper_name = get_identifier (tmp); + } + + while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL)) + { + c_kind = cp_parser_omp_clause_name (parser); + if (c_kind != PRAGMA_OMP_CLAUSE_MAP) + { + if (c_kind != PRAGMA_OMP_CLAUSE_NONE) + cp_parser_error (parser, "unexpected clause"); + finish_omp_structured_block (block); + return false; + } + maplist = cp_parser_omp_clause_map (parser, maplist, GOMP_MAP_UNSET); + if (maplist == NULL_TREE) + { + finish_omp_structured_block (block); + return false; + } + } + + if (maplist == NULL_TREE) + { + cp_parser_error (parser, "missing % clause"); + finish_omp_structured_block (block); + return false; + } + + 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); + + block = finish_omp_structured_block (block); + + add_stmt (block); + + return true; +} + +/* OpenMP 5.0 + #pragma omp declare mapper([mapper-identifier:]type var) \ + [clause[[,] clause] ... ] new-line */ + +static void +cp_parser_omp_declare_mapper (cp_parser *parser, cp_token *pragma_tok, + enum pragma_context) +{ + cp_token *token = NULL; + cp_token *first_token = NULL; + cp_token_cache *cp = NULL; + tree type = NULL_TREE, fndecl = NULL_TREE, block = NULL_TREE; + bool block_scope = false; + /* Don't create location wrapper nodes within "declare mapper" + directives. */ + auto_suppress_location_wrappers sentinel; + tree mapper_name = NULL_TREE; + tree mapper_id, fntype; + + if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + goto fail; + + if (current_function_decl) + block_scope = true; + + token = cp_lexer_peek_token (parser->lexer); + + if (cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON)) + { + switch (token->type) + { + case CPP_NAME: + { + cp_expr e = cp_parser_identifier (parser); + if (e != error_mark_node) + mapper_name = e; + else + goto fail; + } + break; + + case CPP_KEYWORD: + if (token->keyword == RID_DEFAULT) + { + mapper_name = NULL_TREE; + cp_lexer_consume_token (parser->lexer); + break; + } + /* Fallthrough. */ + + default: + cp_parser_error (parser, "expected identifier or %"); + } + + if (!cp_parser_require (parser, CPP_COLON, RT_COLON)) + goto fail; + } + + { + const char *saved_message = parser->type_definition_forbidden_message; + parser->type_definition_forbidden_message + = G_("types may not be defined within %"); + type_id_in_expr_sentinel s (parser); + type = cp_parser_type_id (parser); + parser->type_definition_forbidden_message = saved_message; + } + + if (dependent_type_p (type)) + mapper_id = omp_mapper_id (mapper_name, NULL_TREE); + else + mapper_id = omp_mapper_id (mapper_name, type); + + fntype = build_function_type_list (void_type_node, + cp_build_reference_type (type, false), + NULL_TREE); + fndecl = build_lang_decl (FUNCTION_DECL, mapper_id, fntype); + DECL_SOURCE_LOCATION (fndecl) = pragma_tok->location; + DECL_ARTIFICIAL (fndecl) = 1; + DECL_EXTERNAL (fndecl) = 1; + DECL_DECLARED_INLINE_P (fndecl) = 1; + DECL_IGNORED_P (fndecl) = 1; + DECL_OMP_DECLARE_MAPPER_P (fndecl) = 1; + SET_DECL_ASSEMBLER_NAME (fndecl, get_identifier ("")); + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("gnu_inline"), NULL_TREE, + DECL_ATTRIBUTES (fndecl)); + + if (block_scope) + block = begin_omp_structured_block (); + + first_token = cp_lexer_peek_token (parser->lexer); + + if (processing_template_decl) + fndecl = push_template_decl (fndecl); + + if (block_scope) + { + DECL_CONTEXT (fndecl) = current_function_decl; + DECL_LOCAL_DECL_P (fndecl) = 1; + } + else if (current_class_type) + { + while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL)) + cp_lexer_consume_token (parser->lexer); + cp = cp_token_cache_new (first_token, + cp_lexer_peek_nth_token (parser->lexer, 2)); + DECL_STATIC_FUNCTION_P (fndecl) = 1; + finish_member_declaration (fndecl); + DECL_PENDING_INLINE_INFO (fndecl) = cp; + DECL_PENDING_INLINE_P (fndecl) = 1; + vec_safe_push (unparsed_funs_with_definitions, fndecl); + cp_parser_require_pragma_eol (parser, pragma_tok); + return; + } + else + { + DECL_CONTEXT (fndecl) = current_namespace; + tree d = pushdecl (fndecl); + gcc_checking_assert (d == error_mark_node || d == fndecl); + + start_preparsed_function (fndecl, NULL_TREE, SF_PRE_PARSED); + } + + if (!cp_parser_omp_declare_mapper_maplist (fndecl, parser)) + { + if (block_scope) + finish_omp_structured_block (block); + else + finish_function (false); + goto fail; + } + + if (!block_scope) + { + tree fn = finish_function (/*inline_p=*/false); + expand_or_defer_fn (fn); + } + else + { + DECL_CONTEXT (fndecl) = current_function_decl; + if (DECL_TEMPLATE_INFO (fndecl)) + DECL_CONTEXT (DECL_TI_TEMPLATE (fndecl)) = current_function_decl; + + block = finish_omp_structured_block (block); + if (TREE_CODE (block) == BIND_EXPR) + DECL_SAVED_TREE (fndecl) = BIND_EXPR_BODY (block); + else if (TREE_CODE (block) == STATEMENT_LIST) + DECL_SAVED_TREE (fndecl) = block; + if (processing_template_decl) + add_decl_expr (fndecl); + else + pushdecl (fndecl); + } + + cp_check_omp_declare_mapper (fndecl); + + cp_parser_require_pragma_eol (parser, pragma_tok); + return; + +fail: + cp_parser_skip_to_pragma_eol (parser, pragma_tok); +} + /* OpenMP 4.0 #pragma omp declare simd declare-simd-clauses[optseq] new-line #pragma omp declare reduction (reduction-id : typename-list : expression) \ @@ -46994,6 +47357,12 @@ cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok, context); return false; } + if (strcmp (p, "mapper") == 0) + { + cp_lexer_consume_token (parser->lexer); + cp_parser_omp_declare_mapper (parser, pragma_tok, context); + return false; + } if (!flag_openmp) /* flag_openmp_simd */ { cp_parser_skip_to_pragma_eol (parser, pragma_tok); @@ -47007,7 +47376,7 @@ cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok, } } cp_parser_error (parser, "expected %, %, " - "% or %"); + "%, % or %"); cp_parser_require_pragma_eol (parser, pragma_tok); return false; } diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index a920d86b354..ca7547ec9b0 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -12256,9 +12256,13 @@ instantiate_class_template (tree type) /* Instantiate members marked with attribute used. */ if (r != error_mark_node && DECL_PRESERVE_P (r)) used.safe_push (r); - if (TREE_CODE (r) == FUNCTION_DECL - && DECL_OMP_DECLARE_REDUCTION_P (r)) - cp_check_omp_declare_reduction (r); + if (TREE_CODE (r) == FUNCTION_DECL) + { + if (DECL_OMP_DECLARE_REDUCTION_P (r)) + cp_check_omp_declare_reduction (r); + else if (DECL_OMP_DECLARE_MAPPER_P (r)) + cp_check_omp_declare_mapper (r); + } } else if ((DECL_CLASS_TEMPLATE_P (t) || DECL_IMPLICIT_TYPEDEF_P (t)) && LAMBDA_TYPE_P (TREE_TYPE (t))) @@ -14343,6 +14347,14 @@ tsubst_function_decl (tree t, tree args, tsubst_flags_t complain, DECL_NAME (r) = omp_reduction_id (ERROR_MARK, DECL_NAME (t), argtype); } + else if (DECL_OMP_DECLARE_MAPPER_P (t)) + { + tree argtype + = TREE_TYPE (TREE_VALUE (TYPE_ARG_TYPES (TREE_TYPE (t)))); + argtype = tsubst (argtype, args, complain, in_decl); + if (strchr (IDENTIFIER_POINTER (DECL_NAME (t)), '~') == NULL) + DECL_NAME (r) = omp_mapper_id (DECL_NAME (t), argtype); + } if (member && DECL_CONV_FN_P (r)) /* Type-conversion operator. Reconstruct the name, in @@ -18138,6 +18150,8 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort, new_clauses = nreverse (new_clauses); if (ort != C_ORT_OMP_DECLARE_SIMD) { + if (ort == C_ORT_OMP_TARGET) + new_clauses = 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)) @@ -18877,6 +18891,10 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, && DECL_OMP_DECLARE_REDUCTION_P (decl) && cp_check_omp_declare_reduction (decl)) instantiate_body (pattern_decl, args, decl, true); + else if (TREE_CODE (decl) == FUNCTION_DECL + && DECL_OMP_DECLARE_MAPPER_P (decl) + && cp_check_omp_declare_mapper (decl)) + instantiate_body (pattern_decl, args, decl, true); } else { @@ -19827,6 +19845,66 @@ tsubst_omp_udr (tree t, tree args, tsubst_flags_t complain, tree in_decl) } } +static void +tsubst_omp_udm (tree t, tree args, tsubst_flags_t complain, tree in_decl) +{ + if (t == NULL_TREE || t == error_mark_node) + return; + + gcc_assert ((TREE_CODE (t) == STATEMENT_LIST + || TREE_CODE (t) == OMP_DECLARE_MAPPER) + && current_function_decl); + + tree decl = NULL_TREE, mapper; + + /* The function body is: + + statement-list: + TYPE t; + #pragma omp declare mapper (TYPE t) map(...) + */ + + if (TREE_CODE (t) == STATEMENT_LIST) + { + tree_stmt_iterator tsi = tsi_start (t); + decl = tsi_stmt (tsi); + tsi_next (&tsi); + mapper = tsi_stmt (tsi); + tsi_next (&tsi); + gcc_assert (tsi_end_p (tsi)); + + gcc_assert (TREE_CODE (decl) == DECL_EXPR); + + decl = tsubst (DECL_EXPR_DECL (decl), args, complain, in_decl); + pushdecl (decl); + } + else + fatal_error (input_location, "malformed OpenMP user-defined mapper"); + + if (TREE_CODE (mapper) == DECL_EXPR) + mapper = DECL_EXPR_DECL (mapper); + + gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER); + + tree id = OMP_DECLARE_MAPPER_ID (mapper); + tree type = OMP_DECLARE_MAPPER_TYPE (mapper); + tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper); + + type = tsubst (type, args, complain, in_decl); + /* The _DECLARE_SIMD variant prevents calling finish_omp_clauses on the + substituted OMP clauses just yet. */ + clauses = tsubst_omp_clauses (clauses, C_ORT_OMP_DECLARE_SIMD, args, + complain, in_decl); + + mapper = make_node (OMP_DECLARE_MAPPER); + OMP_DECLARE_MAPPER_ID (mapper) = id; + OMP_DECLARE_MAPPER_TYPE (mapper) = type; + OMP_DECLARE_MAPPER_DECL (mapper) = decl; + OMP_DECLARE_MAPPER_CLAUSES (mapper) = clauses; + SET_EXPR_LOCATION (mapper, EXPR_LOCATION (t)); + add_stmt (mapper); +} + /* T is a postfix-expression that is not being used in a function call. Return the substituted version of T. */ @@ -26602,7 +26680,8 @@ instantiate_body (tree pattern, tree args, tree d, bool nested_p) } else /* Only OMP reductions are nested. */ - gcc_checking_assert (DECL_OMP_DECLARE_REDUCTION_P (code_pattern)); + gcc_checking_assert (DECL_OMP_DECLARE_REDUCTION_P (code_pattern) + || DECL_OMP_DECLARE_MAPPER_P (code_pattern)); vec omp_privatization_save; if (current_function_decl) @@ -26701,6 +26780,9 @@ instantiate_body (tree pattern, tree args, tree d, bool nested_p) if (DECL_OMP_DECLARE_REDUCTION_P (code_pattern)) tsubst_omp_udr (DECL_SAVED_TREE (code_pattern), args, tf_warning_or_error, d); + else if (DECL_OMP_DECLARE_MAPPER_P (code_pattern)) + tsubst_omp_udm (DECL_SAVED_TREE (code_pattern), args, + tf_warning_or_error, d); else { tsubst_expr (DECL_SAVED_TREE (code_pattern), args, @@ -26728,6 +26810,8 @@ instantiate_body (tree pattern, tree args, tree d, bool nested_p) if (DECL_OMP_DECLARE_REDUCTION_P (code_pattern)) cp_check_omp_declare_reduction (d); + else if (DECL_OMP_DECLARE_MAPPER_P (code_pattern)) + cp_check_omp_declare_mapper (d); } /* We're not deferring instantiation any more. */ diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index aa8239c4090..001e2c8eaf9 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -45,6 +45,7 @@ along with GCC; see the file COPYING3. If not see #include "gomp-constants.h" #include "predict.h" #include "memmodel.h" +#include "gimplify.h" /* There routines provide a modular interface to perform many parsing operations. They may therefore be used during actual parsing, or @@ -4762,7 +4763,8 @@ expand_or_defer_fn_1 (tree fn) be handled. */; else if (!at_eof || DECL_IMMEDIATE_FUNCTION_P (fn) - || DECL_OMP_DECLARE_REDUCTION_P (fn)) + || DECL_OMP_DECLARE_REDUCTION_P (fn) + || DECL_OMP_DECLARE_MAPPER_P (fn)) tentative_decl_linkage (fn); else import_export_decl (fn); @@ -4775,6 +4777,7 @@ expand_or_defer_fn_1 (tree fn) && !DECL_REALLY_EXTERN (fn) && !DECL_IMMEDIATE_FUNCTION_P (fn) && !DECL_OMP_DECLARE_REDUCTION_P (fn) + && !DECL_OMP_DECLARE_MAPPER_P (fn) && (flag_keep_inline_functions || (flag_keep_inline_dllexport && lookup_attribute ("dllexport", DECL_ATTRIBUTES (fn))))) @@ -4808,7 +4811,8 @@ expand_or_defer_fn_1 (tree fn) return false; } - if (DECL_OMP_DECLARE_REDUCTION_P (fn)) + if (DECL_OMP_DECLARE_REDUCTION_P (fn) + || DECL_OMP_DECLARE_MAPPER_P (fn)) return false; return true; @@ -5926,6 +5930,76 @@ omp_reduction_lookup (location_t loc, tree id, tree type, tree *baselinkp, return id; } +/* Return identifier to look up for omp declare mapper. */ + +tree +omp_mapper_id (tree mapper_id, tree type) +{ + const char *p = NULL; + const char *m = NULL; + + if (mapper_id == NULL_TREE) + p = ""; + else if (TREE_CODE (mapper_id) == IDENTIFIER_NODE) + p = IDENTIFIER_POINTER (mapper_id); + else + return error_mark_node; + + if (type != NULL_TREE) + m = mangle_type_string (TYPE_MAIN_VARIANT (type)); + + const char prefix[] = "omp declare mapper "; + size_t lenp = sizeof (prefix); + if (strncmp (p, prefix, lenp - 1) == 0) + lenp = 1; + size_t len = strlen (p); + size_t lenm = m ? strlen (m) + 1 : 0; + char *name = XALLOCAVEC (char, lenp + len + lenm); + memcpy (name, prefix, lenp - 1); + memcpy (name + lenp - 1, p, len + 1); + if (m) + { + name[lenp + len - 1] = '~'; + memcpy (name + lenp + len, m, lenm); + } + return get_identifier (name); +} + +static tree +omp_mapper_lookup (tree id, tree type) +{ + if (TREE_CODE (type) != RECORD_TYPE + && TREE_CODE (type) != UNION_TYPE) + return NULL_TREE; + id = omp_mapper_id (id, type); + return lookup_name (id); +} + +static tree +omp_extract_mapper_directive (tree fndecl) +{ + if (BASELINK_P (fndecl)) + /* See through BASELINK nodes to the underlying function. */ + fndecl = BASELINK_FUNCTIONS (fndecl); + + tree body = DECL_SAVED_TREE (fndecl); + + if (TREE_CODE (body) == BIND_EXPR) + body = BIND_EXPR_BODY (body); + + if (TREE_CODE (body) == STATEMENT_LIST) + { + tree_stmt_iterator tsi = tsi_start (body); + gcc_assert (TREE_CODE (tsi_stmt (tsi)) == DECL_EXPR); + tsi_next (&tsi); + body = tsi_stmt (tsi); + } + + gcc_assert (TREE_CODE (body) == OMP_DECLARE_MAPPER); + + return body; +} + /* Helper function for cp_parser_omp_declare_reduction_exprs and tsubst_omp_udr. Remove CLEANUP_STMT for data (omp_priv variable). @@ -6407,6 +6481,31 @@ finish_omp_reduction_clause (tree c, bool *need_default_ctor, bool *need_dtor) return false; } +/* Check an instance of an "omp declare mapper" function. */ + +bool +cp_check_omp_declare_mapper (tree udm) +{ + tree type = TREE_VALUE (TYPE_ARG_TYPES (TREE_TYPE (udm))); + location_t loc = DECL_SOURCE_LOCATION (udm); + gcc_assert (TYPE_REF_P (type)); + type = TREE_TYPE (type); + + if (type == error_mark_node) + return false; + + if (!processing_template_decl + && TREE_CODE (type) != RECORD_TYPE + && TREE_CODE (type) != UNION_TYPE) + { + error_at (loc, "%qT is not a struct, union or class type in " + "%<#pragma omp declare mapper%>", type); + return false; + } + + return true; +} + /* Called from finish_struct_1. linear(this) or linear(this:step) clauses might not be finalized yet because the class has been incomplete when parsing #pragma omp declare simd methods. Fix those up now. */ @@ -6667,6 +6766,242 @@ 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. */ @@ -7949,6 +8284,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_MAP: if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved) goto move_implicit; + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME) + { + remove = true; + break; + } /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: @@ -9378,6 +9719,108 @@ 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 @@ -9403,6 +9846,8 @@ struct omp_target_walk_data /* Local variables declared inside a BIND_EXPR, used to filter out such variables when recording lambda_objects_accessed. */ hash_set local_decls; + + mapper_list *mappers; }; /* Helper function of finish_omp_target_clauses, called via @@ -9416,6 +9861,8 @@ 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; + tree aggr_type = NULL_TREE; /* References inside of these expression codes shouldn't incur any form of mapping, so return early. */ @@ -9429,6 +9876,22 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr) 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 = omp_mapper_lookup (NULL_TREE, aggr_type); + if (mapper_fn) + mlist->add_mapper (NULL_TREE, aggr_type, mapper_fn); + } + if (current_object) { tree this_expr = TREE_OPERAND (current_object, 0); @@ -9532,10 +9995,38 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr) else data.current_closure = NULL_TREE; + hash_set seen_types; + auto_vec mapper_fns; + mapper_list mlist (&seen_types, &mapper_fns); + data.mappers = &mlist; + cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r, &data); + unsigned int i; + tree mapper_fn; + FOR_EACH_VEC_ELT (mapper_fns, i, mapper_fn) + 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); + if (mapper == error_mark_node) + continue; + tree mapper_name = OMP_DECLARE_MAPPER_ID (mapper); + tree decl = OMP_DECLARE_MAPPER_DECL (mapper); + if (BASELINK_P (mapper_fn)) + mapper_fn = BASELINK_FUNCTIONS (mapper_fn); + + tree c = build_omp_clause (loc, OMP_CLAUSE__MAPPER_BINDING_); + OMP_CLAUSE__MAPPER_BINDING__ID (c) = mapper_name; + OMP_CLAUSE__MAPPER_BINDING__DECL (c) = decl; + OMP_CLAUSE__MAPPER_BINDING__MAPPER (c) = mapper_fn; + + new_clauses.safe_push (c); + } + tree omp_target_this_expr = NULL_TREE; tree *explicit_this_deref_map = NULL; if (data.this_expr_accessed) diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc index 80492c952aa..34245e106b7 100644 --- a/gcc/fortran/parse.cc +++ b/gcc/fortran/parse.cc @@ -27,6 +27,9 @@ along with GCC; see the file COPYING3. If not see #include "match.h" #include "parse.h" #include "tree-core.h" +#include "tree.h" +#include "fold-const.h" +#include "tree-hash-traits.h" #include "omp-general.h" /* Current statement label. Zero means no statement label. Because new_st diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index a08e3468788..e14efe780cf 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -219,6 +219,7 @@ struct gimplify_omp_ctx { struct gimplify_omp_ctx *outer_context; splay_tree variables; + hash_map *implicit_mappers; hash_set *privatized_types; tree clauses; /* Iteration variables in an OMP_FOR. */ @@ -452,6 +453,7 @@ new_omp_context (enum omp_region_type region_type) c = XCNEW (struct gimplify_omp_ctx); c->outer_context = gimplify_omp_ctxp; c->variables = splay_tree_new (splay_tree_compare_decl_uid, 0, 0); + c->implicit_mappers = new hash_map; c->privatized_types = new hash_set; c->location = input_location; c->region_type = region_type; @@ -475,6 +477,7 @@ delete_omp_context (struct gimplify_omp_ctx *c) { splay_tree_delete (c->variables); delete c->privatized_types; + delete c->implicit_mappers; c->loop_iter_var.release (); XDELETE (c); } @@ -11328,21 +11331,190 @@ error_out: return success; } -static void -gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, - enum omp_region_type region_type, - enum tree_code code) +struct instantiate_mapper_info { - using namespace omp_addr_tokenizer; - struct gimplify_omp_ctx *ctx, *outer_ctx; - tree c; - tree *prev_list_p = NULL, *orig_list_p = list_p; - int handled_depend_iterators = -1; - int nowait = -1; + hash_set *handled_structs; + tree *mapper_clauses_p; + struct gimplify_omp_ctx *omp_ctx; +}; - ctx = new_omp_context (region_type); +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 = unshare_expr (map_info->expr); + *walk_subtrees = 0; + } + + return NULL_TREE; +} + +static tree * +omp_instantiate_mapper (hash_map *implicit_mappers, + tree mapper, tree expr, enum gomp_map_kind outer_kind, + tree *mapper_clauses_p) +{ + tree mapper_name = NULL_TREE; + gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER); + + tree clause = OMP_DECLARE_MAPPER_CLAUSES (mapper); + tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper); + + remap_mapper_decl_info map_info; + + map_info.dummy_var = dummy_var; + map_info.expr = expr; + + for (; clause; clause = OMP_CLAUSE_CHAIN (clause)) + { + enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (clause); + + if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME) + { + mapper_name = OMP_CLAUSE_DECL (clause); + continue; + } + else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME) + { + mapper_name = NULL_TREE; + continue; + } + + tree decl = OMP_CLAUSE_DECL (clause), unshared; + + if (TREE_CODE (decl) == OMP_ARRAY_SECTION + && TREE_OPERAND (decl, 2) + && integer_onep (TREE_OPERAND (decl, 2))) + { + 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)); + else + OMP_CLAUSE_DECL (unshared) = decl; + OMP_CLAUSE_SIZE (unshared) = OMP_CLAUSE_SIZE (clause); + } + else + unshared = unshare_expr (clause); + + 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); + + decl = OMP_CLAUSE_DECL (unshared); + tree type = TYPE_MAIN_VARIANT (TREE_TYPE (decl)); + + tree *nested_mapper_p = implicit_mappers->get ({ mapper_name, type }); + + if (nested_mapper_p && *nested_mapper_p != mapper) + { + if (clause_kind == GOMP_MAP_UNSET) + clause_kind = outer_kind; + + mapper_clauses_p + = omp_instantiate_mapper (implicit_mappers, *nested_mapper_p, + decl, clause_kind, mapper_clauses_p); + continue; + } + + *mapper_clauses_p = unshared; + mapper_clauses_p = &OMP_CLAUSE_CHAIN (unshared); + } + + return mapper_clauses_p; +} + +/* Scan GROUPS for mappings that map all or part of a struct or union decl, + and add to HANDLED. Implicit user-defined mappers will then not be invoked + for those decls. */ + +static void +omp_find_explicitly_mapped_structs (hash_set *handled, + vec *groups) +{ + if (!groups) + return; + + for (auto &i : *groups) + { + tree node = *i.grp_start; + + if (OMP_CLAUSE_CODE (node) == OMP_CLAUSE_MAP + && DECL_P (OMP_CLAUSE_DECL (node))) + switch (OMP_CLAUSE_MAP_KIND (node) & ~GOMP_MAP_FLAG_FORCE) + { + case GOMP_MAP_ALLOC: + case GOMP_MAP_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_STRUCT: + handled->add (OMP_CLAUSE_DECL (node)); + break; + default: + ; + } + } +} + +static int +omp_instantiate_implicit_mappers (splay_tree_node n, void *data) +{ + tree decl = (tree) n->key; + instantiate_mapper_info *im_info = (instantiate_mapper_info *) data; + gimplify_omp_ctx *ctx = im_info->omp_ctx; + tree *mapper_p = NULL; + tree type = TREE_TYPE (decl); + bool ref_p = false; + + if (TREE_CODE (type) == REFERENCE_TYPE) + { + ref_p = true; + type = TREE_TYPE (type); + } + + type = TYPE_MAIN_VARIANT (type); + + if (DECL_P (decl) && type && AGGREGATE_TYPE_P (type)) + { + gcc_assert (ctx); + mapper_p = ctx->implicit_mappers->get ({ NULL_TREE, type }); + } + + bool handled_p = im_info->handled_structs->contains (decl); + + if (mapper_p && !handled_p) + { + /* If we have a reference, map the pointed-to object rather than the + reference itself. */ + if (ref_p) + decl = build_fold_indirect_ref (decl); + + im_info->mapper_clauses_p + = omp_instantiate_mapper (ctx->implicit_mappers, *mapper_p, decl, + GOMP_MAP_TOFROM, im_info->mapper_clauses_p); + } + + return 0; +} + +static struct gimplify_omp_ctx * +new_omp_context_for_scan (enum omp_region_type region_type, + enum tree_code code) +{ + struct gimplify_omp_ctx *ctx = new_omp_context (region_type); ctx->code = code; - outer_ctx = ctx->outer_context; if (code == OMP_TARGET) { if (!lang_GNU_Fortran ()) @@ -11366,19 +11538,48 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, default: break; } + gimplify_omp_ctxp = ctx; + return ctx; +} - if (code == OMP_TARGET - || code == OMP_TARGET_DATA - || code == OMP_TARGET_ENTER_DATA - || code == OMP_TARGET_EXIT_DATA) +/* Scan the OMP clauses in *LIST_P, installing mappings into a new + and previous omp contexts. */ + +static void +gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, + enum omp_region_type region_type, + enum tree_code code, + struct gimplify_omp_ctx *existing_ctx = NULL, + bool suppress_gimplification = false) +{ + struct gimplify_omp_ctx *ctx, *outer_ctx; + tree c; + tree *prev_list_p = NULL, *orig_list_p = list_p; + int handled_depend_iterators = -1; + int nowait = -1; + bool second_pass = existing_ctx != NULL && !suppress_gimplification; + + if (existing_ctx) + ctx = existing_ctx; + else + ctx = new_omp_context_for_scan (region_type, code); + + outer_ctx = ctx->outer_context; + gimplify_omp_ctxp = outer_ctx; + + if (!suppress_gimplification + && (code == OMP_TARGET + || code == OMP_TARGET_DATA + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA)) { vec *groups; groups = omp_gather_mapping_groups (list_p); + hash_map *grpmap = NULL; + if (groups) { - hash_map *grpmap; grpmap = omp_index_mapping_groups (groups); - omp_resolve_clause_dependencies (code, groups, grpmap); omp_build_struct_sibling_lists (code, region_type, groups, &grpmap, list_p); @@ -11433,6 +11634,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, tree decl; auto_vec addr_tokens; + if (second_pass && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) + { + list_p = &OMP_CLAUSE_CHAIN (c); + continue; + } + switch (OMP_CLAUSE_CODE (c)) { case OMP_CLAUSE_PRIVATE: @@ -11834,8 +12041,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (OMP_CLAUSE_SIZE (c) == NULL_TREE) OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl) : TYPE_SIZE_UNIT (TREE_TYPE (decl)); - if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, - NULL, is_gimple_val, fb_rvalue) == GS_ERROR) + if (!suppress_gimplification + && gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, + NULL, is_gimple_val, fb_rvalue) == GS_ERROR) { remove = true; break; @@ -11854,9 +12062,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, GOVD_FIRSTPRIVATE | GOVD_SEEN); } - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT - && (addr_tokens[0]->type == STRUCTURE_BASE - || addr_tokens[0]->type == ARRAY_BASE) + /* Adding the decl for a struct access: if we're suppressing + gimplification (on the first pass before "declare mapper" + instantiation), we haven't built sibling lists yet, so process + decls that will be turned into GOMP_MAP_STRUCT in the second + pass. */ + if (((!suppress_gimplification + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT + && (addr_tokens[0]->type == STRUCTURE_BASE + || addr_tokens[0]->type == ARRAY_BASE)) + || (suppress_gimplification + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH + && addr_tokens[0]->type == STRUCTURE_BASE)) && addr_tokens[0]->u.structure_base_kind == BASE_DECL) { gcc_assert (addr_tokens[1]->type == ACCESS_METHOD); @@ -11865,6 +12082,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (omp_access_chain_p (addr_tokens, 1)) break; decl = addr_tokens[1]->expr; + if (suppress_gimplification + && splay_tree_lookup (ctx->variables, (splay_tree_key) decl)) + break; flags = GOVD_MAP | GOVD_EXPLICIT; gcc_assert (addr_tokens[1]->u.access_kind != ACCESS_DIRECT @@ -11874,9 +12094,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (TREE_CODE (decl) == TARGET_EXPR) { - if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL, - is_gimple_lvalue, fb_lvalue) - == GS_ERROR) + if (!suppress_gimplification + && (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL, + is_gimple_lvalue, fb_lvalue) + == GS_ERROR)) remove = true; } else if (!DECL_P (decl)) @@ -11954,7 +12175,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } } - if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c)) + if (!suppress_gimplification + && code == OMP_TARGET + && OMP_CLAUSE_MAP_IN_REDUCTION (c)) { /* Don't gimplify *pd fully at this point, as the base will need to be adjusted during omp lowering. */ @@ -12062,8 +12285,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } } } - else if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, - fb_lvalue) == GS_ERROR) + else if (!suppress_gimplification + && gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, + fb_lvalue) == GS_ERROR) { remove = true; break; @@ -12213,6 +12437,29 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } goto do_notice; + case OMP_CLAUSE__MAPPER_BINDING_: + { + tree name = OMP_CLAUSE__MAPPER_BINDING__ID (c); + tree var = OMP_CLAUSE__MAPPER_BINDING__DECL (c); + tree type = TYPE_MAIN_VARIANT (TREE_TYPE (var)); + tree fndecl = OMP_CLAUSE__MAPPER_BINDING__MAPPER (c); + 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); + } + gcc_assert (mapper != NULL_TREE + && TREE_CODE (mapper) == OMP_DECLARE_MAPPER); + ctx->implicit_mappers->put ({ name, type }, mapper); + remove = true; + break; + } + case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_ADDR: flags = GOVD_EXPLICIT; @@ -12760,7 +13007,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, list_p = &OMP_CLAUSE_CHAIN (c); } - if ((region_type & ORT_TARGET) != 0) + if ((region_type & ORT_TARGET) != 0 && !suppress_gimplification) /* If we have a target region, we can push all the attaches to the end of the list (we may have standalone "attach" operations synthesized for GOMP_MAP_STRUCT nodes that must be processed after the attachment point @@ -15984,10 +16231,25 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) } bool save_in_omp_construct = in_omp_construct; + if ((ort & ORT_ACC) == 0) in_omp_construct = false; - gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort, - TREE_CODE (expr)); + + struct gimplify_omp_ctx *ctx = NULL; + + if (TREE_CODE (expr) == OMP_TARGET + || TREE_CODE (expr) == OMP_TARGET_DATA + || TREE_CODE (expr) == OMP_TARGET_ENTER_DATA + || TREE_CODE (expr) == OMP_TARGET_EXIT_DATA) + { + ctx = new_omp_context_for_scan (ort, TREE_CODE (expr)); + gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort, + TREE_CODE (expr), ctx, true); + } + else + gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort, + TREE_CODE (expr)); + if (TREE_CODE (expr) == OMP_TARGET) optimize_target_teams (expr, pre_p); if ((ort & (ORT_TARGET | ORT_TARGET_DATA)) != 0 @@ -16025,6 +16287,45 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) } else gimplify_and_add (OMP_BODY (expr), &body); + if (TREE_CODE (expr) == OMP_TARGET + || TREE_CODE (expr) == OMP_TARGET_DATA + || TREE_CODE (expr) == OMP_TARGET_ENTER_DATA + || TREE_CODE (expr) == OMP_TARGET_EXIT_DATA) + { + tree mapper_clauses = NULL_TREE; + instantiate_mapper_info im_info; + hash_set handled_structs; + + /* If we already have clauses pertaining to a struct variable, then + we don't want to implicitly invoke a user-defined mapper. Scan + through the groups to check what we have already. */ + + splay_tree_foreach (ctx->variables, + omp_find_explicitly_mapped_structs_r, + (void *) &handled_structs); + + im_info.handled_structs = &handled_structs; + im_info.mapper_clauses_p = &mapper_clauses; + im_info.omp_ctx = ctx; + im_info.pre_p = pre_p; + + splay_tree_foreach (ctx->variables, + omp_instantiate_implicit_mappers, + (void *) &im_info); + + if (mapper_clauses) + mapper_clauses + = lang_hooks.decls.omp_finish_mapper_clauses (mapper_clauses); + + /* Stick the implicitly-expanded mapper clauses at the end of the + clause list. */ + tree *tail = &OMP_CLAUSES (expr); + while (*tail) + tail = &OMP_CLAUSE_CHAIN (*tail); + *tail = mapper_clauses; + gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort, + TREE_CODE (expr), ctx); + } gimplify_adjust_omp_clauses (pre_p, body, &OMP_CLAUSES (expr), TREE_CODE (expr)); in_omp_construct = save_in_omp_construct; @@ -16684,6 +16985,16 @@ gimplify_omp_ordered (tree expr, gimple_seq body) return gimple_build_omp_ordered (body, OMP_ORDERED_CLAUSES (expr)); } +static enum gimplify_status +gimplify_omp_declare_mapper (tree *expr_p) +{ + /* We don't want assembler output -- this inhibits it. */ + DECL_DECLARED_INLINE_P (current_function_decl) = 1; + + *expr_p = NULL_TREE; + return GS_ALL_DONE; +} + /* Convert the GENERIC expression tree *EXPR_P to GIMPLE. If the expression produces a value to be used as an operand inside a GIMPLE statement, the value will be stored back in *EXPR_P. This value will @@ -17603,6 +17914,10 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, ret = gimplify_omp_atomic (expr_p, pre_p); break; + case OMP_DECLARE_MAPPER: + ret = gimplify_omp_declare_mapper (expr_p); + break; + case TRANSACTION_EXPR: ret = gimplify_transaction (expr_p, pre_p); break; diff --git a/gcc/langhooks-def.h b/gcc/langhooks-def.h index 0f1bfbf255c..dcb1e8c383e 100644 --- a/gcc/langhooks-def.h +++ b/gcc/langhooks-def.h @@ -85,6 +85,7 @@ extern enum omp_clause_defaultmap_kind lhd_omp_predetermined_mapping (tree); extern tree lhd_omp_assignment (tree, tree, tree); extern void lhd_omp_finish_clause (tree, gimple_seq *, bool); extern tree lhd_omp_array_size (tree, gimple_seq *); +extern tree lhd_omp_finish_mapper_clauses (tree); struct gimplify_omp_ctx; extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *, tree); @@ -273,6 +274,7 @@ extern tree lhd_unit_size_without_reusable_padding (tree); #define LANG_HOOKS_OMP_CLAUSE_LINEAR_CTOR NULL #define LANG_HOOKS_OMP_CLAUSE_DTOR hook_tree_tree_tree_null #define LANG_HOOKS_OMP_FINISH_CLAUSE lhd_omp_finish_clause +#define LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES lhd_omp_finish_mapper_clauses #define LANG_HOOKS_OMP_ALLOCATABLE_P hook_bool_tree_false #define LANG_HOOKS_OMP_SCALAR_P lhd_omp_scalar_p #define LANG_HOOKS_OMP_SCALAR_TARGET_P hook_bool_tree_false @@ -307,6 +309,7 @@ extern tree lhd_unit_size_without_reusable_padding (tree); LANG_HOOKS_OMP_CLAUSE_LINEAR_CTOR, \ LANG_HOOKS_OMP_CLAUSE_DTOR, \ LANG_HOOKS_OMP_FINISH_CLAUSE, \ + LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES, \ LANG_HOOKS_OMP_ALLOCATABLE_P, \ LANG_HOOKS_OMP_SCALAR_P, \ LANG_HOOKS_OMP_SCALAR_TARGET_P, \ diff --git a/gcc/langhooks.cc b/gcc/langhooks.cc index 7d975675990..7960fe05b24 100644 --- a/gcc/langhooks.cc +++ b/gcc/langhooks.cc @@ -642,6 +642,15 @@ lhd_omp_array_size (tree, gimple_seq *) return NULL_TREE; } +/* Finalize clause list C after expanding custom mappers for implicitly-mapped + variables. */ + +tree +lhd_omp_finish_mapper_clauses (tree c) +{ + return c; +} + /* 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 b1b2b0e10f0..8ccbab83db3 100644 --- a/gcc/langhooks.h +++ b/gcc/langhooks.h @@ -313,6 +313,10 @@ struct lang_hooks_for_decls /* Do language specific checking on an implicitly determined clause. */ void (*omp_finish_clause) (tree clause, gimple_seq *pre_p, bool); + /* Finish language-specific processing on mapping nodes after expanding + user-defined mappers. */ + tree (*omp_finish_mapper_clauses) (tree clauses); + /* 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 6fc0b11996a..80325e5e4d7 100644 --- a/gcc/omp-general.h +++ b/gcc/omp-general.h @@ -209,4 +209,56 @@ typedef omp_addr_tokenizer::omp_addr_token omp_addr_token; extern bool omp_parse_expr (vec &, tree); +struct omp_name_type +{ + tree name; + tree type; +}; + +template <> +struct default_hash_traits + : typed_noop_remove +{ + GTY((skip)) typedef omp_name_type value_type; + GTY((skip)) typedef omp_name_type compare_type; + + static hashval_t + hash (omp_name_type p) + { + return p.name ? iterative_hash_expr (p.name, TYPE_UID (p.type)) + : TYPE_UID (p.type); + } + + static const bool empty_zero_p = true; + + static bool + is_empty (omp_name_type p) + { + return p.type == NULL; + } + + static bool + is_deleted (omp_name_type) + { + return false; + } + + static bool + equal (const omp_name_type &a, const omp_name_type &b) + { + if (a.name == NULL_TREE && b.name == NULL_TREE) + return a.type == b.type; + else if (a.name == NULL_TREE || b.name == NULL_TREE) + return false; + else + return a.name == b.name && a.type == b.type; + } + + static void + mark_empty (omp_name_type &e) + { + e.type = NULL; + } +}; + #endif /* GCC_OMP_GENERAL_H */ diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c index c749db845b0..77054b6b56a 100644 --- a/gcc/testsuite/c-c++-common/gomp/map-6.c +++ b/gcc/testsuite/c-c++-common/gomp/map-6.c @@ -13,10 +13,12 @@ foo (void) #pragma omp target map (to:a) ; - #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" } */ + #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" "" { target c } } */ +/* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'mapper'" "" { target c++ } .-1 } */ ; - #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" } */ + #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" "" { target c } } */ +/* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'mapper'" "" { target c++ } .-1 } */ ; #pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */ diff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C new file mode 100644 index 00000000000..3177d20adbc --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C @@ -0,0 +1,58 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } + +// "omp declare mapper" support -- check expansion in gimple. + +struct S { + int *ptr; + int size; +}; + +#define N 64 + +#pragma omp declare mapper (S w) map(w.size, w.ptr, w.ptr[:w.size]) +#pragma omp declare mapper (foo:S w) map(to:w.size, w.ptr) map(w.ptr[:w.size]) + +int main (int argc, char *argv[]) +{ + S s; + s.ptr = new int[N]; + s.size = N; + +#pragma omp declare mapper (bar:S w) map(w.size, w.ptr, w.ptr[:w.size]) + +#pragma omp target + { + for (int i = 0; i < N; i++) + s.ptr[i]++; + } + +#pragma omp target map(tofrom: s) + { + for (int i = 0; i < N; i++) + s.ptr[i]++; + } + +#pragma omp target map(mapper(default), tofrom: s) + { + for (int i = 0; i < N; i++) + s.ptr[i]++; + } + +#pragma omp target map(mapper(foo), alloc: s) + { + for (int i = 0; i < N; i++) + s.ptr[i]++; + } + +#pragma omp target map(mapper(bar), tofrom: s) + { + for (int i = 0; i < N; i++) + s.ptr[i]++; + } + + return 0; +} + +// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(tofrom:s\.size \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 4 "gimple" } } +// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 1 "gimple" } } diff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-2.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-2.C new file mode 100644 index 00000000000..06d999ea654 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/declare-mapper-2.C @@ -0,0 +1,30 @@ +// { dg-do compile } + +// Error-checking tests for "omp declare mapper". + +struct S { + int *ptr; + int size; +}; + +struct Z { + int z; +}; + +int main (int argc, char *argv[]) +{ +#pragma omp declare mapper (S v) map(v.size, v.ptr[:v.size]) + + /* 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'" } + + /* ...and this one doesn't use a "base language identifier" for the mapper + name. */ +#pragma omp declare mapper (case: S v) map (to: v.size) // { dg-error "expected identifier or 'default' before 'case'" } + // { dg-error "expected ':' before 'case'" "" { target *-*-* } .-1 } + + /* A non-struct/class/union type isn't supposed to work. */ +#pragma omp declare mapper (name:Z [5]foo) map (foo[0].z) // { dg-error "'Z \\\[5\\\]' is not a struct, union or class type in '#pragma omp declare mapper'" } + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-3.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-3.C new file mode 100644 index 00000000000..92212fd0dbd --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/declare-mapper-3.C @@ -0,0 +1,27 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } + +// Test named mapper invocation. + +struct S { + int *ptr; + int size; +}; + +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]) + + S s; + s.ptr = new 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" } } + { + for (int i = 0; i < N; i++) + s.ptr[i]++; + } + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-4.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-4.C new file mode 100644 index 00000000000..85bef470332 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/declare-mapper-4.C @@ -0,0 +1,74 @@ +// { 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/g++.dg/gomp/ind-base-3.C b/gcc/testsuite/g++.dg/gomp/ind-base-3.C index dbabaf7680c..7695b1f907e 100644 --- a/gcc/testsuite/g++.dg/gomp/ind-base-3.C +++ b/gcc/testsuite/g++.dg/gomp/ind-base-3.C @@ -24,7 +24,6 @@ int main (int argc, char *argv[]) { #pragma omp target map(choose(&a, &b, i)->x[:10]) /* { dg-message {sorry, unimplemented: unsupported map expression 'choose\(\(& a\), \(& b\), i\)->S::x\[0\]'} "" { target *-*-* } .-1 } */ -/* { dg-message {sorry, unimplemented: unsupported map expression 'choose\(\(& a\), \(& b\), i\)'} "" { target *-*-* } .-2 } */ for (int j = 0; j < 10; j++) choose (&a, &b, i)->x[j]++; } diff --git a/gcc/testsuite/g++.dg/gomp/member-array-2.C b/gcc/testsuite/g++.dg/gomp/member-array-2.C index e60bb5585a1..caf8ece4262 100644 --- a/gcc/testsuite/g++.dg/gomp/member-array-2.C +++ b/gcc/testsuite/g++.dg/gomp/member-array-2.C @@ -70,7 +70,6 @@ main (int argc, char *argv[]) it for now. */ #pragma omp target map(c.get_arr()[:100]) /* { dg-message {sorry, unimplemented: unsupported map expression 'c\.C::get_arr\(\)\[0\]'} "" { target *-*-* } .-1 } */ - /* { dg-message {sorry, unimplemented: unsupported map expression 'c\.C::get_arr\(\)'} "" { target *-*-* } .-2 } */ #pragma omp teams distribute parallel for for (int i = 0; i < 100; i++) c.get_arr()[i] += 2; diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 86a07c282af..c13564b85f5 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -345,6 +345,10 @@ enum omp_clause_code { /* OpenMP clause: has_device_addr (variable-list). */ OMP_CLAUSE_HAS_DEVICE_ADDR, + /* OpenMP mapper binding: record implicit mappers in scope for aggregate + types used within an offload region. */ + OMP_CLAUSE__MAPPER_BINDING_, + /* Internal structure to hold OpenACC cache directive's variable-list. #pragma acc cache (variable-list). */ OMP_CLAUSE__CACHE_, diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index ccbafa98699..beb0073800c 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -973,6 +973,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: pp_string (pp, "attach_zero_length_array_section"); break; + case GOMP_MAP_UNSET: + pp_string (pp, "unset"); + break; + case GOMP_MAP_PUSH_MAPPER_NAME: + pp_string (pp, "push_mapper"); + break; + case GOMP_MAP_POP_MAPPER_NAME: + pp_string (pp, "pop_mapper"); + break; default: gcc_unreachable (); } @@ -1036,6 +1045,23 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) spc, flags, false); goto print_clause_size; + case OMP_CLAUSE__MAPPER_BINDING_: + pp_string (pp, "mapper_binding("); + if (OMP_CLAUSE__MAPPER_BINDING__ID (clause)) + { + dump_generic_node (pp, OMP_CLAUSE__MAPPER_BINDING__ID (clause), spc, + flags, false); + pp_comma (pp); + } + dump_generic_node (pp, + TREE_TYPE (OMP_CLAUSE__MAPPER_BINDING__DECL (clause)), + spc, flags, false); + pp_comma (pp); + dump_generic_node (pp, OMP_CLAUSE__MAPPER_BINDING__MAPPER (clause), spc, + flags, false); + pp_right_paren (pp); + break; + case OMP_CLAUSE_NUM_TEAMS: pp_string (pp, "num_teams("); if (OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (clause)) @@ -3806,6 +3832,22 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags, is_expr = false; break; + case OMP_DECLARE_MAPPER: + pp_string (pp, "#pragma omp declare mapper ("); + if (OMP_DECLARE_MAPPER_ID (node)) + { + dump_generic_node (pp, OMP_DECLARE_MAPPER_ID (node), spc, flags, + false); + pp_colon (pp); + } + dump_generic_node (pp, OMP_DECLARE_MAPPER_TYPE (node), spc, flags, false); + pp_space (pp); + dump_generic_node (pp, OMP_DECLARE_MAPPER_DECL (node), spc, flags, false); + pp_right_paren (pp); + dump_omp_clauses (pp, OMP_DECLARE_MAPPER_CLAUSES (node), spc, flags); + is_expr = false; + break; + case TRANSACTION_EXPR: if (TRANSACTION_EXPR_OUTER (node)) pp_string (pp, "__transaction_atomic [[outer]]"); diff --git a/gcc/tree.cc b/gcc/tree.cc index 007c9325b17..f28570d5e16 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -294,6 +294,7 @@ unsigned const char omp_clause_num_ops[] = 2, /* OMP_CLAUSE_TO */ 2, /* OMP_CLAUSE_MAP */ 1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */ + 3, /* OMP_CLAUSE__MAPPER_BINDING_ */ 2, /* OMP_CLAUSE__CACHE_ */ 2, /* OMP_CLAUSE_GANG */ 1, /* OMP_CLAUSE_ASYNC */ @@ -384,6 +385,7 @@ const char * const omp_clause_code_name[] = "to", "map", "has_device_addr", + "_mapper_binding_", "_cache_", "gang", "async", diff --git a/gcc/tree.def b/gcc/tree.def index f015021e9dc..d093927a0b5 100644 --- a/gcc/tree.def +++ b/gcc/tree.def @@ -1247,6 +1247,13 @@ DEFTREECODE (OMP_SECTION, "omp_section", tcc_statement, 1) Operand 0: OMP_MASTER_BODY: Master section body. */ DEFTREECODE (OMP_MASTER, "omp_master", tcc_statement, 1) +/* OpenMP - #pragma omp declare mapper ([id:] type var) [clause1 ... clauseN] + Operand 0: Identifier. + Operand 1: Type. + Operand 2: Variable decl. + Operand 3: List of clauses. */ +DEFTREECODE (OMP_DECLARE_MAPPER, "omp_declare_mapper", tcc_statement, 4) + /* OpenACC - #pragma acc cache (variable1 ... variableN) Operand 0: OACC_CACHE_CLAUSES: List of variables (transformed into OMP_CLAUSE__CACHE_ clauses). */ diff --git a/gcc/tree.h b/gcc/tree.h index 31c7a129f3c..d580e9bcdc1 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1532,6 +1532,15 @@ class auto_suppress_location_wrappers #define OMP_TARGET_EXIT_DATA_CLAUSES(NODE)\ TREE_OPERAND (OMP_TARGET_EXIT_DATA_CHECK (NODE), 0) +#define OMP_DECLARE_MAPPER_ID(NODE) \ + TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 0) +#define OMP_DECLARE_MAPPER_TYPE(NODE) \ + TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 1) +#define OMP_DECLARE_MAPPER_DECL(NODE) \ + TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 2) +#define OMP_DECLARE_MAPPER_CLAUSES(NODE) \ + TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 3) + #define OMP_SCAN_BODY(NODE) TREE_OPERAND (OMP_SCAN_CHECK (NODE), 0) #define OMP_SCAN_CLAUSES(NODE) TREE_OPERAND (OMP_SCAN_CHECK (NODE), 1) @@ -1960,6 +1969,18 @@ class auto_suppress_location_wrappers #define OMP_CLAUSE__SCANTEMP__CONTROL(NODE) \ TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__SCANTEMP_)) +#define OMP_CLAUSE__MAPPER_BINDING__ID(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \ + OMP_CLAUSE__MAPPER_BINDING_), 0) + +#define OMP_CLAUSE__MAPPER_BINDING__DECL(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \ + OMP_CLAUSE__MAPPER_BINDING_), 1) + +#define OMP_CLAUSE__MAPPER_BINDING__MAPPER(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \ + OMP_CLAUSE__MAPPER_BINDING_), 2) + /* SSA_NAME accessors. */ /* Whether SSA_NAME NODE is a virtual operand. This simply caches the diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 84316f953d0..df706830551 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -182,7 +182,13 @@ enum gomp_map_kind /* An attach or detach operation. Rewritten to the appropriate type during gimplification, depending on directive (i.e. "enter data" or parallel/kernels region vs. "exit data"). */ - GOMP_MAP_ATTACH_DETACH = (GOMP_MAP_LAST | 3) + GOMP_MAP_ATTACH_DETACH = (GOMP_MAP_LAST | 3), + /* Unset, used for "declare mapper" maps with no explicit data movement + specified. These use the movement specified at the invocation site. */ + GOMP_MAP_UNSET = (GOMP_MAP_LAST | 4), + /* Used to record the name of a named mapper. */ + GOMP_MAP_PUSH_MAPPER_NAME = (GOMP_MAP_LAST | 5), + GOMP_MAP_POP_MAPPER_NAME = (GOMP_MAP_LAST | 6) }; #define GOMP_MAP_COPY_TO_P(X) \ diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-1.C b/libgomp/testsuite/libgomp.c++/declare-mapper-1.C new file mode 100644 index 00000000000..aba4f426539 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-1.C @@ -0,0 +1,87 @@ +// { dg-do run } + +#include +#include + +#define N 64 + +struct points +{ + double *x; + double *y; + double *z; + size_t len; +}; + +#pragma omp declare mapper(points p) map(to:p.x, p.y, p.z) \ + map(p.x[0:p.len]) \ + map(p.y[0:p.len]) \ + map(p.z[0:p.len]) + +struct shape +{ + points tmp; + points *pts; + int metadata[128]; +}; + +#pragma omp declare mapper(shape s) map(tofrom:s.pts, *s.pts) map(alloc:s.tmp) + +void +alloc_points (points *pts, size_t sz) +{ + pts->x = new double[sz]; + pts->y = new double[sz]; + pts->z = new double[sz]; + pts->len = sz; + for (int i = 0; i < sz; i++) + pts->x[i] = pts->y[i] = pts->z[i] = 0; +} + +int main (int argc, char *argv[]) +{ + shape myshape; + points mypts; + + myshape.pts = &mypts; + + alloc_points (&myshape.tmp, N); + myshape.pts = new points; + alloc_points (myshape.pts, N); + + #pragma omp target map(myshape) + { + for (int i = 0; i < N; i++) + { + myshape.pts->x[i]++; + myshape.pts->y[i]++; + myshape.pts->z[i]++; + } + } + + for (int i = 0; i < N; i++) + { + assert (myshape.pts->x[i] == 1); + assert (myshape.pts->y[i] == 1); + assert (myshape.pts->z[i] == 1); + } + + #pragma omp target + { + for (int i = 0; i < N; i++) + { + myshape.pts->x[i]++; + myshape.pts->y[i]++; + myshape.pts->z[i]++; + } + } + + for (int i = 0; i < N; i++) + { + assert (myshape.pts->x[i] == 2); + assert (myshape.pts->y[i] == 2); + assert (myshape.pts->z[i] == 2); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-2.C b/libgomp/testsuite/libgomp.c++/declare-mapper-2.C new file mode 100644 index 00000000000..d848fdb7369 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-2.C @@ -0,0 +1,55 @@ +// { dg-do run } + +#include + +#define N 256 + +struct doublebuf +{ + int buf_a[N][N]; + int buf_b[N][N]; +}; + +#pragma omp declare mapper(lo:doublebuf b) map(b.buf_a[0:N/2][0:N]) \ + map(b.buf_b[0:N/2][0:N]) + +#pragma omp declare mapper(hi:doublebuf b) map(b.buf_a[N/2:N/2][0:N]) \ + map(b.buf_b[N/2:N/2][0:N]) + +int main (int argc, char *argv[]) +{ + doublebuf db; + + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + db.buf_a[i][j] = db.buf_b[i][j] = 0; + + #pragma omp target map(mapper(lo), tofrom:db) + { + for (int i = 0; i < N / 2; i++) + for (int j = 0; j < N; j++) + { + db.buf_a[i][j]++; + db.buf_b[i][j]++; + } + } + + #pragma omp target map(mapper(hi), tofrom:db) + { + for (int i = N / 2; i < N; i++) + for (int j = 0; j < N; j++) + { + db.buf_a[i][j]++; + db.buf_b[i][j]++; + } + } + + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + { + assert (db.buf_a[i][j] == 1); + assert (db.buf_b[i][j] == 1); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-3.C b/libgomp/testsuite/libgomp.c++/declare-mapper-3.C new file mode 100644 index 00000000000..ea9b7ded75b --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-3.C @@ -0,0 +1,63 @@ +// { dg-do run } + +#include +#include + +struct S { + int *myarr; +}; + +#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:20]) + +namespace A { +#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:100]) +} + +namespace B { +#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[100:100]) +} + +namespace A +{ + void incr_a (S my_s) + { +#pragma omp target + { + for (int i = 0; i < 100; i++) + my_s.myarr[i]++; + } + } +} + +namespace B +{ + void incr_b (S my_s) + { +#pragma omp target + { + for (int i = 100; i < 200; i++) + my_s.myarr[i]++; + } + } +} + +int main (int argc, char *argv[]) +{ + S my_s; + + my_s.myarr = (int *) calloc (200, sizeof (int)); + +#pragma omp target + { + for (int i = 0; i < 20; i++) + my_s.myarr[i]++; + } + + A::incr_a (my_s); + B::incr_b (my_s); + + for (int i = 0; i < 200; i++) + assert (my_s.myarr[i] == (i < 20) ? 2 : 1); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-4.C b/libgomp/testsuite/libgomp.c++/declare-mapper-4.C new file mode 100644 index 00000000000..f194e63b5b7 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-4.C @@ -0,0 +1,63 @@ +// { dg-do run } + +#include +#include + +struct S { + int *myarr; +}; + +#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:20]) + +namespace A { +#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:100]) +} + +namespace B { +#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[100:100]) +} + +namespace A +{ + void incr_a (S &my_s) + { +#pragma omp target + { + for (int i = 0; i < 100; i++) + my_s.myarr[i]++; + } + } +} + +namespace B +{ + void incr_b (S &my_s) + { +#pragma omp target + { + for (int i = 100; i < 200; i++) + my_s.myarr[i]++; + } + } +} + +int main (int argc, char *argv[]) +{ + S my_s; + + my_s.myarr = (int *) calloc (200, sizeof (int)); + +#pragma omp target + { + for (int i = 0; i < 20; i++) + my_s.myarr[i]++; + } + + A::incr_a (my_s); + B::incr_b (my_s); + + for (int i = 0; i < 200; i++) + assert (my_s.myarr[i] == (i < 20) ? 2 : 1); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-5.C b/libgomp/testsuite/libgomp.c++/declare-mapper-5.C new file mode 100644 index 00000000000..0030de8791a --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-5.C @@ -0,0 +1,52 @@ +// { dg-do run } + +#include + +struct S +{ + int *myarr; + int len; +}; + +class C +{ + S smemb; +#pragma omp declare mapper (custom:S s) map(to:s.myarr) \ + map(tofrom:s.myarr[0:s.len]) + +public: + C(int l) + { + smemb.myarr = new int[l]; + smemb.len = l; + for (int i = 0; i < l; i++) + smemb.myarr[i] = 0; + } + void bump(); + void check(); +}; + +void +C::bump () +{ +#pragma omp target map(mapper(custom), tofrom: smemb) + { + for (int i = 0; i < smemb.len; i++) + smemb.myarr[i]++; + } +} + +void +C::check () +{ + for (int i = 0; i < smemb.len; i++) + assert (smemb.myarr[i] == 1); +} + +int main (int argc, char *argv[]) +{ + C test (100); + test.bump (); + test.check (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-6.C b/libgomp/testsuite/libgomp.c++/declare-mapper-6.C new file mode 100644 index 00000000000..14ed10df702 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-6.C @@ -0,0 +1,37 @@ +// { dg-do run } + +#include + +template +void adjust (T param) +{ +#pragma omp declare mapper (T x) map(to:x.len, x.base) \ + map(tofrom:x.base[0:x.len]) + +#pragma omp target + for (int i = 0; i < param.len; i++) + param.base[i]++; +} + +struct S { + int len; + int *base; +}; + +int main (int argc, char *argv[]) +{ + S a; + + a.len = 100; + a.base = new int[a.len]; + + for (int i = 0; i < a.len; i++) + a.base[i] = 0; + + adjust (a); + + for (int i = 0; i < a.len; i++) + assert (a.base[i] == 1); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-7.C b/libgomp/testsuite/libgomp.c++/declare-mapper-7.C new file mode 100644 index 00000000000..ab632099714 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-7.C @@ -0,0 +1,48 @@ +// { dg-do run } + +#include + +struct S +{ + int *myarr; +}; + +struct T +{ + S *s; +}; + +#pragma omp declare mapper (s100: S x) map(to: x.myarr) \ + map(tofrom: x.myarr[0:100]) + +void +bump (T t) +{ + /* Here we have an implicit/default mapper invoking a named mapper. We + need to make sure that can be located properly at gimplification + time. */ +#pragma omp declare mapper (T t) map(to:t.s) map(mapper(s100), tofrom: t.s[0]) + +#pragma omp target + for (int i = 0; i < 100; i++) + t.s->myarr[i]++; +} + +int main (int argc, char *argv[]) +{ + S my_s; + T my_t; + + my_s.myarr = new int[100]; + my_t.s = &my_s; + + for (int i = 0; i < 100; i++) + my_s.myarr[i] = 0; + + bump (my_t); + + for (int i = 0; i < 100; i++) + assert (my_s.myarr[i] == 1); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-8.C b/libgomp/testsuite/libgomp.c++/declare-mapper-8.C new file mode 100644 index 00000000000..3818e5264d3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/declare-mapper-8.C @@ -0,0 +1,61 @@ +// { dg-do run } + +#include + +struct S +{ + int *myarr; + int len; +}; + +template +class C +{ + T memb; +#pragma omp declare mapper (T t) map(to:t.len, t.myarr) \ + map(tofrom:t.myarr[0:t.len]) + +public: + C(int sz); + ~C(); + void bump(); + void check(); +}; + +template +C::C(int sz) +{ + memb.myarr = new int[sz]; + for (int i = 0; i < sz; i++) + memb.myarr[i] = 0; + memb.len = sz; +} + +template +C::~C() +{ + delete[] memb.myarr; +} + +template +void C::bump() +{ +#pragma omp target map(memb) + for (int i = 0; i < memb.len; i++) + memb.myarr[i]++; +} + +template +void C::check() +{ + for (int i = 0; i < memb.len; i++) + assert (memb.myarr[i] == 1); +} + +int main(int argc, char *argv[]) +{ + C c_int(100); + c_int.bump(); + c_int.check(); + return 0; +}