From patchwork Wed Aug 11 16:58:24 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1516001 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+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) 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 4GlGK86bg6z9sWl for ; Thu, 12 Aug 2021 03:00:12 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 7F8E93985811 for ; Wed, 11 Aug 2021 17:00:10 +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 9EF36398580F for ; Wed, 11 Aug 2021 16:58:45 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 9EF36398580F Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: 9EiPvWkMzEY7+tnujPGrOqTiEpNAc22/WQUiatw1gc9BhMAOziaTN5YDMNazQ9alSb2jDB706k 3f7eei8JZdk0ZEAOeC1+6fCnd69fPPI17nsW8igkpf2yZIxlnpVjHEmpEwPGnPRviz9yniHiXg 9YrGJoB7iJ4DZ8MEWP4btujslKY3YJqApvdzzb8VmllpPB8fOvAm5mSGb1HKWYBNr4qIPY1XiB CgkVlFmTHNZ27/BFrOp9VEwKP/hm6bfgaQxWQQlG9BDBgftEcw47q8Sjg9XRj8B968IxM3gjvN Q5Vwyvz7x+5YwcRFDqaxoFyz X-IronPort-AV: E=Sophos;i="5.84,313,1620720000"; d="scan'208";a="64742300" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 11 Aug 2021 08:58:44 -0800 IronPort-SDR: rxQYY/c/kITmLCHNtamfbtc79cJt2dyUokRkC+/7+89OUIaF1JOLXlJ/bjf3KxwSa31XYHDZ2+ 6T3fcwZ6MbukleY5ojKpCHCZblZEfCsevTiY7iXfn+01Bh7OxVyvzmY5B+IYkounvzG1O/qZl/ eIdFfhIhmabt/ymJpEiCsVSsFqNxlzXpddG+qPgd8/2Ya7+vGOOJN6mFkYQcWBxh9kr84Qabkx gtzayW73ieluf314Mj59VPgJHWOZsHPdo/KyXGe3ga7oiloZL9I2QJMGhJkC7CHQfunpZAnLXR cmk= From: Julian Brown To: Subject: [PATCH 1/8] Improve OpenMP target support for C++ [PR92120 v4b] Date: Wed, 11 Aug 2021 09:58:24 -0700 Message-ID: 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-04.mgc.mentorg.com (139.181.222.4) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: 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 Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" From: Chung-Lin Tang This is a version "v4b" of a patch by Chung-Lin, merged to current mainline. All errors introduced are my own! Previously posted here: https://gcc.gnu.org/pipermail/gcc-patches/2021-June/573166.html Chung-Lin's description from the last submission follows. This patch is the "v4" version of my PR92120 patch, v3 was here: https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570886.html (there I listed the various patches from devel/omp/gcc-10 branch that was combined, which I won't repeat here). Basically this v4 adds fixes for lambda capture, which was already pushed to devel/omp/gcc-11 yesterday: https://gcc.gnu.org/pipermail/gcc-patches/2021-June/572988.html Thanks, Chung-Lin gcc/cp/ * cp-tree.h (finish_omp_target): New declaration. (finish_omp_target_clauses): Likewise. * parser.c (cp_parser_omp_clause_map): Adjust call to cp_parser_omp_var_list_no_open to set 'allow_deref' argument to true. (cp_parser_omp_target): Factor out code, adjust into calls to new function finish_omp_target. * pt.c (tsubst_expr): Add call to finish_omp_target_clauses for OMP_TARGET case. * semantics.c (handle_omp_array_sections_1): Add handling to create 'this->member' from 'member' FIELD_DECL. (handle_omp_array_sections): Likewise. (finish_omp_clauses): Likewise. Adjust to allow 'this[]' in OpenMP map clauses. Handle 'A->member' case in map clauses. (struct omp_target_walk_data): New struct for walking over target-directive tree body. (finish_omp_target_clauses_r): New function for tree walk. (finish_omp_target_clauses): New function. (finish_omp_target): New function. gcc/c/ * c-parser.c (c_parser_omp_clause_map): Set 'allow_deref' argument in call to c_parser_omp_variable_list to 'true'. * c-typeck.c (handle_omp_array_sections_1): Add strip of MEM_REF in array base handling. (c_finish_omp_clauses): Handle 'A->member' case in map clauses. gcc/ * gimplify.c ("tree-hash-traits.h"): Add include. (gimplify_scan_omp_clauses): Change struct_map_to_clause to type hash_map *. Adjust struct map handling to handle cases of *A and A->B expressions. Under !DECL_P case of GOMP_CLAUSE_MAP handling, add STRIP_NOPS for indir_p case, add to struct_deref_set for map(*ptr_to_struct) cases. Add MEM_REF case when handling component_ref_p case. Add unshare_expr and gimplification when created GOMP_MAP_STRUCT is not a DECL. Add code to add firstprivate pointer for *pointer-to-struct case. (gimplify_adjust_omp_clauses): Move GOMP_MAP_STRUCT removal code for exit data directives code to earlier position. * omp-low.c (lower_omp_target): Handle GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds. * tree-pretty-print.c (dump_omp_clause): Likewise. gcc/testsuite/ * gcc.dg/gomp/target-3.c: New testcase. * g++.dg/gomp/target-3.C: New testcase. * g++.dg/gomp/target-lambda-1.C: New testcase. * g++.dg/gomp/target-lambda-2.C: New testcase. * g++.dg/gomp/target-this-1.C: New testcase. * g++.dg/gomp/target-this-2.C: New testcase. * g++.dg/gomp/target-this-3.C: New testcase. * g++.dg/gomp/target-this-4.C: New testcase. * g++.dg/gomp/target-this-5.C: New testcase. * g++.dg/gomp/this-2.C: Adjust testcase. include/ * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds. (GOMP_MAP_POINTER_P): Include GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION. libgomp/ * libgomp.h (gomp_attach_pointer): Add bool parameter. * oacc-mem.c (acc_attach_async): Update call to gomp_attach_pointer. (goacc_enter_data_internal): Likewise. * target.c (gomp_map_vars_existing): Update assert condition to include GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION. (gomp_map_pointer): Add 'bool allow_zero_length_array_sections' parameter, add support for mapping a pointer with NULL target. (gomp_attach_pointer): Add 'bool allow_zero_length_array_sections' parameter, add support for attaching a pointer with NULL target. (gomp_map_vars_internal): Update calls to gomp_map_pointer and gomp_attach_pointer, add handling for GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION cases. * testsuite/libgomp.c++/target-23.C: New testcase. * testsuite/libgomp.c++/target-lambda-1.C: New testcase. * testsuite/libgomp.c++/target-lambda-2.C: New testcase. * testsuite/libgomp.c++/target-this-1.C: New testcase. * testsuite/libgomp.c++/target-this-2.C: New testcase. * testsuite/libgomp.c++/target-this-3.C: New testcase. * testsuite/libgomp.c++/target-this-4.C: New testcase. * testsuite/libgomp.c++/target-this-5.C: New testcase. -------------- next part -------------- more more more wip --- gcc/c/c-parser.c | 3 +- gcc/c/c-typeck.c | 23 +- gcc/cp/cp-tree.h | 2 + gcc/cp/parser.c | 70 +-- gcc/cp/pt.c | 5 + gcc/cp/semantics.c | 540 +++++++++++++++++- gcc/gimplify.c | 118 +++- gcc/omp-low.c | 2 + gcc/testsuite/g++.dg/gomp/target-3.C | 36 ++ gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 94 +++ gcc/testsuite/g++.dg/gomp/target-lambda-2.C | 35 ++ gcc/testsuite/g++.dg/gomp/target-this-1.C | 33 ++ gcc/testsuite/g++.dg/gomp/target-this-2.C | 49 ++ gcc/testsuite/g++.dg/gomp/target-this-3.C | 105 ++++ gcc/testsuite/g++.dg/gomp/target-this-4.C | 107 ++++ gcc/testsuite/g++.dg/gomp/target-this-5.C | 34 ++ gcc/testsuite/g++.dg/gomp/this-2.C | 24 +- gcc/testsuite/gcc.dg/gomp/target-3.c | 16 + gcc/tree-pretty-print.c | 8 + include/gomp-constants.h | 14 +- libgomp/libgomp.h | 2 +- libgomp/oacc-mem.c | 7 +- libgomp/target.c | 76 ++- libgomp/testsuite/libgomp.c++/target-23.C | 34 ++ .../testsuite/libgomp.c++/target-lambda-1.C | 86 +++ .../testsuite/libgomp.c++/target-lambda-2.C | 30 + libgomp/testsuite/libgomp.c++/target-this-1.C | 29 + libgomp/testsuite/libgomp.c++/target-this-2.C | 47 ++ libgomp/testsuite/libgomp.c++/target-this-3.C | 99 ++++ libgomp/testsuite/libgomp.c++/target-this-4.C | 104 ++++ libgomp/testsuite/libgomp.c++/target-this-5.C | 30 + 31 files changed, 1734 insertions(+), 128 deletions(-) create mode 100644 gcc/testsuite/g++.dg/gomp/target-3.C create mode 100644 gcc/testsuite/g++.dg/gomp/target-lambda-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/target-lambda-2.C create mode 100644 gcc/testsuite/g++.dg/gomp/target-this-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/target-this-2.C create mode 100644 gcc/testsuite/g++.dg/gomp/target-this-3.C create mode 100644 gcc/testsuite/g++.dg/gomp/target-this-4.C create mode 100644 gcc/testsuite/g++.dg/gomp/target-this-5.C create mode 100644 gcc/testsuite/gcc.dg/gomp/target-3.c create mode 100644 libgomp/testsuite/libgomp.c++/target-23.C create mode 100644 libgomp/testsuite/libgomp.c++/target-lambda-1.C create mode 100644 libgomp/testsuite/libgomp.c++/target-lambda-2.C create mode 100644 libgomp/testsuite/libgomp.c++/target-this-1.C create mode 100644 libgomp/testsuite/libgomp.c++/target-this-2.C create mode 100644 libgomp/testsuite/libgomp.c++/target-this-3.C create mode 100644 libgomp/testsuite/libgomp.c++/target-this-4.C create mode 100644 libgomp/testsuite/libgomp.c++/target-this-5.C diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 9a56e0c04c6..e8aaec75677 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -15854,7 +15854,8 @@ c_parser_omp_clause_map (c_parser *parser, tree list) c_parser_consume_token (parser); } - nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list); + nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list, + true); for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_SET_MAP_KIND (c, kind); diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 5349ef1f392..773cd2f8703 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13111,6 +13111,11 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == MEM_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + } if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF) { if (maybe_ne (mem_ref_offset (t), 0)) @@ -13956,6 +13961,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) tree ordered_clause = NULL_TREE; tree schedule_clause = NULL_TREE; bool oacc_async = false; + bool indir_component_ref_p = false; tree last_iterators = NULL_TREE; bool last_iterators_remove = false; tree *nogroup_seen = NULL; @@ -14757,6 +14763,11 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { while (TREE_CODE (t) == COMPONENT_REF) t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == MEM_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + } if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_IMPLICIT (c) && (bitmap_bit_p (&map_head, DECL_UID (t)) @@ -14823,6 +14834,14 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bias) to zero here, so it is not set erroneously to the pointer size later on in gimplify.c. */ OMP_CLAUSE_SIZE (c) = size_zero_node; + indir_component_ref_p = false; + if (TREE_CODE (t) == COMPONENT_REF + && TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF) + { + t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); + indir_component_ref_p = true; + STRIP_NOPS (t); + } if (TREE_CODE (t) == COMPONENT_REF && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) { @@ -14895,6 +14914,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP || (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER)) + && !indir_component_ref_p && !c_mark_addressable (t)) remove = true; else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -14951,8 +14971,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_set_bit (&map_firstprivate_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), diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index d4810c0c986..83592107170 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -7601,6 +7601,8 @@ extern tree start_lambda_function (tree fn, tree lambda_expr); extern void finish_lambda_function (tree body); extern bool regenerated_lambda_fn_p (tree); extern tree most_general_lambda (tree); +extern tree finish_omp_target (location_t, tree, tree, bool); +extern void finish_omp_target_clauses (location_t, tree, tree *); /* in tree.c */ extern int cp_tree_operand_length (const_tree); diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 93698aa14c9..d90408aa3a1 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -38527,7 +38527,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) } nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list, - NULL); + NULL, true); for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_SET_MAP_KIND (c, kind); @@ -42583,8 +42583,6 @@ static bool cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, enum pragma_context context, bool *if_p) { - tree *pc = NULL, stmt; - if (flag_openmp) omp_requires_mask = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED); @@ -42688,16 +42686,10 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc; } } - tree stmt = make_node (OMP_TARGET); - TREE_TYPE (stmt) = void_type_node; - OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; - c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); - OMP_TARGET_BODY (stmt) = body; - OMP_TARGET_COMBINED (stmt) = 1; - SET_EXPR_LOCATION (stmt, pragma_tok->location); - add_stmt (stmt); - pc = &OMP_TARGET_CLAUSES (stmt); - goto check_clauses; + c_omp_adjust_map_clauses (cclauses[C_OMP_CLAUSE_SPLIT_TARGET], true); + finish_omp_target (pragma_tok->location, + cclauses[C_OMP_CLAUSE_SPLIT_TARGET], body, true); + return true; } else if (!flag_openmp) /* flag_openmp_simd */ { @@ -42734,13 +42726,10 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, return false; } - 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, false); - for (tree c = OMP_TARGET_CLAUSES (stmt); c; c = OMP_CLAUSE_CHAIN (c)) + tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, + "#pragma omp target", pragma_tok, + false); + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION) { tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); @@ -42749,45 +42738,12 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = nc; } - OMP_TARGET_CLAUSES (stmt) - = finish_omp_clauses (OMP_TARGET_CLAUSES (stmt), C_ORT_OMP_TARGET); - c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); - - pc = &OMP_TARGET_CLAUSES (stmt); + clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET); + c_omp_adjust_map_clauses (clauses, true); keep_next_level (true); - OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p); + tree body = cp_parser_omp_structured_block (parser, if_p); - SET_EXPR_LOCATION (stmt, pragma_tok->location); - add_stmt (stmt); - -check_clauses: - while (*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_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_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); - continue; - } - pc = &OMP_CLAUSE_CHAIN (*pc); - } + finish_omp_target (pragma_tok->location, clauses, body, false); return true; } diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index 7e56ccfc45f..f88333e0733 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -18889,6 +18889,11 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, t = copy_node (t); OMP_BODY (t) = stmt; OMP_CLAUSES (t) = tmp; + + if (TREE_CODE (t) == OMP_TARGET) + finish_omp_target_clauses (EXPR_LOCATION (t), OMP_BODY (t), + &OMP_CLAUSES (t)); + if (TREE_CODE (t) == OMP_TARGET && OMP_TARGET_COMBINED (t)) { tree teams = cp_walk_tree (&stmt, tsubst_find_omp_teams, NULL, NULL); diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index b080259083e..6778efae606 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5019,15 +5019,16 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } t = TREE_OPERAND (t, 0); - if (ort == C_ORT_ACC && TREE_CODE (t) == INDIRECT_REF) - t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == INDIRECT_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + } } if (REFERENCE_REF_P (t)) t = TREE_OPERAND (t, 0); } - if (TREE_CODE (t) == FIELD_DECL - && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY - || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND)) + if (TREE_CODE (t) == FIELD_DECL) ret = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); else if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) { @@ -5044,6 +5045,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } else if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_TO + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FROM && TREE_CODE (t) == PARM_DECL && DECL_ARTIFICIAL (t) && DECL_NAME (t) == this_identifier @@ -5568,6 +5572,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) } OMP_CLAUSE_DECL (c) = first; OMP_CLAUSE_SIZE (c) = size; + if (TREE_CODE (t) == FIELD_DECL) + t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP || (TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)) @@ -6581,6 +6587,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bool order_seen = false; bool schedule_seen = false; bool oacc_async = false; + bool indir_component_ref_p = false; tree last_iterators = NULL_TREE; bool last_iterators_remove = false; /* 1 if normal/task reduction has been seen, -1 if inscan reduction @@ -7738,6 +7745,11 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); if (REFERENCE_REF_P (t)) t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == INDIRECT_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + } if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_IMPLICIT (c) && (bitmap_bit_p (&map_head, DECL_UID (t)) @@ -7810,9 +7822,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); OMP_CLAUSE_DECL (c) = t; } + indir_component_ref_p = false; if (TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF) - t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); + { + t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); + indir_component_ref_p = true; + STRIP_NOPS (t); + } if (TREE_CODE (t) == COMPONENT_REF && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) { @@ -7859,6 +7876,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) goto handle_map_references; } } + if (!processing_template_decl && TREE_CODE (t) == FIELD_DECL) + { + OMP_CLAUSE_DECL (c) = finish_non_static_data_member (t, NULL_TREE, + NULL_TREE); + break; + } if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) { if (processing_template_decl && TREE_CODE (t) != OVERLOAD) @@ -7885,19 +7908,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } - else if (ort != C_ORT_ACC && t == current_class_ptr) - { - error_at (OMP_CLAUSE_LOCATION (c), - "% allowed in OpenMP only in %" - " clauses"); - remove = true; - break; - } else if (!processing_template_decl && !TYPE_REF_P (TREE_TYPE (t)) && (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP || (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER)) + && !indir_component_ref_p && !cxx_mark_addressable (t)) remove = true; else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -9029,6 +9045,502 @@ finish_omp_construct (enum tree_code code, tree body, tree clauses) return add_stmt (stmt); } +/* Used to walk OpenMP target directive body. */ + +struct omp_target_walk_data +{ + tree current_object; + bool this_expr_accessed; + + hash_map ptr_members_accessed; + hash_set lambda_objects_accessed; + + tree current_closure; + hash_set closure_vars_accessed; + + hash_set local_decls; +}; + +static tree +finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr) +{ + tree t = *tp; + struct omp_target_walk_data *data = (struct omp_target_walk_data *) ptr; + tree current_object = data->current_object; + tree current_closure = data->current_closure; + + if (current_object) + { + tree this_expr = TREE_OPERAND (current_object, 0); + + if (operand_equal_p (t, this_expr)) + { + data->this_expr_accessed = true; + *walk_subtrees = 0; + return NULL_TREE; + } + + if (TREE_CODE (t) == COMPONENT_REF + && POINTER_TYPE_P (TREE_TYPE (t)) + && operand_equal_p (TREE_OPERAND (t, 0), current_object) + && TREE_CODE (TREE_OPERAND (t, 1)) == FIELD_DECL) + { + data->this_expr_accessed = true; + tree fld = TREE_OPERAND (t, 1); + if (data->ptr_members_accessed.get (fld) == NULL) + { + if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE) + t = convert_from_reference (t); + data->ptr_members_accessed.put (fld, t); + } + *walk_subtrees = 0; + return NULL_TREE; + } + } + + /* When the current_function_decl is a lambda function, the closure object + argument's type seems to not yet have fields layed out, so a recording + of DECL_VALUE_EXPRs during the target body walk seems the only way to + find them. */ + if (current_closure + && (TREE_CODE (t) == VAR_DECL + || TREE_CODE (t) == PARM_DECL + || TREE_CODE (t) == RESULT_DECL) + && DECL_HAS_VALUE_EXPR_P (t) + && TREE_CODE (DECL_VALUE_EXPR (t)) == COMPONENT_REF + && operand_equal_p (current_closure, + TREE_OPERAND (DECL_VALUE_EXPR (t), 0))) + { + if (!data->closure_vars_accessed.contains (t)) + data->closure_vars_accessed.add (t); + *walk_subtrees = 0; + return NULL_TREE; + } + + if (TREE_CODE (t) == BIND_EXPR) + { + tree block = BIND_EXPR_BLOCK (t); + for (tree var = BLOCK_VARS (block); var; var = DECL_CHAIN (var)) + if (!data->local_decls.contains (var)) + data->local_decls.add (var); + return NULL_TREE; + } + + if (TREE_CODE (t) == BIND_EXPR) + { + tree block = BIND_EXPR_BLOCK (t); + for (tree var = BLOCK_VARS (block); var; var = DECL_CHAIN (var)) + if (!data->local_decls.contains (var)) + data->local_decls.add (var); + return NULL_TREE; + } + + if (TREE_TYPE(t) && LAMBDA_TYPE_P (TREE_TYPE (t))) + { + tree lt = TREE_TYPE (t); + gcc_assert (CLASS_TYPE_P (lt)); + + if (!data->lambda_objects_accessed.contains (t) + /* Do not prepare to create target maps for locally declared + lambdas or anonymous ones. */ + && !data->local_decls.contains (t) + && TREE_CODE (t) != TARGET_EXPR) + data->lambda_objects_accessed.add (t); + *walk_subtrees = 0; + return NULL_TREE; + } + + return NULL_TREE; +} + +void +finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr) +{ + omp_target_walk_data data; + data.this_expr_accessed = false; + + tree ct = current_nonlambda_class_type (); + if (ct) + { + tree object = maybe_dummy_object (ct, NULL); + object = maybe_resolve_dummy (object, true); + data.current_object = object; + } + else + data.current_object = NULL_TREE; + + if (DECL_LAMBDA_FUNCTION_P (current_function_decl)) + { + tree closure = DECL_ARGUMENTS (current_function_decl); + data.current_closure = build_indirect_ref (loc, closure, RO_UNARY_STAR); + } + else + data.current_closure = NULL_TREE; + + cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r, &data); + + auto_vec new_clauses; + + if (data.this_expr_accessed) + { + tree omp_target_this_expr = TREE_OPERAND (data.current_object, 0); + + /* See if explicit user-specified map(this[:]) clause already exists. + If not, we create an implicit map(tofrom:this[:1]) clause. */ + tree *explicit_this_deref_map = NULL; + for (tree *c = clauses_ptr; *c; c = &OMP_CLAUSE_CHAIN (*c)) + if (OMP_CLAUSE_CODE (*c) == OMP_CLAUSE_MAP + && TREE_CODE (OMP_CLAUSE_DECL (*c)) == INDIRECT_REF + && operand_equal_p (TREE_OPERAND (OMP_CLAUSE_DECL (*c), 0), + omp_target_this_expr)) + { + explicit_this_deref_map = c; + break; + } + + if (DECL_LAMBDA_FUNCTION_P (current_function_decl)) + { + /* For lambda functions, we need to first create a copy of the + __closure object. */ + tree closure = DECL_ARGUMENTS (current_function_decl); + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO); + OMP_CLAUSE_DECL (c) + = build_indirect_ref (loc, closure, RO_UNARY_STAR); + OMP_CLAUSE_SIZE (c) + = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure))); + new_clauses.safe_push (c); + + tree closure_obj = OMP_CLAUSE_DECL (c); + tree closure_type = TREE_TYPE (closure_obj); + + gcc_assert (LAMBDA_TYPE_P (closure_type) + && CLASS_TYPE_P (closure_type)); + + tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); + OMP_CLAUSE_DECL (c2) = closure; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + new_clauses.safe_push (c2); + + STRIP_NOPS (omp_target_this_expr); + gcc_assert (DECL_HAS_VALUE_EXPR_P (omp_target_this_expr)); + omp_target_this_expr = DECL_VALUE_EXPR (omp_target_this_expr); + + for (hash_set::iterator i = data.closure_vars_accessed.begin (); + i != data.closure_vars_accessed.end (); ++i) + { + tree orig_decl = *i; + tree closure_expr = DECL_VALUE_EXPR (orig_decl); + + if (TREE_CODE (TREE_TYPE (orig_decl)) == POINTER_TYPE) + { + /* this-pointer is processed outside this loop. */ + if (operand_equal_p (closure_expr, omp_target_this_expr)) + continue; + + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC); + OMP_CLAUSE_DECL (c) + = build_indirect_ref (loc, closure_expr, RO_UNARY_STAR); + OMP_CLAUSE_SIZE (c) = size_zero_node; + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + new_clauses.safe_push (c); + + c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND + (c, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); + OMP_CLAUSE_DECL (c) = closure_expr; + OMP_CLAUSE_SIZE (c) = size_zero_node; + new_clauses.safe_push (c); + } + else if (TREE_CODE (TREE_TYPE (orig_decl)) == REFERENCE_TYPE) + { + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO); + OMP_CLAUSE_DECL (c) + = build1 (INDIRECT_REF, + TREE_TYPE (TREE_TYPE (closure_expr)), + closure_expr); + OMP_CLAUSE_SIZE (c) + = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure_expr))); + new_clauses.safe_push (c); + + c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); + OMP_CLAUSE_DECL (c) = closure_expr; + OMP_CLAUSE_SIZE (c) = size_zero_node; + new_clauses.safe_push (c); + } + } + + if (explicit_this_deref_map) + { + /* Transform *this into *__closure->this in maps. */ + tree this_map = *explicit_this_deref_map; + OMP_CLAUSE_DECL (this_map) + = build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR); + + tree nc = OMP_CLAUSE_CHAIN (this_map); + gcc_assert (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nc) + == GOMP_MAP_FIRSTPRIVATE_POINTER)); + OMP_CLAUSE_DECL (nc) = omp_target_this_expr; + OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_ALWAYS_POINTER); + + /* Unlink this two-map sequence away from the chain. */ + *explicit_this_deref_map = OMP_CLAUSE_CHAIN (nc); + + /* Move map(*__closure->this) map(always_pointer:__closure->this) + sequence to right after __closure map. */ + new_clauses.safe_push (this_map); + new_clauses.safe_push (nc); + } + else + { + tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_TOFROM); + OMP_CLAUSE_DECL (c3) + = build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR); + OMP_CLAUSE_SIZE (c3) + = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr))); + + tree c4 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c4, GOMP_MAP_ALWAYS_POINTER); + + OMP_CLAUSE_DECL (c4) = omp_target_this_expr; + OMP_CLAUSE_SIZE (c4) = size_zero_node; + + new_clauses.safe_push (c3); + new_clauses.safe_push (c4); + } + } + else + { + /* For the non-lambda case, we only need to create map(this[:1]) when + it's not present, no transforming needed. */ + if (!explicit_this_deref_map) + { + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM); + OMP_CLAUSE_DECL (c) + = build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR); + OMP_CLAUSE_SIZE (c) + = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr))); + + tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); + STRIP_NOPS (omp_target_this_expr); + OMP_CLAUSE_DECL (c2) = omp_target_this_expr; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + + new_clauses.safe_push (c); + new_clauses.safe_push (c2); + } + } + + if (!data.ptr_members_accessed.is_empty ()) + for (hash_map::iterator i + = data.ptr_members_accessed.begin (); + i != data.ptr_members_accessed.end (); ++i) + { + /* For each referenced member that is of pointer or + reference-to-pointer type, create the equivalent of + map(alloc:this->ptr[:0]). */ + tree field_decl = (*i).first; + tree ptr_member = (*i).second; + + for (tree c = *clauses_ptr; c; c = OMP_CLAUSE_CHAIN (c)) + { + /* If map(this->ptr[:N] already exists, avoid creating another + such map. */ + tree decl = OMP_CLAUSE_DECL (c); + if ((TREE_CODE (decl) == INDIRECT_REF + || TREE_CODE (decl) == MEM_REF) + && operand_equal_p (TREE_OPERAND (decl, 0), + ptr_member)) + goto next_ptr_member; + } + + if (!cxx_mark_addressable (ptr_member)) + gcc_unreachable (); + + if (TREE_CODE (TREE_TYPE (field_decl)) == REFERENCE_TYPE) + { + /* For reference to pointers, we need to map the referenced + pointer first for things to be correct. */ + tree ptr_member_type = TREE_TYPE (ptr_member); + + /* Map pointer target as zero-length array section. */ + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC); + OMP_CLAUSE_DECL (c) + = build1 (INDIRECT_REF, TREE_TYPE (ptr_member_type), ptr_member); + OMP_CLAUSE_SIZE (c) = size_zero_node; + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + + /* Map pointer to zero-length array section. */ + tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND + (c2, GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION); + OMP_CLAUSE_DECL (c2) = ptr_member; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + + /* Attach reference-to-pointer field to pointer. */ + tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH); + OMP_CLAUSE_DECL (c3) = TREE_OPERAND (ptr_member, 0); + OMP_CLAUSE_SIZE (c3) = size_zero_node; + + new_clauses.safe_push (c); + new_clauses.safe_push (c2); + new_clauses.safe_push (c3); + } + else if (TREE_CODE (TREE_TYPE (field_decl)) == POINTER_TYPE) + { + /* Map pointer target as zero-length array section. */ + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC); + OMP_CLAUSE_DECL (c) + = build_indirect_ref (loc, ptr_member, RO_UNARY_STAR); + OMP_CLAUSE_SIZE (c) = size_zero_node; + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + + /* Attach zero-length array section to pointer. */ + tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND + (c2, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); + OMP_CLAUSE_DECL (c2) = ptr_member; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + + new_clauses.safe_push (c); + new_clauses.safe_push (c2); + } + else + gcc_unreachable (); + + next_ptr_member: + ; + } + } + + if (!data.lambda_objects_accessed.is_empty ()) + { + for (hash_set::iterator i = data.lambda_objects_accessed.begin (); + i != data.lambda_objects_accessed.end (); ++i) + { + tree lobj = *i; + if (TREE_CODE (lobj) == TARGET_EXPR) + lobj = TREE_OPERAND (lobj, 0); + + tree lt = TREE_TYPE (lobj); + gcc_assert (LAMBDA_TYPE_P (lt) && CLASS_TYPE_P (lt)); + + tree lc = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (lc, GOMP_MAP_TO); + OMP_CLAUSE_DECL (lc) = lobj; + OMP_CLAUSE_SIZE (lc) = TYPE_SIZE_UNIT (lt); + new_clauses.truncate (0); + new_clauses.safe_push (lc); + + for (tree fld = TYPE_FIELDS (lt); fld; fld = DECL_CHAIN (fld)) + { + if (TREE_CODE (TREE_TYPE (fld)) == POINTER_TYPE) + { + tree exp = build3 (COMPONENT_REF, TREE_TYPE (fld), + lobj, fld, NULL_TREE); + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC); + OMP_CLAUSE_DECL (c) + = build_indirect_ref (loc, exp, RO_UNARY_STAR); + OMP_CLAUSE_SIZE (c) = size_zero_node; + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + new_clauses.safe_push (c); + + c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND + (c, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); + OMP_CLAUSE_DECL (c) = exp; + OMP_CLAUSE_SIZE (c) = size_zero_node; + new_clauses.safe_push (c); + } + else if (TREE_CODE (TREE_TYPE (fld)) == REFERENCE_TYPE) + { + tree exp = build3 (COMPONENT_REF, TREE_TYPE (fld), + lobj, fld, NULL_TREE); + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM); + OMP_CLAUSE_DECL (c) + = build1 (INDIRECT_REF, TREE_TYPE (TREE_TYPE (exp)), exp); + OMP_CLAUSE_SIZE (c) + = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (exp))); + new_clauses.safe_push (c); + + c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); + OMP_CLAUSE_DECL (c) = exp; + OMP_CLAUSE_SIZE (c) = size_zero_node; + new_clauses.safe_push (c); + } + } + } + } + + tree c = *clauses_ptr; + for (int i = new_clauses.length () - 1; i >= 0; i--) + { + OMP_CLAUSE_CHAIN (new_clauses[i]) = c; + c = new_clauses[i]; + } + *clauses_ptr = c; +} + +tree +finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p) +{ + if (!processing_template_decl) + finish_omp_target_clauses (loc, body, &clauses); + + tree stmt = make_node (OMP_TARGET); + TREE_TYPE (stmt) = void_type_node; + OMP_TARGET_CLAUSES (stmt) = clauses; + OMP_TARGET_BODY (stmt) = body; + OMP_TARGET_COMBINED (stmt) = combined_p; + SET_EXPR_LOCATION (stmt, loc); + + tree c = clauses; + while (c) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_TO: + case GOMP_MAP_ALWAYS_TO: + 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_FIRSTPRIVATE_REFERENCE: + case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: + case GOMP_MAP_ATTACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + break; + default: + error_at (OMP_CLAUSE_LOCATION (c), + "%<#pragma omp target%> with map-type other " + "than %, %, % or % " + "on % clause"); + break; + } + c = OMP_CLAUSE_CHAIN (c); + } + return add_stmt (stmt); +} + tree finish_omp_parallel (tree clauses, tree body) { diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 75a4a9d59fd..e508ef362a1 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -53,6 +53,7 @@ along with GCC; see the file COPYING3. If not see #include "langhooks.h" #include "tree-cfg.h" #include "tree-ssa.h" +#include "tree-hash-traits.h" #include "omp-general.h" #include "omp-low.h" #include "gimple-low.h" @@ -8742,7 +8743,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, { struct gimplify_omp_ctx *ctx, *outer_ctx; tree c; - hash_map *struct_map_to_clause = NULL; + hash_map *struct_map_to_clause = NULL; hash_set *struct_deref_set = NULL; tree *prev_list_p = NULL, *orig_list_p = list_p; int handled_depend_iterators = -1; @@ -9180,7 +9181,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, GOVD_FIRSTPRIVATE | GOVD_SEEN); } - if (!DECL_P (decl)) + if (TREE_CODE (decl) == TARGET_EXPR) + { + if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL, + is_gimple_lvalue, fb_lvalue) + == GS_ERROR) + remove = true; + } + else if (!DECL_P (decl)) { tree d = decl, *pd; if (TREE_CODE (d) == ARRAY_REF) @@ -9196,12 +9204,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, && TREE_CODE (decl) == INDIRECT_REF && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == REFERENCE_TYPE)) + == REFERENCE_TYPE) + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)) { pd = &TREE_OPERAND (decl, 0); decl = TREE_OPERAND (decl, 0); } bool indir_p = false; + bool component_ref_p = false; tree orig_decl = decl; tree decl_ref = NULL_TREE; if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0 @@ -9212,6 +9223,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, while (TREE_CODE (decl) == COMPONENT_REF) { decl = TREE_OPERAND (decl, 0); + component_ref_p = true; if (((TREE_CODE (decl) == MEM_REF && integer_zerop (TREE_OPERAND (decl, 1))) || INDIRECT_REF_P (decl)) @@ -9220,6 +9232,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, { indir_p = true; decl = TREE_OPERAND (decl, 0); + STRIP_NOPS (decl); } if (TREE_CODE (decl) == INDIRECT_REF && DECL_P (TREE_OPERAND (decl, 0)) @@ -9231,8 +9244,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } } } - else if (TREE_CODE (decl) == COMPONENT_REF) + else if (TREE_CODE (decl) == COMPONENT_REF + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)) { + component_ref_p = true; while (TREE_CODE (decl) == COMPONENT_REF) decl = TREE_OPERAND (decl, 0); if (TREE_CODE (decl) == INDIRECT_REF @@ -9302,7 +9318,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (code == OACC_UPDATE && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); - if (DECL_P (decl) + if ((DECL_P (decl) + || (component_ref_p + && (INDIRECT_REF_P (decl) + || TREE_CODE (decl) == MEM_REF))) && 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 @@ -9359,7 +9378,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, gcc_assert (base == decl); splay_tree_node n - = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); + = (DECL_P (decl) + ? splay_tree_lookup (ctx->variables, + (splay_tree_key) decl) + : NULL); bool ptr = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER); bool attach_detach = (OMP_CLAUSE_MAP_KIND (c) @@ -9385,7 +9407,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, OMP_CLAUSE_SET_MAP_KIND (c, k); has_attachments = true; } - if (n == NULL || (n->value & GOVD_MAP) == 0) + if ((DECL_P (decl) + && (n == NULL || (n->value & GOVD_MAP) == 0)) + || (!DECL_P (decl) + && (!struct_map_to_clause + || struct_map_to_clause->get (decl) == NULL))) { tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); @@ -9396,7 +9422,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (base_ref) OMP_CLAUSE_DECL (l) = unshare_expr (base_ref); else - OMP_CLAUSE_DECL (l) = decl; + { + OMP_CLAUSE_DECL (l) = unshare_expr (decl); + if (!DECL_P (OMP_CLAUSE_DECL (l)) + && (gimplify_expr (&OMP_CLAUSE_DECL (l), + pre_p, NULL, is_gimple_lvalue, + fb_lvalue) + == GS_ERROR)) + { + remove = true; + break; + } + } OMP_CLAUSE_SIZE (l) = (!attach ? size_int (1) @@ -9404,7 +9441,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l)) : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l)))); if (struct_map_to_clause == NULL) - struct_map_to_clause = new hash_map; + struct_map_to_clause + = new hash_map; struct_map_to_clause->put (decl, l); if (ptr || attach_detach) { @@ -9438,7 +9476,32 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, flags |= GOVD_SEEN; if (has_attachments) flags |= GOVD_MAP_HAS_ATTACHMENTS; - goto do_add_decl; + + /* If this is a *pointer-to-struct expression, make sure a + firstprivate map of the base-pointer exists. */ + if (component_ref_p + && ((TREE_CODE (decl) == MEM_REF + && integer_zerop (TREE_OPERAND (decl, 1))) + || INDIRECT_REF_P (decl)) + && DECL_P (TREE_OPERAND (decl, 0)) + && !splay_tree_lookup (ctx->variables, + ((splay_tree_key) + TREE_OPERAND (decl, 0)))) + { + decl = TREE_OPERAND (decl, 0); + tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + enum gomp_map_kind mkind + = GOMP_MAP_FIRSTPRIVATE_POINTER; + OMP_CLAUSE_SET_MAP_KIND (c2, mkind); + OMP_CLAUSE_DECL (c2) = decl; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = c2; + } + + if (DECL_P (decl)) + goto do_add_decl; } else if (struct_map_to_clause) { @@ -9547,6 +9610,13 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } else if (*sc != c) { + if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, + fb_lvalue) + == GS_ERROR) + { + remove = true; + break; + } *list_p = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = *sc; *sc = c; @@ -9682,6 +9752,24 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } + /* If this was of the form map(*pointer_to_struct), then the + 'pointer_to_struct' DECL should be considered deref'ed. */ + if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC + || GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (c)) + || GOMP_MAP_COPY_FROM_P (OMP_CLAUSE_MAP_KIND (c))) + && INDIRECT_REF_P (orig_decl) + && DECL_P (TREE_OPERAND (orig_decl, 0)) + && TREE_CODE (TREE_TYPE (orig_decl)) == RECORD_TYPE) + { + tree ptr = TREE_OPERAND (orig_decl, 0); + if (!struct_deref_set || !struct_deref_set->contains (ptr)) + { + if (!struct_deref_set) + struct_deref_set = new hash_set (); + struct_deref_set->add (ptr); + } + } + if (!remove && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH @@ -10962,6 +11050,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, } } } + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT + && (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA)) + { + remove = true; + break; + } if (!DECL_P (decl)) { if ((ctx->region_type & ORT_TARGET) != 0 @@ -11008,10 +11102,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, = OMP_CLAUSE_CHAIN (OMP_CLAUSE_CHAIN (c)); } } - else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT - && (code == OMP_TARGET_EXIT_DATA - || code == OACC_EXIT_DATA)) - remove = true; else if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER diff --git a/gcc/omp-low.c b/gcc/omp-low.c index e7049c825a4..3f21c47c0ba 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -12367,6 +12367,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_ALWAYS_POINTER: case GOMP_MAP_ATTACH: case GOMP_MAP_DETACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: break; case GOMP_MAP_IF_PRESENT: case GOMP_MAP_FORCE_ALLOC: diff --git a/gcc/testsuite/g++.dg/gomp/target-3.C b/gcc/testsuite/g++.dg/gomp/target-3.C new file mode 100644 index 00000000000..f4d40ec8e4b --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-3.C @@ -0,0 +1,36 @@ +// { dg-do compile } +// { dg-options "-fopenmp -fdump-tree-gimple" } + +struct S +{ + int a, b; + void bar (int); +}; + +void +S::bar (int x) +{ + #pragma omp target map (alloc: a, b) + ; + #pragma omp target enter data map (alloc: a, b) +} + +template +struct T +{ + int a, b; + void bar (int); +}; + +template +void +T::bar (int x) +{ + #pragma omp target map (alloc: a, b) + ; + #pragma omp target enter data map (alloc: a, b) +} + +template struct T<0>; + +/* { dg-final { scan-tree-dump-times "map\\(struct:\\*this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C new file mode 100644 index 00000000000..7dceef80f47 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C @@ -0,0 +1,94 @@ +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14 -fdump-tree-gimple" } +#include +#include + +template +void +omp_target_loop (int begin, int end, L loop) +{ + #pragma omp target teams distribute parallel for + for (int i = begin; i < end; i++) + loop (i); +} + +struct S +{ + int a, len; + int *ptr; + + auto merge_data_func (int *iptr, int &b) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + mapped = (ptr != NULL && iptr != NULL); + if (mapped) + { + for (int i = 0; i < len; i++) + ptr[i] += a + b + iptr[i]; + } + } + return mapped; + }; + return fn; + } +}; + +int x = 1; + +int main (void) +{ + const int N = 10; + int *data1 = new int[N]; + int *data2 = new int[N]; + memset (data1, 0xab, sizeof (int) * N); + memset (data1, 0xcd, sizeof (int) * N); + + int val = 1; + int &valref = val; + #pragma omp target enter data map(alloc: data1[:N], data2[:N]) + + omp_target_loop (0, N, [=](int i) { data1[i] = val; }); + omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; }); + + #pragma omp target update from(data1[:N], data2[:N]) + + for (int i = 0; i < N; i++) + { + if (data1[i] != 1) abort (); + if (data2[i] != 2) abort (); + } + + #pragma omp target exit data map(delete: data1[:N], data2[:N]) + + int b = 8; + S s = { 4, N, data1 }; + auto f = s.merge_data_func (data2, b); + + if (f ()) abort (); + + #pragma omp target enter data map(to: data1[:N]) + if (f ()) abort (); + + #pragma omp target enter data map(to: data2[:N]) + if (!f ()) abort (); + + #pragma omp target exit data map(from: data1[:N], data2[:N]) + + for (int i = 0; i < N; i++) + { + if (data1[i] != 0xf) abort (); + if (data2[i] != 2) abort (); + } + + return 0; +} + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-2.C b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C new file mode 100644 index 00000000000..bdf2564cd04 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C @@ -0,0 +1,35 @@ +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14 -fdump-tree-gimple" } +#include + +#define N 10 +int main (void) +{ + int X, Y; + #pragma omp target map(from: X, Y) + { + int x = 0, y = 0; + + for (int i = 0; i < N; i++) + [&] (int v) { x += v; } (i); + + auto yinc = [&y] { y++; }; + for (int i = 0; i < N; i++) + yinc (); + + X = x; + Y = y; + } + + int Xs = 0; + for (int i = 0; i < N; i++) + Xs += i; + if (X != Xs) + abort (); + + if (Y != N) + abort (); +} + +/* Make sure lambda objects do NOT appear in target maps. */ +/* { dg-final { scan-tree-dump {(?n)#pragma omp target num_teams.* map\(from:Y \[len: [0-9]+\]\) map\(from:X \[len: [0-9]+\]\)$} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-1.C b/gcc/testsuite/g++.dg/gomp/target-this-1.C new file mode 100644 index 00000000000..de93a3e5e57 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-1.C @@ -0,0 +1,33 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } +extern "C" void abort (); + +struct S +{ + int a, b, c, d; + + int sum (void) + { + int val = 0; + val += a + b + this->c + this->d; + return val; + } + + int sum_offload (void) + { + int val = 0; + #pragma omp target map(val) + val += a + b + this->c + this->d; + return val; + } +}; + +int main (void) +{ + S s = { 1, 2, 3, 4 }; + if (s.sum () != s.sum_offload ()) + abort (); + return 0; +} + +/* { dg-final { scan-tree-dump {map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C new file mode 100644 index 00000000000..679c85a54dd --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C @@ -0,0 +1,49 @@ +// We use 'auto' without a function return type, so specify dialect here +// { dg-do compile } +// { dg-additional-options "-std=c++14 -fdump-tree-gimple" } + +extern "C" void abort (); + +struct T +{ + int x, y; + + auto sum_func (int n) + { + auto fn = [=](int m) -> int + { + int v; + v = (x + y) * n + m; + return v; + }; + return fn; + } + + auto sum_func_offload (int n) + { + auto fn = [=](int m) -> int + { + int v; + #pragma omp target map(from:v) + v = (x + y) * n + m; + return v; + }; + return fn; + } + +}; + +int main (void) +{ + T a = { 1, 2 }; + + auto s1 = a.sum_func (3); + auto s2 = a.sum_func_offload (3); + + if (s1 (1) != s2 (1)) + abort (); + + return 0; +} + +/* { dg-final { scan-tree-dump {map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C new file mode 100644 index 00000000000..08568f9284c --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C @@ -0,0 +1,105 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } +#include +#include +extern "C" void abort (); + +struct S +{ + int * ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + bool set_ptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr != NULL) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + } + + bool set_refptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr != NULL) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + S s = { ptr1, N, ptr2, N }; + + bool mapped; + int val = 123; + + mapped = s.set_ptr (val); + if (mapped) + abort (); + if (s.ptr != ptr1) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + mapped = s.set_refptr (val); + if (mapped) + abort (); + if (s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N]) + mapped = s.set_ptr (val); + + if (!mapped) + abort (); + if (s.set_refptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != val) + abort (); + + #pragma omp target data map(ptr2[:N]) + mapped = s.set_refptr (val); + + if (!mapped) + abort (); + if (s.set_ptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != val) + abort (); + + return 0; +} + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\) firstprivate\(n\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C new file mode 100644 index 00000000000..3b2d5811350 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C @@ -0,0 +1,107 @@ +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14 -fdump-tree-gimple" } +#include +#include + +struct T +{ + int *ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + auto set_ptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + }; + return fn; + } + + auto set_refptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + }; + return fn; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + T a = { ptr1, N, ptr2, N }; + + auto p1 = a.set_ptr_func (1); + auto r2 = a.set_refptr_func (2); + + if (p1 ()) + abort (); + if (r2 ()) + abort (); + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N], ptr2[:N]) + { + if (!p1 ()) + abort (); + if (!r2 ()) + abort (); + } + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 1) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 2) + abort (); + + return 0; +} + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-5.C b/gcc/testsuite/g++.dg/gomp/target-this-5.C new file mode 100644 index 00000000000..a9ac74bcf1f --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-5.C @@ -0,0 +1,34 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } +extern "C" void abort (); + +template +struct S +{ + T a, b, c, d; + + T sum (void) + { + T val = 0; + val += a + b + this->c + this->d; + return val; + } + + T sum_offload (void) + { + T val = 0; + #pragma omp target map(val) + val += a + b + this->c + this->d; + return val; + } +}; + +int main (void) +{ + S s = { 1, 2, 3, 4 }; + if (s.sum () != s.sum_offload ()) + abort (); + return 0; +} + +/* { dg-final { scan-tree-dump {map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/this-2.C b/gcc/testsuite/g++.dg/gomp/this-2.C index d03b8a0728e..b521a4faf5e 100644 --- a/gcc/testsuite/g++.dg/gomp/this-2.C +++ b/gcc/testsuite/g++.dg/gomp/this-2.C @@ -9,14 +9,14 @@ struct S void S::bar (int x) { - #pragma omp target map (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } ; - #pragma omp target map (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this[0], x) ; - #pragma omp target update to (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update to (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this[1], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target update to (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update to (this[0], x) + #pragma omp target update from (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update from (this[1], x) } template @@ -29,14 +29,14 @@ template void T::bar (int x) { - #pragma omp target map (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } ; - #pragma omp target map (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this[0], x) ; - #pragma omp target update to (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update to (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this[1], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target update to (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update to (this[0], x) + #pragma omp target update from (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update from (this[1], x) } template struct T<0>; diff --git a/gcc/testsuite/gcc.dg/gomp/target-3.c b/gcc/testsuite/gcc.dg/gomp/target-3.c new file mode 100644 index 00000000000..3e7921270c9 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/target-3.c @@ -0,0 +1,16 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp -fdump-tree-gimple" } */ + +struct S +{ + int a, b; +}; + +void foo (struct S *s) +{ + #pragma omp target map (alloc: s->a, s->b) + ; + #pragma omp target enter data map (alloc: s->a, s->b) +} + +/* { dg-final { scan-tree-dump-times "map\\(struct:\\*s \\\[len: 2\\\]\\) map\\(alloc:s->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:s->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index fde07dfd0e1..26e51d4b9f7 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -836,6 +836,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) { case GOMP_MAP_ALLOC: case GOMP_MAP_POINTER: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: pp_string (pp, "alloc"); break; case GOMP_MAP_IF_PRESENT: @@ -914,6 +915,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_ATTACH_DETACH: pp_string (pp, "attach_detach"); break; + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + pp_string (pp, "attach_zero_length_array_section"); + break; default: gcc_unreachable (); } @@ -932,6 +936,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_ALWAYS_POINTER: pp_string (pp, " [pointer assign, bias: "); break; + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + pp_string (pp, " [pointer assign, zero-length array section, bias: "); + break; case GOMP_MAP_TO_PSET: pp_string (pp, " [pointer set, len: "); break; @@ -939,6 +946,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_DETACH: case GOMP_MAP_FORCE_DETACH: case GOMP_MAP_ATTACH_DETACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: pp_string (pp, " [bias: "); break; default: diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 6e163b02560..4a4a14393d2 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -132,6 +132,11 @@ enum gomp_map_kind No refcount is bumped by this, and the store is done unconditionally. */ GOMP_MAP_ALWAYS_POINTER = (GOMP_MAP_FLAG_SPECIAL_2 | GOMP_MAP_FLAG_SPECIAL | 1), + /* Like GOMP_MAP_POINTER, but allow zero-length array section, i.e. set to + NULL if target is not mapped. */ + GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION + = (GOMP_MAP_FLAG_SPECIAL_2 + | GOMP_MAP_FLAG_SPECIAL | 2), /* Forced deallocation of zero length array section. */ GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION = (GOMP_MAP_FLAG_SPECIAL_2 @@ -152,6 +157,12 @@ enum gomp_map_kind GOMP_MAP_FORCE_DETACH = (GOMP_MAP_DEEP_COPY | GOMP_MAP_FLAG_FORCE | 1), + /* Like GOMP_MAP_ATTACH, but allow attaching to zero-length array sections + (i.e. set to NULL when array section is not mapped) Currently only used + by OpenMP. */ + GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION + = (GOMP_MAP_DEEP_COPY | 2), + /* Internal to GCC, not used in libgomp. */ /* Do not map, but pointer assign a pointer instead. */ GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1), @@ -175,7 +186,8 @@ enum gomp_map_kind ((X) == GOMP_MAP_ALWAYS_POINTER) #define GOMP_MAP_POINTER_P(X) \ - ((X) == GOMP_MAP_POINTER) + ((X) == GOMP_MAP_POINTER \ + || (X) == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION) #define GOMP_MAP_ALWAYS_TO_P(X) \ (((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM)) diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 8d25dc8e2a8..235e919f191 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1234,7 +1234,7 @@ extern uintptr_t gomp_map_val (struct target_mem_desc *, void **, size_t); extern void gomp_attach_pointer (struct gomp_device_descr *, struct goacc_asyncqueue *, splay_tree, splay_tree_key, uintptr_t, size_t, - struct gomp_coalesce_buf *); + struct gomp_coalesce_buf *, bool); extern void gomp_detach_pointer (struct gomp_device_descr *, struct goacc_asyncqueue *, splay_tree_key, uintptr_t, bool, struct gomp_coalesce_buf *); diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index c21508f3739..5ee8647cddf 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -937,7 +937,7 @@ acc_attach_async (void **hostaddr, int async) } gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr, - 0, NULL); + 0, NULL, false); gomp_mutex_unlock (&acc_dev->lock); } @@ -1141,7 +1141,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH) { gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, - (uintptr_t) h, s, NULL); + (uintptr_t) h, s, NULL, false); /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic reference counts ('n->refcount', 'n->dynamic_refcount'). */ } @@ -1159,7 +1159,8 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, splay_tree_key m = lookup_host (acc_dev, hostaddrs[j], sizeof (void *)); gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, m, - (uintptr_t) hostaddrs[j], sizes[j], NULL); + (uintptr_t) hostaddrs[j], sizes[j], NULL, + false); } bool processed = false; diff --git a/libgomp/target.c b/libgomp/target.c index bb09d501dd6..1ff6a503fe4 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -496,7 +496,8 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, struct gomp_coalesce_buf *cbuf, htab_t *refcount_set) { - assert (kind != GOMP_MAP_ATTACH); + assert (kind != GOMP_MAP_ATTACH + || kind != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); tgt_var->key = oldn; tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind); @@ -536,7 +537,8 @@ get_kind (bool short_mapkind, void *kinds, int idx) static void gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq, uintptr_t host_ptr, uintptr_t target_offset, uintptr_t bias, - struct gomp_coalesce_buf *cbuf) + struct gomp_coalesce_buf *cbuf, + bool allow_zero_length_array_sections) { struct gomp_device_descr *devicep = tgt->device_descr; struct splay_tree_s *mem_map = &devicep->mem_map; @@ -558,16 +560,24 @@ gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq, splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); if (n == NULL) { - gomp_mutex_unlock (&devicep->lock); - gomp_fatal ("Pointer target of array section wasn't mapped"); + if (allow_zero_length_array_sections) + cur_node.tgt_offset = 0; + else + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("Pointer target of array section wasn't mapped"); + } + } + else + { + cur_node.host_start -= n->host_start; + cur_node.tgt_offset + = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start; + /* At this point tgt_offset is target address of the + array section. Now subtract bias to get what we want + to initialize the pointer with. */ + cur_node.tgt_offset -= bias; } - cur_node.host_start -= n->host_start; - cur_node.tgt_offset - = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start; - /* At this point tgt_offset is target address of the - array section. Now subtract bias to get what we want - to initialize the pointer with. */ - cur_node.tgt_offset -= bias; gomp_copy_host2dev (devicep, aq, (void *) (tgt->tgt_start + target_offset), (void *) &cur_node.tgt_offset, sizeof (void *), cbuf); } @@ -638,7 +648,8 @@ attribute_hidden void gomp_attach_pointer (struct gomp_device_descr *devicep, struct goacc_asyncqueue *aq, splay_tree mem_map, splay_tree_key n, uintptr_t attach_to, size_t bias, - struct gomp_coalesce_buf *cbufp) + struct gomp_coalesce_buf *cbufp, + bool allow_zero_length_array_sections) { struct splay_tree_key_s s; size_t size, idx; @@ -690,11 +701,21 @@ gomp_attach_pointer (struct gomp_device_descr *devicep, if (!tn) { - gomp_mutex_unlock (&devicep->lock); - gomp_fatal ("pointer target not mapped for attach"); + if (allow_zero_length_array_sections) + { + /* When allowing attachment to zero-length array sections, we + allow attaching to NULL pointers when the target region is not + mapped. */ + data = 0; + } + else + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("pointer target not mapped for attach"); + } } - - data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start; + else + data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start; gomp_debug (1, "%s: attaching host %p, target %p (struct base %p) to %p\n", @@ -950,7 +971,9 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, has_firstprivate = true; continue; } - else if ((kind & typemask) == GOMP_MAP_ATTACH) + else if ((kind & typemask) == GOMP_MAP_ATTACH + || ((kind & typemask) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)) { tgt->list[i].key = NULL; has_firstprivate = true; @@ -1197,7 +1220,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, (uintptr_t) *(void **) hostaddrs[j], k->tgt_offset + ((uintptr_t) hostaddrs[j] - k->host_start), - sizes[j], cbufp); + sizes[j], cbufp, false); } } i = j - 1; @@ -1325,6 +1348,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, ++i; continue; case GOMP_MAP_ATTACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: { cur_node.host_start = (uintptr_t) hostaddrs[i]; cur_node.host_end = cur_node.host_start + sizeof (void *); @@ -1341,9 +1365,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, structured/dynamic reference counts ('n->refcount', 'n->dynamic_refcount'). */ + bool zlas + = ((kind & typemask) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); gomp_attach_pointer (devicep, aq, mem_map, n, (uintptr_t) hostaddrs[i], sizes[i], - cbufp); + cbufp, zlas); } else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0) { @@ -1453,9 +1480,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, k->host_end - k->host_start, cbufp); break; case GOMP_MAP_POINTER: - gomp_map_pointer (tgt, aq, - (uintptr_t) *(void **) k->host_start, - k->tgt_offset, sizes[i], cbufp); + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + gomp_map_pointer + (tgt, aq, (uintptr_t) *(void **) k->host_start, + k->tgt_offset, sizes[i], cbufp, + ((kind & typemask) + == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)); break; case GOMP_MAP_TO_PSET: gomp_copy_host2dev (devicep, aq, @@ -1496,7 +1526,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, k->tgt_offset + ((uintptr_t) hostaddrs[j] - k->host_start), - sizes[j], cbufp); + sizes[j], cbufp, false); } } i = j - 1; diff --git a/libgomp/testsuite/libgomp.c++/target-23.C b/libgomp/testsuite/libgomp.c++/target-23.C new file mode 100644 index 00000000000..d4f9ff3e983 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-23.C @@ -0,0 +1,34 @@ +extern "C" void abort (); + +struct S +{ + int *data; +}; + +int +main (void) +{ + #define SZ 10 + S *s = new S (); + s->data = new int[SZ]; + + for (int i = 0; i < SZ; i++) + s->data[i] = 0; + + #pragma omp target enter data map(to: s) + #pragma omp target enter data map(to: s->data[:SZ]) + #pragma omp target + { + for (int i = 0; i < SZ; i++) + s->data[i] = i; + } + #pragma omp target exit data map(from: s->data[:SZ]) + #pragma omp target exit data map(from: s) + + for (int i = 0; i < SZ; i++) + if (s->data[i] != i) + abort (); + + return 0; +} + diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-1.C b/libgomp/testsuite/libgomp.c++/target-lambda-1.C new file mode 100644 index 00000000000..06c6470b4ff --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-lambda-1.C @@ -0,0 +1,86 @@ +#include +#include + +template +void +omp_target_loop (int begin, int end, L loop) +{ + #pragma omp target teams distribute parallel for + for (int i = begin; i < end; i++) + loop (i); +} + +struct S +{ + int a, len; + int *ptr; + + auto merge_data_func (int *iptr, int &b) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + mapped = (ptr != NULL && iptr != NULL); + if (mapped) + { + for (int i = 0; i < len; i++) + ptr[i] += a + b + iptr[i]; + } + } + return mapped; + }; + return fn; + } +}; + +int x = 1; + +int main (void) +{ + const int N = 10; + int *data1 = new int[N]; + int *data2 = new int[N]; + memset (data1, 0xab, sizeof (int) * N); + memset (data1, 0xcd, sizeof (int) * N); + + int val = 1; + int &valref = val; + #pragma omp target enter data map(alloc: data1[:N], data2[:N]) + + omp_target_loop (0, N, [=](int i) { data1[i] = val; }); + omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; }); + + #pragma omp target update from(data1[:N], data2[:N]) + + for (int i = 0; i < N; i++) + { + if (data1[i] != 1) abort (); + if (data2[i] != 2) abort (); + } + + #pragma omp target exit data map(delete: data1[:N], data2[:N]) + + int b = 8; + S s = { 4, N, data1 }; + auto f = s.merge_data_func (data2, b); + + if (f ()) abort (); + + #pragma omp target enter data map(to: data1[:N]) + if (f ()) abort (); + + #pragma omp target enter data map(to: data2[:N]) + if (!f ()) abort (); + + #pragma omp target exit data map(from: data1[:N], data2[:N]) + + for (int i = 0; i < N; i++) + { + if (data1[i] != 0xf) abort (); + if (data2[i] != 2) abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-2.C b/libgomp/testsuite/libgomp.c++/target-lambda-2.C new file mode 100644 index 00000000000..1d3561ffbd7 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-lambda-2.C @@ -0,0 +1,30 @@ +#include + +#define N 10 +int main (void) +{ + int X, Y; + #pragma omp target map(from: X, Y) + { + int x = 0, y = 0; + + for (int i = 0; i < N; i++) + [&] (int v) { x += v; } (i); + + auto yinc = [&y] { y++; }; + for (int i = 0; i < N; i++) + yinc (); + + X = x; + Y = y; + } + + int Xs = 0; + for (int i = 0; i < N; i++) + Xs += i; + if (X != Xs) + abort (); + + if (Y != N) + abort (); +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-1.C b/libgomp/testsuite/libgomp.c++/target-this-1.C new file mode 100644 index 00000000000..a591ea4c564 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-1.C @@ -0,0 +1,29 @@ +extern "C" void abort (); + +struct S +{ + int a, b, c, d; + + int sum (void) + { + int val = 0; + val += a + b + this->c + this->d; + return val; + } + + int sum_offload (void) + { + int val = 0; + #pragma omp target map(val) + val += a + b + this->c + this->d; + return val; + } +}; + +int main (void) +{ + S s = { 1, 2, 3, 4 }; + if (s.sum () != s.sum_offload ()) + abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-2.C b/libgomp/testsuite/libgomp.c++/target-this-2.C new file mode 100644 index 00000000000..8119be8c2c5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-2.C @@ -0,0 +1,47 @@ + +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14" } + +extern "C" void abort (); + +struct T +{ + int x, y; + + auto sum_func (int n) + { + auto fn = [=](int m) -> int + { + int v; + v = (x + y) * n + m; + return v; + }; + return fn; + } + + auto sum_func_offload (int n) + { + auto fn = [=](int m) -> int + { + int v; + #pragma omp target map(from:v) + v = (x + y) * n + m; + return v; + }; + return fn; + } + +}; + +int main (void) +{ + T a = { 1, 2 }; + + auto s1 = a.sum_func (3); + auto s2 = a.sum_func_offload (3); + + if (s1 (1) != s2 (1)) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-3.C b/libgomp/testsuite/libgomp.c++/target-this-3.C new file mode 100644 index 00000000000..e15f69a1623 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-3.C @@ -0,0 +1,99 @@ +#include +#include +extern "C" void abort (); + +struct S +{ + int * ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + bool set_ptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr != NULL) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + } + + bool set_refptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr != NULL) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + S s = { ptr1, N, ptr2, N }; + + bool mapped; + int val = 123; + + mapped = s.set_ptr (val); + if (mapped) + abort (); + if (s.ptr != ptr1) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + mapped = s.set_refptr (val); + if (mapped) + abort (); + if (s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N]) + mapped = s.set_ptr (val); + + if (!mapped) + abort (); + if (s.set_refptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != val) + abort (); + + #pragma omp target data map(ptr2[:N]) + mapped = s.set_refptr (val); + + if (!mapped) + abort (); + if (s.set_ptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != val) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-4.C b/libgomp/testsuite/libgomp.c++/target-this-4.C new file mode 100644 index 00000000000..9f53677a240 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-4.C @@ -0,0 +1,104 @@ + +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14" } +#include +#include + +struct T +{ + int *ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + auto set_ptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + }; + return fn; + } + + auto set_refptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + }; + return fn; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + T a = { ptr1, N, ptr2, N }; + + auto p1 = a.set_ptr_func (1); + auto r2 = a.set_refptr_func (2); + + if (p1 ()) + abort (); + if (r2 ()) + abort (); + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N], ptr2[:N]) + { + if (!p1 ()) + abort (); + if (!r2 ()) + abort (); + } + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 1) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 2) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-5.C b/libgomp/testsuite/libgomp.c++/target-this-5.C new file mode 100644 index 00000000000..e71c566687d --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-5.C @@ -0,0 +1,30 @@ +extern "C" void abort (); + +template +struct S +{ + T a, b, c, d; + + T sum (void) + { + T val = 0; + val += a + b + this->c + this->d; + return val; + } + + T sum_offload (void) + { + T val = 0; + #pragma omp target map(val) + val += a + b + this->c + this->d; + return val; + } +}; + +int main (void) +{ + S s = { 1, 2, 3, 4 }; + if (s.sum () != s.sum_offload ()) + abort (); + return 0; +}