From patchwork Wed Sep 6 09:34:33 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1830287 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=8.43.85.97; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (ip-8-43-85-97.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4RgcjT03jCz1yh1 for ; Wed, 6 Sep 2023 19:37:33 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 174433882050 for ; Wed, 6 Sep 2023 09:37:31 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 7F3423856243; Wed, 6 Sep 2023 09:35:12 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 7F3423856243 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.02,231,1688457600"; d="scan'208";a="16228991" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 06 Sep 2023 01:35:11 -0800 IronPort-SDR: qxv+7OfINASedqVpLWYSpRK4lZ4hJrhdr87rJBg3KYd0NKuj+AZu/9xSR3EnIuzr26nWqIAyzP uljncmcXu8agieyMnCDJjoHiWlKmh+pjZXjy8Z04BOxv7+Y/AxIwBL/2QoQgiJEXcgx9vXUPbD +EeK/ONQY4vin5lrpR9JVqKPRFADerB3GAESWmcmv97IZeXpTXT8SzDmq/R8U8CHpmLhlRfLtl UVAwCr20KKdNBqm5MggR40JPB+PddZYCkI9bdZrsrYr9jxv8deg56YZiMn3BEVV+9uOFh2yFXl gco= From: Julian Brown To: CC: , , Subject: [PATCH 4/5] OpenMP: Array shaping operator and strided "target update" for C Date: Wed, 6 Sep 2023 02:34:33 -0700 Message-ID: X-Mailer: git-send-email 2.41.0 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-15.mgc.mentorg.com (139.181.222.15) 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 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.30 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++, 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 by the previous patch. Some details of parsing necessarily differ for C, but the general ideas are the same. This version of the patch has been rebased and contains a couple of minor fixes relative to versions posted previously. 2023-09-05 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 | 300 +++++++++++++++++- gcc/c/c-tree.h | 4 + gcc/c/c-typeck.cc | 235 ++++++++++++-- .../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, 1097 insertions(+), 43 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 be80c2723ef0..1c28f763c8a0 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -5768,7 +5768,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) @@ -5808,6 +5810,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) { @@ -8198,6 +8201,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 ()); @@ -8206,6 +8210,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 @@ -8259,6 +8264,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; } { @@ -8306,6 +8312,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; } @@ -8687,6 +8694,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 @@ -8695,6 +8704,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; @@ -8706,6 +8719,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; @@ -8716,9 +8731,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) @@ -8735,10 +8756,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). @@ -9758,6 +9830,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); @@ -9775,6 +9848,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; @@ -9786,6 +9860,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 { @@ -11271,17 +11346,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 %<]%>"); 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; @@ -11292,7 +11376,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; @@ -11581,7 +11678,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) @@ -11636,6 +11735,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; } @@ -13835,6 +13935,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 @@ -13962,12 +14091,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 @@ -13980,6 +14121,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); @@ -13988,18 +14154,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, @@ -14012,6 +14223,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); } } @@ -14039,6 +14279,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); @@ -17803,7 +18051,7 @@ c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind, c_parser_consume_token (parser); } - tree nl = c_parser_omp_variable_list (parser, loc, kind, list); + tree nl = c_parser_omp_variable_list (parser, loc, kind, list, true); parens.skip_until_found_close (parser); if (present) @@ -22240,8 +22488,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 10f1cf26dd28..25830711e55a 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 c7409045abf3..79ca45401606 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; @@ -2939,6 +2946,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 @@ -13696,7 +13743,7 @@ 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, - enum c_omp_region_type ort) + enum c_omp_region_type ort, int *discontiguous) { tree ret, low_bound, length, stride, type; bool openacc = (ort & C_ORT_ACC) != 0; @@ -13776,11 +13823,15 @@ 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, ort); + maybe_zero_len, first_non_one, 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); @@ -13821,8 +13872,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)) @@ -13941,12 +13999,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) @@ -13954,7 +14029,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)) { @@ -14026,13 +14101,19 @@ 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))) { - 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; + } } } } @@ -14044,7 +14125,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) @@ -14052,14 +14133,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; @@ -14074,7 +14183,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, - ort); + ort, discontiguous); if (first == error_mark_node) return true; if (first == NULL_TREE) @@ -14112,11 +14221,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 @@ -14129,12 +14241,56 @@ 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; + 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]) @@ -14148,12 +14304,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 @@ -14265,6 +14426,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) @@ -14273,7 +14436,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; @@ -14286,7 +14450,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; @@ -14633,7 +14798,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; @@ -15255,7 +15420,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 @@ -15362,6 +15527,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: @@ -15376,7 +15544,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 { @@ -15773,7 +15944,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 000000000000..42d584fa6240 --- /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 000000000000..6be3e009ecb1 --- /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 000000000000..1715b8ff9edd --- /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 000000000000..cebefd36d189 --- /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 000000000000..e1c4991f5c34 --- /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 000000000000..d282d8598b22 --- /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 000000000000..233d8da6f445 --- /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 000000000000..808c5f9ceae6 --- /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 000000000000..42a6e0ca7d82 --- /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 000000000000..5dda2e328328 --- /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 000000000000..2b9e6949b602 --- /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 000000000000..1034682e4ca2 --- /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 000000000000..593882322443 --- /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; +}