@@ -5764,7 +5764,9 @@ c_parser_braced_init (c_parser *parser, tree type, bool nested_p,
gcc_obstack_init (&braced_init_obstack);
gcc_assert (c_parser_next_token_is (parser, CPP_OPEN_BRACE));
bool save_c_omp_array_section_p = c_omp_array_section_p;
+ bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
c_omp_array_section_p = false;
+ c_omp_array_shaping_op_p = false;
matching_braces braces;
braces.consume_open (parser);
if (nested_p)
@@ -5804,6 +5806,7 @@ c_parser_braced_init (c_parser *parser, tree type, bool nested_p,
}
}
c_omp_array_section_p = save_c_omp_array_section_p;
+ c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
c_token *next_tok = c_parser_peek_token (parser);
if (next_tok->type != CPP_CLOSE_BRACE)
{
@@ -8310,6 +8313,7 @@ c_parser_conditional_expression (c_parser *parser, struct c_expr *after,
struct c_expr cond, exp1, exp2, ret;
location_t start, cond_loc, colon_loc;
bool save_c_omp_array_section_p = c_omp_array_section_p;
+ bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
gcc_assert (!after || c_dialect_objc ());
@@ -8318,6 +8322,7 @@ c_parser_conditional_expression (c_parser *parser, struct c_expr *after,
if (c_parser_next_token_is_not (parser, CPP_QUERY))
return cond;
c_omp_array_section_p = false;
+ c_omp_array_shaping_op_p = false;
if (cond.value != error_mark_node)
start = cond.get_start ();
else
@@ -8371,6 +8376,7 @@ c_parser_conditional_expression (c_parser *parser, struct c_expr *after,
ret.original_code = ERROR_MARK;
ret.original_type = NULL;
c_omp_array_section_p = save_c_omp_array_section_p;
+ c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
return ret;
}
{
@@ -8418,6 +8424,7 @@ c_parser_conditional_expression (c_parser *parser, struct c_expr *after,
set_c_expr_source_range (&ret, start, exp2.get_finish ());
ret.m_decimal = 0;
c_omp_array_section_p = save_c_omp_array_section_p;
+ c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
return ret;
}
@@ -8799,6 +8806,8 @@ c_parser_cast_expression (c_parser *parser, struct c_expr *after)
if (after)
return c_parser_postfix_expression_after_primary (parser,
cast_loc, *after);
+ bool save_c_omp_has_array_shape_p = c_omp_has_array_shape_p;
+ c_omp_has_array_shape_p = false;
/* If the expression begins with a parenthesized type name, it may
be either a cast or a compound literal; we need to see whether
the next character is '{' to tell the difference. If not, it is
@@ -8807,6 +8816,10 @@ c_parser_cast_expression (c_parser *parser, struct c_expr *after)
if (c_parser_next_token_is (parser, CPP_OPEN_PAREN)
&& c_token_starts_compound_literal (c_parser_peek_2nd_token (parser)))
{
+ bool save_c_omp_array_section_p = c_omp_array_section_p;
+ bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
+ c_omp_array_section_p = false;
+ c_omp_array_shaping_op_p = false;
struct c_declspecs *scspecs;
struct c_type_name *type_name;
struct c_expr ret;
@@ -8818,6 +8831,8 @@ c_parser_cast_expression (c_parser *parser, struct c_expr *after)
parens.skip_until_found_close (parser);
if (type_name == NULL)
{
+ c_omp_array_section_p = save_c_omp_array_section_p;
+ c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
ret.set_error ();
ret.original_code = ERROR_MARK;
ret.original_type = NULL;
@@ -8828,9 +8843,15 @@ c_parser_cast_expression (c_parser *parser, struct c_expr *after)
used_types_insert (type_name->specs->type);
if (c_parser_next_token_is (parser, CPP_OPEN_BRACE))
- return c_parser_postfix_expression_after_paren_type (parser, scspecs,
- type_name,
- cast_loc);
+ {
+ c_expr r = c_parser_postfix_expression_after_paren_type (parser,
+ scspecs,
+ type_name,
+ cast_loc);
+ c_omp_array_section_p = save_c_omp_array_section_p;
+ c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
+ return r;
+ }
if (scspecs)
error_at (cast_loc, "storage class specifier in cast");
if (type_name->specs->alignas_p)
@@ -8847,10 +8868,61 @@ c_parser_cast_expression (c_parser *parser, struct c_expr *after)
ret.original_code = ERROR_MARK;
ret.original_type = NULL;
ret.m_decimal = 0;
+ c_omp_array_section_p = save_c_omp_array_section_p;
+ c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
+ return ret;
+ }
+ else if (c_omp_array_shaping_op_p
+ && c_parser_next_token_is (parser, CPP_OPEN_PAREN)
+ && c_parser_peek_2nd_token (parser)->type == CPP_OPEN_SQUARE)
+ {
+ bool save_c_omp_array_section_p = c_omp_array_section_p;
+ bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
+ c_omp_array_section_p = false;
+ c_omp_array_shaping_op_p = false;
+ auto_vec<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).
@@ -9860,6 +9932,7 @@ c_parser_postfix_expression (c_parser *parser)
tree stmt;
location_t brace_loc;
bool save_c_omp_array_section_p = c_omp_array_section_p;
+ bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
c_parser_consume_token (parser);
brace_loc = c_parser_peek_token (parser)->location;
c_parser_consume_token (parser);
@@ -9877,6 +9950,7 @@ c_parser_postfix_expression (c_parser *parser)
break;
}
c_omp_array_section_p = false;
+ c_omp_array_shaping_op_p = false;
stmt = c_begin_stmt_expr ();
c_parser_compound_statement_nostart (parser);
location_t close_loc = c_parser_peek_token (parser)->location;
@@ -9888,6 +9962,7 @@ c_parser_postfix_expression (c_parser *parser)
set_c_expr_source_range (&expr, loc, close_loc);
mark_exp_read (expr.value);
c_omp_array_section_p = save_c_omp_array_section_p;
+ c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
}
else
{
@@ -11373,20 +11448,26 @@ c_parser_postfix_expression_after_primary (c_parser *parser,
if (c_omp_array_section_p
&& c_parser_next_token_is (parser, CPP_COLON))
{
+ tree stride = NULL_TREE;
+
c_parser_consume_token (parser);
if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE))
len = c_parser_expression (parser).value;
+ if (c_parser_next_token_is (parser, CPP_COLON))
+ {
+ c_parser_consume_token (parser);
+ if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE))
+ stride = c_parser_expression (parser).value;
+ }
+
c_parser_skip_until_found (parser, CPP_CLOSE_SQUARE,
"expected %<]%>");
- /* NOTE: We are reusing using the type of the whole array as the
- type of the array section here, which isn't necessarily
- entirely correct. Might need revisiting. */
start = expr.get_start ();
finish = parser->tokens_buf[0].location;
expr.value = build_omp_array_section (op_loc, expr.value, idx,
- len, NULL_TREE /* fixme */);
+ len, stride);
set_c_expr_source_range (&expr, start, finish);
expr.original_code = ERROR_MARK;
expr.original_type = NULL;
@@ -11397,7 +11478,20 @@ c_parser_postfix_expression_after_primary (c_parser *parser,
"expected %<]%>");
start = expr.get_start ();
finish = parser->tokens_buf[0].location;
- expr.value = build_array_ref (op_loc, expr.value, idx);
+ if (c_omp_has_array_shape_p)
+ /* If we have an array-shaping operator, we may not be able to
+ represent a well-formed ARRAY_REF here, because we are
+ coercing the type of the innermost array base and the
+ original type may not be compatible. Use the
+ OMP_ARRAY_SECTION code instead. We also want to explicitly
+ avoid creating INDIRECT_REFs for pointer bases, because
+ that can lead to parsing ambiguities (see
+ c_parser_omp_variable_list). */
+ expr.value
+ = build_omp_array_section (op_loc, expr.value, idx,
+ size_one_node, NULL_TREE);
+ else
+ expr.value = build_array_ref (op_loc, expr.value, idx);
set_c_expr_source_range (&expr, start, finish);
expr.original_code = ERROR_MARK;
expr.original_type = NULL;
@@ -11694,7 +11788,9 @@ c_parser_expr_list (c_parser *parser, bool convert_p, bool fold_p,
struct c_expr expr;
unsigned int idx = 0;
bool save_c_omp_array_section_p = c_omp_array_section_p;
+ bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
c_omp_array_section_p = false;
+ c_omp_array_shaping_op_p = false;
ret = make_tree_vector ();
if (p_orig_types == NULL)
@@ -11749,6 +11845,7 @@ c_parser_expr_list (c_parser *parser, bool convert_p, bool fold_p,
if (orig_types)
*p_orig_types = orig_types;
c_omp_array_section_p = save_c_omp_array_section_p;
+ c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
return ret;
}
@@ -13939,6 +14036,35 @@ c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list)
return list;
}
+/* Return, as an INTEGER_CST node, the number of elements for TYPE
+ (which is an ARRAY_TYPE). This counts only elements of the top
+ array. (From cp/tree.cc). */
+
+static tree
+c_array_type_nelts_top (tree type)
+{
+ return fold_build2_loc (input_location, PLUS_EXPR, sizetype,
+ array_type_nelts (type), size_one_node);
+}
+
+/* Return, as an INTEGER_CST node, the number of elements for TYPE
+ (which is an ARRAY_TYPE). This one is a recursive count of all
+ ARRAY_TYPEs that are clumped together. (From cp/tree.cc). */
+
+static tree
+c_array_type_nelts_total (tree type)
+{
+ tree sz = c_array_type_nelts_top (type);
+ type = TREE_TYPE (type);
+ while (TREE_CODE (type) == ARRAY_TYPE)
+ {
+ tree n = c_array_type_nelts_top (type);
+ sz = fold_build2_loc (input_location, MULT_EXPR, sizetype, sz, n);
+ type = TREE_TYPE (type);
+ }
+ return sz;
+}
+
/* OpenACC 2.0, OpenMP 2.5:
variable-list:
identifier
@@ -14067,12 +14193,24 @@ c_parser_omp_variable_list (c_parser *parser,
{
location_t loc = c_parser_peek_token (parser)->location;
bool save_c_omp_array_section_p = c_omp_array_section_p;
+ bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
c_omp_array_section_p = true;
+ c_omp_array_shaping_op_p
+ = (kind == OMP_CLAUSE_TO || kind == OMP_CLAUSE_FROM);
c_expr expr = c_parser_expr_no_commas (parser, NULL);
if (expr.value != error_mark_node)
mark_exp_read (expr.value);
c_omp_array_section_p = save_c_omp_array_section_p;
+ c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
tree decl = expr.value;
+ tree reshaped_to = NULL_TREE;
+
+ if (TREE_CODE (decl) == VIEW_CONVERT_EXPR
+ && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+ {
+ reshaped_to = TREE_TYPE (decl);
+ decl = TREE_OPERAND (decl, 0);
+ }
/* This code rewrites a parsed expression containing various tree
codes used to represent array accesses into a more uniform nest of
@@ -14085,6 +14223,31 @@ c_parser_omp_variable_list (c_parser *parser,
dims.truncate (0);
if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
{
+ size_t sections = 0;
+ tree orig_decl = decl;
+ bool update_p = (kind == OMP_CLAUSE_TO
+ || kind == OMP_CLAUSE_FROM);
+ bool maybe_ptr_based_noncontig_update = false;
+
+ while (update_p
+ && !reshaped_to
+ && (TREE_CODE (decl) == OMP_ARRAY_SECTION
+ || TREE_CODE (decl) == ARRAY_REF
+ || TREE_CODE (decl) == COMPOUND_EXPR))
+ {
+ if (TREE_CODE (decl) == COMPOUND_EXPR)
+ decl = TREE_OPERAND (decl, 1);
+ else
+ {
+ if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
+ maybe_ptr_based_noncontig_update = true;
+ decl = TREE_OPERAND (decl, 0);
+ sections++;
+ }
+ }
+
+ decl = orig_decl;
+
while (TREE_CODE (decl) == OMP_ARRAY_SECTION)
{
tree low_bound = TREE_OPERAND (decl, 1);
@@ -14093,18 +14256,63 @@ c_parser_omp_variable_list (c_parser *parser,
dims.safe_push (omp_dim (low_bound, length, stride, loc,
false));
decl = TREE_OPERAND (decl, 0);
+ if (sections > 0)
+ sections--;
}
+ /* The handling of INDIRECT_REF here in the presence of
+ array-shaping operations is a little tricky. We need to
+ avoid treating a pointer dereference as a unit-sized array
+ section when we have an array shaping operation, because we
+ don't want an indirection to consume one of the user's
+ requested array dimensions. E.g. if we have a
+ double-indirect pointer like:
+
+ int **foopp;
+ #pragma omp target update from(([N][N]) (*foopp)[0:X][0:Y])
+
+ We don't want to interpret this as:
+
+ foopp[0:1][0:X][0:Y]
+
+ else the array shape [N][N] won't match. Also we can't match
+ the array sections right-to-left instead, else this:
+
+ #pragma omp target update from(([N][N]) (*foopp)[0:X])
+
+ would not copy the dimensions:
+
+ (*foopp)[0:X][0:N]
+
+ as required. So, avoid descending through INDIRECT_REFs if
+ we have an array-shaping op.
+
+ If we *don't* have an array-shaping op, but we have a
+ multiply-indirected pointer and an array section like this:
+
+ int ***fooppp;
+ #pragma omp target update from((**fooppp)[0:X:S]
+
+ also avoid descending through more indirections than we have
+ array sections, since the noncontiguous update processing code
+ won't understand them (and doesn't need to traverse them
+ anyway). */
+
while (TREE_CODE (decl) == ARRAY_REF
- || TREE_CODE (decl) == INDIRECT_REF
+ || (TREE_CODE (decl) == INDIRECT_REF
+ && !reshaped_to)
|| TREE_CODE (decl) == COMPOUND_EXPR)
{
+ if (maybe_ptr_based_noncontig_update && sections == 0)
+ break;
+
if (TREE_CODE (decl) == COMPOUND_EXPR)
{
decl = TREE_OPERAND (decl, 1);
STRIP_NOPS (decl);
}
- else if (TREE_CODE (decl) == INDIRECT_REF)
+ else if (TREE_CODE (decl) == INDIRECT_REF
+ && !reshaped_to)
{
dims.safe_push (omp_dim (integer_zero_node,
integer_one_node, NULL_TREE, loc,
@@ -14117,6 +14325,35 @@ c_parser_omp_variable_list (c_parser *parser,
dims.safe_push (omp_dim (index, integer_one_node,
NULL_TREE, loc, true));
decl = TREE_OPERAND (decl, 0);
+ if (sections > 0)
+ sections--;
+ }
+ }
+
+ if (reshaped_to)
+ {
+ unsigned reshaped_dims = 0;
+
+ for (tree t = reshaped_to;
+ TREE_CODE (t) == ARRAY_TYPE;
+ t = TREE_TYPE (t))
+ reshaped_dims++;
+
+ if (dims.length () > reshaped_dims)
+ {
+ error_at (loc, "too many array section specifiers "
+ "for %qT", reshaped_to);
+ decl = error_mark_node;
+ }
+ else
+ {
+ /* We have a pointer DECL whose target should be
+ interpreted as an array with particular dimensions,
+ not "the pointer itself". So, add an indirection
+ here. */
+ decl = build_indirect_ref (loc, decl, RO_UNARY_STAR);
+ decl = build1_loc (loc, VIEW_CONVERT_EXPR, reshaped_to,
+ decl);
}
}
@@ -14144,6 +14381,14 @@ c_parser_omp_variable_list (c_parser *parser,
decl = build_omp_array_section (loc, decl, idx, integer_one_node,
NULL_TREE);
}
+ else if (reshaped_to)
+ {
+ /* We're copying the whole of a reshaped array, originally a
+ base pointer. Rewrite as an array section. */
+ tree elems = c_array_type_nelts_total (reshaped_to);
+ decl = build_omp_array_section (loc, decl, size_zero_node, elems,
+ NULL_TREE);
+ }
else if (TREE_CODE (decl) == NON_LVALUE_EXPR
|| CONVERT_EXPR_P (decl))
decl = TREE_OPERAND (decl, 0);
@@ -22777,8 +23022,38 @@ c_parser_omp_target_update (location_t loc, c_parser *parser,
tree clauses
= c_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK,
"#pragma omp target update");
- if (omp_find_clause (clauses, OMP_CLAUSE_TO) == NULL_TREE
- && omp_find_clause (clauses, OMP_CLAUSE_FROM) == NULL_TREE)
+ bool to_clause = false, from_clause = false;
+ for (tree c = clauses;
+ c && !to_clause && !from_clause;
+ c = OMP_CLAUSE_CHAIN (c))
+ {
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_TO:
+ to_clause = true;
+ break;
+ case OMP_CLAUSE_FROM:
+ from_clause = true;
+ break;
+ case OMP_CLAUSE_MAP:
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case GOMP_MAP_TO_GRID:
+ to_clause = true;
+ break;
+ case GOMP_MAP_FROM_GRID:
+ from_clause = true;
+ break;
+ default:
+ ;
+ }
+ break;
+ default:
+ ;
+ }
+ }
+
+ if (!to_clause && !from_clause)
{
error_at (loc,
"%<#pragma omp target update%> must contain at least one "
@@ -727,6 +727,8 @@ extern int in_sizeof;
extern int in_typeof;
extern bool c_in_omp_for;
extern bool c_omp_array_section_p;
+extern bool c_omp_array_shaping_op_p;
+extern bool c_omp_has_array_shape_p;
extern tree c_last_sizeof_arg;
extern location_t c_last_sizeof_loc;
@@ -766,6 +768,8 @@ extern tree build_component_ref (location_t, tree, tree, location_t,
location_t);
extern tree build_array_ref (location_t, tree, tree);
extern tree build_omp_array_section (location_t, tree, tree, tree, tree);
+extern tree create_omp_arrayshape_type (tree expr,
+ vec<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);
@@ -79,6 +79,13 @@ bool c_in_omp_for;
/* True when parsing OpenMP map clause. */
bool c_omp_array_section_p;
+/* True when parsing OpenMP to/from clause. */
+bool c_omp_array_shaping_op_p;
+
+/* True if we have an OpenMP array-shaping "cast" expression. This adjusts
+ the parsed representation for e.g. array refs. */
+bool c_omp_has_array_shape_p;
+
/* The argument of last parsed sizeof expression, only to be tested
if expr.original_code == SIZEOF_EXPR. */
tree c_last_sizeof_arg;
@@ -2930,8 +2937,8 @@ build_omp_array_section (location_t loc, tree array, tree index, tree length,
if (index != NULL_TREE
&& length != NULL_TREE
- && TREE_CODE (index) == INTEGER_CST
- && TREE_CODE (length) == INTEGER_CST)
+ && INTEGRAL_TYPE_P (TREE_TYPE (index))
+ && INTEGRAL_TYPE_P (TREE_TYPE (length)))
{
tree low = fold_convert (sizetype, index);
tree high = fold_convert (sizetype, length);
@@ -2941,7 +2948,7 @@ build_omp_array_section (location_t loc, tree array, tree index, tree length,
}
else if ((index == NULL_TREE || integer_zerop (index))
&& length != NULL_TREE
- && TREE_CODE (length) == INTEGER_CST)
+ && INTEGRAL_TYPE_P (TREE_TYPE (length)))
idxtype = build_index_type (length);
else
idxtype = NULL_TREE;
@@ -2965,6 +2972,46 @@ build_omp_array_section (location_t loc, tree array, tree index, tree length,
stride);
}
+/* Build an array type whose dimensions are given by OMP_SHAPE_DIMS and whose
+ elements are of the type pointed to by the "base" node of EXPR with outer
+ OMP_ARRAY_SECTIONs and ARRAY_REFs stripped off, e.g. the type of "*myptr"
+ in "myptr[0:2:3][4][5:6]". */
+
+tree
+create_omp_arrayshape_type (tree expr, vec<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
@@ -13714,7 +13761,8 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses)
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<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;
@@ -13795,11 +13843,14 @@ 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);
@@ -13840,8 +13891,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))
@@ -13960,12 +14018,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)
@@ -13973,7 +14048,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))
{
@@ -14047,7 +14122,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
{
tree d_length = TREE_OPERAND (d, 2);
tree d_stride = TREE_OPERAND (d, 3);
- if (d_length == NULL_TREE || !integer_onep (d_length)
+ if (d_length == NULL_TREE
+ || !integer_onep (d_length)
|| (d_stride && !integer_onep (d_stride)))
{
if (openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
@@ -14068,10 +14144,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec<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;
+ }
}
}
}
@@ -14083,7 +14164,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)
@@ -14091,14 +14172,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, enum c_omp_region_type ort)
+handle_omp_array_sections (tree *pc, enum c_omp_region_type ort,
+ int *discontiguous)
{
tree c = *pc;
bool maybe_zero_len = false;
@@ -14114,7 +14223,7 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort)
tp = &TREE_VALUE (*tp);
tree first = handle_omp_array_sections_1 (c, *tp, types,
maybe_zero_len, first_non_one,
- non_contiguous, ort);
+ non_contiguous, ort, discontiguous);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -14153,11 +14262,14 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort)
if (int_size_in_bytes (TREE_TYPE (first)) <= 0)
maybe_zero_len = true;
+ bool higher_discontiguous = false;
+
for (i = num, t = OMP_CLAUSE_DECL (c); i > 0;
t = TREE_OPERAND (t, 0))
{
tree low_bound = TREE_OPERAND (t, 1);
tree length = TREE_OPERAND (t, 2);
+ tree stride = TREE_OPERAND (t, 3);
i--;
if (low_bound
@@ -14170,6 +14282,11 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort)
&& TYPE_PRECISION (TREE_TYPE (length))
> TYPE_PRECISION (sizetype))
length = fold_convert (sizetype, length);
+ if (stride
+ && TREE_CODE (stride) == INTEGER_CST
+ && TYPE_PRECISION (TREE_TYPE (stride))
+ > TYPE_PRECISION (sizetype))
+ stride = fold_convert (sizetype, stride);
if (low_bound == NULL_TREE)
low_bound = integer_zero_node;
@@ -14179,10 +14296,49 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort)
continue;
}
+ if (stride == NULL_TREE)
+ stride = size_one_node;
+ if (discontiguous && *discontiguous)
+ {
+ /* This condition is similar to the error check below, but
+ whereas that checks for a definitely-discontiguous array
+ section in order to report an error (where such a section is
+ illegal), here we instead need to know if the array section
+ *may be* discontiguous so we can handle that case
+ appropriately (i.e. for rectangular "target update"
+ operations). */
+ bool full_span = false;
+ if (length != NULL_TREE
+ && TREE_CODE (length) == INTEGER_CST
+ && TREE_CODE (types[i]) == ARRAY_TYPE
+ && TYPE_DOMAIN (types[i])
+ && TYPE_MAX_VALUE (TYPE_DOMAIN (types[i]))
+ && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])))
+ == INTEGER_CST)
+ {
+ tree size;
+ size = size_binop (PLUS_EXPR,
+ TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])),
+ size_one_node);
+ if (tree_int_cst_equal (length, size))
+ full_span = true;
+ }
+
+ if (!integer_onep (stride)
+ || (higher_discontiguous
+ && (!integer_zerop (low_bound)
+ || !full_span)))
+ *discontiguous = 2;
+
+ if (!integer_onep (stride)
+ || !integer_zerop (low_bound)
+ || !full_span)
+ higher_discontiguous = true;
+ }
if (!maybe_zero_len && i > first_non_one)
{
if (integer_nonzerop (low_bound))
- goto do_warn_noncontiguous;
+ goto is_noncontiguous;
if (length != NULL_TREE
&& TREE_CODE (length) == INTEGER_CST
&& TYPE_DOMAIN (types[i])
@@ -14196,12 +14352,17 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort)
size_one_node);
if (!tree_int_cst_equal (length, size))
{
- do_warn_noncontiguous:
- error_at (OMP_CLAUSE_LOCATION (c),
- "array section is not contiguous in %qs "
- "clause",
- omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
- return true;
+ is_noncontiguous:
+ if (discontiguous && *discontiguous)
+ *discontiguous = 2;
+ else
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "array section is not contiguous in %qs "
+ "clause",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ return true;
+ }
}
}
if (length != NULL_TREE
@@ -14321,6 +14482,8 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort)
OMP_CLAUSE_DECL (c) = t;
return false;
}
+ if (discontiguous && *discontiguous != 2)
+ first = omp_array_section_low_bound (OMP_CLAUSE_LOCATION (c), first);
first = c_fully_fold (first, false, NULL);
OMP_CLAUSE_DECL (c) = first;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
@@ -14329,7 +14492,8 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort)
size = c_fully_fold (size, false, NULL);
OMP_CLAUSE_SIZE (c) = size;
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ && !(discontiguous && *discontiguous == 2))
return false;
auto_vec<omp_addr_token *, 10> addr_tokens;
@@ -14342,7 +14506,8 @@ handle_omp_array_sections (tree *pc, enum c_omp_region_type ort)
tree *npc = ai.expand_map_clause (pc, first, addr_tokens, ort);
if (npc != NULL)
{
- if (ai.maybe_zero_length_array_section (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && ai.maybe_zero_length_array_section (c))
OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
return false;
@@ -14689,7 +14854,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
- if (handle_omp_array_sections (pc, ort))
+ if (handle_omp_array_sections (pc, ort, NULL))
{
remove = true;
break;
@@ -15428,7 +15593,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
last_iterators = NULL_TREE;
if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
- if (handle_omp_array_sections (pc, ort))
+ if (handle_omp_array_sections (pc, ort, NULL))
remove = true;
else if ((c = *pc)
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
@@ -15535,6 +15700,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
remove = true;
break;
}
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE)
+ break;
/* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
@@ -15549,7 +15717,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
grp_start_p = pc;
grp_sentinel = OMP_CLAUSE_CHAIN (c);
- if (handle_omp_array_sections (pc, ort))
+ int discontiguous
+ = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM);
+ if (handle_omp_array_sections (pc, ort, &discontiguous))
remove = true;
else
{
@@ -15946,7 +16117,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
- if (handle_omp_array_sections (pc, ort))
+ if (handle_omp_array_sections (pc, ort, NULL))
remove = true;
else
{
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}