From patchwork Tue Sep 1 13:16:23 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 1354991 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4BgnfD4J9fz9sTS for ; Tue, 1 Sep 2020 23:16:50 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 018FE38708B1; Tue, 1 Sep 2020 13:16:45 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 2DEA3386F01D for ; Tue, 1 Sep 2020 13:16:41 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 2DEA3386F01D Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=ChungLin_Tang@mentor.com IronPort-SDR: lB+xg4aFPbZcojOY6KYB8hClAiuT3RRSbBWvUk7l+fW1wINMlLVSn6EKmiGuEvof2HgVh2xsqq XMYuzcRkQ2hmS9fw8x2Ux47g9OoBCg2lvDDHP/Pc2Pqy2tpW16ehVAmNrKj2ls4P1fmC9Zkbgj PemUkj2LHEoDBz1uRQPT936KEp0PUvXJcE7Ie6ll8tODPFcMwj+1lbXqs3mN5N4aO3ljkGXPhq o/MvI3/ufKHhlD9Xjv0QBgPavi2+XW239w3W7nkN+wYjzvDeqkMjBW4GI12la6IKuN7LXX4Hl0 Fts= X-IronPort-AV: E=Sophos;i="5.76,379,1592899200"; d="scan'208";a="52442519" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 01 Sep 2020 05:16:39 -0800 IronPort-SDR: 4aBeTgzjwX2A0TTwxC4Uz+Qm+F0KPOVtMvLPEzvKN+Ay0PN0wddEYeP4MGFr+JOjzHJyD6IpZu 7jkLRuZsL18swzGK7nLqyjlaZ8sMKWNRYtAJX/85nTR2ddldyj3p4xFod1Z7saLb6izb57YSDC jXcBpWRfGGdGgIEgG+mJg5jAh3tYzC23eea9/l1hlzNpem/gYqWnfRuL5ed3c19vt6lkE/rzOG FpzmhIbkpeRLbBBLdCJHnR3dQswLBh8TS9mtc/7LkZq7RIuX6B2UE8OiX30Gnq8gofgor9wuoA V4o= From: Chung-Lin Tang Subject: [PATCH, 1/3, OpenMP] Target mapping changes for OpenMP 5.0, front-end parts To: gcc-patches , Jakub Jelinek , Tobias Burnus , Catherine Moore , Thomas Schwinge Message-ID: <639a56ef-eeed-eb38-8a19-f5cf8d082973@codesourcery.com> Date: Tue, 1 Sep 2020 21:16:23 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:68.0) Gecko/20100101 Thunderbird/68.12.0 MIME-Version: 1.0 Content-Language: en-US X-ClientProxiedBy: SVR-ORW-MBX-05.mgc.mentorg.com (147.34.90.205) To svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) X-Spam-Status: No, score=-9.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: , Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" Hi Jakub, this patch set implements parts of the target mapping changes introduced in OpenMP 5.0, mainly the attachment requirements for pointer-based list items, and the clause ordering. The first patch here are the C/C++ front-end changes. The entire set of changes has been tested for without regressions for the compiler and libgomp. Hope this is ready to commit to master. Thanks, Chung-Lin gcc/c-family/ * c-common.h (c_omp_adjust_clauses): New declaration. * c-omp.c (c_omp_adjust_clauses): New function. gcc/c/ * c-parser.c (c_parser_omp_target_data): Add use of new c_omp_adjust_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (c_parser_omp_target_enter_data): Likewise. (c_parser_omp_target_exit_data): Likewise. (c_parser_omp_target): Likewise. * c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. (c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/cp/ * parser.c (cp_parser_omp_target_data): Add use of new c_omp_adjust_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (cp_parser_omp_target_enter_data): Likewise. (cp_parser_omp_target_exit_data): Likewise. (cp_parser_omp_target): Likewise. * semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix interaction between reference case and attach/detach. (finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index 4fc64bc4aa6..9ef85b401f0 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1208,14 +1208,15 @@ extern tree c_omp_declare_simd_clauses_to_numbers (tree, tree); extern void c_omp_declare_simd_clauses_to_decls (tree, tree); extern bool c_omp_predefined_variable (tree); extern enum omp_clause_default_kind c_omp_predetermined_sharing (tree); extern enum omp_clause_defaultmap_kind c_omp_predetermined_mapping (tree); extern tree c_omp_check_context_selector (location_t, tree); extern void c_omp_mark_declare_variant (location_t, tree, tree); extern const char *c_omp_map_clause_name (tree, bool); +extern void c_omp_adjust_clauses (tree, bool); /* Return next tree in the chain for chain_next walking of tree nodes. */ static inline tree c_tree_chain_next (tree t) { /* TREE_CHAIN of a type is TYPE_STUB_DECL, which is different kind of object, never a long chain of nodes. Prefer diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index d7cff0f4cca..596f33cebfb 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -2575,7 +2575,51 @@ c_omp_map_clause_name (tree clause, bool oacc) case GOMP_MAP_DEVICE_RESIDENT: return "device_resident"; case GOMP_MAP_LINK: return "link"; case GOMP_MAP_FORCE_DEVICEPTR: return "deviceptr"; default: break; } return omp_clause_code_name[OMP_CLAUSE_CODE (clause)]; } + +/* Adjust map clauses after normal clause parsing, mainly to turn specific + base-pointer map cases into attach/detach and mark them addressable. */ +void +c_omp_adjust_clauses (tree clauses, bool is_target) +{ + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) != ARRAY_TYPE) + { + tree ptr = OMP_CLAUSE_DECL (c); + bool ptr_mapped = false; + if (is_target) + { + for (tree m = clauses; m; m = OMP_CLAUSE_CHAIN (m)) + if (OMP_CLAUSE_CODE (m) == OMP_CLAUSE_MAP + && OMP_CLAUSE_DECL (m) == ptr + && (OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALLOC + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TO + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_FROM + || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TOFROM)) + { + ptr_mapped = true; + break; + } + + if (!ptr_mapped + && DECL_P (ptr) + && is_global_var (ptr) + && lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (ptr))) + ptr_mapped = true; + } + + /* If the pointer variable was mapped, or if this is not an offloaded + target region, adjust the map kind to attach/detach. */ + if (ptr_mapped || !is_target) + { + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH_DETACH); + c_common_mark_addressable_vec (ptr); + } + } +} diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index a8bc301ffad..92dfe3b6a4a 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -19452,14 +19452,15 @@ c_parser_omp_teams (location_t loc, c_parser *parser, static tree c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) { tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, "#pragma omp target data"); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_TO: @@ -19469,14 +19470,15 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) case GOMP_MAP_TOFROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target data%> with map-type other " "than %, %, % or % " "on % clause"); @@ -19592,27 +19594,29 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser, c_parser_skip_to_pragma_eol (parser, false); return NULL_TREE; } tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK, "#pragma omp target enter data"); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_TO: case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALLOC: map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target enter data%> with map-type other " "than % or % on % clause"); *pc = OMP_CLAUSE_CHAIN (*pc); @@ -19676,29 +19680,30 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser, c_parser_skip_to_pragma_eol (parser, false); return NULL_TREE; } tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK, "#pragma omp target exit data"); - + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_FROM: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target exit data%> with map-type other " "than %, % or % on %" " clause"); @@ -19900,14 +19905,16 @@ 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"); + c_omp_adjust_clauses (OMP_TARGET_CLAUSES (stmt), 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)); OMP_TARGET_BODY (stmt) = c_end_compound_stmt (loc, block, true); SET_EXPR_LOCATION (stmt, loc); @@ -19924,14 +19931,15 @@ check_clauses: case GOMP_MAP_FROM: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_TOFROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target%> with map-type other " "than %, %, % or % " "on % clause"); *pc = OMP_CLAUSE_CHAIN (*pc); diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 0d639b60ea3..17ac2f566da 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13580,16 +13580,17 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) break; } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); if (ort != C_ORT_OMP && ort != C_ORT_ACC) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); else if (TREE_CODE (t) == COMPONENT_REF) { - gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; + gomp_map_kind k + = ((ort == C_ORT_ACC || ort == C_ORT_OMP) + ? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER); OMP_CLAUSE_SET_MAP_KIND (c2, k); } else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER && !c_mark_addressable (t)) return false; @@ -14682,15 +14683,16 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); } } if (remove) break; if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) { - if (bitmap_bit_p (&map_field_head, DECL_UID (t))) + if (bitmap_bit_p (&map_field_head, DECL_UID (t)) + || bitmap_bit_p (&map_head, DECL_UID (t))) break; } } if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) { error_at (OMP_CLAUSE_LOCATION (c), "%qE is not a variable in %qs clause", t, @@ -14751,29 +14753,36 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears both in data and map clauses", t); remove = true; } else bitmap_set_bit (&generic_head, DECL_UID (t)); } - else if (bitmap_bit_p (&map_head, DECL_UID (t))) + else if (bitmap_bit_p (&map_head, DECL_UID (t)) + && !bitmap_bit_p (&map_field_head, DECL_UID (t))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in motion clauses", t); else if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in map clauses", t); remove = true; } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + && ort == C_ORT_ACC) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD appears more than once in data clauses", t); + remove = true; + } + else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else error_at (OMP_CLAUSE_LOCATION (c), "%qD appears both in data and map clauses", t); diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 7cc2dbed5fe..7773f9d4f79 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -40449,14 +40449,15 @@ cp_parser_omp_teams (cp_parser *parser, cp_token *pragma_tok, static tree cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) { tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, "#pragma omp target data", pragma_tok); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_TO: @@ -40467,14 +40468,15 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target data%> with map-type other " "than %, %, % or % " "on % clause"); @@ -40550,28 +40552,30 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok, cp_parser_skip_to_pragma_eol (parser, pragma_tok); return NULL_TREE; } tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK, "#pragma omp target enter data", pragma_tok); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_TO: case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALLOC: map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target enter data%> with map-type other " "than % or % on % clause"); *pc = OMP_CLAUSE_CHAIN (*pc); @@ -40638,14 +40642,15 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok, cp_parser_skip_to_pragma_eol (parser, pragma_tok); return NULL_TREE; } tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK, "#pragma omp target exit data", pragma_tok); + c_omp_adjust_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_FROM: @@ -40653,14 +40658,15 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok, case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target exit data%> with map-type other " "than %, % or % on %" " clause"); @@ -40901,14 +40907,16 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, stmt = make_node (OMP_TARGET); TREE_TYPE (stmt) = void_type_node; OMP_TARGET_CLAUSES (stmt) = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, "#pragma omp target", pragma_tok); + c_omp_adjust_clauses (OMP_TARGET_CLAUSES (stmt), true); + pc = &OMP_TARGET_CLAUSES (stmt); keep_next_level (true); OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p); SET_EXPR_LOCATION (stmt, pragma_tok->location); add_stmt (stmt); @@ -40924,14 +40932,15 @@ check_clauses: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_TOFROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target%> with map-type other " "than %, %, % or % " "on % clause"); *pc = OMP_CLAUSE_CHAIN (*pc); diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index b71ca0729a8..0f6b36f2dab 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5373,16 +5373,17 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); else if (TREE_CODE (t) == COMPONENT_REF) { - gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; + gomp_map_kind k + = ((ort == C_ORT_ACC || ort == C_ORT_OMP) + ? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER); OMP_CLAUSE_SET_MAP_KIND (c2, k); } else if (REFERENCE_REF_P (t) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) { t = TREE_OPERAND (t, 0); gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH @@ -5414,16 +5415,20 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) && TYPE_REF_P (TREE_TYPE (ptr)) && INDIRECT_TYPE_P (TREE_TYPE (TREE_TYPE (ptr)))) { tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2)); OMP_CLAUSE_DECL (c3) = ptr; - if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER) - OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr); + if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER + || OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ATTACH_DETACH) + { + OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); + } else OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr); OMP_CLAUSE_SIZE (c3) = size_zero_node; OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (c2); OMP_CLAUSE_CHAIN (c2) = c3; } } @@ -7400,15 +7405,15 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) OMP_CLAUSE_SIZE (c) = size_zero_node; if (REFERENCE_REF_P (t) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) { t = TREE_OPERAND (t, 0); OMP_CLAUSE_DECL (c) = t; } - if (ort == C_ORT_ACC + if ((ort == C_ORT_ACC || ort == C_ORT_OMP) && TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF) t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); if (TREE_CODE (t) == COMPONENT_REF && ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP || ort == C_ORT_ACC) && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) @@ -7446,15 +7451,16 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } if (remove) break; if (REFERENCE_REF_P (t)) t = TREE_OPERAND (t, 0); if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) { - if (bitmap_bit_p (&map_field_head, DECL_UID (t))) + if (bitmap_bit_p (&map_field_head, DECL_UID (t)) + || bitmap_bit_p (&map_head, DECL_UID (t))) goto handle_map_references; } } if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) { if (processing_template_decl && TREE_CODE (t) != OVERLOAD) break; @@ -7540,30 +7546,35 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) "%qD appears both in data and map clauses", t); remove = true; } else bitmap_set_bit (&generic_head, DECL_UID (t)); } else if (bitmap_bit_p (&map_head, DECL_UID (t)) - && (ort != C_ORT_ACC - || !bitmap_bit_p (&map_field_head, DECL_UID (t)))) + && !bitmap_bit_p (&map_field_head, DECL_UID (t))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in motion clauses", t); - if (ort == C_ORT_ACC) + else if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in map clauses", t); remove = true; } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + && ort == C_ORT_ACC) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD appears more than once in data clauses", t); + remove = true; + } + else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { if (ort == C_ORT_ACC) error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); else error_at (OMP_CLAUSE_LOCATION (c), "%qD appears both in data and map clauses", t); @@ -7591,23 +7602,25 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (t))); } else if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_REFERENCE) && (OMP_CLAUSE_MAP_KIND (c) - != GOMP_MAP_ALWAYS_POINTER)) + != GOMP_MAP_ALWAYS_POINTER) + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_ATTACH_DETACH)) { tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); if (TREE_CODE (t) == COMPONENT_REF) { gomp_map_kind k - = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH - : GOMP_MAP_ALWAYS_POINTER; + = ((ort == C_ORT_ACC || ort == C_ORT_OMP) + ? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER); OMP_CLAUSE_SET_MAP_KIND (c2, k); } else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_REFERENCE); OMP_CLAUSE_DECL (c2) = t; OMP_CLAUSE_SIZE (c2) = size_zero_node; From patchwork Tue Sep 1 13:16:48 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 1354992 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4BgnfV5YCgz9sTC for ; Tue, 1 Sep 2020 23:17:06 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id D9360386F01D; Tue, 1 Sep 2020 13:17:03 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id BF324384607A for ; Tue, 1 Sep 2020 13:16:57 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org BF324384607A Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=ChungLin_Tang@mentor.com IronPort-SDR: FOdP5zt8lNfDJ/ph4B0EtStGF3kqwXYX7+/2soE/ou1r3j7HnUvo52yieL4+vyvZmwE1RH9pF+ yT6hfRCbn+cQ3+OQIN0MsZAhVt+I5rbyX6n4WhLPPLgZDWnE7wZ9q3f75Z/QZsP4kyR5dXPGML wosjS7K7CVQHc2lNPeLNiUkzUZWU/7/hzia9PznyUnfuGQnKlnPBOcA6xaqxMgg5ELLHZdwck/ +i3qP+NizwRPwuCki2lI1UmvPRWP44gzBmRVSs9617LXSKhgfFoOVKbIIfSp2uIWJ6jkisZFDl NkA= X-IronPort-AV: E=Sophos;i="5.76,379,1592899200"; d="scan'208";a="54638912" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 01 Sep 2020 05:16:56 -0800 IronPort-SDR: 4xL+S3AvyonBa3D0emhGyCztG9oYHghaGR15ALy/9p0FnqvI8cT+OpCtB3cbkrk4iPxEbS0ft7 MnEmPC4yceb7YcdJDPp0/3gYfE6SfsPc/94qwIDcCYl/XQT2dMnOskpDZdvppWYgE3Av37ri2A TM3vlAR0eWm3175g7Vu6bStd8IleE4wgXIqiEyO4VZ1ROwyA1LUm9eUuB/cKIQ+9od1fwoDPZ1 tUYUkm3v930sY0UBJ0k2yk9/1oHoEaAtIoViiTif/ld29d2PYiRl3+GaBcNy4S1ze0L0XATxOJ Meg= From: Chung-Lin Tang Subject: [PATCH, 2/3, OpenMP] Target mapping changes for OpenMP 5.0, middle-end parts and compiler testcases To: gcc-patches , Jakub Jelinek , Tobias Burnus , Catherine Moore , Thomas Schwinge Message-ID: Date: Tue, 1 Sep 2020 21:16:48 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:68.0) Gecko/20100101 Thunderbird/68.12.0 MIME-Version: 1.0 Content-Language: en-US X-ClientProxiedBy: SVR-ORW-MBX-06.mgc.mentorg.com (147.34.90.206) To svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) X-Spam-Status: No, score=-9.2 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: , Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" These are the middle-end gimplifier and omp-low changes. Compiler testcases are also included in this patch. For attach/detach clauses, I'm currently using the clause tree expression itself as the key for lookup, to solve the "same-decl" problem when multiple clauses have the same OMP_CLAUSE_DECL. This is just a special case right now, have yet to see if this can expand to more general use between all map clauses. Thanks, Chung-Lin gcc/ * gimplify.c (is_or_contains_p): New static helper function. (omp_target_reorder_clauses): New function. (gimplify_scan_omp_clauses): Add use of omp_target_reorder_clauses to reorder clause list according to OpenMP 5.0 rules. Add handling of GOMP_MAP_ATTACH_DETACH for OpenMP cases. * omp-low.c (is_omp_target): New static helper function. (scan_sharing_clauses): Add scan phase handling of GOMP_MAP_ATTACH/DETACH for OpenMP cases. (lower_omp_target): Add lowering handling of GOMP_MAP_ATTACH/DETACH for OpenMP cases. gcc/testsuite/ * c-c++-common/goacc/finalize-1.c: Adjust gimple scanning. * c-c++-common/goacc/mdc-1.c: Likewise. * c-c++-common/goacc/struct-enter-exit-data-1.c: Likewise. * gfortran.dg/goacc/attach-descriptor.f90: Likewise. * gfortran.dg/goacc/finalize-1.f: Likewise. * c-c++-common/gomp/clauses-2.c: Remove dg-error cases now valid. * gfortran.dg/gomp/map-2.f90: Likewise. * c-c++-common/gomp/map-5.c: New testcase. diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 23d0e2511f7..0ad141c5b3f 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8350,14 +8350,126 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp, /* Set *BASE_REF if BASE was a dereferenced reference variable. */ if (base_ref && orig_base != base) *base_ref = orig_base; return base; } +/* Returns true if EXPR is or contains (as a sub-component) BASE_PTR. */ + +static bool +is_or_contains_p (tree expr, tree base_ptr) +{ + while (expr != base_ptr) + if (TREE_CODE (base_ptr) == COMPONENT_REF) + base_ptr = TREE_OPERAND (base_ptr, 0); + else + break; + return expr == base_ptr; +} + +/* Implement OpenMP 5.x map ordering rules for target directives. There are + several rules, and with some level of ambiguity, hopefully we can at least + collect the complexity here in one place. */ + +static void +omp_target_reorder_clauses (tree *list_p) +{ + vec clauses = vNULL; + for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp)) + clauses.safe_push (*cp); + + /* Collect refs to alloc/release/delete maps. */ + vec ard = vNULL; + for (unsigned int i = 0; i < clauses.length (); i++) + if (OMP_CLAUSE_CODE (clauses[i]) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (clauses[i]) == GOMP_MAP_ALLOC + || OMP_CLAUSE_MAP_KIND (clauses[i]) == GOMP_MAP_RELEASE + || OMP_CLAUSE_MAP_KIND (clauses[i]) == GOMP_MAP_DELETE)) + { + ard.safe_push (clauses[i]); + clauses[i] = NULL_TREE; + + unsigned int j; + for (j = i + 1; j < clauses.length (); j++) + { + /* Any associated pointer type maps should move along. */ + tree nc = clauses[j]; + if (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_POINTER + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ATTACH_DETACH + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_POINTER + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ALWAYS_POINTER + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_TO_PSET)) + { + ard.safe_push (nc); + clauses[j] = NULL_TREE; + } + else + break; + } + i = j - 1; + } + + tree *cp = list_p; + for (unsigned int i = 0; i < clauses.length (); i++) + if (clauses[i]) + { + *cp = clauses[i]; + cp = &OMP_CLAUSE_CHAIN (clauses[i]); + } + for (unsigned int i = 0; i < ard.length (); i++) + { + *cp = ard[i]; + cp = &OMP_CLAUSE_CHAIN (ard[i]); + } + *cp = NULL_TREE; + + /* OpenMP 5.0 requires that pointer variables are mapped before + its use as a base-pointer. */ + for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp)) + if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP) + { + tree decl = OMP_CLAUSE_DECL (*cp); + gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp); + if ((k == GOMP_MAP_ALLOC + || k == GOMP_MAP_TO + || k == GOMP_MAP_FROM + || k == GOMP_MAP_TOFROM) + && (TREE_CODE (decl) == INDIRECT_REF + || TREE_CODE (decl) == MEM_REF)) + { + tree base_ptr = TREE_OPERAND (decl, 0); + STRIP_TYPE_NOPS (base_ptr); + for (tree *cp2 = &OMP_CLAUSE_CHAIN (*cp); *cp2; + cp2 = &OMP_CLAUSE_CHAIN (*cp2)) + if (OMP_CLAUSE_CODE (*cp2) == OMP_CLAUSE_MAP) + { + tree decl2 = OMP_CLAUSE_DECL (*cp2); + gomp_map_kind k2 = OMP_CLAUSE_MAP_KIND (*cp2); + if ((k2 == GOMP_MAP_ALLOC + || k2 == GOMP_MAP_TO + || k2 == GOMP_MAP_FROM + || k2 == GOMP_MAP_TOFROM) + && is_or_contains_p (decl2, base_ptr)) + { + /* Move *cp2 to before *cp. */ + tree c = *cp2; + *cp2 = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = *cp; + *cp = c; + if (*cp2 == NULL_TREE) + break; + } + } + } + } +} + /* 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) @@ -8391,14 +8503,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OACC_PARALLEL: case OACC_KERNELS: ctx->target_firstprivatize_array_bases = true; default: break; } + if (code == OMP_TARGET + || code == OMP_TARGET_DATA + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA) + omp_target_reorder_clauses (list_p); + while ((c = *list_p) != NULL) { bool remove = false; bool notice_outer = true; const char *check_non_private = NULL; unsigned int flags; tree decl; @@ -8831,23 +8949,26 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) { remove = true; break; } else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) + == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST) { OMP_CLAUSE_SIZE (c) = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL, false); - omp_add_variable (ctx, OMP_CLAUSE_SIZE (c), - GOVD_FIRSTPRIVATE | GOVD_SEEN); + if ((region_type & ORT_TARGET) != 0) + omp_add_variable (ctx, OMP_CLAUSE_SIZE (c), + GOVD_FIRSTPRIVATE | GOVD_SEEN); } + if (!DECL_P (decl)) { tree d = decl, *pd; if (TREE_CODE (d) == ARRAY_REF) { while (TREE_CODE (d) == ARRAY_REF) d = TREE_OPERAND (d, 0); @@ -8864,25 +8985,27 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, { pd = &TREE_OPERAND (decl, 0); decl = TREE_OPERAND (decl, 0); } bool indir_p = false; tree orig_decl = decl; tree decl_ref = NULL_TREE; - if ((region_type & ORT_ACC) != 0 + if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0 && TREE_CODE (*pd) == COMPONENT_REF && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH && code != OACC_UPDATE) { while (TREE_CODE (decl) == COMPONENT_REF) { decl = TREE_OPERAND (decl, 0); - if ((TREE_CODE (decl) == MEM_REF - && integer_zerop (TREE_OPERAND (decl, 1))) - || INDIRECT_REF_P (decl)) + if (((TREE_CODE (decl) == MEM_REF + && integer_zerop (TREE_OPERAND (decl, 1))) + || INDIRECT_REF_P (decl)) + && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) + == POINTER_TYPE)) { indir_p = true; decl = TREE_OPERAND (decl, 0); } if (TREE_CODE (decl) == INDIRECT_REF && DECL_P (TREE_OPERAND (decl, 0)) && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) @@ -8901,24 +9024,26 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, && DECL_P (TREE_OPERAND (decl, 0)) && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) == REFERENCE_TYPE)) decl = TREE_OPERAND (decl, 0); } if (decl != orig_decl && DECL_P (decl) && indir_p) { - gomp_map_kind k = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH - : GOMP_MAP_ATTACH; + gomp_map_kind k + = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA) + ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH); /* We have a dereference of a struct member. Make this an attach/detach operation, and ensure the base pointer is mapped as a FIRSTPRIVATE_POINTER. */ OMP_CLAUSE_SET_MAP_KIND (c, k); flags = GOVD_MAP | GOVD_SEEN | GOVD_EXPLICIT; tree next_clause = OMP_CLAUSE_CHAIN (c); if (k == GOMP_MAP_ATTACH && code != OACC_ENTER_DATA + && code != OMP_TARGET_ENTER_DATA && (!next_clause || (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP) || (OMP_CLAUSE_MAP_KIND (next_clause) != GOMP_MAP_POINTER) || OMP_CLAUSE_DECL (next_clause) != decl) && (!struct_deref_set || !struct_deref_set->contains (decl))) @@ -8958,25 +9083,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, /* An "attach/detach" operation on an update directive should behave as a GOMP_MAP_ALWAYS_POINTER. Beware that unlike attach or detach map kinds, GOMP_MAP_ALWAYS_POINTER depends on the previous mapping. */ if (code == OACC_UPDATE && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); - if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue) - == GS_ERROR) - { - remove = true; - break; - } if (DECL_P (decl) && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH - && code != OACC_UPDATE) + && code != OACC_UPDATE + && code != OMP_TARGET_UPDATE) { if (error_operand_p (decl)) { remove = true; break; } @@ -9030,23 +9150,27 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, bool attach_detach = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH); bool attach = OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH; bool has_attachments = false; /* For OpenACC, pointers in structs should trigger an attach action. */ - if (attach_detach && (region_type & ORT_ACC) != 0) + if (attach_detach + && ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA)) + { /* Turn a GOMP_MAP_ATTACH_DETACH clause into a GOMP_MAP_ATTACH or GOMP_MAP_DETACH clause after we have detected a case that needs a GOMP_MAP_STRUCT mapping added. */ gomp_map_kind k - = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH - : GOMP_MAP_ATTACH; + = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA) + ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH); OMP_CLAUSE_SET_MAP_KIND (c, k); has_attachments = true; } if (n == NULL || (n->value & GOVD_MAP) == 0) { tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); @@ -9134,41 +9258,49 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, tree base = extract_base_bit_offset (sc_decl, NULL, &bitposn, &offsetn); if (base != decl) break; if (scp) continue; - tree d1 = OMP_CLAUSE_DECL (*sc); - tree d2 = OMP_CLAUSE_DECL (c); - while (TREE_CODE (d1) == ARRAY_REF) - d1 = TREE_OPERAND (d1, 0); - while (TREE_CODE (d2) == ARRAY_REF) - d2 = TREE_OPERAND (d2, 0); - if (TREE_CODE (d1) == INDIRECT_REF) - d1 = TREE_OPERAND (d1, 0); - if (TREE_CODE (d2) == INDIRECT_REF) - d2 = TREE_OPERAND (d2, 0); - while (TREE_CODE (d1) == COMPONENT_REF) - if (TREE_CODE (d2) == COMPONENT_REF - && TREE_OPERAND (d1, 1) - == TREE_OPERAND (d2, 1)) - { + if (! (code == OMP_TARGET + || code == OMP_TARGET_DATA + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA)) + { + /* This duplicate checking code is currently only + enabled for OpenACC. */ + tree d1 = OMP_CLAUSE_DECL (*sc); + tree d2 = OMP_CLAUSE_DECL (c); + while (TREE_CODE (d1) == ARRAY_REF) d1 = TREE_OPERAND (d1, 0); + while (TREE_CODE (d2) == ARRAY_REF) d2 = TREE_OPERAND (d2, 0); - } - else - break; - if (d1 == d2) - { - error_at (OMP_CLAUSE_LOCATION (c), - "%qE appears more than once in map " - "clauses", OMP_CLAUSE_DECL (c)); - remove = true; - break; + if (TREE_CODE (d1) == INDIRECT_REF) + d1 = TREE_OPERAND (d1, 0); + if (TREE_CODE (d2) == INDIRECT_REF) + d2 = TREE_OPERAND (d2, 0); + while (TREE_CODE (d1) == COMPONENT_REF) + if (TREE_CODE (d2) == COMPONENT_REF + && TREE_OPERAND (d1, 1) + == TREE_OPERAND (d2, 1)) + { + d1 = TREE_OPERAND (d1, 0); + d2 = TREE_OPERAND (d2, 0); + } + else + break; + if (d1 == d2) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears more than once in map " + "clauses", OMP_CLAUSE_DECL (c)); + remove = true; + break; + } } if (maybe_lt (offset1, offsetn) || (known_eq (offset1, offsetn) && maybe_lt (bitpos1, bitposn))) { if (ptr || attach_detach) scp = sc; @@ -9222,18 +9354,68 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, == GOMP_MAP_ATTACH_DETACH) || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) == GOMP_MAP_TO_PSET))) prev_list_p = list_p; break; } + else + { + /* DECL_P (decl) == true */ + tree *sc; + if (struct_map_to_clause + && (sc = struct_map_to_clause->get (decl)) != NULL + && OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_STRUCT + && decl == OMP_CLAUSE_DECL (*sc)) + { + /* We have found a map of the whole structure after a + leading GOMP_MAP_STRUCT has been created, so refill the + leading clause into a map of the whole structure + variable, and remove the current one. + TODO: we should be able to remove some maps of the + following structure element maps if they are of + compatible TO/FROM/ALLOC type. */ + OMP_CLAUSE_SET_MAP_KIND (*sc, OMP_CLAUSE_MAP_KIND (c)); + OMP_CLAUSE_SIZE (*sc) = unshare_expr (OMP_CLAUSE_SIZE (c)); + remove = true; + break; + } + } flags = GOVD_MAP | GOVD_EXPLICIT; if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM) flags |= GOVD_MAP_ALWAYS_TO; + + if ((code == OMP_TARGET + || code == OMP_TARGET_DATA + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA) + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + { + for (struct gimplify_omp_ctx *octx = outer_ctx; octx; + octx = octx->outer_context) + { + splay_tree_node n + = splay_tree_lookup (octx->variables, + (splay_tree_key) OMP_CLAUSE_DECL (c)); + /* If this is contained in an outer OpenMP region as a + firstprivate value, remove the attach/detach. */ + if (n && (n->value & GOVD_FIRSTPRIVATE)) + { + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FIRSTPRIVATE_POINTER); + goto do_add; + } + } + + enum gomp_map_kind map_kind = (code == OMP_TARGET_EXIT_DATA + ? GOMP_MAP_DETACH + : GOMP_MAP_ATTACH); + OMP_CLAUSE_SET_MAP_KIND (c, map_kind); + } + goto do_add; case OMP_CLAUSE_DEPEND: if (OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SINK) { tree deps = OMP_CLAUSE_DECL (c); while (deps && TREE_CODE (deps) == TREE_LIST) diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 53efe5f750c..8d50774384a 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -210,14 +210,29 @@ is_oacc_kernels (omp_context *ctx) { enum gimple_code outer_type = gimple_code (ctx->stmt); return ((outer_type == GIMPLE_OMP_TARGET) && (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_OACC_KERNELS)); } +/* Return true if STMT corresponds to an OpenMP target region. */ +static bool +is_omp_target (gimple *stmt) +{ + if (gimple_code (stmt) == GIMPLE_OMP_TARGET) + { + int kind = gimple_omp_target_kind (stmt); + return (kind == GF_OMP_TARGET_KIND_REGION + || kind == GF_OMP_TARGET_KIND_DATA + || kind == GF_OMP_TARGET_KIND_ENTER_DATA + || kind == GF_OMP_TARGET_KIND_EXIT_DATA); + } + return false; +} + /* If DECL is the artificial dummy VAR_DECL created for non-static data member privatization, return the underlying "this" parameter, otherwise return NULL. */ tree omp_member_access_dummy_var (tree decl) { @@ -1342,15 +1357,17 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) don't need to be copied, the receiver side will use them directly. However, global variables with "omp declare target link" attribute need to be copied. Or when ALWAYS modifier is used. */ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && DECL_P (decl) && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER && (OMP_CLAUSE_MAP_KIND (c) - != GOMP_MAP_FIRSTPRIVATE_REFERENCE)) + != GOMP_MAP_FIRSTPRIVATE_REFERENCE) + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH) || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TO && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_FROM && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TOFROM && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) && varpool_node::get_create (decl)->offloadable && !lookup_attribute ("omp declare target link", @@ -1362,14 +1379,48 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) /* Ignore GOMP_MAP_POINTER kind for arrays in regions that are not offloaded; there is nothing to map for those. */ if (!is_gimple_omp_offloaded (ctx->stmt) && !POINTER_TYPE_P (TREE_TYPE (decl)) && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)) break; } + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && DECL_P (decl) + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + && is_omp_target (ctx->stmt)) + { + /* If this is an offloaded region, an attach operation should + only exist when the pointer variable is mapped in a prior + clause. */ + if (is_gimple_omp_offloaded (ctx->stmt)) + gcc_assert + (maybe_lookup_decl (decl, ctx) + || (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) + && lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (decl)))); + + /* By itself, attach/detach is generated as part of pointer + variable mapping and should not create new variables in the + offloaded region, however sender refs for it must be created + for its address to be passed to the runtime. */ + tree field + = build_decl (OMP_CLAUSE_LOCATION (c), + FIELD_DECL, NULL_TREE, ptr_type_node); + SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node)); + insert_field_into_struct (ctx->record_type, field); + /* To not clash with a map of the pointer variable itself, + attach/detach maps have their field looked up by the *clause* + tree expression, not the decl. */ + gcc_assert (!splay_tree_lookup (ctx->field_map, + (splay_tree_key) c)); + splay_tree_insert (ctx->field_map, (splay_tree_key) c, + (splay_tree_value) field); + break; + } if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE))) { if (TREE_CODE (decl) == COMPONENT_REF || (TREE_CODE (decl) == INDIRECT_REF @@ -1601,14 +1652,19 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_REFERENCE)) || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) && varpool_node::get_create (decl)->offloadable) break; + if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + && is_omp_target (ctx->stmt) + && !is_gimple_omp_offloaded (ctx->stmt)) + break; if (DECL_P (decl)) { if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE && !COMPLETE_TYPE_P (TREE_TYPE (decl))) { @@ -11405,26 +11461,26 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_STRUCT: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH: + case GOMP_MAP_DETACH: break; case GOMP_MAP_IF_PRESENT: case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_TO: case GOMP_MAP_FORCE_FROM: case GOMP_MAP_FORCE_TOFROM: case GOMP_MAP_FORCE_PRESENT: case GOMP_MAP_FORCE_DEVICEPTR: case GOMP_MAP_DEVICE_RESIDENT: case GOMP_MAP_LINK: - case GOMP_MAP_ATTACH: - case GOMP_MAP_DETACH: case GOMP_MAP_FORCE_DETACH: gcc_assert (is_gimple_omp_oacc (stmt)); break; default: gcc_unreachable (); } #endif @@ -11471,14 +11527,24 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) x = build_simple_mem_ref (x); SET_DECL_VALUE_EXPR (new_var, x); DECL_HAS_VALUE_EXPR_P (new_var) = 1; } continue; } + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + && is_omp_target (stmt)) + { + gcc_assert (maybe_lookup_field (c, ctx)); + map_cnt++; + continue; + } + if (!maybe_lookup_field (var, ctx)) continue; /* Don't remap compute constructs' reduction variables, because the intermediate result must be local to each gang. */ if (offloaded && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_IN_REDUCTION (c))) @@ -11703,22 +11769,36 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) { tree ovar2 = DECL_VALUE_EXPR (ovar); gcc_assert (TREE_CODE (ovar2) == INDIRECT_REF); ovar2 = TREE_OPERAND (ovar2, 0); gcc_assert (DECL_P (ovar2)); ovar = ovar2; } - if (!maybe_lookup_field (ovar, ctx)) + if (!maybe_lookup_field (ovar, ctx) + && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))) continue; } talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar)); if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign) talign = DECL_ALIGN_UNIT (ovar); - if (nc) + + if (nc + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + && is_omp_target (stmt)) + { + var = lookup_decl_in_outer_ctx (ovar, ctx); + x = build_sender_ref (c, ctx); + gimplify_assign (x, build_fold_addr_expr (var), &ilist); + } + else if (nc) { var = lookup_decl_in_outer_ctx (ovar, ctx); x = build_sender_ref (ovar, ctx); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) diff --git a/gcc/testsuite/c-c++-common/goacc/finalize-1.c b/gcc/testsuite/c-c++-common/goacc/finalize-1.c index 3d64b2e7cb3..679b0505e19 100644 --- a/gcc/testsuite/c-c++-common/goacc/finalize-1.c +++ b/gcc/testsuite/c-c++-common/goacc/finalize-1.c @@ -17,21 +17,21 @@ void f () #pragma acc exit data finalize delete (del_f) /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } } { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } */ #pragma acc exit data finalize delete (del_f_p[2:5]) /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(del_f_p \\+ 2\\) \\\[len: 5\\\]\\) map\\(firstprivate:del_f_p \\\[pointer assign, bias: 2\\\]\\) finalize;$" 1 "original" } } - { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\[^ \]+ \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */ + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\\*\\(del_f_p \\+ 2\\) \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */ #pragma acc exit data copyout (cpo_r) /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } } { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ #pragma acc exit data copyout (cpo_f) finalize /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:cpo_f\\);$" 1 "original" } } { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ #pragma acc exit data copyout (cpo_f_p[4:10]) finalize /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:\\*\\(cpo_f_p \\+ 4\\) \\\[len: 10\\\]\\) map\\(firstprivate:cpo_f_p \\\[pointer assign, bias: 4\\\]\\);$" 1 "original" } } - { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:\[^ \]+ \\\[len: 10\\\]\\)$" 1 "gimple" } } */ + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:\\*\\(cpo_f_p \\+ 4\\) \\\[len: 10\\\]\\)$" 1 "gimple" } } */ } diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c index 337c1f7cc77..839269eb62b 100644 --- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c @@ -40,15 +40,15 @@ t1 () #pragma acc exit data detach(a) finalize #pragma acc exit data detach(s.a) finalize } } /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:\*.*z.? .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:\*.*s\.a.? .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 0.." 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 0.." 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 0.." 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c index df405e448b2..9f702ba76f2 100644 --- a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c +++ b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c @@ -16,12 +16,12 @@ struct str { void test (int *b, int *c, int *e) { struct str s = { .a = 0, .b = b, .c = c, .d = 0, .e = e, .f = 0 }; #pragma acc enter data copyin(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f) - /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */ + /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*.*s\.b.? \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*.*s\.c.? \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */ #pragma acc exit data copyout(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f) - /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */ + /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*.*s\.b.? \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*.*s\.c.? \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */ } diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-2.c b/gcc/testsuite/c-c++-common/gomp/clauses-2.c index ded1d74ccde..bbc8fb4e32b 100644 --- a/gcc/testsuite/c-c++-common/gomp/clauses-2.c +++ b/gcc/testsuite/c-c++-common/gomp/clauses-2.c @@ -9,46 +9,46 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l) bar (&q); #pragma omp target map (p[0]) firstprivate (p) /* { dg-error "appears more than once in data clauses" } */ bar (p); #pragma omp target firstprivate (p), map (p[0]) /* { dg-error "appears more than once in data clauses" } */ bar (p); #pragma omp target map (p[0]) map (p) /* { dg-error "appears both in data and map clauses" } */ bar (p); - #pragma omp target map (p) , map (p[0]) /* { dg-error "appears both in data and map clauses" } */ + #pragma omp target map (p) , map (p[0]) bar (p); #pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */ bar (&q); #pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than once in data clauses" } */ bar (p); - #pragma omp target map (t) map (t.r) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t) map (t.r) bar (&t.r); - #pragma omp target map (t.r) map (t) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t.r) map (t) bar (&t.r); - #pragma omp target map (t.r) map (t.r) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t.r) map (t.r) bar (&t.r); #pragma omp target firstprivate (t), map (t.r) /* { dg-error "appears both in data and map clauses" } */ bar (&t.r); #pragma omp target map (t.r) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */ bar (&t.r); - #pragma omp target map (t.s[0]) map (t) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t.s[0]) map (t) bar (t.s); - #pragma omp target map (t) map(t.s[0]) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t) map(t.s[0]) bar (t.s); #pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */ bar (t.s); #pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */ bar (t.s); - #pragma omp target map (t.s[0]) map (t.s[2]) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t.s[0]) map (t.s[2]) bar (t.s); - #pragma omp target map (t.t[0:2]) map (t.t[4:6]) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t.t[0:2]) map (t.t[4:6]) bar (t.t); - #pragma omp target map (t.t[i:j]) map (t.t[k:l]) /* { dg-error "appears more than once in map clauses" } */ + #pragma omp target map (t.t[i:j]) map (t.t[k:l]) bar (t.t); #pragma omp target map (t.s[0]) map (t.r) bar (t.s); #pragma omp target map (t.r) ,map (t.s[0]) bar (t.s); #pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */ bar (t.s); #pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */ - bar (t.s); /* { dg-error "appears more than once in map clauses" "" { target *-*-* } .-1 } */ + bar (t.s); } diff --git a/gcc/testsuite/c-c++-common/gomp/map-5.c b/gcc/testsuite/c-c++-common/gomp/map-5.c new file mode 100644 index 00000000000..1d9d9252864 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/map-5.c @@ -0,0 +1,24 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void foo (void) +{ + /* Basic test to ensure to,from,tofrom is ordered before alloc,release,delete clauses. */ + int a, b, c; + #pragma omp target enter data map(alloc:a) map(to:b) map(alloc:c) + #pragma omp target exit data map(from:a) map(release:b) map(from:c) + + #pragma omp target map(alloc:a) map(tofrom:b) map(alloc:c) + a = b = c = 1; + + #pragma omp target enter data map(to:a) map(alloc:b) map(to:c) + #pragma omp target exit data map(from:a) map(delete:b) map(from:c) +} + +/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(release:.*" "gimple" } } */ + +/* { dg-final { scan-tree-dump "pragma omp target num_teams.* map\\(tofrom:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */ + +/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(to:.* map\\(alloc:.*" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(delete:.*" "gimple" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 index 373bdcb2114..c5ac06943eb 100644 --- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 @@ -8,22 +8,22 @@ program att end type t type(t) :: myvar integer, target :: tarr(10) integer, pointer :: myptr(:) !$acc enter data attach(myvar%arr2, myptr) ! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } } -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } } !$acc exit data detach(myvar%arr2, myptr) ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } } -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } } ! Test valid usage and processing of the finalize clause. !$acc exit data detach(myvar%arr2, myptr) finalize ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize;$" 1 "original" } } ! For array-descriptor detaches, we no longer generate a "release" mapping ! for the pointed-to data for gimplify.c to turn into "delete". Make sure ! the mapping still isn't there. -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } } end program att diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f index a7788580819..0ff2e471180 100644 --- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f +++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f @@ -17,21 +17,21 @@ !$ACC EXIT DATA FINALIZE DELETE (del_f) ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } } ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } !$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5)) ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } } -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } !$ACC EXIT DATA COPYOUT (cpo_r) ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } } ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } !$ACC EXIT DATA COPYOUT (cpo_f) FINALIZE ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_f\\) finalize;$" 1 "original" } } ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } !$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.1\\.data - \\(.*int.*\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } } -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } END SUBROUTINE f diff --git a/gcc/testsuite/gfortran.dg/gomp/map-2.f90 b/gcc/testsuite/gfortran.dg/gomp/map-2.f90 index 73c4f5a87d0..79bab726dea 100644 --- a/gcc/testsuite/gfortran.dg/gomp/map-2.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/map-2.f90 @@ -1,6 +1,6 @@ type t integer :: i end type t type(t) v -!$omp target enter data map(to:v%i, v%i) ! { dg-error "appears more than once in map clauses" } +!$omp target enter data map(to:v%i, v%i) end