From patchwork Mon Jul 3 21:33:19 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1802944 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=server2.sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Received: from server2.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 4QvzkT163bz20Pf for ; Tue, 4 Jul 2023 07:36:01 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 34F5D382BC3C for ; Mon, 3 Jul 2023 21:35:59 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 12E08385770F; Mon, 3 Jul 2023 21:34:38 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 12E08385770F 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="6.01,178,1684828800"; d="scan'208";a="10838083" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 03 Jul 2023 13:34:36 -0800 IronPort-SDR: 6QV3luBQv3cRA/YSlvFsGSAQtqGbuK99egxAsfxMYyrExIMC+O4wj6Y5m+Uh3wC3wF0NGrCoSU B9tuungMYIuqcwHOcOezpRdkHQIyURRhwY2Z3fYzeyCMw071NOfejEErgl5ia7tlrMGpnlyU0Y SHpmZLfmFvl24jqdW+EIYGGrg5VdvATb2HPOlPS+SXiSVshV8VMTcEqDmN0TyTMr7WxvTaQNN8 G9JjhpSuoPVNM2CuqDZrFgq/X0xeXKh4uMfaPlKDfzF+7h1IyxpaXACnMTOTWmUb4RA+TmFgdE kek= From: Julian Brown To: CC: , , Subject: [PATCH 5/5] OpenMP: Array shaping operator and strided "target update" for C Date: Mon, 3 Jul 2023 21:33:19 +0000 Message-ID: <8e7811a8e3679d60fb804ac4df7e7523e0c27364.1688418868.git.julian@codesourcery.com> X-Mailer: git-send-email 2.25.1 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" Following the similar support for C++ and Fortran, here is the C implementation for the OpenMP 5.0 array-shaping operator, and for strided and rectangular updates for "target update". Much of the implementation is shared with the C++ support added earlier in this patch series. Some details of parsing necessarily differ for C, but the general ideas are the same. 2023-07-03 Julian Brown gcc/c/ * c-parser.cc (c_parser_braced_init): Disallow array-shaping operator in braced init. (c_parser_conditional_expression): Disallow array-shaping operator in conditional expression. (c_parser_cast_expression): Add array-shaping operator support. (c_parser_postfix_expression): Disallow array-shaping operator in statement expressions. (c_parser_postfix_expression_after_primary): Add OpenMP array section stride support. (c_parser_expr_list): Disallow array-shaping operator in expression lists. (c_array_type_nelts_top, c_array_type_nelts_total): New functions. (c_parser_omp_variable_list): Support array-shaping operator. (c_parser_omp_target_update): Recognize GOMP_MAP_TO_GRID and GOMP_MAP_FROM_GRID map kinds as well as OMP_CLAUSE_TO/OMP_CLAUSE_FROM. * c-tree.h (c_omp_array_shaping_op_p, c_omp_has_array_shape_p): New extern declarations. (create_omp_arrayshape_type): Add prototype. * c-typeck.cc (c_omp_array_shaping_op_p, c_omp_has_array_shape_p): New globals. (build_omp_array_section): Permit integral types, not just integer constants, when creating array types for array sections. (create_omp_arrayshape_type): New function. (handle_omp_array_sections_1): Add DISCONTIGUOUS parameter. Add strided/rectangular array section support. (omp_array_section_low_bound): New function. (handle_omp_array_sections): Add DISCONTIGUOUS parameter. Add strided/rectangular array section support. (c_finish_omp_clauses): Update calls to handle_omp_array_sections. Handle discontiguous updates. gcc/testsuite/ * gcc.dg/gomp/bad-array-shaping-c-1.c: New test. * gcc.dg/gomp/bad-array-shaping-c-2.c: New test. * gcc.dg/gomp/bad-array-shaping-c-3.c: New test. * gcc.dg/gomp/bad-array-shaping-c-4.c: New test. * gcc.dg/gomp/bad-array-shaping-c-5.c: New test. * gcc.dg/gomp/bad-array-shaping-c-6.c: New test. * gcc.dg/gomp/bad-array-shaping-c-7.c: New test. libgomp/ * testsuite/libgomp.c/array-shaping-1.c: New test. * testsuite/libgomp.c/array-shaping-2.c: New test. * testsuite/libgomp.c/array-shaping-3.c: New test. * testsuite/libgomp.c/array-shaping-4.c: New test. * testsuite/libgomp.c/array-shaping-5.c: New test. * testsuite/libgomp.c/array-shaping-6.c: New test. --- gcc/c/c-parser.cc | 301 +++++++++++++++++- gcc/c/c-tree.h | 4 + gcc/c/c-typeck.cc | 241 ++++++++++++-- .../gcc.dg/gomp/bad-array-shaping-c-1.c | 26 ++ .../gcc.dg/gomp/bad-array-shaping-c-2.c | 24 ++ .../gcc.dg/gomp/bad-array-shaping-c-3.c | 30 ++ .../gcc.dg/gomp/bad-array-shaping-c-4.c | 27 ++ .../gcc.dg/gomp/bad-array-shaping-c-5.c | 17 + .../gcc.dg/gomp/bad-array-shaping-c-6.c | 26 ++ .../gcc.dg/gomp/bad-array-shaping-c-7.c | 15 + libgomp/testsuite/libgomp.c/array-shaping-1.c | 236 ++++++++++++++ libgomp/testsuite/libgomp.c/array-shaping-2.c | 39 +++ libgomp/testsuite/libgomp.c/array-shaping-3.c | 42 +++ libgomp/testsuite/libgomp.c/array-shaping-4.c | 36 +++ libgomp/testsuite/libgomp.c/array-shaping-5.c | 38 +++ libgomp/testsuite/libgomp.c/array-shaping-6.c | 45 +++ 16 files changed, 1099 insertions(+), 48 deletions(-) create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-2.c create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-3.c create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-4.c create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-5.c create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-6.c create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-7.c create mode 100644 libgomp/testsuite/libgomp.c/array-shaping-1.c create mode 100644 libgomp/testsuite/libgomp.c/array-shaping-2.c create mode 100644 libgomp/testsuite/libgomp.c/array-shaping-3.c create mode 100644 libgomp/testsuite/libgomp.c/array-shaping-4.c create mode 100644 libgomp/testsuite/libgomp.c/array-shaping-5.c create mode 100644 libgomp/testsuite/libgomp.c/array-shaping-6.c diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 280426ddf10..7e895e11da2 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -5764,7 +5764,9 @@ c_parser_braced_init (c_parser *parser, tree type, bool nested_p, gcc_obstack_init (&braced_init_obstack); gcc_assert (c_parser_next_token_is (parser, CPP_OPEN_BRACE)); bool save_c_omp_array_section_p = c_omp_array_section_p; + bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p; c_omp_array_section_p = false; + c_omp_array_shaping_op_p = false; matching_braces braces; braces.consume_open (parser); if (nested_p) @@ -5804,6 +5806,7 @@ c_parser_braced_init (c_parser *parser, tree type, bool nested_p, } } c_omp_array_section_p = save_c_omp_array_section_p; + c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p; c_token *next_tok = c_parser_peek_token (parser); if (next_tok->type != CPP_CLOSE_BRACE) { @@ -8310,6 +8313,7 @@ c_parser_conditional_expression (c_parser *parser, struct c_expr *after, struct c_expr cond, exp1, exp2, ret; location_t start, cond_loc, colon_loc; bool save_c_omp_array_section_p = c_omp_array_section_p; + bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p; gcc_assert (!after || c_dialect_objc ()); @@ -8318,6 +8322,7 @@ c_parser_conditional_expression (c_parser *parser, struct c_expr *after, if (c_parser_next_token_is_not (parser, CPP_QUERY)) return cond; c_omp_array_section_p = false; + c_omp_array_shaping_op_p = false; if (cond.value != error_mark_node) start = cond.get_start (); else @@ -8371,6 +8376,7 @@ c_parser_conditional_expression (c_parser *parser, struct c_expr *after, ret.original_code = ERROR_MARK; ret.original_type = NULL; c_omp_array_section_p = save_c_omp_array_section_p; + c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p; return ret; } { @@ -8418,6 +8424,7 @@ c_parser_conditional_expression (c_parser *parser, struct c_expr *after, set_c_expr_source_range (&ret, start, exp2.get_finish ()); ret.m_decimal = 0; c_omp_array_section_p = save_c_omp_array_section_p; + c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p; return ret; } @@ -8799,6 +8806,8 @@ c_parser_cast_expression (c_parser *parser, struct c_expr *after) if (after) return c_parser_postfix_expression_after_primary (parser, cast_loc, *after); + bool save_c_omp_has_array_shape_p = c_omp_has_array_shape_p; + c_omp_has_array_shape_p = false; /* If the expression begins with a parenthesized type name, it may be either a cast or a compound literal; we need to see whether the next character is '{' to tell the difference. If not, it is @@ -8807,6 +8816,10 @@ c_parser_cast_expression (c_parser *parser, struct c_expr *after) if (c_parser_next_token_is (parser, CPP_OPEN_PAREN) && c_token_starts_compound_literal (c_parser_peek_2nd_token (parser))) { + bool save_c_omp_array_section_p = c_omp_array_section_p; + bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p; + c_omp_array_section_p = false; + c_omp_array_shaping_op_p = false; struct c_declspecs *scspecs; struct c_type_name *type_name; struct c_expr ret; @@ -8818,6 +8831,8 @@ c_parser_cast_expression (c_parser *parser, struct c_expr *after) parens.skip_until_found_close (parser); if (type_name == NULL) { + c_omp_array_section_p = save_c_omp_array_section_p; + c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p; ret.set_error (); ret.original_code = ERROR_MARK; ret.original_type = NULL; @@ -8828,9 +8843,15 @@ c_parser_cast_expression (c_parser *parser, struct c_expr *after) used_types_insert (type_name->specs->type); if (c_parser_next_token_is (parser, CPP_OPEN_BRACE)) - return c_parser_postfix_expression_after_paren_type (parser, scspecs, - type_name, - cast_loc); + { + c_expr r = c_parser_postfix_expression_after_paren_type (parser, + scspecs, + type_name, + cast_loc); + c_omp_array_section_p = save_c_omp_array_section_p; + c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p; + return r; + } if (scspecs) error_at (cast_loc, "storage class specifier in cast"); if (type_name->specs->alignas_p) @@ -8847,10 +8868,61 @@ c_parser_cast_expression (c_parser *parser, struct c_expr *after) ret.original_code = ERROR_MARK; ret.original_type = NULL; ret.m_decimal = 0; + c_omp_array_section_p = save_c_omp_array_section_p; + c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p; + return ret; + } + else if (c_omp_array_shaping_op_p + && c_parser_next_token_is (parser, CPP_OPEN_PAREN) + && c_parser_peek_2nd_token (parser)->type == CPP_OPEN_SQUARE) + { + bool save_c_omp_array_section_p = c_omp_array_section_p; + bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p; + c_omp_array_section_p = false; + c_omp_array_shaping_op_p = false; + auto_vec omp_shape_dims; + struct c_expr expr, ret; + matching_parens parens; + parens.consume_open (parser); + while (c_parser_next_token_is (parser, CPP_OPEN_SQUARE)) + { + c_parser_consume_token (parser); + c_expr e = c_parser_expression (parser); + if (e.value == error_mark_node) + break; + omp_shape_dims.safe_push (e.value); + if (!c_parser_require (parser, CPP_CLOSE_SQUARE, + "expected %<]%>")) + break; + } + parens.require_close (parser); + c_omp_array_section_p = save_c_omp_array_section_p; + c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p; + { + location_t expr_loc = c_parser_peek_token (parser)->location; + bool save_c_omp_has_array_shape_p = c_omp_has_array_shape_p; + c_omp_has_array_shape_p = true; + expr = c_parser_cast_expression (parser, NULL); + c_omp_has_array_shape_p = save_c_omp_has_array_shape_p; + /* NOTE: We don't want to introduce conversions here. */ + expr = convert_lvalue_to_rvalue (expr_loc, expr, false, true); + } + tree arrtype + = create_omp_arrayshape_type (expr.value, &omp_shape_dims); + ret.value = build1_loc (cast_loc, VIEW_CONVERT_EXPR, arrtype, + expr.value); + if (ret.value && expr.value) + set_c_expr_source_range (&ret, cast_loc, expr.get_finish ()); + ret.original_code = ERROR_MARK; + ret.original_type = NULL; + ret.m_decimal = 0; return ret; } else - return c_parser_unary_expression (parser); + { + c_omp_has_array_shape_p = save_c_omp_has_array_shape_p; + return c_parser_unary_expression (parser); + } } /* Parse an unary expression (C90 6.3.3, C99 6.5.3, C11 6.5.3). @@ -9860,6 +9932,7 @@ c_parser_postfix_expression (c_parser *parser) tree stmt; location_t brace_loc; bool save_c_omp_array_section_p = c_omp_array_section_p; + bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p; c_parser_consume_token (parser); brace_loc = c_parser_peek_token (parser)->location; c_parser_consume_token (parser); @@ -9877,6 +9950,7 @@ c_parser_postfix_expression (c_parser *parser) break; } c_omp_array_section_p = false; + c_omp_array_shaping_op_p = false; stmt = c_begin_stmt_expr (); c_parser_compound_statement_nostart (parser); location_t close_loc = c_parser_peek_token (parser)->location; @@ -9888,6 +9962,7 @@ c_parser_postfix_expression (c_parser *parser) set_c_expr_source_range (&expr, loc, close_loc); mark_exp_read (expr.value); c_omp_array_section_p = save_c_omp_array_section_p; + c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p; } else { @@ -11373,20 +11448,26 @@ c_parser_postfix_expression_after_primary (c_parser *parser, if (c_omp_array_section_p && c_parser_next_token_is (parser, CPP_COLON)) { + tree stride = NULL_TREE; + c_parser_consume_token (parser); if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE)) len = c_parser_expression (parser).value; + if (c_parser_next_token_is (parser, CPP_COLON)) + { + c_parser_consume_token (parser); + if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE)) + stride = c_parser_expression (parser).value; + } + c_parser_skip_until_found (parser, CPP_CLOSE_SQUARE, "expected %<]%>"); - /* NOTE: We are reusing using the type of the whole array as the - type of the array section here, which isn't necessarily - entirely correct. Might need revisiting. */ start = expr.get_start (); finish = parser->tokens_buf[0].location; expr.value = build_omp_array_section (op_loc, expr.value, idx, - len, NULL_TREE /* fixme */); + len, stride); set_c_expr_source_range (&expr, start, finish); expr.original_code = ERROR_MARK; expr.original_type = NULL; @@ -11397,7 +11478,20 @@ c_parser_postfix_expression_after_primary (c_parser *parser, "expected %<]%>"); start = expr.get_start (); finish = parser->tokens_buf[0].location; - expr.value = build_array_ref (op_loc, expr.value, idx); + if (c_omp_has_array_shape_p) + /* If we have an array-shaping operator, we may not be able to + represent a well-formed ARRAY_REF here, because we are + coercing the type of the innermost array base and the + original type may not be compatible. Use the + OMP_ARRAY_SECTION code instead. We also want to explicitly + avoid creating INDIRECT_REFs for pointer bases, because + that can lead to parsing ambiguities (see + c_parser_omp_variable_list). */ + expr.value + = build_omp_array_section (op_loc, expr.value, idx, + size_one_node, NULL_TREE); + else + expr.value = build_array_ref (op_loc, expr.value, idx); set_c_expr_source_range (&expr, start, finish); expr.original_code = ERROR_MARK; expr.original_type = NULL; @@ -11694,7 +11788,9 @@ c_parser_expr_list (c_parser *parser, bool convert_p, bool fold_p, struct c_expr expr; unsigned int idx = 0; bool save_c_omp_array_section_p = c_omp_array_section_p; + bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p; c_omp_array_section_p = false; + c_omp_array_shaping_op_p = false; ret = make_tree_vector (); if (p_orig_types == NULL) @@ -11749,6 +11845,7 @@ c_parser_expr_list (c_parser *parser, bool convert_p, bool fold_p, if (orig_types) *p_orig_types = orig_types; c_omp_array_section_p = save_c_omp_array_section_p; + c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p; return ret; } @@ -13939,6 +14036,35 @@ c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list) return list; } +/* Return, as an INTEGER_CST node, the number of elements for TYPE + (which is an ARRAY_TYPE). This counts only elements of the top + array. (From cp/tree.cc). */ + +static tree +c_array_type_nelts_top (tree type) +{ + return fold_build2_loc (input_location, PLUS_EXPR, sizetype, + array_type_nelts (type), size_one_node); +} + +/* Return, as an INTEGER_CST node, the number of elements for TYPE + (which is an ARRAY_TYPE). This one is a recursive count of all + ARRAY_TYPEs that are clumped together. (From cp/tree.cc). */ + +static tree +c_array_type_nelts_total (tree type) +{ + tree sz = c_array_type_nelts_top (type); + type = TREE_TYPE (type); + while (TREE_CODE (type) == ARRAY_TYPE) + { + tree n = c_array_type_nelts_top (type); + sz = fold_build2_loc (input_location, MULT_EXPR, sizetype, sz, n); + type = TREE_TYPE (type); + } + return sz; +} + /* OpenACC 2.0, OpenMP 2.5: variable-list: identifier @@ -14067,12 +14193,24 @@ c_parser_omp_variable_list (c_parser *parser, { location_t loc = c_parser_peek_token (parser)->location; bool save_c_omp_array_section_p = c_omp_array_section_p; + bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p; c_omp_array_section_p = true; + c_omp_array_shaping_op_p + = (kind == OMP_CLAUSE_TO || kind == OMP_CLAUSE_FROM); c_expr expr = c_parser_expr_no_commas (parser, NULL); if (expr.value != error_mark_node) mark_exp_read (expr.value); c_omp_array_section_p = save_c_omp_array_section_p; + c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p; tree decl = expr.value; + tree reshaped_to = NULL_TREE; + + if (TREE_CODE (decl) == VIEW_CONVERT_EXPR + && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) + { + reshaped_to = TREE_TYPE (decl); + decl = TREE_OPERAND (decl, 0); + } /* This code rewrites a parsed expression containing various tree codes used to represent array accesses into a more uniform nest of @@ -14085,6 +14223,31 @@ c_parser_omp_variable_list (c_parser *parser, dims.truncate (0); if (TREE_CODE (decl) == OMP_ARRAY_SECTION) { + size_t sections = 0; + tree orig_decl = decl; + bool update_p = (kind == OMP_CLAUSE_TO + || kind == OMP_CLAUSE_FROM); + bool maybe_ptr_based_noncontig_update = false; + + while (update_p + && !reshaped_to + && (TREE_CODE (decl) == OMP_ARRAY_SECTION + || TREE_CODE (decl) == ARRAY_REF + || TREE_CODE (decl) == COMPOUND_EXPR)) + { + if (TREE_CODE (decl) == COMPOUND_EXPR) + decl = TREE_OPERAND (decl, 1); + else + { + if (TREE_CODE (decl) == OMP_ARRAY_SECTION) + maybe_ptr_based_noncontig_update = true; + decl = TREE_OPERAND (decl, 0); + sections++; + } + } + + decl = orig_decl; + while (TREE_CODE (decl) == OMP_ARRAY_SECTION) { tree low_bound = TREE_OPERAND (decl, 1); @@ -14093,18 +14256,63 @@ c_parser_omp_variable_list (c_parser *parser, dims.safe_push (omp_dim (low_bound, length, stride, loc, false)); decl = TREE_OPERAND (decl, 0); + if (sections > 0) + sections--; } + /* The handling of INDIRECT_REF here in the presence of + array-shaping operations is a little tricky. We need to + avoid treating a pointer dereference as a unit-sized array + section when we have an array shaping operation, because we + don't want an indirection to consume one of the user's + requested array dimensions. E.g. if we have a + double-indirect pointer like: + + int **foopp; + #pragma omp target update from(([N][N]) (*foopp)[0:X][0:Y]) + + We don't want to interpret this as: + + foopp[0:1][0:X][0:Y] + + else the array shape [N][N] won't match. Also we can't match + the array sections right-to-left instead, else this: + + #pragma omp target update from(([N][N]) (*foopp)[0:X]) + + would not copy the dimensions: + + (*foopp)[0:X][0:N] + + as required. So, avoid descending through INDIRECT_REFs if + we have an array-shaping op. + + If we *don't* have an array-shaping op, but we have a + multiply-indirected pointer and an array section like this: + + int ***fooppp; + #pragma omp target update from((**fooppp)[0:X:S] + + also avoid descending through more indirections than we have + array sections, since the noncontiguous update processing code + won't understand them (and doesn't need to traverse them + anyway). */ + while (TREE_CODE (decl) == ARRAY_REF - || TREE_CODE (decl) == INDIRECT_REF + || (TREE_CODE (decl) == INDIRECT_REF + && !reshaped_to) || TREE_CODE (decl) == COMPOUND_EXPR) { + if (maybe_ptr_based_noncontig_update && sections == 0) + break; + if (TREE_CODE (decl) == COMPOUND_EXPR) { decl = TREE_OPERAND (decl, 1); STRIP_NOPS (decl); } - else if (TREE_CODE (decl) == INDIRECT_REF) + else if (TREE_CODE (decl) == INDIRECT_REF + && !reshaped_to) { dims.safe_push (omp_dim (integer_zero_node, integer_one_node, NULL_TREE, loc, @@ -14117,6 +14325,35 @@ c_parser_omp_variable_list (c_parser *parser, dims.safe_push (omp_dim (index, integer_one_node, NULL_TREE, loc, true)); decl = TREE_OPERAND (decl, 0); + if (sections > 0) + sections--; + } + } + + if (reshaped_to) + { + unsigned reshaped_dims = 0; + + for (tree t = reshaped_to; + TREE_CODE (t) == ARRAY_TYPE; + t = TREE_TYPE (t)) + reshaped_dims++; + + if (dims.length () > reshaped_dims) + { + error_at (loc, "too many array section specifiers " + "for %qT", reshaped_to); + decl = error_mark_node; + } + else + { + /* We have a pointer DECL whose target should be + interpreted as an array with particular dimensions, + not "the pointer itself". So, add an indirection + here. */ + decl = build_indirect_ref (loc, decl, RO_UNARY_STAR); + decl = build1_loc (loc, VIEW_CONVERT_EXPR, reshaped_to, + decl); } } @@ -14144,6 +14381,14 @@ c_parser_omp_variable_list (c_parser *parser, decl = build_omp_array_section (loc, decl, idx, integer_one_node, NULL_TREE); } + else if (reshaped_to) + { + /* We're copying the whole of a reshaped array, originally a + base pointer. Rewrite as an array section. */ + tree elems = c_array_type_nelts_total (reshaped_to); + decl = build_omp_array_section (loc, decl, size_zero_node, elems, + NULL_TREE); + } else if (TREE_CODE (decl) == NON_LVALUE_EXPR || CONVERT_EXPR_P (decl)) decl = TREE_OPERAND (decl, 0); @@ -22777,8 +23022,38 @@ c_parser_omp_target_update (location_t loc, c_parser *parser, tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK, "#pragma omp target update"); - if (omp_find_clause (clauses, OMP_CLAUSE_TO) == NULL_TREE - && omp_find_clause (clauses, OMP_CLAUSE_FROM) == NULL_TREE) + bool to_clause = false, from_clause = false; + for (tree c = clauses; + c && !to_clause && !from_clause; + c = OMP_CLAUSE_CHAIN (c)) + { + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_TO: + to_clause = true; + break; + case OMP_CLAUSE_FROM: + from_clause = true; + break; + case OMP_CLAUSE_MAP: + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_TO_GRID: + to_clause = true; + break; + case GOMP_MAP_FROM_GRID: + from_clause = true; + break; + default: + ; + } + break; + default: + ; + } + } + + if (!to_clause && !from_clause) { error_at (loc, "%<#pragma omp target update%> must contain at least one " diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h index 37790bab640..ee29f9de2cc 100644 --- a/gcc/c/c-tree.h +++ b/gcc/c/c-tree.h @@ -727,6 +727,8 @@ extern int in_sizeof; extern int in_typeof; extern bool c_in_omp_for; extern bool c_omp_array_section_p; +extern bool c_omp_array_shaping_op_p; +extern bool c_omp_has_array_shape_p; extern tree c_last_sizeof_arg; extern location_t c_last_sizeof_loc; @@ -766,6 +768,8 @@ extern tree build_component_ref (location_t, tree, tree, location_t, location_t); extern tree build_array_ref (location_t, tree, tree); extern tree build_omp_array_section (location_t, tree, tree, tree, tree); +extern tree create_omp_arrayshape_type (tree expr, + vec *omp_shape_dims); extern tree build_external_ref (location_t, tree, bool, tree *); extern void pop_maybe_used (bool); extern struct c_expr c_expr_sizeof_expr (location_t, struct c_expr); diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 0eff41e7567..f6fb68cb491 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -79,6 +79,13 @@ bool c_in_omp_for; /* True when parsing OpenMP map clause. */ bool c_omp_array_section_p; +/* True when parsing OpenMP to/from clause. */ +bool c_omp_array_shaping_op_p; + +/* True if we have an OpenMP array-shaping "cast" expression. This adjusts + the parsed representation for e.g. array refs. */ +bool c_omp_has_array_shape_p; + /* The argument of last parsed sizeof expression, only to be tested if expr.original_code == SIZEOF_EXPR. */ tree c_last_sizeof_arg; @@ -2930,8 +2937,8 @@ build_omp_array_section (location_t loc, tree array, tree index, tree length, if (index != NULL_TREE && length != NULL_TREE - && TREE_CODE (index) == INTEGER_CST - && TREE_CODE (length) == INTEGER_CST) + && INTEGRAL_TYPE_P (TREE_TYPE (index)) + && INTEGRAL_TYPE_P (TREE_TYPE (length))) { tree low = fold_convert (sizetype, index); tree high = fold_convert (sizetype, length); @@ -2941,7 +2948,7 @@ build_omp_array_section (location_t loc, tree array, tree index, tree length, } else if ((index == NULL_TREE || integer_zerop (index)) && length != NULL_TREE - && TREE_CODE (length) == INTEGER_CST) + && INTEGRAL_TYPE_P (TREE_TYPE (length))) idxtype = build_index_type (length); else idxtype = NULL_TREE; @@ -2965,6 +2972,46 @@ build_omp_array_section (location_t loc, tree array, tree index, tree length, stride); } +/* Build an array type whose dimensions are given by OMP_SHAPE_DIMS and whose + elements are of the type pointed to by the "base" node of EXPR with outer + OMP_ARRAY_SECTIONs and ARRAY_REFs stripped off, e.g. the type of "*myptr" + in "myptr[0:2:3][4][5:6]". */ + +tree +create_omp_arrayshape_type (tree expr, vec *omp_shape_dims) +{ + tree strip_sections = expr; + + while (TREE_CODE (strip_sections) == OMP_ARRAY_SECTION + || TREE_CODE (strip_sections) == ARRAY_REF) + strip_sections = TREE_OPERAND (strip_sections, 0); + + tree type = TREE_TYPE (strip_sections); + + if (TREE_CODE (type) == REFERENCE_TYPE) + type = TREE_TYPE (type); + + if (TREE_CODE (type) != POINTER_TYPE) + { + error ("OpenMP array shaping operator with non-pointer argument"); + return error_mark_node; + } + + type = TREE_TYPE (type); + + int i; + tree dim; + FOR_EACH_VEC_ELT_REVERSE (*omp_shape_dims, i, dim) + { + tree maxidx = fold_convert (sizetype, dim); + maxidx = size_binop (MINUS_EXPR, maxidx, size_one_node); + tree index = build_index_type (maxidx); + type = build_array_type (type, index); + } + + return type; +} + /* Build an external reference to identifier ID. FUN indicates whether this will be used for a function call. LOC is the source @@ -13714,7 +13761,8 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses) static tree handle_omp_array_sections_1 (tree c, tree t, vec &types, bool &maybe_zero_len, unsigned int &first_non_one, - bool &non_contiguous, enum c_omp_region_type ort) + bool &non_contiguous, enum c_omp_region_type ort, + int *discontiguous) { tree ret, low_bound, length, stride, type; bool openacc = (ort & C_ORT_ACC) != 0; @@ -13795,11 +13843,14 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, ret = handle_omp_array_sections_1 (c, TREE_OPERAND (t, 0), types, maybe_zero_len, first_non_one, - non_contiguous, ort); + non_contiguous, ort, discontiguous); if (ret == error_mark_node || ret == NULL_TREE) return ret; - type = TREE_TYPE (ret); + if (TREE_CODE (ret) == OMP_ARRAY_SECTION) + type = TREE_TYPE (TREE_TYPE (TREE_OPERAND (ret, 0))); + else + type = TREE_TYPE (ret); low_bound = TREE_OPERAND (t, 1); length = TREE_OPERAND (t, 2); stride = TREE_OPERAND (t, 3); @@ -13840,8 +13891,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, && TYPE_PRECISION (TREE_TYPE (length)) > TYPE_PRECISION (sizetype)) length = fold_convert (sizetype, length); + if (stride + && TREE_CODE (stride) == INTEGER_CST + && TYPE_PRECISION (TREE_TYPE (stride)) + > TYPE_PRECISION (sizetype)) + stride = fold_convert (sizetype, stride); if (low_bound == NULL_TREE) low_bound = integer_zero_node; + if (stride == NULL_TREE) + stride = size_one_node; if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)) @@ -13960,12 +14018,29 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, } if (length && TREE_CODE (length) == INTEGER_CST) { - if (tree_int_cst_lt (size, length)) + tree slength = length; + if (stride && TREE_CODE (stride) == INTEGER_CST) { - error_at (OMP_CLAUSE_LOCATION (c), - "length %qE above array section size " - "in %qs clause", length, - omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + slength = size_binop (MULT_EXPR, + fold_convert (sizetype, length), + fold_convert (sizetype, stride)); + slength = size_binop (MINUS_EXPR, + slength, + fold_convert (sizetype, stride)); + slength = size_binop (PLUS_EXPR, slength, size_one_node); + } + if (tree_int_cst_lt (size, slength)) + { + if (stride && !integer_onep (stride)) + error_at (OMP_CLAUSE_LOCATION (c), + "length %qE with stride %qE above array " + "section size in %qs clause", length, stride, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + else + error_at (OMP_CLAUSE_LOCATION (c), + "length %qE above array section size " + "in %qs clause", length, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); return error_mark_node; } if (TREE_CODE (low_bound) == INTEGER_CST) @@ -13973,7 +14048,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, tree lbpluslen = size_binop (PLUS_EXPR, fold_convert (sizetype, low_bound), - fold_convert (sizetype, length)); + fold_convert (sizetype, slength)); if (TREE_CODE (lbpluslen) == INTEGER_CST && tree_int_cst_lt (size, lbpluslen)) { @@ -14047,7 +14122,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, { tree d_length = TREE_OPERAND (d, 2); tree d_stride = TREE_OPERAND (d, 3); - if (d_length == NULL_TREE || !integer_onep (d_length) + if (d_length == NULL_TREE + || !integer_onep (d_length) || (d_stride && !integer_onep (d_stride))) { if (openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) @@ -14068,10 +14144,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } - error_at (OMP_CLAUSE_LOCATION (c), - "array section is not contiguous in %qs clause", - omp_clause_code_name[OMP_CLAUSE_CODE (c)]); - return error_mark_node; + if (discontiguous && *discontiguous) + *discontiguous = 2; + else + { + error_at (OMP_CLAUSE_LOCATION (c), + "array section is not contiguous in %qs clause", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + return error_mark_node; + } } } } @@ -14083,7 +14164,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND) - types.safe_push (TREE_TYPE (ret)); + types.safe_push (type); /* We will need to evaluate lb more than once. */ tree lb = save_expr (low_bound); if (lb != low_bound) @@ -14091,14 +14172,42 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, TREE_OPERAND (t, 1) = lb; low_bound = lb; } - ret = build_array_ref (OMP_CLAUSE_LOCATION (c), ret, low_bound); + /* NOTE: Stride/length are discarded for affinity/depend here. */ + if (discontiguous + && *discontiguous + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND) + ret = build_omp_array_section (OMP_CLAUSE_LOCATION (c), ret, low_bound, + length, stride); + else + ret = build_array_ref (OMP_CLAUSE_LOCATION (c), ret, low_bound); return ret; } -/* Handle array sections for clause C. */ +/* We built a reference to an array section, but it turns out we only need a + set of ARRAY_REFs to the lower bound. Rewrite the node. */ + +static tree +omp_array_section_low_bound (location_t loc, tree node) +{ + if (TREE_CODE (node) == OMP_ARRAY_SECTION) + { + tree low_bound = TREE_OPERAND (node, 1); + tree ret = omp_array_section_low_bound (loc, TREE_OPERAND (node, 0)); + return build_array_ref (loc, ret, low_bound); + } + + return node; +} + +/* Handle array sections for clause C. On entry *DISCONTIGUOUS is 0 if array + section must be contiguous, 1 if it can be discontiguous, and in the latter + case it is set to 2 on exit if it is determined to be discontiguous during + the function's execution. */ static bool -handle_omp_array_sections (tree *pc, enum c_omp_region_type ort) +handle_omp_array_sections (tree *pc, enum c_omp_region_type ort, + int *discontiguous) { tree c = *pc; bool maybe_zero_len = false; @@ -14114,7 +14223,7 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort) tp = &TREE_VALUE (*tp); tree first = handle_omp_array_sections_1 (c, *tp, types, maybe_zero_len, first_non_one, - non_contiguous, ort); + non_contiguous, ort, discontiguous); if (first == error_mark_node) return true; if (first == NULL_TREE) @@ -14153,11 +14262,14 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort) if (int_size_in_bytes (TREE_TYPE (first)) <= 0) maybe_zero_len = true; + bool higher_discontiguous = false; + for (i = num, t = OMP_CLAUSE_DECL (c); i > 0; t = TREE_OPERAND (t, 0)) { tree low_bound = TREE_OPERAND (t, 1); tree length = TREE_OPERAND (t, 2); + tree stride = TREE_OPERAND (t, 3); i--; if (low_bound @@ -14170,6 +14282,11 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort) && TYPE_PRECISION (TREE_TYPE (length)) > TYPE_PRECISION (sizetype)) length = fold_convert (sizetype, length); + if (stride + && TREE_CODE (stride) == INTEGER_CST + && TYPE_PRECISION (TREE_TYPE (stride)) + > TYPE_PRECISION (sizetype)) + stride = fold_convert (sizetype, stride); if (low_bound == NULL_TREE) low_bound = integer_zero_node; @@ -14179,10 +14296,49 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort) continue; } + if (stride == NULL_TREE) + stride = size_one_node; + if (discontiguous && *discontiguous) + { + /* This condition is similar to the error check below, but + whereas that checks for a definitely-discontiguous array + section in order to report an error (where such a section is + illegal), here we instead need to know if the array section + *may be* discontiguous so we can handle that case + appropriately (i.e. for rectangular "target update" + operations). */ + bool full_span = false; + if (length != NULL_TREE + && TREE_CODE (length) == INTEGER_CST + && TREE_CODE (types[i]) == ARRAY_TYPE + && TYPE_DOMAIN (types[i]) + && TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])) + && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (types[i]))) + == INTEGER_CST) + { + tree size; + size = size_binop (PLUS_EXPR, + TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])), + size_one_node); + if (tree_int_cst_equal (length, size)) + full_span = true; + } + + if (!integer_onep (stride) + || (higher_discontiguous + && (!integer_zerop (low_bound) + || !full_span))) + *discontiguous = 2; + + if (!integer_onep (stride) + || !integer_zerop (low_bound) + || !full_span) + higher_discontiguous = true; + } if (!maybe_zero_len && i > first_non_one) { if (integer_nonzerop (low_bound)) - goto do_warn_noncontiguous; + goto is_noncontiguous; if (length != NULL_TREE && TREE_CODE (length) == INTEGER_CST && TYPE_DOMAIN (types[i]) @@ -14196,12 +14352,17 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort) size_one_node); if (!tree_int_cst_equal (length, size)) { - do_warn_noncontiguous: - error_at (OMP_CLAUSE_LOCATION (c), - "array section is not contiguous in %qs " - "clause", - omp_clause_code_name[OMP_CLAUSE_CODE (c)]); - return true; + is_noncontiguous: + if (discontiguous && *discontiguous) + *discontiguous = 2; + else + { + error_at (OMP_CLAUSE_LOCATION (c), + "array section is not contiguous in %qs " + "clause", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + return true; + } } } if (length != NULL_TREE @@ -14321,6 +14482,8 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort) OMP_CLAUSE_DECL (c) = t; return false; } + if (discontiguous && *discontiguous != 2) + first = omp_array_section_low_bound (OMP_CLAUSE_LOCATION (c), first); first = c_fully_fold (first, false, NULL); OMP_CLAUSE_DECL (c) = first; if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) @@ -14329,7 +14492,8 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort) size = c_fully_fold (size, false, NULL); OMP_CLAUSE_SIZE (c) = size; - if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + && !(discontiguous && *discontiguous == 2)) return false; auto_vec addr_tokens; @@ -14342,7 +14506,8 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort) tree *npc = ai.expand_map_clause (pc, first, addr_tokens, ort); if (npc != NULL) { - if (ai.maybe_zero_length_array_section (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && ai.maybe_zero_length_array_section (c)) OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; return false; @@ -14689,7 +14854,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == OMP_ARRAY_SECTION) { - if (handle_omp_array_sections (pc, ort)) + if (handle_omp_array_sections (pc, ort, NULL)) { remove = true; break; @@ -15428,7 +15593,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) last_iterators = NULL_TREE; if (TREE_CODE (t) == OMP_ARRAY_SECTION) { - if (handle_omp_array_sections (pc, ort)) + if (handle_omp_array_sections (pc, ort, NULL)) remove = true; else if ((c = *pc) && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND @@ -15535,6 +15700,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; break; } + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE) + break; /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: @@ -15549,7 +15717,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) grp_start_p = pc; grp_sentinel = OMP_CLAUSE_CHAIN (c); - if (handle_omp_array_sections (pc, ort)) + int discontiguous + = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM); + if (handle_omp_array_sections (pc, ort, &discontiguous)) remove = true; else { @@ -15946,7 +16117,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == OMP_ARRAY_SECTION) { - if (handle_omp_array_sections (pc, ort)) + if (handle_omp_array_sections (pc, ort, NULL)) remove = true; else { diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c new file mode 100644 index 00000000000..42d584fa624 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c @@ -0,0 +1,26 @@ +// { dg-do compile } + +#include +#include +#include + +int main (void) +{ + float *arr = calloc (100, sizeof (float)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = i + j * 3; + +#pragma omp target update to(([10][10]) arr[3:2][1:8][0:5]) +// { dg-error "too many array section specifiers for" "" { target *-*-* } .-1 } +// { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } + +#pragma omp target exit data map(from: arr[:100]) + + free (arr); + + return 0; +} diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-2.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-2.c new file mode 100644 index 00000000000..6be3e009ecb --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-2.c @@ -0,0 +1,24 @@ +// { dg-do compile } + +#include +#include +#include + +int main (void) +{ + float *arr = calloc (100, sizeof (float)); + + /* This isn't allowed. */ +#pragma omp target enter data map(to: ([10][10]) arr[:100]) +/* { dg-error {expected expression before '\[' token} "" { target *-*-* } .-1 } */ +/* { dg-error {'#pragma omp target enter data' must contain at least one 'map' clause} "" { target *-*-* } .-2 } */ + + /* Nor this. */ +#pragma omp target exit data map(from: ([10][10]) arr[:100]) +/* { dg-error {expected expression before '\[' token} "" { target *-*-* } .-1 } */ +/* { dg-error {'#pragma omp target exit data' must contain at least one 'map' clause} "" { target *-*-* } .-2 } */ + + free (arr); + + return 0; +} diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-3.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-3.c new file mode 100644 index 00000000000..1715b8ff9ed --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-3.c @@ -0,0 +1,30 @@ +// { dg-do compile } + +#include +#include +#include + +extern float* baz(void*); + +int main (void) +{ + float *arr = calloc (100, sizeof (float)); + int c = 50; + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = i + j * 3; + + /* No array shaping inside a function call. */ +#pragma omp target update to(baz(([10][10]) arr)) +/* { dg-error {expected expression before '\[' token} "" { target *-*-* } .-1 } */ +/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */ + +#pragma omp target exit data map(from: arr[:100]) + + free (arr); + + return 0; +} diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-4.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-4.c new file mode 100644 index 00000000000..cebefd36d18 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-4.c @@ -0,0 +1,27 @@ +// { dg-do compile } + +#include +#include +#include + +int main (void) +{ + float *arr = calloc (100, sizeof (float)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = i + j * 3; + + /* No array shaping inside a statement expression. */ +#pragma omp target update to( ({ int d = 10; ([d][d]) arr; }) ) +/* { dg-error {expected expression before '\[' token} "" { target *-*-* } .-1 } */ +/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */ + +#pragma omp target exit data map(from: arr[:100]) + + free (arr); + + return 0; +} diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-5.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-5.c new file mode 100644 index 00000000000..e1c4991f5c3 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-5.c @@ -0,0 +1,17 @@ +// { dg-do compile } + +struct S { + void *pp; +}; + +int main() +{ + int *sub1; + + /* No array section inside compound literal. */ +#pragma omp target update to( (struct S) { .pp = ([10][10]) sub1 } ) +/* { dg-error {expected expression before '\[' token} "" { target *-*-* } .-1 } */ +/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */ + + return 0; +} diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-6.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-6.c new file mode 100644 index 00000000000..d282d8598b2 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-6.c @@ -0,0 +1,26 @@ +// { dg-do compile } + +int main (void) +{ + char *ptr; + +#pragma omp target update to(([5][6][7]) ptr[0:4][0:7][0:7]) +/* { dg-error {length '7' above array section size in 'to' clause} "" { target *-*-* } .-1 } */ +/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */ + +#pragma omp target update to(([5][6][7]) ptr[1:5][0:6][0:7]) +/* { dg-error {high bound '6' above array section size in 'to' clause} "" { target *-*-* } .-1 } */ +/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */ + +#pragma omp target update from(([100]) ptr[3:33:3]) + +#pragma omp target update from(([100]) ptr[4:33:3]) +/* { dg-error {high bound '101' above array section size in 'from' clause} "" { target *-*-* } .-1 } */ +/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */ + +#pragma omp target update to(([10][10]) ptr[0:9:-1][0:9]) +/* { dg-error {length '9' with stride '-1' above array section size in 'to' clause} "" { target *-*-* } .-1 } */ +/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */ + + return 0; +} diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-7.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-7.c new file mode 100644 index 00000000000..233d8da6f44 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-7.c @@ -0,0 +1,15 @@ +/* { dg-do compile } */ + +int cond; + +int main (void) +{ + int *arr; + + /* No array shaping inside conditional operator. */ +#pragma omp target update to(cond ? ([3][9]) arr : ([2][7]) arr) +/* { dg-error {expected expression before '\[' token} "" { target *-*-* } .-1 } */ +/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */ + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/array-shaping-1.c b/libgomp/testsuite/libgomp.c/array-shaping-1.c new file mode 100644 index 00000000000..808c5f9ceae --- /dev/null +++ b/libgomp/testsuite/libgomp.c/array-shaping-1.c @@ -0,0 +1,236 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include +#include + +volatile int yy = 4, zz = 2, str_str = 2; + +int main() +{ + int *arr; + int x = 5; + int arr2d[10][10]; + + arr = calloc (100, sizeof (int)); + + /* Update whole reshaped array. */ + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < x; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = i ^ j; + +#pragma omp target update to(([10][x]) arr) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (j < x) + assert (arr[j * 10 + i] == i ^ j); + else + assert (arr[j * 10 + i] == 0); + + + /* Strided update. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 20; j++) + for (int i = 0; i < 5; i++) + arr[j * 5 + i] = i + j; + +#pragma omp target update to(([5][5]) arr[0:3][0:3:2]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 20; j++) + for (int i = 0; i < 5; i++) + if (j < 3 && (i & 1) == 0 && i < 6) + assert (arr[j * 5 + i] == i + j); + else + assert (arr[j * 5 + i] == 0); + + + /* Reshaped update, contiguous. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 20; j++) + for (int i = 0; i < 5; i++) + arr[j * 5 + i] = 2 * j + i; + +#pragma omp target update to(([5][5]) arr[0:5][0:5]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 20; j++) + for (int i = 0; i < 5; i++) + if (j < 5 && i < 5) + assert (arr[j * 5 + i] == 2 * j + i); + else + assert (arr[j * 5 + i] == 0); + + + /* Strided update on actual array. */ + + memset (arr2d, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr2d) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr2d[j][i] = j + 2 * i; + +#pragma omp target update to(arr2d[0:5:2][5:2]) + +#pragma omp target exit data map(from: arr2d) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if ((j & 1) == 0 && i >= 5 && i < 7) + assert (arr2d[j][i] == j + 2 * i); + else + assert (arr2d[j][i] == 0); + + + /* Update with non-constant bounds. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = (2 * j) ^ i; + + x = 3; + int y = yy, z = zz, str = str_str; + /* This is actually [0:3:2] [4:2:2]. */ +#pragma omp target update to(([10][10]) arr[0:x:2][y:z:str]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if ((j & 1) == 0 && j < 6 && (i & 1) == 0 && i >= 4 && i < 8) + assert (arr[j * 10 + i] == (2 * j) ^ i); + else + assert (arr[j * 10 + i] == 0); + + + /* Update with full "major" dimension. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = i + j; + +#pragma omp target update to(([10][10]) arr[0:10][3:1]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (i == 3) + assert (arr[j * 10 + i] == i + j); + else + assert (arr[j * 10 + i] == 0); + + + /* Update with full "minor" dimension. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = 3 * (i + j); + +#pragma omp target update to(([10][10]) arr[3:2][0:10]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (j >= 3 && j < 5) + assert (arr[j * 10 + i] == 3 * (i + j)); + else + assert (arr[j * 10 + i] == 0); + + + /* Rectangle update. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = 5 * (i + j); + +#pragma omp target update to(([10][10]) arr[3:2][0:9]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (j >= 3 && j < 5 && i < 9) + assert (arr[j * 10 + i] == 5 * (i + j)); + else + assert (arr[j * 10 + i] == 0); + + + /* One-dimensional strided update. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int i = 0; i < 100; i++) + arr[i] = i + 99; + +#pragma omp target update to(([100]) arr[3:33:3]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int i = 0; i < 100; i++) + if (i >= 3 && ((i - 3) % 3) == 0) + assert (arr[i] == i + 99); + else + assert (arr[i] == 0); + + + /* One-dimensional strided update without explicit array shape. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int i = 0; i < 100; i++) + arr[i] = i + 121; + +#pragma omp target update to(arr[3:33:3]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int i = 0; i < 100; i++) + if (i >= 3 && ((i - 3) % 3) == 0) + assert (arr[i] == i + 121); + else + assert (arr[i] == 0); + + free (arr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/array-shaping-2.c b/libgomp/testsuite/libgomp.c/array-shaping-2.c new file mode 100644 index 00000000000..42a6e0ca7d8 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/array-shaping-2.c @@ -0,0 +1,39 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +typedef struct { + int *aptr; +} C; + +int main() +{ + C cvar; + + cvar.aptr = calloc (100, sizeof (float)); + +#pragma omp target enter data map(to: cvar.aptr, cvar.aptr[:100]) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + cvar.aptr[i * 10 + j] = i + j; + } + +#pragma omp target update from(([10][10]) cvar.aptr[4:3][4:3]) + + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + if (i >= 4 && i < 7 && j >= 4 && j < 7) + assert (cvar.aptr[i * 10 + j] == i + j); + else + assert (cvar.aptr[i * 10 + j] == 0); + +#pragma omp target exit data map(delete: cvar.aptr, cvar.aptr[:100]) + + free (cvar.aptr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/array-shaping-3.c b/libgomp/testsuite/libgomp.c/array-shaping-3.c new file mode 100644 index 00000000000..5dda2e32832 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/array-shaping-3.c @@ -0,0 +1,42 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include +#include + +#define N 10 + +typedef struct { + int arr[N][N]; +} B; + +int main() +{ + B *bvar = malloc (sizeof (B)); + + memset (bvar, 0, sizeof (B)); + +#pragma omp target enter data map(to: bvar->arr) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + bvar->arr[i][j] = i + j; + } + +#pragma omp target update from(bvar->arr[4:3][4:3]) + + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + if (i >= 4 && i < 7 && j >= 4 && j < 7) + assert (bvar->arr[i][j] == i + j); + else + assert (bvar->arr[i][j] == 0); + +#pragma omp target exit data map(delete: bvar->arr) + + free (bvar); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/array-shaping-4.c b/libgomp/testsuite/libgomp.c/array-shaping-4.c new file mode 100644 index 00000000000..2b9e6949b60 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/array-shaping-4.c @@ -0,0 +1,36 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +#define N 10 + +int main () +{ + int iarr[N * N]; + + memset (iarr, 0, N * N * sizeof (int)); + +#pragma omp target enter data map(to: iarr) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + iarr[i * 10 + j] = i + j; + } + + /* An array, but cast to a pointer, then reshaped. */ +#pragma omp target update from(([10][10]) ((int *) &iarr[0])[4:3][4:3]) + + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + if (i >= 4 && i < 7 && j >= 4 && j < 7) + assert (iarr[i * 10 + j] == i + j); + else + assert (iarr[i * 10 + j] == 0); + +#pragma omp target exit data map(delete: iarr) + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/array-shaping-5.c b/libgomp/testsuite/libgomp.c/array-shaping-5.c new file mode 100644 index 00000000000..1034682e4ca --- /dev/null +++ b/libgomp/testsuite/libgomp.c/array-shaping-5.c @@ -0,0 +1,38 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +#define N 10 + +int main () +{ + int iarr_real[N * N]; + int *iarrp = &iarr_real[0]; + int **iarrpp = &iarrp; + + memset (iarrp, 0, N * N * sizeof (int)); + +#pragma omp target enter data map(to: iarr_real) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + iarrp[i * 10 + j] = i + j; + } + + /* A pointer with an extra indirection. */ +#pragma omp target update from(([10][10]) (*iarrpp)[4:3][4:3]) + + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + if (i >= 4 && i < 7 && j >= 4 && j < 7) + assert (iarrp[i * 10 + j] == i + j); + else + assert (iarrp[i * 10 + j] == 0); + +#pragma omp target exit data map(delete: iarr_real) + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/array-shaping-6.c b/libgomp/testsuite/libgomp.c/array-shaping-6.c new file mode 100644 index 00000000000..59388232244 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/array-shaping-6.c @@ -0,0 +1,45 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include +#include + +#define N 10 + +int main () +{ + int *iptr = calloc (N * N * N, sizeof (int)); + +#pragma omp target enter data map(to: iptr[0:N*N*N]) + +#pragma omp target + { + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + iptr[i * N * N + 4 * N + j] = i + j; + } + + /* An array ref between two array sections. */ +#pragma omp target update from(([N][N][N]) iptr[2:3][4][6:3]) + + for (int i = 2; i < 5; i++) + for (int j = 6; j < 9; j++) + assert (iptr[i * N * N + 4 * N + j] == i + j); + + memset (iptr, 0, N * N * N * sizeof (int)); + + for (int i = 0; i < N; i++) + iptr[2 * N * N + i * N + 4] = 3 * i; + + /* Array section between two array refs. */ +#pragma omp target update to(([N][N][N]) iptr[2][3:6][4]) + +#pragma omp target exit data map(from: iptr[0:N*N*N]) + + for (int i = 3; i < 9; i++) + assert (iptr[2 * N * N + i * N + 4] == 3 * i); + + free (iptr); + + return 0; +}