From patchwork Fri Dec 23 12:13:00 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1719124 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.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 ECDSA (P-384) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4NdmN71Fjqz1ydd for ; Fri, 23 Dec 2022 23:15:19 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 28CE03889E03 for ; Fri, 23 Dec 2022 12:15:17 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 4BE2E385E004; Fri, 23 Dec 2022 12:14:12 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 4BE2E385E004 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.96,268,1665475200"; d="scan'208";a="91397251" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 23 Dec 2022 04:14:11 -0800 IronPort-SDR: vyr0Kh4FIYiYWhmaW7IMTHacNgo8UvecdCICAAnc5KkB7+PSPjSok/+tcD1CfxuxIjU4w5Kty3 Rb6155HOu5frbN423onnCWWGwk9o4gam2vDp9LnM0lZDDxV5I8cdRSA2TKEVwCe2WNLyRrtzOA kJPVl6TbcStxVLalTLpc281d707J9EVKS1hMjr+lRYhrKuinZ8tpdqvTsFZKI93c6SVDFR5M5T 5C8MsAZnyZeuZBitlkQWyUCOxWHY6hJJkH20szJuEYgPKCTRc53ap1id1ZDDce01XZCk8pKftg kic= From: Julian Brown To: CC: , Tobias Burnus , "Jakub Jelinek" , Thomas Schwinge Subject: [PATCH v6 07/11] OpenMP: lvalue parsing for map/to/from clauses (C++) Date: Fri, 23 Dec 2022 04:13:00 -0800 Message-ID: <3bbe7d5add6431aac1fcf8d076c412907cfcd856.1671796516.git.julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-13.mgc.mentorg.com (139.181.222.13) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, RCVD_IN_MSPIKE_H2, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org 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+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" This is another iteration of the patch to support "lvalue" parsing (or "locator list item type" parsing), for several OpenMP clause types for C++. This version incorporates the "Use OMP_ARRAY_SECTION instead of TREE_LIST in C++ FE" patch posted previously and fixes a couple of parsing/semantic issues pointed out in the last round of review (thanks Jakub!). OpenMP array section nodes are now built via new functions in decl2.cc/semantics.cc and template handling has been improved. New diagnostics have been added for attempts to use C++23 multidimensional array index syntax with OMP array sections. Several new tests have been added, and some of the existing tests have been augmented with variants using templates. 2022-11-08 Julian Brown gcc/c-family/ * c-omp.cc (c_omp_split_clauses): Support OMP_ARRAY_SECTION. (c_omp_address_inspector::map_supported_p): Handle OMP_ARRAY_SECTION. (c_omp_address_inspector::get_origin): Avoid crash with missing type. (c_omp_address_inspector::maybe_unconvert_ref): Likewise. gcc/cp/ * constexpr.cc (potential_consant_expression_1): Handle OMP_ARRAY_SECTION. * cp-tree.h (grok_omp_array_section, build_omp_array_section): Add prototypes. * decl2.cc (grok_omp_array_section): New function. * error.cc (dump_expr): Handle OMP_ARRAY_SECTION. * parser.cc (cp_parser_new): Initialize parser->omp_array_section_p. (cp_parser_statement_expr): Disallow array sections. (cp_parser_postfix_open_square_expression): Support OMP_ARRAY_SECTION parsing. (cp_parser_parenthesized_expression_list, cp_parser_lambda_expression, cp_parser_braced_list): Disallow array sections. (cp_parser_omp_var_list_no_open): Remove ALLOW_DEREF parameter, add MAP_LVALUE in its place. Supported generalised lvalue parsing for OpenMP map, to and from clauses. Use OMP_ARRAY_SECTION code instead of TREE_LIST to represent OpenMP array sections. (cp_parser_omp_var_list): Remove ALLOW_DEREF parameter, add MAP_LVALUE. Pass to cp_parser_omp_var_list_no_open. (cp_parser_oacc_data_clause): Update call to cp_parser_omp_var_list. (cp_parser_omp_clause_map): Add sk_omp scope around cp_parser_omp_var_list_no_open call. * parser.h (cp_parser): Add omp_array_section_p field. * pt.cc (tsubst, tsubst_copy, tsubst_omp_clause_decl, tsubst_copy_and_build): Add OMP_ARRAY_SECTION support. * semantics.cc (handle_omp_array_sections_1, handle_omp_array_sections, cp_oacc_check_attachments, finish_omp_clauses): Use OMP_ARRAY_SECTION instead of TREE_LIST where appropriate. Handle more types of map expression. * typeck.cc (build_omp_array_section): New function. gcc/ * gimplify.cc (gimplify_expr): Ensure OMP_ARRAY_SECTION has been processed out before gimplification. * tree-pretty-print.cc (dump_generic_node): Support OMP_ARRAY_SECTION. * tree.def (OMP_ARRAY_SECTION): New tree code. gcc/testsuite/ * c-c++-common/gomp/map-6.c: Update expected output. * g++.dg/gomp/array-section-1.C: New test. * g++.dg/gomp/array-section-2.C: New test. * g++.dg/gomp/bad-array-section-1.C: New test. * g++.dg/gomp/bad-array-section-2.C: New test. * g++.dg/gomp/bad-array-section-3.C: New test. * g++.dg/gomp/bad-array-section-4.C: New test. * g++.dg/gomp/bad-array-section-5.C: New test. * g++.dg/gomp/bad-array-section-6.C: New test. * g++.dg/gomp/bad-array-section-7.C: New test. * g++.dg/gomp/bad-array-section-8.C: New test. * g++.dg/gomp/bad-array-section-9.C: New test. * g++.dg/gomp/bad-array-section-10.C: New test. * g++.dg/gomp/bad-array-section-11.C: New test. * g++.dg/gomp/has_device_addr-non-lvalue-1.C: New test. * g++.dg/gomp/pr67522.C: Update expected output. * g++.dg/gomp/ind-base-3.C: New test. * g++.dg/gomp/map-assignment-1.C: New test. * g++.dg/gomp/map-inc-1.C: New test. * g++.dg/gomp/map-lvalue-ref-1.C: New test. * g++.dg/gomp/map-ptrmem-1.C: New test. * g++.dg/gomp/map-ptrmem-2.C: New test. * g++.dg/gomp/map-static-cast-lvalue-1.C: New test. * g++.dg/gomp/map-ternary-1.C: New test. * g++.dg/gomp/member-array-2.C: New test. libgomp/ * testsuite/libgomp.c++/baseptrs-4.C: Remove commented-out cases that now work. * testsuite/libgomp.c++/baseptrs-6.C: New test. * testsuite/libgomp.c++/ind-base-1.C: New test. * testsuite/libgomp.c++/ind-base-2.C: New test. * testsuite/libgomp.c++/lvalue-tofrom-1.C: New test. * testsuite/libgomp.c++/lvalue-tofrom-2.C: New test. * testsuite/libgomp.c++/map-comma-1.C: New test. * testsuite/libgomp.c++/map-rvalue-ref-1.C: New test. * testsuite/libgomp.c++/struct-ref-1.C: New test. * testsuite/libgomp.c-c++-common/array-field-1.c: New test. * testsuite/libgomp.c-c++-common/array-of-struct-1.c: New test. * testsuite/libgomp.c-c++-common/array-of-struct-2.c: New test. --- gcc/c-family/c-common.h | 4 +- gcc/c-family/c-omp.cc | 23 +- gcc/cp/constexpr.cc | 1 + gcc/cp/cp-tree.h | 2 + gcc/cp/decl2.cc | 45 + gcc/cp/error.cc | 9 + gcc/cp/parser.cc | 209 +- gcc/cp/parser.h | 3 + gcc/cp/pt.cc | 49 + gcc/cp/semantics.cc | 69 +- gcc/cp/typeck.cc | 50 + gcc/gimplify.cc | 3 + gcc/testsuite/c-c++-common/gomp/map-6.c | 4 +- gcc/testsuite/g++.dg/gomp/array-section-1.C | 38 + gcc/testsuite/g++.dg/gomp/array-section-2.C | 63 + .../g++.dg/gomp/bad-array-section-1.C | 35 + .../g++.dg/gomp/bad-array-section-10.C | 35 + .../g++.dg/gomp/bad-array-section-11.C | 36 + .../g++.dg/gomp/bad-array-section-2.C | 33 + .../g++.dg/gomp/bad-array-section-3.C | 28 + .../g++.dg/gomp/bad-array-section-4.C | 50 + .../g++.dg/gomp/bad-array-section-5.C | 50 + .../g++.dg/gomp/bad-array-section-6.C | 24 + .../g++.dg/gomp/bad-array-section-7.C | 36 + .../g++.dg/gomp/bad-array-section-8.C | 53 + .../g++.dg/gomp/bad-array-section-9.C | 39 + .../gomp/has_device_addr-non-lvalue-1.C | 36 + gcc/testsuite/g++.dg/gomp/ind-base-3.C | 37 + gcc/testsuite/g++.dg/gomp/map-assignment-1.C | 12 + gcc/testsuite/g++.dg/gomp/map-inc-1.C | 10 + gcc/testsuite/g++.dg/gomp/map-lvalue-ref-1.C | 19 + gcc/testsuite/g++.dg/gomp/map-ptrmem-1.C | 37 + gcc/testsuite/g++.dg/gomp/map-ptrmem-2.C | 40 + .../g++.dg/gomp/map-static-cast-lvalue-1.C | 17 + gcc/testsuite/g++.dg/gomp/map-ternary-1.C | 20 + gcc/testsuite/g++.dg/gomp/member-array-2.C | 91 + gcc/testsuite/g++.dg/gomp/pr67522.C | 2 +- gcc/tree-pretty-print.cc | 14 + gcc/tree.def | 3 + libgomp/testsuite/libgomp.c++/baseptrs-4.C | 26 +- libgomp/testsuite/libgomp.c++/baseptrs-6.C | 3199 +++++++++++++++++ libgomp/testsuite/libgomp.c++/ind-base-1.C | 162 + libgomp/testsuite/libgomp.c++/ind-base-2.C | 93 + .../testsuite/libgomp.c++/lvalue-tofrom-1.C | 75 + .../testsuite/libgomp.c++/lvalue-tofrom-2.C | 71 + libgomp/testsuite/libgomp.c++/map-comma-1.C | 15 + .../testsuite/libgomp.c++/map-rvalue-ref-1.C | 22 + libgomp/testsuite/libgomp.c++/struct-ref-1.C | 97 + .../libgomp.c-c++-common/array-field-1.c | 35 + .../libgomp.c-c++-common/array-of-struct-1.c | 65 + .../libgomp.c-c++-common/array-of-struct-2.c | 65 + 51 files changed, 5187 insertions(+), 67 deletions(-) create mode 100644 gcc/testsuite/g++.dg/gomp/array-section-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/array-section-2.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-section-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-section-10.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-section-11.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-section-2.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-section-3.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-section-4.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-section-5.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-section-6.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-section-7.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-section-8.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-section-9.C create mode 100644 gcc/testsuite/g++.dg/gomp/has_device_addr-non-lvalue-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/ind-base-3.C create mode 100644 gcc/testsuite/g++.dg/gomp/map-assignment-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/map-inc-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/map-lvalue-ref-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/map-ptrmem-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/map-ptrmem-2.C create mode 100644 gcc/testsuite/g++.dg/gomp/map-static-cast-lvalue-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/map-ternary-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/member-array-2.C create mode 100644 libgomp/testsuite/libgomp.c++/baseptrs-6.C create mode 100644 libgomp/testsuite/libgomp.c++/ind-base-1.C create mode 100644 libgomp/testsuite/libgomp.c++/ind-base-2.C create mode 100644 libgomp/testsuite/libgomp.c++/lvalue-tofrom-1.C create mode 100644 libgomp/testsuite/libgomp.c++/lvalue-tofrom-2.C create mode 100644 libgomp/testsuite/libgomp.c++/map-comma-1.C create mode 100644 libgomp/testsuite/libgomp.c++/map-rvalue-ref-1.C create mode 100644 libgomp/testsuite/libgomp.c++/struct-ref-1.C create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-field-1.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-of-struct-1.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-of-struct-2.c diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index 06674e769bd4..510ee2d24e0c 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1340,8 +1340,8 @@ public: bool map_supported_p (); - static tree get_origin (tree); - static tree maybe_unconvert_ref (tree); + tree get_origin (tree); + tree maybe_unconvert_ref (tree); bool maybe_zero_length_array_section (tree); diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index 74c01d8f2a52..2de3fc2be759 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -2680,6 +2680,9 @@ c_omp_split_clauses (location_t loc, enum tree_code code, } else if (TREE_CODE (OMP_CLAUSE_DECL (c)) == TREE_LIST) { + /* TODO: This can go away once we transition all uses of + TREE_LIST for representing OMP array sections to + OMP_ARRAY_SECTION. */ tree t; for (t = OMP_CLAUSE_DECL (c); TREE_CODE (t) == TREE_LIST; t = TREE_CHAIN (t)) @@ -2688,6 +2691,17 @@ c_omp_split_clauses (location_t loc, enum tree_code code, bitmap_clear_bit (&allocate_head, DECL_UID (t)); break; } + else if (TREE_CODE (OMP_CLAUSE_DECL (c)) == OMP_ARRAY_SECTION) + { + tree t; + for (t = OMP_CLAUSE_DECL (c); + TREE_CODE (t) == OMP_ARRAY_SECTION; + t = TREE_OPERAND (t, 0)) + ; + if (DECL_P (t)) + bitmap_clear_bit (&allocate_head, DECL_UID (t)); + break; + } /* FALLTHRU */ case OMP_CLAUSE_PRIVATE: case OMP_CLAUSE_FIRSTPRIVATE: @@ -3229,6 +3243,7 @@ c_omp_address_inspector::map_supported_p () || TREE_CODE (t) == SAVE_EXPR || TREE_CODE (t) == POINTER_PLUS_EXPR || TREE_CODE (t) == NON_LVALUE_EXPR + || TREE_CODE (t) == OMP_ARRAY_SECTION || TREE_CODE (t) == NOP_EXPR) if (TREE_CODE (t) == COMPOUND_EXPR) t = TREE_OPERAND (t, 1); @@ -3258,7 +3273,8 @@ c_omp_address_inspector::get_origin (tree t) else if (TREE_CODE (t) == POINTER_PLUS_EXPR || TREE_CODE (t) == SAVE_EXPR) t = TREE_OPERAND (t, 0); - else if (TREE_CODE (t) == INDIRECT_REF + else if (!processing_template_decl_p () + && TREE_CODE (t) == INDIRECT_REF && TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == REFERENCE_TYPE) t = TREE_OPERAND (t, 0); else @@ -3275,7 +3291,10 @@ c_omp_address_inspector::get_origin (tree t) tree c_omp_address_inspector::maybe_unconvert_ref (tree t) { - if (TREE_CODE (t) == INDIRECT_REF + /* Be careful not to dereference the type if we're processing a + template decl, else it might be NULL. */ + if (!processing_template_decl_p () + && TREE_CODE (t) == INDIRECT_REF && TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == REFERENCE_TYPE) return TREE_OPERAND (t, 0); diff --git a/gcc/cp/constexpr.cc b/gcc/cp/constexpr.cc index d99c49bdbe28..6994d2b25e04 100644 --- a/gcc/cp/constexpr.cc +++ b/gcc/cp/constexpr.cc @@ -9488,6 +9488,7 @@ potential_constant_expression_1 (tree t, bool want_rval, bool strict, bool now, case OACC_ENTER_DATA: case OACC_EXIT_DATA: case OACC_UPDATE: + case OMP_ARRAY_SECTION: /* GCC internal stuff. */ case VA_ARG_EXPR: case TRANSACTION_EXPR: diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index 541d760492d0..8157d8fd0e5f 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -6956,6 +6956,7 @@ extern void grokclassfn (tree, tree, enum overload_flags); extern tree grok_array_decl (location_t, tree, tree, vec **, tsubst_flags_t); +extern tree grok_omp_array_section (location_t, tree, tree, tree); extern tree delete_sanity (location_t, tree, tree, bool, int, tsubst_flags_t); extern tree check_classfn (tree, tree, tree); @@ -8050,6 +8051,7 @@ inline tree build_x_binary_op (const op_location_t &loc, } extern tree build_x_array_ref (location_t, tree, tree, tsubst_flags_t); +extern tree build_omp_array_section (location_t, tree, tree, tree); extern tree build_x_unary_op (location_t, enum tree_code, cp_expr, tree, tsubst_flags_t); diff --git a/gcc/cp/decl2.cc b/gcc/cp/decl2.cc index f95529a5c9a7..dac74198b9f1 100644 --- a/gcc/cp/decl2.cc +++ b/gcc/cp/decl2.cc @@ -591,6 +591,51 @@ grok_array_decl (location_t loc, tree array_expr, tree index_exp, return expr; } +/* Build an OMP_ARRAY_SECTION expression, handling usage in template + definitions, etc. */ + +tree +grok_omp_array_section (location_t loc, tree array_expr, tree index, + tree length) +{ + tree orig_array_expr = array_expr; + tree orig_index = index; + tree orig_length = length; + + if (error_operand_p (array_expr) + || error_operand_p (index) + || error_operand_p (length)) + return error_mark_node; + + if (processing_template_decl) + { + if (type_dependent_expression_p (array_expr) + || type_dependent_expression_p (index) + || type_dependent_expression_p (length)) + return build_min_nt_loc (loc, OMP_ARRAY_SECTION, array_expr, index, + length); + array_expr = build_non_dependent_expr (array_expr); + if (index) + index = build_non_dependent_expr (index); + if (length) + length = build_non_dependent_expr (length); + } + + index = fold_non_dependent_expr (index); + length = fold_non_dependent_expr (length); + + /* NOTE: We can pass through invalidly-typed index/length fields + here (e.g. if the user tries to use a floating-point index/length). + This is diagnosed later in semantics.cc:handle_omp_array_sections_1. */ + + tree expr = build_omp_array_section (loc, array_expr, index, length); + + if (processing_template_decl) + expr = build_min_non_dep (OMP_ARRAY_SECTION, expr, orig_array_expr, + orig_index, orig_length); + return expr; +} + /* Given the cast expression EXP, checking out its validity. Either return an error_mark_node if there was an unavoidable error, return a cast to void for trying to delete a pointer w/ the value 0, or return the diff --git a/gcc/cp/error.cc b/gcc/cp/error.cc index e7f60335cc06..6262028da272 100644 --- a/gcc/cp/error.cc +++ b/gcc/cp/error.cc @@ -2497,6 +2497,15 @@ dump_expr (cxx_pretty_printer *pp, tree t, int flags) pp_cxx_right_bracket (pp); break; + case OMP_ARRAY_SECTION: + dump_expr (pp, TREE_OPERAND (t, 0), flags); + pp_cxx_left_bracket (pp); + dump_expr (pp, TREE_OPERAND (t, 1), flags); + pp_colon (pp); + dump_expr (pp, TREE_OPERAND (t, 2), flags); + pp_cxx_right_bracket (pp); + break; + case UNARY_PLUS_EXPR: dump_unary_op (pp, "+", t, flags); break; diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 3905d561152b..1a45dad6b8e4 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -4344,6 +4344,9 @@ cp_parser_new (cp_lexer *lexer) parser->omp_declare_simd = NULL; parser->oacc_routine = NULL; + /* Disallow OpenMP array sections in expressions. */ + parser->omp_array_section_p = false; + /* Not declaring an implicit function template. */ parser->auto_is_implicit_function_template_parm_p = false; parser->fully_implicit_function_template_p = false; @@ -5279,6 +5282,7 @@ static cp_expr cp_parser_statement_expr (cp_parser *parser) { cp_token_position start = cp_parser_start_tentative_firewall (parser); + auto oas = make_temp_override (parser->omp_array_section_p, false); /* Consume the '('. */ location_t start_loc = cp_lexer_peek_token (parser->lexer)->location; @@ -8096,6 +8100,7 @@ cp_parser_postfix_open_square_expression (cp_parser *parser, releasing_vec expression_list = NULL; location_t loc = cp_lexer_peek_token (parser->lexer)->location; bool saved_greater_than_is_operator_p; + bool saved_colon_corrects_to_scope_p; /* Consume the `[' token. */ cp_lexer_consume_token (parser->lexer); @@ -8103,6 +8108,10 @@ cp_parser_postfix_open_square_expression (cp_parser *parser, saved_greater_than_is_operator_p = parser->greater_than_is_operator_p; parser->greater_than_is_operator_p = true; + saved_colon_corrects_to_scope_p = parser->colon_corrects_to_scope_p; + if (parser->omp_array_section_p) + parser->colon_corrects_to_scope_p = false; + /* Parse the index expression. */ /* ??? For offsetof, there is a question of what to allow here. If offsetof is not being used in an integral constant expression context, @@ -8113,7 +8122,8 @@ cp_parser_postfix_open_square_expression (cp_parser *parser, constant expressions here. */ if (for_offsetof) index = cp_parser_constant_expression (parser); - else + else if (!parser->omp_array_section_p + || cp_lexer_next_token_is_not (parser->lexer, CPP_COLON)) { if (cxx_dialect >= cxx23 && cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_SQUARE)) @@ -8170,6 +8180,68 @@ cp_parser_postfix_open_square_expression (cp_parser *parser, parser->greater_than_is_operator_p = saved_greater_than_is_operator_p; + if (cxx_dialect >= cxx23 + && parser->omp_array_section_p + && expression_list.get () != NULL + && vec_safe_length (expression_list) > 1) + { + error_at (loc, "cannot use multidimensional subscript in OpenMP array " + "section"); + index = error_mark_node; + } + if (parser->omp_array_section_p + && cp_lexer_next_token_is (parser->lexer, CPP_COLON)) + { + cp_lexer_consume_token (parser->lexer); + tree length = NULL_TREE; + if (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_SQUARE)) + { + if (cxx_dialect >= cxx23) + { + cp_expr expr + = cp_parser_parenthesized_expression_list_elt (parser, + /*cast_p=*/ + false, + /*allow_exp_p=*/ + true, + /*non_cst_p=*/ + NULL); + + if (expr == error_mark_node) + length = error_mark_node; + else + length = expr.get_value (); + + if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)) + { + error_at (loc, "cannot use multidimensional subscript in " + "OpenMP array section"); + length = error_mark_node; + } + } + else + length + = cp_parser_expression (parser, NULL, /*cast_p=*/false, + /*decltype_p=*/false, + /*warn_comma_p=*/warn_comma_subscript); + } + + parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p; + + if (index == error_mark_node || length == error_mark_node) + { + cp_parser_skip_to_closing_square_bracket (parser); + return error_mark_node; + } + else + cp_parser_require (parser, CPP_CLOSE_SQUARE, RT_CLOSE_SQUARE); + + return grok_omp_array_section (input_location, postfix_expression, index, + length); + } + + parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p; + /* Look for the closing `]'. */ cp_parser_require (parser, CPP_CLOSE_SQUARE, RT_CLOSE_SQUARE); @@ -8498,6 +8570,7 @@ cp_parser_parenthesized_expression_list (cp_parser* parser, { vec *expression_list; bool saved_greater_than_is_operator_p; + bool saved_omp_array_section_p; /* Assume all the expressions will be constant. */ if (non_constant_p) @@ -8515,6 +8588,9 @@ cp_parser_parenthesized_expression_list (cp_parser* parser, = parser->greater_than_is_operator_p; parser->greater_than_is_operator_p = true; + saved_omp_array_section_p = parser->omp_array_section_p; + parser->omp_array_section_p = false; + cp_expr expr (NULL_TREE); /* Consume expressions until there are no more. */ @@ -8579,12 +8655,14 @@ cp_parser_parenthesized_expression_list (cp_parser* parser, { parser->greater_than_is_operator_p = saved_greater_than_is_operator_p; + parser->omp_array_section_p = saved_omp_array_section_p; return NULL; } } parser->greater_than_is_operator_p = saved_greater_than_is_operator_p; + parser->omp_array_section_p = saved_omp_array_section_p; return expression_list; } @@ -11069,6 +11147,7 @@ cp_parser_lambda_expression (cp_parser* parser) cp_binding_level* implicit_template_scope = parser->implicit_template_scope; bool auto_is_implicit_function_template_parm_p = parser->auto_is_implicit_function_template_parm_p; + bool saved_omp_array_section_p = parser->omp_array_section_p; parser->num_template_parameter_lists = 0; parser->in_statement = 0; @@ -11077,6 +11156,7 @@ cp_parser_lambda_expression (cp_parser* parser) parser->implicit_template_parms = 0; parser->implicit_template_scope = 0; parser->auto_is_implicit_function_template_parm_p = false; + parser->omp_array_section_p = false; /* The body of a lambda in a discarded statement is not discarded. */ bool discarded = in_discarded_stmt; @@ -11127,6 +11207,7 @@ cp_parser_lambda_expression (cp_parser* parser) parser->implicit_template_scope = implicit_template_scope; parser->auto_is_implicit_function_template_parm_p = auto_is_implicit_function_template_parm_p; + parser->omp_array_section_p = saved_omp_array_section_p; } /* This field is only used during parsing of the lambda. */ @@ -25460,6 +25541,7 @@ cp_parser_braced_list (cp_parser* parser, bool* non_constant_p) { tree initializer; location_t start_loc = cp_lexer_peek_token (parser->lexer)->location; + auto oas = make_temp_override (parser->omp_array_section_p, false); /* Consume the `{' token. */ matching_braces braces; @@ -37324,7 +37406,7 @@ struct omp_dim static tree cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, tree list, bool *colon, - bool allow_deref = false) + bool map_lvalue = false) { auto_vec dims; bool array_section_p; @@ -37341,6 +37423,105 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, if (kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY) cp_parser_parse_tentatively (parser); + /* This condition doesn't include OMP_CLAUSE_DEPEND or + OMP_CLAUSE_AFFINITY since lvalue ("locator list") parsing for those is + handled further down the function. */ + else if (map_lvalue + && (kind == OMP_CLAUSE_MAP + || kind == OMP_CLAUSE_TO + || kind == OMP_CLAUSE_FROM)) + { + auto s = make_temp_override (parser->omp_array_section_p, true); + token = cp_lexer_peek_token (parser->lexer); + location_t loc = token->location; + decl = cp_parser_assignment_expression (parser); + + /* This code rewrites a parsed expression containing various tree + codes used to represent array accesses into a more uniform nest of + OMP_ARRAY_SECTION nodes before it is processed by + semantics.cc:handle_omp_array_sections_1. It might be more + efficient to move this logic to that function instead, analysing + the parsed expression directly rather than this preprocessed + form. */ + dims.truncate (0); + if (TREE_CODE (decl) == OMP_ARRAY_SECTION) + { + while (TREE_CODE (decl) == OMP_ARRAY_SECTION) + { + tree low_bound = TREE_OPERAND (decl, 1); + tree length = TREE_OPERAND (decl, 2); + dims.safe_push (omp_dim (low_bound, length, loc, false)); + decl = TREE_OPERAND (decl, 0); + } + + while (TREE_CODE (decl) == ARRAY_REF + || TREE_CODE (decl) == INDIRECT_REF + || TREE_CODE (decl) == COMPOUND_EXPR) + { + if (REFERENCE_REF_P (decl)) + break; + + if (TREE_CODE (decl) == COMPOUND_EXPR) + { + decl = TREE_OPERAND (decl, 1); + STRIP_NOPS (decl); + } + else if (TREE_CODE (decl) == INDIRECT_REF) + { + dims.safe_push (omp_dim (integer_zero_node, + integer_one_node, loc, true)); + decl = TREE_OPERAND (decl, 0); + } + else /* ARRAY_REF. */ + { + tree index = TREE_OPERAND (decl, 1); + dims.safe_push (omp_dim (index, integer_one_node, loc, + true)); + decl = TREE_OPERAND (decl, 0); + } + } + + /* Bare references have their own special handling, so remove + the explicit dereference added by convert_from_reference. */ + if (REFERENCE_REF_P (decl)) + decl = TREE_OPERAND (decl, 0); + + for (int i = dims.length () - 1; i >= 0; i--) + decl = grok_omp_array_section (loc, decl, dims[i].low_bound, + dims[i].length); + } + else if (TREE_CODE (decl) == INDIRECT_REF) + { + bool ref_p = REFERENCE_REF_P (decl); + + /* Turn *foo into foo[0:1]. */ + decl = TREE_OPERAND (decl, 0); + STRIP_NOPS (decl); + + /* If we have "*foo" and + - it's an indirection of a reference, "unconvert" it, i.e. + strip the indirection (to just "foo"). + - it's an indirection of a pointer, turn it into + "foo[0:1]". */ + if (!ref_p) + decl = grok_omp_array_section (loc, decl, integer_zero_node, + integer_one_node); + } + else if (TREE_CODE (decl) == ARRAY_REF) + { + tree idx = TREE_OPERAND (decl, 1); + + decl = TREE_OPERAND (decl, 0); + STRIP_NOPS (decl); + + decl = grok_omp_array_section (loc, decl, idx, integer_one_node); + } + else if (TREE_CODE (decl) == NON_LVALUE_EXPR + || CONVERT_EXPR_P (decl)) + decl = TREE_OPERAND (decl, 0); + + goto build_clause; + } token = cp_lexer_peek_token (parser->lexer); if (kind != 0 && cp_parser_is_keyword (token, RID_THIS)) @@ -37419,8 +37600,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, case OMP_CLAUSE_TO: start_component_ref: while (cp_lexer_next_token_is (parser->lexer, CPP_DOT) - || (allow_deref - && cp_lexer_next_token_is (parser->lexer, CPP_DEREF))) + || cp_lexer_next_token_is (parser->lexer, CPP_DEREF)) { cpp_ttype ttype = cp_lexer_next_token_is (parser->lexer, CPP_DOT) @@ -37506,9 +37686,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, || kind == OMP_CLAUSE_TO) && !array_section_p && (cp_lexer_next_token_is (parser->lexer, CPP_DOT) - || (allow_deref - && cp_lexer_next_token_is (parser->lexer, - CPP_DEREF)))) + || cp_lexer_next_token_is (parser->lexer, CPP_DEREF))) { for (unsigned i = 0; i < dims.length (); i++) { @@ -37520,8 +37698,9 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, } else for (unsigned i = 0; i < dims.length (); i++) - decl = tree_cons (dims[i].low_bound, dims[i].length, decl); - + decl = build_omp_array_section (input_location, decl, + dims[i].low_bound, + dims[i].length); break; default: break; @@ -37542,6 +37721,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, cp_parser_parse_definitely (parser); } + build_clause: tree u = build_omp_clause (token->location, kind); OMP_CLAUSE_DECL (u) = decl; OMP_CLAUSE_CHAIN (u) = list; @@ -37591,11 +37771,11 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, static tree cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list, - bool allow_deref = false) + bool map_lvalue = false) { if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) return cp_parser_omp_var_list_no_open (parser, kind, list, NULL, - allow_deref); + map_lvalue); return list; } @@ -37662,7 +37842,7 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, gcc_unreachable (); } tree nl, c; - nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true); + nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, false); for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_SET_MAP_KIND (c, kind); @@ -40486,8 +40666,13 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) cp_lexer_consume_token (parser->lexer); } + /* We introduce a scope here so that errors parsing e.g. "always", "close" + tokens do not propagate to later directives that might use them + legally. */ + begin_scope (sk_omp, NULL); nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list, NULL, true); + finish_scope (); for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_SET_MAP_KIND (c, kind); diff --git a/gcc/cp/parser.h b/gcc/cp/parser.h index 5737146dd424..bb8d7b31db82 100644 --- a/gcc/cp/parser.h +++ b/gcc/cp/parser.h @@ -407,6 +407,9 @@ struct GTY(()) cp_parser { /* TRUE if omp::directive or omp::sequence attributes may not appear. */ bool omp_attrs_forbidden_p; + /* TRUE if an OpenMP array section is allowed. */ + bool omp_array_section_p; + /* Tracks the function's template parameter list when declaring a function using generic type parameters. This is either a new chain in the case of a fully implicit function template or an extension of the function's existing diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index 63772bd64e73..6cf3ed8d63ad 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -16579,6 +16579,7 @@ tsubst (tree t, tree args, tsubst_flags_t complain, tree in_decl) case CALL_EXPR: case ARRAY_REF: case SCOPE_REF: + case OMP_ARRAY_SECTION: /* We should use one of the expression tsubsts for these codes. */ gcc_unreachable (); @@ -17584,6 +17585,17 @@ tsubst_copy (tree t, tree args, tsubst_flags_t complain, tree in_decl) return build_nt (ARRAY_REF, op0, op1, NULL_TREE, NULL_TREE); } + case OMP_ARRAY_SECTION: + { + tree op0 = tsubst_copy (TREE_OPERAND (t, 0), args, complain, in_decl); + tree op1 = NULL_TREE, op2 = NULL_TREE; + if (TREE_OPERAND (t, 1)) + op1 = tsubst_copy (TREE_OPERAND (t, 1), args, complain, in_decl); + if (TREE_OPERAND (t, 2)) + op2 = tsubst_copy (TREE_OPERAND (t, 2), args, complain, in_decl); + return build_nt (OMP_ARRAY_SECTION, op0, op1, op2); + } + case CALL_EXPR: { int n = VL_EXP_OPERAND_LENGTH (t); @@ -17869,6 +17881,22 @@ tsubst_omp_clause_decl (tree decl, tree args, tsubst_flags_t complain, = OMP_CLAUSE_DOACROSS_SINK_NEGATIVE (decl); return ret; } + else if (TREE_CODE (decl) == OMP_ARRAY_SECTION) + { + tree low_bound + = tsubst_expr (TREE_OPERAND (decl, 1), args, complain, in_decl); + tree length = tsubst_expr (TREE_OPERAND (decl, 2), args, complain, + in_decl); + tree base = tsubst_omp_clause_decl (TREE_OPERAND (decl, 0), args, + complain, in_decl, NULL); + if (TREE_OPERAND (decl, 0) == base + && TREE_OPERAND (decl, 1) == low_bound + && TREE_OPERAND (decl, 2) == length) + return decl; + tree ret = build3 (OMP_ARRAY_SECTION, TREE_TYPE (base), base, low_bound, + length); + return ret; + } tree ret = tsubst_expr (decl, args, complain, in_decl); /* Undo convert_from_reference tsubst_expr could have called. */ if (decl @@ -20612,6 +20640,27 @@ tsubst_copy_and_build (tree t, RECUR (TREE_OPERAND (t, 1)), complain|decltype_flag)); + case OMP_ARRAY_SECTION: + { + tree op0 = RECUR (TREE_OPERAND (t, 0)); + tree op1 = NULL_TREE, op2 = NULL_TREE; + if (op0 == error_mark_node) + RETURN (error_mark_node); + if (TREE_OPERAND (t, 1)) + { + op1 = RECUR (TREE_OPERAND (t, 1)); + if (op1 == error_mark_node) + RETURN (error_mark_node); + } + if (TREE_OPERAND (t, 2)) + { + op2 = RECUR (TREE_OPERAND (t, 2)); + if (op2 == error_mark_node) + RETURN (error_mark_node); + } + RETURN (build_omp_array_section (EXPR_LOCATION (t), op0, op1, op2)); + } + case SIZEOF_EXPR: if (PACK_EXPANSION_P (TREE_OPERAND (t, 0)) || ARGUMENT_PACK_P (TREE_OPERAND (t, 0))) diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index f8f8299bbf64..f43816dafff0 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -5168,7 +5168,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, { tree ret, low_bound, length, type; bool openacc = (ort & C_ORT_ACC) != 0; - if (TREE_CODE (t) != TREE_LIST) + if (TREE_CODE (t) != OMP_ARRAY_SECTION) { if (error_operand_p (t)) return error_mark_node; @@ -5190,7 +5190,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, ret = t_refto; 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) + else if (!VAR_P (t) + && (openacc || !EXPR_P (t)) + && TREE_CODE (t) != PARM_DECL) { if (processing_template_decl && TREE_CODE (t) != OVERLOAD) return NULL_TREE; @@ -5223,16 +5225,16 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TASK_REDUCTION) - && TREE_CODE (TREE_CHAIN (t)) == FIELD_DECL) - TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t), false); - ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types, + && TREE_CODE (TREE_OPERAND (t, 0)) == FIELD_DECL) + TREE_OPERAND (t, 0) = omp_privatize_field (TREE_OPERAND (t, 0), false); + ret = handle_omp_array_sections_1 (c, TREE_OPERAND (t, 0), types, maybe_zero_len, first_non_one, ort); if (ret == error_mark_node || ret == NULL_TREE) return ret; type = TREE_TYPE (ret); - low_bound = TREE_PURPOSE (t); - length = TREE_VALUE (t); + low_bound = TREE_OPERAND (t, 1); + length = TREE_OPERAND (t, 2); if ((low_bound && type_dependent_expression_p (low_bound)) || (length && type_dependent_expression_p (length))) return NULL_TREE; @@ -5438,7 +5440,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, tree lb = cp_save_expr (low_bound); if (lb != low_bound) { - TREE_PURPOSE (t) = lb; + TREE_OPERAND (t, 1) = lb; low_bound = lb; } } @@ -5469,14 +5471,14 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, array-section-subscript, the array section could be non-contiguous. */ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND - && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST) + && TREE_CODE (TREE_OPERAND (t, 0)) == OMP_ARRAY_SECTION) { /* If any prior dimension has a non-one length, then deem this array section as non-contiguous. */ - for (tree d = TREE_CHAIN (t); TREE_CODE (d) == TREE_LIST; - d = TREE_CHAIN (d)) + for (tree d = TREE_OPERAND (t, 0); TREE_CODE (d) == OMP_ARRAY_SECTION; + d = TREE_OPERAND (d, 0)) { - tree d_length = TREE_VALUE (d); + tree d_length = TREE_OPERAND (d, 2); if (d_length == NULL_TREE || !integer_onep (d_length)) { error_at (OMP_CLAUSE_LOCATION (c), @@ -5499,7 +5501,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, tree lb = cp_save_expr (low_bound); if (lb != low_bound) { - TREE_PURPOSE (t) = lb; + TREE_OPERAND (t, 1) = lb; low_bound = lb; } /* Temporarily disable -fstrong-eval-order for array reductions. @@ -5577,10 +5579,12 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort) return false; for (i = num, t = OMP_CLAUSE_DECL (c); i > 0; - t = TREE_CHAIN (t)) + t = TREE_OPERAND (t, 0)) { - tree low_bound = TREE_PURPOSE (t); - tree length = TREE_VALUE (t); + gcc_assert (TREE_CODE (t) == OMP_ARRAY_SECTION); + + tree low_bound = TREE_OPERAND (t, 1); + tree length = TREE_OPERAND (t, 2); i--; if (low_bound @@ -6690,8 +6694,8 @@ cp_oacc_check_attachments (tree c) tree t = OMP_CLAUSE_DECL (c); tree type; - while (TREE_CODE (t) == TREE_LIST) - t = TREE_CHAIN (t); + while (TREE_CODE (t) == OMP_ARRAY_SECTION) + t = TREE_OPERAND (t, 0); type = TREE_TYPE (t); @@ -6798,7 +6802,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_TASK_REDUCTION: field_ok = ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP); t = OMP_CLAUSE_DECL (c); - if (TREE_CODE (t) == TREE_LIST) + if (TREE_CODE (t) == OMP_ARRAY_SECTION) { if (handle_omp_array_sections (c, ort)) { @@ -6814,10 +6818,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; break; } - if (TREE_CODE (t) == TREE_LIST) + if (TREE_CODE (t) == OMP_ARRAY_SECTION) { - while (TREE_CODE (t) == TREE_LIST) - t = TREE_CHAIN (t); + while (TREE_CODE (t) == OMP_ARRAY_SECTION) + t = TREE_OPERAND (t, 0); } else { @@ -7838,7 +7842,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else last_iterators = NULL_TREE; - if (TREE_CODE (t) == TREE_LIST) + if (TREE_CODE (t) == OMP_ARRAY_SECTION) { if (handle_omp_array_sections (c, ort)) remove = true; @@ -7998,7 +8002,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) auto_vec addr_tokens; t = OMP_CLAUSE_DECL (c); - if (TREE_CODE (t) == TREE_LIST) + if (TREE_CODE (t) == OMP_ARRAY_SECTION) { grp_start_p = pc; grp_sentinel = OMP_CLAUSE_CHAIN (c); @@ -8008,7 +8012,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else { t = OMP_CLAUSE_DECL (c); - if (TREE_CODE (t) != TREE_LIST + if (TREE_CODE (t) != OMP_ARRAY_SECTION && !type_dependent_expression_p (t) && !omp_mappable_type (TREE_TYPE (t))) { @@ -8189,7 +8193,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)) + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH + || (!openacc && EXPR_P (t)))) break; if (DECL_P (t)) error_at (OMP_CLAUSE_LOCATION (c), @@ -8584,15 +8589,15 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_HAS_DEVICE_ADDR: t = OMP_CLAUSE_DECL (c); - if (TREE_CODE (t) == TREE_LIST) + if (TREE_CODE (t) == OMP_ARRAY_SECTION) { if (handle_omp_array_sections (c, ort)) remove = true; else { t = OMP_CLAUSE_DECL (c); - while (TREE_CODE (t) == TREE_LIST) - t = TREE_CHAIN (t); + while (TREE_CODE (t) == OMP_ARRAY_SECTION) + t = TREE_OPERAND (t, 0); while (TREE_CODE (t) == INDIRECT_REF || TREE_CODE (t) == ARRAY_REF) t = TREE_OPERAND (t, 0); @@ -8963,10 +8968,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (DECL_P (t)) bitmap_clear_bit (&aligned_head, DECL_UID (t)); } - else if (TREE_CODE (t) == TREE_LIST) + else if (TREE_CODE (t) == OMP_ARRAY_SECTION) { - while (TREE_CODE (t) == TREE_LIST) - t = TREE_CHAIN (t); + while (TREE_CODE (t) == OMP_ARRAY_SECTION) + t = TREE_OPERAND (t, 0); if (DECL_P (t)) bitmap_clear_bit (&aligned_head, DECL_UID (t)); t = OMP_CLAUSE_DECL (c); diff --git a/gcc/cp/typeck.cc b/gcc/cp/typeck.cc index 69b1268cfec9..73890c0c27bc 100644 --- a/gcc/cp/typeck.cc +++ b/gcc/cp/typeck.cc @@ -4786,6 +4786,56 @@ build_x_array_ref (location_t loc, tree arg1, tree arg2, return expr; } +/* Build an OpenMP array section reference, creating an exact type for the + resulting expression based on the element type and bounds if possible. If + we have variable bounds, create an incomplete array type for the result + instead. */ + +tree +build_omp_array_section (location_t loc, tree array_expr, tree index, + tree length) +{ + tree idxtype; + + /* If we know the integer bounds, create an index type with exact + low/high (or zero/length) bounds. Otherwise, create an incomplete + array type. (This mostly only affects diagnostics.) */ + if (index != NULL_TREE + && length != NULL_TREE + && TREE_CODE (index) == INTEGER_CST + && TREE_CODE (length) == INTEGER_CST) + { + tree low = fold_convert (sizetype, index); + tree high = fold_convert (sizetype, length); + high = size_binop (PLUS_EXPR, low, high); + high = size_binop (MINUS_EXPR, high, size_one_node); + idxtype = build_range_type (sizetype, low, high); + } + else if ((index == NULL_TREE || integer_zerop (index)) + && length != NULL_TREE + && TREE_CODE (length) == INTEGER_CST) + idxtype = build_index_type (length); + else + idxtype = NULL_TREE; + + tree type = TREE_TYPE (array_expr); + gcc_assert (type); + type = non_reference (type); + + tree sectype, eltype = TREE_TYPE (type); + + /* It's not an array or pointer type. Just reuse the type of the + original expression as the type of the array section (an error will be + raised anyway, later). */ + if (eltype == NULL_TREE) + sectype = TREE_TYPE (array_expr); + else + sectype = build_array_type (eltype, idxtype); + + return build3_loc (loc, OMP_ARRAY_SECTION, sectype, array_expr, index, + length); +} + /* Return whether OP is an expression of enum type cast to integer type. In C++ even unsigned enum types are cast to signed integer types. We do not want to issue warnings about comparisons between diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index d086ab8e4455..97f69a9fa935 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -17246,6 +17246,9 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, case TREE_LIST: gcc_unreachable (); + case OMP_ARRAY_SECTION: + gcc_unreachable (); + case COMPOUND_EXPR: ret = gimplify_compound_expr (expr_p, pre_p, fallback != fb_none); break; diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c index 6ee597148472..c749db845b0a 100644 --- a/gcc/testsuite/c-c++-common/gomp/map-6.c +++ b/gcc/testsuite/c-c++-common/gomp/map-6.c @@ -20,12 +20,12 @@ foo (void) ; #pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */ - /* { dg-error "'close' has not been declared" "" { target c++ } .-1 } */ + /* { dg-error "'close' was not declared in this scope" "" { target c++ } .-1 } */ /* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */ ; #pragma omp target map (always a) /* { dg-error "'always' undeclared" "" { target c } } */ - /* { dg-error "'always' has not been declared" "" { target c++ } .-1 } */ + /* { dg-error "'always' was not declared in this scope" "" { target c++ } .-1 } */ /* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */ ; diff --git a/gcc/testsuite/g++.dg/gomp/array-section-1.C b/gcc/testsuite/g++.dg/gomp/array-section-1.C new file mode 100644 index 000000000000..023706b15c5f --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/array-section-1.C @@ -0,0 +1,38 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-original" } + +int x; + +template +void foo() +{ + int arr1[40]; +#pragma omp target map(arr1[x ? C : D]) +// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR \] \[len: [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR \] - \(long int\) &arr1\]\)} "original" } } + { } +#pragma omp target map(arr1[x ? C : D : D]) +// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR \] \[len: [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR \] - \(long int\) &arr1\]\)} "original" } } + { } +#pragma omp target map(arr1[1 : x ? C : D]) +// { dg-final { scan-tree-dump {map\(tofrom:arr1\[1\] \[len: x != 0 \? [0-9]+ : [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: [0-9]+\]\)} "original" } } + { } +} + +int main() +{ + int arr1[40]; +#pragma omp target map(arr1[x ? 3 : 5]) +// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR \] \[len: [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR \] - \(long int\) &arr1\]\)} "original" } } + { } +#pragma omp target map(arr1[x ? 3 : 5 : 5]) +// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR \] \[len: [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR \] - \(long int\) &arr1\]\)} "original" } } + { } +#pragma omp target map(arr1[1 : x ? 3 : 5]) +// { dg-final { scan-tree-dump {map\(tofrom:arr1\[1\] [len: x != 0 ? [0-9]+ : [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: [0-9]+\]\)} "original" } } + { } + + foo<3, 5> (); + + return 0; +} + diff --git a/gcc/testsuite/g++.dg/gomp/array-section-2.C b/gcc/testsuite/g++.dg/gomp/array-section-2.C new file mode 100644 index 000000000000..072108d1f894 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/array-section-2.C @@ -0,0 +1,63 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-original" } + +int x, y; + +class C { + int x, y; + +public: + int foo(); +}; + +int C::foo() +{ + int arr1[40]; + /* There is a parsing ambiguity here without the space. We don't try to + resolve that automatically (though maybe we could, in theory). */ +#pragma omp target map(arr1[::x: ::y]) +// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR \] \[len: \(sizetype\) y \* [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR \] - \(long int\) &arr1\]\)} "original" } } + { } +#pragma omp target map(arr1[::x:]) +// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR \] \[len: \(40 - \(sizetype\) SAVE_EXPR \) \* [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR \] - \(long int\) &arr1\]\)} "original" } } + { } +#pragma omp target map(arr1[: ::y]) +// { dg-final { scan-tree-dump {map\(tofrom:arr1\[0\] \[len: \(sizetype\) y \* [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: 0\]\)} "original" } } + { } + return ::x + ::y; +} + +template +class Ct { + T x, y; + +public: + void foo(); +}; + +template +void Ct::foo() +{ + int arr1[40]; +#pragma omp target map(arr1[::x: ::y]) +// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR \] \[len: \(sizetype\) y \* [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR \] - \(long int\) &arr1\]\)} "original" } } + { } +#pragma omp target map(arr1[::x:]) +// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR \] \[len: \(40 - \(sizetype\) SAVE_EXPR \) \* [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR \] - \(long int\) &arr1\]\)} "original" } } + { } +#pragma omp target map(arr1[: ::y]) +// { dg-final { scan-tree-dump {map\(tofrom:arr1\[0\] \[len: \(sizetype\) y \* [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: 0\]\)} "original" } } + { } +} + +int main() +{ + C c; + Ct ct; + + c.foo (); + ct.foo (); + + return 0; +} + diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-section-1.C b/gcc/testsuite/g++.dg/gomp/bad-array-section-1.C new file mode 100644 index 000000000000..0ea0c0e66f58 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-section-1.C @@ -0,0 +1,35 @@ +// { dg-do compile } + +int foo (int *ptr); + +template +T baz (T *ptr); + +template +void bar() +{ + T arr[20]; + +#pragma omp target map(baz(arr[3:5])) +// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 } +// { dg-error {expected '\)' before ':' token} "" { target *-*-* } .-2 } +// { dg-error {expected '\)' before '\]' token} "" { target *-*-* } .-3 } +// { dg-error {expected '#pragma omp' clause before '\]' token} "" { target *-*-* } .-4 } + { } +} + +int main() +{ + int arr[20]; + // Reject array section as function argument. +#pragma omp target map(foo(arr[3:5])) +// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 } +// { dg-error {expected '\)' before ':' token} "" { target *-*-* } .-2 } +// { dg-error {expected '\)' before '\]' token} "" { target *-*-* } .-3 } +// { dg-error {expected '#pragma omp' clause before '\]' token} "" { target *-*-* } .-4 } + { } + + bar (); + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-section-10.C b/gcc/testsuite/g++.dg/gomp/bad-array-section-10.C new file mode 100644 index 000000000000..393b0fefe512 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-section-10.C @@ -0,0 +1,35 @@ +// { dg-do compile } + +template +void foo() +{ + int arr1[40]; +#pragma omp target map(arr1[4,C:]) +// { dg-warning "top-level comma expression in array subscript is deprecated" "" { target c++20_only } .-1 } + { } +#pragma omp target map(arr1[4,5:C,7]) +// { dg-warning "top-level comma expression in array subscript is deprecated" "" { target c++20_only } .-1 } + { } +#pragma omp target map(arr1[:8,C,10]) +// { dg-warning "top-level comma expression in array subscript is deprecated" "" { target c++20_only } .-1 } + { } +} + +int main() +{ + int arr1[40]; +#pragma omp target map(arr1[4,5:]) +// { dg-warning "top-level comma expression in array subscript is deprecated" "" { target c++20_only } .-1 } + { } +#pragma omp target map(arr1[4,5:6,7]) +// { dg-warning "top-level comma expression in array subscript is deprecated" "" { target c++20_only } .-1 } + { } +#pragma omp target map(arr1[:8,9,10]) +// { dg-warning "top-level comma expression in array subscript is deprecated" "" { target c++20_only } .-1 } + { } + + foo<6, 9> (); + + return 0; +} + diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-section-11.C b/gcc/testsuite/g++.dg/gomp/bad-array-section-11.C new file mode 100644 index 000000000000..dea3b4428f07 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-section-11.C @@ -0,0 +1,36 @@ +// { dg-do compile } +// { dg-additional-options "-std=c++23" } + +template +void foo() +{ + int arr1[40]; +#pragma omp target map(arr1[4,C:]) +// { dg-error "cannot use multidimensional subscript in OpenMP array section" "" { target *-*-* } .-1 } + { } +#pragma omp target map(arr1[4,5:C,7]) +// { dg-error "cannot use multidimensional subscript in OpenMP array section" "" { target *-*-* } .-1 } + { } +#pragma omp target map(arr1[:8,C,10]) +// { dg-error "cannot use multidimensional subscript in OpenMP array section" "" { target *-*-* } .-1 } + { } +} + +int main() +{ + int arr1[40]; +#pragma omp target map(arr1[4,5:]) +// { dg-error "cannot use multidimensional subscript in OpenMP array section" "" { target *-*-* } .-1 } + { } +#pragma omp target map(arr1[4,5:6,7]) +// { dg-error "cannot use multidimensional subscript in OpenMP array section" "" { target *-*-* } .-1 } + { } +#pragma omp target map(arr1[:8,9,10]) +// { dg-error "cannot use multidimensional subscript in OpenMP array section" "" { target *-*-* } .-1 } + { } + + foo<6, 9> (); + + return 0; +} + diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-section-2.C b/gcc/testsuite/g++.dg/gomp/bad-array-section-2.C new file mode 100644 index 000000000000..811d1fee5a0b --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-section-2.C @@ -0,0 +1,33 @@ +// { dg-do compile } +// { dg-additional-options "-std=c++11" } + +template +void foo() +{ + T arr[20]; + // Reject array section in lambda function. +#pragma omp target map([&](const int x) -> T* { return arr[0:x]; } (5)) +// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 } +// { dg-error {invalid conversion from 'int' to 'int\*'} "" { target *-*-* } .-2 } +// { dg-error {expected ';' before ':' token} "" { target *-*-* } .-3 } +// { dg-error {expected primary-expression before ':' token} "" { target *-*-* } .-4 } +// { dg-message {sorry, unimplemented: unsupported map expression 'foo\(\)::\{arr\}.foo\(\)::\(5\)'} "" { target *-*-* } .-5 } + { } +} + +int main() +{ + int arr[20]; + // Reject array section in lambda function. +#pragma omp target map([&](const int x) -> int* { return arr[0:x]; } (5)) +// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 } +// { dg-error {invalid conversion from 'int' to 'int\*'} "" { target *-*-* } .-2 } +// { dg-error {expected ';' before ':' token} "" { target *-*-* } .-3 } +// { dg-error {expected primary-expression before ':' token} "" { target *-*-* } .-4 } +// { dg-message {sorry, unimplemented: unsupported map expression 'main\(\)::\{arr\}.main\(\)::\(5\)'} "" { target *-*-* } .-5 } + { } + + foo (); + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-section-3.C b/gcc/testsuite/g++.dg/gomp/bad-array-section-3.C new file mode 100644 index 000000000000..d1f067af2e98 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-section-3.C @@ -0,0 +1,28 @@ +// { dg-do compile } + +template +void foo() +{ + T arr[20]; + // Reject array section in statement expression. +#pragma omp target map( ({ int x = 5; arr[0:x]; }) ) +// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 } +// { dg-error {expected ';' before ':' token} "" { target *-*-* } .-2 } +// { dg-message {sorry, unimplemented: unsupported map expression '\(\{\.\.\.\}\)'} "" { target *-*-* } .-3 } + { } +} + +int main() +{ + int arr[20]; + // Reject array section in statement expression. +#pragma omp target map( ({ int x = 5; arr[0:x]; }) ) +// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 } +// { dg-error {expected ';' before ':' token} "" { target *-*-* } .-2 } +// { dg-message {sorry, unimplemented: unsupported map expression '\(\{\.\.\.\}\)'} "" { target *-*-* } .-3 } + { } + + foo (); + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-section-4.C b/gcc/testsuite/g++.dg/gomp/bad-array-section-4.C new file mode 100644 index 000000000000..707c2c31cb2a --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-section-4.C @@ -0,0 +1,50 @@ +// { dg-do compile } + +template +struct St { + T *ptr; +}; + +template +void foo() +{ + T arr[20]; + + // Reject array section in compound initialiser. +#pragma omp target map( (struct St) { .ptr = (T *) arr[5:5] } ) +// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 } +// { dg-error {expected primary-expression before 'struct'} "" { target *-*-* } .-2 } +// { dg-error {expected '\)' before 'struct'} "" { target *-*-* } .-3 } + { } + + // ...and this is unsupported too (probably not useful anyway). +#pragma omp target map( (struct St) { .ptr = &arr[5] } ) +// { dg-message {sorry, unimplemented: unsupported map expression 'St\{\(\& arr\[5\]\)\}'} "" { target *-*-* } .-1 } + { } +} + +struct S { + int *ptr; +}; + +int main() +{ + int arr[20]; + + // Reject array section in compound initialiser. +#pragma omp target map( (struct S) { .ptr = (int *) arr[5:5] } ) +// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 } +// { dg-warning {cast to pointer from integer of different size} "" { target *-*-* } .-2 } +// { dg-error {expected primary-expression before 'struct'} "" { target *-*-* } .-3 } +// { dg-error {expected '\)' before 'struct'} "" { target *-*-* } .-4 } + { } + + // ...and this is unsupported too (probably not useful anyway). +#pragma omp target map( (struct S) { .ptr = &arr[5] } ) +// { dg-message {sorry, unimplemented: unsupported map expression 'S\{\(\& arr\[5\]\)\}'} "" { target *-*-* } .-1 } + { } + + foo (); + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-section-5.C b/gcc/testsuite/g++.dg/gomp/bad-array-section-5.C new file mode 100644 index 000000000000..f9c27d48f0e3 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-section-5.C @@ -0,0 +1,50 @@ +// { dg-do compile } + +int x; + +template +void foo() +{ + T arr[20]; + T *ptr; + /* "arr[1:10]" looks like it might be an expression of array type, hence + able to be indexed (again). This isn't allowed, though. */ +#pragma omp target map(arr[1:10][2]) +// { dg-error {'arr\[1\]' does not have pointer or array type} "" { target *-*-* } .-1 } + { } +#pragma omp target map(arr[1:x][2]) +// { dg-error {'arr\[1\]' does not have pointer or array type} "" { target *-*-* } .-1 } + { } + // ...and nor is this. +#pragma omp target map(ptr[1:10][2]) +// { dg-error {'\*\(ptr \+ [0-9]+\)' does not have pointer or array type} "" { target *-*-* } .-1 } + { } +#pragma omp target map(ptr[1:x][2]) +// { dg-error {'\*\(ptr \+ [0-9]+\)' does not have pointer or array type} "" { target *-*-* } .-1 } + { } +} + +int main() +{ + int arr[20]; + int *ptr; + /* "arr[1:10]" looks like it might be an expression of array type, hence + able to be indexed (again). This isn't allowed, though. */ +#pragma omp target map(arr[1:10][2]) +// { dg-error {'arr\[1\]' does not have pointer or array type} "" { target *-*-* } .-1 } + { } +#pragma omp target map(arr[1:x][2]) +// { dg-error {'arr\[1\]' does not have pointer or array type} "" { target *-*-* } .-1 } + { } + // ...and nor is this. +#pragma omp target map(ptr[1:10][2]) +// { dg-error {'\*\(ptr \+ [0-9]+\)' does not have pointer or array type} "" { target *-*-* } .-1 } + { } +#pragma omp target map(ptr[1:x][2]) +// { dg-error {'\*\(ptr \+ [0-9]+\)' does not have pointer or array type} "" { target *-*-* } .-1 } + { } + + foo (); + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-section-6.C b/gcc/testsuite/g++.dg/gomp/bad-array-section-6.C new file mode 100644 index 000000000000..418ee80431f3 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-section-6.C @@ -0,0 +1,24 @@ +// { dg-do compile } + +bool partly = false; + +template +void foo() +{ + T arr[20]; +#pragma omp target map(partly ? arr[5:5] : arr) +// { dg-message {sorry, unimplemented: unsupported map expression '\(partly \? \(\(int\*\)\(\& arr\[5:5\]\)\) : \(\(int\*\)\(\& arr\)\)\)'} "" { target *-*-* } .-1 } + { } +} + +int main() +{ + int arr[20]; +#pragma omp target map(partly ? arr[5:5] : arr) +// { dg-message {sorry, unimplemented: unsupported map expression '\(partly \? \(\(int\*\)\(\& arr\[5:5\]\)\) : \(\(int\*\)\(\& arr\)\)\)'} "" { target *-*-* } .-1 } + { } + + foo (); + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-section-7.C b/gcc/testsuite/g++.dg/gomp/bad-array-section-7.C new file mode 100644 index 000000000000..24ac165e2bdb --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-section-7.C @@ -0,0 +1,36 @@ +// { dg-do compile } + +int x; + +template +void foo() +{ + T arr[20]; + // Here we know the type of the array section (the upper bound is reported)... +#pragma omp target map(arr[5:5] * 2) +// { dg-error {invalid operands of types 'int \[10\]' and 'int'} "" { target *-*-* } .-1 } + { } + /* ...but here, we have an incomplete array type because of the variable + low bound 'x'. */ +#pragma omp target map(arr[x:5] * 2) +// { dg-error {invalid operands of types 'int \[\]' and 'int'} "" { target *-*-* } .-1 } + { } +} + +int main() +{ + int arr[20]; + // Here we know the type of the array section (the upper bound is reported)... +#pragma omp target map(arr[5:5] * 2) +// { dg-error {invalid operands of types 'int \[10\]' and 'int'} "" { target *-*-* } .-1 } + { } + /* ...but here, we have an incomplete array type because of the variable + low bound 'x'. */ +#pragma omp target map(arr[x:5] * 2) +// { dg-error {invalid operands of types 'int \[\]' and 'int'} "" { target *-*-* } .-1 } + { } + + foo (); + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-section-8.C b/gcc/testsuite/g++.dg/gomp/bad-array-section-8.C new file mode 100644 index 000000000000..2353722a581d --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-section-8.C @@ -0,0 +1,53 @@ +// { dg-do compile } + +int x; + +template +struct Tt { + X arr[20]; +}; + +template +struct St { + X *tvec; +}; + +template +void foo() +{ + struct St > *s; + // You can't use an array section like this. Make sure sensible errors are + // reported. +#pragma omp target map(s->tvec[3:5].arr[0:20]) +// { dg-error {request for member 'arr' in 's->St >::tvec\[3:5\]', which is of non-class type 'Tt \[8\]'} "" { target *-*-* } .-1 } + { } +#pragma omp target map(s->tvec[5:x].arr[0:20]) +// { dg-error {invalid use of array with unspecified bounds} "" { target *-*-* } .-1 } + { } +} + +struct T { + int arr[20]; +}; + +struct S { + struct T *tvec; +}; + +int main() +{ + struct S *s; + // You can't use an array section like this. Make sure sensible errors are + // reported. +#pragma omp target map(s->tvec[3:5].arr[0:20]) +// { dg-error {request for member 'arr' in 's->S::tvec\[3:5\]', which is of non-class type 'T \[8\]'} "" { target *-*-* } .-1 } + { } +#pragma omp target map(s->tvec[5:x].arr[0:20]) +// { dg-error {invalid use of array with unspecified bounds} "" { target *-*-* } .-1 } +// { dg-error {expected '\)' before 'arr'} "" { target *-*-* } .-2 } + { } + + foo (); + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-section-9.C b/gcc/testsuite/g++.dg/gomp/bad-array-section-9.C new file mode 100644 index 000000000000..bba7772a3c90 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-section-9.C @@ -0,0 +1,39 @@ +// { dg-do compile } + +int x; + +template +void foo() +{ + T arr1[40]; + T arr2[40]; +#pragma omp target map(arr1[arr2[4:5]:arr2[6:7]]) +// { dg-error {low bound 'arr2\[4:5\]' of array section does not have integral type} "" { target *-*-* } .-1 } + { } +#pragma omp target map(arr1[arr2[:1]:arr2[6:1]]) +// { dg-error {low bound 'arr2\[:1\]' of array section does not have integral type} "" { target *-*-* } .-1 } + { } +#pragma omp target map(arr1[x:arr2[6:1]]) +// { dg-error {length 'arr2\[6:1\]' of array section does not have integral type} "" { target *-*-* } .-1 } + { } +} + +int main() +{ + int arr1[40]; + int arr2[40]; +#pragma omp target map(arr1[arr2[4:5]:arr2[6:7]]) +// { dg-error {low bound 'arr2\[4:5\]' of array section does not have integral type} "" { target *-*-* } .-1 } + { } +#pragma omp target map(arr1[arr2[:1]:arr2[6:1]]) +// { dg-error {low bound 'arr2\[:1\]' of array section does not have integral type} "" { target *-*-* } .-1 } + { } +#pragma omp target map(arr1[x:arr2[6:1]]) +// { dg-error {length 'arr2\[6:1\]' of array section does not have integral type} "" { target *-*-* } .-1 } + { } + + foo (); + + return 0; +} + diff --git a/gcc/testsuite/g++.dg/gomp/has_device_addr-non-lvalue-1.C b/gcc/testsuite/g++.dg/gomp/has_device_addr-non-lvalue-1.C new file mode 100644 index 000000000000..3d778538d3ab --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/has_device_addr-non-lvalue-1.C @@ -0,0 +1,36 @@ +// { dg-do compile } + +#include +#include +#include + +typedef struct { + int arr[100]; +} S; + +int main() +{ + S *s = new S; + + memset (s->arr, '\0', sizeof s->arr); + +#pragma omp target enter data map(to: (*s).arr) + /* You can't do this, at least as of OpenMP 5.2. "has_device_addr" takes + a "variable list" item type + (OpenMP 5.2, "5.4.9 has_device_addr Clause"). */ +#pragma omp target has_device_addr((*s).arr[5:20]) +// { dg-error {expected unqualified-id before '\(' token} "" { target *-*-* } .-1 } + { + for (int i = 5; i < 25; i++) + s->arr[i] = i; + } + +#pragma omp target exit data map(from: (*s).arr) + + for (int i = 0; i < 100; i++) + assert (i >= 5 && i < 25 ? s->arr[i] == i : s->arr[i] == 0); + + delete s; + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/ind-base-3.C b/gcc/testsuite/g++.dg/gomp/ind-base-3.C new file mode 100644 index 000000000000..7695b1f907e1 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/ind-base-3.C @@ -0,0 +1,37 @@ +#include + +struct S { + int x[10]; +}; + +S * +choose (S *a, S *b, int c) +{ + if (c < 5) + return a; + else + return b; +} + +int main (int argc, char *argv[]) +{ + S a, b; + + for (int i = 0; i < 10; i++) + a.x[i] = b.x[i] = 0; + + for (int i = 0; i < 10; i++) + { +#pragma omp target map(choose(&a, &b, i)->x[:10]) +/* { dg-message {sorry, unimplemented: unsupported map expression 'choose\(\(& a\), \(& b\), i\)->S::x\[0\]'} "" { target *-*-* } .-1 } */ + for (int j = 0; j < 10; j++) + choose (&a, &b, i)->x[j]++; + } + + for (int i = 0; i < 10; i++) + assert (a.x[i] == 5 && b.x[i] == 5); + + return 0; +} + + diff --git a/gcc/testsuite/g++.dg/gomp/map-assignment-1.C b/gcc/testsuite/g++.dg/gomp/map-assignment-1.C new file mode 100644 index 000000000000..5979ec379f19 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/map-assignment-1.C @@ -0,0 +1,12 @@ +#include + +int main (int argc, char *argv[]) +{ + int a = 5, b = 2; +#pragma omp target map(a += b) + /* { dg-message {sorry, unimplemented: unsupported map expression '\(a = \(a \+ b\)\)'} "" { target *-*-* } .-1 } */ + { + a++; + } + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/map-inc-1.C b/gcc/testsuite/g++.dg/gomp/map-inc-1.C new file mode 100644 index 000000000000..b469a4bd5485 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/map-inc-1.C @@ -0,0 +1,10 @@ +int main (int argc, char *argv[]) +{ + int a = 5; +#pragma omp target map(++a) + /* { dg-message {sorry, unimplemented: unsupported map expression '\+\+ a'} "" { target *-*-* } .-1 } */ + { + a++; + } + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/map-lvalue-ref-1.C b/gcc/testsuite/g++.dg/gomp/map-lvalue-ref-1.C new file mode 100644 index 000000000000..d720d4318ae4 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/map-lvalue-ref-1.C @@ -0,0 +1,19 @@ +#include + +int glob = 10; + +int& foo () +{ + return glob; +} + +int main (int argc, char *argv[]) +{ +#pragma omp target map(foo()) + /* { dg-message {sorry, unimplemented: unsupported map expression 'foo\(\)'} "" { target *-*-* } .-1 } */ + { + foo()++; + } + assert (glob == 11); + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/map-ptrmem-1.C b/gcc/testsuite/g++.dg/gomp/map-ptrmem-1.C new file mode 100644 index 000000000000..c4023f59fc60 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/map-ptrmem-1.C @@ -0,0 +1,37 @@ +#include + +struct S { + int x; + int *ptr; +}; + +int +main (int argc, char *argv[]) +{ + S s; + int S::* xp = &S::x; + int* S::* ptrp = &S::ptr; + + s.ptr = new int[64]; + + s.*xp = 6; + for (int i = 0; i < 64; i++) + (s.*ptrp)[i] = i; + +#pragma omp target map(s.*xp, s.*ptrp, (s.*ptrp)[:64]) + /* { dg-message {sorry, unimplemented: pointer-to-member mapping '\*\(\*\(\(\(int\*\*\)\(& s\)\) \+ \(\(sizetype\)ptrp\)\)\)' not supported} "" { target *-*-* } .-1 } */ + /* { dg-message {sorry, unimplemented: pointer-to-member mapping '\*\(\(\(int\*\*\)\(& s\)\) \+ \(\(sizetype\)ptrp\)\)' not supported} "" { target *-*-* } .-2 } */ + /* { dg-message {sorry, unimplemented: pointer-to-member mapping '\*\(\(\(int\*\)\(& s\)\) \+ \(\(sizetype\)xp\)\)' not supported} "" { target *-*-* } .-3 } */ +#pragma omp teams distribute parallel for + for (int i = 0; i < 64; i++) + { + (s.*xp)++; + (s.*ptrp)[i]++; + } + + assert (s.*xp == 70); + for (int i = 0; i < 64; i++) + assert ((s.*ptrp)[i] == i + 1); + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/map-ptrmem-2.C b/gcc/testsuite/g++.dg/gomp/map-ptrmem-2.C new file mode 100644 index 000000000000..fbf379da0eb2 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/map-ptrmem-2.C @@ -0,0 +1,40 @@ +#include + +struct S { + int x; + int *ptr; +}; + +int +main (int argc, char *argv[]) +{ + S *s = new S; + int S::* xp = &S::x; + int* S::* ptrp = &S::ptr; + + s->ptr = new int[64]; + + s->*xp = 4; + for (int i = 0; i < 64; i++) + (s->*ptrp)[i] = i; + +#pragma omp target map(s->*xp, s->*ptrp, (s->*ptrp)[:64]) + /* { dg-message {sorry, unimplemented: pointer-to-member mapping '\*\(\(\(int\*\*\)s\) \+ \(\(sizetype\)ptrp\)\)' not supported} "" { target *-*-* } .-1 } */ + /* { dg-message {sorry, unimplemented: pointer-to-member mapping '\*\(\(\(int\*\)s\) \+ \(\(sizetype\)xp\)\)' not supported} "" { target *-*-* } .-2 } */ + /* { dg-message {sorry, unimplemented: pointer-to-member mapping '\*\(\*\(\(\(int\*\*\)s\) \+ \(\(sizetype\)ptrp\)\)\)' not supported} "" { target *-*-* } .-3 } */ +#pragma omp teams distribute parallel for + for (int i = 0; i < 64; i++) + { + (s->*xp)++; + (s->*ptrp)[i]++; + } + + assert (s->*xp == 68); + for (int i = 0; i < 64; i++) + assert ((s->*ptrp)[i] == i + 1); + + delete s->ptr; + delete s; + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/map-static-cast-lvalue-1.C b/gcc/testsuite/g++.dg/gomp/map-static-cast-lvalue-1.C new file mode 100644 index 000000000000..3af9668202cb --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/map-static-cast-lvalue-1.C @@ -0,0 +1,17 @@ +#include + +int foo (int x) +{ +#pragma omp target map(static_cast(x)) + /* { dg-message {sorry, unimplemented: unsupported map expression '& x'} "" { target *-*-* } .-1 } */ + { + x += 3; + } + return x; +} + +int main (int argc, char *argv[]) +{ + assert (foo (5) == 8); + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/map-ternary-1.C b/gcc/testsuite/g++.dg/gomp/map-ternary-1.C new file mode 100644 index 000000000000..7b365a909bbb --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/map-ternary-1.C @@ -0,0 +1,20 @@ +#include + +int foo (bool yesno) +{ + int x = 5, y = 7; +#pragma omp target map(yesno ? x : y) + /* { dg-message {sorry, unimplemented: unsupported map expression '\(yesno \? x : y\)'} "" { target *-*-* } .-1 } */ + { + x += 3; + y += 5; + } + return yesno ? x : y; +} + +int main (int argc, char *argv[]) +{ + assert (foo (true) == 8); + assert (foo (false) == 12); + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/member-array-2.C b/gcc/testsuite/g++.dg/gomp/member-array-2.C new file mode 100644 index 000000000000..caf8ece42624 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/member-array-2.C @@ -0,0 +1,91 @@ +#include + +typedef int intarr100[100]; + +class C { + int arr[100]; + int *ptr; + +public: + C(); + ~C(); + void zero (); + void do_operation (); + void check (int, int); + intarr100 &get_arr () { return arr; } + int *get_ptr() { return ptr; } +}; + +C::C() +{ + ptr = new int[100]; + for (int i = 0; i < 100; i++) + arr[i] = 0; +} + +C::~C() +{ + delete ptr; +} + +void +C::zero () +{ + for (int i = 0; i < 100; i++) + arr[i] = ptr[i] = 0; +} + +void +C::do_operation () +{ +#pragma omp target map(arr, ptr, ptr[:100]) +#pragma omp teams distribute parallel for + for (int i = 0; i < 100; i++) + { + arr[i] = arr[i] + 3; + ptr[i] = ptr[i] + 5; + } +} + +void +C::check (int arrval, int ptrval) +{ + for (int i = 0; i < 100; i++) + { + assert (arr[i] == arrval); + assert (ptr[i] == ptrval); + } +} + +int +main (int argc, char *argv[]) +{ + C c; + + c.zero (); + c.do_operation (); + c.check (3, 5); + + /* It might sort of make sense to be able to do this, but we don't support + it for now. */ + #pragma omp target map(c.get_arr()[:100]) + /* { dg-message {sorry, unimplemented: unsupported map expression 'c\.C::get_arr\(\)\[0\]'} "" { target *-*-* } .-1 } */ + #pragma omp teams distribute parallel for + for (int i = 0; i < 100; i++) + c.get_arr()[i] += 2; + + c.check (5, 5); + + /* Same for this. */ + #pragma omp target map(c.get_ptr(), c.get_ptr()[:100]) + /* { dg-message {sorry, unimplemented: unsupported map expression 'c\.C::get_ptr\(\)'} "" { target *-*-* } .-1 } */ + /* { dg-message {sorry, unimplemented: unsupported map expression '\* c\.C::get_ptr\(\)'} "" { target *-*-* } .-2 } */ + #pragma omp teams distribute parallel for + for (int i = 0; i < 100; i++) + c.get_ptr()[i] += 3; + + c.check (5, 8); + + return 0; +} + diff --git a/gcc/testsuite/g++.dg/gomp/pr67522.C b/gcc/testsuite/g++.dg/gomp/pr67522.C index da8cb74d1fa2..4a901ba68c7a 100644 --- a/gcc/testsuite/g++.dg/gomp/pr67522.C +++ b/gcc/testsuite/g++.dg/gomp/pr67522.C @@ -12,7 +12,7 @@ foo (void) for (int i = 0; i < 16; i++) ; - #pragma omp target map (S[0:10]) // { dg-error "is not a variable in" } + #pragma omp target map (S[0:10]) // { dg-error "expected primary-expression before '\\\[' token" } ; #pragma omp task depend (inout: S[0:10]) // { dg-error "is not a variable in" } diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index f55f0ffd3db6..1399667cbdb6 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -2549,6 +2549,20 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags, } break; + case OMP_ARRAY_SECTION: + op0 = TREE_OPERAND (node, 0); + if (op_prio (op0) < op_prio (node)) + pp_left_paren (pp); + dump_generic_node (pp, op0, spc, flags, false); + if (op_prio (op0) < op_prio (node)) + pp_right_paren (pp); + pp_left_bracket (pp); + dump_generic_node (pp, TREE_OPERAND (node, 1), spc, flags, false); + pp_colon (pp); + dump_generic_node (pp, TREE_OPERAND (node, 2), spc, flags, false); + pp_right_bracket (pp); + break; + case CONSTRUCTOR: { unsigned HOST_WIDE_INT ix; diff --git a/gcc/tree.def b/gcc/tree.def index 69783bd44fcc..f82bf257aedc 100644 --- a/gcc/tree.def +++ b/gcc/tree.def @@ -1313,6 +1313,9 @@ DEFTREECODE (OMP_ATOMIC_CAPTURE_NEW, "omp_atomic_capture_new", tcc_statement, 2) /* OpenMP clauses. */ DEFTREECODE (OMP_CLAUSE, "omp_clause", tcc_exceptional, 0) +/* An OpenMP array section. */ +DEFTREECODE (OMP_ARRAY_SECTION, "omp_array_section", tcc_expression, 3) + /* TRANSACTION_EXPR tree code. Operand 0: BODY: contains body of the transaction. */ DEFTREECODE (TRANSACTION_EXPR, "transaction_expr", tcc_expression, 1) diff --git a/libgomp/testsuite/libgomp.c++/baseptrs-4.C b/libgomp/testsuite/libgomp.c++/baseptrs-4.C index 196029ac1868..d5ca79c3344d 100644 --- a/libgomp/testsuite/libgomp.c++/baseptrs-4.C +++ b/libgomp/testsuite/libgomp.c++/baseptrs-4.C @@ -11,11 +11,9 @@ #define REF2PTR_DECL_BASE #define ARRAY_DECL_BASE -// Needs map clause "lvalue"-parsing support. -//#define REF2ARRAY_DECL_BASE +#define REF2ARRAY_DECL_BASE #define PTR_OFFSET_DECL_BASE -// Needs map clause "lvalue"-parsing support. -//#define REF2PTR_OFFSET_DECL_BASE +#define REF2PTR_OFFSET_DECL_BASE #define MAP_SECTIONS @@ -30,25 +28,21 @@ #define ARRAY_DECL_MEMBER_SLICE #define ARRAY_DECL_MEMBER_SLICE_BASEPTR -// Needs map clause "lvalue"-parsing support. -//#define REF2ARRAY_DECL_MEMBER_SLICE -//#define REF2ARRAY_DECL_MEMBER_SLICE_BASEPTR +#define REF2ARRAY_DECL_MEMBER_SLICE +#define REF2ARRAY_DECL_MEMBER_SLICE_BASEPTR #define PTR_OFFSET_DECL_MEMBER_SLICE #define PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR -// Needs map clause "lvalue"-parsing support. -//#define REF2PTR_OFFSET_DECL_MEMBER_SLICE -//#define REF2PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR +#define REF2PTR_OFFSET_DECL_MEMBER_SLICE +#define REF2PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR #define PTRARRAY_DECL_MEMBER_SLICE #define PTRARRAY_DECL_MEMBER_SLICE_BASEPTR -// Needs map clause "lvalue"-parsing support. -//#define REF2PTRARRAY_DECL_MEMBER_SLICE -//#define REF2PTRARRAY_DECL_MEMBER_SLICE_BASEPTR +#define REF2PTRARRAY_DECL_MEMBER_SLICE +#define REF2PTRARRAY_DECL_MEMBER_SLICE_BASEPTR #define PTRPTR_OFFSET_DECL_MEMBER_SLICE #define PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR -// Needs map clause "lvalue"-parsing support. -//#define REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE -//#define REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR +#define REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE +#define REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR #define NONREF_COMPONENT_BASE #define NONREF_COMPONENT_MEMBER_SLICE diff --git a/libgomp/testsuite/libgomp.c++/baseptrs-6.C b/libgomp/testsuite/libgomp.c++/baseptrs-6.C new file mode 100644 index 000000000000..0fc93d286645 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/baseptrs-6.C @@ -0,0 +1,3199 @@ +// { dg-do run } + +// This is essentially baseptrs-4.C with templates. + +#include +#include + +#define MAP_DECLS + +#define NONREF_DECL_BASE +#define REF_DECL_BASE +#define PTR_DECL_BASE +#define REF2PTR_DECL_BASE + +#define ARRAY_DECL_BASE +#define REF2ARRAY_DECL_BASE +#define PTR_OFFSET_DECL_BASE +#define REF2PTR_OFFSET_DECL_BASE + +#define MAP_SECTIONS + +#define NONREF_DECL_MEMBER_SLICE +#define NONREF_DECL_MEMBER_SLICE_BASEPTR +#define REF_DECL_MEMBER_SLICE +#define REF_DECL_MEMBER_SLICE_BASEPTR +#define PTR_DECL_MEMBER_SLICE +#define PTR_DECL_MEMBER_SLICE_BASEPTR +#define REF2PTR_DECL_MEMBER_SLICE +#define REF2PTR_DECL_MEMBER_SLICE_BASEPTR + +#define ARRAY_DECL_MEMBER_SLICE +#define ARRAY_DECL_MEMBER_SLICE_BASEPTR +#define REF2ARRAY_DECL_MEMBER_SLICE +#define REF2ARRAY_DECL_MEMBER_SLICE_BASEPTR +#define PTR_OFFSET_DECL_MEMBER_SLICE +#define PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR +#define REF2PTR_OFFSET_DECL_MEMBER_SLICE +#define REF2PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR + +#define PTRARRAY_DECL_MEMBER_SLICE +#define PTRARRAY_DECL_MEMBER_SLICE_BASEPTR +#define REF2PTRARRAY_DECL_MEMBER_SLICE +#define REF2PTRARRAY_DECL_MEMBER_SLICE_BASEPTR +#define PTRPTR_OFFSET_DECL_MEMBER_SLICE +#define PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR +#define REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE +#define REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR + +#define NONREF_COMPONENT_BASE +#define NONREF_COMPONENT_MEMBER_SLICE +#define NONREF_COMPONENT_MEMBER_SLICE_BASEPTR + +#define REF_COMPONENT_BASE +#define REF_COMPONENT_MEMBER_SLICE +#define REF_COMPONENT_MEMBER_SLICE_BASEPTR + +#define PTR_COMPONENT_BASE +#define PTR_COMPONENT_MEMBER_SLICE +#define PTR_COMPONENT_MEMBER_SLICE_BASEPTR + +#define REF2PTR_COMPONENT_BASE +#define REF2PTR_COMPONENT_MEMBER_SLICE +#define REF2PTR_COMPONENT_MEMBER_SLICE_BASEPTR + +#ifdef MAP_DECLS +template +void +map_decls (void) +{ + int x = 0; + int &y = x; + int arr[4]; + int (&arrref)[N] = arr; + int *z = &arr[0]; + int *&t = z; + + memset (arr, 0, sizeof arr); + + #pragma omp target map(x) + { + x++; + } + + #pragma omp target map(y) + { + y++; + } + + assert (x == 2); + assert (y == 2); + + /* "A variable that is of type pointer is treated as if it is the base + pointer of a zero-length array section that appeared as a list item in a + map clause." */ + #pragma omp target map(z) + { + z++; + } + + /* "A variable that is of type reference to pointer is treated as if it had + appeared in a map clause as a zero-length array section." + + The pointer here is *not* associated with a target address, so we're not + disallowed from modifying it. */ + #pragma omp target map(t) + { + t++; + } + + assert (z == &arr[2]); + assert (t == &arr[2]); + + #pragma omp target map(arr) + { + arr[2]++; + } + + #pragma omp target map(arrref) + { + arrref[2]++; + } + + assert (arr[2] == 2); + assert (arrref[2] == 2); +} +#endif + +template +struct S { + T a; + T &b; + T *c; + T *&d; + T e[4]; + T (&f)[4]; + + S(T a1, T &b1, T *c1, T *&d1) : + a(a1), b(b1), c(c1), d(d1), f(e) + { + memset (e, 0, sizeof e); + } +}; + +#ifdef NONREF_DECL_BASE +template +void +nonref_decl_base (void) +{ + T a = 0, b = 0, c, *d = &c; + S mys(a, b, &c, d); + + #pragma omp target map(mys.a) + { + mys.a++; + } + + #pragma omp target map(mys.b) + { + mys.b++; + } + + assert (mys.a == 1); + assert (mys.b == 1); + + #pragma omp target map(mys.c) + { + mys.c++; + } + + #pragma omp target map(mys.d) + { + mys.d++; + } + + assert (mys.c == &c + 1); + assert (mys.d == &c + 1); + + #pragma omp target map(mys.e) + { + mys.e[0]++; + } + + #pragma omp target map(mys.f) + { + mys.f[0]++; + } + + assert (mys.e[0] == 2); + assert (mys.f[0] == 2); +} +#endif + +#ifdef REF_DECL_BASE +template +void +ref_decl_base (void) +{ + T a = 0, b = 0, c, *d = &c; + S mys_orig(a, b, &c, d); + S &mys = mys_orig; + + #pragma omp target map(mys.a) + { + mys.a++; + } + + #pragma omp target map(mys.b) + { + mys.b++; + } + + assert (mys.a == 1); + assert (mys.b == 1); + + #pragma omp target map(mys.c) + { + mys.c++; + } + + #pragma omp target map(mys.d) + { + mys.d++; + } + + assert (mys.c == &c + 1); + assert (mys.d == &c + 1); + + #pragma omp target map(mys.e) + { + mys.e[0]++; + } + + #pragma omp target map(mys.f) + { + mys.f[0]++; + } + + assert (mys.e[0] == 2); + assert (mys.f[0] == 2); +} +#endif + +#ifdef PTR_DECL_BASE +template +void +ptr_decl_base (void) +{ + A a = 0, b = 0, c, *d = &c; + S mys_orig(a, b, &c, d); + S *mys = &mys_orig; + + #pragma omp target map(mys->a) + { + mys->a++; + } + + #pragma omp target map(mys->b) + { + mys->b++; + } + + assert (mys->a == 1); + assert (mys->b == 1); + + #pragma omp target map(mys->c) + { + mys->c++; + } + + #pragma omp target map(mys->d) + { + mys->d++; + } + + assert (mys->c == &c + 1); + assert (mys->d == &c + 1); + + #pragma omp target map(mys->e) + { + mys->e[0]++; + } + + #pragma omp target map(mys->f) + { + mys->f[0]++; + } + + assert (mys->e[0] == 2); + assert (mys->f[0] == 2); +} +#endif + +#ifdef REF2PTR_DECL_BASE +template +void +ref2ptr_decl_base (void) +{ + A a = 0, b = 0, c, *d = &c; + S mys_orig(a, b, &c, d); + S *mysp = &mys_orig; + S *&mys = mysp; + + #pragma omp target map(mys->a) + { + mys->a++; + } + + #pragma omp target map(mys->b) + { + mys->b++; + } + + assert (mys->a == 1); + assert (mys->b == 1); + + #pragma omp target map(mys->c) + { + mys->c++; + } + + #pragma omp target map(mys->d) + { + mys->d++; + } + + assert (mys->c == &c + 1); + assert (mys->d == &c + 1); + + #pragma omp target map(mys->e) + { + mys->e[0]++; + } + + #pragma omp target map(mys->f) + { + mys->f[0]++; + } + + assert (mys->e[0] == 2); + assert (mys->f[0] == 2); +} +#endif + +#ifdef ARRAY_DECL_BASE +template +void +array_decl_base (void) +{ + A a = 0, b = 0, c, *d = &c; + S mys[4] = + { + S(a, b, &c, d), + S(a, b, &c, d), + S(a, b, &c, d), + S(a, b, &c, d) + }; + + #pragma omp target map(mys[2].a) + { + mys[2].a++; + } + + #pragma omp target map(mys[2].b) + { + mys[2].b++; + } + + assert (mys[2].a == 1); + assert (mys[2].b == 1); + + #pragma omp target map(mys[2].c) + { + mys[2].c++; + } + + #pragma omp target map(mys[2].d) + { + mys[2].d++; + } + + assert (mys[2].c == &c + 1); + assert (mys[2].d == &c + 1); + + #pragma omp target map(mys[2].e) + { + mys[2].e[0]++; + } + + #pragma omp target map(mys[2].f) + { + mys[2].f[0]++; + } + + assert (mys[2].e[0] == 2); + assert (mys[2].f[0] == 2); +} +#endif + +#ifdef REF2ARRAY_DECL_BASE +template +void +ref2array_decl_base (void) +{ + A a = 0, b = 0, c, *d = &c; + S mys_orig[4] = + { + S(a, b, &c, d), + S(a, b, &c, d), + S(a, b, &c, d), + S(a, b, &c, d) + }; + S (&mys)[4] = mys_orig; + + #pragma omp target map(mys[2].a) + { + mys[2].a++; + } + + #pragma omp target map(mys[2].b) + { + mys[2].b++; + } + + assert (mys[2].a == 1); + assert (mys[2].b == 1); + + #pragma omp target map(mys[2].c) + { + mys[2].c++; + } + + #pragma omp target map(mys[2].d) + { + mys[2].d++; + } + + assert (mys[2].c == &c + 1); + assert (mys[2].d == &c + 1); + + #pragma omp target map(mys[2].e) + { + mys[2].e[0]++; + } + + #pragma omp target map(mys[2].f) + { + mys[2].f[0]++; + } + + assert (mys[2].e[0] == 2); + assert (mys[2].f[0] == 2); +} +#endif + +#ifdef PTR_OFFSET_DECL_BASE +template +void +ptr_offset_decl_base (void) +{ + A a = 0, b = 0, c, *d = &c; + S mys_orig[4] = + { + S(a, b, &c, d), + S(a, b, &c, d), + S(a, b, &c, d), + S(a, b, &c, d) + }; + S *mys = &mys_orig[0]; + + #pragma omp target map(mys[2].a) + { + mys[2].a++; + } + + #pragma omp target map(mys[2].b) + { + mys[2].b++; + } + + assert (mys[2].a == 1); + assert (mys[2].b == 1); + + #pragma omp target map(mys[2].c) + { + mys[2].c++; + } + + #pragma omp target map(mys[2].d) + { + mys[2].d++; + } + + assert (mys[2].c == &c + 1); + assert (mys[2].d == &c + 1); + + #pragma omp target map(mys[2].e) + { + mys[2].e[0]++; + } + + #pragma omp target map(mys[2].f) + { + mys[2].f[0]++; + } + + assert (mys[2].e[0] == 2); + assert (mys[2].f[0] == 2); +} +#endif + +#ifdef REF2PTR_OFFSET_DECL_BASE +template +void +ref2ptr_offset_decl_base (void) +{ + A a = 0, b = 0, c, *d = &c; + S mys_orig[4] = + { + S(a, b, &c, d), + S(a, b, &c, d), + S(a, b, &c, d), + S(a, b, &c, d) + }; + S *mys_ptr = &mys_orig[0]; + S *&mys = mys_ptr; + + #pragma omp target map(mys[2].a) + { + mys[2].a++; + } + + #pragma omp target map(mys[2].b) + { + mys[2].b++; + } + + assert (mys[2].a == 1); + assert (mys[2].b == 1); + + #pragma omp target map(mys[2].c) + { + mys[2].c++; + } + + #pragma omp target map(mys[2].d) + { + mys[2].d++; + } + + assert (mys[2].c == &c + 1); + assert (mys[2].d == &c + 1); + + #pragma omp target map(mys[2].e) + { + mys[2].e[0]++; + } + + #pragma omp target map(mys[2].f) + { + mys[2].f[0]++; + } + + assert (mys[2].e[0] == 2); + assert (mys[2].f[0] == 2); +} +#endif + +#ifdef MAP_SECTIONS +template +void +map_sections (void) +{ + A arr[B]; + A *ptr; + A (&arrref)[B] = arr; + A *&ptrref = ptr; + + ptr = new int[B]; + memset (ptr, 0, sizeof (int) * B); + memset (arr, 0, sizeof (int) * B); + + #pragma omp target map(arr[0:B]) + { + arr[2]++; + } + + #pragma omp target map(ptr[0:B]) + { + ptr[2]++; + } + + #pragma omp target map(arrref[0:B]) + { + arrref[2]++; + } + + #pragma omp target map(ptrref[0:B]) + { + ptrref[2]++; + } + + assert (arr[2] == 2); + assert (ptr[2] == 2); + + delete ptr; +} +#endif + +template +struct T { + A a[10]; + A (&b)[10]; + A *c; + A *&d; + + T(A (&b1)[10], A *c1, A *&d1) : b(b1), c(c1), d(d1) + { + memset (a, 0, sizeof a); + } +}; + +#ifdef NONREF_DECL_MEMBER_SLICE +template +void +nonref_decl_member_slice (void) +{ + A c[B]; + A *d = &c[0]; + T myt(c, &c[0], d); + + memset (c, 0, sizeof c); + + #pragma omp target map(myt.a[0:B]) + { + myt.a[2]++; + } + + #pragma omp target map(myt.b[0:B]) + { + myt.b[2]++; + } + + #pragma omp target enter data map(to: myt.c) + + #pragma omp target map(myt.c[0:B]) + { + myt.c[2]++; + } + + #pragma omp target exit data map(release: myt.c) + + #pragma omp target enter data map(to: myt.d) + + #pragma omp target map(myt.d[0:B]) + { + myt.d[2]++; + } + + #pragma omp target exit data map(from: myt.d) + + assert (myt.a[2] == 1); + assert (myt.b[2] == 3); + assert (myt.c[2] == 3); + assert (myt.d[2] == 3); +} +#endif + +#ifdef NONREF_DECL_MEMBER_SLICE_BASEPTR +template +void +nonref_decl_member_slice_baseptr (void) +{ + A c[B]; + A *d = &c[0]; + T myt(c, &c[0], d); + + memset (c, 0, sizeof c); + + #pragma omp target map(to:myt.c) map(myt.c[0:B]) + { + myt.c[2]++; + } + + #pragma omp target map(to:myt.d) map(myt.d[0:B]) + { + myt.d[2]++; + } + + assert (myt.c[2] == 2); + assert (myt.d[2] == 2); +} +#endif + +#ifdef REF_DECL_MEMBER_SLICE +template +void +ref_decl_member_slice (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real(c, &c[0], d); + T &myt = myt_real; + + memset (c, 0, sizeof c); + + #pragma omp target map(myt.a[0:B]) + { + myt.a[2]++; + } + + #pragma omp target map(myt.b[0:B]) + { + myt.b[2]++; + } + + #pragma omp target enter data map(to: myt.c) + + #pragma omp target map(myt.c[0:B]) + { + myt.c[2]++; + } + + #pragma omp target exit data map(release: myt.c) + + #pragma omp target enter data map(to: myt.d) + + #pragma omp target map(myt.d[0:B]) + { + myt.d[2]++; + } + + #pragma omp target exit data map(release: myt.d) + + assert (myt.a[2] == 1); + assert (myt.b[2] == 3); + assert (myt.c[2] == 3); + assert (myt.d[2] == 3); +} +#endif + +#ifdef REF_DECL_MEMBER_SLICE_BASEPTR +template +void +ref_decl_member_slice_baseptr (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real(c, &c[0], d); + T &myt = myt_real; + + memset (c, 0, sizeof c); + + #pragma omp target map(to:myt.c) map(myt.c[0:B]) + { + myt.c[2]++; + } + + #pragma omp target map(to:myt.d) map(myt.d[0:B]) + { + myt.d[2]++; + } + + assert (myt.c[2] == 2); + assert (myt.d[2] == 2); +} +#endif + +#ifdef PTR_DECL_MEMBER_SLICE +template +void +ptr_decl_member_slice (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real(c, &c[0], d); + T *myt = &myt_real; + + memset (c, 0, sizeof c); + + #pragma omp target enter data map(to: myt) + + #pragma omp target map(myt->a[0:B]) + { + myt->a[2]++; + } + + #pragma omp target map(myt->b[0:B]) + { + myt->b[2]++; + } + + #pragma omp target enter data map(to: myt->c) + + #pragma omp target map(myt->c[0:B]) + { + myt->c[2]++; + } + + #pragma omp target exit data map(release: myt->c) + + #pragma omp target enter data map(to: myt->d) + + #pragma omp target map(myt->d[0:B]) + { + myt->d[2]++; + } + + #pragma omp target exit data map(release: myt, myt->d) + + assert (myt->a[2] == 1); + assert (myt->b[2] == 3); + assert (myt->c[2] == 3); + assert (myt->d[2] == 3); +} +#endif + +#ifdef PTR_DECL_MEMBER_SLICE_BASEPTR +template +void +ptr_decl_member_slice_baseptr (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real(c, &c[0], d); + T *myt = &myt_real; + + memset (c, 0, sizeof c); + + // These ones have an implicit firstprivate for 'myt'. + #pragma omp target map(to:myt->c) map(myt->c[0:B]) + { + myt->c[2]++; + } + + #pragma omp target map(to:myt->d) map(myt->d[0:B]) + { + myt->d[2]++; + } + + // These ones have an explicit "TO" mapping for 'myt'. + #pragma omp target map(to:myt) map(to:myt->c) map(myt->c[0:B]) + { + myt->c[2]++; + } + + #pragma omp target map(to:myt) map(to:myt->d) map(myt->d[0:B]) + { + myt->d[2]++; + } + + assert (myt->c[2] == 4); + assert (myt->d[2] == 4); +} +#endif + +#ifdef REF2PTR_DECL_MEMBER_SLICE +template +void +ref2ptr_decl_member_slice (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real(c, &c[0], d); + T *myt_ptr = &myt_real; + T *&myt = myt_ptr; + + memset (c, 0, sizeof c); + + #pragma omp target enter data map(to: myt) + + #pragma omp target map(myt->a[0:B]) + { + myt->a[2]++; + } + + #pragma omp target map(myt->b[0:B]) + { + myt->b[2]++; + } + + #pragma omp target enter data map(to: myt->c) + + #pragma omp target map(myt->c[0:B]) + { + myt->c[2]++; + } + + #pragma omp target exit data map(release: myt->c) + + #pragma omp target enter data map(to: myt->d) + + #pragma omp target map(myt->d[0:B]) + { + myt->d[2]++; + } + + #pragma omp target exit data map(from: myt, myt->d) + + assert (myt->a[2] == 1); + assert (myt->b[2] == 3); + assert (myt->c[2] == 3); + assert (myt->d[2] == 3); +} +#endif + +#ifdef REF2PTR_DECL_MEMBER_SLICE_BASEPTR +template +void +ref2ptr_decl_member_slice_baseptr (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real(c, &c[0], d); + T *myt_ptr = &myt_real; + T *&myt = myt_ptr; + + memset (c, 0, sizeof c); + + // These ones have an implicit firstprivate for 'myt'. + #pragma omp target map(to:myt->c) map(myt->c[0:B]) + { + myt->c[2]++; + } + + #pragma omp target map(to:myt->d) map(myt->d[0:B]) + { + myt->d[2]++; + } + + // These ones have an explicit "TO" mapping for 'myt'. + #pragma omp target map(to:myt) map(to:myt->c) map(myt->c[0:B]) + { + myt->c[2]++; + } + + #pragma omp target map(to:myt) map(to:myt->d) map(myt->d[0:B]) + { + myt->d[2]++; + } + + assert (myt->c[2] == 4); + assert (myt->d[2] == 4); +} +#endif + +#ifdef ARRAY_DECL_MEMBER_SLICE +template +void +array_decl_member_slice (void) +{ + A c[B]; + A *d = &c[0]; + T myt[4] = + { + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d) + }; + + memset (c, 0, sizeof c); + + #pragma omp target map(myt[2].a[0:B]) + { + myt[2].a[2]++; + } + + #pragma omp target map(myt[2].b[0:B]) + { + myt[2].b[2]++; + } + + #pragma omp target enter data map(to: myt[2].c) + + #pragma omp target map(myt[2].c[0:B]) + { + myt[2].c[2]++; + } + + #pragma omp target exit data map(release: myt[2].c) + + #pragma omp target enter data map(to: myt[2].d) + + #pragma omp target map(myt[2].d[0:B]) + { + myt[2].d[2]++; + } + + #pragma omp target exit data map(release: myt[2].d) + + assert (myt[2].a[2] == 1); + assert (myt[2].b[2] == 3); + assert (myt[2].c[2] == 3); + assert (myt[2].d[2] == 3); +} +#endif + +#ifdef ARRAY_DECL_MEMBER_SLICE_BASEPTR +template +void +array_decl_member_slice_baseptr (void) +{ + A c[10]; + A *d = &c[0]; + T myt[4] = + { + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d) + }; + + memset (c, 0, sizeof c); + + #pragma omp target map(to:myt[2].c) map(myt[2].c[0:B]) + { + myt[2].c[2]++; + } + + #pragma omp target map(to:myt[2].d) map(myt[2].d[0:B]) + { + myt[2].d[2]++; + } + + assert (myt[2].c[2] == 2); + assert (myt[2].d[2] == 2); +} +#endif + +#ifdef REF2ARRAY_DECL_MEMBER_SLICE +template +void +ref2array_decl_member_slice (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real[4] = + { + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d) + }; + T (&myt)[4] = myt_real; + + memset (c, 0, sizeof c); + + #pragma omp target map(myt[2].a[0:B]) + { + myt[2].a[2]++; + } + + #pragma omp target map(myt[2].b[0:B]) + { + myt[2].b[2]++; + } + + #pragma omp target enter data map(to: myt[2].c) + + #pragma omp target map(myt[2].c[0:B]) + { + myt[2].c[2]++; + } + + #pragma omp target exit data map(release: myt[2].c) + + #pragma omp target enter data map(to: myt[2].d) + + #pragma omp target map(myt[2].d[0:B]) + { + myt[2].d[2]++; + } + + #pragma omp target exit data map(release: myt[2].d) + + assert (myt[2].a[2] == 1); + assert (myt[2].b[2] == 3); + assert (myt[2].c[2] == 3); + assert (myt[2].d[2] == 3); +} +#endif + +#ifdef REF2ARRAY_DECL_MEMBER_SLICE_BASEPTR +template +void +ref2array_decl_member_slice_baseptr (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real[4] = + { + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d) + }; + T (&myt)[4] = myt_real; + + memset (c, 0, sizeof c); + + #pragma omp target map(to:myt[2].c) map(myt[2].c[0:B]) + { + myt[2].c[2]++; + } + + #pragma omp target map(to:myt[2].d) map(myt[2].d[0:B]) + { + myt[2].d[2]++; + } + + assert (myt[2].c[2] == 2); + assert (myt[2].d[2] == 2); +} +#endif + +#ifdef PTR_OFFSET_DECL_MEMBER_SLICE +template +void +ptr_offset_decl_member_slice (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real[4] = + { + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d) + }; + T *myt = &myt_real[0]; + + memset (c, 0, sizeof c); + + #pragma omp target map(myt[2].a[0:B]) + { + myt[2].a[2]++; + } + + #pragma omp target map(myt[2].b[0:B]) + { + myt[2].b[2]++; + } + + #pragma omp target enter data map(to: myt[2].c) + + #pragma omp target map(myt[2].c[0:B]) + { + myt[2].c[2]++; + } + + #pragma omp target exit data map(release: myt[2].c) + + #pragma omp target enter data map(to: myt[2].d) + + #pragma omp target map(myt[2].d[0:B]) + { + myt[2].d[2]++; + } + + #pragma omp target exit data map(release: myt[2].d) + + assert (myt[2].a[2] == 1); + assert (myt[2].b[2] == 3); + assert (myt[2].c[2] == 3); + assert (myt[2].d[2] == 3); +} +#endif + +#ifdef PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR +template +void +ptr_offset_decl_member_slice_baseptr (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real[4] = + { + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d) + }; + T *myt = &myt_real[0]; + + memset (c, 0, sizeof c); + + /* Implicit 'myt'. */ + #pragma omp target map(to:myt[2].c) map(myt[2].c[0:B]) + { + myt[2].c[2]++; + } + + #pragma omp target map(to:myt[2].d) map(myt[2].d[0:B]) + { + myt[2].d[2]++; + } + + /* Explicit 'to'-mapped 'myt'. */ + #pragma omp target map(to:myt) map(to:myt[2].c) map(myt[2].c[0:B]) + { + myt[2].c[2]++; + } + + #pragma omp target map(to:myt) map(to:myt[2].d) map(myt[2].d[0:B]) + { + myt[2].d[2]++; + } + + assert (myt[2].c[2] == 4); + assert (myt[2].d[2] == 4); +} +#endif + +#ifdef REF2PTR_OFFSET_DECL_MEMBER_SLICE +template +void +ref2ptr_offset_decl_member_slice (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real[4] = + { + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d) + }; + T *myt_ptr = &myt_real[0]; + T *&myt = myt_ptr; + + memset (c, 0, sizeof c); + + #pragma omp target map(myt[2].a[0:B]) + { + myt[2].a[2]++; + } + + #pragma omp target map(myt[2].b[0:B]) + { + myt[2].b[2]++; + } + + #pragma omp target enter data map(to: myt[2].c) + + #pragma omp target map(myt[2].c[0:B]) + { + myt[2].c[2]++; + } + + #pragma omp target exit data map(release: myt[2].c) + + #pragma omp target enter data map(to: myt[2].d) + + #pragma omp target map(myt[2].d[0:B]) + { + myt[2].d[2]++; + } + + #pragma omp target exit data map(release: myt[2].d) + + assert (myt[2].a[2] == 1); + assert (myt[2].b[2] == 3); + assert (myt[2].c[2] == 3); + assert (myt[2].d[2] == 3); +} +#endif + +#ifdef REF2PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR +template +void +ref2ptr_offset_decl_member_slice_baseptr (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real[4] = + { + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d), + T (c, &c[0], d) + }; + T *myt_ptr = &myt_real[0]; + T *&myt = myt_ptr; + + memset (c, 0, sizeof c); + + /* Implicit 'myt'. */ + #pragma omp target map(to:myt[2].c) map(myt[2].c[0:B]) + { + myt[2].c[2]++; + } + + #pragma omp target map(to:myt[2].d) map(myt[2].d[0:B]) + { + myt[2].d[2]++; + } + + /* Explicit 'to'-mapped 'myt'. */ + #pragma omp target map(to:myt) map(to:myt[2].c) map(myt[2].c[0:B]) + { + myt[2].c[2]++; + } + + #pragma omp target map(to:myt) map(to:myt[2].d) map(myt[2].d[0:B]) + { + myt[2].d[2]++; + } + + assert (myt[2].c[2] == 4); + assert (myt[2].d[2] == 4); +} +#endif + +#ifdef PTRARRAY_DECL_MEMBER_SLICE +template +void +ptrarray_decl_member_slice (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real(c, &c[0], d); + T *myt[4] = + { + &myt_real, + &myt_real, + &myt_real, + &myt_real + }; + + memset (c, 0, sizeof c); + + #pragma omp target enter data map(to: myt[2]) + + #pragma omp target map(myt[2]->a[0:B]) + { + myt[2]->a[2]++; + } + + #pragma omp target map(myt[2]->b[0:B]) + { + myt[2]->b[2]++; + } + + #pragma omp target enter data map(to: myt[2]->c) + + #pragma omp target map(myt[2]->c[0:B]) + { + myt[2]->c[2]++; + } + + #pragma omp target exit data map(from: myt[2]->c) + + #pragma omp target enter data map(to: myt[2]->d) + + #pragma omp target map(myt[2]->d[0:B]) + { + myt[2]->d[2]++; + } + + #pragma omp target exit data map(from: myt[2]->d) + + #pragma omp target exit data map(release: myt[2]) + + assert (myt[2]->a[2] == 1); + assert (myt[2]->b[2] == 3); + assert (myt[2]->c[2] == 3); + assert (myt[2]->d[2] == 3); +} +#endif + +#ifdef PTRARRAY_DECL_MEMBER_SLICE_BASEPTR +template +void +ptrarray_decl_member_slice_baseptr (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real(c, &c[0], d); + T *myt[4] = + { + &myt_real, + &myt_real, + &myt_real, + &myt_real + }; + + memset (c, 0, sizeof c); + + // Implicit 'myt' + #pragma omp target map(to: myt[2]->c) map(myt[2]->c[0:B]) + { + myt[2]->c[2]++; + } + + #pragma omp target map(to: myt[2]->d) map(myt[2]->d[0:B]) + { + myt[2]->d[2]++; + } + + // One element of 'myt' + #pragma omp target map(to:myt[2], myt[2]->c) map(myt[2]->c[0:B]) + { + myt[2]->c[2]++; + } + + #pragma omp target map(to:myt[2], myt[2]->d) map(myt[2]->d[0:B]) + { + myt[2]->d[2]++; + } + + // Explicit map of all of 'myt' + #pragma omp target map(to:myt, myt[2]->c) map(myt[2]->c[0:B]) + { + myt[2]->c[2]++; + } + + #pragma omp target map(to:myt, myt[2]->d) map(myt[2]->d[0:B]) + { + myt[2]->d[2]++; + } + + // Explicit map slice of 'myt' + #pragma omp target map(to:myt[1:3], myt[2]->c) map(myt[2]->c[0:B]) + { + myt[2]->c[2]++; + } + + #pragma omp target map(to:myt[1:3], myt[2]->d) map(myt[2]->d[0:B]) + { + myt[2]->d[2]++; + } + + assert (myt[2]->c[2] == 8); + assert (myt[2]->d[2] == 8); +} +#endif + +#ifdef REF2PTRARRAY_DECL_MEMBER_SLICE +template +void +ref2ptrarray_decl_member_slice (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real(c, &c[0], d); + T *myt_ptrarr[4] = + { + &myt_real, + &myt_real, + &myt_real, + &myt_real + }; + T *(&myt)[4] = myt_ptrarr; + + memset (c, 0, sizeof c); + + #pragma omp target enter data map(to: myt[2]) + + #pragma omp target map(myt[2]->a[0:B]) + { + myt[2]->a[2]++; + } + + #pragma omp target map(myt[2]->b[0:B]) + { + myt[2]->b[2]++; + } + + #pragma omp target enter data map(to: myt[2]->c) + + #pragma omp target map(myt[2]->c[0:B]) + { + myt[2]->c[2]++; + } + + #pragma omp target exit data map(release: myt[2]->c) + + #pragma omp target enter data map(to: myt[2]->d) + + #pragma omp target map(myt[2]->d[0:B]) + { + myt[2]->d[2]++; + } + + #pragma omp target exit data map(release: myt[2]->d) + + #pragma omp target exit data map(release: myt[2]) + + assert (myt[2]->a[2] == 1); + assert (myt[2]->b[2] == 3); + assert (myt[2]->c[2] == 3); + assert (myt[2]->d[2] == 3); +} +#endif + +#ifdef REF2PTRARRAY_DECL_MEMBER_SLICE_BASEPTR +template +void +ref2ptrarray_decl_member_slice_baseptr (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real(c, &c[0], d); + T *myt_ptrarr[4] = + { + &myt_real, + &myt_real, + &myt_real, + &myt_real + }; + T *(&myt)[4] = myt_ptrarr; + + memset (c, 0, sizeof c); + + #pragma omp target map(to:myt[2], myt[2]->c) map(myt[2]->c[0:B]) + { + myt[2]->c[2]++; + } + + #pragma omp target map(to:myt[2], myt[2]->d) map(myt[2]->d[0:B]) + { + myt[2]->d[2]++; + } + + #pragma omp target map(to:myt, myt[2]->c) map(myt[2]->c[0:B]) + { + myt[2]->c[2]++; + } + + #pragma omp target map(to:myt, myt[2]->d) map(myt[2]->d[0:B]) + { + myt[2]->d[2]++; + } + + assert (myt[2]->c[2] == 4); + assert (myt[2]->d[2] == 4); +} +#endif + +#ifdef PTRPTR_OFFSET_DECL_MEMBER_SLICE +template +void +ptrptr_offset_decl_member_slice (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real(c, &c[0], d); + T *myt_ptrarr[4] = + { + &myt_real, + &myt_real, + &myt_real, + &myt_real + }; + T **myt = &myt_ptrarr[0]; + + memset (c, 0, sizeof c); + + #pragma omp target enter data map(to: myt[0:3]) + + /* NOTE: For the implicit firstprivate 'myt' to work, the zeroth element of + myt[] must be mapped above -- otherwise the zero-length array section + lookup fails. */ + #pragma omp target map(myt[2]->a[0:B]) + { + myt[2]->a[2]++; + } + + #pragma omp target map(myt[2]->b[0:B]) + { + myt[2]->b[2]++; + } + + #pragma omp target enter data map(to: myt[2]->c) + + #pragma omp target map(myt[2]->c[0:B]) + { + myt[2]->c[2]++; + } + + #pragma omp target exit data map(from: myt[2]->c) + + #pragma omp target enter data map(to: myt[2]->d) + + #pragma omp target map(myt[2]->d[0:B]) + { + myt[2]->d[2]++; + } + + #pragma omp target exit data map(from: myt[0:3], myt[2]->d) + + assert (myt[2]->a[2] == 1); + assert (myt[2]->b[2] == 3); + assert (myt[2]->c[2] == 3); + assert (myt[2]->d[2] == 3); +} +#endif + +#ifdef PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR +template +void +ptrptr_offset_decl_member_slice_baseptr (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real(c, &c[0], d); + T *myt_ptrarr[4] = + { + 0, + 0, + 0, + &myt_real + }; + T **myt = &myt_ptrarr[0]; + + memset (c, 0, sizeof c); + + #pragma omp target map(to:myt[3], myt[3]->c) map(myt[3]->c[0:B]) + { + myt[3]->c[2]++; + } + + #pragma omp target map(to:myt[3], myt[3]->d) map(myt[3]->d[0:B]) + { + myt[3]->d[2]++; + } + + #pragma omp target map(to:myt, myt[3], myt[3]->c) map(myt[3]->c[0:B]) + { + myt[3]->c[2]++; + } + + #pragma omp target map(to:myt, myt[3], myt[3]->d) map(myt[3]->d[0:B]) + { + myt[3]->d[2]++; + } + + assert (myt[3]->c[2] == 4); + assert (myt[3]->d[2] == 4); +} +#endif + +#ifdef REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE +template +void +ref2ptrptr_offset_decl_member_slice (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real(c, &c[0], d); + T *myt_ptrarr[4] = + { + 0, + 0, + &myt_real, + 0 + }; + T **myt_ptrptr = &myt_ptrarr[0]; + T **&myt = myt_ptrptr; + + memset (c, 0, sizeof c); + + #pragma omp target enter data map(to: myt[0:3]) + + #pragma omp target map(myt[2]->a[0:B]) + { + myt[2]->a[2]++; + } + + #pragma omp target map(myt[2]->b[0:B]) + { + myt[2]->b[2]++; + } + + #pragma omp target enter data map(to:myt[2]->c) + + #pragma omp target map(myt[2]->c[0:B]) + { + myt[2]->c[2]++; + } + + #pragma omp target exit data map(release:myt[2]->c) + + #pragma omp target enter data map(to:myt[2]->d) + + #pragma omp target map(myt[2]->d[0:B]) + { + myt[2]->d[2]++; + } + + #pragma omp target exit data map(release: myt[0:3], myt[2]->d) + + assert (myt[2]->a[2] == 1); + assert (myt[2]->b[2] == 3); + assert (myt[2]->c[2] == 3); + assert (myt[2]->d[2] == 3); +} +#endif + +#ifdef REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR +template +void +ref2ptrptr_offset_decl_member_slice_baseptr (void) +{ + A c[B]; + A *d = &c[0]; + T myt_real(c, &c[0], d); + T *myt_ptrarr[4] = + { + 0, + 0, + &myt_real, + 0 + }; + T **myt_ptrptr = &myt_ptrarr[0]; + T **&myt = myt_ptrptr; + + memset (c, 0, sizeof c); + + #pragma omp target map(to:myt[2], myt[2]->c) map(myt[2]->c[0:B]) + { + myt[2]->c[2]++; + } + + #pragma omp target map(to:myt[2], myt[2]->d) map(myt[2]->d[0:B]) + { + myt[2]->d[2]++; + } + + #pragma omp target map(to:myt, myt[2], myt[2]->c) map(myt[2]->c[0:B]) + { + myt[2]->c[2]++; + } + + #pragma omp target map(to:myt, myt[2], myt[2]->d) map(myt[2]->d[0:B]) + { + myt[2]->d[2]++; + } + + assert (myt[2]->c[2] == 4); + assert (myt[2]->d[2] == 4); +} +#endif + +template +struct U +{ + S s1; + T t1; + S &s2; + T &t2; + S *s3; + T *t3; + S *&s4; + T *&t4; + + U(S &sptr1, T &tptr1, S &sptr2, T &tptr2, + S *sptr3, T *tptr3, S *&sptr4, T *&tptr4) + : s1(sptr1), t1(tptr1), s2(sptr2), t2(tptr2), s3(sptr3), t3(tptr3), + s4(sptr4), t4(tptr4) + { + } +}; + +#define INIT_S(N) \ + A a##N = 0, b##N = 0, c##N = 0, d##N = 0; \ + A *d##N##ptr = &d##N; \ + S s##N(a##N, b##N, &c##N, d##N##ptr) + +#define INIT_T(N) \ + A arr##N[10]; \ + A *ptr##N = &arr##N[0]; \ + T t##N(arr##N, &arr##N[0], ptr##N); \ + memset (arr##N, 0, sizeof arr##N) + +#define INIT_ST \ + INIT_S(1); \ + INIT_T(1); \ + INIT_S(2); \ + INIT_T(2); \ + INIT_S(3); \ + INIT_T(3); \ + A a4 = 0, b4 = 0, c4 = 0, d4 = 0; \ + A *d4ptr = &d4; \ + S *s4 = new S(a4, b4, &c4, d4ptr); \ + A arr4[10]; \ + A *ptr4 = &arr4[0]; \ + T *t4 = new T(arr4, &arr4[0], ptr4); \ + memset (arr4, 0, sizeof arr4) + +#ifdef NONREF_COMPONENT_BASE +template +void +nonref_component_base (void) +{ + INIT_ST; + U myu(s1, t1, s2, t2, &s3, &t3, s4, t4); + + #pragma omp target map(myu.s1.a, myu.s1.b, myu.s1.c, myu.s1.d) + { + myu.s1.a++; + myu.s1.b++; + myu.s1.c++; + myu.s1.d++; + } + + assert (myu.s1.a == 1); + assert (myu.s1.b == 1); + assert (myu.s1.c == &c1 + 1); + assert (myu.s1.d == &d1 + 1); + + #pragma omp target map(myu.s2.a, myu.s2.b, myu.s2.c, myu.s2.d) + { + myu.s2.a++; + myu.s2.b++; + myu.s2.c++; + myu.s2.d++; + } + + assert (myu.s2.a == 1); + assert (myu.s2.b == 1); + assert (myu.s2.c == &c2 + 1); + assert (myu.s2.d == &d2 + 1); + + #pragma omp target map(to:myu.s3) \ + map(myu.s3->a, myu.s3->b, myu.s3->c, myu.s3->d) + { + myu.s3->a++; + myu.s3->b++; + myu.s3->c++; + myu.s3->d++; + } + + assert (myu.s3->a == 1); + assert (myu.s3->b == 1); + assert (myu.s3->c == &c3 + 1); + assert (myu.s3->d == &d3 + 1); + + #pragma omp target map(to:myu.s4) \ + map(myu.s4->a, myu.s4->b, myu.s4->c, myu.s4->d) + { + myu.s4->a++; + myu.s4->b++; + myu.s4->c++; + myu.s4->d++; + } + + assert (myu.s4->a == 1); + assert (myu.s4->b == 1); + assert (myu.s4->c == &c4 + 1); + assert (myu.s4->d == &d4 + 1); + + delete s4; + delete t4; +} +#endif + +#ifdef NONREF_COMPONENT_MEMBER_SLICE +template +void +nonref_component_member_slice (void) +{ + INIT_ST; + U myu(s1, t1, s2, t2, &s3, &t3, s4, t4); + + #pragma omp target map(myu.t1.a[2:5]) + { + myu.t1.a[2]++; + } + + #pragma omp target map(myu.t1.b[2:5]) + { + myu.t1.b[2]++; + } + + #pragma omp target enter data map(to: myu.t1.c) + + #pragma omp target map(myu.t1.c[2:5]) + { + myu.t1.c[2]++; + } + + #pragma omp target exit data map(release: myu.t1.c) + + #pragma omp target enter data map(to: myu.t1.d) + + #pragma omp target map(myu.t1.d[2:5]) + { + myu.t1.d[2]++; + } + + #pragma omp target exit data map(from: myu.t1.d) + + assert (myu.t1.a[2] == 1); + assert (myu.t1.b[2] == 3); + assert (myu.t1.c[2] == 3); + assert (myu.t1.d[2] == 3); + + #pragma omp target map(myu.t2.a[2:5]) + { + myu.t2.a[2]++; + } + + #pragma omp target map(myu.t2.b[2:5]) + { + myu.t2.b[2]++; + } + + #pragma omp target enter data map(to: myu.t2.c) + + #pragma omp target map(myu.t2.c[2:5]) + { + myu.t2.c[2]++; + } + + #pragma omp target exit data map(release: myu.t2.c) + + #pragma omp target enter data map(to: myu.t2.d) + + #pragma omp target map(myu.t2.d[2:5]) + { + myu.t2.d[2]++; + } + + #pragma omp target exit data map(release: myu.t2.d) + + assert (myu.t2.a[2] == 1); + assert (myu.t2.b[2] == 3); + assert (myu.t2.c[2] == 3); + assert (myu.t2.d[2] == 3); + + #pragma omp target enter data map(to: myu.t3) + + #pragma omp target map(myu.t3->a[2:5]) + { + myu.t3->a[2]++; + } + + #pragma omp target map(myu.t3->b[2:5]) + { + myu.t3->b[2]++; + } + + #pragma omp target enter data map(to: myu.t3->c) + + #pragma omp target map(myu.t3->c[2:5]) + { + myu.t3->c[2]++; + } + + #pragma omp target exit data map(release: myu.t3->c) + + #pragma omp target enter data map(to: myu.t3->d) + + #pragma omp target map(myu.t3->d[2:5]) + { + myu.t3->d[2]++; + } + + #pragma omp target exit data map(release: myu.t3, myu.t3->d) + + assert (myu.t3->a[2] == 1); + assert (myu.t3->b[2] == 3); + assert (myu.t3->c[2] == 3); + assert (myu.t3->d[2] == 3); + + #pragma omp target enter data map(to: myu.t4) + + #pragma omp target map(myu.t4->a[2:5]) + { + myu.t4->a[2]++; + } + + #pragma omp target map(myu.t4->b[2:5]) + { + myu.t4->b[2]++; + } + + #pragma omp target enter data map(to: myu.t4->c) + + #pragma omp target map(myu.t4->c[2:5]) + { + myu.t4->c[2]++; + } + + #pragma omp target exit data map(release: myu.t4->c) + + #pragma omp target enter data map(to: myu.t4->d) + + #pragma omp target map(myu.t4->d[2:5]) + { + myu.t4->d[2]++; + } + + #pragma omp target exit data map(release: myu.t4, myu.t4->d) + + assert (myu.t4->a[2] == 1); + assert (myu.t4->b[2] == 3); + assert (myu.t4->c[2] == 3); + assert (myu.t4->d[2] == 3); + + delete s4; + delete t4; +} +#endif + +#ifdef NONREF_COMPONENT_MEMBER_SLICE_BASEPTR +template +void +nonref_component_member_slice_baseptr (void) +{ + INIT_ST; + U myu(s1, t1, s2, t2, &s3, &t3, s4, t4); + + #pragma omp target map(to: myu.t1.c) map(myu.t1.c[2:5]) + { + myu.t1.c[2]++; + } + + #pragma omp target map(to: myu.t1.d) map(myu.t1.d[2:5]) + { + myu.t1.d[2]++; + } + + assert (myu.t1.c[2] == 2); + assert (myu.t1.d[2] == 2); + + #pragma omp target map(to: myu.t2.c) map(myu.t2.c[2:5]) + { + myu.t2.c[2]++; + } + + #pragma omp target map(to: myu.t2.d) map(myu.t2.d[2:5]) + { + myu.t2.d[2]++; + } + + assert (myu.t2.c[2] == 2); + assert (myu.t2.d[2] == 2); + + #pragma omp target map(to: myu.t3, myu.t3->c) map(myu.t3->c[2:5]) + { + myu.t3->c[2]++; + } + + #pragma omp target map(to: myu.t3, myu.t3->d) map(myu.t3->d[2:5]) + { + myu.t3->d[2]++; + } + + assert (myu.t3->c[2] == 2); + assert (myu.t3->d[2] == 2); + + #pragma omp target map(to: myu.t4, myu.t4->c) map(myu.t4->c[2:5]) + { + myu.t4->c[2]++; + } + + #pragma omp target map(to: myu.t4, myu.t4->d) map(myu.t4->d[2:5]) + { + myu.t4->d[2]++; + } + + assert (myu.t4->c[2] == 2); + assert (myu.t4->d[2] == 2); + + delete s4; + delete t4; +} +#endif + +#ifdef REF_COMPONENT_BASE +template +void +ref_component_base (void) +{ + INIT_ST; + U myu_real(s1, t1, s2, t2, &s3, &t3, s4, t4); + U &myu = myu_real; + + #pragma omp target map(myu.s1.a, myu.s1.b, myu.s1.c, myu.s1.d) + { + myu.s1.a++; + myu.s1.b++; + myu.s1.c++; + myu.s1.d++; + } + + assert (myu.s1.a == 1); + assert (myu.s1.b == 1); + assert (myu.s1.c == &c1 + 1); + assert (myu.s1.d == &d1 + 1); + + #pragma omp target map(myu.s2.a, myu.s2.b, myu.s2.c, myu.s2.d) + { + myu.s2.a++; + myu.s2.b++; + myu.s2.c++; + myu.s2.d++; + } + + assert (myu.s2.a == 1); + assert (myu.s2.b == 1); + assert (myu.s2.c == &c2 + 1); + assert (myu.s2.d == &d2 + 1); + + #pragma omp target map(to:myu.s3) \ + map(myu.s3->a, myu.s3->b, myu.s3->c, myu.s3->d) + { + myu.s3->a++; + myu.s3->b++; + myu.s3->c++; + myu.s3->d++; + } + + assert (myu.s3->a == 1); + assert (myu.s3->b == 1); + assert (myu.s3->c == &c3 + 1); + assert (myu.s3->d == &d3 + 1); + + #pragma omp target map(to:myu.s4) \ + map(myu.s4->a, myu.s4->b, myu.s4->c, myu.s4->d) + { + myu.s4->a++; + myu.s4->b++; + myu.s4->c++; + myu.s4->d++; + } + + assert (myu.s4->a == 1); + assert (myu.s4->b == 1); + assert (myu.s4->c == &c4 + 1); + assert (myu.s4->d == &d4 + 1); + + delete s4; + delete t4; +} +#endif + +#ifdef REF_COMPONENT_MEMBER_SLICE +template +void +ref_component_member_slice (void) +{ + INIT_ST; + U myu_real(s1, t1, s2, t2, &s3, &t3, s4, t4); + U &myu = myu_real; + + #pragma omp target map(myu.t1.a[2:5]) + { + myu.t1.a[2]++; + } + + #pragma omp target map(myu.t1.b[2:5]) + { + myu.t1.b[2]++; + } + + #pragma omp target enter data map(to: myu.t1.c) + + #pragma omp target map(myu.t1.c[2:5]) + { + myu.t1.c[2]++; + } + + #pragma omp target exit data map(release: myu.t1.c) + + #pragma omp target enter data map(to: myu.t1.d) + + #pragma omp target map(myu.t1.d[2:5]) + { + myu.t1.d[2]++; + } + + #pragma omp target exit data map(release: myu.t1.d) + + assert (myu.t1.a[2] == 1); + assert (myu.t1.b[2] == 3); + assert (myu.t1.c[2] == 3); + assert (myu.t1.d[2] == 3); + + #pragma omp target map(myu.t2.a[2:5]) + { + myu.t2.a[2]++; + } + + #pragma omp target map(myu.t2.b[2:5]) + { + myu.t2.b[2]++; + } + + #pragma omp target enter data map(to: myu.t2.c) + + #pragma omp target map(myu.t2.c[2:5]) + { + myu.t2.c[2]++; + } + + #pragma omp target exit data map(release: myu.t2.c) + + #pragma omp target enter data map(to: myu.t2.d) + + #pragma omp target map(myu.t2.d[2:5]) + { + myu.t2.d[2]++; + } + + #pragma omp target exit data map(release: myu.t2.d) + + assert (myu.t2.a[2] == 1); + assert (myu.t2.b[2] == 3); + assert (myu.t2.c[2] == 3); + assert (myu.t2.d[2] == 3); + + #pragma omp target enter data map(to: myu.t3) + + #pragma omp target map(myu.t3->a[2:5]) + { + myu.t3->a[2]++; + } + + #pragma omp target map(myu.t3->b[2:5]) + { + myu.t3->b[2]++; + } + + #pragma omp target enter data map(to: myu.t3->c) + + #pragma omp target map(myu.t3->c[2:5]) + { + myu.t3->c[2]++; + } + + #pragma omp target exit data map(release: myu.t3->c) + + #pragma omp target enter data map(to: myu.t3->d) + + #pragma omp target map(myu.t3->d[2:5]) + { + myu.t3->d[2]++; + } + + #pragma omp target exit data map(release: myu.t3, myu.t3->d) + + assert (myu.t3->a[2] == 1); + assert (myu.t3->b[2] == 3); + assert (myu.t3->c[2] == 3); + assert (myu.t3->d[2] == 3); + + #pragma omp target enter data map(to: myu.t4) + + #pragma omp target map(myu.t4->a[2:5]) + { + myu.t4->a[2]++; + } + + #pragma omp target map(myu.t4->b[2:5]) + { + myu.t4->b[2]++; + } + + #pragma omp target enter data map(to: myu.t4->c) + + #pragma omp target map(myu.t4->c[2:5]) + { + myu.t4->c[2]++; + } + + #pragma omp target exit data map(release: myu.t4->c) + + #pragma omp target enter data map(to: myu.t4->d) + + #pragma omp target map(myu.t4->d[2:5]) + { + myu.t4->d[2]++; + } + + #pragma omp target exit data map(release: myu.t4, myu.t4->d) + + assert (myu.t4->a[2] == 1); + assert (myu.t4->b[2] == 3); + assert (myu.t4->c[2] == 3); + assert (myu.t4->d[2] == 3); + + delete s4; + delete t4; +} +#endif + +#ifdef REF_COMPONENT_MEMBER_SLICE_BASEPTR +template +void +ref_component_member_slice_baseptr (void) +{ + INIT_ST; + U myu_real(s1, t1, s2, t2, &s3, &t3, s4, t4); + U &myu = myu_real; + + #pragma omp target map(to: myu.t1.c) map(myu.t1.c[2:5]) + { + myu.t1.c[2]++; + } + + #pragma omp target map(to: myu.t1.d) map(myu.t1.d[2:5]) + { + myu.t1.d[2]++; + } + + assert (myu.t1.c[2] == 2); + assert (myu.t1.d[2] == 2); + + #pragma omp target map(to: myu.t2.c) map(myu.t2.c[2:5]) + { + myu.t2.c[2]++; + } + + #pragma omp target map(to: myu.t2.d) map(myu.t2.d[2:5]) + { + myu.t2.d[2]++; + } + + assert (myu.t2.c[2] == 2); + assert (myu.t2.d[2] == 2); + + #pragma omp target map(to: myu.t3, myu.t3->c) map(myu.t3->c[2:5]) + { + myu.t3->c[2]++; + } + + #pragma omp target map(to: myu.t3, myu.t3->d) map(myu.t3->d[2:5]) + { + myu.t3->d[2]++; + } + + assert (myu.t3->c[2] == 2); + assert (myu.t3->d[2] == 2); + + #pragma omp target map(to: myu.t4, myu.t4->c) map(myu.t4->c[2:5]) + { + myu.t4->c[2]++; + } + + #pragma omp target map(to: myu.t4, myu.t4->d) map(myu.t4->d[2:5]) + { + myu.t4->d[2]++; + } + + assert (myu.t4->c[2] == 2); + assert (myu.t4->d[2] == 2); + + delete s4; + delete t4; +} +#endif + +#ifdef PTR_COMPONENT_BASE +template +void +ptr_component_base (void) +{ + INIT_ST; + U *myu = new U(s1, t1, s2, t2, &s3, &t3, s4, t4); + + #pragma omp target map(myu->s1.a, myu->s1.b, myu->s1.c, myu->s1.d) + { + myu->s1.a++; + myu->s1.b++; + myu->s1.c++; + myu->s1.d++; + } + + assert (myu->s1.a == 1); + assert (myu->s1.b == 1); + assert (myu->s1.c == &c1 + 1); + assert (myu->s1.d == &d1 + 1); + + #pragma omp target map(myu->s2.a, myu->s2.b, myu->s2.c, myu->s2.d) + { + myu->s2.a++; + myu->s2.b++; + myu->s2.c++; + myu->s2.d++; + } + + assert (myu->s2.a == 1); + assert (myu->s2.b == 1); + assert (myu->s2.c == &c2 + 1); + assert (myu->s2.d == &d2 + 1); + + #pragma omp target map(to:myu->s3) \ + map(myu->s3->a, myu->s3->b, myu->s3->c, myu->s3->d) + { + myu->s3->a++; + myu->s3->b++; + myu->s3->c++; + myu->s3->d++; + } + + assert (myu->s3->a == 1); + assert (myu->s3->b == 1); + assert (myu->s3->c == &c3 + 1); + assert (myu->s3->d == &d3 + 1); + + #pragma omp target map(to:myu->s4) \ + map(myu->s4->a, myu->s4->b, myu->s4->c, myu->s4->d) + { + myu->s4->a++; + myu->s4->b++; + myu->s4->c++; + myu->s4->d++; + } + + assert (myu->s4->a == 1); + assert (myu->s4->b == 1); + assert (myu->s4->c == &c4 + 1); + assert (myu->s4->d == &d4 + 1); + + delete s4; + delete t4; + delete myu; +} +#endif + +#ifdef PTR_COMPONENT_MEMBER_SLICE +template +void +ptr_component_member_slice (void) +{ + INIT_ST; + U *myu = new U(s1, t1, s2, t2, &s3, &t3, s4, t4); + + #pragma omp target map(myu->t1.a[2:5]) + { + myu->t1.a[2]++; + } + + #pragma omp target map(myu->t1.b[2:5]) + { + myu->t1.b[2]++; + } + + #pragma omp target enter data map(to: myu->t1.c) + + #pragma omp target map(myu->t1.c[2:5]) + { + myu->t1.c[2]++; + } + + #pragma omp target exit data map(release: myu->t1.c) + + #pragma omp target enter data map(to: myu->t1.d) + + #pragma omp target map(myu->t1.d[2:5]) + { + myu->t1.d[2]++; + } + + #pragma omp target exit data map(release: myu->t1.d) + + assert (myu->t1.a[2] == 1); + assert (myu->t1.b[2] == 3); + assert (myu->t1.c[2] == 3); + assert (myu->t1.d[2] == 3); + + #pragma omp target map(myu->t2.a[2:5]) + { + myu->t2.a[2]++; + } + + #pragma omp target map(myu->t2.b[2:5]) + { + myu->t2.b[2]++; + } + + #pragma omp target enter data map(to: myu->t2.c) + + #pragma omp target map(myu->t2.c[2:5]) + { + myu->t2.c[2]++; + } + + #pragma omp target exit data map(release: myu->t2.c) + + #pragma omp target enter data map(to: myu->t2.d) + + #pragma omp target map(myu->t2.d[2:5]) + { + myu->t2.d[2]++; + } + + #pragma omp target exit data map(release: myu->t2.d) + + assert (myu->t2.a[2] == 1); + assert (myu->t2.b[2] == 3); + assert (myu->t2.c[2] == 3); + assert (myu->t2.d[2] == 3); + + #pragma omp target enter data map(to: myu->t3) + + #pragma omp target map(myu->t3->a[2:5]) + { + myu->t3->a[2]++; + } + + #pragma omp target map(myu->t3->b[2:5]) + { + myu->t3->b[2]++; + } + + #pragma omp target enter data map(to: myu->t3->c) + + #pragma omp target map(myu->t3->c[2:5]) + { + myu->t3->c[2]++; + } + + #pragma omp target exit data map(release: myu->t3->c) + + #pragma omp target enter data map(to: myu->t3->d) + + #pragma omp target map(myu->t3->d[2:5]) + { + myu->t3->d[2]++; + } + + #pragma omp target exit data map(release: myu->t3, myu->t3->d) + + assert (myu->t3->a[2] == 1); + assert (myu->t3->b[2] == 3); + assert (myu->t3->c[2] == 3); + assert (myu->t3->d[2] == 3); + + #pragma omp target enter data map(to: myu->t4) + + #pragma omp target map(myu->t4->a[2:5]) + { + myu->t4->a[2]++; + } + + #pragma omp target map(myu->t4->b[2:5]) + { + myu->t4->b[2]++; + } + + #pragma omp target enter data map(to: myu->t4->c) + + #pragma omp target map(myu->t4->c[2:5]) + { + myu->t4->c[2]++; + } + + #pragma omp target exit data map(release: myu->t4->c) + + #pragma omp target enter data map(to: myu->t4->d) + + #pragma omp target map(myu->t4->d[2:5]) + { + myu->t4->d[2]++; + } + + #pragma omp target exit data map(release: myu->t4, myu->t4->d) + + assert (myu->t4->a[2] == 1); + assert (myu->t4->b[2] == 3); + assert (myu->t4->c[2] == 3); + assert (myu->t4->d[2] == 3); + + delete s4; + delete t4; + delete myu; +} +#endif + +#ifdef PTR_COMPONENT_MEMBER_SLICE_BASEPTR +template +void +ptr_component_member_slice_baseptr (void) +{ + INIT_ST; + U *myu = new U(s1, t1, s2, t2, &s3, &t3, s4, t4); + + /* Implicit firstprivate 'myu'. */ + #pragma omp target map(to: myu->t1.c) map(myu->t1.c[2:5]) + { + myu->t1.c[2]++; + } + + #pragma omp target map(to: myu->t1.d) map(myu->t1.d[2:5]) + { + myu->t1.d[2]++; + } + + assert (myu->t1.c[2] == 2); + assert (myu->t1.d[2] == 2); + + /* Explicitly-mapped 'myu'. */ + #pragma omp target map(to: myu, myu->t1.c) map(myu->t1.c[2:5]) + { + myu->t1.c[2]++; + } + + #pragma omp target map(to: myu, myu->t1.d) map(myu->t1.d[2:5]) + { + myu->t1.d[2]++; + } + + assert (myu->t1.c[2] == 4); + assert (myu->t1.d[2] == 4); + + /* Implicit firstprivate 'myu'. */ + #pragma omp target map(to: myu->t2.c) map(myu->t2.c[2:5]) + { + myu->t2.c[2]++; + } + + #pragma omp target map(to: myu->t2.d) map(myu->t2.d[2:5]) + { + myu->t2.d[2]++; + } + + assert (myu->t2.c[2] == 2); + assert (myu->t2.d[2] == 2); + + /* Explicitly-mapped 'myu'. */ + #pragma omp target map(to: myu, myu->t2.c) map(myu->t2.c[2:5]) + { + myu->t2.c[2]++; + } + + #pragma omp target map(to: myu, myu->t2.d) map(myu->t2.d[2:5]) + { + myu->t2.d[2]++; + } + + assert (myu->t2.c[2] == 4); + assert (myu->t2.d[2] == 4); + + /* Implicit firstprivate 'myu'. */ + #pragma omp target map(to: myu->t3, myu->t3->c) map(myu->t3->c[2:5]) + { + myu->t3->c[2]++; + } + + #pragma omp target map(to: myu->t3, myu->t3->d) map(myu->t3->d[2:5]) + { + myu->t3->d[2]++; + } + + assert (myu->t3->c[2] == 2); + assert (myu->t3->d[2] == 2); + + /* Explicitly-mapped 'myu'. */ + #pragma omp target map(to: myu, myu->t3, myu->t3->c) map(myu->t3->c[2:5]) + { + myu->t3->c[2]++; + } + + #pragma omp target map(to: myu, myu->t3, myu->t3->d) map(myu->t3->d[2:5]) + { + myu->t3->d[2]++; + } + + assert (myu->t3->c[2] == 4); + assert (myu->t3->d[2] == 4); + + /* Implicit firstprivate 'myu'. */ + #pragma omp target map(to: myu->t4, myu->t4->c) map(myu->t4->c[2:5]) + { + myu->t4->c[2]++; + } + + #pragma omp target map(to: myu->t4, myu->t4->d) map(myu->t4->d[2:5]) + { + myu->t4->d[2]++; + } + + assert (myu->t4->c[2] == 2); + assert (myu->t4->d[2] == 2); + + /* Explicitly-mapped 'myu'. */ + #pragma omp target map(to: myu, myu->t4, myu->t4->c) map(myu->t4->c[2:5]) + { + myu->t4->c[2]++; + } + + #pragma omp target map(to: myu, myu->t4, myu->t4->d) map(myu->t4->d[2:5]) + { + myu->t4->d[2]++; + } + + assert (myu->t4->c[2] == 4); + assert (myu->t4->d[2] == 4); + + delete s4; + delete t4; + delete myu; +} +#endif + +#ifdef REF2PTR_COMPONENT_BASE +template +void +ref2ptr_component_base (void) +{ + INIT_ST; + U *myu_ptr = new U(s1, t1, s2, t2, &s3, &t3, s4, t4); + U *&myu = myu_ptr; + + #pragma omp target map(myu->s1.a, myu->s1.b, myu->s1.c, myu->s1.d) + { + myu->s1.a++; + myu->s1.b++; + myu->s1.c++; + myu->s1.d++; + } + + assert (myu->s1.a == 1); + assert (myu->s1.b == 1); + assert (myu->s1.c == &c1 + 1); + assert (myu->s1.d == &d1 + 1); + + #pragma omp target map(myu->s2.a, myu->s2.b, myu->s2.c, myu->s2.d) + { + myu->s2.a++; + myu->s2.b++; + myu->s2.c++; + myu->s2.d++; + } + + assert (myu->s2.a == 1); + assert (myu->s2.b == 1); + assert (myu->s2.c == &c2 + 1); + assert (myu->s2.d == &d2 + 1); + + #pragma omp target map(to:myu->s3) \ + map(myu->s3->a, myu->s3->b, myu->s3->c, myu->s3->d) + { + myu->s3->a++; + myu->s3->b++; + myu->s3->c++; + myu->s3->d++; + } + + assert (myu->s3->a == 1); + assert (myu->s3->b == 1); + assert (myu->s3->c == &c3 + 1); + assert (myu->s3->d == &d3 + 1); + + #pragma omp target map(to:myu->s4) \ + map(myu->s4->a, myu->s4->b, myu->s4->c, myu->s4->d) + { + myu->s4->a++; + myu->s4->b++; + myu->s4->c++; + myu->s4->d++; + } + + assert (myu->s4->a == 1); + assert (myu->s4->b == 1); + assert (myu->s4->c == &c4 + 1); + assert (myu->s4->d == &d4 + 1); + + delete s4; + delete t4; + delete myu_ptr; +} +#endif + +#ifdef REF2PTR_COMPONENT_MEMBER_SLICE +template +void +ref2ptr_component_member_slice (void) +{ + INIT_ST; + U *myu_ptr = new U(s1, t1, s2, t2, &s3, &t3, s4, t4); + U *&myu = myu_ptr; + + #pragma omp target map(myu->t1.a[2:5]) + { + myu->t1.a[2]++; + } + + #pragma omp target map(myu->t1.b[2:5]) + { + myu->t1.b[2]++; + } + + #pragma omp target enter data map(to: myu->t1.c) + + #pragma omp target map(myu->t1.c[2:5]) + { + myu->t1.c[2]++; + } + + #pragma omp target exit data map(release: myu->t1.c) + + #pragma omp target enter data map(to: myu->t1.d) + + #pragma omp target map(myu->t1.d[2:5]) + { + myu->t1.d[2]++; + } + + #pragma omp target exit data map(release: myu->t1.d) + + assert (myu->t1.a[2] == 1); + assert (myu->t1.b[2] == 3); + assert (myu->t1.c[2] == 3); + assert (myu->t1.d[2] == 3); + + #pragma omp target map(myu->t2.a[2:5]) + { + myu->t2.a[2]++; + } + + #pragma omp target map(myu->t2.b[2:5]) + { + myu->t2.b[2]++; + } + + #pragma omp target enter data map(to: myu->t2.c) + + #pragma omp target map(myu->t2.c[2:5]) + { + myu->t2.c[2]++; + } + + #pragma omp target exit data map(release: myu->t2.c) + + #pragma omp target enter data map(to: myu->t2.d) + + #pragma omp target map(myu->t2.d[2:5]) + { + myu->t2.d[2]++; + } + + #pragma omp target exit data map(release: myu->t2.d) + + assert (myu->t2.a[2] == 1); + assert (myu->t2.b[2] == 3); + assert (myu->t2.c[2] == 3); + assert (myu->t2.d[2] == 3); + + #pragma omp target enter data map(to: myu->t3) + + #pragma omp target map(myu->t3->a[2:5]) + { + myu->t3->a[2]++; + } + + #pragma omp target map(myu->t3->b[2:5]) + { + myu->t3->b[2]++; + } + + #pragma omp target enter data map(to: myu->t3->c) + + #pragma omp target map(myu->t3->c[2:5]) + { + myu->t3->c[2]++; + } + + #pragma omp target exit data map(release: myu->t3->c) + + #pragma omp target enter data map(to: myu->t3->d) + + #pragma omp target map(myu->t3->d[2:5]) + { + myu->t3->d[2]++; + } + + #pragma omp target exit data map(release: myu->t3, myu->t3->d) + + assert (myu->t3->a[2] == 1); + assert (myu->t3->b[2] == 3); + assert (myu->t3->c[2] == 3); + assert (myu->t3->d[2] == 3); + + #pragma omp target enter data map(to: myu->t4) + + #pragma omp target map(myu->t4->a[2:5]) + { + myu->t4->a[2]++; + } + + #pragma omp target map(myu->t4->b[2:5]) + { + myu->t4->b[2]++; + } + + #pragma omp target enter data map(to: myu->t4->c) + + #pragma omp target map(myu->t4->c[2:5]) + { + myu->t4->c[2]++; + } + + #pragma omp target exit data map(release: myu->t4->c) + + #pragma omp target enter data map(to: myu->t4->d) + + #pragma omp target map(myu->t4->d[2:5]) + { + myu->t4->d[2]++; + } + + #pragma omp target exit data map(release: myu->t4, myu->t4->d) + + assert (myu->t4->a[2] == 1); + assert (myu->t4->b[2] == 3); + assert (myu->t4->c[2] == 3); + assert (myu->t4->d[2] == 3); + + delete s4; + delete t4; + delete myu_ptr; +} +#endif + +#ifdef REF2PTR_COMPONENT_MEMBER_SLICE_BASEPTR +template +void +ref2ptr_component_member_slice_baseptr (void) +{ + INIT_ST; + U *myu_ptr = new U(s1, t1, s2, t2, &s3, &t3, s4, t4); + U *&myu = myu_ptr; + + /* Implicit firstprivate 'myu'. */ + #pragma omp target map(to: myu->t1.c) map(myu->t1.c[2:5]) + { + myu->t1.c[2]++; + } + + #pragma omp target map(to: myu->t1.d) map(myu->t1.d[2:5]) + { + myu->t1.d[2]++; + } + + assert (myu->t1.c[2] == 2); + assert (myu->t1.d[2] == 2); + + /* Explicitly-mapped 'myu'. */ + #pragma omp target map(to: myu, myu->t1.c) map(myu->t1.c[2:5]) + { + myu->t1.c[2]++; + } + + #pragma omp target map(to: myu, myu->t1.d) map(myu->t1.d[2:5]) + { + myu->t1.d[2]++; + } + + assert (myu->t1.c[2] == 4); + assert (myu->t1.d[2] == 4); + + /* Implicit firstprivate 'myu'. */ + #pragma omp target map(to: myu->t2.c) map(myu->t2.c[2:5]) + { + myu->t2.c[2]++; + } + + #pragma omp target map(to: myu->t2.d) map(myu->t2.d[2:5]) + { + myu->t2.d[2]++; + } + + assert (myu->t2.c[2] == 2); + assert (myu->t2.d[2] == 2); + + /* Explicitly-mapped 'myu'. */ + #pragma omp target map(to: myu, myu->t2.c) map(myu->t2.c[2:5]) + { + myu->t2.c[2]++; + } + + #pragma omp target map(to: myu, myu->t2.d) map(myu->t2.d[2:5]) + { + myu->t2.d[2]++; + } + + assert (myu->t2.c[2] == 4); + assert (myu->t2.d[2] == 4); + + /* Implicit firstprivate 'myu'. */ + #pragma omp target map(to: myu->t3, myu->t3->c) map(myu->t3->c[2:5]) + { + myu->t3->c[2]++; + } + + #pragma omp target map(to: myu->t3, myu->t3->d) map(myu->t3->d[2:5]) + { + myu->t3->d[2]++; + } + + assert (myu->t3->c[2] == 2); + assert (myu->t3->d[2] == 2); + + /* Explicitly-mapped 'myu'. */ + #pragma omp target map(to: myu, myu->t3, myu->t3->c) map(myu->t3->c[2:5]) + { + myu->t3->c[2]++; + } + + #pragma omp target map(to: myu, myu->t3, myu->t3->d) map(myu->t3->d[2:5]) + { + myu->t3->d[2]++; + } + + assert (myu->t3->c[2] == 4); + assert (myu->t3->d[2] == 4); + + /* Implicit firstprivate 'myu'. */ + #pragma omp target map(to: myu->t4, myu->t4->c) map(myu->t4->c[2:5]) + { + myu->t4->c[2]++; + } + + #pragma omp target map(to: myu->t4, myu->t4->d) map(myu->t4->d[2:5]) + { + myu->t4->d[2]++; + } + + assert (myu->t4->c[2] == 2); + assert (myu->t4->d[2] == 2); + + /* Explicitly-mapped 'myu'. */ + #pragma omp target map(to: myu, myu->t4, myu->t4->c) map(myu->t4->c[2:5]) + { + myu->t4->c[2]++; + } + + #pragma omp target map(to: myu, myu->t4, myu->t4->d) map(myu->t4->d[2:5]) + { + myu->t4->d[2]++; + } + + assert (myu->t4->c[2] == 4); + assert (myu->t4->d[2] == 4); + + delete s4; + delete t4; + delete myu_ptr; +} +#endif + +int main (int argc, char *argv[]) +{ +#ifdef MAP_DECLS + map_decls<4> (); +#endif + +#ifdef NONREF_DECL_BASE + nonref_decl_base (); +#endif +#ifdef REF_DECL_BASE + ref_decl_base (); +#endif +#ifdef PTR_DECL_BASE + ptr_decl_base (); +#endif +#ifdef REF2PTR_DECL_BASE + ref2ptr_decl_base (); +#endif + +#ifdef ARRAY_DECL_BASE + array_decl_base (); +#endif +#ifdef REF2ARRAY_DECL_BASE + ref2array_decl_base (); +#endif +#ifdef PTR_OFFSET_DECL_BASE + ptr_offset_decl_base (); +#endif +#ifdef REF2PTR_OFFSET_DECL_BASE + ref2ptr_offset_decl_base (); +#endif + +#ifdef MAP_SECTIONS + map_sections (); +#endif + +#ifdef NONREF_DECL_MEMBER_SLICE + nonref_decl_member_slice (); +#endif +#ifdef NONREF_DECL_MEMBER_SLICE_BASEPTR + nonref_decl_member_slice_baseptr (); +#endif +#ifdef REF_DECL_MEMBER_SLICE + ref_decl_member_slice (); +#endif +#ifdef REF_DECL_MEMBER_SLICE_BASEPTR + ref_decl_member_slice_baseptr (); +#endif +#ifdef PTR_DECL_MEMBER_SLICE + ptr_decl_member_slice (); +#endif +#ifdef PTR_DECL_MEMBER_SLICE_BASEPTR + ptr_decl_member_slice_baseptr (); +#endif +#ifdef REF2PTR_DECL_MEMBER_SLICE + ref2ptr_decl_member_slice (); +#endif +#ifdef REF2PTR_DECL_MEMBER_SLICE_BASEPTR + ref2ptr_decl_member_slice_baseptr (); +#endif + +#ifdef ARRAY_DECL_MEMBER_SLICE + array_decl_member_slice (); +#endif +#ifdef ARRAY_DECL_MEMBER_SLICE_BASEPTR + array_decl_member_slice_baseptr (); +#endif +#ifdef REF2ARRAY_DECL_MEMBER_SLICE + ref2array_decl_member_slice (); +#endif +#ifdef REF2ARRAY_DECL_MEMBER_SLICE_BASEPTR + ref2array_decl_member_slice_baseptr (); +#endif +#ifdef PTR_OFFSET_DECL_MEMBER_SLICE + ptr_offset_decl_member_slice (); +#endif +#ifdef PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR + ptr_offset_decl_member_slice_baseptr (); +#endif +#ifdef REF2PTR_OFFSET_DECL_MEMBER_SLICE + ref2ptr_offset_decl_member_slice (); +#endif +#ifdef REF2PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR + ref2ptr_offset_decl_member_slice_baseptr (); +#endif + +#ifdef PTRARRAY_DECL_MEMBER_SLICE + ptrarray_decl_member_slice (); +#endif +#ifdef PTRARRAY_DECL_MEMBER_SLICE_BASEPTR + ptrarray_decl_member_slice_baseptr (); +#endif +#ifdef REF2PTRARRAY_DECL_MEMBER_SLICE + ref2ptrarray_decl_member_slice (); +#endif +#ifdef REF2PTRARRAY_DECL_MEMBER_SLICE_BASEPTR + ref2ptrarray_decl_member_slice_baseptr (); +#endif +#ifdef PTRPTR_OFFSET_DECL_MEMBER_SLICE + ptrptr_offset_decl_member_slice (); +#endif +#ifdef PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR + ptrptr_offset_decl_member_slice_baseptr (); +#endif +#ifdef REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE + ref2ptrptr_offset_decl_member_slice (); +#endif +#ifdef REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR + ref2ptrptr_offset_decl_member_slice_baseptr (); +#endif + +#ifdef NONREF_COMPONENT_BASE + nonref_component_base (); +#endif +#ifdef NONREF_COMPONENT_MEMBER_SLICE + nonref_component_member_slice (); +#endif +#ifdef NONREF_COMPONENT_MEMBER_SLICE_BASEPTR + nonref_component_member_slice_baseptr (); +#endif + +#ifdef REF_COMPONENT_BASE + ref_component_base (); +#endif +#ifdef REF_COMPONENT_MEMBER_SLICE + ref_component_member_slice (); +#endif +#ifdef REF_COMPONENT_MEMBER_SLICE_BASEPTR + ref_component_member_slice_baseptr (); +#endif + +#ifdef PTR_COMPONENT_BASE + ptr_component_base (); +#endif +#ifdef PTR_COMPONENT_MEMBER_SLICE + ptr_component_member_slice (); +#endif +#ifdef PTR_COMPONENT_MEMBER_SLICE_BASEPTR + ptr_component_member_slice_baseptr (); +#endif + +#ifdef REF2PTR_COMPONENT_BASE + ref2ptr_component_base (); +#endif +#ifdef REF2PTR_COMPONENT_MEMBER_SLICE + ref2ptr_component_member_slice (); +#endif +#ifdef REF2PTR_COMPONENT_MEMBER_SLICE_BASEPTR + ref2ptr_component_member_slice_baseptr (); +#endif + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/ind-base-1.C b/libgomp/testsuite/libgomp.c++/ind-base-1.C new file mode 100644 index 000000000000..4566854e60ae --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/ind-base-1.C @@ -0,0 +1,162 @@ +// { dg-do run } +// { dg-options "-fopenmp" } + +#include + +struct S +{ + int x[10]; +}; + +struct T +{ + struct S *s; +}; + +struct U +{ + struct T *t; +}; + +void +foo_siblist (void) +{ + U *u = new U; + u->t = new T; + u->t->s = new S; + for (int i = 0; i < 10; i++) + u->t->s->x[i] = 0; +#pragma omp target map(u->t, *(u->t), u->t->s, *u->t->s) + for (int i = 0; i < 10; i++) + u->t->s->x[i] = i * 3; + for (int i = 0; i < 10; i++) + assert (u->t->s->x[i] == i * 3); + delete u->t->s; + delete u->t; + delete u; +} + +void +foo (void) +{ + U *u = new U; + u->t = new T; + u->t->s = new S; + for (int i = 0; i < 10; i++) + u->t->s->x[i] = 0; +#pragma omp target map(*u, u->t, *(u->t), u->t->s, *u->t->s) + for (int i = 0; i < 10; i++) + u->t->s->x[i] = i * 3; + for (int i = 0; i < 10; i++) + assert (u->t->s->x[i] == i * 3); + delete u->t->s; + delete u->t; + delete u; +} + +void +foo_tofrom (void) +{ + U *u = new U; + u->t = new T; + u->t->s = new S; + for (int i = 0; i < 10; i++) + u->t->s->x[i] = 0; +#pragma omp target map(u, *u, u->t, *(u->t), u->t->s, *u->t->s) + for (int i = 0; i < 10; i++) + u->t->s->x[i] = i * 3; + for (int i = 0; i < 10; i++) + assert (u->t->s->x[i] == i * 3); + delete u->t->s; + delete u->t; + delete u; +} + +void +bar (void) +{ + U *u = new U; + U **up = &u; + u->t = new T; + u->t->s = new S; + for (int i = 0; i < 10; i++) + (*up)->t->s->x[i] = 0; +#pragma omp target map(*up, (*up)->t, *(*up)->t, (*up)->t->s, *(*up)->t->s) + for (int i = 0; i < 10; i++) + (*up)->t->s->x[i] = i * 3; + for (int i = 0; i < 10; i++) + assert ((*up)->t->s->x[i] == i * 3); + delete u->t->s; + delete u->t; + delete u; +} + +void +bar_pp (void) +{ + U *u = new U; + U **up = &u; + u->t = new T; + u->t->s = new S; + for (int i = 0; i < 10; i++) + (*up)->t->s->x[i] = 0; +#pragma omp target map(*up, **up, (*up)->t, *(*up)->t, (*up)->t->s, *(*up)->t->s) + for (int i = 0; i < 10; i++) + (*up)->t->s->x[i] = i * 3; + for (int i = 0; i < 10; i++) + assert ((*up)->t->s->x[i] == i * 3); + delete u->t->s; + delete u->t; + delete u; +} + +void +bar_tofrom (void) +{ + U *u = new U; + U **up = &u; + u->t = new T; + u->t->s = new S; + for (int i = 0; i < 10; i++) + (*up)->t->s->x[i] = 0; +#pragma omp target map(*up, up, (*up)->t, *(*up)->t, (*up)->t->s, *(*up)->t->s) + for (int i = 0; i < 10; i++) + (*up)->t->s->x[i] = i * 3; + for (int i = 0; i < 10; i++) + assert ((*up)->t->s->x[i] == i * 3); + delete u->t->s; + delete u->t; + delete u; +} + +void +bar_tofrom_pp (void) +{ + U *u = new U; + U **up = &u; + u->t = new T; + u->t->s = new S; + for (int i = 0; i < 10; i++) + (*up)->t->s->x[i] = 0; +#pragma omp target map(**up, *up, up, (*up)->t, *(*up)->t, (*up)->t->s, \ + *(*up)->t->s) + for (int i = 0; i < 10; i++) + (*up)->t->s->x[i] = i * 3; + for (int i = 0; i < 10; i++) + assert ((*up)->t->s->x[i] == i * 3); + delete u->t->s; + delete u->t; + delete u; +} + +int main (int argc, char *argv[]) +{ + foo_siblist (); + foo (); + foo_tofrom (); + bar (); + bar_pp (); + bar_tofrom (); + bar_tofrom_pp (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/ind-base-2.C b/libgomp/testsuite/libgomp.c++/ind-base-2.C new file mode 100644 index 000000000000..4c05c2ef8f61 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/ind-base-2.C @@ -0,0 +1,93 @@ +// { dg-do run } +// { dg-options "-fopenmp" } + +#include + +struct S +{ + int x[10]; +}; + +struct T +{ + struct S ***s; +}; + +struct U +{ + struct T **t; +}; + +void +foo (void) +{ + U *u = new U; + T *real_t = new T; + S *real_s = new S; + T **t_pp = &real_t; + S **s_pp = &real_s; + S ***s_ppp = &s_pp; + u->t = t_pp; + (*u->t)->s = s_ppp; + for (int i = 0; i < 10; i++) + (**((*u->t)->s))->x[i] = 0; +#pragma omp target map(u->t, *u->t, (*u->t)->s, *(*u->t)->s, **(*u->t)->s, \ + (**(*u->t)->s)->x[0:10]) + for (int i = 0; i < 10; i++) + (**((*u->t)->s))->x[i] = i * 3; + for (int i = 0; i < 10; i++) + assert ((**((*u->t)->s))->x[i] == i * 3); + delete real_s; + delete real_t; + delete u; +} + +template +struct St +{ + X x[10]; +}; + +template +struct Tt +{ + X ***s; +}; + +template +struct Ut +{ + X **t; +}; + +template +void +tfoo (void) +{ + Ut > > *u = new Ut > >; + Tt > *real_t = new Tt >; + St *real_s = new St; + Tt > **t_pp = &real_t; + St **s_pp = &real_s; + St ***s_ppp = &s_pp; + u->t = t_pp; + (*u->t)->s = s_ppp; + for (int i = 0; i < 10; i++) + (**((*u->t)->s))->x[i] = 0; +#pragma omp target map(u->t, *u->t, (*u->t)->s, *(*u->t)->s, **(*u->t)->s, \ + (**(*u->t)->s)->x[0:10]) + for (int i = 0; i < 10; i++) + (**((*u->t)->s))->x[i] = i * 3; + for (int i = 0; i < 10; i++) + assert ((**((*u->t)->s))->x[i] == i * 3); + delete real_s; + delete real_t; + delete u; +} + +int main (int argc, char *argv[]) +{ + foo (); + tfoo (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/lvalue-tofrom-1.C b/libgomp/testsuite/libgomp.c++/lvalue-tofrom-1.C new file mode 100644 index 000000000000..643cfdb6e28b --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/lvalue-tofrom-1.C @@ -0,0 +1,75 @@ +#include +#include + +static int lo() +{ + return 30; +} + +static int len() +{ + return 10; +} + +template +void foo () +{ + T arr[100]; + T *ptr; + + memset (arr, '\0', sizeof arr); + +#pragma omp target enter data map(to: arr[0:100]) + + for (int i = 0; i < 100; i++) + arr[i] = i; + + ptr = &arr[10]; + +#pragma omp target update to(*ptr) + + for (int i = lo (); i < lo () + len (); i++) + arr[i] = i * 2; + +#pragma omp target update to(arr[lo():len()]) + +#pragma omp target exit data map(from: arr[0:100]) + + assert (arr[10] == 10); + for (int i = lo (); i < lo () + len (); i++) + assert (arr[i] == i * 2); +} + +int +main () +{ + char arr[100]; + char *ptr; + + memset (arr, '\0', sizeof arr); + +#pragma omp target enter data map(to: arr[0:100]) + + for (int i = 0; i < 100; i++) + arr[i] = i; + + ptr = &arr[10]; + +#pragma omp target update to(*ptr) + + for (int i = lo (); i < lo () + len (); i++) + arr[i] = i * 2; + +#pragma omp target update to(arr[lo():len()]) + +#pragma omp target exit data map(from: arr[0:100]) + + assert (arr[10] == 10); + for (int i = lo (); i < lo () + len (); i++) + assert (arr[i] == i * 2); + + foo (); + + return 0; +} + diff --git a/libgomp/testsuite/libgomp.c++/lvalue-tofrom-2.C b/libgomp/testsuite/libgomp.c++/lvalue-tofrom-2.C new file mode 100644 index 000000000000..adc493b1315c --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/lvalue-tofrom-2.C @@ -0,0 +1,71 @@ +#include +#include +#include + +template +struct t_array_wrapper { + T *data; + unsigned int length; +}; + +template +void foo() +{ + struct t_array_wrapper aw; + + aw.data = new T[100]; + aw.length = 100; + +#pragma omp target enter data map(to: aw.data, aw.length) \ + map(alloc: aw.data[0:aw.length]) + +#pragma omp target + for (int i = 0; i < aw.length; i++) + aw.data[i] = i; + +#pragma omp target update from(aw.data[:aw.length]) + +#pragma omp target exit data map(delete: aw.data, aw.length, \ + aw.data[0:aw.length]) + + for (int i = 0; i < aw.length; i++) + assert (aw.data[i] == i); + + delete[] aw.data; +} + +struct array_wrapper { + int *data; + unsigned int length; +}; + +int +main () +{ + struct array_wrapper aw; + + aw.data = new int[100]; + aw.length = 100; + +#pragma omp target enter data map(to: aw.data, aw.length) \ + map(alloc: aw.data[0:aw.length]) + +#pragma omp target + for (int i = 0; i < aw.length; i++) + aw.data[i] = i; + +#pragma omp target update from(aw.data[:aw.length]) + +#pragma omp target exit data map(delete: aw.data, aw.length, \ + aw.data[0:aw.length]) + + for (int i = 0; i < aw.length; i++) + assert (aw.data[i] == i); + + delete[] aw.data; + + foo (); + + return 0; +} + diff --git a/libgomp/testsuite/libgomp.c++/map-comma-1.C b/libgomp/testsuite/libgomp.c++/map-comma-1.C new file mode 100644 index 000000000000..ee03c5ac1aa3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/map-comma-1.C @@ -0,0 +1,15 @@ +/* { dg-do run } */ + +#include + +int main (int argc, char *argv[]) +{ + int a = 5, b = 7; +#pragma omp target map((a, b)) + { + a++; + b++; + } + assert (a == 5 && b == 8); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/map-rvalue-ref-1.C b/libgomp/testsuite/libgomp.c++/map-rvalue-ref-1.C new file mode 100644 index 000000000000..93811da40007 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/map-rvalue-ref-1.C @@ -0,0 +1,22 @@ +/* { dg-do run } */ + +#include + +int foo (int &&x) +{ + int y; +#pragma omp target map(x, y) + { + x++; + y = x; + } + return y; +} + +int main (int argc, char *argv[]) +{ + int y = 5; + y = foo (y + 3); + assert (y == 9); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/struct-ref-1.C b/libgomp/testsuite/libgomp.c++/struct-ref-1.C new file mode 100644 index 000000000000..d38746500178 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/struct-ref-1.C @@ -0,0 +1,97 @@ +// { dg-do run } +// { dg-options "-fopenmp" } + +#include + +struct S +{ + int x[10]; +}; + +void +foo (S *s, int x) +{ + S *&r = s; + for (int i = 0; i < x; i++) + s[0].x[i] = s[1].x[i] = 0; + #pragma omp target map (s, x) + ; + #pragma omp target map (s[0], x) + for (int i = 0; i < x; i++) + s[0].x[i] = i; + #pragma omp target map (s[1], x) + for (int i = 0; i < x; i++) + s[1].x[i] = i * 2; + for (int i = 0; i < x; i++) + { + assert (s[0].x[i] == i); + assert (s[1].x[i] == i * 2); + s[0].x[i] = 0; + s[1].x[i] = 0; + } + #pragma omp target map (r, x) + ; + #pragma omp target map (r[0], x) + for (int i = 0; i < x; i++) + r[0].x[i] = i; + #pragma omp target map (r[1], x) + for (int i = 0; i < x; i++) + r[1].x[i] = i * 2; + for (int i = 0; i < x; i++) + { + assert (r[0].x[i] == i); + assert (r[1].x[i] == i * 2); + } +} + +template +struct T +{ + int x[N]; +}; + +template +void +bar (T *t, int x) +{ + T *&r = t; + for (int i = 0; i < x; i++) + t[0].x[i] = t[1].x[i] = 0; + #pragma omp target map (t, x) + ; + #pragma omp target map (t[0], x) + for (int i = 0; i < x; i++) + t[0].x[i] = i; + #pragma omp target map (t[1], x) + for (int i = 0; i < x; i++) + t[1].x[i] = i * 2; + for (int i = 0; i < x; i++) + { + assert (t[0].x[i] == i); + assert (t[1].x[i] == i * 2); + t[0].x[i] = 0; + t[1].x[i] = 0; + } + #pragma omp target map (r, x) + ; + #pragma omp target map (r[0], x) + for (int i = 0; i < x; i++) + r[0].x[i] = i; + #pragma omp target map (r[1], x) + for (int i = 0; i < x; i++) + r[1].x[i] = i * 2; + for (int i = 0; i < x; i++) + { + assert (r[0].x[i] == i); + assert (r[1].x[i] == i * 2); + } +} + +int main (int argc, char *argv[]) +{ + S s[2]; + foo (s, 10); + T<10> t[2]; + bar (t, 10); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-field-1.c b/libgomp/testsuite/libgomp.c-c++-common/array-field-1.c new file mode 100644 index 000000000000..6dd8b5c48e1e --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-field-1.c @@ -0,0 +1,35 @@ +/* { dg-do run } */ + +#include +#include +#include + +#define N 16 + +struct Z { + int *ptr; + int arr[N]; + int c; +}; + +int main (int argc, char *argv[]) +{ + struct Z *myz; + myz = (struct Z *) calloc (1, sizeof *myz); + +#pragma omp target map(tofrom:myz->arr[0:N], myz->c) + { + for (int i = 0; i < N; i++) + myz->arr[i]++; + myz->c++; + } + + for (int i = 0; i < N; i++) + assert (myz->arr[i] == 1); + assert (myz->c == 1); + + free (myz); + + return 0; +} + diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-of-struct-1.c b/libgomp/testsuite/libgomp.c-c++-common/array-of-struct-1.c new file mode 100644 index 000000000000..11215b1df7a3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-of-struct-1.c @@ -0,0 +1,65 @@ +/* { dg-do run } */ + +#include +#include +#include + +#define N 16 + +/* NOTE: This test is the same as array-of-struct-2.c, except the fields of + this struct are in a different order. */ + +struct Z { + int arr[N]; + int *ptr; + int c; +}; + +void +foo (struct Z *zarr, int len) +{ +#pragma omp target map(to:zarr, zarr[5].ptr) map(tofrom:zarr[5].ptr[0:len]) + { + for (int i = 0; i < len; i++) + zarr[5].ptr[i]++; + } + +#pragma omp target map(to:zarr) map(tofrom:zarr[4].arr[0:len]) + { + for (int i = 0; i < len; i++) + zarr[4].arr[i]++; + } + +#pragma omp target map (to:zarr[3].ptr) map(tofrom:zarr[3].ptr[0:len]) + { + for (int i = 0; i < len; i++) + zarr[3].ptr[i]++; + } + +#pragma omp target map(tofrom:zarr[2].arr[0:len]) + { + for (int i = 0; i < len; i++) + zarr[2].arr[i]++; + } +} + +int main (int argc, char *argv[]) +{ + struct Z zs[10]; + memset (zs, 0, sizeof zs); + + for (int i = 0; i < 10; i++) + zs[i].ptr = (int *) calloc (N, sizeof (int)); + + foo (zs, N); + + for (int i = 0; i < N; i++) + { + assert (zs[2].arr[i] == 1); + assert (zs[4].arr[i] == 1); + assert (zs[3].ptr[i] == 1); + assert (zs[5].ptr[i] == 1); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-of-struct-2.c b/libgomp/testsuite/libgomp.c-c++-common/array-of-struct-2.c new file mode 100644 index 000000000000..d5d74b8c07d6 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-of-struct-2.c @@ -0,0 +1,65 @@ +/* { dg-do run } */ + +#include +#include +#include + +#define N 16 + +/* NOTE: This test is the same as array-of-struct-1.c, except the fields of + this struct are in a different order. */ + +struct Z { + int *ptr; + int arr[N]; + int c; +}; + +void +foo (struct Z *zarr, int len) +{ +#pragma omp target map(to:zarr, zarr[5].ptr) map(tofrom:zarr[5].ptr[0:len]) + { + for (int i = 0; i < len; i++) + zarr[5].ptr[i]++; + } + +#pragma omp target map(to:zarr) map(tofrom:zarr[4].arr[0:len]) + { + for (int i = 0; i < len; i++) + zarr[4].arr[i]++; + } + +#pragma omp target map (to:zarr[3].ptr) map(tofrom:zarr[3].ptr[0:len]) + { + for (int i = 0; i < len; i++) + zarr[3].ptr[i]++; + } + +#pragma omp target map(tofrom:zarr[2].arr[0:len]) + { + for (int i = 0; i < len; i++) + zarr[2].arr[i]++; + } +} + +int main (int argc, char *argv[]) +{ + struct Z zs[10]; + memset (zs, 0, sizeof zs); + + for (int i = 0; i < 10; i++) + zs[i].ptr = (int *) calloc (N, sizeof (int)); + + foo (zs, N); + + for (int i = 0; i < N; i++) + { + assert (zs[2].arr[i] == 1); + assert (zs[4].arr[i] == 1); + assert (zs[3].ptr[i] == 1); + assert (zs[5].ptr[i] == 1); + } + + return 0; +}