[gcc/devel/omp/gcc-14] OpenMP: Array shaping operator and strided "target update" for C
Paul-Antoine Arras
parras@gcc.gnu.org
Fri Jun 28 09:54:53 GMT 2024
https://gcc.gnu.org/g:16968aa143637e56a3834b8615b18d0a915498de
commit 16968aa143637e56a3834b8615b18d0a915498de
Author: Julian Brown <julian@codesourcery.com>
Date: Fri May 3 11:26:08 2024 +0200
OpenMP: Array shaping operator and strided "target update" for C
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 <julian@codesourcery.com>
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.
Diff:
---
gcc/ChangeLog.omp | 11 +
gcc/c/ChangeLog.omp | 33 ++
gcc/c/c-parser.cc | 305 ++++++++++++-
gcc/c/c-tree.h | 4 +
gcc/c/c-typeck.cc | 232 ++++++++--
gcc/fortran/ChangeLog.omp | 10 +
gcc/fortran/trans-openmp.cc | 500 +++++++++++++++++++++
gcc/gimplify.cc | 10 +
gcc/omp-low.cc | 73 ++-
gcc/testsuite/ChangeLog.omp | 17 +
gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c | 26 ++
gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-2.c | 24 +
gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-3.c | 30 ++
gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-4.c | 27 ++
gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-5.c | 17 +
gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-6.c | 26 ++
gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-7.c | 15 +
.../gfortran.dg/gomp/noncontig-updates-1.f90 | 19 +
.../gfortran.dg/gomp/noncontig-updates-2.f90 | 16 +
.../gfortran.dg/gomp/noncontig-updates-3.f90 | 16 +
.../gfortran.dg/gomp/noncontig-updates-4.f90 | 15 +
libgomp/ChangeLog.omp | 30 ++
libgomp/libgomp.h | 1 +
libgomp/target.c | 42 +-
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 ++
.../libgomp.fortran/noncontig-updates-1.f90 | 54 +++
.../libgomp.fortran/noncontig-updates-10.f90 | 29 ++
.../libgomp.fortran/noncontig-updates-11.f90 | 51 +++
.../libgomp.fortran/noncontig-updates-12.f90 | 59 +++
.../libgomp.fortran/noncontig-updates-13.f90 | 42 ++
.../libgomp.fortran/noncontig-updates-2.f90 | 101 +++++
.../libgomp.fortran/noncontig-updates-3.f90 | 47 ++
.../libgomp.fortran/noncontig-updates-4.f90 | 78 ++++
.../libgomp.fortran/noncontig-updates-5.f90 | 55 +++
.../libgomp.fortran/noncontig-updates-6.f90 | 34 ++
.../libgomp.fortran/noncontig-updates-7.f90 | 36 ++
.../libgomp.fortran/noncontig-updates-8.f90 | 39 ++
.../libgomp.fortran/noncontig-updates-9.f90 | 34 ++
43 files changed, 2522 insertions(+), 72 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 56fb41704cc..1d8f67992e8 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,14 @@
+2023-07-03 Julian Brown <julian@codesourcery.com>
+
+ * gimplify.cc (gimplify_adjust_omp_clauses): Don't gimplify
+ VIEW_CONVERT_EXPR away in GOMP_MAP_TO_GRID/GOMP_MAP_FROM_GRID clauses.
+ * omp-low.cc (omp_noncontig_descriptor_type): Add SPAN field.
+ (scan_sharing_clauses): Don't store descriptor size in its
+ OMP_CLAUSE_SIZE field.
+ (lower_omp_target): Add missing OMP_CLAUSE_MAP check. Add special-case
+ string handling. Handle span and bias. Use low bound instead of zero
+ as index for trailing full dimensions.
+
2023-07-03 Julian Brown <julian@codesourcery.com>
* gimplify.cc (omp_group_last, omp_group_base): Add GOMP_MAP_TO_GRID,
diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp
index 125f7e550d2..f0b91cb5403 100644
--- a/gcc/c/ChangeLog.omp
+++ b/gcc/c/ChangeLog.omp
@@ -1,3 +1,36 @@
+2023-09-05 Julian Brown <julian@codesourcery.com>
+
+ * 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.
+
2023-09-05 Julian Brown <julian@codesourcery.com>
* c-parser.cc (c_parser_postfix_expression_after_primary): Dummy stride
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 20126243bb0..899dc901589 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -6178,7 +6178,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)
@@ -6218,6 +6220,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)
{
@@ -9192,6 +9195,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 ());
@@ -9200,6 +9204,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
@@ -9253,6 +9258,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;
}
{
@@ -9300,6 +9306,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;
}
@@ -9681,6 +9688,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
@@ -9689,6 +9698,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;
@@ -9700,6 +9713,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;
@@ -9710,9 +9725,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)
@@ -9729,10 +9750,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<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_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).
@@ -10752,6 +10824,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);
@@ -10769,6 +10842,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;
@@ -10780,6 +10854,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
{
@@ -12577,15 +12652,38 @@ 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;
+ }
+
expr.value = build_omp_array_section (op_loc, expr.value, idx,
- len, NULL_TREE /* fixme */);
+ len, stride);
}
else
- 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);
+ }
c_parser_skip_until_found (parser, CPP_CLOSE_SQUARE,
"expected %<]%>");
@@ -12783,8 +12881,8 @@ c_parser_postfix_expression_after_primary (c_parser *parser,
finish = c_parser_peek_token (parser)->get_finish ();
c_parser_consume_token (parser);
expr = default_function_array_read_conversion (expr_loc, expr);
- expr.value = build_unary_op (op_loc, POSTINCREMENT_EXPR,
- expr.value, false);
+ expr.value
+ = build_unary_op (op_loc, POSTINCREMENT_EXPR, expr.value, false);
set_c_expr_source_range (&expr, start, finish);
expr.original_code = ERROR_MARK;
expr.original_type = NULL;
@@ -12917,7 +13015,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)
@@ -12972,6 +13072,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;
}
@@ -15194,6 +15295,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
@@ -15323,12 +15453,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
@@ -15341,6 +15483,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);
@@ -15349,18 +15516,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,
@@ -15373,6 +15585,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);
}
}
@@ -15399,6 +15640,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);
@@ -19518,7 +19767,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, C_ORT_OMP);
+ tree nl = c_parser_omp_variable_list (parser, loc, kind, list, C_ORT_OMP, true);
parens.skip_until_found_close (parser);
if (present)
@@ -24527,8 +24776,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 bf33a9c61cb..9238831913e 100644
--- a/gcc/c/c-tree.h
+++ b/gcc/c/c-tree.h
@@ -739,6 +739,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;
@@ -781,6 +783,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<tree> *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 5efb7e9ef4f..a1a2f10c05c 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;
@@ -2971,6 +2978,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<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 ("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
@@ -13865,7 +13912,7 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses)
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
bool &maybe_zero_len, unsigned int &first_non_one,
- bool &non_contiguous, enum c_omp_region_type ort)
+ bool &non_contiguous, enum c_omp_region_type ort, int *discontiguous)
{
tree ret, low_bound, length, stride, type;
bool openacc = (ort & C_ORT_ACC) != 0;
@@ -13946,11 +13993,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,
- non_contiguous, ort);
+ non_contiguous, ort,
+ discontiguous);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
- type = TREE_TYPE (ret);
+ if (TREE_CODE (ret) == OMP_ARRAY_SECTION)
+ type = TREE_TYPE (TREE_TYPE (TREE_OPERAND (ret, 0)));
+ else
+ type = TREE_TYPE (ret);
low_bound = TREE_OPERAND (t, 1);
length = TREE_OPERAND (t, 2);
stride = TREE_OPERAND (t, 3);
@@ -13991,8 +14042,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))
@@ -14111,12 +14169,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)
@@ -14124,7 +14199,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))
{
@@ -14220,10 +14295,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
- error_at (OMP_CLAUSE_LOCATION (c),
- "array section is not contiguous in %qs clause",
- omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
- return error_mark_node;
+ if (discontiguous && *discontiguous)
+ *discontiguous = 2;
+ else
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "array section is not contiguous in %qs clause",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ return error_mark_node;
+ }
}
}
}
@@ -14235,7 +14315,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)
@@ -14243,14 +14323,42 @@ 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);
+ /* 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, tree **pnext, enum c_omp_region_type ort)
+handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort,
+ int *discontiguous)
{
tree c = *pc;
bool maybe_zero_len = false;
@@ -14266,7 +14374,7 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
tp = &TREE_VALUE (*tp);
tree first = handle_omp_array_sections_1 (c, *tp, types,
maybe_zero_len, first_non_one,
- non_contiguous, ort);
+ non_contiguous, ort, discontiguous);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -14305,11 +14413,14 @@ handle_omp_array_sections (tree *pc, tree **pnext, 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
@@ -14322,6 +14433,11 @@ handle_omp_array_sections (tree *pc, tree **pnext, 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;
@@ -14331,10 +14447,49 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
continue;
}
+ if (stride == NULL_TREE)
+ stride = size_one_node;
+ if (discontiguous && *discontiguous)
+ {
+ /* This condition is similar to the error check below, but
+ whereas that checks for a definitely-discontiguous array
+ section in order to report an error (where such a section is
+ illegal), here we instead need to know if the array section
+ *may be* discontiguous so we can handle that case
+ appropriately (i.e. for rectangular "target update"
+ operations). */
+ bool full_span = false;
+ if (length != NULL_TREE
+ && TREE_CODE (length) == INTEGER_CST
+ && TREE_CODE (types[i]) == ARRAY_TYPE
+ && TYPE_DOMAIN (types[i])
+ && TYPE_MAX_VALUE (TYPE_DOMAIN (types[i]))
+ && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])))
+ == INTEGER_CST)
+ {
+ tree size;
+ size = size_binop (PLUS_EXPR,
+ TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])),
+ size_one_node);
+ if (tree_int_cst_equal (length, size))
+ full_span = true;
+ }
+
+ if (!integer_onep (stride)
+ || (higher_discontiguous
+ && (!integer_zerop (low_bound)
+ || !full_span)))
+ *discontiguous = 2;
+
+ if (!integer_onep (stride)
+ || !integer_zerop (low_bound)
+ || !full_span)
+ higher_discontiguous = true;
+ }
if (!maybe_zero_len && i > first_non_one)
{
if (integer_nonzerop (low_bound))
- goto do_warn_noncontiguous;
+ goto is_noncontiguous;
if (length != NULL_TREE
&& TREE_CODE (length) == INTEGER_CST
&& TYPE_DOMAIN (types[i])
@@ -14348,12 +14503,17 @@ handle_omp_array_sections (tree *pc, tree **pnext, 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
@@ -14473,6 +14633,8 @@ handle_omp_array_sections (tree *pc, tree **pnext, 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)
@@ -14488,7 +14650,8 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
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<omp_addr_token *, 10> addr_tokens;
@@ -14505,7 +14668,8 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
c = *pc;
- 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
@@ -14866,7 +15030,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, NULL, ort))
+ if (handle_omp_array_sections (pc, NULL, ort, NULL))
{
remove = true;
break;
@@ -15596,7 +15760,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, NULL, ort))
+ if (handle_omp_array_sections (pc, NULL, ort, NULL))
remove = true;
else if ((c = *pc)
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
@@ -15703,6 +15867,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:
@@ -15718,7 +15885,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
grp_sentinel = OMP_CLAUSE_CHAIN (c);
tree *pnext = NULL;
- if (handle_omp_array_sections (pc, &pnext, ort))
+ int discontiguous
+ = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM);
+ if (handle_omp_array_sections (pc, &pnext, ort, &discontiguous))
remove = true;
else
{
@@ -16121,7 +16291,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, NULL, ort))
+ if (handle_omp_array_sections (pc, NULL, ort, NULL))
remove = true;
else
{
diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index 30610d7e699..3794560b606 100644
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,13 @@
+2023-07-03 Julian Brown <julian@codesourcery.com>
+
+ * trans-openmp.cc (gfc_omp_deep_map_kind_p): Handle
+ GOMP_MAP_{TO,FROM}_GRID, GOMP_MAP_GRID_{DIM,STRIDE}.
+ (gfc_trans_omp_arrayshape_type, gfc_omp_calculate_gcd,
+ gfc_desc_to_omp_noncontig_array, gfc_omp_contiguous_update_p): New
+ functions.
+ (gfc_trans_omp_clauses): Handle noncontiguous to/from clauses for OMP
+ "target update" directives.
+
2023-08-10 Julian Brown <julian@codesourcery.com>
* gfortran.h (gfc_omp_namelist_udm): Add MAPPER_ID field to store the
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index d23359e6dce..e4475cfc403 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -2991,6 +2991,10 @@ gfc_omp_deep_map_kind_p (tree clause)
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ATTACH_DETACH:
+ case GOMP_MAP_TO_GRID:
+ case GOMP_MAP_FROM_GRID:
+ case GOMP_MAP_GRID_DIM:
+ case GOMP_MAP_GRID_STRIDE:
break;
default:
gcc_unreachable ();
@@ -4248,6 +4252,346 @@ get_symbol_rooted_namelist (hash_map<gfc_symbol *,
return NULL;
}
+/* We build an "un-Fortrannish" array-of-arrays here to pass our calculated
+ array bounds to the middle end for strided/rectangular OpenMP
+ "target update" operations. */
+
+static tree
+gfc_trans_omp_arrayshape_type (tree type, vec<tree> *dims)
+{
+ gcc_assert (dims->length () > 0);
+
+ for (int i = dims->length () - 1; i >= 0; i--)
+ {
+ tree dim = fold_convert (sizetype, (*dims)[i]);
+ /* We need the index of the last element, not the array size. */
+ dim = size_binop (MINUS_EXPR, dim, size_one_node);
+ tree idxtype = build_index_type (dim);
+ type = build_array_type (type, idxtype);
+ }
+
+ return type;
+}
+
+/* Emit code to find the greatest common divisor of two (gfc_array_index_type)
+ trees to BLOCK. This is Euclid's algorithm:
+
+ int
+ gcd (int a, int b)
+ {
+ int tmp;
+ while (b != 0)
+ {
+ tmp = b;
+ b = a % b;
+ a = tmp;
+ }
+ return a;
+ }
+*/
+
+static void
+gfc_omp_calculate_gcd (stmtblock_t *block, tree dst, tree a, tree b)
+{
+ tree tmp = gfc_create_var (gfc_array_index_type, "tmp");
+ tree avar = gfc_create_var (gfc_array_index_type, "a");
+ tree bvar = gfc_create_var (gfc_array_index_type, "b");
+
+ /* Avoid clobbering the inputs. */
+ gfc_add_modify (block, avar, a);
+ gfc_add_modify (block, bvar, b);
+
+ tree label_cond = gfc_build_label_decl (NULL_TREE);
+ tree label_loop = gfc_build_label_decl (NULL_TREE);
+ TREE_USED (label_cond) = 1;
+ TREE_USED (label_loop) = 1;
+
+ gfc_add_expr_to_block (block, build1_v (GOTO_EXPR, label_cond));
+ gfc_add_expr_to_block (block, build1_v (LABEL_EXPR, label_loop));
+
+ gfc_add_modify (block, tmp, bvar);
+ gfc_add_modify (block, bvar,
+ fold_build2_loc (input_location, TRUNC_MOD_EXPR,
+ gfc_array_index_type, avar, bvar));
+ gfc_add_modify (block, avar, tmp);
+
+ gfc_add_expr_to_block (block, build1_v (LABEL_EXPR, label_cond));
+
+ tmp = fold_build2_loc (input_location, NE_EXPR, boolean_type_node, bvar,
+ gfc_index_zero_node);
+ tmp = build3_v (COND_EXPR, tmp, build1_v (GOTO_EXPR, label_loop),
+ build_empty_stmt (input_location));
+ gfc_add_expr_to_block (block, tmp);
+
+ gfc_add_modify (block, dst, avar);
+}
+
+/* Convert a gfortran array descriptor -- specifically the per-dimension
+ strides -- into a form that can be easily translated to a noncontiguous
+ OpenMP "target update" operation. We emit a specialized version of a
+ function like this inline:
+
+ void
+ gfc_desc_to_omp_noncontig_array (int *dims, int *strides, int ndims,
+ int *fstrides, int *flo, int *fhi)
+ {
+ dims[ndims - 1] = (fhi[ndims - 1] - flo[ndims - 1] + 1);
+ strides[0] = fstrides[0];
+ if (ndims > 1)
+ strides[ndims - 1] = 1;
+ if (ndims == 2)
+ dims[0] = fstrides[1];
+ else if (ndims > 2)
+ {
+ int grains[ndims - 2];
+
+ int bigger_grain = fstrides[ndims - 1];
+ for (int i = ndims - 2; i > 0; i--)
+ {
+ grains[i - 1] = gcd (fstrides[i], bigger_grain);
+ bigger_grain = grains[i - 1];
+ }
+
+ int volume = 1;
+ for (int i = 0; i < ndims - 2; i++)
+ {
+ int g = grains[i];
+ dims[i] = g / volume;
+ strides[i + 1] = fstrides[i + 1] / g;
+ volume = volume * dims[i];
+ }
+ dims[ndims - 2] = fstrides[ndims - 1] / volume;
+ }
+ }
+
+ where "fstrides", "flo" and "fhi" represent the stride, low bound and upper
+ bound of each dimension in the Fortran array descriptor.
+
+ (Note that most of the complexity only applies to arrays with more than two
+ dimensions, and the final stanza won't be emitted at all for lower-ranked
+ arrays.)
+
+ The output of the algorithm is a set of dimensions dims[] = { D, C, B, A }
+ "as if" the array was declared like this (in C):
+
+ type arr[A][B][C][D];
+
+ i.e. with the innermost dimension first, and a set of strides (in terms of
+ the step size along each dimension, without previous dimensions multiplied
+ in).
+
+ As an example, if we have an array:
+
+ allocate (arr(18,19,20,21,22))
+
+ and an update operation:
+
+ !$omp target update to(arr(1:3:2,1:4:3,1:5:4,1:6:5,1:7:6))
+
+ the strides we see in the Fortran array descriptor will be:
+
+ 2 54 1368 34200 861840
+
+ as given by:
+
+ 2 = stride0
+ 54 = dim0 * stride1
+ 1368 = dim0 * dim1 * stride2
+ 34200 = dim0 * dim1 * dim2 * stride3
+ 861840 = dim0 * dim1 * dim2 * dim3 * stride4
+
+ where "dimN" are the extents of each dimension (18,19,20,21,22), and
+ "strideN" are the strides in terms of step length along each dimension
+ (2,3,4,5,6).
+
+ We'd like to figure out what the original dimN, strideN were from the
+ Fortran array descriptor, but that's in general impossible. Furthermore,
+ if we naively divide a stride by the preceding stride, the result isn't
+ necessarily an integer, as for e.g.:
+
+ 861840/34200 = 25.2
+
+ What we can do though is figure out the greatest common divisor of
+ each stride and the preceding one, from the largest down, and use those as
+ units of granularity, i.e. the size of the corresponding dimension we pass
+ to the middle-end/runtime. The stepwise stride is then the number of
+ times each "grain" fits into the Fortran array descriptor stride.
+
+ The output of the algorithm will be:
+
+ dims strides
+ 18 2
+ 76 3
+ 5 1
+ 126 5
+ 9 1
+
+ These numbers work fine for libgomp target.c:omp_target_memcpy_rect_worker.
+ Multiplying them through also gives the same numbers as the source Fortran
+ array strides, i.e. dim0*dim1*dim2*stride3 (18*76*5*5) = 34200. */
+
+static void
+gfc_desc_to_omp_noncontig_array (stmtblock_t *block, tree *ompdimsp,
+ tree *ompstridesp, tree desc, int ndims)
+{
+ tree lastdim = build_int_cst (gfc_array_index_type, ndims - 1);
+ tree dimrange = build_index_type (lastdim);
+ tree ndimarrtype = build_array_type (gfc_array_index_type, dimrange);
+ tree ompdims = gfc_create_var (ndimarrtype, "dims");
+ tree ompstrides = gfc_create_var (ndimarrtype, "strides");
+
+ *ompdimsp = ompdims;
+ *ompstridesp = ompstrides;
+
+ /* dims[ndims - 1] = (fhi[ndims - 1] - flo[ndims - 1] + 1); */
+ tree lastlbound = gfc_conv_array_lbound (desc, ndims - 1);
+ tree lastubound = gfc_conv_array_ubound (desc, ndims - 1);
+ tree lastrange = fold_build2_loc (input_location, MINUS_EXPR,
+ gfc_array_index_type, lastubound,
+ lastlbound);
+ lastrange = fold_build2_loc (input_location, PLUS_EXPR, gfc_array_index_type,
+ lastrange, gfc_index_one_node);
+
+ gfc_add_modify (block,
+ gfc_build_array_ref (ompdims, lastdim, NULL_TREE, true),
+ lastrange);
+
+ /* strides[0] = fstrides[0]; */
+ tree stride0 = gfc_conv_array_stride (desc, 0);
+ gfc_add_modify (block,
+ gfc_build_array_ref (ompstrides, gfc_index_zero_node,
+ NULL_TREE, true),
+ stride0);
+
+ if (ndims > 1)
+ /* strides[ndims - 1] = 1; */
+ gfc_add_modify (block,
+ gfc_build_array_ref (ompstrides, lastdim, NULL_TREE, true),
+ gfc_index_one_node);
+
+ if (ndims == 2)
+ /* dims[0] = fstrides[1]; */
+ gfc_add_modify (block,
+ gfc_build_array_ref (ompdims, gfc_index_zero_node,
+ NULL_TREE, true),
+ gfc_conv_array_stride (desc, 1));
+ else if (ndims > 2)
+ {
+ /* int grains[ndims - 2]; */
+ tree lastgrain = build_int_cst (gfc_array_index_type, ndims - 3);
+ tree grainrange = build_index_type (lastgrain);
+ tree grainarrtype = build_array_type (gfc_array_index_type, grainrange);
+ tree grains = gfc_create_var (grainarrtype, "grains");
+
+ /* int bigger_grain = fstrides[ndims - 1]; */
+ tree bigger_grain = gfc_create_var (gfc_array_index_type, "bigger_grain");
+ tree fstridem1 = gfc_conv_array_stride (desc, ndims - 1);
+ gfc_add_modify (block, bigger_grain, fstridem1);
+
+ /*
+ for (int i = ndims - 2; i > 0; i--)
+ {
+ grains[i - 1] = gcd (fstrides[i], bigger_grain);
+ bigger_grain = grains[i - 1];
+ }
+ */
+ stmtblock_t loop_body;
+ gfc_init_block (&loop_body);
+
+ tree idx = gfc_create_var (gfc_array_index_type, "idx");
+
+ tree gcdtmp = gfc_create_var (gfc_array_index_type, "tmp");
+ gfc_omp_calculate_gcd (&loop_body, gcdtmp,
+ gfc_conv_descriptor_stride_get (desc, idx),
+ bigger_grain);
+ tree idxm1 = fold_build2_loc (input_location, MINUS_EXPR,
+ gfc_array_index_type, idx,
+ gfc_index_one_node);
+ gfc_add_modify (&loop_body,
+ gfc_build_array_ref (grains, idxm1, NULL_TREE, true),
+ gcdtmp);
+ gfc_add_modify (&loop_body, bigger_grain, gcdtmp);
+
+ gfc_simple_for_loop (block, idx,
+ build_int_cst (gfc_array_index_type, ndims - 2),
+ gfc_index_zero_node, GT_EXPR,
+ build_int_cst (gfc_array_index_type, -1),
+ gfc_finish_block (&loop_body));
+ /*
+ int volume = 1;
+ for (int i = 0; i < ndims - 2; i++)
+ {
+ int g = grains[i];
+ dims[i] = g / volume;
+ strides[i + 1] = fstrides[i + 1] / g;
+ volume = volume * dims[i];
+ }
+ */
+ tree volume = gfc_create_var (gfc_array_index_type, "volume");
+ gfc_add_modify (block, volume, gfc_index_one_node);
+
+ gfc_init_block (&loop_body);
+ tree grain = gfc_create_var (gfc_array_index_type, "grain");
+ gfc_add_modify (&loop_body, grain,
+ gfc_build_array_ref (grains, idx, NULL_TREE, true));
+ tree dims_i = gfc_build_array_ref (ompdims, idx, NULL_TREE, true);
+ gfc_add_modify (&loop_body, dims_i,
+ fold_build2_loc (input_location, TRUNC_DIV_EXPR,
+ gfc_array_index_type, grain, volume));
+ tree nidx = fold_build2_loc (input_location, PLUS_EXPR,
+ gfc_array_index_type, idx,
+ gfc_index_one_node);
+ tree strides_ni = gfc_build_array_ref (ompstrides, nidx, NULL_TREE, true);
+ tree fstrides_ni = gfc_conv_descriptor_stride_get (desc, nidx);
+ gfc_add_modify (&loop_body, strides_ni,
+ fold_build2_loc (input_location, TRUNC_DIV_EXPR,
+ gfc_array_index_type, fstrides_ni,
+ grain));
+ gfc_add_modify (&loop_body, volume,
+ fold_build2_loc (input_location, MULT_EXPR,
+ gfc_array_index_type, volume, dims_i));
+
+ gfc_simple_for_loop (block, idx, gfc_index_zero_node,
+ build_int_cst (gfc_array_index_type, ndims - 2),
+ LT_EXPR, gfc_index_one_node,
+ gfc_finish_block (&loop_body));
+
+ /* dims[ndims - 2] = fstrides[ndims - 1] / volume; */
+ tree dimsm2
+ = gfc_build_array_ref (ompdims,
+ build_int_cst (gfc_array_index_type, ndims - 2),
+ NULL_TREE, true);
+ gfc_add_modify (block, dimsm2,
+ fold_build2_loc (input_location, TRUNC_DIV_EXPR,
+ gfc_array_index_type, fstridem1,
+ volume));
+ }
+}
+
+/* Return TRUE if update for N can definitely be done with a single contiguous
+ transfer. If no or if we can't tell, return FALSE. */
+
+static bool
+gfc_omp_contiguous_update_p (gfc_omp_namelist *n)
+{
+ gfc_expr *contig_expr = n->expr;
+
+ if (!n->expr)
+ {
+ if (n->sym->attr.contiguous)
+ return true;
+
+ tree desc = gfc_trans_omp_variable (n->sym, false);
+ tree type = TREE_TYPE (desc);
+ if (!GFC_ARRAY_TYPE_P (type) && !GFC_DESCRIPTOR_TYPE_P (type))
+ return true;
+
+ contig_expr = gfc_lval_expr_from_sym (n->sym);
+ }
+
+ return gfc_is_simply_contiguous (contig_expr, false, true);
+}
+
static tree
gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
locus where, toc_directive cd = TOC_OPENMP)
@@ -5830,6 +6174,162 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
default:
gcc_unreachable ();
}
+
+ if ((list == OMP_LIST_TO || list == OMP_LIST_FROM)
+ && (!n->expr
+ || (n->expr
+ && n->expr->ref
+ && n->expr->ref->type == REF_ARRAY))
+ && !gfc_omp_contiguous_update_p (n))
+ {
+ int ndims;
+ gfc_se se;
+ gfc_init_se (&se, NULL);
+
+ tree desc, span = NULL_TREE;
+
+ if (n->expr)
+ {
+ if (n->expr->rank)
+ gfc_conv_expr_descriptor (&se, n->expr);
+ else
+ gfc_conv_expr (&se, n->expr);
+
+ desc = se.expr;
+ /* The span is the distance between two array elements
+ along the innermost dimension (there may be padding
+ or other data between elements, e.g. of a derived-type
+ array). */
+ span = gfc_get_array_span (desc, n->expr);
+ ndims = n->expr->ref->u.ar.dimen;
+ }
+ else
+ {
+ desc = gfc_trans_omp_variable (n->sym, false);
+ tree type = TREE_TYPE (desc);
+ if (GFC_DESCRIPTOR_TYPE_P (type))
+ span = gfc_conv_descriptor_span_get (desc);
+ ndims = GFC_TYPE_ARRAY_RANK (type);
+ }
+
+ gfc_add_block_to_block (block, &se.pre);
+
+ tree ompdims, ompstrides;
+
+ gfc_desc_to_omp_noncontig_array (block, &ompdims,
+ &ompstrides, desc, ndims);
+
+ tree type = TREE_TYPE (desc);
+ tree etype = gfc_get_element_type (type);
+ tree elsize = fold_convert (gfc_array_index_type,
+ size_in_bytes (etype));
+
+ tree ptr = gfc_conv_array_data (desc);
+ tree offset = gfc_conv_array_offset (desc);
+
+ if (!span)
+ /* The span is the element size. */
+ span = elsize;
+
+ tree node = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+
+ switch (list)
+ {
+ case OMP_LIST_TO:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO_GRID);
+ break;
+ case OMP_LIST_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FROM_GRID);
+ break;
+ default:
+ gcc_unreachable ();
+ }
+
+ gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
+ tree byte_offset = fold_convert (sizetype, offset);
+ byte_offset = size_binop (MULT_EXPR, byte_offset,
+ fold_convert (sizetype, span));
+ tree origin
+ = fold_build2_loc (input_location, POINTER_PLUS_EXPR,
+ TREE_TYPE (ptr), ptr, byte_offset);
+
+ OMP_CLAUSE_SIZE (node) = elsize;
+
+ omp_clauses = gfc_trans_add_clause (node, omp_clauses);
+
+ auto_vec<tree, 5> dims;
+
+ for (int r = 0; r < ndims; r++)
+ {
+ tree d
+ = gfc_build_array_ref (ompdims,
+ build_int_cst
+ (gfc_array_index_type, r),
+ NULL_TREE, true);
+ d = gfc_evaluate_now (d, block);
+ dims.safe_push (d);
+ }
+
+ for (int r = ndims - 1; r >= 0; r--)
+ {
+ tree stride_r, len_r, lowbound_r;
+
+ tree rcst = build_int_cst (gfc_array_index_type, r);
+
+ stride_r = gfc_build_array_ref (ompstrides, rcst,
+ NULL_TREE, true);
+ lowbound_r = gfc_conv_array_lbound (desc, r);
+ len_r
+ = fold_build2_loc (input_location, MINUS_EXPR,
+ gfc_array_index_type,
+ gfc_conv_array_ubound (desc, r),
+ lowbound_r);
+ len_r
+ = fold_build2_loc (input_location, PLUS_EXPR,
+ gfc_array_index_type,
+ len_r, gfc_index_one_node);
+
+ lowbound_r
+ = fold_build2_loc (input_location, MULT_EXPR,
+ gfc_array_index_type, lowbound_r,
+ stride_r);
+
+ stride_r = gfc_evaluate_now (stride_r, block);
+ lowbound_r = gfc_evaluate_now (lowbound_r, block);
+ len_r = gfc_evaluate_now (len_r, block);
+
+ tree dim = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (dim, GOMP_MAP_GRID_DIM);
+ OMP_CLAUSE_DECL (dim) = lowbound_r;
+ OMP_CLAUSE_SIZE (dim) = len_r;
+
+ omp_clauses = gfc_trans_add_clause (dim, omp_clauses);
+
+ if (!integer_onep (stride_r)
+ || (r == 0 && !operand_equal_p (span, elsize)))
+ {
+ tree snode = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (snode,
+ GOMP_MAP_GRID_STRIDE);
+ OMP_CLAUSE_DECL (snode) = stride_r;
+ if (r == 0 && !operand_equal_p (span, elsize))
+ OMP_CLAUSE_SIZE (snode) = span;
+ omp_clauses = gfc_trans_add_clause (snode,
+ omp_clauses);
+ }
+ }
+ origin = build_fold_indirect_ref (origin);
+ tree eltype = gfc_get_element_type (TREE_TYPE (desc));
+ tree arrtype
+ = gfc_trans_omp_arrayshape_type (eltype, &dims);
+ OMP_CLAUSE_DECL (node)
+ = build1_loc (input_location, VIEW_CONVERT_EXPR,
+ arrtype, origin);
+ continue;
+ }
+
tree node = build_omp_clause (input_location, clause_code);
if (n->expr == NULL
|| (n->expr->ref->type == REF_ARRAY
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index bc781f99f5a..e04f9af0a24 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -15072,6 +15072,16 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT)
break;
+ /* If we have a non-contiguous (strided/rectangular) update
+ operation with a VIEW_CONVERT_EXPR, we need to be careful not
+ to gimplify the conversion away, because we need it during
+ omp-low.cc in order to retrieve the array's dimensions. Just
+ gimplify partially instead. */
+ if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_GRID
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM_GRID)
+ && TREE_CODE (*pd) == VIEW_CONVERT_EXPR)
+ pd = &TREE_OPERAND (*pd, 0);
+
/* We've already partly gimplified this in
gimplify_scan_omp_clauses. Don't do any more. */
if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c))
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 7165349fa84..69da3660e68 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1292,6 +1292,11 @@ omp_noncontig_descriptor_type (location_t loc)
TREE_CHAIN (field) = fields;
fields = field;
+ field = build_decl (loc, FIELD_DECL, get_identifier ("__span"),
+ size_type_node);
+ TREE_CHAIN (field) = fields;
+ fields = field;
+
tree ptr_size_type = build_pointer_type (size_type_node);
field = build_decl (loc, FIELD_DECL, get_identifier ("__dim"), ptr_size_type);
@@ -1956,7 +1961,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (dn, GOMP_MAP_TO_PSET);
OMP_CLAUSE_DECL (dn) = desc;
- OMP_CLAUSE_SIZE (dn) = TYPE_SIZE_UNIT (desc_type);
OMP_CLAUSE_CHAIN (dn) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = dn;
@@ -13375,6 +13379,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
&& OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_TO_PSET);
c = nc;
while ((nc = OMP_CLAUSE_CHAIN (c))
+ && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM
|| OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE))
c = nc;
@@ -13815,7 +13820,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
int i, dims = 0;
auto_vec<tree> tdims;
bool pointer_based = false, handled_pointer_section = false;
- tree arrsize = fold_convert (sizetype, elsize);
+ tree arrsize = size_one_node;
/* Allow a single (maybe strided) array section if we have a
pointer base. */
@@ -13827,8 +13832,12 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
dims = 1;
}
else
+ /* NOTE: Don't treat (e.g. Fortran, fixed-length) strings as
+ array types here; array section syntax isn't applicable to
+ strings. */
for (tree itype = type;
- TREE_CODE (itype) == ARRAY_TYPE;
+ TREE_CODE (itype) == ARRAY_TYPE
+ && !TYPE_STRING_FLAG (itype);
itype = TREE_TYPE (itype))
{
tdims.safe_push (itype);
@@ -13869,13 +13878,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
oc = c;
c = dn;
+ tree span = NULL_TREE;
+
for (i = 0; i < dims; i++)
{
nc = OMP_CLAUSE_CHAIN (c);
tree dim = NULL_TREE, index = NULL_TREE, len = NULL_TREE,
stride = size_one_node;
- if (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+ if (nc
+ && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM)
{
index = OMP_CLAUSE_DECL (nc);
@@ -13892,6 +13904,18 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
stride = OMP_CLAUSE_DECL (nc2);
stride = fold_convert (sizetype, stride);
+ if (OMP_CLAUSE_SIZE (nc2))
+ {
+ /* If the element size is not the same as the
+ distance between two adjacent array
+ elements (in the innermost dimension),
+ retrieve the latter value ("span") from the
+ size field of the stride. We only expect to
+ see one such field per array. */
+ gcc_assert (!span);
+ span = OMP_CLAUSE_SIZE (nc2);
+ span = fold_convert (sizetype, span);
+ }
nc = nc2;
}
@@ -13949,7 +13973,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
dim = size_binop (MINUS_EXPR, maxval, minval);
dim = size_binop (PLUS_EXPR, dim, size_one_node);
len = dim;
- index = size_zero_node;
+ index = minval;
+ nc = c;
}
if (TREE_CODE (dim) != INTEGER_CST)
@@ -13971,10 +13996,40 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
CONSTRUCTOR_APPEND_ELT (vstride, cidx, stride);
}
+ tree bias = size_zero_node;
+ tree volume = size_one_node;
+ for (i = dims - 1; i >= 0; i--)
+ {
+ tree dim = (*vdim)[i].value;
+ tree index = (*vindex)[i].value;
+ tree stride = (*vstride)[i].value;
+
+ /* For the bias we want, e.g.:
+
+ index[0] * stride[0] * dim[1] * dim[2]
+ + index[1] * stride[1] * dim[2]
+ + index[2] * stride[2]
+
+ All multiplied by "span" (or "elsize"). */
+
+ tree index_stride = size_binop (MULT_EXPR, index, stride);
+ bias = size_binop (PLUS_EXPR, bias,
+ size_binop (MULT_EXPR, volume,
+ index_stride));
+ volume = size_binop (MULT_EXPR, volume, dim);
+ }
+
+ /* If we don't have a separate span size, use the element size
+ instead. */
+ if (!span)
+ span = fold_convert (sizetype, elsize);
+
/* The size of the whole array -- to make sure we find any
part of the array via splay-tree lookup that might be
mapped on the target at runtime. */
- OMP_CLAUSE_SIZE (oc) = arrsize;
+ OMP_CLAUSE_SIZE (oc) = size_binop (MULT_EXPR, arrsize, span);
+ /* And the bias of the first element we will update. */
+ OMP_CLAUSE_SIZE (dn) = size_binop (MULT_EXPR, bias, span);
tree cdim = build_constructor (size_arr_type, vdim);
tree cindex = build_constructor (size_arr_type, vindex);
@@ -14005,13 +14060,14 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
tree ndims_field = TYPE_FIELDS (desc_type);
tree elemsize_field = DECL_CHAIN (ndims_field);
- tree dim_field = DECL_CHAIN (elemsize_field);
+ tree span_field = DECL_CHAIN (elemsize_field);
+ tree dim_field = DECL_CHAIN (span_field);
tree index_field = DECL_CHAIN (dim_field);
tree len_field = DECL_CHAIN (index_field);
tree stride_field = DECL_CHAIN (len_field);
vec<constructor_elt, va_gc> *v;
- vec_alloc (v, 6);
+ vec_alloc (v, 7);
bool all_static = (TREE_STATIC (dim_tmp)
&& TREE_STATIC (index_tmp)
@@ -14041,6 +14097,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
CONSTRUCTOR_APPEND_ELT (v, ndims_field, ndims);
CONSTRUCTOR_APPEND_ELT (v, elemsize_field, elsize);
+ CONSTRUCTOR_APPEND_ELT (v, span_field, span);
CONSTRUCTOR_APPEND_ELT (v, dim_field, dim_tmp);
CONSTRUCTOR_APPEND_ELT (v, index_field, index_tmp);
CONSTRUCTOR_APPEND_ELT (v, len_field, len_tmp);
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index bf369cc36b0..29500d60ae3 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,20 @@
+2023-09-05 Julian Brown <julian@codesourcery.com>
+
+ * 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.
+
+2023-07-03 Julian Brown <julian@codesourcery.com>
+
+ * gfortran.dg/gomp/noncontig-updates-1.f90: New test.
+ * gfortran.dg/gomp/noncontig-updates-2.f90: New test.
+ * gfortran.dg/gomp/noncontig-updates-3.f90: New test.
+ * gfortran.dg/gomp/noncontig-updates-4.f90: New test.
+
2023-07-03 Julian Brown <julian@codesourcery.com>
* g++.dg/gomp/array-shaping-1.C: New test.
diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c
new file mode 100644
index 00000000000..42d584fa624
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c
@@ -0,0 +1,26 @@
+// { dg-do compile }
+
+#include <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;
+
+#pragma omp target update to(([10][10]) arr[3:2][1:8][0:5])
+// { dg-error "too many array section specifiers for" "" { target *-*-* } .-1 }
+// { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 }
+
+#pragma omp target exit data map(from: arr[:100])
+
+ free (arr);
+
+ return 0;
+}
diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-2.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-2.c
new file mode 100644
index 00000000000..6be3e009ecb
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-2.c
@@ -0,0 +1,24 @@
+// { dg-do compile }
+
+#include <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-6.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-6.c
new file mode 100644
index 00000000000..d282d8598b2
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-6.c
@@ -0,0 +1,26 @@
+// { dg-do compile }
+
+int main (void)
+{
+ char *ptr;
+
+#pragma omp target update to(([5][6][7]) ptr[0:4][0:7][0:7])
+/* { dg-error {length '7' above array section size in 'to' clause} "" { target *-*-* } .-1 } */
+/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */
+
+#pragma omp target update to(([5][6][7]) ptr[1:5][0:6][0:7])
+/* { dg-error {high bound '6' above array section size in 'to' clause} "" { target *-*-* } .-1 } */
+/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */
+
+#pragma omp target update from(([100]) ptr[3:33:3])
+
+#pragma omp target update from(([100]) ptr[4:33:3])
+/* { dg-error {high bound '101' above array section size in 'from' clause} "" { target *-*-* } .-1 } */
+/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */
+
+#pragma omp target update to(([10][10]) ptr[0:9:-1][0:9])
+/* { dg-error {length '9' with stride '-1' above array section size in 'to' clause} "" { target *-*-* } .-1 } */
+/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */
+
+ return 0;
+}
diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-7.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-7.c
new file mode 100644
index 00000000000..233d8da6f44
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-7.c
@@ -0,0 +1,15 @@
+/* { dg-do compile } */
+
+int cond;
+
+int main (void)
+{
+ int *arr;
+
+ /* No array shaping inside conditional operator. */
+#pragma omp target update to(cond ? ([3][9]) arr : ([2][7]) arr)
+/* { dg-error {expected expression before '\[' token} "" { target *-*-* } .-1 } */
+/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */
+
+ return 0;
+}
diff --git a/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-1.f90 b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-1.f90
new file mode 100644
index 00000000000..5c60f5cac62
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-1.f90
@@ -0,0 +1,19 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+integer :: basicarray(100)
+integer, allocatable :: allocarray(:)
+
+allocate(allocarray(1:20))
+
+!$omp target update to(basicarray)
+
+!$omp target update from(basicarray(:))
+
+!$omp target update to(allocarray)
+
+!$omp target update from(allocarray(:))
+
+end
+
+! { dg-final { scan-tree-dump-times {omp target update from\(} 2 "original" } }
+! { dg-final { scan-tree-dump-times {omp target update to\(} 2 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-2.f90 b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-2.f90
new file mode 100644
index 00000000000..f5a52736b0c
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-2.f90
@@ -0,0 +1,16 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+integer, allocatable :: allocarray(:)
+integer, allocatable :: allocarray2(:,:)
+
+allocate(allocarray(1:20))
+allocate(allocarray2(1:20,1:20))
+
+! This one must be noncontiguous
+!$omp target update to(allocarray(::2))
+! { dg-final { scan-tree-dump {omp target update map\(to_grid:} "original" } }
+
+!$omp target update from(allocarray2(:,5:15))
+! { dg-final { scan-tree-dump {omp target update from\(} "original" } }
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-3.f90 b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-3.f90
new file mode 100644
index 00000000000..5cbfe7c7be5
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-3.f90
@@ -0,0 +1,16 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+integer, allocatable :: allocarray(:,:)
+
+allocate(allocarray(1:20,1:20))
+
+! This one could possibly be handled as a contiguous update - but isn't,
+! for now.
+!$omp target update to(allocarray(1:20,5:15))
+! { dg-final { scan-tree-dump {omp target update map\(to_grid:} "original" } }
+
+!$omp target update from(allocarray(:,5:15:2))
+! { dg-final { scan-tree-dump {omp target update map\(from_grid:} "original" } }
+
+end
+
diff --git a/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-4.f90 b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-4.f90
new file mode 100644
index 00000000000..53152aacbb4
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-4.f90
@@ -0,0 +1,15 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+integer, target :: tgtarray(20)
+integer, pointer, contiguous :: arrayptr(:)
+
+arrayptr => tgtarray
+
+!$omp target update from(arrayptr)
+! { dg-final { scan-tree-dump {omp target update from\(} "original" } }
+
+!$omp target update to(arrayptr(::2))
+! { dg-final { scan-tree-dump {omp target update map\(to_grid:} "original" } }
+
+end
+
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 286707575c7..e00cb42ef53 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,33 @@
+2023-09-05 Julian Brown <julian@codesourcery.com>
+
+ * 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.
+
+2023-07-03 Julian Brown <julian@codesourcery.com>
+
+ * libgomp.h (omp_noncontig_array_desc): Add span field.
+ * target.c (omp_target_memcpy_rect_worker): Add span parameter. Update
+ forward declaration. Handle span != element_size.
+ (gomp_update): Handle bias in descriptor's size slot. Update calls to
+ omp_target_memcpy_rect_worker.
+ * testsuite/libgomp.fortran/noncontig-updates-1.f90: New test.
+ * testsuite/libgomp.fortran/noncontig-updates-2.f90: New test.
+ * testsuite/libgomp.fortran/noncontig-updates-3.f90: New test.
+ * testsuite/libgomp.fortran/noncontig-updates-4.f90: New test.
+ * testsuite/libgomp.fortran/noncontig-updates-5.f90: New test.
+ * testsuite/libgomp.fortran/noncontig-updates-6.f90: New test.
+ * testsuite/libgomp.fortran/noncontig-updates-7.f90: New test.
+ * testsuite/libgomp.fortran/noncontig-updates-8.f90: New test.
+ * testsuite/libgomp.fortran/noncontig-updates-9.f90: New test.
+ * testsuite/libgomp.fortran/noncontig-updates-10.f90: New test.
+ * testsuite/libgomp.fortran/noncontig-updates-11.f90: New test.
+ * testsuite/libgomp.fortran/noncontig-updates-12.f90: New test.
+ * testsuite/libgomp.fortran/noncontig-updates-13.f90: New test.
+
2023-07-03 Julian Brown <julian@codesourcery.com>
* libgomp.h (omp_noncontig_array_desc): New struct.
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 8caf8a0ee5d..ea7e06b6d7b 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1333,6 +1333,7 @@ struct target_mem_desc {
typedef struct {
size_t ndims;
size_t elemsize;
+ size_t span;
size_t *dim;
size_t *index;
size_t *length;
diff --git a/libgomp/target.c b/libgomp/target.c
index 17de0e82fd2..a75cd751b39 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2312,7 +2312,7 @@ goacc_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom,
}
static int
-omp_target_memcpy_rect_worker (void *, const void *, size_t, int,
+omp_target_memcpy_rect_worker (void *, const void *, size_t, size_t, int,
const size_t *, const size_t *, const size_t *,
const size_t *, const size_t *, const size_t *,
struct gomp_device_descr *,
@@ -2348,9 +2348,9 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
{
omp_noncontig_array_desc *desc
= (omp_noncontig_array_desc *) hostaddrs[i + 1];
- cur_node.host_start = (uintptr_t) hostaddrs[i];
+ size_t bias = sizes[i + 1];
+ cur_node.host_start = (uintptr_t) hostaddrs[i] + bias;
cur_node.host_end = cur_node.host_start + sizes[i];
- assert (sizes[i + 1] == sizeof (omp_noncontig_array_desc));
splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
if (n)
{
@@ -2362,22 +2362,25 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
}
void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset
+ cur_node.host_start
- - n->host_start);
+ - n->host_start
+ - bias);
size_t tmp_size = 0;
void *tmp = NULL;
if ((kind & typemask) == GOMP_MAP_TO_GRID)
omp_target_memcpy_rect_worker (devaddr, hostaddrs[i],
- desc->elemsize, desc->ndims,
- desc->length, desc->stride,
- desc->index, desc->index,
- desc->dim, desc->dim, devicep,
+ desc->elemsize, desc->span,
+ desc->ndims, desc->length,
+ desc->stride, desc->index,
+ desc->index, desc->dim,
+ desc->dim, devicep,
NULL, &tmp_size, &tmp);
else
omp_target_memcpy_rect_worker (hostaddrs[i], devaddr,
- desc->elemsize, desc->ndims,
- desc->length, desc->stride,
- desc->index, desc->index,
- desc->dim, desc->dim, NULL,
+ desc->elemsize, desc->span,
+ desc->ndims, desc->length,
+ desc->stride, desc->index,
+ desc->index, desc->dim,
+ desc->dim, NULL,
devicep, &tmp_size, &tmp);
}
i++;
@@ -4965,7 +4968,7 @@ omp_target_memcpy_async (void *dst, const void *src, size_t length,
static int
omp_target_memcpy_rect_worker (void *dst, const void *src, size_t element_size,
- int num_dims, const size_t *volume,
+ size_t span, int num_dims, const size_t *volume,
const size_t *strides,
const size_t *dst_offsets,
const size_t *src_offsets,
@@ -4980,7 +4983,7 @@ omp_target_memcpy_rect_worker (void *dst, const void *src, size_t element_size,
size_t j, dst_off, src_off, length;
int i, ret;
- if (num_dims == 1 && (!strides || strides[0] == 1))
+ if (num_dims == 1 && (!strides || (strides[0] == 1 && element_size == span)))
{
if (__builtin_mul_overflow (element_size, volume[0], &length)
|| __builtin_mul_overflow (element_size, dst_offsets[0], &dst_off)
@@ -5041,12 +5044,11 @@ omp_target_memcpy_rect_worker (void *dst, const void *src, size_t element_size,
assert ((src_devicep == NULL || dst_devicep == NULL)
&& (src_devicep != NULL || dst_devicep != NULL));
- if (__builtin_mul_overflow (element_size, dst_offsets[0], &dst_off)
- || __builtin_mul_overflow (element_size, src_offsets[0], &src_off))
+ if (__builtin_mul_overflow (span, dst_offsets[0], &dst_off)
+ || __builtin_mul_overflow (span, src_offsets[0], &src_off))
return EINVAL;
- if (strides
- && __builtin_mul_overflow (element_size, strides[0], &stride))
+ if (__builtin_mul_overflow (span, strides[0], &stride))
return EINVAL;
for (i = 0, ret = 1; i < volume[0] && ret; i++)
@@ -5138,7 +5140,7 @@ omp_target_memcpy_rect_worker (void *dst, const void *src, size_t element_size,
{
ret = omp_target_memcpy_rect_worker ((char *) dst + dst_off,
(const char *) src + src_off,
- element_size, num_dims - 1,
+ element_size, span, num_dims - 1,
volume + 1,
strides ? strides + 1 : NULL,
dst_offsets + 1, src_offsets + 1,
@@ -5191,7 +5193,7 @@ omp_target_memcpy_rect_copy (void *dst, const void *src,
gomp_mutex_lock (&src_devicep->lock);
if (lock_dst)
gomp_mutex_lock (&dst_devicep->lock);
- int ret = omp_target_memcpy_rect_worker (dst, src, element_size, num_dims,
+ int ret = omp_target_memcpy_rect_worker (dst, src, element_size, element_size, num_dims,
volume, NULL, dst_offsets, src_offsets,
dst_dimensions, src_dimensions,
dst_devicep, src_devicep,
diff --git a/libgomp/testsuite/libgomp.c/array-shaping-1.c b/libgomp/testsuite/libgomp.c/array-shaping-1.c
new file mode 100644
index 00000000000..808c5f9ceae
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/array-shaping-1.c
@@ -0,0 +1,236 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <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 = calloc (100, sizeof (int));
+
+ /* Update whole reshaped array. */
+
+#pragma omp target enter data map(to: arr[:100])
+
+ for (int j = 0; j < x; j++)
+ for (int i = 0; i < 10; i++)
+ arr[j * 10 + i] = i ^ j;
+
+#pragma omp target update to(([10][x]) arr)
+
+#pragma omp target exit data map(from: arr[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ if (j < x)
+ assert (arr[j * 10 + i] == i ^ j);
+ else
+ assert (arr[j * 10 + i] == 0);
+
+
+ /* Strided update. */
+
+ memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+ for (int j = 0; j < 20; j++)
+ for (int i = 0; i < 5; i++)
+ arr[j * 5 + i] = i + j;
+
+#pragma omp target update to(([5][5]) arr[0:3][0:3:2])
+
+#pragma omp target exit data map(from: arr[:100])
+
+ for (int j = 0; j < 20; j++)
+ for (int i = 0; i < 5; i++)
+ if (j < 3 && (i & 1) == 0 && i < 6)
+ assert (arr[j * 5 + i] == i + j);
+ else
+ assert (arr[j * 5 + i] == 0);
+
+
+ /* Reshaped update, contiguous. */
+
+ memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+ for (int j = 0; j < 20; j++)
+ for (int i = 0; i < 5; i++)
+ arr[j * 5 + i] = 2 * j + i;
+
+#pragma omp target update to(([5][5]) arr[0:5][0:5])
+
+#pragma omp target exit data map(from: arr[:100])
+
+ for (int j = 0; j < 20; j++)
+ for (int i = 0; i < 5; i++)
+ if (j < 5 && i < 5)
+ assert (arr[j * 5 + i] == 2 * j + i);
+ else
+ assert (arr[j * 5 + i] == 0);
+
+
+ /* Strided update on actual array. */
+
+ memset (arr2d, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr2d)
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ arr2d[j][i] = j + 2 * i;
+
+#pragma omp target update to(arr2d[0:5:2][5:2])
+
+#pragma omp target exit data map(from: arr2d)
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ if ((j & 1) == 0 && i >= 5 && i < 7)
+ assert (arr2d[j][i] == j + 2 * i);
+ else
+ assert (arr2d[j][i] == 0);
+
+
+ /* Update with non-constant bounds. */
+
+ memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ arr[j * 10 + i] = (2 * j) ^ i;
+
+ x = 3;
+ int y = yy, z = zz, str = str_str;
+ /* This is actually [0:3:2] [4:2:2]. */
+#pragma omp target update to(([10][10]) arr[0:x:2][y:z:str])
+
+#pragma omp target exit data map(from: arr[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ if ((j & 1) == 0 && j < 6 && (i & 1) == 0 && i >= 4 && i < 8)
+ assert (arr[j * 10 + i] == (2 * j) ^ i);
+ else
+ assert (arr[j * 10 + i] == 0);
+
+
+ /* Update with full "major" dimension. */
+
+ memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ arr[j * 10 + i] = i + j;
+
+#pragma omp target update to(([10][10]) arr[0:10][3:1])
+
+#pragma omp target exit data map(from: arr[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ if (i == 3)
+ assert (arr[j * 10 + i] == i + j);
+ else
+ assert (arr[j * 10 + i] == 0);
+
+
+ /* Update with full "minor" dimension. */
+
+ memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ arr[j * 10 + i] = 3 * (i + j);
+
+#pragma omp target update to(([10][10]) arr[3:2][0:10])
+
+#pragma omp target exit data map(from: arr[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ if (j >= 3 && j < 5)
+ assert (arr[j * 10 + i] == 3 * (i + j));
+ else
+ assert (arr[j * 10 + i] == 0);
+
+
+ /* Rectangle update. */
+
+ memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ arr[j * 10 + i] = 5 * (i + j);
+
+#pragma omp target update to(([10][10]) arr[3:2][0:9])
+
+#pragma omp target exit data map(from: arr[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ if (j >= 3 && j < 5 && i < 9)
+ assert (arr[j * 10 + i] == 5 * (i + j));
+ else
+ assert (arr[j * 10 + i] == 0);
+
+
+ /* One-dimensional strided update. */
+
+ memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+ for (int i = 0; i < 100; i++)
+ arr[i] = i + 99;
+
+#pragma omp target update to(([100]) arr[3:33:3])
+
+#pragma omp target exit data map(from: arr[:100])
+
+ for (int i = 0; i < 100; i++)
+ if (i >= 3 && ((i - 3) % 3) == 0)
+ assert (arr[i] == i + 99);
+ else
+ assert (arr[i] == 0);
+
+
+ /* One-dimensional strided update without explicit array shape. */
+
+ memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+ for (int i = 0; i < 100; i++)
+ arr[i] = i + 121;
+
+#pragma omp target update to(arr[3:33:3])
+
+#pragma omp target exit data map(from: arr[:100])
+
+ for (int i = 0; i < 100; i++)
+ if (i >= 3 && ((i - 3) % 3) == 0)
+ assert (arr[i] == i + 121);
+ else
+ assert (arr[i] == 0);
+
+ free (arr);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/array-shaping-2.c b/libgomp/testsuite/libgomp.c/array-shaping-2.c
new file mode 100644
index 00000000000..42a6e0ca7d8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/array-shaping-2.c
@@ -0,0 +1,39 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <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;
+}
diff --git a/libgomp/testsuite/libgomp.c/array-shaping-4.c b/libgomp/testsuite/libgomp.c/array-shaping-4.c
new file mode 100644
index 00000000000..2b9e6949b60
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/array-shaping-4.c
@@ -0,0 +1,36 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <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/array-shaping-5.c b/libgomp/testsuite/libgomp.c/array-shaping-5.c
new file mode 100644
index 00000000000..1034682e4ca
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/array-shaping-5.c
@@ -0,0 +1,38 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <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/array-shaping-6.c b/libgomp/testsuite/libgomp.c/array-shaping-6.c
new file mode 100644
index 00000000000..59388232244
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/array-shaping-6.c
@@ -0,0 +1,45 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <stdlib.h>
+#include <string.h>
+
+#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;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-1.f90 b/libgomp/testsuite/libgomp.fortran/noncontig-updates-1.f90
new file mode 100644
index 00000000000..6ee87e8043b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-1.f90
@@ -0,0 +1,54 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+implicit none
+integer, allocatable, target :: arr(:), arr2(:,:)
+integer, pointer :: ap(:), ap2(:,:)
+integer :: i, j
+
+allocate(arr(1:20))
+
+arr = 0
+
+!$omp target enter data map(to: arr)
+
+ap => arr(1:20:2)
+ap = 5
+
+!$omp target update to(ap)
+
+!$omp target exit data map(from: arr)
+
+do i=1,20
+ if (mod(i,2).eq.1.and.arr(i).ne.5) stop 1
+ if (mod(i,2).eq.0.and.arr(i).ne.0) stop 2
+end do
+
+allocate(arr2(1:20,1:20))
+
+ap2 => arr2(2:10:2,3:12:3)
+
+arr2 = 1
+
+!$omp target enter data map(to: arr2)
+
+!$omp target
+ap2 = 5
+!$omp end target
+
+!$omp target update from(ap2)
+
+do i=1,20
+ do j=1,20
+ if (i.ge.2.and.i.le.10.and.mod(i-2,2).eq.0.and.&
+ &j.ge.3.and.j.le.12.and.mod(j-3,3).eq.0) then
+ if (arr2(i,j).ne.5) stop 3
+ else
+ if (arr2(i,j).ne.1) stop 4
+ end if
+ end do
+end do
+
+!$omp target exit data map(delete: arr2)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-10.f90 b/libgomp/testsuite/libgomp.fortran/noncontig-updates-10.f90
new file mode 100644
index 00000000000..c47ce38918d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-10.f90
@@ -0,0 +1,29 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+character(len=8), allocatable, dimension(:) :: lines
+integer :: i
+
+allocate(lines(10))
+
+lines = "OMPHELLO"
+
+!$omp target enter data map(to: lines)
+
+!$omp target
+lines = "NEWVALUE"
+!$omp end target
+
+!$omp target update from(lines(5:7:2))
+
+do i=1,10
+ if (i.eq.5.or.i.eq.7) then
+ if (lines(i).ne."NEWVALUE") stop 1
+ else
+ if (lines(i).ne."OMPHELLO") stop 2
+ end if
+end do
+
+!$omp target exit data map(delete: lines)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-11.f90 b/libgomp/testsuite/libgomp.fortran/noncontig-updates-11.f90
new file mode 100644
index 00000000000..a93acf21d77
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-11.f90
@@ -0,0 +1,51 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+program p
+implicit none
+real(kind=4) :: arr(10,10,10,10)
+
+call s(arr,9,9,9,9)
+
+contains
+
+subroutine s(arr,m,n,o,p)
+implicit none
+integer :: i,m,n,o,p
+integer :: a,b,c,d
+real(kind=4) :: arr(0:m,0:n,0:o,0:p)
+
+arr = 0
+
+!$omp target enter data map(to: arr)
+
+!$omp target
+do i=0,9
+ arr(i,i,i,i) = i
+end do
+!$omp end target
+
+!$omp target update from(arr(0:2,0:2,0:2,0:2))
+
+do a=0,9
+ do b=0,9
+ do c=0,9
+ do d=0,9
+ if (a.le.2.and.b.le.2.and.c.le.2.and.d.le.2) then
+ if (a.eq.b.and.b.eq.c.and.c.eq.d) then
+ if (arr(a,b,c,d).ne.a) stop 1
+ else
+ if (arr(a,b,c,d).ne.0) stop 2
+ end if
+ else
+ if (arr(a,b,c,d).ne.0) stop 3
+ end if
+ end do
+ end do
+ end do
+end do
+
+!$omp target exit data map(delete: arr)
+
+end subroutine s
+end program p
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-12.f90 b/libgomp/testsuite/libgomp.fortran/noncontig-updates-12.f90
new file mode 100644
index 00000000000..c47fbdb0d11
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-12.f90
@@ -0,0 +1,59 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! Test plain, fixed-size arrays, and also pointers to same.
+
+implicit none
+integer(kind=8) :: arr(10,30)
+integer, target :: arr2(9,11,13)
+integer, pointer :: parr(:,:,:)
+integer :: i, j, k
+
+arr = 0
+!$omp target enter data map(to: arr)
+
+!$omp target
+arr = 99
+!$omp end target
+
+!$omp target update from(arr(1:10:3,5:30:7))
+
+do i=1,10
+ do j=1,30
+ if (mod(i-1,3).eq.0.and.mod(j-5,7).eq.0) then
+ if (arr(i,j).ne.99) stop 1
+ else
+ if (arr(i,j).ne.0) stop 2
+ endif
+ end do
+end do
+
+!$omp target exit data map(delete: arr)
+
+arr2 = 0
+parr => arr2
+!$omp target enter data map(to: parr)
+
+!$omp target
+parr = 99
+!$omp end target
+
+!$omp target update from(parr(7:9:2,5:7:2,3:6:3))
+
+do i=1,9
+ do j=1,11
+ do k=1,13
+ if (i.ge.7.and.j.ge.5.and.k.ge.3.and.&
+ &i.le.9.and.j.le.7.and.k.le.6.and.&
+ &mod(i-7,2).eq.0.and.mod(j-5,2).eq.0.and.mod(k-3,3).eq.0) then
+ if (parr(i,j,k).ne.99) stop 3
+ else
+ if (parr(i,j,k).ne.0) stop 4
+ end if
+ end do
+ end do
+end do
+
+!$omp target exit data map(delete: parr)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-13.f90 b/libgomp/testsuite/libgomp.fortran/noncontig-updates-13.f90
new file mode 100644
index 00000000000..42f867efefc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-13.f90
@@ -0,0 +1,42 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+implicit none
+integer, allocatable :: arr(:,:,:,:,:)
+integer :: i, j, k, l, m
+
+allocate (arr(18,19,20,21,22))
+
+arr = 0
+
+!$omp target enter data map(to: arr)
+
+arr = 10
+
+!$omp target update to(arr(1:3:2,1:4:3,1:5:4,1:6:5,1:7:6))
+
+!$omp target
+do i=1,18
+ do j=1,19
+ do k=1,20
+ do l=1,21
+ do m=1,22
+ if ((i.eq.1.or.i.eq.3).and.&
+ &(j.eq.1.or.j.eq.4).and.&
+ &(k.eq.1.or.k.eq.5).and.&
+ &(l.eq.1.or.l.eq.6).and.&
+ &(m.eq.1.or.m.eq.7)) then
+ if (arr(i,j,k,l,m).ne.10) stop 1
+ else
+ if (arr(i,j,k,l,m).ne.0) stop 2
+ end if
+ end do
+ end do
+ end do
+ end do
+end do
+!$omp end target
+
+!$omp target exit data map(delete: arr)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-2.f90 b/libgomp/testsuite/libgomp.fortran/noncontig-updates-2.f90
new file mode 100644
index 00000000000..2d3efb8bfcc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-2.f90
@@ -0,0 +1,101 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+program p
+implicit none
+integer, allocatable, target :: arr3(:,:,:)
+integer, pointer :: ap3(:,:,:)
+integer :: i, j, k
+
+allocate(arr3(1:10,1:10,1:10))
+
+! CHECK 1
+
+arr3 = 0
+ap3 => arr3(1:10,1:10,1:10:2)
+
+!$omp target enter data map(to: arr3)
+
+!$omp target
+ap3 = 5
+!$omp end target
+
+!$omp target update from(ap3)
+
+call check(arr3, 0, 1, 1, 2)
+
+!$omp target exit data map(delete: arr3)
+
+! CHECK 2
+
+arr3 = 0
+ap3 => arr3(1:10,1:10:2,1:10)
+
+!$omp target enter data map(to: arr3)
+
+!$omp target
+ap3 = 5
+!$omp end target
+
+!$omp target update from(ap3)
+
+call check(arr3, 2, 1, 2, 1)
+
+!$omp target exit data map(delete: arr3)
+
+! CHECK 3
+
+arr3 = 0
+ap3 => arr3(1:10:2,1:10,1:10)
+
+!$omp target enter data map(to: arr3)
+
+!$omp target
+ap3 = 5
+!$omp end target
+
+!$omp target update from(ap3)
+
+call check(arr3, 4, 2, 1, 1)
+
+!$omp target exit data map(delete: arr3)
+
+! CHECK 4
+
+arr3 = 0
+ap3 => arr3(1:10:2,1:10:2,1:10:2)
+
+!$omp target enter data map(to: arr3)
+
+!$omp target
+ap3 = 5
+!$omp end target
+
+!$omp target update from(ap3)
+
+call check(arr3, 6, 2, 2, 2)
+
+!$omp target exit data map(delete: arr3)
+
+contains
+
+subroutine check(arr,cb,s1,s2,s3)
+implicit none
+integer :: arr(:,:,:)
+integer :: cb, s1, s2, s3
+
+do i=1,10
+ do j=1,10
+ do k=1,10
+ if (mod(k-1,s1).eq.0.and.mod(j-1,s2).eq.0.and.mod(i-1,s3).eq.0) then
+ if (arr(k,j,i).ne.5) stop cb+1
+ else
+ if (arr(k,j,i).ne.0) stop cb+2
+ end if
+ end do
+ end do
+end do
+
+end subroutine check
+
+end program p
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-3.f90 b/libgomp/testsuite/libgomp.fortran/noncontig-updates-3.f90
new file mode 100644
index 00000000000..14f1288a697
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-3.f90
@@ -0,0 +1,47 @@
+program p
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+integer :: A(200)
+A = [(i, i=1,200)]
+!$omp target enter data map(to: A(40:200))
+call foo(A(101:))
+
+contains
+
+subroutine foo(x)
+integer, target :: x(100)
+integer, pointer :: p(:,:)
+integer :: i, j
+
+p(0:5,-5:-1) => x(::2)
+
+!$omp target
+x = x * 2
+!$omp end target
+
+!$omp target update from(x(1:20:2))
+
+do i=1,20
+if (mod(i,2).eq.1 .and. x(i).ne.(100+i)*2) stop 1
+if (mod(i,2).eq.0 .and. x(i).ne.100+i) stop 2
+end do
+
+!$omp target
+p = 0
+!$omp end target
+
+!$omp target update from(p(::3,::2))
+
+do i=0,5
+ do j=-5,-1
+ if (mod(i,3).eq.0 .and. mod(j+5,2).eq.0) then
+ if (p(i,j).ne.0) stop 3
+ else
+ if (p(i,j).eq.0) stop 4
+ end if
+ end do
+end do
+
+end subroutine foo
+end program p
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-4.f90 b/libgomp/testsuite/libgomp.fortran/noncontig-updates-4.f90
new file mode 100644
index 00000000000..46e8c23d285
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-4.f90
@@ -0,0 +1,78 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+type t
+ complex(kind=8) :: c
+ integer :: i
+end type t
+
+type u
+ integer :: i, j
+ complex(kind=8) :: c
+ integer :: k
+end type u
+
+type(t), target :: var(10)
+type(u), target :: var2(10)
+complex(kind=8), pointer :: ptr(:)
+integer :: i
+
+do i=1,10
+ var(i)%c = dcmplx(i,0)
+ var(i)%i = i
+end do
+
+ptr => var(:)%c
+
+!$omp target enter data map(to: var)
+
+!$omp target
+var(:)%c = dcmplx(0,0)
+var(:)%i = 0
+!$omp end target
+
+!$omp target update from(ptr)
+
+do i=1,10
+ if (var(i)%c.ne.dcmplx(0,0)) stop 1
+ if (var(i)%i.ne.i) stop 2
+end do
+
+!$omp target exit data map(delete: var)
+
+! Now do it again with a differently-ordered derived type.
+
+do i=1,10
+ var2(i)%c = dcmplx(0,i)
+ var2(i)%i = i
+ var2(i)%j = i * 2
+ var2(i)%k = i * 3
+end do
+
+ptr => var2(::2)%c
+
+!$omp target enter data map(to: var2)
+
+!$omp target
+var2(:)%c = dcmplx(0,0)
+var2(:)%i = 0
+var2(:)%j = 0
+var2(:)%k = 0
+!$omp end target
+
+!$omp target update from(ptr)
+
+do i=1,10
+ if (mod(i,2).eq.1) then
+ if (var2(i)%c.ne.dcmplx(0,0)) stop 3
+ else
+ if (var2(i)%c.ne.dcmplx(0,i)) stop 4
+ end if
+ if (var2(i)%i.ne.i) stop 5
+ if (var2(i)%j.ne.i * 2) stop 6
+ if (var2(i)%k.ne.i * 3) stop 7
+end do
+
+!$omp target exit data map(delete: var2)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-5.f90 b/libgomp/testsuite/libgomp.fortran/noncontig-updates-5.f90
new file mode 100644
index 00000000000..9cc20fa321e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-5.f90
@@ -0,0 +1,55 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! Only some of an array mapped on the target
+
+integer, target :: arr(100)
+integer, pointer :: ptr(:)
+
+arr = [(i * 2, i=1,100)]
+
+!$omp target enter data map(to: arr(51:100))
+
+!$omp target
+arr(51:100) = arr(51:100) + 1
+!$omp end target
+
+!$omp target update from(arr(51:100:2))
+
+do i=1,100
+ if (i.le.50) then
+ if (arr(i).ne.i*2) stop 1
+ else
+ if (mod(i,2).eq.1 .and. arr(i).ne.i*2+1) stop 2
+ if (mod(i,2).eq.0 .and. arr(i).ne.i*2) stop 3
+ end if
+end do
+
+!$omp target exit data map(delete: arr)
+
+arr = [(i * 2, i=1,100)]
+
+! Similar, but update via pointer.
+
+ptr => arr(51:100)
+
+!$omp target enter data map(to: ptr(1:50))
+
+!$omp target
+ptr = ptr + 1
+!$omp end target
+
+!$omp target update from(ptr(::2))
+
+do i=1,100
+ if (i.le.50) then
+ if (arr(i).ne.i*2) stop 1
+ else
+ if (mod(i,2).eq.1 .and. arr(i).ne.i*2+1) stop 2
+ if (mod(i,2).eq.0 .and. arr(i).ne.i*2) stop 3
+ end if
+end do
+
+!$omp target exit data map(delete: ptr)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-6.f90 b/libgomp/testsuite/libgomp.fortran/noncontig-updates-6.f90
new file mode 100644
index 00000000000..5c42b9077b3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-6.f90
@@ -0,0 +1,34 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+program p
+implicit none
+integer, dimension(100) :: parr
+integer :: i
+
+parr = [(i, i=1,100)]
+
+!$omp target enter data map(to: parr)
+
+call s(parr)
+
+do i=1,100
+ if (mod(i,3).eq.1 .and. parr(i).ne.999) stop 1
+ if (mod(i,3).ne.1 .and. parr(i).ne.i) stop 2
+end do
+
+!$omp target exit data map(delete: parr)
+
+contains
+subroutine s(arr)
+implicit none
+integer, intent(inout) :: arr(*)
+
+!$omp target map(alloc: arr(1:100))
+arr(1:100) = 999
+!$omp end target
+
+!$omp target update from(arr(1:100:3))
+
+end subroutine s
+end program p
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-7.f90 b/libgomp/testsuite/libgomp.fortran/noncontig-updates-7.f90
new file mode 100644
index 00000000000..120fd9c90ed
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-7.f90
@@ -0,0 +1,36 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! Assumed-shape arrays
+
+program p
+implicit none
+integer, dimension(100) :: parr
+integer :: i
+
+parr = [(i, i=1,100)]
+
+!$omp target enter data map(to: parr)
+
+call s(parr)
+
+do i=1,100
+ if (mod(i,3).eq.1 .and. parr(i).ne.999) stop 1
+ if (mod(i,3).ne.1 .and. parr(i).ne.i) stop 2
+end do
+
+!$omp target exit data map(delete: parr)
+
+contains
+subroutine s(arr)
+implicit none
+integer, intent(inout) :: arr(:)
+
+!$omp target
+arr = 999
+!$omp end target
+
+!$omp target update from(arr(1:100:3))
+
+end subroutine s
+end program p
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-8.f90 b/libgomp/testsuite/libgomp.fortran/noncontig-updates-8.f90
new file mode 100644
index 00000000000..d9b3c9ca896
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-8.f90
@@ -0,0 +1,39 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! Test biasing for target-region lookup.
+
+implicit none
+integer, allocatable, target :: var(:,:,:)
+integer, pointer :: p(:,:,:)
+integer :: i, j, k
+
+allocate(var(1:20,5:25,10:30))
+
+var = 0
+
+!$omp target enter data map(to: var)
+
+!$omp target
+var = 99
+!$omp end target
+
+p => var(1:3:2,5:5,10:10)
+
+!$omp target update from(p)
+
+do i=1,20
+ do j=5,25
+ do k=10,30
+ if ((i.eq.1.or.i.eq.3).and.j.eq.5.and.k.eq.10) then
+ if (var(i,j,k).ne.99) stop 1
+ else
+ if (var(i,j,k).ne.0) stop 2
+ end if
+ end do
+ end do
+end do
+
+!$omp target exit data map(delete: var)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-9.f90 b/libgomp/testsuite/libgomp.fortran/noncontig-updates-9.f90
new file mode 100644
index 00000000000..689a46a91f0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-9.f90
@@ -0,0 +1,34 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! This test case hits the problem described in:
+! https://gcc.gnu.org/pipermail/gcc-patches/2023-February/612219.html
+
+! { dg-xfail-run-if "'enter data' bug" { offload_device_nonshared_as } }
+
+character(len=:), allocatable, dimension(:) :: lines
+integer :: i
+
+allocate(character(len=8) :: lines(10))
+
+lines = "OMPHELLO"
+
+!$omp target enter data map(to: lines)
+
+!$omp target
+lines = "NEWVALUE"
+!$omp end target
+
+!$omp target update from(lines(5:7:2))
+
+do i=1,10
+ if (i.eq.5.or.i.eq.7) then
+ if (lines(i).ne."NEWVALUE") stop 1
+ else
+ if (lines(i).ne."OMPHELLO") stop 2
+ end if
+end do
+
+!$omp target exit data map(delete: lines)
+
+end
More information about the Gcc-cvs
mailing list