@@ -1376,7 +1376,7 @@ public:
bool maybe_zero_length_array_section (tree);
tree * expand_array_base (tree *, vec<omp_addr_token *> &, tree, unsigned *,
- c_omp_region_type, bool);
+ c_omp_region_type);
tree * expand_component_selector (tree *, vec<omp_addr_token *> &, tree,
unsigned *);
tree * expand_map_clause (tree *, tree, vec<omp_addr_token *> &,
@@ -4042,7 +4042,9 @@ c_omp_address_inspector::map_supported_p ()
|| TREE_CODE (t) == POINTER_PLUS_EXPR
|| TREE_CODE (t) == NON_LVALUE_EXPR
|| TREE_CODE (t) == OMP_ARRAY_SECTION
- || TREE_CODE (t) == NOP_EXPR)
+ || TREE_CODE (t) == NOP_EXPR
+ || TREE_CODE (t) == VIEW_CONVERT_EXPR
+ || TREE_CODE (t) == ADDR_EXPR)
if (TREE_CODE (t) == COMPOUND_EXPR)
t = TREE_OPERAND (t, 1);
else
@@ -4192,21 +4194,95 @@ omp_expand_access_chain (tree *pc, tree expr,
return pc;
}
+static tree *
+omp_expand_grid_dim (location_t loc, tree *pc, tree decl)
+{
+ if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
+ pc = omp_expand_grid_dim (loc, pc, TREE_OPERAND (decl, 0));
+ else
+ return pc;
+
+ tree c = *pc;
+ tree low_bound = TREE_OPERAND (decl, 1);
+ tree length = TREE_OPERAND (decl, 2);
+ tree stride = TREE_OPERAND (decl, 3);
+
+ tree cd = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (cd, GOMP_MAP_GRID_DIM);
+ OMP_CLAUSE_DECL (cd) = unshare_expr (low_bound);
+ OMP_CLAUSE_SIZE (cd) = unshare_expr (length);
+
+ if (stride && !integer_onep (stride))
+ {
+ tree cs = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (cs, GOMP_MAP_GRID_STRIDE);
+ OMP_CLAUSE_DECL (cs) = unshare_expr (stride);
+
+ OMP_CLAUSE_CHAIN (cs) = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (cd) = cs;
+ OMP_CLAUSE_CHAIN (c) = cd;
+ pc = &OMP_CLAUSE_CHAIN (cd);
+ }
+ else
+ {
+ OMP_CLAUSE_CHAIN (cd) = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = cd;
+ pc = &OMP_CLAUSE_CHAIN (c);
+ }
+
+ return pc;
+}
+
+tree *
+omp_handle_noncontig_array (location_t loc, tree *pc, tree c, tree base)
+{
+ tree type;
+
+ if (POINTER_TYPE_P (TREE_TYPE (base)))
+ type = TREE_TYPE (TREE_TYPE (base));
+ else
+ type = strip_array_types (TREE_TYPE (base));
+
+ tree c_map = build_omp_clause (loc, OMP_CLAUSE_MAP);
+
+ OMP_CLAUSE_DECL (c_map) = unshare_expr (base);
+ /* Use the element size (or pointed-to type size) here. */
+ OMP_CLAUSE_SIZE (c_map) = TYPE_SIZE_UNIT (type);
+
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_TO:
+ OMP_CLAUSE_SET_MAP_KIND (c_map, GOMP_MAP_TO_GRID);
+ break;
+ case OMP_CLAUSE_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (c_map, GOMP_MAP_FROM_GRID);
+ break;
+ default:
+ gcc_unreachable ();
+ }
+
+ OMP_CLAUSE_CHAIN (c_map) = OMP_CLAUSE_CHAIN (c);
+
+ *pc = c_map;
+
+ return omp_expand_grid_dim (loc, pc, OMP_CLAUSE_DECL (c));
+}
+
/* Translate "array_base_decl access_method" to OMP mapping clauses. */
tree *
c_omp_address_inspector::expand_array_base (tree *pc,
vec<omp_addr_token *> &addr_tokens,
tree expr, unsigned *idx,
- c_omp_region_type ort,
- bool decl_p)
+ c_omp_region_type ort)
{
using namespace omp_addr_tokenizer;
tree c = *pc;
location_t loc = OMP_CLAUSE_LOCATION (c);
int i = *idx;
tree decl = addr_tokens[i + 1]->expr;
- bool declare_target_p = (decl_p
+ bool decl_p = DECL_P (decl);
+ bool declare_target_p = (DECL_P (decl)
&& is_global_var (decl)
&& lookup_attribute ("omp declare target",
DECL_ATTRIBUTES (decl)));
@@ -4218,6 +4294,7 @@ c_omp_address_inspector::expand_array_base (tree *pc,
unsigned consume_tokens = 2;
bool target = (ort & C_ORT_TARGET) != 0;
bool openmp = (ort & C_ORT_OMP) != 0;
+ unsigned acc = i + 1;
gcc_assert (i == 0);
@@ -4230,7 +4307,15 @@ c_omp_address_inspector::expand_array_base (tree *pc,
return pc;
}
- switch (addr_tokens[i + 1]->u.access_kind)
+ if (!map_p && chain_p)
+ {
+ /* See comment in c_omp_address_inspector::expand_component_selector. */
+ while (acc + 1 < addr_tokens.length ()
+ && addr_tokens[acc + 1]->type == ACCESS_METHOD)
+ acc++;
+ }
+
+ switch (addr_tokens[acc]->u.access_kind)
{
case ACCESS_DIRECT:
if (decl_p && !target)
@@ -4474,6 +4559,40 @@ c_omp_address_inspector::expand_array_base (tree *pc,
}
break;
+ case ACCESS_NONCONTIG_ARRAY:
+ {
+ gcc_assert (!map_p);
+
+ tree base = addr_tokens[acc]->expr;
+
+ if (decl_p)
+ c_common_mark_addressable_vec (base);
+
+ pc = omp_handle_noncontig_array (loc, pc, c, base);
+ consume_tokens = (acc + 1) - i;
+ chain_p = false;
+ }
+ break;
+
+ case ACCESS_NONCONTIG_REF_TO_ARRAY:
+ {
+ gcc_assert (!map_p);
+
+ if (decl_p)
+ c_common_mark_addressable_vec (addr_tokens[acc]->expr);
+
+ /* Or here. */
+ gcc_assert (!chain_p);
+
+ tree base = addr_tokens[i + 1]->expr;
+ base = convert_from_reference (base);
+
+ pc = omp_handle_noncontig_array (loc, pc, c, base);
+ consume_tokens = (acc + 1) - i;
+ chain_p = false;
+ }
+ break;
+
default:
*idx = i + consume_tokens;
return NULL;
@@ -4524,8 +4643,27 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
tree c2 = NULL_TREE, c3 = NULL_TREE;
bool chain_p = omp_access_chain_p (addr_tokens, i + 1);
bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP;
+ unsigned acc = i + 1;
- switch (addr_tokens[i + 1]->u.access_kind)
+ if (!map_p && chain_p)
+ {
+ /* We have a non-map clause (i.e. to/from for an "update" directive),
+ and we might have a noncontiguous array section at the end of a
+ chain of other accesses, e.g. pointer indirections like this:
+
+ struct_base_decl access_pointer access_pointer component_selector
+ access_pointer access_pointer access_noncontig_array
+
+ We only need to process the last access in this case, so skip
+ over previous accesses. */
+
+ while (acc + 1 < addr_tokens.length ()
+ && addr_tokens[acc + 1]->type == ACCESS_METHOD)
+ acc++;
+ chain_p = false;
+ }
+
+ switch (addr_tokens[acc]->u.access_kind)
{
case ACCESS_DIRECT:
case ACCESS_INDEXED_ARRAY:
@@ -4535,7 +4673,7 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
{
/* Copy the referenced object. Note that we also do this for !MAP_P
clauses. */
- tree obj = convert_from_reference (addr_tokens[i + 1]->expr);
+ tree obj = convert_from_reference (addr_tokens[acc]->expr);
OMP_CLAUSE_DECL (c) = obj;
OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (obj));
@@ -4544,7 +4682,7 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
- OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
+ OMP_CLAUSE_DECL (c2) = addr_tokens[acc]->expr;
OMP_CLAUSE_SIZE (c2) = size_zero_node;
}
break;
@@ -4555,15 +4693,15 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
break;
tree virtual_origin
- = convert_from_reference (addr_tokens[i + 1]->expr);
+ = convert_from_reference (addr_tokens[acc]->expr);
virtual_origin = build_fold_addr_expr (virtual_origin);
virtual_origin = fold_convert_loc (loc, ptrdiff_type_node,
virtual_origin);
- tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
+ tree data_addr = omp_accessed_addr (addr_tokens, acc, expr);
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
- OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
+ OMP_CLAUSE_DECL (c2) = addr_tokens[acc]->expr;
OMP_CLAUSE_SIZE (c2)
= fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node,
fold_convert_loc (loc, ptrdiff_type_node,
@@ -4580,12 +4718,12 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
tree virtual_origin
= fold_convert_loc (loc, ptrdiff_type_node,
- addr_tokens[i + 1]->expr);
- tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
+ addr_tokens[acc]->expr);
+ tree data_addr = omp_accessed_addr (addr_tokens, acc, expr);
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
- OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
+ OMP_CLAUSE_DECL (c2) = addr_tokens[acc]->expr;
OMP_CLAUSE_SIZE (c2)
= fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node,
fold_convert_loc (loc, ptrdiff_type_node,
@@ -4600,10 +4738,10 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
if (!map_p)
break;
- tree ptr = convert_from_reference (addr_tokens[i + 1]->expr);
+ tree ptr = convert_from_reference (addr_tokens[acc]->expr);
tree virtual_origin = fold_convert_loc (loc, ptrdiff_type_node,
ptr);
- tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
+ tree data_addr = omp_accessed_addr (addr_tokens, acc, expr);
/* Attach the pointer... */
c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
@@ -4618,13 +4756,38 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
/* ...and also the reference. */
c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH_DETACH);
- OMP_CLAUSE_DECL (c3) = addr_tokens[i + 1]->expr;
+ OMP_CLAUSE_DECL (c3) = addr_tokens[acc]->expr;
OMP_CLAUSE_SIZE (c3) = size_zero_node;
}
break;
+ case ACCESS_NONCONTIG_ARRAY:
+ {
+ gcc_assert (!map_p);
+
+ /* We don't expect to see further accesses here. */
+ gcc_assert (!chain_p);
+
+ pc = omp_handle_noncontig_array (loc, pc, c, addr_tokens[acc]->expr);
+ }
+ break;
+
+ case ACCESS_NONCONTIG_REF_TO_ARRAY:
+ {
+ gcc_assert (!map_p);
+
+ /* Or here. */
+ gcc_assert (!chain_p);
+
+ tree base = addr_tokens[acc]->expr;
+ base = convert_from_reference (base);
+
+ pc = omp_handle_noncontig_array (loc, pc, c, base);
+ }
+ break;
+
default:
- *idx = i + 2;
+ *idx = acc + 1;
return NULL;
}
@@ -4642,8 +4805,7 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
pc = &OMP_CLAUSE_CHAIN (c);
}
- i += 2;
- *idx = i;
+ *idx = acc + 1;
if (chain_p && map_p)
return omp_expand_access_chain (pc, expr, addr_tokens, idx);
@@ -4671,7 +4833,7 @@ c_omp_address_inspector::expand_map_clause (tree *pc, tree expr,
&& addr_tokens[i]->u.structure_base_kind == BASE_DECL
&& addr_tokens[i + 1]->type == ACCESS_METHOD)
{
- pc = expand_array_base (pc, addr_tokens, expr, &i, ort, true);
+ pc = expand_array_base (pc, addr_tokens, expr, &i, ort);
if (pc == NULL)
return NULL;
}
@@ -4680,7 +4842,7 @@ c_omp_address_inspector::expand_map_clause (tree *pc, tree expr,
&& addr_tokens[i]->u.structure_base_kind == BASE_ARBITRARY_EXPR
&& addr_tokens[i + 1]->type == ACCESS_METHOD)
{
- pc = expand_array_base (pc, addr_tokens, expr, &i, ort, false);
+ pc = expand_array_base (pc, addr_tokens, expr, &i, ort);
if (pc == NULL)
return NULL;
}
@@ -1623,6 +1623,11 @@ c_pretty_printer::postfix_expression (tree e)
pp_colon (this);
if (TREE_OPERAND (e, 2))
expression (TREE_OPERAND (e, 2));
+ if (TREE_OPERAND (e, 3))
+ {
+ pp_colon (this);
+ expression (TREE_OPERAND (e, 3));
+ }
pp_c_right_bracket (this);
break;
@@ -11386,7 +11386,7 @@ c_parser_postfix_expression_after_primary (c_parser *parser,
start = expr.get_start ();
finish = parser->tokens_buf[0].location;
expr.value = build_omp_array_section (op_loc, expr.value, idx,
- len);
+ len, NULL_TREE /* fixme */);
set_c_expr_source_range (&expr, start, finish);
expr.original_code = ERROR_MARK;
expr.original_type = NULL;
@@ -13956,11 +13956,11 @@ c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list)
struct omp_dim
{
- tree low_bound, length;
+ tree low_bound, length, stride;
location_t loc;
bool no_colon;
- omp_dim (tree lb, tree len, location_t lo, bool nc)
- : low_bound (lb), length (len), loc (lo), no_colon (nc) {}
+ omp_dim (tree lb, tree len, tree str, location_t lo, bool nc)
+ : low_bound (lb), length (len), stride (str), loc (lo), no_colon (nc) {}
};
static tree
@@ -14089,7 +14089,9 @@ c_parser_omp_variable_list (c_parser *parser,
{
tree low_bound = TREE_OPERAND (decl, 1);
tree length = TREE_OPERAND (decl, 2);
- dims.safe_push (omp_dim (low_bound, length, loc, false));
+ tree stride = TREE_OPERAND (decl, 3);
+ dims.safe_push (omp_dim (low_bound, length, stride, loc,
+ false));
decl = TREE_OPERAND (decl, 0);
}
@@ -14105,21 +14107,22 @@ c_parser_omp_variable_list (c_parser *parser,
else if (TREE_CODE (decl) == INDIRECT_REF)
{
dims.safe_push (omp_dim (integer_zero_node,
- integer_one_node, loc, true));
+ integer_one_node, NULL_TREE, loc,
+ true));
decl = TREE_OPERAND (decl, 0);
}
else /* ARRAY_REF. */
{
tree index = TREE_OPERAND (decl, 1);
- dims.safe_push (omp_dim (index, integer_one_node, loc,
- true));
+ dims.safe_push (omp_dim (index, integer_one_node,
+ NULL_TREE, loc, true));
decl = TREE_OPERAND (decl, 0);
}
}
for (int i = dims.length () - 1; i >= 0; i--)
decl = build_omp_array_section (loc, decl, dims[i].low_bound,
- dims[i].length);
+ dims[i].length, dims[i].stride);
}
else if (TREE_CODE (decl) == INDIRECT_REF)
{
@@ -14129,7 +14132,7 @@ c_parser_omp_variable_list (c_parser *parser,
STRIP_NOPS (decl);
decl = build_omp_array_section (loc, decl, integer_zero_node,
- integer_one_node);
+ integer_one_node, NULL_TREE);
}
else if (TREE_CODE (decl) == ARRAY_REF)
{
@@ -14138,7 +14141,8 @@ c_parser_omp_variable_list (c_parser *parser,
decl = TREE_OPERAND (decl, 0);
STRIP_NOPS (decl);
- decl = build_omp_array_section (loc, decl, idx, integer_one_node);
+ decl = build_omp_array_section (loc, decl, idx, integer_one_node,
+ NULL_TREE);
}
else if (TREE_CODE (decl) == NON_LVALUE_EXPR
|| CONVERT_EXPR_P (decl))
@@ -14293,7 +14297,8 @@ c_parser_omp_variable_list (c_parser *parser,
break;
}
- dims.safe_push (omp_dim (low_bound, length, loc, no_colon));
+ dims.safe_push (omp_dim (low_bound, length, NULL_TREE, loc,
+ no_colon));
}
if (t != error_mark_node)
@@ -14317,7 +14322,8 @@ c_parser_omp_variable_list (c_parser *parser,
for (unsigned i = 0; i < dims.length (); i++)
t = build_omp_array_section (clause_loc, t,
dims[i].low_bound,
- dims[i].length);
+ dims[i].length,
+ dims[i].stride);
}
if ((kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY)
@@ -765,7 +765,7 @@ extern tree composite_type (tree, tree);
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);
+extern tree build_omp_array_section (location_t, tree, tree, tree, tree);
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);
@@ -2036,6 +2036,8 @@ mark_exp_read (tree exp)
mark_exp_read (TREE_OPERAND (exp, 1));
if (TREE_OPERAND (exp, 2))
mark_exp_read (TREE_OPERAND (exp, 2));
+ if (TREE_OPERAND (exp, 3))
+ mark_exp_read (TREE_OPERAND (exp, 3));
break;
default:
break;
@@ -2921,7 +2923,8 @@ build_array_ref (location_t loc, tree array, tree index)
instead. */
tree
-build_omp_array_section (location_t loc, tree array, tree index, tree length)
+build_omp_array_section (location_t loc, tree array, tree index, tree length,
+ tree stride)
{
tree idxtype;
@@ -2958,7 +2961,8 @@ build_omp_array_section (location_t loc, tree array, tree index, tree length)
else
sectype = build_array_type (eltype, idxtype);
- return build3_loc (loc, OMP_ARRAY_SECTION, sectype, array, index, length);
+ return build4_loc (loc, OMP_ARRAY_SECTION, sectype, array, index, length,
+ stride);
}
@@ -13712,7 +13716,7 @@ 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)
{
- tree ret, low_bound, length, type;
+ tree ret, low_bound, length, stride, type;
bool openacc = (ort & C_ORT_ACC) != 0;
if (TREE_CODE (t) != OMP_ARRAY_SECTION)
{
@@ -13798,8 +13802,11 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
type = TREE_TYPE (ret);
low_bound = TREE_OPERAND (t, 1);
length = TREE_OPERAND (t, 2);
+ stride = TREE_OPERAND (t, 3);
- if (low_bound == error_mark_node || length == error_mark_node)
+ if (low_bound == error_mark_node
+ || length == error_mark_node
+ || stride == error_mark_node)
return error_mark_node;
if (low_bound && !INTEGRAL_TYPE_P (TREE_TYPE (low_bound)))
@@ -13816,6 +13823,13 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
length);
return error_mark_node;
}
+ if (stride && !INTEGRAL_TYPE_P (TREE_TYPE (stride)))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "stride %qE of array section does not have integral type",
+ stride);
+ return error_mark_node;
+ }
if (low_bound
&& TREE_CODE (low_bound) == INTEGER_CST
&& TYPE_PRECISION (TREE_TYPE (low_bound))
@@ -14032,7 +14046,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
d = TREE_OPERAND (d, 0))
{
tree d_length = TREE_OPERAND (d, 2);
- if (d_length == NULL_TREE || !integer_onep (d_length))
+ tree d_stride = TREE_OPERAND (d, 3);
+ if (d_length == NULL_TREE || !integer_onep (d_length)
+ || (d_stride && !integer_onep (d_stride)))
{
if (openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
{
@@ -530,6 +530,7 @@ cp_common_init_ts (void)
MARK_TS_EXP (OFFSET_REF);
MARK_TS_EXP (PSEUDO_DTOR_EXPR);
MARK_TS_EXP (REINTERPRET_CAST_EXPR);
+ MARK_TS_EXP (OMP_ARRAYSHAPE_CAST_EXPR);
MARK_TS_EXP (SCOPE_REF);
MARK_TS_EXP (STATIC_CAST_EXPR);
MARK_TS_EXP (STMT_EXPR);
@@ -256,6 +256,7 @@ DEFTREECODE (REINTERPRET_CAST_EXPR, "reinterpret_cast_expr", tcc_unary, 1)
DEFTREECODE (CONST_CAST_EXPR, "const_cast_expr", tcc_unary, 1)
DEFTREECODE (STATIC_CAST_EXPR, "static_cast_expr", tcc_unary, 1)
DEFTREECODE (DYNAMIC_CAST_EXPR, "dynamic_cast_expr", tcc_unary, 1)
+DEFTREECODE (OMP_ARRAYSHAPE_CAST_EXPR, "omp_arrayshape_cast_expr", tcc_unary, 1)
DEFTREECODE (IMPLICIT_CONV_EXPR, "implicit_conv_expr", tcc_unary, 1)
DEFTREECODE (DOTSTAR_EXPR, "dotstar_expr", tcc_expression, 2)
DEFTREECODE (TYPEID_EXPR, "typeid_expr", tcc_expression, 1)
@@ -504,6 +504,7 @@ extern GTY(()) tree cp_global_trees[CPTI_MAX];
OVL_LOOKUP_P (in OVERLOAD)
LOOKUP_FOUND_P (in RECORD_TYPE, UNION_TYPE, ENUMERAL_TYPE, NAMESPACE_DECL)
FNDECL_MANIFESTLY_CONST_EVALUATED (in FUNCTION_DECL)
+ DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (in DECLTYPE_TYPE)
5: IDENTIFIER_VIRTUAL_P (in IDENTIFIER_NODE)
FUNCTION_RVALUE_QUALIFIED (in FUNCTION_TYPE, METHOD_TYPE)
CALL_EXPR_REVERSE_ARGS (in CALL_EXPR, AGGR_INIT_EXPR)
@@ -4852,6 +4853,8 @@ get_vec_init_expr (tree t)
TREE_LANG_FLAG_2 (DECLTYPE_TYPE_CHECK (NODE))
#define DECLTYPE_FOR_REF_CAPTURE(NODE) \
TREE_LANG_FLAG_3 (DECLTYPE_TYPE_CHECK (NODE))
+#define DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST(NODE) \
+ TREE_LANG_FLAG_4 (DECLTYPE_TYPE_CHECK (NODE))
/* Nonzero for VAR_DECL and FUNCTION_DECL node means that `extern' was
specified in its declaration. This can also be set for an
@@ -6924,6 +6927,8 @@ extern tree cxx_comdat_group (tree);
extern bool cp_missing_noreturn_ok_p (tree);
extern bool is_direct_enum_init (tree, tree);
extern void initialize_artificial_var (tree, vec<constructor_elt, va_gc> *);
+extern tree cp_omp_create_arrayshape_type (location_t, tree,
+ vec<cp_expr> *);
extern tree check_var_type (tree, tree, location_t);
extern tree reshape_init (tree, tree, tsubst_flags_t);
extern tree next_aggregate_field (tree);
@@ -6957,7 +6962,8 @@ extern void grokclassfn (tree, tree,
enum overload_flags);
extern tree grok_array_decl (location_t, tree, tree,
vec<tree, va_gc> **, tsubst_flags_t);
-extern tree grok_omp_array_section (location_t, tree, tree, tree);
+extern tree grok_omp_array_section (location_t, tree, tree, tree,
+ tree);
extern tree delete_sanity (location_t, tree, tree, bool,
int, tsubst_flags_t);
extern tree check_classfn (tree, tree, tree);
@@ -7809,6 +7815,8 @@ extern tree cp_build_vec_convert (tree, location_t, tree,
tsubst_flags_t);
extern tree cp_build_bit_cast (location_t, tree, tree,
tsubst_flags_t);
+extern tree cp_build_omp_arrayshape_cast (location_t, tree, tree,
+ tsubst_flags_t);
extern void start_lambda_scope (tree decl);
extern void finish_lambda_scope (void);
extern void record_lambda_scope (tree lambda);
@@ -8061,7 +8069,8 @@ inline tree build_x_binary_op (const op_location_t &loc,
}
extern tree build_x_array_ref (location_t, tree, tree,
tsubst_flags_t);
-extern tree build_omp_array_section (location_t, tree, tree, tree);
+extern tree build_omp_array_section (location_t, tree, tree, tree,
+ tree);
extern tree build_x_unary_op (location_t,
enum tree_code, cp_expr,
tree, tsubst_flags_t);
@@ -11591,6 +11591,81 @@ create_array_type_for_decl (tree name, tree type, tree size, location_t loc)
return build_cplus_array_type (type, itype);
}
+/* Build an anonymous array of SIZE elements of ELTYPE. */
+
+static tree
+create_anon_array_type (location_t loc, tree eltype, tree size)
+{
+ if (eltype == error_mark_node || size == error_mark_node)
+ return error_mark_node;
+
+ tree itype = compute_array_index_type_loc (loc, NULL_TREE, size,
+ tf_warning_or_error);
+
+ if (type_uses_auto (eltype)
+ && variably_modified_type_p (itype, /*fn=*/NULL_TREE))
+ {
+ sorry_at (loc, "variable-length array of %<auto%>");
+ return error_mark_node;
+ }
+
+ return build_cplus_array_type (eltype, itype);
+}
+
+/* Derive an array type for an OpenMP array-shaping operator given EXPR, which
+ is an expression that might have array refs or array sections postfixed
+ (e.g. "ptr[0:3:2][3:4]"), and OMP_SHAPE_DIMS, a vector of dimensions. */
+
+tree
+cp_omp_create_arrayshape_type (location_t loc, tree expr,
+ vec<cp_expr> *omp_shape_dims)
+{
+ tree type, strip_sections = expr;
+
+ while (TREE_CODE (strip_sections) == OMP_ARRAY_SECTION
+ || TREE_CODE (strip_sections) == ARRAY_REF)
+ strip_sections = TREE_OPERAND (strip_sections, 0);
+
+ /* Determine the element type, either directly or by using
+ "decltype" of an expression representing an element to
+ figure it out later during template instantiation. */
+ if (type_dependent_expression_p (expr))
+ {
+ type = cxx_make_type (DECLTYPE_TYPE);
+
+ DECLTYPE_TYPE_EXPR (type)
+ = build_min_nt_loc (loc, INDIRECT_REF, strip_sections);
+ DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (type) = true;
+ SET_TYPE_STRUCTURAL_EQUALITY (type);
+ }
+ else
+ {
+ 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;
+ cp_expr dim;
+ FOR_EACH_VEC_ELT_REVERSE (*omp_shape_dims, i, dim)
+ {
+ if (!type_dependent_expression_p (dim))
+ dim = fold_convert (sizetype, dim);
+ type = create_anon_array_type (loc, type, dim);
+ }
+
+ return type;
+}
+
/* Returns the smallest location that is not UNKNOWN_LOCATION. */
static location_t
@@ -620,43 +620,49 @@ grok_array_decl (location_t loc, tree array_expr, tree index_exp,
tree
grok_omp_array_section (location_t loc, tree array_expr, tree index,
- tree length)
+ tree length, tree stride)
{
tree orig_array_expr = array_expr;
tree orig_index = index;
tree orig_length = length;
+ tree orig_stride = stride;
if (error_operand_p (array_expr)
|| error_operand_p (index)
- || error_operand_p (length))
+ || error_operand_p (length)
+ || error_operand_p (stride))
return error_mark_node;
if (processing_template_decl)
{
if (type_dependent_expression_p (array_expr)
|| type_dependent_expression_p (index)
- || type_dependent_expression_p (length))
+ || type_dependent_expression_p (length)
+ || type_dependent_expression_p (stride))
return build_min_nt_loc (loc, OMP_ARRAY_SECTION, array_expr, index,
- length);
+ length, stride);
array_expr = build_non_dependent_expr (array_expr);
if (index)
index = build_non_dependent_expr (index);
if (length)
length = build_non_dependent_expr (length);
+ if (stride)
+ stride = build_non_dependent_expr (stride);
}
index = fold_non_dependent_expr (index);
length = fold_non_dependent_expr (length);
+ stride = fold_non_dependent_expr (stride);
/* NOTE: We can pass through invalidly-typed index/length fields
here (e.g. if the user tries to use a floating-point index/length).
This is diagnosed later in semantics.cc:handle_omp_array_sections_1. */
- tree expr = build_omp_array_section (loc, array_expr, index, length);
+ tree expr = build_omp_array_section (loc, array_expr, index, length, stride);
if (processing_template_decl)
expr = build_min_non_dep (OMP_ARRAY_SECTION, expr, orig_array_expr,
- orig_index, orig_length);
+ orig_index, orig_length, orig_stride);
return expr;
}
@@ -2642,6 +2648,7 @@ min_vis_expr_r (tree *tp, int */*walk_subtrees*/, void *data)
case REINTERPRET_CAST_EXPR:
case CONST_CAST_EXPR:
case DYNAMIC_CAST_EXPR:
+ case OMP_ARRAYSHAPE_CAST_EXPR:
case NEW_EXPR:
case CONSTRUCTOR:
case LAMBDA_EXPR:
@@ -2528,6 +2528,11 @@ dump_expr (cxx_pretty_printer *pp, tree t, int flags)
dump_expr (pp, TREE_OPERAND (t, 1), flags);
pp_colon (pp);
dump_expr (pp, TREE_OPERAND (t, 2), flags);
+ if (TREE_OPERAND (t, 3))
+ {
+ pp_colon (pp);
+ dump_expr (pp, TREE_OPERAND (t, 3), flags);
+ }
pp_cxx_right_bracket (pp);
break;
@@ -3611,6 +3611,7 @@ write_expression (tree expr)
case REINTERPRET_CAST_EXPR:
case STATIC_CAST_EXPR:
case CONST_CAST_EXPR:
+ case OMP_ARRAYSHAPE_CAST_EXPR:
write_type (TREE_TYPE (expr));
write_expression (TREE_OPERAND (expr, 0));
break;
@@ -134,6 +134,7 @@ DEF_OPERATOR (NULL, DYNAMIC_CAST_EXPR, "dc", OVL_OP_FLAG_UNARY)
DEF_OPERATOR (NULL, REINTERPRET_CAST_EXPR, "rc", OVL_OP_FLAG_UNARY)
DEF_OPERATOR (NULL, CONST_CAST_EXPR, "cc", OVL_OP_FLAG_UNARY)
DEF_OPERATOR (NULL, STATIC_CAST_EXPR, "sc", OVL_OP_FLAG_UNARY)
+DEF_OPERATOR (NULL, OMP_ARRAYSHAPE_CAST_EXPR, "oc", OVL_OP_FLAG_UNARY)
DEF_OPERATOR (NULL, SCOPE_REF, "sr", OVL_OP_FLAG_NONE)
DEF_OPERATOR (NULL, EXPR_PACK_EXPANSION, "sp", OVL_OP_FLAG_NONE)
DEF_OPERATOR (NULL, UNARY_LEFT_FOLD_EXPR, "fl", OVL_OP_FLAG_NONE)
@@ -4380,6 +4380,12 @@ cp_parser_new (cp_lexer *lexer)
/* Disallow OpenMP array sections in expressions. */
parser->omp_array_section_p = false;
+ /* Disallow OpenMP array-shaping operator in expressions. */
+ parser->omp_array_shaping_op_p = false;
+
+ /* We don't have an OpenMP array shape here. */
+ parser->omp_has_array_shape_p = false;
+
/* Not declaring an implicit function template. */
parser->auto_is_implicit_function_template_parm_p = false;
parser->fully_implicit_function_template_p = false;
@@ -5365,6 +5371,7 @@ cp_parser_statement_expr (cp_parser *parser)
{
cp_token_position start = cp_parser_start_tentative_firewall (parser);
auto oas = make_temp_override (parser->omp_array_section_p, false);
+ auto aso = make_temp_override (parser->omp_array_shaping_op_p, false);
/* Consume the '('. */
location_t start_loc = cp_lexer_peek_token (parser->lexer)->location;
@@ -8290,7 +8297,7 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
&& cp_lexer_next_token_is (parser->lexer, CPP_COLON))
{
cp_lexer_consume_token (parser->lexer);
- tree length = NULL_TREE;
+ tree length = NULL_TREE, stride = NULL_TREE;
if (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_SQUARE))
{
if (cxx_dialect >= cxx23)
@@ -8323,9 +8330,23 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
/*warn_comma_p=*/warn_comma_subscript);
}
+ if (cp_lexer_next_token_is (parser->lexer, CPP_COLON))
+ {
+ cp_lexer_consume_token (parser->lexer);
+ /* We could check for C++-23 multidimensional/comma-separated
+ subscripts here, or not bother. */
+ if (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_SQUARE))
+ stride
+ = cp_parser_expression (parser, NULL, /*cast_p=*/false,
+ /*decltype_p=*/false,
+ /*warn_comma_p=*/warn_comma_subscript);
+ }
+
parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p;
- if (index == error_mark_node || length == error_mark_node)
+ if (index == error_mark_node
+ || length == error_mark_node
+ || stride == error_mark_node)
{
cp_parser_skip_to_closing_square_bracket (parser);
return error_mark_node;
@@ -8334,7 +8355,7 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
cp_parser_require (parser, CPP_CLOSE_SQUARE, RT_CLOSE_SQUARE);
return grok_omp_array_section (input_location, postfix_expression, index,
- length);
+ length, stride);
}
parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p;
@@ -8342,11 +8363,23 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
/* Look for the closing `]'. */
cp_parser_require (parser, CPP_CLOSE_SQUARE, RT_CLOSE_SQUARE);
- /* Build the ARRAY_REF. */
- postfix_expression = grok_array_decl (loc, postfix_expression,
- index, &expression_list,
- tf_warning_or_error
- | (decltype_p ? tf_decltype : 0));
+ if (parser->omp_has_array_shape_p
+ && (expression_list.get () == NULL
+ || vec_safe_length (expression_list) == 1))
+ /* 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 cp_parser_omp_var_list_no_open). */
+ return grok_omp_array_section (loc, postfix_expression, index,
+ size_one_node, NULL_TREE);
+ else
+ /* Build the ARRAY_REF. */
+ postfix_expression = grok_array_decl (loc, postfix_expression,
+ index, &expression_list,
+ tf_warning_or_error
+ | (decltype_p ? tf_decltype : 0));
/* When not doing offsetof, array references are not permitted in
constant-expressions. */
@@ -8668,6 +8701,7 @@ cp_parser_parenthesized_expression_list (cp_parser* parser,
vec<tree, va_gc> *expression_list;
bool saved_greater_than_is_operator_p;
bool saved_omp_array_section_p;
+ bool saved_omp_array_shaping_op_p;
/* Assume all the expressions will be constant. */
if (non_constant_p)
@@ -8686,7 +8720,9 @@ cp_parser_parenthesized_expression_list (cp_parser* parser,
parser->greater_than_is_operator_p = true;
saved_omp_array_section_p = parser->omp_array_section_p;
+ saved_omp_array_shaping_op_p = parser->omp_array_shaping_op_p;
parser->omp_array_section_p = false;
+ parser->omp_array_shaping_op_p = false;
cp_expr expr (NULL_TREE);
@@ -8753,6 +8789,7 @@ cp_parser_parenthesized_expression_list (cp_parser* parser,
parser->greater_than_is_operator_p
= saved_greater_than_is_operator_p;
parser->omp_array_section_p = saved_omp_array_section_p;
+ parser->omp_array_shaping_op_p = saved_omp_array_shaping_op_p;
return NULL;
}
}
@@ -8760,6 +8797,7 @@ cp_parser_parenthesized_expression_list (cp_parser* parser,
parser->greater_than_is_operator_p
= saved_greater_than_is_operator_p;
parser->omp_array_section_p = saved_omp_array_section_p;
+ parser->omp_array_shaping_op_p = saved_omp_array_shaping_op_p;
return expression_list;
}
@@ -10028,6 +10066,8 @@ cp_parser_cast_expression (cp_parser *parser, bool address_p, bool cast_p,
cp_expr expr (NULL_TREE);
int cast_expression = 0;
const char *saved_message;
+ auto_vec<cp_expr, 4> omp_shape_dims;
+ bool omp_array_shape_p = false;
/* There's no way to know yet whether or not this is a cast.
For example, `(int (3))' is a unary-expression, while `(int)
@@ -10097,6 +10137,28 @@ cp_parser_cast_expression (cp_parser *parser, bool address_p, bool cast_p,
that the call to cp_parser_error_occurred below returns true. */
if (!cast_expression)
cp_parser_simulate_error (parser);
+ else if (parser->omp_array_shaping_op_p
+ && cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE))
+ {
+ auto oas = make_temp_override (parser->omp_array_section_p, false);
+ auto aso = make_temp_override (parser->omp_array_shaping_op_p, false);
+
+ while (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE))
+ {
+ cp_lexer_consume_token (parser->lexer);
+ cp_expr e = cp_parser_expression (parser);
+ if (e.get_value () == error_mark_node)
+ break;
+ omp_shape_dims.safe_push (e);
+ if (!cp_parser_require (parser, CPP_CLOSE_SQUARE,
+ RT_CLOSE_SQUARE))
+ break;
+ }
+ cp_token *close_paren = parens.require_close (parser);
+ if (close_paren)
+ close_paren_loc = close_paren->location;
+ omp_array_shape_p = true;
+ }
else
{
bool saved_in_type_id_in_expr_p = parser->in_type_id_in_expr_p;
@@ -10118,6 +10180,10 @@ cp_parser_cast_expression (cp_parser *parser, bool address_p, bool cast_p,
function returning T. */
if (!cp_parser_error_occurred (parser))
{
+ auto aso = make_temp_override (parser->omp_array_shaping_op_p, false);
+ auto as = make_temp_override (parser->omp_has_array_shape_p,
+ omp_array_shape_p);
+
/* Only commit if the cast-expression doesn't start with
'++', '--', or '[' in C++11. */
if (cast_expression > 0)
@@ -10131,6 +10197,24 @@ cp_parser_cast_expression (cp_parser *parser, bool address_p, bool cast_p,
if (cp_parser_parse_definitely (parser))
{
+ if (omp_array_shape_p)
+ {
+ location_t cast_loc = make_location (open_paren_loc,
+ open_paren_loc,
+ expr.get_finish ());
+
+ type = cp_omp_create_arrayshape_type (cast_loc, expr,
+ &omp_shape_dims);
+
+ /* Things rapidly get worse below if we carry on from here
+ with an erroneous type... */
+ if (error_operand_p (type))
+ return error_mark_node;
+
+ return cp_build_omp_arrayshape_cast (cast_loc, type, expr,
+ tf_warning_or_error);
+ }
+
/* Warn about old-style casts, if so requested. */
if (warn_old_style_cast
&& !in_system_header_at (input_location)
@@ -11258,6 +11342,7 @@ cp_parser_lambda_expression (cp_parser* parser)
bool auto_is_implicit_function_template_parm_p
= parser->auto_is_implicit_function_template_parm_p;
bool saved_omp_array_section_p = parser->omp_array_section_p;
+ bool saved_omp_array_shaping_op_p = parser->omp_array_shaping_op_p;
parser->num_template_parameter_lists = 0;
parser->in_statement = 0;
@@ -11267,6 +11352,7 @@ cp_parser_lambda_expression (cp_parser* parser)
parser->implicit_template_scope = 0;
parser->auto_is_implicit_function_template_parm_p = false;
parser->omp_array_section_p = false;
+ parser->omp_array_shaping_op_p = false;
/* The body of a lambda in a discarded statement is not discarded. */
bool discarded = in_discarded_stmt;
@@ -11318,6 +11404,7 @@ cp_parser_lambda_expression (cp_parser* parser)
parser->auto_is_implicit_function_template_parm_p
= auto_is_implicit_function_template_parm_p;
parser->omp_array_section_p = saved_omp_array_section_p;
+ parser->omp_array_shaping_op_p = saved_omp_array_shaping_op_p;
}
/* This field is only used during parsing of the lambda. */
@@ -25745,6 +25832,7 @@ cp_parser_braced_list (cp_parser* parser, bool* non_constant_p)
tree initializer;
location_t start_loc = cp_lexer_peek_token (parser->lexer)->location;
auto oas = make_temp_override (parser->omp_array_section_p, false);
+ auto aso = make_temp_override (parser->omp_array_shaping_op_p, false);
/* Consume the `{' token. */
matching_braces braces;
@@ -37643,11 +37731,11 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code,
struct omp_dim
{
- tree low_bound, length;
+ tree low_bound, length, stride;
location_t loc;
bool no_colon;
- omp_dim (tree lb, tree len, location_t lo, bool nc)
- : low_bound (lb), length (len), loc (lo), no_colon (nc) {}
+ omp_dim (tree lb, tree len, tree str, location_t lo, bool nc)
+ : low_bound (lb), length (len), stride (str), loc (lo), no_colon (nc) {}
};
static tree
@@ -37680,10 +37768,22 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
|| kind == OMP_CLAUSE_FROM))
{
auto s = make_temp_override (parser->omp_array_section_p, true);
+ auto o = make_temp_override (parser->omp_array_shaping_op_p,
+ (kind == OMP_CLAUSE_TO
+ || kind == OMP_CLAUSE_FROM));
+ tree reshaped_to = NULL_TREE;
token = cp_lexer_peek_token (parser->lexer);
location_t loc = token->location;
decl = cp_parser_assignment_expression (parser);
+ if ((TREE_CODE (decl) == VIEW_CONVERT_EXPR
+ && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+ || TREE_CODE (decl) == OMP_ARRAYSHAPE_CAST_EXPR)
+ {
+ 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
OMP_ARRAY_SECTION nodes before it is processed by
@@ -37694,49 +37794,159 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
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);
tree length = TREE_OPERAND (decl, 2);
- dims.safe_push (omp_dim (low_bound, length, loc, false));
+ tree stride = TREE_OPERAND (decl, 3);
+ 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 (REFERENCE_REF_P (decl))
break;
+ if (maybe_ptr_based_noncontig_update && sections == 0)
+ break;
+
if (TREE_CODE (decl) == COMPOUND_EXPR)
{
decl = TREE_OPERAND (decl, 1);
STRIP_NOPS (decl);
+ continue;
}
- 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, loc, true));
+ integer_one_node, NULL_TREE, loc,
+ true));
decl = TREE_OPERAND (decl, 0);
}
else /* ARRAY_REF. */
{
tree index = TREE_OPERAND (decl, 1);
- dims.safe_push (omp_dim (index, integer_one_node, loc,
- true));
+ 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. */
+ if (type_dependent_expression_p (decl))
+ decl = build_min_nt_loc (loc, INDIRECT_REF, decl);
+ else
+ {
+ /* We're interested in the reference target. */
+ decl = convert_from_reference (decl);
+ decl = cp_build_fold_indirect_ref (decl);
+ }
+ decl
+ = cp_build_omp_arrayshape_cast (loc, reshaped_to, decl,
+ tf_warning_or_error);
+ }
+ }
/* Bare references have their own special handling, so remove
the explicit dereference added by convert_from_reference. */
- if (REFERENCE_REF_P (decl))
+ else if (REFERENCE_REF_P (decl))
decl = TREE_OPERAND (decl, 0);
for (int i = dims.length () - 1; i >= 0; i--)
decl = grok_omp_array_section (loc, decl, dims[i].low_bound,
- dims[i].length);
+ dims[i].length, dims[i].stride);
}
else if (TREE_CODE (decl) == INDIRECT_REF)
{
@@ -37753,7 +37963,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
"foo[0:1]". */
if (!ref_p)
decl = grok_omp_array_section (loc, decl, integer_zero_node,
- integer_one_node);
+ integer_one_node, NULL_TREE);
}
else if (TREE_CODE (decl) == ARRAY_REF)
{
@@ -37762,7 +37972,16 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
decl = TREE_OPERAND (decl, 0);
STRIP_NOPS (decl);
- decl = grok_omp_array_section (loc, decl, idx, integer_one_node);
+ decl = grok_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 = array_type_nelts_total (reshaped_to);
+ decl = grok_omp_array_section (loc, decl, size_zero_node, elems,
+ NULL_TREE);
}
else if (TREE_CODE (decl) == NON_LVALUE_EXPR
|| CONVERT_EXPR_P (decl))
@@ -37927,7 +38146,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
goto skip_comma;
}
- dims.safe_push (omp_dim (low_bound, length, loc, no_colon));
+ dims.safe_push (omp_dim (low_bound, length, NULL_TREE, loc,
+ no_colon));
}
if ((kind == OMP_CLAUSE_MAP
@@ -37949,7 +38169,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
for (unsigned i = 0; i < dims.length (); i++)
decl = build_omp_array_section (input_location, decl,
dims[i].low_bound,
- dims[i].length);
+ dims[i].length,
+ dims[i].stride);
break;
default:
break;
@@ -37962,6 +38183,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
&& cp_parser_simulate_error (parser))
{
depend_lvalue:
+ auto o = make_temp_override (parser->omp_array_shaping_op_p,
+ true);
cp_parser_abort_tentative_parse (parser);
decl = cp_parser_assignment_expression (parser, NULL,
false, false);
@@ -46698,8 +46921,38 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok,
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK,
"#pragma omp target update", pragma_tok);
- 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 (pragma_tok->location,
"%<#pragma omp target update%> must contain at least one "
@@ -410,6 +410,13 @@ struct GTY(()) cp_parser {
/* TRUE if an OpenMP array section is allowed. */
bool omp_array_section_p;
+ /* TRUE if an OpenMP array-shaping operator is allowed. */
+ bool omp_array_shaping_op_p;
+
+ /* TRUE if we are parsing an expression with an OpenMP array-shaping
+ operator. */
+ bool omp_has_array_shape_p;
+
/* Tracks the function's template parameter list when declaring a function
using generic type parameters. This is either a new chain in the case of a
fully implicit function template or an extension of the function's existing
@@ -16700,6 +16700,10 @@ tsubst (tree t, tree args, tsubst_flags_t complain, tree in_decl)
member access. */
id = false;
type = finish_decltype_type (type, id, complain);
+
+ if (DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (t)
+ && TYPE_REF_P (type))
+ type = TREE_TYPE (type);
}
return cp_build_qualified_type (type,
cp_type_quals (t)
@@ -17531,6 +17535,7 @@ tsubst_copy (tree t, tree args, tsubst_flags_t complain, tree in_decl)
case STATIC_CAST_EXPR:
case DYNAMIC_CAST_EXPR:
case IMPLICIT_CONV_EXPR:
+ case OMP_ARRAYSHAPE_CAST_EXPR:
CASE_CONVERT:
{
tsubst_flags_t tcomplain = complain;
@@ -17756,12 +17761,14 @@ tsubst_copy (tree t, tree args, tsubst_flags_t complain, tree in_decl)
case OMP_ARRAY_SECTION:
{
tree op0 = tsubst_copy (TREE_OPERAND (t, 0), args, complain, in_decl);
- tree op1 = NULL_TREE, op2 = NULL_TREE;
+ tree op1 = NULL_TREE, op2 = NULL_TREE, op3 = NULL_TREE;
if (TREE_OPERAND (t, 1))
op1 = tsubst_copy (TREE_OPERAND (t, 1), args, complain, in_decl);
if (TREE_OPERAND (t, 2))
op2 = tsubst_copy (TREE_OPERAND (t, 2), args, complain, in_decl);
- return build_nt (OMP_ARRAY_SECTION, op0, op1, op2);
+ if (TREE_OPERAND (t, 3))
+ op3 = tsubst_copy (TREE_OPERAND (t, 3), args, complain, in_decl);
+ return build_nt (OMP_ARRAY_SECTION, op0, op1, op2, op3);
}
case CALL_EXPR:
@@ -18038,14 +18045,17 @@ tsubst_omp_clause_decl (tree decl, tree args, tsubst_flags_t complain,
= tsubst_expr (TREE_OPERAND (decl, 1), args, complain, in_decl);
tree length = tsubst_expr (TREE_OPERAND (decl, 2), args, complain,
in_decl);
+ tree stride = tsubst_expr (TREE_OPERAND (decl, 3), args, complain,
+ in_decl);
tree base = tsubst_omp_clause_decl (TREE_OPERAND (decl, 0), args,
complain, in_decl, NULL);
if (TREE_OPERAND (decl, 0) == base
&& TREE_OPERAND (decl, 1) == low_bound
- && TREE_OPERAND (decl, 2) == length)
+ && TREE_OPERAND (decl, 2) == length
+ && TREE_OPERAND (decl, 3) == stride)
return decl;
- tree ret = build3 (OMP_ARRAY_SECTION, TREE_TYPE (base), base, low_bound,
- length);
+ tree ret = build4 (OMP_ARRAY_SECTION, TREE_TYPE (base), base, low_bound,
+ length, stride);
return ret;
}
tree ret = tsubst_expr (decl, args, complain, in_decl);
@@ -20685,6 +20695,14 @@ tsubst_copy_and_build (tree t,
RETURN (cp_build_bit_cast (EXPR_LOCATION (t), type, op0, complain));
}
+ case OMP_ARRAYSHAPE_CAST_EXPR:
+ {
+ tree type = tsubst (TREE_TYPE (t), args, complain, in_decl);
+ tree op0 = RECUR (TREE_OPERAND (t, 0));
+ RETURN (cp_build_omp_arrayshape_cast (EXPR_LOCATION (t), type, op0,
+ complain));
+ }
+
case POSTDECREMENT_EXPR:
case POSTINCREMENT_EXPR:
op1 = tsubst_non_call_postfix_expression (TREE_OPERAND (t, 0),
@@ -20851,7 +20869,7 @@ tsubst_copy_and_build (tree t,
case OMP_ARRAY_SECTION:
{
tree op0 = RECUR (TREE_OPERAND (t, 0));
- tree op1 = NULL_TREE, op2 = NULL_TREE;
+ tree op1 = NULL_TREE, op2 = NULL_TREE, op3 = NULL_TREE;
if (op0 == error_mark_node)
RETURN (error_mark_node);
if (TREE_OPERAND (t, 1))
@@ -20866,7 +20884,14 @@ tsubst_copy_and_build (tree t,
if (op2 == error_mark_node)
RETURN (error_mark_node);
}
- RETURN (build_omp_array_section (EXPR_LOCATION (t), op0, op1, op2));
+ if (TREE_OPERAND (t, 3))
+ {
+ op3 = RECUR (TREE_OPERAND (t, 3));
+ if (op3 == error_mark_node)
+ RETURN (error_mark_node);
+ }
+ RETURN (build_omp_array_section (EXPR_LOCATION (t), op0, op1, op2,
+ op3));
}
case SIZEOF_EXPR:
@@ -5180,9 +5180,10 @@ public:
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, type;
+ tree ret, low_bound, length, stride, type;
bool openacc = (ort & C_ORT_ACC) != 0;
if (TREE_CODE (t) != OMP_ARRAY_SECTION)
{
@@ -5245,18 +5246,25 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
TREE_OPERAND (t, 0) = omp_privatize_field (TREE_OPERAND (t, 0), false);
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);
if ((low_bound && type_dependent_expression_p (low_bound))
- || (length && type_dependent_expression_p (length)))
+ || (length && type_dependent_expression_p (length))
+ || (stride && type_dependent_expression_p (stride)))
return NULL_TREE;
- if (low_bound == error_mark_node || length == error_mark_node)
+ if (low_bound == error_mark_node
+ || length == error_mark_node
+ || stride == error_mark_node)
return error_mark_node;
if (low_bound && !INTEGRAL_TYPE_P (TREE_TYPE (low_bound)))
@@ -5273,10 +5281,19 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
length);
return error_mark_node;
}
+ if (stride && !INTEGRAL_TYPE_P (TREE_TYPE (stride)))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "stride %qE of array section does not have integral type",
+ stride);
+ return error_mark_node;
+ }
if (low_bound)
low_bound = mark_rvalue_use (low_bound);
if (length)
length = mark_rvalue_use (length);
+ if (stride)
+ stride = mark_rvalue_use (stride);
/* We need to reduce to real constant-values for checks below. */
if (length)
STRIP_NOPS (length);
@@ -5286,6 +5303,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
length = fold_simple (length);
if (low_bound)
low_bound = fold_simple (low_bound);
+ if (stride)
+ stride = fold_simple (stride);
if (low_bound
&& TREE_CODE (low_bound) == INTEGER_CST
&& TYPE_PRECISION (TREE_TYPE (low_bound))
@@ -5296,9 +5315,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))
@@ -5417,12 +5442,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)
+ 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)
@@ -5430,7 +5472,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))
{
@@ -5500,7 +5542,10 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
d = TREE_OPERAND (d, 0))
{
tree d_length = TREE_OPERAND (d, 2);
- if (d_length == NULL_TREE || !integer_onep (d_length))
+ tree d_stride = TREE_OPERAND (d, 3);
+ if (d_length == NULL_TREE
+ || !integer_onep (d_length)
+ || (d_stride && !integer_onep (d_stride)))
{
if (openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
{
@@ -5520,10 +5565,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;
+ }
}
}
}
@@ -5535,7 +5585,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 = cp_save_expr (low_bound);
if (lb != low_bound)
@@ -5554,15 +5604,45 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TASK_REDUCTION);
- ret = grok_array_decl (OMP_CLAUSE_LOCATION (c), ret, low_bound, NULL,
- tf_warning_or_error);
+ /* 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 = grok_omp_array_section (OMP_CLAUSE_LOCATION (c), ret, low_bound,
+ length, stride);
+ else
+ ret = grok_array_decl (OMP_CLAUSE_LOCATION (c), ret, low_bound, NULL,
+ tf_warning_or_error);
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 grok_array_decl (loc, ret, low_bound, NULL, tf_warning_or_error);
+ }
+
+ 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. PC points to the clause to be processed, and
+ *PNEXT to the last mapping node created, if passed as non-NULL. */
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;
@@ -5578,7 +5658,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)
@@ -5620,6 +5700,8 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
if (processing_template_decl && maybe_zero_len)
return false;
+ bool higher_discontiguous = false;
+
for (i = num, t = OMP_CLAUSE_DECL (c); i > 0;
t = TREE_OPERAND (t, 0))
{
@@ -5627,6 +5709,7 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
tree low_bound = TREE_OPERAND (t, 1);
tree length = TREE_OPERAND (t, 2);
+ tree stride = TREE_OPERAND (t, 3);
if (length)
STRIP_NOPS (length);
@@ -5644,6 +5727,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;
@@ -5653,10 +5741,50 @@ 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])
@@ -5670,12 +5798,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 (!processing_template_decl
@@ -5792,6 +5925,9 @@ 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);
OMP_CLAUSE_DECL (c) = first;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
return false;
@@ -5799,9 +5935,6 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
if (TREE_CODE (t) == FIELD_DECL)
t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
- return false;
-
if (TREE_CODE (first) == INDIRECT_REF)
{
/* Detect and skip adding extra nodes for pointer-to-member
@@ -5828,6 +5961,10 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
}
}
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ && !(discontiguous && *discontiguous == 2))
+ return false;
+
/* FIRST represents the first item of data that we are mapping.
E.g. if we're mapping an array, FIRST might resemble
"foo.bar.myarray[0]". */
@@ -5846,7 +5983,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
@@ -6988,7 +7126,7 @@ 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;
@@ -8136,7 +8274,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
- if (handle_omp_array_sections (pc, NULL, ort))
+ int discontiguous = 1;
+ if (handle_omp_array_sections (pc, NULL, ort, &discontiguous))
remove = true;
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
&& (OMP_CLAUSE_DEPEND_KIND (c)
@@ -8291,6 +8430,9 @@ 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:
@@ -8305,8 +8447,11 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
grp_start_p = pc;
grp_sentinel = OMP_CLAUSE_CHAIN (c);
+ int discontiguous
+ = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM);
tree *pnext = NULL;
- if (handle_omp_array_sections (pc, &pnext, ort))
+ if (handle_omp_array_sections (pc, &pnext, ort, &discontiguous))
remove = true;
else
{
@@ -8897,7 +9042,7 @@ 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
{
@@ -13037,4 +13182,43 @@ cp_build_bit_cast (location_t loc, tree type, tree arg,
return ret;
}
+/* Build an OpenMP array-shape cast of ARG to TYPE. */
+
+tree
+cp_build_omp_arrayshape_cast (location_t loc, tree type, tree arg,
+ tsubst_flags_t complain)
+{
+ if (error_operand_p (type))
+ return error_mark_node;
+
+ if (!dependent_type_p (type)
+ && !complete_type_or_maybe_complain (type, NULL_TREE, complain))
+ return error_mark_node;
+
+ if (error_operand_p (arg))
+ return error_mark_node;
+
+ if (!type_dependent_expression_p (arg) && !dependent_type_p (type))
+ {
+ if (!trivially_copyable_p (TREE_TYPE (arg)))
+ {
+ error_at (cp_expr_loc_or_loc (arg, loc),
+ "OpenMP array shape source type %qT "
+ "is not trivially copyable", TREE_TYPE (arg));
+ return error_mark_node;
+ }
+
+ /* A pointer to multi-dimensional array conversion isn't normally
+ allowed, but we force it here for array shape operators by creating
+ the node directly. We also want to avoid any overloaded conversions
+ the user might have defined, not that there are likely to be any. */
+ return build1_loc (loc, VIEW_CONVERT_EXPR, type, arg);
+ }
+
+ tree ret = build_min (OMP_ARRAYSHAPE_CAST_EXPR, type, arg);
+ SET_EXPR_LOCATION (ret, loc);
+
+ return ret;
+}
+
#include "gt-cp-semantics.h"
@@ -1625,6 +1625,9 @@ structural_comptypes (tree t1, tree t2, int strict)
return false;
if (DECLTYPE_FOR_LAMBDA_PROXY (t1) != DECLTYPE_FOR_LAMBDA_PROXY (t2))
return false;
+ if (DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (t1)
+ != DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (t2))
+ return false;
if (!cp_tree_equal (DECLTYPE_TYPE_EXPR (t1), DECLTYPE_TYPE_EXPR (t2)))
return false;
break;
@@ -4793,7 +4796,7 @@ build_x_array_ref (location_t loc, tree arg1, tree arg2,
tree
build_omp_array_section (location_t loc, tree array_expr, tree index,
- tree length)
+ tree length, tree stride)
{
tree idxtype;
@@ -4832,8 +4835,8 @@ build_omp_array_section (location_t loc, tree array_expr, tree index,
else
sectype = build_array_type (eltype, idxtype);
- return build3_loc (loc, OMP_ARRAY_SECTION, sectype, array_expr, index,
- length);
+ return build4_loc (loc, OMP_ARRAY_SECTION, sectype, array_expr, index,
+ length, stride);
}
/* Return whether OP is an expression of enum type cast to integer
@@ -8151,6 +8154,9 @@ check_for_casting_away_constness (location_t loc, tree src_type,
src_type, dest_type);
return true;
+ case OMP_ARRAYSHAPE_CAST_EXPR:
+ return true;
+
default:
gcc_unreachable();
}
@@ -9307,6 +9307,19 @@ omp_group_last (tree *start_p)
grp_last_p = &OMP_CLAUSE_CHAIN (c);
break;
+ case GOMP_MAP_TO_GRID:
+ case GOMP_MAP_FROM_GRID:
+ while (nc
+ && 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))
+ {
+ grp_last_p = &OMP_CLAUSE_CHAIN (c);
+ c = nc;
+ nc = OMP_CLAUSE_CHAIN (c);
+ }
+ break;
+
case GOMP_MAP_STRUCT:
case GOMP_MAP_STRUCT_UNORD:
{
@@ -9455,6 +9468,10 @@ omp_group_base (omp_mapping_group *grp, unsigned int *chained,
internal_error ("unexpected mapping node");
return error_mark_node;
+ case GOMP_MAP_TO_GRID:
+ case GOMP_MAP_FROM_GRID:
+ return *grp->grp_start;
+
case GOMP_MAP_ATTACH:
case GOMP_MAP_DETACH:
node = OMP_CLAUSE_CHAIN (node);
@@ -14396,7 +14413,9 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
}
if (remove)
break;
- if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
+ if (OMP_CLAUSE_SIZE (c) == NULL_TREE
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_GRID_DIM
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_GRID_STRIDE)
OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
: TYPE_SIZE_UNIT (TREE_TYPE (decl));
gimplify_omp_ctxp = ctx->outer_context;
@@ -14483,6 +14502,20 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
is_gimple_lvalue, fb_lvalue) == GS_ERROR)
remove = true;
}
+ else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE)
+ {
+ /* The OMP_CLAUSE_DECL for GRID_DIM/GRID_STRIDE isn't necessarily
+ an lvalue -- e.g. it might be a constant. So handle it
+ specially here. */
+ if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
+ is_gimple_val, fb_rvalue) == GS_ERROR)
+ {
+ gimplify_omp_ctxp = ctx;
+ remove = true;
+ }
+ break;
+ }
else if (!DECL_P (decl))
{
if ((ctx->region_type & ORT_TARGET) != 0
@@ -14575,8 +14608,13 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
gimplify_omp_ctxp = ctx->outer_context;
if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
- fb_lvalue) == GS_ERROR)
- remove = true;
+ fb_lvalue | fb_mayfail) == GS_ERROR)
+ {
+ sorry_at (OMP_CLAUSE_LOCATION (c),
+ "unsupported map expression %qE",
+ OMP_CLAUSE_DECL (c));
+ remove = true;
+ }
gimplify_omp_ctxp = ctx;
break;
}
@@ -3702,6 +3702,32 @@ omp_parse_pointer (tree *expr0, bool *has_offset)
return false;
}
+static bool
+omp_parse_noncontiguous_array (tree *expr0)
+{
+ tree expr = *expr0;
+ bool noncontig = false;
+
+ while (TREE_CODE (expr) == OMP_ARRAY_SECTION
+ || TREE_CODE (expr) == ARRAY_REF)
+ {
+ /* Contiguous arrays use ARRAY_REF. By the time we reach here,
+ OMP_ARRAY_SECTION is only used for noncontiguous arrays. */
+ if (TREE_CODE (expr) == OMP_ARRAY_SECTION)
+ noncontig = true;
+
+ expr = TREE_OPERAND (expr, 0);
+ }
+
+ if (noncontig)
+ {
+ *expr0 = expr;
+ return true;
+ }
+
+ return false;
+}
+
static bool
omp_parse_access_method (tree *expr0, enum access_method_kinds *kind)
{
@@ -3710,6 +3736,13 @@ omp_parse_access_method (tree *expr0, enum access_method_kinds *kind)
if (omp_parse_ref (&expr))
*kind = ACCESS_REF;
+ else if (omp_parse_noncontiguous_array (&expr))
+ {
+ if (omp_parse_ref (&expr))
+ *kind = ACCESS_NONCONTIG_REF_TO_ARRAY;
+ else
+ *kind = ACCESS_NONCONTIG_ARRAY;
+ }
else if (omp_parse_pointer (&expr, &has_offset))
{
if (omp_parse_ref (&expr))
@@ -3783,6 +3816,14 @@ omp_parse_structure_base (vec<omp_addr_token *> &addr_tokens,
return true;
}
+ if (TREE_CODE (expr) == VIEW_CONVERT_EXPR
+ && TREE_CODE (TREE_TYPE (expr)) == ARRAY_TYPE)
+ {
+ *kind = BASE_DECL;
+ *expr0 = TREE_OPERAND (expr, 0);
+ return true;
+ }
+
*kind = BASE_ARBITRARY_EXPR;
*expr0 = expr;
return true;
@@ -3932,6 +3973,12 @@ debug_omp_tokenized_addr (vec<omp_addr_token *> &addr_tokens,
case ACCESS_INDEXED_REF_TO_ARRAY:
fputs ("access_indexed_ref_to_array", stderr);
break;
+ case ACCESS_NONCONTIG_ARRAY:
+ fputs ("access_noncontig_array", stderr);
+ break;
+ case ACCESS_NONCONTIG_REF_TO_ARRAY:
+ fputs ("access_noncontig_ref_to_array", stderr);
+ break;
}
break;
case ARRAY_BASE:
@@ -273,7 +273,9 @@ enum access_method_kinds
ACCESS_POINTER_OFFSET,
ACCESS_REF_TO_POINTER_OFFSET,
ACCESS_INDEXED_ARRAY,
- ACCESS_INDEXED_REF_TO_ARRAY
+ ACCESS_INDEXED_REF_TO_ARRAY,
+ ACCESS_NONCONTIG_ARRAY,
+ ACCESS_NONCONTIG_REF_TO_ARRAY
};
/* These are the kinds that a STRUCTURE_BASE or ARRAY_BASE (except
@@ -1388,6 +1388,55 @@ oacc_record_private_scalars (omp_context *ctx, tree clauses)
}
}
+/* Build record type for noncontiguous target update operations. Must be kept
+ in sync with libgomp/libgomp.h omp_noncontig_array_desc. */
+
+static tree
+omp_noncontig_descriptor_type (location_t loc)
+{
+ static tree cached = NULL_TREE;
+
+ if (cached)
+ return cached;
+
+ tree t = make_node (RECORD_TYPE);
+
+ tree fields = build_decl (loc, FIELD_DECL, get_identifier ("__ndims"),
+ size_type_node);
+
+ tree field = build_decl (loc, FIELD_DECL, get_identifier ("__elemsize"),
+ 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);
+ TREE_CHAIN (field) = fields;
+ fields = field;
+
+ field = build_decl (loc, FIELD_DECL, get_identifier ("__index"),
+ ptr_size_type);
+ TREE_CHAIN (field) = fields;
+ fields = field;
+
+ field = build_decl (loc, FIELD_DECL, get_identifier ("__length"),
+ ptr_size_type);
+ TREE_CHAIN (field) = fields;
+ fields = field;
+
+ field = build_decl (loc, FIELD_DECL, get_identifier ("__stride"),
+ ptr_size_type);
+ TREE_CHAIN (field) = fields;
+ fields = field;
+
+ finish_builtin_struct (t, "__omp_noncontig_desc_type", fields, ptr_type_node);
+
+ cached = t;
+
+ return t;
+}
+
/* Instantiate decls as necessary in CTX to satisfy the data sharing
specified by CLAUSES. If BASE_POINTERS_RESTRICT, install var field with
restrict. */
@@ -1949,8 +1998,74 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
install_var_local (array_decl, ctx);
break;
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_GRID
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM_GRID))
+ {
+ tree desc_type = omp_noncontig_descriptor_type (UNKNOWN_LOCATION);
- if (DECL_P (decl))
+ tree bare = decl;
+ if (TREE_CODE (bare) == VIEW_CONVERT_EXPR)
+ bare = TREE_OPERAND (bare, 0);
+
+ const char *desc_name = ".omp_noncontig_desc";
+ /* Try (but not too hard) to make a friendly name for the
+ descriptor. */
+ if (DECL_P (bare))
+ desc_name = ACONCAT ((".omp_nc_desc_",
+ IDENTIFIER_POINTER (DECL_NAME (bare)),
+ NULL));
+ tree desc = create_tmp_var (desc_type, desc_name);
+ DECL_NAMELESS (desc) = 1;
+ TREE_ADDRESSABLE (desc) = 1;
+
+ /* Adjust DECL so it refers to the first element of the array:
+ either by indirecting a pointer, or by selecting the zero'th
+ index of each dimension of an array. (We don't have a "bias"
+ as such for this type of noncontiguous update operation, just
+ the volume specified in the descriptor we build in
+ lower_omp_target.) */
+
+ if (TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE)
+ {
+ decl = build_fold_indirect_ref (decl);
+ OMP_CLAUSE_DECL (c) = decl;
+ }
+
+ tree field
+ = build_decl (OMP_CLAUSE_LOCATION (c), FIELD_DECL, NULL_TREE,
+ ptr_type_node);
+ SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node));
+ insert_field_into_struct (ctx->record_type, field);
+ splay_tree_insert (ctx->field_map, (splay_tree_key) c,
+ (splay_tree_value) field);
+
+ tree dn = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ 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;
+
+ field = build_decl (OMP_CLAUSE_LOCATION (c), FIELD_DECL,
+ NULL_TREE, ptr_type_node);
+ SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node));
+ insert_field_into_struct (ctx->record_type, field);
+ splay_tree_insert (ctx->field_map, (splay_tree_key) dn,
+ (splay_tree_value) field);
+
+ c = dn;
+ tree 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;
+ }
+ else if (DECL_P (decl))
{
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
@@ -2192,6 +2307,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
&& is_omp_target (ctx->stmt)
&& !is_gimple_omp_offloaded (ctx->stmt))
break;
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_GRID
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM_GRID
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE)
+ break;
if (DECL_P (decl))
{
if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
@@ -13666,6 +13786,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GOMP_MAP_DETACH:
case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+ case GOMP_MAP_TO_GRID:
+ case GOMP_MAP_FROM_GRID:
+ case GOMP_MAP_GRID_DIM:
+ case GOMP_MAP_GRID_STRIDE:
break;
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_FORCE_ALLOC:
@@ -13693,6 +13817,20 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gcc_unreachable ();
}
#endif
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_GRID
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM_GRID)
+ {
+ tree nc = OMP_CLAUSE_CHAIN (c);
+ gcc_assert (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_TO_PSET);
+ c = nc;
+ while ((nc = OMP_CLAUSE_CHAIN (c))
+ && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE))
+ c = nc;
+ map_cnt += 2;
+ continue;
+ }
/* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
@@ -14108,7 +14246,267 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
deep_map_offset_data,
deep_map_offset, &ilist);
}
- if (!DECL_P (ovar))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_GRID
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM_GRID))
+ {
+ tree decl = OMP_CLAUSE_DECL (c);
+ tree dn = OMP_CLAUSE_CHAIN (c);
+ gcc_assert (OMP_CLAUSE_CODE (dn) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (dn) == GOMP_MAP_TO_PSET);
+ tree desc = OMP_CLAUSE_DECL (dn);
+
+ tree oc, elsize = OMP_CLAUSE_SIZE (c);
+ tree type = TREE_TYPE (decl);
+ int i, dims = 0;
+ auto_vec<tree> tdims;
+ bool pointer_based = false, handled_pointer_section = false;
+ tree arrsize = fold_convert (sizetype, elsize);
+
+ /* Allow a single (maybe strided) array section if we have a
+ pointer base. */
+ if (TREE_CODE (decl) == INDIRECT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == POINTER_TYPE))
+ {
+ pointer_based = true;
+ dims = 1;
+ }
+ else
+ for (tree itype = type;
+ TREE_CODE (itype) == ARRAY_TYPE;
+ itype = TREE_TYPE (itype))
+ {
+ tdims.safe_push (itype);
+ dims++;
+ }
+
+ int tdim = tdims.length () - 1;
+
+ vec<constructor_elt, va_gc> *vdim;
+ vec<constructor_elt, va_gc> *vindex;
+ vec<constructor_elt, va_gc> *vlen;
+ vec<constructor_elt, va_gc> *vstride;
+ vec_alloc (vdim, dims);
+ vec_alloc (vindex, dims);
+ vec_alloc (vlen, dims);
+ vec_alloc (vstride, dims);
+
+ tree size_arr_type
+ = build_array_type_nelts (size_type_node, dims);
+
+ tree dim_tmp = create_tmp_var (size_arr_type, ".omp_dim");
+ DECL_NAMELESS (dim_tmp) = 1;
+ TREE_ADDRESSABLE (dim_tmp) = 1;
+ TREE_STATIC (dim_tmp) = 1;
+ tree index_tmp = create_tmp_var (size_arr_type, ".omp_index");
+ DECL_NAMELESS (index_tmp) = 1;
+ TREE_ADDRESSABLE (index_tmp) = 1;
+ TREE_STATIC (index_tmp) = 1;
+ tree len_tmp = create_tmp_var (size_arr_type, ".omp_len");
+ DECL_NAMELESS (len_tmp) = 1;
+ TREE_ADDRESSABLE (len_tmp) = 1;
+ TREE_STATIC (len_tmp) = 1;
+ tree stride_tmp = create_tmp_var (size_arr_type, ".omp_stride");
+ DECL_NAMELESS (stride_tmp) = 1;
+ TREE_ADDRESSABLE (stride_tmp) = 1;
+ TREE_STATIC (stride_tmp) = 1;
+
+ oc = c;
+ c = dn;
+
+ 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
+ && OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM)
+ {
+ index = OMP_CLAUSE_DECL (nc);
+ len = OMP_CLAUSE_SIZE (nc);
+
+ index = fold_convert (sizetype, index);
+ len = fold_convert (sizetype, len);
+
+ tree nc2 = OMP_CLAUSE_CHAIN (nc);
+ if (nc2
+ && OMP_CLAUSE_CODE (nc2) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (nc2)
+ == GOMP_MAP_GRID_STRIDE))
+ {
+ stride = OMP_CLAUSE_DECL (nc2);
+ stride = fold_convert (sizetype, stride);
+ nc = nc2;
+ }
+
+ if (tdim >= 0)
+ {
+ /* We have an array shape -- use that to find the
+ total size of the data on the target to look up
+ in libgomp. */
+ tree dtype = TYPE_DOMAIN (tdims[tdim]);
+ tree minval = TYPE_MIN_VALUE (dtype);
+ tree maxval = TYPE_MAX_VALUE (dtype);
+ minval = fold_convert (sizetype, minval);
+ maxval = fold_convert (sizetype, maxval);
+ dim = size_binop (MINUS_EXPR, maxval, minval);
+ dim = size_binop (PLUS_EXPR, dim,
+ size_one_node);
+ arrsize = size_binop (MULT_EXPR, arrsize, dim);
+ }
+ else if (pointer_based && !handled_pointer_section)
+ {
+ /* Use the selected array section to determine the
+ size of the array. */
+ tree tmp = size_binop (MULT_EXPR, len, stride);
+ tmp = size_binop (MINUS_EXPR, tmp, stride);
+ tmp = size_binop (PLUS_EXPR, tmp, size_one_node);
+ dim = size_binop (PLUS_EXPR, index, tmp);
+ arrsize = size_binop (MULT_EXPR, arrsize, dim);
+ handled_pointer_section = true;
+ }
+ else
+ {
+ if (pointer_based)
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "too many array section specifiers "
+ "for pointer-based array");
+ else
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "too many array section specifiers "
+ "for array");
+ dim = index = len = stride = error_mark_node;
+ }
+ tdim--;
+
+ c = nc;
+ }
+ else
+ {
+ /* We have more array dimensions than array section
+ specifiers. Copy the whole span. */
+ tree dtype = TYPE_DOMAIN (tdims[tdim]);
+ tree minval = TYPE_MIN_VALUE (dtype);
+ tree maxval = TYPE_MAX_VALUE (dtype);
+ minval = fold_convert (sizetype, minval);
+ maxval = fold_convert (sizetype, maxval);
+ dim = size_binop (MINUS_EXPR, maxval, minval);
+ dim = size_binop (PLUS_EXPR, dim, size_one_node);
+ len = dim;
+ index = size_zero_node;
+ }
+
+ if (TREE_CODE (dim) != INTEGER_CST)
+ TREE_STATIC (dim_tmp) = 0;
+
+ if (TREE_CODE (index) != INTEGER_CST)
+ TREE_STATIC (index_tmp) = 0;
+
+ if (TREE_CODE (len) != INTEGER_CST)
+ TREE_STATIC (len_tmp) = 0;
+
+ if (TREE_CODE (stride) != INTEGER_CST)
+ TREE_STATIC (stride_tmp) = 0;
+
+ tree cidx = size_int (i);
+ CONSTRUCTOR_APPEND_ELT (vdim, cidx, dim);
+ CONSTRUCTOR_APPEND_ELT (vindex, cidx, index);
+ CONSTRUCTOR_APPEND_ELT (vlen, cidx, len);
+ CONSTRUCTOR_APPEND_ELT (vstride, cidx, stride);
+ }
+
+ /* 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;
+
+ tree cdim = build_constructor (size_arr_type, vdim);
+ tree cindex = build_constructor (size_arr_type, vindex);
+ tree clen = build_constructor (size_arr_type, vlen);
+ tree cstride = build_constructor (size_arr_type, vstride);
+
+ if (TREE_STATIC (dim_tmp))
+ DECL_INITIAL (dim_tmp) = cdim;
+ else
+ gimplify_assign (dim_tmp, cdim, &ilist);
+
+ if (TREE_STATIC (index_tmp))
+ DECL_INITIAL (index_tmp) = cindex;
+ else
+ gimplify_assign (index_tmp, cindex, &ilist);
+
+ if (TREE_STATIC (len_tmp))
+ DECL_INITIAL (len_tmp) = clen;
+ else
+ gimplify_assign (len_tmp, clen, &ilist);
+
+ if (TREE_STATIC (stride_tmp))
+ DECL_INITIAL (stride_tmp) = cstride;
+ else
+ gimplify_assign (stride_tmp, cstride, &ilist);
+
+ tree desc_type = TREE_TYPE (desc);
+
+ tree ndims_field = TYPE_FIELDS (desc_type);
+ tree elemsize_field = DECL_CHAIN (ndims_field);
+ tree dim_field = DECL_CHAIN (elemsize_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);
+
+ bool all_static = (TREE_STATIC (dim_tmp)
+ && TREE_STATIC (index_tmp)
+ && TREE_STATIC (len_tmp)
+ && TREE_STATIC (stride_tmp));
+
+ dim_tmp = build4 (ARRAY_REF, sizetype, dim_tmp, size_zero_node,
+ NULL_TREE, NULL_TREE);
+ dim_tmp = build_fold_addr_expr (dim_tmp);
+
+ /* TODO: we could skip all-zeros index. */
+ index_tmp = build4 (ARRAY_REF, sizetype, index_tmp,
+ size_zero_node, NULL_TREE, NULL_TREE);
+ index_tmp = build_fold_addr_expr (index_tmp);
+
+ len_tmp = build4 (ARRAY_REF, sizetype, len_tmp, size_zero_node,
+ NULL_TREE, NULL_TREE);
+ len_tmp = build_fold_addr_expr (len_tmp);
+
+ /* TODO: we could skip all-ones stride. */
+ stride_tmp = build4 (ARRAY_REF, sizetype, stride_tmp,
+ size_zero_node, NULL_TREE, NULL_TREE);
+ stride_tmp = build_fold_addr_expr (stride_tmp);
+
+ elsize = fold_convert (sizetype, elsize);
+ tree ndims = size_int (dims);
+
+ CONSTRUCTOR_APPEND_ELT (v, ndims_field, ndims);
+ CONSTRUCTOR_APPEND_ELT (v, elemsize_field, elsize);
+ CONSTRUCTOR_APPEND_ELT (v, dim_field, dim_tmp);
+ CONSTRUCTOR_APPEND_ELT (v, index_field, index_tmp);
+ CONSTRUCTOR_APPEND_ELT (v, len_field, len_tmp);
+ CONSTRUCTOR_APPEND_ELT (v, stride_field, stride_tmp);
+
+ tree desc_ctor = build_constructor (desc_type, v);
+
+ if (all_static)
+ {
+ TREE_STATIC (desc) = 1;
+ DECL_INITIAL (desc) = desc_ctor;
+ }
+ else
+ gimplify_assign (desc, desc_ctor, &ilist);
+
+ OMP_CLAUSE_CHAIN (dn) = OMP_CLAUSE_CHAIN (nc);
+ c = oc;
+ nc = c;
+ }
+ else if (!DECL_P (ovar))
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
new file mode 100644
@@ -0,0 +1,22 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-original" }
+
+template<typename T, typename E, int A, int B, int C, int D>
+void foo ()
+{
+ T *ptr;
+ E a = A, b = B, c = C, d = D;
+
+ /* Dependent types for indices. */
+#pragma omp target update from(([a][b+1][c][d]) ptr[1:a-2][1:b][1:c-2][1:d-2])
+// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR.*\(\*ptr\) \[len: 1\]\) map\(grid_dim:1 \[len: [^\]]+\]\) map\(grid_dim:1 \[len: [^\]]+\]\) map\(grid_dim:1 \[len: [^\]]+\]\) map\(grid_dim:1 \[len: [^]]+\]\)} "original" } }
+}
+
+int main()
+{
+ char *ptr;
+
+ foo<char, short, 3, 4, 5, 6> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,134 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-original" }
+
+template<typename T>
+struct St
+{
+ T ***ppptr;
+ T ***&rppptr;
+
+ St(T ***p, T ***&rp) : ppptr(p), rppptr(rp) { }
+};
+
+template<typename A, typename B>
+void foo()
+{
+ A *ptr;
+ A **pptr = &ptr;
+ A ***ppptr = &pptr;
+ A ***&rppptr = ppptr;
+
+#pragma omp target update to(([10]) (**ppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(to_grid:VIEW_CONVERT_EXPR<int\[10\]>\(\*\*\*ppptr\) \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update to(([10]) (**rppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(to_grid:VIEW_CONVERT_EXPR<int\[10\]>\(\*\*\*\*rppptr\) \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update to((**ppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(to_grid:\*\*ppptr \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update to((**rppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(to_grid:\*\*\*rppptr \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+ B *ptr2;
+ B **pptr2 = &ptr2;
+ B ***ppptr2 = &pptr2;
+ St<B> *s = new St<B>(ppptr2, ppptr2);
+ St<B> **ps = &s;
+ St<B> **&rps = ps;
+
+#pragma omp target update from(([10]) (**(*ps)->ppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR<long int\[10\]>\(\*\*\*\(\*ps\)->ppptr\) \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from(([10]) (**(*rps)->rppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR<long int\[10\]>\(\*\*\*\*\(\*\*rps\)->rppptr\) \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from((**(*ps)->ppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(from_grid:\*\*\(\*ps\)->ppptr \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from((**(*rps)->rppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(from_grid:\*\*\*\(\*\*rps\)->rppptr \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+ B arr[10][10];
+ B (*parr)[10][10] = &arr;
+ B (**pparr2)[10][10] = &parr;
+ B (**&rpparr2)[10][10] = pparr2;
+
+#pragma omp target update from(**pparr2)
+// { dg-final { scan-tree-dump {from\(\*NON_LVALUE_EXPR <\*pparr2> \[len: [0-9]+\]\)} "original" } }
+
+#pragma omp target update to((**pparr2)[1:5:2][3:4:2])
+// { dg-final { scan-tree-dump {map\(to_grid:\*\*pparr2 \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from((**rpparr2)[1:5:2][3:4:2])
+// { dg-final { scan-tree-dump {map\(from_grid:\*\*\*rpparr2 \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+ delete s;
+}
+
+struct S
+{
+ short ***ppptr;
+ short ***&rppptr;
+
+ S(short ***p, short ***&rp) : ppptr(p), rppptr(rp) { }
+};
+
+int main()
+{
+ char *ptr;
+ char **pptr = &ptr;
+ char ***ppptr = &pptr;
+ char ***&rppptr = ppptr;
+
+#pragma omp target update to(([10]) (**ppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(to_grid:VIEW_CONVERT_EXPR<char\[10\]>\(\*\*\*ppptr\) \[len: 1\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update to(([10]) (**rppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(to_grid:VIEW_CONVERT_EXPR<char\[10\]>\(\*\*\*\*rppptr\) \[len: 1\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update to((**ppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(to_grid:\*\*ppptr \[len: 1\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update to((**rppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(to_grid:\*\*\*rppptr \[len: 1\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+ short *ptr2;
+ short **pptr2 = &ptr2;
+ short ***ppptr2 = &pptr2;
+ S *s = new S(ppptr2, ppptr2);
+ S **ps = &s;
+ S **&rps = ps;
+
+#pragma omp target update from(([10]) (**(*ps)->ppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR<short int\[10\]>\(\*\*\*\(\*ps\)->ppptr\) \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from(([10]) (**(*rps)->rppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR<short int\[10\]>\(\*\*\*\*\(\*\*rps\)->rppptr\) \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from((**(*ps)->ppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(from_grid:\*\*\(\*ps\)->ppptr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from((**(*rps)->rppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(from_grid:\*\*\*\(\*\*rps\)->rppptr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+ delete s;
+
+ short arr[10][10];
+ short (*parr)[10][10] = &arr;
+ short (**pparr)[10][10] = &parr;
+ short (**&rpparr)[10][10] = pparr;
+
+#pragma omp target update from(**pparr)
+// { dg-final { scan-tree-dump {from\(\*NON_LVALUE_EXPR <\*pparr> \[len: [0-9]+\]\)} "original" } }
+
+#pragma omp target update to((**pparr)[1:5:2][1:5:2])
+// { dg-final { scan-tree-dump {map\(to_grid:\*\*pparr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from((**rpparr)[1:5:2][1:5:2])
+// { dg-final { scan-tree-dump {map\(from_grid:\*\*\*rpparr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+ foo<int, long> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,47 @@
+// { dg-do compile }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T, int C, int D>
+void foo (T *w)
+{
+ memset (w, 0, sizeof (T) * 100);
+
+#pragma omp target enter data map(to: w[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ w[j * 10 + i] = i + j * 3;
+
+#pragma omp target update to(([C][D]) w[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: w[:100])
+}
+
+int main()
+{
+ float *arr = new float[100];
+
+ memset (arr, 0, sizeof (float) * 100);
+
+#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])
+
+ foo<float, 5, 20> (arr);
+
+ delete[] arr;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,52 @@
+// { dg-do compile }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T, int C, int D>
+void foo (T *w)
+{
+ /* This isn't allowed. We get a cascade of errors because it looks a bit
+ like lambda-definition syntax */
+#pragma omp target enter data map(to: ([C][D]) w[:100])
+ // { dg-error {capture of non-variable 'C'} "" { target *-*-* } .-1 }
+ // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+ // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 }
+ // { dg-error {expected '\)' before 'w'} "" { target *-*-* } .-4 }
+ // { dg-error {does not have pointer or array type} "" { target *-*-* } .-5 }
+
+#pragma omp target exit data map(from: ([C][D]) w[:100])
+ // { dg-error {capture of non-variable 'C'} "" { target *-*-* } .-1 }
+ // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+ // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 }
+ // { dg-error {expected '\)' before 'w'} "" { target *-*-* } .-4 }
+ // { dg-error {does not have pointer or array type} "" { target *-*-* } .-5 }
+}
+
+int main()
+{
+ float *arr = new float[100];
+
+ /* This isn't allowed (as above). */
+#pragma omp target enter data map(to: ([10][10]) arr[:100])
+ // { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 }
+ // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+ // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 }
+ // { dg-error {expected '\)' before 'arr'} "" { target *-*-* } .-4 }
+ // { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-5 }
+ // { dg-error {'#pragma omp target enter data' must contain at least one 'map' clause} "" { target *-*-*} .-6 }
+
+#pragma omp target exit data map(from: ([10][10]) arr[:100])
+ // { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 }
+ // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+ // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 }
+ // { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-4 }
+ // { dg-error {expected '\)' before 'arr'} "" { target *-*-* } .-5 }
+ // { dg-error {'#pragma omp target exit data' must contain at least one 'map' clause} "" { target *-*-* } .-6 }
+
+ foo<float, 5, 20> (arr);
+
+ delete[] arr;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,53 @@
+// { dg-do compile }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T>
+void foo (T *w)
+{
+ memset (w, 0, sizeof (T) * 100);
+ int c = 50;
+
+#pragma omp target enter data map(to: w[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ w[j * 10 + i] = i + j * 3;
+
+ /* This starts out looking like an array-shape cast. Make sure it's still
+ parsed as a lambda. */
+#pragma omp target update to(([c] (T *v) -> T { return v[c]; } (w)))
+ // { dg-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-1 }
+ // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-2 }
+
+#pragma omp target exit data map(from: w[:100])
+}
+
+int main()
+{
+ float *arr = new float[100];
+ int c = 50;
+
+ memset (arr, 0, sizeof (float) * 100);
+
+#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;
+
+ /* As above. */
+#pragma omp target update to(([c] (float *v) -> float { return v[c]; } (arr)))
+ // { dg-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-1 }
+ // { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 }
+ // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 }
+
+#pragma omp target exit data map(from: arr[:100])
+
+ foo<float> (arr);
+
+ delete[] arr;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,60 @@
+// { dg-do compile }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T>
+extern T* baz(T*);
+
+template<typename T>
+void foo (T *w)
+{
+ memset (w, 0, sizeof (T) * 100);
+ int c = 50;
+
+#pragma omp target enter data map(to: w[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ w[j * 10 + i] = i + j * 3;
+
+ /* No array-shaping inside a function call. */
+#pragma omp target update to(baz(([10][10]) w))
+ // { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 }
+ // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+ // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 }
+ // { dg-error {expected '\)' before 'w'} "" { target *-*-* } .-4 }
+ // { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-5 }
+
+#pragma omp target exit data map(from: w[:100])
+}
+
+int main()
+{
+ float *arr = new float[100];
+ int c = 50;
+
+ memset (arr, 0, sizeof (float) * 100);
+
+#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;
+
+ /* As above. */
+#pragma omp target update to(baz(([10][10]) arr))
+ // { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 }
+ // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+ // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 }
+ // { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-4 }
+ // { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-5 }
+
+#pragma omp target exit data map(from: arr[:100])
+
+ foo<float> (arr);
+
+ delete[] arr;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,55 @@
+// { dg-do compile }
+// { dg-additional-options "-std=c++14" }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T>
+void foo (T *w)
+{
+ memset (w, 0, sizeof (T) * 100);
+ int c = 50;
+
+#pragma omp target enter data map(to: w[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ w[j * 10 + i] = i + j * 3;
+
+ /* No array-shaping inside a lambda body. */
+#pragma omp target update to([&](const int d) -> auto& { return ([d][d]) w; } (10))
+// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-1 }
+// { dg-error {expected ';' before 'w'} "" { target *-*-* } .-2 }
+// { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-3 }
+
+#pragma omp target exit data map(from: w[:100])
+}
+
+int main()
+{
+ float *arr = new float[100];
+ int c = 50;
+
+ memset (arr, 0, sizeof (float) * 100);
+
+#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;
+
+ /* As above. */
+#pragma omp target update to([&](const int d) -> auto& { return ([d][d]) arr; } (10))
+// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-1 }
+// { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-2 }
+// { dg-error {expected ';' before 'arr'} "" { target *-*-* } .-3 }
+// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-4 }
+
+#pragma omp target exit data map(from: arr[:100])
+
+ foo<float> (arr);
+
+ delete[] arr;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,59 @@
+// { dg-do compile }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T>
+void foo (T *w)
+{
+ memset (w, 0, sizeof (T) * 100);
+
+#pragma omp target enter data map(to: w[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ w[j * 10 + i] = i + j * 3;
+
+ /* No array-shaping inside a statement expression. */
+#pragma omp target update to( ({ int d = 10; ([d][d]) w; )} )
+// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-1 }
+// { dg-warning {lambda expressions only available with} "" { target c++98_only } .-2 }
+// { dg-error {no match for 'operator\[\]'} "" { target *-*-* } .-3 }
+// { dg-error {expected ';' before 'w'} "" { target *-*-* } .-4 }
+// { dg-error {expected primary-expression before '\)' token} "" { target *-*-* } .-5 }
+// { dg-error {expected '\)' before end of line} "" { target *-*-* } .-6 }
+// { dg-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-7 }
+
+#pragma omp target exit data map(from: w[:100])
+}
+
+int main()
+{
+ float *arr = new float[100];
+
+ memset (arr, 0, sizeof (float) * 100);
+
+#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;
+
+ /* As above. */
+#pragma omp target update to( ({ int d = 10; ([d][d]) arr; )} )
+// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-1 }
+// { dg-warning {lambda expressions only available with} "" { target c++98_only } .-2 }
+// { dg-error {no match for 'operator\[\]'} "" { target *-*-* } .-3 }
+// { dg-error {expected primary-expression before '\)' token} "" { target *-*-* } .-4 }
+// { dg-error {expected '\)' before end of line} "" { target *-*-* } .-5 }
+// { dg-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-6 }
+// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-7 }
+
+#pragma omp target exit data map(from: arr[:100])
+
+ foo<float> (arr);
+
+ delete[] arr;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,48 @@
+// { dg-do compile }
+// { dg-additional-options "-std=c++11" }
+
+#include <new>
+
+template<typename T>
+struct St {
+ T *pp;
+};
+
+template<typename T>
+void foo (T *w)
+{
+ alignas (St<T>) unsigned char buf[sizeof (St<T>)];
+ T *sub1;
+
+ /* No array shaping op in brace initialiser (nonsensical anyway, but make
+ sure it doesn't parse). */
+#pragma omp target update to( new (buf) St<T> { ([10][10]) sub1 } )
+// { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 }
+// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+// { dg-error {expected '\}' before 'sub1'} "" { target *-*-* } .-3 }
+// { dg-error {expected '\)' before 'sub1'} "" { target *-*-* } .-4 }
+// { dg-error {expected an OpenMP clause before '\}' token} "" { target *-*-* } .-5 }
+}
+
+struct S {
+ int *pp;
+};
+
+int main()
+{
+ alignas (S) unsigned char buf[sizeof (S)];
+ int *sub1;
+
+ // As above.
+#pragma omp target update to( new (buf) S { ([10][10]) sub1 } )
+// { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 }
+// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+// { dg-error {expected '\}' before 'sub1'} "" { target *-*-* } .-3 }
+// { dg-error {expected '\)' before 'sub1'} "" { target *-*-* } .-4 }
+// { dg-error {expected an OpenMP clause before '\}' token} "" { target *-*-* } .-5 }
+// { dg-error {no match for 'operator\[\]'} "" { target *-*-* } .-6 }
+// { dg-error {could not convert} "" { target *-*-* } .-7 }
+// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-8 }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,50 @@
+// { dg-do compile }
+
+template<typename T>
+void foo ()
+{
+ T *ptr;
+
+#pragma omp target update to(([5][6][7]) ptr[0:4][0:7][0:7])
+// { dg-error {length '7' with stride '1' above array section size in 'to' clause} "" { target *-*-* } .-1 }
+
+#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 }
+
+ // This one's OK...
+#pragma omp target update from(([100]) ptr[3:33:3])
+
+ // But this is one element out of bounds.
+#pragma omp target update from(([100]) ptr[4:33:3])
+// { dg-error {high bound '101' above array section size in 'from' clause} "" { target *-*-* } .-1 }
+
+#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 }
+}
+
+int main()
+{
+ char *ptr;
+
+#pragma omp target update to(([5][6][7]) ptr[0:4][0:7][0:7])
+// { dg-error {length '7' 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 }
+
+#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 }
+
+ foo<char> ();
+
+ return 0;
+}
@@ -1124,6 +1124,18 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT:
pp_string (pp, "force_present,noncontig_array");
break;
+ case GOMP_MAP_TO_GRID:
+ pp_string (pp, "to_grid");
+ break;
+ case GOMP_MAP_FROM_GRID:
+ pp_string (pp, "from_grid");
+ break;
+ case GOMP_MAP_GRID_DIM:
+ pp_string (pp, "grid_dim");
+ break;
+ case GOMP_MAP_GRID_STRIDE:
+ pp_string (pp, "grid_stride");
+ break;
case GOMP_MAP_UNSET:
pp_string (pp, "unset");
break;
@@ -2750,6 +2762,11 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags,
dump_generic_node (pp, TREE_OPERAND (node, 1), spc, flags, false);
pp_colon (pp);
dump_generic_node (pp, TREE_OPERAND (node, 2), spc, flags, false);
+ if (TREE_OPERAND (node, 3))
+ {
+ pp_colon (pp);
+ dump_generic_node (pp, TREE_OPERAND (node, 3), spc, flags, false);
+ }
pp_right_bracket (pp);
break;
@@ -1372,7 +1372,7 @@ DEFTREECODE (OMP_ALLOCATE, "omp allocate", tcc_statement, 1)
DEFTREECODE (OMP_CLAUSE, "omp_clause", tcc_exceptional, 0)
/* An OpenMP array section. */
-DEFTREECODE (OMP_ARRAY_SECTION, "omp_array_section", tcc_expression, 3)
+DEFTREECODE (OMP_ARRAY_SECTION, "omp_array_section", tcc_expression, 4)
/* TRANSACTION_EXPR tree code.
Operand 0: BODY: contains body of the transaction. */
@@ -220,6 +220,9 @@ enum gomp_map_kind
GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION
= (GOMP_MAP_DEEP_COPY | 2),
+ GOMP_MAP_TO_GRID = (GOMP_MAP_DEEP_COPY | 4),
+ GOMP_MAP_FROM_GRID = (GOMP_MAP_DEEP_COPY | 5),
+
/* Internal to GCC, not used in libgomp. */
/* Do not map, but pointer assign a pointer instead. */
GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1),
@@ -243,7 +246,9 @@ enum gomp_map_kind
GOMP_MAP_POP_MAPPER_NAME = (GOMP_MAP_LAST | 10),
/* Used to hold a TREE_LIST of grouped nodes in an 'omp declare mapper'
definition (only for Fortran at present). */
- GOMP_MAP_MAPPING_GROUP = (GOMP_MAP_LAST | 11)
+ GOMP_MAP_MAPPING_GROUP = (GOMP_MAP_LAST | 11),
+ GOMP_MAP_GRID_DIM = (GOMP_MAP_LAST | 12),
+ GOMP_MAP_GRID_STRIDE = (GOMP_MAP_LAST | 13)
};
#define GOMP_MAP_COPY_TO_P(X) \
@@ -1313,6 +1313,20 @@ struct target_mem_desc {
};
+/* A rectangular section of an array, for noncontiguous target update
+ operations. Must be kept in sync with
+ omp-low.cc:omp_noncontig_descriptor_type. */
+
+typedef struct {
+ size_t ndims;
+ size_t elemsize;
+ size_t *dim;
+ size_t *index;
+ size_t *length;
+ size_t *stride;
+} omp_noncontig_array_desc;
+
+
typedef struct acc_dispatch_t
{
/* Execute. */
@@ -2604,6 +2604,13 @@ goacc_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom,
gomp_unmap_vars_internal (tgt, do_copyfrom, NULL, aq);
}
+static int omp_target_memcpy_rect_worker (void *, const void *, 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 *,
+ struct gomp_device_descr *);
+
static void
gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
size_t *sizes, void *kinds, bool short_mapkind)
@@ -2626,90 +2633,129 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
}
for (i = 0; i < mapnum; i++)
- if (sizes[i])
- {
- cur_node.host_start = (uintptr_t) hostaddrs[i];
- cur_node.host_end = cur_node.host_start + sizes[i];
- splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
- if (n)
- {
- int kind = get_kind (short_mapkind, kinds, i);
- if (n->host_start > cur_node.host_start
- || n->host_end < cur_node.host_end)
- {
- gomp_mutex_unlock (&devicep->lock);
- gomp_fatal ("Trying to update [%p..%p) object when "
- "only [%p..%p) is mapped",
- (void *) cur_node.host_start,
- (void *) cur_node.host_end,
- (void *) n->host_start,
- (void *) n->host_end);
- }
+ {
+ int kind = get_kind (short_mapkind, kinds, i);
+ if ((kind & typemask) == GOMP_MAP_TO_GRID
+ || (kind & typemask) == GOMP_MAP_FROM_GRID)
+ {
+ omp_noncontig_array_desc *desc
+ = (omp_noncontig_array_desc *) hostaddrs[i + 1];
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ 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)
+ {
+ if (n->aux && n->aux->attach_count)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_error ("noncontiguous update with attached pointers");
+ return;
+ }
+ void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset
+ + cur_node.host_start
+ - n->host_start);
+ 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,
+ NULL);
+ 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,
+ devicep);
+ }
+ i++;
+ }
+ else if (sizes[i])
+ {
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start + sizes[i];
+ splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
+ if (n)
+ {
+ if (n->host_start > cur_node.host_start
+ || n->host_end < cur_node.host_end)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("Trying to update [%p..%p) object when "
+ "only [%p..%p) is mapped",
+ (void *) cur_node.host_start,
+ (void *) cur_node.host_end,
+ (void *) n->host_start,
+ (void *) n->host_end);
+ }
- if (n->aux && n->aux->attach_count)
- {
- uintptr_t addr = cur_node.host_start;
- while (addr < cur_node.host_end)
- {
- /* We have to be careful not to overwrite still attached
- pointers during host<->device updates. */
- size_t i = (addr - cur_node.host_start) / sizeof (void *);
- if (n->aux->attach_count[i] == 0)
- {
- void *devaddr = (void *) (n->tgt->tgt_start
- + n->tgt_offset
- + addr - n->host_start);
- if (GOMP_MAP_COPY_TO_P (kind & typemask))
- gomp_copy_host2dev (devicep, NULL,
- devaddr, (void *) addr,
- sizeof (void *), false, NULL);
- if (GOMP_MAP_COPY_FROM_P (kind & typemask))
- gomp_copy_dev2host (devicep, NULL,
- (void *) addr, devaddr,
- sizeof (void *));
- }
- addr += sizeof (void *);
- }
- }
- else
- {
- void *hostaddr = (void *) cur_node.host_start;
- void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset
- + cur_node.host_start
- - n->host_start);
- size_t size = cur_node.host_end - cur_node.host_start;
+ if (n->aux && n->aux->attach_count)
+ {
+ uintptr_t addr = cur_node.host_start;
+ while (addr < cur_node.host_end)
+ {
+ /* We have to be careful not to overwrite still attached
+ pointers during host<->device updates. */
+ size_t i = (addr - cur_node.host_start) / sizeof (void *);
+ if (n->aux->attach_count[i] == 0)
+ {
+ void *devaddr = (void *) (n->tgt->tgt_start
+ + n->tgt_offset
+ + addr - n->host_start);
+ if (GOMP_MAP_COPY_TO_P (kind & typemask))
+ gomp_copy_host2dev (devicep, NULL,
+ devaddr, (void *) addr,
+ sizeof (void *), false, NULL);
+ if (GOMP_MAP_COPY_FROM_P (kind & typemask))
+ gomp_copy_dev2host (devicep, NULL,
+ (void *) addr, devaddr,
+ sizeof (void *));
+ }
+ addr += sizeof (void *);
+ }
+ }
+ else
+ {
+ void *hostaddr = (void *) cur_node.host_start;
+ void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset
+ + cur_node.host_start
+ - n->host_start);
+ size_t size = cur_node.host_end - cur_node.host_start;
- if (GOMP_MAP_COPY_TO_P (kind & typemask))
- gomp_copy_host2dev (devicep, NULL, devaddr, hostaddr, size,
- false, NULL);
- if (GOMP_MAP_COPY_FROM_P (kind & typemask))
- gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size);
- }
- }
- else
- {
- int kind = get_kind (short_mapkind, kinds, i);
+ if (GOMP_MAP_COPY_TO_P (kind & typemask))
+ gomp_copy_host2dev (devicep, NULL, devaddr, hostaddr, size,
+ false, NULL);
+ if (GOMP_MAP_COPY_FROM_P (kind & typemask))
+ gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size);
+ }
+ }
+ else
+ {
+ int kind = get_kind (short_mapkind, kinds, i);
- if (GOMP_MAP_PRESENT_P (kind))
- {
- /* We already looked up the memory region above and it
- was missing. */
- gomp_mutex_unlock (&devicep->lock);
+ if (GOMP_MAP_PRESENT_P (kind))
+ {
+ /* We already looked up the memory region above and it
+ was missing. */
+ gomp_mutex_unlock (&devicep->lock);
#ifdef HAVE_INTTYPES_H
- gomp_fatal ("present clause: not present on the device "
- "(addr: %p, size: %"PRIu64" (0x%"PRIx64"), "
- "dev: %d)", (void *) hostaddrs[i],
- (uint64_t) sizes[i], (uint64_t) sizes[i],
- devicep->target_id);
+ gomp_fatal ("present clause: not present on the device "
+ "(addr: %p, size: %"PRIu64" (0x%"PRIx64"), "
+ "dev: %d)", (void *) hostaddrs[i],
+ (uint64_t) sizes[i], (uint64_t) sizes[i],
+ devicep->target_id);
#else
- gomp_fatal ("present clause: not present on the device "
- "(addr: %p, size: %lu (0x%lx), dev: %d)",
- (void *) hostaddrs[i], (unsigned long) sizes[i],
- (unsigned long) sizes[i], devicep->target_id);
+ gomp_fatal ("present clause: not present on the device "
+ "(addr: %p, size: %lu (0x%lx), dev: %d)",
+ (void *) hostaddrs[i], (unsigned long) sizes[i],
+ (unsigned long) sizes[i], devicep->target_id);
#endif
- }
- }
- }
+ }
+ }
+ }
+ }
gomp_mutex_unlock (&devicep->lock);
}
@@ -5641,6 +5687,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,
+ const size_t *strides,
const size_t *dst_offsets,
const size_t *src_offsets,
const size_t *dst_dimensions,
@@ -5653,7 +5700,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)
+ if (num_dims == 1 && (!strides || strides[0] == 1))
{
if (__builtin_mul_overflow (element_size, volume[0], &length)
|| __builtin_mul_overflow (element_size, dst_offsets[0], &dst_off)
@@ -5726,6 +5773,38 @@ omp_target_memcpy_rect_worker (void *dst, const void *src, size_t element_size,
ret = 0;
return ret ? 0 : EINVAL;
}
+ else if (num_dims == 1 && strides)
+ {
+ size_t stride;
+
+ 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))
+ return EINVAL;
+
+ if (strides
+ && __builtin_mul_overflow (element_size, strides[0], &stride))
+ return EINVAL;
+
+ for (i = 0, ret = 1; i < volume[0] && ret; i++)
+ {
+ if (src_devicep == NULL)
+ ret = dst_devicep->host2dev_func (dst_devicep->target_id,
+ (char *) dst + dst_off,
+ (const char *) src + src_off,
+ element_size);
+ else if (dst_devicep == NULL)
+ ret = src_devicep->dev2host_func (src_devicep->target_id,
+ (char *) dst + dst_off,
+ (const char *) src + src_off,
+ element_size);
+ dst_off += stride;
+ src_off += stride;
+ }
+ return ret ? 0 : EINVAL;
+ }
/* FIXME: it would be nice to have some plugin function to handle
num_dims == 2 and num_dims == 3 more efficiently. Larger ones can
@@ -5739,13 +5818,19 @@ omp_target_memcpy_rect_worker (void *dst, const void *src, size_t element_size,
if (__builtin_mul_overflow (dst_slice, dst_offsets[0], &dst_off)
|| __builtin_mul_overflow (src_slice, src_offsets[0], &src_off))
return EINVAL;
+ if (strides
+ && (__builtin_mul_overflow (dst_slice, strides[0], &dst_slice)
+ || __builtin_mul_overflow (src_slice, strides[0], &src_slice)))
+ return EINVAL;
for (j = 0; j < volume[0]; j++)
{
ret = omp_target_memcpy_rect_worker ((char *) dst + dst_off,
(const char *) src + src_off,
element_size, num_dims - 1,
- volume + 1, dst_offsets + 1,
- src_offsets + 1, dst_dimensions + 1,
+ volume + 1,
+ strides ? strides + 1 : NULL,
+ dst_offsets + 1, src_offsets + 1,
+ dst_dimensions + 1,
src_dimensions + 1, dst_devicep,
src_devicep);
if (ret)
@@ -5791,9 +5876,10 @@ omp_target_memcpy_rect_copy (void *dst, const void *src,
else if (dst_devicep)
gomp_mutex_lock (&dst_devicep->lock);
int ret = omp_target_memcpy_rect_worker (dst, src, element_size, num_dims,
- volume, dst_offsets, src_offsets,
- dst_dimensions, src_dimensions,
- dst_devicep, src_devicep);
+ volume, NULL, dst_offsets,
+ src_offsets, dst_dimensions,
+ src_dimensions, dst_devicep,
+ src_devicep);
if (src_devicep)
gomp_mutex_unlock (&src_devicep->lock);
else if (dst_devicep)
new file mode 100644
@@ -0,0 +1,469 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <string.h>
+#include <assert.h>
+
+volatile int yy = 4, zz = 2, str_str = 2;
+
+template<typename T>
+void foo()
+{
+ T *arr;
+ int x = 5;
+ T arr2d[10][10];
+
+ arr = new T[100];
+
+ /* Update whole reshaped array. */
+
+ memset (arr, 0, 100 * sizeof (T));
+
+#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 (T));
+
+#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 (T));
+
+#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 (T));
+
+#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 (T));
+
+#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 (T));
+
+#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 (T));
+
+#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 (T));
+
+#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 (T));
+
+#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 (T));
+
+#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);
+
+ delete[] arr;
+}
+
+int main()
+{
+ int *arr;
+ int x = 5;
+ int arr2d[10][10];
+
+ arr = new int[100];
+
+ /* Update whole reshaped array. */
+
+ memset (arr, 0, 100 * sizeof (int));
+
+#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);
+
+ delete[] arr;
+
+ foo<long> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,61 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+#define N 10
+
+template<typename T>
+void foo ()
+{
+ T tarr[N * N];
+
+ memset (tarr, 0, N * N * sizeof (T));
+
+#pragma omp target enter data map(to: tarr)
+
+#pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ tarr[i * N + j] = 2 * (i + j);
+ }
+
+ /* An array, but cast to a pointer, then reshaped. */
+#pragma omp target update from(([N][N]) ((T *) &tarr[0])[4:3][5:3])
+
+ for (int i = 4; i < 7; i++)
+ for (int j = 5; j < 8; j++)
+ assert (tarr[i * N + j] == 2 * (i + j));
+
+#pragma omp target exit data map(delete: tarr)
+}
+
+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 = 4; i < 7; i++)
+ for (int j = 4; j < 7; j++)
+ assert (iarr[i * 10 + j] == i + j);
+
+#pragma omp target exit data map(delete: iarr)
+
+ foo<unsigned short> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,63 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+#define N 10
+
+template<typename T>
+void foo ()
+{
+ T tarr_real[N * N];
+ T (&tarr)[N * N] = tarr_real;
+
+ memset (tarr, 0, N * N * sizeof (T));
+
+#pragma omp target enter data map(to: tarr)
+
+#pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ tarr[i * N + j] = 2 * (i + j);
+ }
+
+ /* A ref to an array, but cast to a pointer, then reshaped. */
+#pragma omp target update from(([N][N]) ((T *) &tarr[0])[4:3][5:3])
+
+ for (int i = 4; i < 7; i++)
+ for (int j = 5; j < 8; j++)
+ assert (tarr[i * N + j] == 2 * (i + j));
+
+#pragma omp target exit data map(delete: tarr)
+}
+
+int main ()
+{
+ int iarr_real[N * N];
+ int (&iarr)[N * N] = iarr_real;
+
+ 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;
+ }
+
+ /* A ref to 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 = 4; i < 7; i++)
+ for (int j = 4; j < 7; j++)
+ assert (iarr[i * 10 + j] == i + j);
+
+#pragma omp target exit data map(delete: iarr)
+
+ foo<unsigned short> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,65 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+#define N 10
+
+template<typename T>
+void foo ()
+{
+ T tarr_real[N * N];
+ T *tarrp = &tarr_real[0];
+ T **tarrpp = &tarrp;
+
+ memset (tarrp, 0, N * N * sizeof (T));
+
+#pragma omp target enter data map(to: tarr_real)
+
+#pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ tarrp[i * N + j] = 2 * (i + j);
+ }
+
+ /* A pointer with an extra indirection. */
+#pragma omp target update from(([N][N]) (*tarrpp)[4:3][5:3])
+
+ for (int i = 4; i < 7; i++)
+ for (int j = 5; j < 8; j++)
+ assert (tarrp[i * N + j] == 2 * (i + j));
+
+#pragma omp target exit data map(delete: tarr_real)
+}
+
+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 = 4; i < 7; i++)
+ for (int j = 4; j < 7; j++)
+ assert (iarrp[i * 10 + j] == i + j);
+
+#pragma omp target exit data map(delete: iarr_real)
+
+ foo<unsigned short> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,89 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+#define N 10
+
+template<typename T>
+void foo ()
+{
+ T *tptr = new T[N * N * N];
+
+ memset (tptr, 0, N * N * N * sizeof (T));
+
+#pragma omp target enter data map(to: tptr[0:N*N*N])
+
+#pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ tptr[i * N * N + 4 * N + j] = 2 * (i + j);
+ }
+
+ /* An array ref between two array sections. */
+#pragma omp target update from(([N][N][N]) tptr[4:3][4][5:3])
+
+ for (int i = 4; i < 7; i++)
+ for (int j = 5; j < 8; j++)
+ assert (tptr[i * N * N + 4 * N + j] == 2 * (i + j));
+
+ memset (tptr, 0, N * N * N * sizeof (T));
+
+ for (int i = 0; i < N; i++)
+ tptr[2 * N * N + i * N + 4] = 4 * i;
+
+ /* Array section between two array refs. */
+#pragma omp target update to(([N][N][N]) tptr[2][3:6][4])
+
+#pragma omp target exit data map(from: tptr[0:N*N*N])
+
+ for (int i = 3; i < 9; i++)
+ assert (tptr[2 * N * N + i * N + 4] == 4 * i);
+
+#pragma omp target exit data map(delete: tptr[0:N*N*N])
+
+ delete[] tptr;
+}
+
+int main ()
+{
+ int *iptr = new int[N * N * N];
+
+ memset (iptr, 0, 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);
+
+ delete[] iptr;
+
+ foo<unsigned long> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,38 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T>
+void foo (T *w)
+{
+ memset (w, 0, sizeof (T) * 100);
+
+#pragma omp target enter data map(to: w[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ w[j * 10 + i] = i + j;
+
+#pragma omp target update to(([10][10]) w[3:2][1:8])
+
+#pragma omp target exit data map(from: w[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ if (j >= 3 && j < 5 && i >= 1 && i < 9)
+ assert (w[j * 10 + i] == i + j);
+ else
+ assert (w[j * 10 + i] == 0);
+}
+
+int main()
+{
+ int *arr = new int[100];
+
+ foo<int> (arr);
+
+ delete[] arr;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,38 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <string.h>
+#include <assert.h>
+
+template<int C, int D>
+void foo (double *w)
+{
+ memset (w, 0, sizeof (double) * 100);
+
+#pragma omp target enter data map(to: w[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ w[j * 10 + i] = i * 3 + j * 2;
+
+#pragma omp target update to(([C][D]) w[3:2][1:8])
+
+#pragma omp target exit data map(from: w[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ if (j >= 3 && j < 5 && i >= 1 && i < 9)
+ assert (w[j * 10 + i] == i * 3 + j * 2);
+ else
+ assert (w[j * 10 + i] == 0.0f);
+}
+
+int main()
+{
+ double *arr = new double[100];
+
+ foo<10, 10> (arr);
+
+ delete[] arr;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,38 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <string.h>
+#include <assert.h>
+
+template<auto C, auto D>
+void foo (double *w)
+{
+ memset (w, 0, sizeof (double) * 100);
+
+#pragma omp target enter data map(to: w[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ w[j * 10 + i] = i * 2 + j * 3;
+
+#pragma omp target update to(([C][D]) w[3:2][1:8])
+
+#pragma omp target exit data map(from: w[:100])
+
+ for (int j = 0; j < 10; j++)
+ for (int i = 0; i < 10; i++)
+ if (j >= 3 && j < 5 && i >= 1 && i < 9)
+ assert (w[j * 10 + i] == i * 2 + j * 3);
+ else
+ assert (w[j * 10 + i] == 0.0f);
+}
+
+int main()
+{
+ double *arr = new double[100];
+
+ foo<10, 10> (arr);
+
+ delete[] arr;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,38 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T, auto C>
+void foo (T *w, int e, int f, int g)
+{
+ memset (w, 0, sizeof (T) * 100);
+
+#pragma omp target enter data map(to: w[:100])
+
+ for (int j = 0; j < e; j++)
+ for (int i = 0; i < C; i++)
+ w[j * C + i] = i + j;
+
+#pragma omp target update to(([e][C]) w[3:2][f:g])
+
+#pragma omp target exit data map(from: w[:100])
+
+ for (int j = 0; j < e; j++)
+ for (int i = 0; i < C; i++)
+ if (j >= 3 && j < 5 && i >= f && i < f + g)
+ assert (w[j * C + i] == i + j);
+ else
+ assert (w[j * C + i] == 0.0f);
+}
+
+int main()
+{
+ float *arr = new float[100];
+
+ foo<float, 10> (arr, 10, 1, 8);
+
+ delete[] arr;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,54 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+template<typename T>
+void foo (T *&aref)
+{
+#pragma omp target enter data map(to: aref[:100])
+
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++)
+ for (int j = 0; j < 10; j++)
+ aref[i * 10 + j] = i + j;
+ }
+
+#pragma omp target update from(([10][10]) aref[2:3:2][7:3])
+
+ for (int i = 2; i < 8; i += 2)
+ for (int j = 7; j < 10; j++)
+ assert (aref[i * 10 + j] == i + j);
+
+#pragma omp target exit data map(delete: aref[:100])
+}
+
+int main()
+{
+ float *arr = new float[100];
+ float *&w = arr;
+
+ memset (arr, 0, 100 * sizeof (float));
+
+#pragma omp target enter data map(to: w[:100])
+
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++)
+ for (int j = 0; j < 10; j++)
+ w[i * 10 + j] = i + j;
+ }
+
+#pragma omp target update from(([10][10]) w[4:3][4:3])
+
+ for (int i = 4; i < 7; i++)
+ for (int j = 4; j < 7; j++)
+ assert (w[i * 10 + j] == i + j);
+
+#pragma omp target exit data map(delete: w[:100])
+
+ foo<float> (arr);
+
+ delete[] arr;
+}
new file mode 100644
@@ -0,0 +1,54 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+template<typename T>
+void foo (T (&aref)[10][10])
+{
+#pragma omp target enter data map(to: aref)
+
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++)
+ for (int j = 0; j < 10; j++)
+ aref[i][j] = i + j;
+ }
+
+#pragma omp target update from(aref[2:3:2][7:3])
+
+ for (int i = 2; i < 8; i += 2)
+ for (int j = 7; j < 10; j++)
+ assert (aref[i][j] == i + j);
+
+#pragma omp target exit data map(delete: aref)
+}
+
+int main()
+{
+ float arr2d[10][10];
+ float (&w)[10][10] = arr2d;
+
+ memset (&arr2d, 0, 100 * sizeof (float));
+
+#pragma omp target enter data map(to: w)
+
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++)
+ for (int j = 0; j < 10; j++)
+ w[i][j] = i + j;
+ }
+
+#pragma omp target update from(w[4:3][4:3])
+
+ for (int i = 4; i < 7; i++)
+ for (int j = 4; j < 7; j++)
+ assert (w[i][j] == i + j);
+
+#pragma omp target exit data map(delete: w)
+
+ foo<float> (arr2d);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,65 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+template<typename T>
+struct C {
+ T *&aptr;
+
+ C(T *&aptr_1) : aptr(aptr_1)
+ {
+ }
+};
+
+template<typename T>
+void foo (T *c)
+{
+#pragma omp target enter data map(to: c->aptr, c->aptr[:100])
+
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++)
+ for (int j = 0; j < 10; j++)
+ c->aptr[i * 10 + j] = i + j;
+ }
+
+#pragma omp target update from(([10][10]) c->aptr[2:3:2][7:3])
+
+ for (int i = 2; i < 8; i += 2)
+ for (int j = 7; j < 10; j++)
+ assert (c->aptr[i * 10 + j] == i + j);
+
+#pragma omp target exit data map(delete: c->aptr, c->aptr[:100])
+}
+
+int main()
+{
+ float *arr = new float[100];
+ C<float> cvar(arr);
+
+ memset (arr, 0, 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 = 4; i < 7; i++)
+ for (int j = 4; j < 7; j++)
+ assert (cvar.aptr[i * 10 + j] == i + j);
+
+#pragma omp target exit data map(delete: cvar.aptr, cvar.aptr[:100])
+
+ foo<C<float> > (&cvar);
+
+ delete[] arr;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,95 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+#define N 10
+
+struct B {
+ int (&aref)[N][N];
+
+ B(int (&aref1)[N][N]) : aref(aref1)
+ {
+ }
+};
+
+template<typename T, int S>
+struct C {
+ T (&aref)[S][S];
+
+ C(T (&aref1)[S][S]) : aref(aref1)
+ {
+ }
+};
+
+template<typename T>
+void foo (T *c)
+{
+#pragma omp target enter data map(to: c->aref)
+
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++)
+ for (int j = 0; j < 10; j++)
+ c->aref[i][j] = 2 * (i + j);
+ }
+
+#pragma omp target update from(c->aref[2:3:2][7:3])
+
+ for (int i = 2; i < 8; i += 2)
+ for (int j = 7; j < 10; j++)
+ assert (c->aref[i][j] == 2 * (i + j));
+
+#pragma omp target exit data map(delete: c->aref)
+}
+
+int main()
+{
+ int iarr[N][N];
+ float farr[N][N];
+ B bvar(iarr);
+ C<float, N> cvar(farr);
+
+ memset (iarr, 0, N * N * sizeof (int));
+ memset (farr, 0, N * N * sizeof (float));
+
+#pragma omp target enter data map(to: bvar.aref)
+
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++)
+ for (int j = 0; j < 10; j++)
+ bvar.aref[i][j] = i + j;
+ }
+
+#pragma omp target update from(bvar.aref[4:3][4:3])
+
+ for (int i = 4; i < 7; i++)
+ for (int j = 4; j < 7; j++)
+ assert (bvar.aref[i][j] == i + j);
+
+#pragma omp target exit data map(delete: bvar.aref)
+
+#pragma omp target enter data map(to: cvar.aref)
+
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++)
+ for (int j = 0; j < 10; j++)
+ cvar.aref[i][j] = i + j;
+ }
+
+#pragma omp target update from(cvar.aref[4:3][4:3])
+
+ for (int i = 4; i < 7; i++)
+ for (int j = 4; j < 7; j++)
+ assert (cvar.aref[i][j] == i + j);
+
+#pragma omp target exit data map(delete: cvar.aref)
+
+ memset (farr, 0, N * N * sizeof (float));
+
+ foo<C<float, N> > (&cvar);
+
+ return 0;
+}
This patch adds support for OpenMP 5.0 strided updates and the array-shaping operator ("([x][y][z]) foo[0:n]..."). This is mostly for C++ only so far, though necessary changes have been made to the C FE to adjust for changes to shared data structures. In terms of the implementation of various bits: - The OMP_ARRAY_SECTION tree code has been extended to take a 'stride' argument, and changes have been made throughout semantics.cc, etc. to take the new field into account -- including bounds checking. - A new type of cast operator has been added to represent the OpenMP array-shaping operator: OMP_ARRAYSHAPE_CAST_EXPR (1). - The address tokenization mechanism from previous patches has been extended with two new access kinds to represent noncontiguous array updates. - New mapping kinds have been added to represent noncontiguous updates: those which may be subject to array shaping, or have non-unit strides. These are processed by omp-low.cc into a kind of descriptor that is passed to the libgomp runtime (2). The current patch reuses an extended version of the helper code for omp_target_memcpy_rect, which may generate very many small host-device or device-host copies. (The "descriptor" has also been designed so reusing that functionality is relatively straightforward.) Optimising those multiple copies, e.g. by packing them into a single transfer when it would be beneficial, is left as the subject of a future patch. This patch has some adjustments to the omp-low.cc code after Chung-Lin's patch "OpenMP 5.0: Allow multiple clauses mapping same variable" (325f085897efca59879a64704ab15f1763ecb807), relative to the version last posted for mainline. Notes: (1) In a bit more detail: the array-shaping operator has the same precedence as a C-style cast, but applies to the whole expression, including array-section specifiers. We parse it initially as if it applies to the "value" of the whole expression: ([x][y]) ptr[0:10:2][1:5:2] i.e., something like: ([x][y]) (ptr[0:10:2][1:5:2]) or as if the cast applies to the innermost/right-hand side array section. Then, a little later in parsing (cp_parser_omp_var_list_no_open), we rewrite it to apply to the inner pointer instead: (([x][y]) ptr)[0:10:2][1:5:2] and that means a genuine multi-dimensional array or an array-shaped pointer can be handled pretty much the same for the rest of compilation. We use VIEW_CONVERT_EXPR for the "cast", unless we're processing a template definition, where we use a new tree code instead. (2) The new map kinds work like this. An update directive starts out with OMP_CLAUSE_TO or OMP_CLAUSE_FROM clauses representing the block in question and the direction of the needed transfer. If we detect a noncontiguous update, we emit a list of mapping nodes (type OMP_CLAUSE_MAP, with new kinds, so the "mapping group" machinery in gimplify.cc can be reused): OMP_CLAUSE_TO --> GOMP_MAP_TO_GRID (VIEW_CONVERT_EXPR<int[x][y]>(ptr) [len: <element-size>]) GOMP_MAP_GRID_DIM 0 [len: 10] (i.e. [0:10:2]) GOMP_MAP_GRID_STRIDE 2 GOMP_MAP_GRID_DIM 1 [len: 5] (i.e. [1:5:2]) GOMP_MAP_GRID_STRIDE 2 During omp-low.cc, this sequence is reformulated into: GOMP_MAP_TO_GRID (ptr) [len: <whole array size>] GOMP_MAP_TO_PSET (&ptr_desc [len: <desc size>]) "ptr_desc" is a struct, stored statically or constructed on the (host) stack, containing arrays representing the size of the whole array, the rectangular subregion to transfer, and the stride with which to walk over elements in each dimension. 2023-07-03 Julian Brown <julian@codesourcery.com> gcc/c-family/ * c-common.h (expand_array_base): Update prototype. * c-omp.cc (c_omp_address_inspector::map_supported_p): Support VIEW_CONVERT_EXPR and ADDR_EXPR codes. (omp_expand_grid_dim): New function. (omp_handle_noncontig_array): New function. (c_omp_address_inspector:expand_array_base): Remove DECL_P parameter. Support noncontiguous array updates. (c_omp_address_inspector::expand_component_selector): Support noncontiguous array updates. (c_omp_address_inspector::expand_map_clause): Update calls to expand_array_base. * c-pretty-print.cc (c_pretty_printer::postfix_expression): Add OMP_ARRAY_SECTION stride support. gcc/c/ * c-parser.cc (c_parser_postfix_expression_after_primary): Dummy stride support (for now). (struct omp_dim): Add stride support. (c_parser_omp_variable_list): Likewise. * c-tree.h (build_omp_array_section): Update prototype. * c-typeck.cc (mark_exp_read): Add stride support for OMP_ARRAY_SECTION. (build_omp_array_section): Add stride support. (handle_omp_array_sections_1): Add minimal stride support. gcc/cp/ * cp-objcp-common.cc (cp_common_init_ts): Add array-shape cast support. * cp-tree.def (OMP_ARRAYSHAPE_CAST_EXPR): Add tree code. * cp-tree.h (DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST): Add flag. (cp_omp_create_arrayshape_type, cp_build_omp_arrayshape_cast): Add prototypes. (grok_omp_array_section, build_omp_array_section): Add stride parameters. * decl.cc (create_anon_array_type): New function. (cp_omp_create_arrayshape_type): New function. * decl2.cc (grok_omp_array_section): Add stride parameter. (min_vis_expr_r): Add OMP_ARRAYSHAPE_CAST_EXPR support. * error.cc (dump_expr): Add stride support for OMP_ARRAY_SECTION. * mangle.cc (write_expression): Add OMP_ARRAYSHAPE_CAST_EXPR support. * operators.def (OMP_ARRAYSHAPE_CAST_EXPR): Add. * parser.cc (cp_parser_new): Initialise omp_array_shaping_op_p and omp_has_array_shape_p fields. (cp_parser_statement_expr): Don't allow array shaping op in statement exprs. (cp_parser_postfix_open_square_expression): Add stride parsing for array sections. Use array section code to represent array refs if we have an array-shaping operator. (cp_parser_parenthesized_expression_list): Don't allow array-shaping op here. (cp_parser_cast_expression): Add array-shaping operator parsing. (cp_parser_lambda_expression): Don't allow array-shaping op in lambda body. (cp_parser_braced_list): Don't allow array-shaping op in braced list. (struct omp_dim): Add stride field. (cp_parser_var_list_no_open): Add stride/array shape support. (cp_parser_omp_target_update): Handle noncontiguous updates. * parser.h (cp_parser): Add omp_array_shaping_op_p and omp_has_array_shape_p fields. * pt.cc (tsubst): Add array-shape cast support. (tsubst_copy, tsubst_copy_and_build): Likewise. Add stride support for OMP_ARRAY_SECTION. (tsubst_omp_clause_decl): Add stride support for OMP_ARRAY_SECTION. * semantics.cc (handle_omp_array_sections_1): Add DISCONTIGUOUS parameter and stride support. (omp_array_section_low_bound): New function. (handle_omp_array_sections): Add DISCONTIGUOUS parameter and stride support. (finish_omp_clauses): Update calls to handle_omp_array_sections, and add noncontiguous array update support. (cp_build_omp_arrayshape_cast): New function. * typeck.cc (structural_comptypes): Add array-shape cast support. (build_omp_array_section): Add stride parameter. (check_for_casting_away_constness): Add OMP_ARRAYSHAPE_CAST_EXPR support. gcc/ * gimplify.cc (omp_group_last, omp_group_base): Add GOMP_MAP_TO_GRID, GOMP_MAP_FROM_GRID support. (gimplify_adjust_omp_clauses): Support new GOMP_MAP_GRID_DIM, GOMP_MAP_GRID_STRIDE mapping nodes. Don't crash on e.g. misuse of ADDR_EXPR in mapping clauses. * omp-general.cc (omp_parse_noncontiguous_array): New function. (omp_parse_access_method): Add noncontiguous array support. (omp_parse_structure_base): Add array-shaping support. (debug_omp_tokenized_addr): Add ACCESS_NONCONTIG_ARRAY, ACCESS_NONCONTIG_REF_TO_ARRAY token support. * omp-general.h (access_method_kinds): Add ACCESS_NONCONTIG_ARRAY and ACCESS_NONCONTIG_REF_TO_ARRAY access kinds. * omp-low.cc (omp_noncontig_descriptor_type): New function. (scan_sharing_clauses): Support noncontiguous array updates. (lower_omp_target): Likewise. * tree-pretty-print.cc (dump_omp_clause): Add GOMP_MAP_TO_GRID, GOMP_MAP_FROM_GRID, GOMP_MAP_GRID_DIM, GOMP_MAP_GRID_STRIDE map kinds. (dump_generic_node): Add stride support for OMP_ARRAY_SECTION. * tree.def (OMP_ARRAY_SECTION): Add stride argument. include/ * gomp-constants.h (gomp_map_kind): Add GOMP_MAP_TO_GRID, GOMP_MAP_FROM_GRID, GOMP_MAP_GRID_DIM, GOMP_MAP_GRID_STRIDE map kinds. gcc/testsuite/ * g++.dg/gomp/array-shaping-1.C: New test. * g++.dg/gomp/array-shaping-2.C: New test. * g++.dg/gomp/bad-array-shaping-1.C: New test. * g++.dg/gomp/bad-array-shaping-2.C: New test. * g++.dg/gomp/bad-array-shaping-3.C: New test. * g++.dg/gomp/bad-array-shaping-4.C: New test. * g++.dg/gomp/bad-array-shaping-5.C: New test. * g++.dg/gomp/bad-array-shaping-6.C: New test. * g++.dg/gomp/bad-array-shaping-7.C: New test. * g++.dg/gomp/bad-array-shaping-8.C: New test. libgomp/ * libgomp.h (omp_noncontig_array_desc): New struct. * target.c (omp_target_memcpy_rect_worker): Add stride array parameter. Forward declare. Add STRIDES parameter and strided update support. (gomp_update): Add noncontiguous (strided/shaped) update support. * 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. * testsuite/libgomp.c++/array-shaping-7.C: New test. * testsuite/libgomp.c++/array-shaping-8.C: New test. * testsuite/libgomp.c++/array-shaping-9.C: New test. * testsuite/libgomp.c++/array-shaping-10.C: New test. * testsuite/libgomp.c++/array-shaping-11.C: New test. * testsuite/libgomp.c++/array-shaping-12.C: New test. * testsuite/libgomp.c++/array-shaping-13.C: New test. --- gcc/c-family/c-common.h | 2 +- gcc/c-family/c-omp.cc | 206 +++++++- gcc/c-family/c-pretty-print.cc | 5 + gcc/c/c-parser.cc | 32 +- gcc/c/c-tree.h | 2 +- gcc/c/c-typeck.cc | 26 +- gcc/cp/cp-objcp-common.cc | 1 + gcc/cp/cp-tree.def | 1 + gcc/cp/cp-tree.h | 13 +- gcc/cp/decl.cc | 75 +++ gcc/cp/decl2.cc | 19 +- gcc/cp/error.cc | 5 + gcc/cp/mangle.cc | 1 + gcc/cp/operators.def | 1 + gcc/cp/parser.cc | 303 ++++++++++- gcc/cp/parser.h | 7 + gcc/cp/pt.cc | 39 +- gcc/cp/semantics.cc | 262 ++++++++-- gcc/cp/typeck.cc | 12 +- gcc/gimplify.cc | 44 +- gcc/omp-general.cc | 47 ++ gcc/omp-general.h | 4 +- gcc/omp-low.cc | 402 ++++++++++++++- gcc/testsuite/g++.dg/gomp/array-shaping-1.C | 22 + gcc/testsuite/g++.dg/gomp/array-shaping-2.C | 134 +++++ .../g++.dg/gomp/bad-array-shaping-1.C | 47 ++ .../g++.dg/gomp/bad-array-shaping-2.C | 52 ++ .../g++.dg/gomp/bad-array-shaping-3.C | 53 ++ .../g++.dg/gomp/bad-array-shaping-4.C | 60 +++ .../g++.dg/gomp/bad-array-shaping-5.C | 55 ++ .../g++.dg/gomp/bad-array-shaping-6.C | 59 +++ .../g++.dg/gomp/bad-array-shaping-7.C | 48 ++ .../g++.dg/gomp/bad-array-shaping-8.C | 50 ++ gcc/tree-pretty-print.cc | 17 + gcc/tree.def | 2 +- include/gomp-constants.h | 7 +- libgomp/libgomp.h | 14 + libgomp/target.c | 254 ++++++---- .../testsuite/libgomp.c++/array-shaping-1.C | 469 ++++++++++++++++++ .../testsuite/libgomp.c++/array-shaping-10.C | 61 +++ .../testsuite/libgomp.c++/array-shaping-11.C | 63 +++ .../testsuite/libgomp.c++/array-shaping-12.C | 65 +++ .../testsuite/libgomp.c++/array-shaping-13.C | 89 ++++ .../testsuite/libgomp.c++/array-shaping-2.C | 38 ++ .../testsuite/libgomp.c++/array-shaping-3.C | 38 ++ .../testsuite/libgomp.c++/array-shaping-4.C | 38 ++ .../testsuite/libgomp.c++/array-shaping-5.C | 38 ++ .../testsuite/libgomp.c++/array-shaping-6.C | 54 ++ .../testsuite/libgomp.c++/array-shaping-7.C | 54 ++ .../testsuite/libgomp.c++/array-shaping-8.C | 65 +++ .../testsuite/libgomp.c++/array-shaping-9.C | 95 ++++ 51 files changed, 3334 insertions(+), 216 deletions(-) create mode 100644 gcc/testsuite/g++.dg/gomp/array-shaping-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/array-shaping-2.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-shaping-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-shaping-2.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-shaping-3.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-shaping-4.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-shaping-5.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-shaping-6.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-shaping-7.C create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C create mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-1.C create mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-10.C create mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-11.C create mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-12.C create mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-13.C create mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-2.C create mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-3.C create mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-4.C create mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-5.C create mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-6.C create mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-7.C create mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-8.C create mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-9.C