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;