@@ -1349,7 +1349,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 *, c_omp_region_type);
tree * expand_map_clause (tree *, tree, vec<omp_addr_token *> &,
@@ -3244,7 +3244,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
@@ -3384,21 +3386,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)));
@@ -3410,6 +3486,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);
@@ -3422,7 +3499,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)
@@ -3663,6 +3748,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;
@@ -3719,8 +3838,27 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
bool chain_p = omp_access_chain_p (addr_tokens, i + 1);
bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP;
bool target = (ort & C_ORT_TARGET) != 0;
+ 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:
@@ -3730,7 +3868,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));
@@ -3739,7 +3877,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;
@@ -3750,15 +3888,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,
@@ -3775,12 +3913,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,
@@ -3795,10 +3933,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);
@@ -3813,13 +3951,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;
}
@@ -3837,8 +4000,7 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
pc = &OMP_CLAUSE_CHAIN (c);
}
- i += 2;
- *idx = i;
+ *idx = acc + 1;
if (target && chain_p && map_p)
return omp_expand_access_chain (pc, expr, addr_tokens, idx);
@@ -3870,7 +4032,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;
}
@@ -3879,7 +4041,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;
@@ -11224,7 +11224,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;
@@ -13760,11 +13760,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
@@ -13892,7 +13892,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);
}
@@ -13908,21 +13910,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)
{
@@ -13932,7 +13935,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)
{
@@ -13941,7 +13944,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))
@@ -14095,7 +14099,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)
@@ -14119,7 +14124,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)
@@ -752,7 +752,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);
@@ -2046,6 +2046,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;
@@ -2931,7 +2933,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;
@@ -2968,7 +2971,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);
}
@@ -13694,7 +13698,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
bool &maybe_zero_len, unsigned int &first_non_one,
enum c_omp_region_type ort)
{
- 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)
{
@@ -13779,8 +13783,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)))
@@ -13797,6 +13804,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))
@@ -14011,7 +14025,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)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"array section is not contiguous in %qs clause",
@@ -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)
@@ -515,6 +515,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)
@@ -4863,6 +4864,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
@@ -6931,6 +6934,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);
@@ -6964,7 +6969,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);
@@ -7812,6 +7818,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);
@@ -8063,7 +8071,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);
@@ -11552,6 +11552,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
@@ -596,43 +596,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;
}
@@ -2608,6 +2614,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;
@@ -3580,6 +3580,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)
@@ -4347,6 +4347,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;
@@ -5283,6 +5289,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;
@@ -8193,7 +8200,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)
@@ -8226,9 +8233,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;
@@ -8237,7 +8258,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;
@@ -8245,11 +8266,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. */
@@ -8571,6 +8604,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)
@@ -8589,7 +8623,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);
@@ -8656,6 +8692,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;
}
}
@@ -8663,6 +8700,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;
}
@@ -9931,6 +9969,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)
@@ -10000,6 +10040,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;
@@ -10021,6 +10083,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)
@@ -10034,6 +10100,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)
@@ -11148,6 +11232,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;
@@ -11157,6 +11242,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;
@@ -11208,6 +11294,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. */
@@ -25542,6 +25629,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;
@@ -37407,11 +37495,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
@@ -37443,10 +37531,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
@@ -37457,49 +37557,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)
{
@@ -37516,7 +37726,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)
{
@@ -37525,7 +37735,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))
@@ -37689,7 +37908,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
@@ -37711,7 +37931,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;
@@ -37724,6 +37945,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);
@@ -45557,8 +45780,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
@@ -16574,6 +16574,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)
@@ -17392,6 +17396,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;
@@ -17617,12 +17622,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:
@@ -17916,14 +17923,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);
@@ -20524,6 +20534,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),
@@ -20690,7 +20708,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))
@@ -20705,7 +20723,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:
@@ -5165,9 +5165,9 @@ public:
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
bool &maybe_zero_len, unsigned int &first_non_one,
- enum c_omp_region_type ort)
+ enum c_omp_region_type ort, int *discontiguous)
{
- tree ret, low_bound, length, type;
+ tree ret, low_bound, length, stride, type;
bool openacc = (ort & C_ORT_ACC) != 0;
if (TREE_CODE (t) != OMP_ARRAY_SECTION)
{
@@ -5229,18 +5229,26 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
&& TREE_CODE (TREE_OPERAND (t, 0)) == FIELD_DECL)
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, ort);
+ maybe_zero_len, first_non_one, ort,
+ discontiguous);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
- type = TREE_TYPE (ret);
+ if (TREE_CODE (ret) == OMP_ARRAY_SECTION)
+ type = TREE_TYPE (TREE_TYPE (TREE_OPERAND (ret, 0)));
+ else
+ type = TREE_TYPE (ret);
low_bound = TREE_OPERAND (t, 1);
length = TREE_OPERAND (t, 2);
+ stride = TREE_OPERAND (t, 3);
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)))
@@ -5257,15 +5265,26 @@ 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)
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))
@@ -5276,9 +5295,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))
@@ -5397,12 +5422,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)
@@ -5410,7 +5452,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))
{
@@ -5480,12 +5522,20 @@ 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)))
{
- 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;
+ }
}
}
}
@@ -5497,7 +5547,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)
@@ -5516,15 +5566,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;
@@ -5539,7 +5619,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,
- ort);
+ ort, discontiguous);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -5580,6 +5660,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))
{
@@ -5587,6 +5669,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);
i--;
if (low_bound
@@ -5599,12 +5682,56 @@ 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;
+ 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])
@@ -5618,12 +5745,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
@@ -5732,6 +5864,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;
@@ -5739,9 +5874,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
@@ -5768,6 +5900,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]". */
@@ -5786,7 +5922,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
@@ -6925,7 +7062,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;
@@ -7965,7 +8102,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)
@@ -8120,6 +8258,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:
@@ -8134,8 +8275,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
{
@@ -8723,7 +8867,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
{
@@ -12719,4 +12863,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
@@ -8150,6 +8153,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();
}
@@ -9203,6 +9203,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:
{
@@ -9346,6 +9359,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);
@@ -14047,7 +14064,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;
@@ -14155,6 +14174,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
@@ -14247,8 +14280,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;
@@ -3174,6 +3174,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)
{
@@ -3182,6 +3208,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))
@@ -3255,6 +3288,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;
@@ -3404,6 +3445,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:
@@ -254,7 +254,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
@@ -1141,6 +1141,55 @@ fixup_child_record_type (omp_context *ctx)
= build_qualified_type (build_reference_type (type), TYPE_QUAL_RESTRICT);
}
+/* 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. */
@@ -1657,7 +1706,74 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_local (decl, ctx);
break;
}
- if (DECL_P (decl))
+ 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);
+
+ 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) decl,
+ (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) desc,
+ (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)
@@ -1890,6 +2006,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
@@ -12809,6 +12930,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:
@@ -12826,6 +12951,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:
@@ -13117,7 +13256,267 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FIRSTPRIVATE_REFERENCE)))
break;
- 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 '#pragma omp' 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 '#pragma omp' 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;
+}
@@ -994,6 +994,18 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
pp_string (pp, "attach_zero_length_array_section");
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;
@@ -2589,6 +2601,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;
@@ -1321,7 +1321,7 @@ DEFTREECODE (OMP_ATOMIC_CAPTURE_NEW, "omp_atomic_capture_new", tcc_statement, 2)
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. */
@@ -180,6 +180,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),
@@ -197,7 +200,9 @@ enum gomp_map_kind
GOMP_MAP_POP_MAPPER_NAME = (GOMP_MAP_LAST | 6),
/* 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 | 7)
+ GOMP_MAP_MAPPING_GROUP = (GOMP_MAP_LAST | 7),
+ GOMP_MAP_GRID_DIM = (GOMP_MAP_LAST | 8),
+ GOMP_MAP_GRID_STRIDE = (GOMP_MAP_LAST | 9),
};
#define GOMP_MAP_COPY_TO_P(X) \
@@ -1303,6 +1303,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. */
@@ -2103,6 +2103,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)
@@ -2125,67 +2132,106 @@ 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);
- }
- }
- }
+ 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);
+ }
+ }
+ }
+ }
gomp_mutex_unlock (&devicep->lock);
}
@@ -4524,6 +4570,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,
@@ -4536,7 +4583,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)
@@ -4567,6 +4614,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
@@ -4580,13 +4659,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)
@@ -4632,9 +4717,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;
+}