From: Julian Brown <[email protected]> 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. 2023-09-05 Julian Brown <[email protected]> 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_total): New function. (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-c++-common/array-shaping-14.c: New test. * 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. 2026-08-13 Paul-Antoine Arras <[email protected]> gcc/c/ * c-tree.h (c_omp_array_section_p): Replace bool with... (c_omp_array_section_kind): ...this 3-state unsigned char. (OMP_ARRAY_SECTION_NONE, OMP_ARRAY_SECTION_UNSTRIDED, OMP_ARRAY_SECTION_STRIDED): Define. * c-parser.cc (c_parser_postfix_expression_after_primary): Only accept a 2rd colon (stride) for to/from clauses, not map. (c_parser_omp_variable_list): Only count OMP_ARRAY_SECTION layers, not every ARRAY_REF layer, in a to/from clause's base decl. * c-typeck.cc (handle_omp_array_sections): Handle a genuinely discontiguous update at any access-chain depth. (c_finish_omp_clauses): Splice the GOMP_MAP_TO_GRID/FROM_GRID replacement clause in at *PC when DISCONTIGUOUS == 2. gcc/testsuite/ * c-c++-common/gomp/target-update-iterators-4.c: New test. 2026-09-25 Paul-Antoine Arras <[email protected]> gcc/testsuite/ * c-c++-common/gomp/array-section-1.c: No longer expect a C++-specific error for 'arr[::1]', now that the C++ front end also accepts it without a disambiguating space. * c-c++-common/gomp/array-section-3.c: New test. * c-c++-common/gomp/array-shaping-1.c: New test. * c-c++-common/gomp/bad-array-shaping-2.c: New test, merged from gcc.dg/gomp/bad-array-shaping-c-1.c. * c-c++-common/gomp/bad-array-shaping-3.c: New test, merged from gcc.dg/gomp/bad-array-shaping-c-6.c. * gcc.dg/gomp/array-shaping-2.c: New test. * gcc.dg/gomp/bad-array-shaping-c-1.c: Remove; merged into c-c++-common/gomp/bad-array-shaping-2.c. * gcc.dg/gomp/bad-array-shaping-c-6.c: Remove; merged into c-c++-common/gomp/bad-array-shaping-3.c. libgomp/ * testsuite/libgomp.c-c++-common/array-shaping-15.c: New test. * testsuite/libgomp.c-c++-common/array-shaping-17.c: New test, merged from testsuite/libgomp.c++/array-shaping-2.C. * testsuite/libgomp.c-c++-common/array-shaping-18.c: New test, merged from testsuite/libgomp.c/array-shaping-1.c. * testsuite/libgomp.c-c++-common/array-shaping-19.c: New test, merged from testsuite/libgomp.c/array-shaping-4.c. * testsuite/libgomp.c-c++-common/array-shaping-20.c: New test, merged from testsuite/libgomp.c/array-shaping-5.c. * testsuite/libgomp.c-c++-common/array-shaping-21.c: New test, merged from testsuite/libgomp.c/array-shaping-6.c. * testsuite/libgomp.c/array-shaping-1.c: Remove; merged into testsuite/libgomp.c-c++-common/array-shaping-18.c. * testsuite/libgomp.c/array-shaping-4.c: Remove; merged into testsuite/libgomp.c-c++-common/array-shaping-19.c. * testsuite/libgomp.c/array-shaping-5.c: Remove; merged into testsuite/libgomp.c-c++-common/array-shaping-20.c. * testsuite/libgomp.c/array-shaping-6.c: Remove; merged into testsuite/libgomp.c-c++-common/array-shaping-21.c. --- gcc/c/c-parser.cc | 408 +++++++++++++++--- gcc/c/c-tree.h | 11 +- gcc/c/c-typeck.cc | 297 +++++++++++-- .../c-c++-common/gomp/array-section-1.c | 18 + .../c-c++-common/gomp/array-section-2.c | 32 ++ .../c-c++-common/gomp/array-section-3.c | 18 + .../c-c++-common/gomp/array-shaping-1.c | 14 + .../c-c++-common/gomp/bad-array-shaping-1.c | 12 + .../c-c++-common/gomp/bad-array-shaping-2.c | 25 ++ .../c-c++-common/gomp/bad-array-shaping-3.c | 27 ++ gcc/testsuite/gcc.dg/gomp/array-shaping-1.c | 17 + gcc/testsuite/gcc.dg/gomp/array-shaping-2.c | 16 + .../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-7.c | 15 + .../libgomp.c-c++-common/array-shaping-14.c | 34 ++ .../libgomp.c-c++-common/array-shaping-15.c | 41 ++ .../libgomp.c-c++-common/array-shaping-17.c | 40 ++ .../libgomp.c-c++-common/array-shaping-18.c | 236 ++++++++++ .../libgomp.c-c++-common/array-shaping-19.c | 36 ++ .../libgomp.c-c++-common/array-shaping-20.c | 38 ++ .../libgomp.c-c++-common/array-shaping-21.c | 45 ++ libgomp/testsuite/libgomp.c/array-shaping-2.c | 39 ++ libgomp/testsuite/libgomp.c/array-shaping-3.c | 42 ++ 26 files changed, 1464 insertions(+), 95 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-3.c create mode 100644 gcc/testsuite/c-c++-common/gomp/array-shaping-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/bad-array-shaping-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/bad-array-shaping-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/bad-array-shaping-3.c create mode 100644 gcc/testsuite/gcc.dg/gomp/array-shaping-1.c create mode 100644 gcc/testsuite/gcc.dg/gomp/array-shaping-2.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-7.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-shaping-14.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-shaping-15.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-shaping-17.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-shaping-18.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-shaping-19.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-shaping-20.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-shaping-21.c create mode 100644 libgomp/testsuite/libgomp.c/array-shaping-2.c create mode 100644 libgomp/testsuite/libgomp.c/array-shaping-3.c diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 0df45e5fc9c..dc32c8d9a2f 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -6636,8 +6636,11 @@ c_parser_braced_init (c_parser *parser, tree type, bool nested_p, location_t brace_loc = c_parser_peek_token (parser)->location; 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; - c_omp_array_section_p = false; + enum c_omp_array_section_kind save_c_omp_array_section_kind + = c_omp_array_section_kind; + bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p; + c_omp_array_section_kind = OMP_ARRAY_SECTION_NONE; + c_omp_array_shaping_op_p = false; bool zero_init_padding_bits = false; matching_braces braces; braces.consume_open (parser); @@ -6698,7 +6701,8 @@ c_parser_braced_init (c_parser *parser, tree type, bool nested_p, break; } } - c_omp_array_section_p = save_c_omp_array_section_p; + c_omp_array_section_kind = save_c_omp_array_section_kind; + 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) { @@ -10103,7 +10107,9 @@ 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; + enum c_omp_array_section_kind save_c_omp_array_section_kind + = c_omp_array_section_kind; + bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p; gcc_assert (!after || c_dialect_objc ()); @@ -10111,7 +10117,8 @@ 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_section_kind = OMP_ARRAY_SECTION_NONE; + c_omp_array_shaping_op_p = false; if (cond.value != error_mark_node) start = cond.get_start (); else @@ -10164,7 +10171,8 @@ c_parser_conditional_expression (c_parser *parser, struct c_expr *after, ret.set_error (); ret.original_code = ERROR_MARK; ret.original_type = NULL; - c_omp_array_section_p = save_c_omp_array_section_p; + c_omp_array_section_kind = save_c_omp_array_section_kind; + c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p; return ret; } { @@ -10211,7 +10219,8 @@ 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_section_kind = save_c_omp_array_section_kind; + c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p; return ret; } @@ -10602,6 +10611,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 @@ -10610,6 +10621,11 @@ 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))) { + enum c_omp_array_section_kind save_c_omp_array_section_kind + = c_omp_array_section_kind; + bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p; + c_omp_array_section_kind = OMP_ARRAY_SECTION_NONE; + c_omp_array_shaping_op_p = false; struct c_declspecs *scspecs; struct c_type_name *type_name; struct c_expr ret; @@ -10621,6 +10637,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_kind = save_c_omp_array_section_kind; + 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; @@ -10631,9 +10649,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_kind = save_c_omp_array_section_kind; + 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) @@ -10650,10 +10674,62 @@ 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_kind = save_c_omp_array_section_kind; + 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) + { + enum c_omp_array_section_kind save_c_omp_array_section_kind + = c_omp_array_section_kind; + bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p; + c_omp_array_section_kind = OMP_ARRAY_SECTION_NONE; + c_omp_array_shaping_op_p = false; + auto_vec<tree, 4> 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_kind = save_c_omp_array_section_kind; + 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 (cast_loc, 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); + } } @@ -11907,7 +11983,9 @@ c_parser_postfix_expression (c_parser *parser) /* A statement expression. */ tree stmt; location_t brace_loc; - bool save_c_omp_array_section_p = c_omp_array_section_p; + enum c_omp_array_section_kind save_c_omp_array_section_kind + = c_omp_array_section_kind; + 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); @@ -11924,7 +12002,8 @@ c_parser_postfix_expression (c_parser *parser) expr.set_error (); break; } - c_omp_array_section_p = false; + c_omp_array_section_kind = OMP_ARRAY_SECTION_NONE; + 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; @@ -11935,7 +12014,8 @@ c_parser_postfix_expression (c_parser *parser) expr.value = c_finish_stmt_expr (brace_loc, stmt); 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_section_kind = save_c_omp_array_section_kind; + c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p; } else { @@ -14043,11 +14123,14 @@ warn_for_abs (location_t loc, tree fndecl, tree arg) "of value", fndecl, atype, ftype); } - /* Parse a postfix expression after the initial primary or compound literal; that is, parse a series of postfix operators. - EXPR_LOC is the location of the primary expression. */ + EXPR_LOC is the location of the primary expression. + + With c_omp_array_section_kind, it also handles OpenMP array sections of the + type [ index : length : stride ] where index, length and stride are optional. + Note that an absent length might be parsed as CPP_SCOPE ('::'). */ static struct c_expr c_parser_postfix_expression_after_primary (c_parser *parser, @@ -14072,36 +14155,81 @@ c_parser_postfix_expression_after_primary (c_parser *parser, switch (c_parser_peek_token (parser)->type) { case CPP_OPEN_SQUARE: - /* Array reference. */ - c_parser_consume_token (parser); - idx = len = NULL_TREE; - if (!c_omp_array_section_p - || c_parser_next_token_is_not (parser, CPP_COLON)) - idx = c_parser_expression (parser).value; + { + /* Array reference. */ + c_parser_consume_token (parser); + idx = len = NULL_TREE; + if (c_omp_array_section_kind == OMP_ARRAY_SECTION_NONE + || (c_parser_next_token_is_not (parser, CPP_COLON) + && c_parser_next_token_is_not (parser, CPP_SCOPE))) + idx = c_parser_expression (parser).value; - if (c_omp_array_section_p - && c_parser_next_token_is (parser, CPP_COLON)) - { + bool double_colon + = c_parser_next_token_is (parser, CPP_SCOPE) + || (c_parser_next_token_is (parser, CPP_COLON) + && c_parser_peek_2nd_token (parser)->type == CPP_COLON); + if (double_colon && c_parser_next_token_is (parser, CPP_COLON)) c_parser_consume_token (parser); - if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE)) - len = c_parser_expression (parser).value; - expr.value = build_omp_array_section (op_loc, expr.value, idx, - len, NULL_TREE /* fixme */); - } - else - expr.value = build_array_ref (op_loc, expr.value, idx); + if (c_omp_array_section_kind != OMP_ARRAY_SECTION_NONE + && c_parser_next_token_is (parser, CPP_COLON) && !double_colon) + { + tree stride = NULL_TREE; - c_parser_skip_until_found (parser, CPP_CLOSE_SQUARE, - "expected %<]%>"); + c_parser_consume_token (parser); + if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE)) + len = c_parser_expression (parser).value; - start = expr.get_start (); - finish = parser->tokens_buf[0].location; - set_c_expr_source_range (&expr, start, finish); - expr.original_code = ERROR_MARK; - expr.original_type = NULL; - expr.m_decimal = 0; - break; + if (c_omp_array_section_kind == OMP_ARRAY_SECTION_STRIDED + && 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; + } + + expr.value = build_omp_array_section (op_loc, expr.value, idx, + len, stride); + } + else if (c_omp_array_section_kind == OMP_ARRAY_SECTION_STRIDED + && double_colon) + { + tree stride = NULL_TREE; + c_parser_consume_token (parser); + if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE)) + stride = c_parser_expression (parser).value; + expr.value = build_omp_array_section (op_loc, expr.value, idx, + len, stride); + } + else + { + 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); + } + + c_parser_skip_until_found (parser, CPP_CLOSE_SQUARE, + "expected %<]%>"); + + start = expr.get_start (); + finish = parser->tokens_buf[0].location; + set_c_expr_source_range (&expr, start, finish); + expr.original_code = ERROR_MARK; + expr.original_type = NULL; + expr.m_decimal = 0; + break; + } case CPP_OPEN_PAREN: /* Function call. */ { @@ -14442,8 +14570,11 @@ c_parser_expr_list (c_parser *parser, bool convert_p, bool fold_p, vec<tree, va_gc> *orig_types; struct c_expr expr; unsigned int idx = 0; - bool save_c_omp_array_section_p = c_omp_array_section_p; - c_omp_array_section_p = false; + enum c_omp_array_section_kind save_c_omp_array_section_kind + = c_omp_array_section_kind; + bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p; + c_omp_array_section_kind = OMP_ARRAY_SECTION_NONE; + c_omp_array_shaping_op_p = false; ret = make_tree_vector (); if (p_orig_types == NULL) @@ -14524,7 +14655,8 @@ 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_section_kind = save_c_omp_array_section_kind; + c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p; return ret; } @@ -16789,6 +16921,24 @@ 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 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 = array_type_nelts_top (type); + type = TREE_TYPE (type); + while (TREE_CODE (type) == ARRAY_TYPE) + { + tree n = 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 @@ -16916,13 +17066,28 @@ c_parser_omp_variable_list (c_parser *parser, || kind == OMP_CLAUSE_FROM)) { location_t loc = c_parser_peek_token (parser)->location; - bool save_c_omp_array_section_p = c_omp_array_section_p; - c_omp_array_section_p = true; + enum c_omp_array_section_kind save_c_omp_array_section_kind + = c_omp_array_section_kind; + bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p; + c_omp_array_section_kind = kind == OMP_CLAUSE_MAP + ? OMP_ARRAY_SECTION_UNSTRIDED + : OMP_ARRAY_SECTION_STRIDED; + 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_section_kind = save_c_omp_array_section_kind; + 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 @@ -16935,6 +17100,33 @@ 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; + sections++; + } + decl = TREE_OPERAND (decl, 0); + } + } + + decl = orig_decl; + while (TREE_CODE (decl) == OMP_ARRAY_SECTION) { tree low_bound = TREE_OPERAND (decl, 1); @@ -16943,18 +17135,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, @@ -16967,6 +17204,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); } } @@ -16993,6 +17259,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); @@ -21471,7 +21745,7 @@ c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind, if (colon_pos) c_parser_require (parser, CPP_COLON, "expected %<:%>"); - 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) @@ -27207,8 +27481,38 @@ c_parser_omp_target_update (location_t loc, c_parser *parser, clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_UPDATE); clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_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 35d44f5ac88..65b02a094de 100644 --- a/gcc/c/c-tree.h +++ b/gcc/c/c-tree.h @@ -795,7 +795,14 @@ extern int in_countof; extern int in_typeof; extern int in_generic; extern bool c_in_omp_for; -extern bool c_omp_array_section_p; +enum c_omp_array_section_kind : unsigned char { + OMP_ARRAY_SECTION_NONE = 0, + OMP_ARRAY_SECTION_UNSTRIDED = 1, + OMP_ARRAY_SECTION_STRIDED = 2 +}; +extern enum c_omp_array_section_kind c_omp_array_section_kind; +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; @@ -875,6 +882,8 @@ extern tree build_component_ref (location_t, tree, tree, location_t, extern tree handle_counted_by_for_component_ref (location_t, tree); 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 (location_t loc, tree expr, + vec<tree> *omp_shape_dims); extern tree build_external_ref (location_t, tree, bool, tree *); extern void pop_maybe_used (bool); extern struct maybe_used_decl *save_maybe_used (); diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 3ed109ff4a0..89654ae03ba 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -84,8 +84,15 @@ int in_generic; /* True when parsing OpenMP loop expressions. */ bool c_in_omp_for; -/* True when parsing OpenMP map clause. */ -bool c_omp_array_section_p; +/* The kind of OpenMP array section being parsed, if any. */ +enum c_omp_array_section_kind c_omp_array_section_kind; + +/* 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. */ @@ -3869,6 +3876,47 @@ mark_decl_used (tree ref, bool address) } } +/* 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 (location_t loc, tree expr, + vec<tree> *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_at (loc, "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 = c_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 @@ -3912,7 +3960,7 @@ build_external_ref (location_t loc, tree id, bool fun, tree *type) unmappable types if we return the decl with an error_mark_node type, rather than returning error_mark_node for the decl itself. */ if (TREE_TYPE (ref) == error_mark_node - && !c_omp_array_section_p) + && c_omp_array_section_kind == OMP_ARRAY_SECTION_NONE) return error_mark_node; if (TREE_UNAVAILABLE (ref)) @@ -15841,7 +15889,7 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses) T current expression (initially OMP_CLAUSE_DECL), which is either a TREE_LIST for array-section-subscript (TREE_PURPOSE is low-bound expression if specified, TREE_VALUE length expression if specified, - TREE_CHAIN is what it has been specified after, or some decl. + TREE_CHAIN is what it has been specified after), or some decl. TYPES vector is populated with array section types, MAYBE_ZERO_LEN set to true if any of the array-section-subscript could have length of zero (explicit or implicit), FIRST_NON_ONE is the index of the @@ -15854,12 +15902,19 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses) <= FIRST_NON_ONE we diagnose non-contiguous arrays if low bound isn't 0 or length isn't the array domain max + 1, for > FIRST_NON_ONE we can if MAYBE_ZERO_LEN is false. MAYBE_ZERO_LEN will be true in the above - case though, as some lengths could be zero. */ + case though, as some lengths could be zero. + On entry *DISCONTIGUOUS is 0 if the array section must be contiguous, and + non-zero if a discontiguous section is permitted; this function sets + *DISCONTIGUOUS to 2 when it determines the section is definitely + discontiguous, and otherwise leaves it unchanged. + Return the array reference or array-section tree built for T, or T itself + (possibly converted to an rvalue) if it is not an array section, or + ERROR_MARK_NODE on error. */ static tree handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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; @@ -15935,11 +15990,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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); @@ -15980,8 +16039,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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)) @@ -16047,6 +16113,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, omp_clause_code_name[OMP_CLAUSE_CODE (c)]); return error_mark_node; } + if (stride != NULL_TREE && TREE_CODE (stride) == INTEGER_CST + && (tree_int_cst_sgn (stride) == -1 + || tree_int_cst_sgn (stride) == 0)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "negative or zero stride in array section in %qs clause", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + return error_mark_node; + } if (TYPE_DOMAIN (type) && TYPE_MAX_VALUE (TYPE_DOMAIN (type)) && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (type))) @@ -16100,12 +16175,29 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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) @@ -16113,7 +16205,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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)) { @@ -16184,14 +16276,17 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, d = TREE_OPERAND (d, 0)) { tree d_length = TREE_OPERAND (d, 2); - tree d_stride = TREE_OPERAND (d, 3); - if (d_length == NULL_TREE || !integer_onep (d_length) - || (d_stride && !integer_onep (d_stride))) + if (d_length == NULL_TREE || !integer_onep (d_length)) { - 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 = 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; + } } } } @@ -16203,7 +16298,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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) @@ -16211,26 +16306,55 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, TREE_OPERAND (t, 1) = lb; low_bound = lb; } - ret = build_array_ref (OMP_CLAUSE_LOCATION (c), ret, low_bound); + if (*discontiguous) + 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. DISCONTIGUOUS_P is NULL, or points to + a bool that is false, if the array section must be contiguous. If + *DISCONTIGUOUS_P is true on entry, a discontiguous section is permitted; + on exit it is left true only if the section was determined to be + (possibly) discontiguous, and reset to false if it turned out to be + contiguous after all. Return true on error. */ static bool -handle_omp_array_sections (tree &c, enum c_omp_region_type ort) +handle_omp_array_sections (tree &c, enum c_omp_region_type ort, + bool *discontiguous_p = NULL) { bool maybe_zero_len = false; unsigned int first_non_one = 0; auto_vec<tree, 10> types; tree *tp = &OMP_CLAUSE_DECL (c); + int discontiguous = discontiguous_p && *discontiguous_p ? 1 : 0; + if (discontiguous_p) + *discontiguous_p = false; if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY) && OMP_ITERATOR_DECL_P (*tp)) 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) @@ -16268,11 +16392,14 @@ handle_omp_array_sections (tree &c, 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 @@ -16285,12 +16412,63 @@ handle_omp_array_sections (tree &c, 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) + { + /* 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; + } + + /* A stride other than one only matters when the length is + also other than one; a length-one dimension touches a + single element regardless of stride. */ + bool length_is_one = length != NULL_TREE && integer_onep (length); + bool strided = !integer_onep (stride) && !length_is_one; + + if (strided + || (higher_discontiguous + && (!integer_zerop (low_bound) + || !full_span))) + discontiguous = 2; + + if (strided + || !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]) @@ -16304,12 +16482,17 @@ handle_omp_array_sections (tree &c, 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 = 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 @@ -16421,6 +16604,10 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort) OMP_CLAUSE_DECL (c) = t; return false; } + if (discontiguous_p) + *discontiguous_p = discontiguous == 2; + if (discontiguous == 1) + 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) @@ -16436,7 +16623,9 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort) OMP_CLAUSE_SIZE (c) = size; } - if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) + /* Done, except for MAP clauses and for array sections that may be + noncontiguous. */ + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP && discontiguous != 2) return false; auto_vec<omp_addr_token *, 10> addr_tokens; @@ -16451,17 +16640,23 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort) { using namespace omp_addr_tokenizer; - 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; /* !!! If we're accessing a base decl via chained access methods (e.g. multiple indirections), duplicate clause - detection won't work properly. Skip it in that case. */ - if ((addr_tokens[0]->type == STRUCTURE_BASE - || addr_tokens[0]->type == ARRAY_BASE) - && addr_tokens[0]->u.structure_base_kind == BASE_DECL - && addr_tokens[1]->type == ACCESS_METHOD - && omp_access_chain_p (addr_tokens, 1)) + detection won't work properly. Skip it in that case. + Noncontiguous (strided/shaped) array updates also replace C + with a differently-kinded clause (regardless of how deep the + access chain is, e.g. for a struct member), so the caller + needs NC for those too. */ + if (((addr_tokens[0]->type == STRUCTURE_BASE + || addr_tokens[0]->type == ARRAY_BASE) + && addr_tokens[0]->u.structure_base_kind == BASE_DECL + && addr_tokens[1]->type == ACCESS_METHOD + && omp_access_chain_p (addr_tokens, 1)) + || discontiguous == 2) c = nc; return false; @@ -17688,6 +17883,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: @@ -17709,10 +17907,25 @@ 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 (c, ort)) + bool discontiguous + = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM); + if (handle_omp_array_sections (c, ort, &discontiguous)) remove = true; else { + /* If we ended up with a genuinely discontiguous update, + handle_omp_array_sections replaced C with a + differently-kinded (GOMP_MAP_TO_GRID/FROM_GRID) + clause; splice that replacement into the actual + clause chain here, since PC still points at the + original clause's slot. (This is unlike the + pre-existing "chained access" case, where C is + reassigned only to advance past appended nodes that + stay chained after the original, unmoved clause.) */ + if (discontiguous) + *pc = c; + t = OMP_CLAUSE_DECL (c); if (!omp_mappable_type (TREE_TYPE (t))) { diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-1.c b/gcc/testsuite/c-c++-common/gomp/array-section-1.c new file mode 100644 index 00000000000..5af3f950dd7 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/array-section-1.c @@ -0,0 +1,18 @@ +/* { dg-do compile } */ + +/* Check that array section syntax with empty length is handled correctly. */ + +int n; + +void f() { + int A[10], B[10][10]; +#pragma omp target update to(A[::], B[::][::1]) + +#pragma omp target update to(A[::], B[0][: : ]) +#pragma omp target update to(A[1:2:], B[:1:][1: : ]) + +#pragma omp target update to(A[:::n], B[:::n][: :: n : ]) +/* { dg-error "expected expression before ':' token" "" { target c } .-1 } */ +/* { dg-error "expected expression before '::' token" "" { target c } .-2 } */ +/* { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target c } .-3 } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-2.c b/gcc/testsuite/c-c++-common/gomp/array-section-2.c new file mode 100644 index 00000000000..4daa0a6b497 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/array-section-2.c @@ -0,0 +1,32 @@ +/* { dg-do compile } */ + +/* A stride specifier is only valid for "to"/"from" clauses (including on + "target update"), not for "map" clauses -- not even a "map" clause + using the "to"/"from" modifier. */ + +int main () +{ + int x[10]; + +#pragma omp target enter data map(to: x) + +#pragma omp target map(x[0:5:2]) + // { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 } + // { dg-error {expected '\)' before ':' token} "" { target c++ } .-2 } + // { dg-error "expected an OpenMP clause before '\\\]' token" "" { target c++ } .-3 } + ; + +#pragma omp target enter data map(to: x[0:5:2]) + // { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 } + // { dg-error {expected '\)' before ':' token} "" { target c++ } .-2 } + // { dg-error "expected an OpenMP clause before '\\\]' token" "" { target c++ } .-3 } + + /* These are fine: "to"/"from" clauses on "target update" do accept a + stride. */ +#pragma omp target update to(x[0:5:2]) +#pragma omp target update from(x[0:5:2]) + +#pragma omp target exit data map(release: x) + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-3.c b/gcc/testsuite/c-c++-common/gomp/array-section-3.c new file mode 100644 index 00000000000..ff2698cb1de --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/array-section-3.c @@ -0,0 +1,18 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original" } */ + +/* An explicit stride makes the section noncontiguous. Check that both a literal + and a variable stride route through the noncontiguous-update lowering: + to_grid, grid_dim and grid_stride. */ + +int arr[128]; + +void +g (int s) +{ +#pragma omp target update to(arr[::2]) + /* { dg-final { scan-tree-dump {map\(to_grid:arr \[len: 4\]\) map\(grid_dim:0 \[len: 64\]\) map\(grid_stride:2\)} "original" } } */ + +#pragma omp target update to(arr[::(s)]) + /* { dg-final { scan-tree-dump {map\(to_grid:arr \[len: 4\]\) map\(grid_dim:0 \[len: 128 /\[cl\] \(sizetype\) s\]\) map\(grid_stride:s\)} "original" } } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/array-shaping-1.c b/gcc/testsuite/c-c++-common/gomp/array-shaping-1.c new file mode 100644 index 00000000000..e8181b1b044 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/array-shaping-1.c @@ -0,0 +1,14 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original" } */ + +/* The reshaped array is contiguous overall and the section has no explicit + stride, but for runtime-variable bounds "f"/"g" the accessed memory may be + noncontiguous. + Check that this routes through the grid/noncontiguous-update lowering. */ + +void +foo (float *w, int e, int f, int g) +{ +#pragma omp target update to(([e][10]) w[3:2][f:g]) + /* { dg-final { scan-tree-dump {map\(to_grid:VIEW_CONVERT_EXPR.*\(\*w\) \[len: 4\]\) map\(grid_dim:3 \[len: 2\]\) map\(grid_dim:[^\[]+\[len: g\]\)} "original" } } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/bad-array-shaping-1.c b/gcc/testsuite/c-c++-common/gomp/bad-array-shaping-1.c new file mode 100644 index 00000000000..fafbfc86086 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/bad-array-shaping-1.c @@ -0,0 +1,12 @@ +/* { dg-do compile } */ + +/* The OpenMP array-shaping operator requires a pointer argument. */ + +int x; + +void f (void) +{ +#pragma omp target update to( ([10])x ) +/* { dg-error "OpenMP array shaping operator with non-pointer argument" "" { target *-*-* } .-1 } */ +/* { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target c++ } .-2 } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/bad-array-shaping-2.c b/gcc/testsuite/c-c++-common/gomp/bad-array-shaping-2.c new file mode 100644 index 00000000000..954545cf9d4 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/bad-array-shaping-2.c @@ -0,0 +1,25 @@ +/* { dg-do compile } */ + +#include <string.h> +#include <stdlib.h> + +int main (void) +{ + float *arr = (float *) 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/c-c++-common/gomp/bad-array-shaping-3.c b/gcc/testsuite/c-c++-common/gomp/bad-array-shaping-3.c new file mode 100644 index 00000000000..bfcd85ae628 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/bad-array-shaping-3.c @@ -0,0 +1,27 @@ +/* { 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 c } .-1 } */ +/* { dg-error {length '7' with stride '1' above array section size in 'to' clause} "" { target c++ } .-2 } */ +/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-3 } */ + +#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 {negative or zero stride in array section 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/array-shaping-1.c b/gcc/testsuite/gcc.dg/gomp/array-shaping-1.c new file mode 100644 index 00000000000..9bb08acf0ae --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/array-shaping-1.c @@ -0,0 +1,17 @@ +// { dg-do compile } + +/* The C front end used to ICE building the type for a variable-bound + dimension of the OpenMP array-shaping operator whenever the + resulting shaped expression was then further subscripted with an array + section. create_omp_arrayshape_type built each dimension with a + build_array_type call instead of *c_*build_array_type, leaving the + C-specific "variable size" bits inconsistent for later reuse. */ + +void +array_shaping (double *a, int nx, int ny) +{ +#pragma omp target update from( (([nx][ny + 2])a)[0:nx][1], \ + (([nx][ny + 2])a)[0:nx][ny] ) +#pragma omp target update to( (([nx][ny + 2])a)[0:nx][0], \ + (([nx][ny + 2])a)[0:nx][ny + 1] ) +} diff --git a/gcc/testsuite/gcc.dg/gomp/array-shaping-2.c b/gcc/testsuite/gcc.dg/gomp/array-shaping-2.c new file mode 100644 index 00000000000..67d5f1b2fc2 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/array-shaping-2.c @@ -0,0 +1,16 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original" } */ + +/* The reshaped array is contiguous overall and the section has no explicit + stride, but for a runtime-variable "len" the accessed memory may be + noncontiguous. + Check that this routes through the grid/noncontiguous-update lowering. + Check in particular that VIEW_CONVERT_EXPR applies to the bare pointer "a", + which is specific to the C front end. */ + +void +array_shaping (double *a, int nx, int ny, int lo, int len) +{ +#pragma omp target update from((([nx][ny]) a)[0:nx][lo:len]) + /* { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR.*\(a\) \[len: 8\]\) map\(grid_dim:0 \[len: nx\]\) map\(grid_dim:[^\[]+\[len: len\]\)} "original" } } */ +} 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 <string.h> +#include <assert.h> +#include <stdlib.h> + +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 <string.h> +#include <assert.h> +#include <stdlib.h> + +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 <string.h> +#include <assert.h> +#include <stdlib.h> + +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-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-c++-common/array-shaping-14.c b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-14.c new file mode 100644 index 00000000000..4ca6f794f93 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-14.c @@ -0,0 +1,34 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ + +#include <stdlib.h> +#include <stdio.h> +#include <assert.h> + +typedef struct { + int *ptr; +} S; + +int main(void) +{ + S q; + q.ptr = (int *) calloc (9 * 11, sizeof (int)); + +#pragma omp target enter data map(to: q.ptr, q.ptr[0:9*11]) + +#pragma omp target + for (int i = 0; i < 9*11; i++) + q.ptr[i] = i; + +#pragma omp target update from(([9][11]) q.ptr[3:3:2][1:4:3]) + + for (int j = 0; j < 9; j++) + for (int i = 0; i < 11; i++) + if (j >= 3 && j <= 7 && ((j - 3) % 2) == 0 + && i >= 1 && i <= 10 && ((i - 1) % 3) == 0) + assert (q.ptr[j * 11 + i] == j * 11 + i); + else + assert (q.ptr[j * 11 + i] == 0); + +#pragma omp target exit data map(release: q.ptr, q.ptr[0:9*11]) + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-shaping-15.c b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-15.c new file mode 100644 index 00000000000..34867e9e0e0 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-15.c @@ -0,0 +1,41 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ + +/* The reshaped array is contiguous overall and the section has no explicit + stride, but for runtime-variable bounds "f"/"g" the accessed memory may be + noncontiguous. */ + +#include <string.h> +#include <assert.h> + +void +foo (float *w, int e, int f, int g) +{ + memset (w, 0, sizeof (float) * 100); + +#pragma omp target enter data map(to: w[:100]) + + for (int j = 0; j < e; j++) + for (int i = 0; i < 10; i++) + w[j * 10 + i] = i + j; + +#pragma omp target update to(([e][10]) w[3:2][f:g]) + +#pragma omp target exit data map(from: w[:100]) + + for (int j = 0; j < e; j++) + for (int i = 0; i < 10; i++) + if (j >= 3 && j < 5 && i >= f && i < f + g) + assert (w[j * 10 + i] == i + j); + else + assert (w[j * 10 + i] == 0.0f); +} + +int +main (void) +{ + float arr[100]; + + foo (arr, 10, 1, 8); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-shaping-17.c b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-17.c new file mode 100644 index 00000000000..f0dc710b65f --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-17.c @@ -0,0 +1,40 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ + +#include <stdlib.h> +#include <string.h> +#include <assert.h> + +void +foo (int *w) +{ + memset (w, 0, sizeof (int) * 100); + +#pragma omp target enter data map(to: w[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + w[j * 10 + i] = i + j; + +#pragma omp target update to(([10][10]) w[3:2][1:8]) + +#pragma omp target exit data map(from: w[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (j >= 3 && j < 5 && i >= 1 && i < 9) + assert (w[j * 10 + i] == i + j); + else + assert (w[j * 10 + i] == 0); +} + +int +main (void) +{ + int *arr = (int *) calloc (100, sizeof (int)); + + foo (arr); + + free (arr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-shaping-18.c b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-18.c new file mode 100644 index 00000000000..64e963d4629 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-18.c @@ -0,0 +1,236 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include <string.h> +#include <assert.h> +#include <stdlib.h> + +volatile int yy = 4, zz = 2, str_str = 2; + +int main() +{ + int *arr; + int x = 5; + int arr2d[10][10]; + + arr = (int *) 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-c++-common/array-shaping-19.c b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-19.c new file mode 100644 index 00000000000..2b9e6949b60 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-19.c @@ -0,0 +1,36 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include <assert.h> +#include <string.h> + +#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-c++-common/array-shaping-20.c b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-20.c new file mode 100644 index 00000000000..1034682e4ca --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-20.c @@ -0,0 +1,38 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include <assert.h> +#include <string.h> + +#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-c++-common/array-shaping-21.c b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-21.c new file mode 100644 index 00000000000..753bf0c2605 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-21.c @@ -0,0 +1,45 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include <assert.h> +#include <stdlib.h> +#include <string.h> + +#define N 10 + +int main () +{ + int *iptr = (int *) 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; +} 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 <assert.h> +#include <stdlib.h> + +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 <assert.h> +#include <stdlib.h> +#include <string.h> + +#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; +} -- 2.55.0
