@@ -1340,8 +1340,8 @@ public:
bool map_supported_p ();
- static tree get_origin (tree);
- static tree maybe_unconvert_ref (tree);
+ tree get_origin (tree);
+ tree maybe_unconvert_ref (tree);
bool maybe_zero_length_array_section (tree);
@@ -2680,6 +2680,9 @@ c_omp_split_clauses (location_t loc, enum tree_code code,
}
else if (TREE_CODE (OMP_CLAUSE_DECL (c)) == TREE_LIST)
{
+ /* TODO: This can go away once we transition all uses of
+ TREE_LIST for representing OMP array sections to
+ OMP_ARRAY_SECTION. */
tree t;
for (t = OMP_CLAUSE_DECL (c);
TREE_CODE (t) == TREE_LIST; t = TREE_CHAIN (t))
@@ -2688,6 +2691,17 @@ c_omp_split_clauses (location_t loc, enum tree_code code,
bitmap_clear_bit (&allocate_head, DECL_UID (t));
break;
}
+ else if (TREE_CODE (OMP_CLAUSE_DECL (c)) == OMP_ARRAY_SECTION)
+ {
+ tree t;
+ for (t = OMP_CLAUSE_DECL (c);
+ TREE_CODE (t) == OMP_ARRAY_SECTION;
+ t = TREE_OPERAND (t, 0))
+ ;
+ if (DECL_P (t))
+ bitmap_clear_bit (&allocate_head, DECL_UID (t));
+ break;
+ }
/* FALLTHRU */
case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_FIRSTPRIVATE:
@@ -3229,6 +3243,7 @@ c_omp_address_inspector::map_supported_p ()
|| TREE_CODE (t) == SAVE_EXPR
|| TREE_CODE (t) == POINTER_PLUS_EXPR
|| TREE_CODE (t) == NON_LVALUE_EXPR
+ || TREE_CODE (t) == OMP_ARRAY_SECTION
|| TREE_CODE (t) == NOP_EXPR)
if (TREE_CODE (t) == COMPOUND_EXPR)
t = TREE_OPERAND (t, 1);
@@ -3258,7 +3273,8 @@ c_omp_address_inspector::get_origin (tree t)
else if (TREE_CODE (t) == POINTER_PLUS_EXPR
|| TREE_CODE (t) == SAVE_EXPR)
t = TREE_OPERAND (t, 0);
- else if (TREE_CODE (t) == INDIRECT_REF
+ else if (!processing_template_decl_p ()
+ && TREE_CODE (t) == INDIRECT_REF
&& TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == REFERENCE_TYPE)
t = TREE_OPERAND (t, 0);
else
@@ -3275,7 +3291,10 @@ c_omp_address_inspector::get_origin (tree t)
tree
c_omp_address_inspector::maybe_unconvert_ref (tree t)
{
- if (TREE_CODE (t) == INDIRECT_REF
+ /* Be careful not to dereference the type if we're processing a
+ template decl, else it might be NULL. */
+ if (!processing_template_decl_p ()
+ && TREE_CODE (t) == INDIRECT_REF
&& TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == REFERENCE_TYPE)
return TREE_OPERAND (t, 0);
@@ -9488,6 +9488,7 @@ potential_constant_expression_1 (tree t, bool want_rval, bool strict, bool now,
case OACC_ENTER_DATA:
case OACC_EXIT_DATA:
case OACC_UPDATE:
+ case OMP_ARRAY_SECTION:
/* GCC internal stuff. */
case VA_ARG_EXPR:
case TRANSACTION_EXPR:
@@ -6956,6 +6956,7 @@ 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 delete_sanity (location_t, tree, tree, bool,
int, tsubst_flags_t);
extern tree check_classfn (tree, tree, tree);
@@ -8050,6 +8051,7 @@ 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_x_unary_op (location_t,
enum tree_code, cp_expr,
tree, tsubst_flags_t);
@@ -591,6 +591,51 @@ grok_array_decl (location_t loc, tree array_expr, tree index_exp,
return expr;
}
+/* Build an OMP_ARRAY_SECTION expression, handling usage in template
+ definitions, etc. */
+
+tree
+grok_omp_array_section (location_t loc, tree array_expr, tree index,
+ tree length)
+{
+ tree orig_array_expr = array_expr;
+ tree orig_index = index;
+ tree orig_length = length;
+
+ if (error_operand_p (array_expr)
+ || error_operand_p (index)
+ || error_operand_p (length))
+ 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))
+ return build_min_nt_loc (loc, OMP_ARRAY_SECTION, array_expr, index,
+ length);
+ array_expr = build_non_dependent_expr (array_expr);
+ if (index)
+ index = build_non_dependent_expr (index);
+ if (length)
+ length = build_non_dependent_expr (length);
+ }
+
+ index = fold_non_dependent_expr (index);
+ length = fold_non_dependent_expr (length);
+
+ /* 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);
+
+ if (processing_template_decl)
+ expr = build_min_non_dep (OMP_ARRAY_SECTION, expr, orig_array_expr,
+ orig_index, orig_length);
+ return expr;
+}
+
/* Given the cast expression EXP, checking out its validity. Either return
an error_mark_node if there was an unavoidable error, return a cast to
void for trying to delete a pointer w/ the value 0, or return the
@@ -2497,6 +2497,15 @@ dump_expr (cxx_pretty_printer *pp, tree t, int flags)
pp_cxx_right_bracket (pp);
break;
+ case OMP_ARRAY_SECTION:
+ dump_expr (pp, TREE_OPERAND (t, 0), flags);
+ pp_cxx_left_bracket (pp);
+ dump_expr (pp, TREE_OPERAND (t, 1), flags);
+ pp_colon (pp);
+ dump_expr (pp, TREE_OPERAND (t, 2), flags);
+ pp_cxx_right_bracket (pp);
+ break;
+
case UNARY_PLUS_EXPR:
dump_unary_op (pp, "+", t, flags);
break;
@@ -4344,6 +4344,9 @@ cp_parser_new (cp_lexer *lexer)
parser->omp_declare_simd = NULL;
parser->oacc_routine = NULL;
+ /* Disallow OpenMP array sections in expressions. */
+ parser->omp_array_section_p = false;
+
/* Not declaring an implicit function template. */
parser->auto_is_implicit_function_template_parm_p = false;
parser->fully_implicit_function_template_p = false;
@@ -5279,6 +5282,7 @@ static cp_expr
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);
/* Consume the '('. */
location_t start_loc = cp_lexer_peek_token (parser->lexer)->location;
@@ -8096,6 +8100,7 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
releasing_vec expression_list = NULL;
location_t loc = cp_lexer_peek_token (parser->lexer)->location;
bool saved_greater_than_is_operator_p;
+ bool saved_colon_corrects_to_scope_p;
/* Consume the `[' token. */
cp_lexer_consume_token (parser->lexer);
@@ -8103,6 +8108,10 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
saved_greater_than_is_operator_p = parser->greater_than_is_operator_p;
parser->greater_than_is_operator_p = true;
+ saved_colon_corrects_to_scope_p = parser->colon_corrects_to_scope_p;
+ if (parser->omp_array_section_p)
+ parser->colon_corrects_to_scope_p = false;
+
/* Parse the index expression. */
/* ??? For offsetof, there is a question of what to allow here. If
offsetof is not being used in an integral constant expression context,
@@ -8113,7 +8122,8 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
constant expressions here. */
if (for_offsetof)
index = cp_parser_constant_expression (parser);
- else
+ else if (!parser->omp_array_section_p
+ || cp_lexer_next_token_is_not (parser->lexer, CPP_COLON))
{
if (cxx_dialect >= cxx23
&& cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_SQUARE))
@@ -8170,6 +8180,68 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
parser->greater_than_is_operator_p = saved_greater_than_is_operator_p;
+ if (cxx_dialect >= cxx23
+ && parser->omp_array_section_p
+ && expression_list.get () != NULL
+ && vec_safe_length (expression_list) > 1)
+ {
+ error_at (loc, "cannot use multidimensional subscript in OpenMP array "
+ "section");
+ index = error_mark_node;
+ }
+ if (parser->omp_array_section_p
+ && cp_lexer_next_token_is (parser->lexer, CPP_COLON))
+ {
+ cp_lexer_consume_token (parser->lexer);
+ tree length = NULL_TREE;
+ if (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_SQUARE))
+ {
+ if (cxx_dialect >= cxx23)
+ {
+ cp_expr expr
+ = cp_parser_parenthesized_expression_list_elt (parser,
+ /*cast_p=*/
+ false,
+ /*allow_exp_p=*/
+ true,
+ /*non_cst_p=*/
+ NULL);
+
+ if (expr == error_mark_node)
+ length = error_mark_node;
+ else
+ length = expr.get_value ();
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+ {
+ error_at (loc, "cannot use multidimensional subscript in "
+ "OpenMP array section");
+ length = error_mark_node;
+ }
+ }
+ else
+ length
+ = 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)
+ {
+ cp_parser_skip_to_closing_square_bracket (parser);
+ return error_mark_node;
+ }
+ else
+ cp_parser_require (parser, CPP_CLOSE_SQUARE, RT_CLOSE_SQUARE);
+
+ return grok_omp_array_section (input_location, postfix_expression, index,
+ length);
+ }
+
+ parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p;
+
/* Look for the closing `]'. */
cp_parser_require (parser, CPP_CLOSE_SQUARE, RT_CLOSE_SQUARE);
@@ -8498,6 +8570,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;
/* Assume all the expressions will be constant. */
if (non_constant_p)
@@ -8515,6 +8588,9 @@ cp_parser_parenthesized_expression_list (cp_parser* parser,
= parser->greater_than_is_operator_p;
parser->greater_than_is_operator_p = true;
+ saved_omp_array_section_p = parser->omp_array_section_p;
+ parser->omp_array_section_p = false;
+
cp_expr expr (NULL_TREE);
/* Consume expressions until there are no more. */
@@ -8579,12 +8655,14 @@ 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;
return NULL;
}
}
parser->greater_than_is_operator_p
= saved_greater_than_is_operator_p;
+ parser->omp_array_section_p = saved_omp_array_section_p;
return expression_list;
}
@@ -11069,6 +11147,7 @@ cp_parser_lambda_expression (cp_parser* parser)
cp_binding_level* implicit_template_scope = parser->implicit_template_scope;
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;
parser->num_template_parameter_lists = 0;
parser->in_statement = 0;
@@ -11077,6 +11156,7 @@ cp_parser_lambda_expression (cp_parser* parser)
parser->implicit_template_parms = 0;
parser->implicit_template_scope = 0;
parser->auto_is_implicit_function_template_parm_p = false;
+ parser->omp_array_section_p = false;
/* The body of a lambda in a discarded statement is not discarded. */
bool discarded = in_discarded_stmt;
@@ -11127,6 +11207,7 @@ cp_parser_lambda_expression (cp_parser* parser)
parser->implicit_template_scope = implicit_template_scope;
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;
}
/* This field is only used during parsing of the lambda. */
@@ -25460,6 +25541,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);
/* Consume the `{' token. */
matching_braces braces;
@@ -37324,7 +37406,7 @@ struct omp_dim
static tree
cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
tree list, bool *colon,
- bool allow_deref = false)
+ bool map_lvalue = false)
{
auto_vec<omp_dim> dims;
bool array_section_p;
@@ -37341,6 +37423,105 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
if (kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY)
cp_parser_parse_tentatively (parser);
+ /* This condition doesn't include OMP_CLAUSE_DEPEND or
+ OMP_CLAUSE_AFFINITY since lvalue ("locator list") parsing for those is
+ handled further down the function. */
+ else if (map_lvalue
+ && (kind == OMP_CLAUSE_MAP
+ || kind == OMP_CLAUSE_TO
+ || kind == OMP_CLAUSE_FROM))
+ {
+ auto s = make_temp_override (parser->omp_array_section_p, true);
+ token = cp_lexer_peek_token (parser->lexer);
+ location_t loc = token->location;
+ decl = cp_parser_assignment_expression (parser);
+
+ /* 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
+ semantics.cc:handle_omp_array_sections_1. It might be more
+ efficient to move this logic to that function instead, analysing
+ the parsed expression directly rather than this preprocessed
+ form. */
+ dims.truncate (0);
+ if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
+ {
+ 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));
+ decl = TREE_OPERAND (decl, 0);
+ }
+
+ while (TREE_CODE (decl) == ARRAY_REF
+ || TREE_CODE (decl) == INDIRECT_REF
+ || TREE_CODE (decl) == COMPOUND_EXPR)
+ {
+ if (REFERENCE_REF_P (decl))
+ break;
+
+ if (TREE_CODE (decl) == COMPOUND_EXPR)
+ {
+ decl = TREE_OPERAND (decl, 1);
+ STRIP_NOPS (decl);
+ }
+ else if (TREE_CODE (decl) == INDIRECT_REF)
+ {
+ dims.safe_push (omp_dim (integer_zero_node,
+ integer_one_node, 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));
+ decl = TREE_OPERAND (decl, 0);
+ }
+ }
+
+ /* Bare references have their own special handling, so remove
+ the explicit dereference added by convert_from_reference. */
+ 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);
+ }
+ else if (TREE_CODE (decl) == INDIRECT_REF)
+ {
+ bool ref_p = REFERENCE_REF_P (decl);
+
+ /* Turn *foo into foo[0:1]. */
+ decl = TREE_OPERAND (decl, 0);
+ STRIP_NOPS (decl);
+
+ /* If we have "*foo" and
+ - it's an indirection of a reference, "unconvert" it, i.e.
+ strip the indirection (to just "foo").
+ - it's an indirection of a pointer, turn it into
+ "foo[0:1]". */
+ if (!ref_p)
+ decl = grok_omp_array_section (loc, decl, integer_zero_node,
+ integer_one_node);
+ }
+ else if (TREE_CODE (decl) == ARRAY_REF)
+ {
+ tree idx = TREE_OPERAND (decl, 1);
+
+ decl = TREE_OPERAND (decl, 0);
+ STRIP_NOPS (decl);
+
+ decl = grok_omp_array_section (loc, decl, idx, integer_one_node);
+ }
+ else if (TREE_CODE (decl) == NON_LVALUE_EXPR
+ || CONVERT_EXPR_P (decl))
+ decl = TREE_OPERAND (decl, 0);
+
+ goto build_clause;
+ }
token = cp_lexer_peek_token (parser->lexer);
if (kind != 0
&& cp_parser_is_keyword (token, RID_THIS))
@@ -37419,8 +37600,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
case OMP_CLAUSE_TO:
start_component_ref:
while (cp_lexer_next_token_is (parser->lexer, CPP_DOT)
- || (allow_deref
- && cp_lexer_next_token_is (parser->lexer, CPP_DEREF)))
+ || cp_lexer_next_token_is (parser->lexer, CPP_DEREF))
{
cpp_ttype ttype
= cp_lexer_next_token_is (parser->lexer, CPP_DOT)
@@ -37506,9 +37686,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
|| kind == OMP_CLAUSE_TO)
&& !array_section_p
&& (cp_lexer_next_token_is (parser->lexer, CPP_DOT)
- || (allow_deref
- && cp_lexer_next_token_is (parser->lexer,
- CPP_DEREF))))
+ || cp_lexer_next_token_is (parser->lexer, CPP_DEREF)))
{
for (unsigned i = 0; i < dims.length (); i++)
{
@@ -37520,8 +37698,9 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
}
else
for (unsigned i = 0; i < dims.length (); i++)
- decl = tree_cons (dims[i].low_bound, dims[i].length, decl);
-
+ decl = build_omp_array_section (input_location, decl,
+ dims[i].low_bound,
+ dims[i].length);
break;
default:
break;
@@ -37542,6 +37721,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
cp_parser_parse_definitely (parser);
}
+ build_clause:
tree u = build_omp_clause (token->location, kind);
OMP_CLAUSE_DECL (u) = decl;
OMP_CLAUSE_CHAIN (u) = list;
@@ -37591,11 +37771,11 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
static tree
cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list,
- bool allow_deref = false)
+ bool map_lvalue = false)
{
if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
return cp_parser_omp_var_list_no_open (parser, kind, list, NULL,
- allow_deref);
+ map_lvalue);
return list;
}
@@ -37662,7 +37842,7 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
gcc_unreachable ();
}
tree nl, c;
- nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true);
+ nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, false);
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_SET_MAP_KIND (c, kind);
@@ -40486,8 +40666,13 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
cp_lexer_consume_token (parser->lexer);
}
+ /* We introduce a scope here so that errors parsing e.g. "always", "close"
+ tokens do not propagate to later directives that might use them
+ legally. */
+ begin_scope (sk_omp, NULL);
nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list,
NULL, true);
+ finish_scope ();
for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_SET_MAP_KIND (c, kind);
@@ -407,6 +407,9 @@ struct GTY(()) cp_parser {
/* TRUE if omp::directive or omp::sequence attributes may not appear. */
bool omp_attrs_forbidden_p;
+ /* TRUE if an OpenMP array section is allowed. */
+ bool omp_array_section_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
@@ -16579,6 +16579,7 @@ tsubst (tree t, tree args, tsubst_flags_t complain, tree in_decl)
case CALL_EXPR:
case ARRAY_REF:
case SCOPE_REF:
+ case OMP_ARRAY_SECTION:
/* We should use one of the expression tsubsts for these codes. */
gcc_unreachable ();
@@ -17584,6 +17585,17 @@ tsubst_copy (tree t, tree args, tsubst_flags_t complain, tree in_decl)
return build_nt (ARRAY_REF, op0, op1, NULL_TREE, NULL_TREE);
}
+ case OMP_ARRAY_SECTION:
+ {
+ tree op0 = tsubst_copy (TREE_OPERAND (t, 0), args, complain, in_decl);
+ tree op1 = NULL_TREE, op2 = 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);
+ }
+
case CALL_EXPR:
{
int n = VL_EXP_OPERAND_LENGTH (t);
@@ -17869,6 +17881,22 @@ tsubst_omp_clause_decl (tree decl, tree args, tsubst_flags_t complain,
= OMP_CLAUSE_DOACROSS_SINK_NEGATIVE (decl);
return ret;
}
+ else if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
+ {
+ tree low_bound
+ = tsubst_expr (TREE_OPERAND (decl, 1), args, complain, in_decl);
+ tree length = tsubst_expr (TREE_OPERAND (decl, 2), 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)
+ return decl;
+ tree ret = build3 (OMP_ARRAY_SECTION, TREE_TYPE (base), base, low_bound,
+ length);
+ return ret;
+ }
tree ret = tsubst_expr (decl, args, complain, in_decl);
/* Undo convert_from_reference tsubst_expr could have called. */
if (decl
@@ -20612,6 +20640,27 @@ tsubst_copy_and_build (tree t,
RECUR (TREE_OPERAND (t, 1)),
complain|decltype_flag));
+ case OMP_ARRAY_SECTION:
+ {
+ tree op0 = RECUR (TREE_OPERAND (t, 0));
+ tree op1 = NULL_TREE, op2 = NULL_TREE;
+ if (op0 == error_mark_node)
+ RETURN (error_mark_node);
+ if (TREE_OPERAND (t, 1))
+ {
+ op1 = RECUR (TREE_OPERAND (t, 1));
+ if (op1 == error_mark_node)
+ RETURN (error_mark_node);
+ }
+ if (TREE_OPERAND (t, 2))
+ {
+ op2 = RECUR (TREE_OPERAND (t, 2));
+ if (op2 == error_mark_node)
+ RETURN (error_mark_node);
+ }
+ RETURN (build_omp_array_section (EXPR_LOCATION (t), op0, op1, op2));
+ }
+
case SIZEOF_EXPR:
if (PACK_EXPANSION_P (TREE_OPERAND (t, 0))
|| ARGUMENT_PACK_P (TREE_OPERAND (t, 0)))
@@ -5168,7 +5168,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
{
tree ret, low_bound, length, type;
bool openacc = (ort & C_ORT_ACC) != 0;
- if (TREE_CODE (t) != TREE_LIST)
+ if (TREE_CODE (t) != OMP_ARRAY_SECTION)
{
if (error_operand_p (t))
return error_mark_node;
@@ -5190,7 +5190,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
ret = t_refto;
if (TREE_CODE (t) == FIELD_DECL)
ret = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
- else if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
+ else if (!VAR_P (t)
+ && (openacc || !EXPR_P (t))
+ && TREE_CODE (t) != PARM_DECL)
{
if (processing_template_decl && TREE_CODE (t) != OVERLOAD)
return NULL_TREE;
@@ -5223,16 +5225,16 @@ 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)
- && TREE_CODE (TREE_CHAIN (t)) == FIELD_DECL)
- TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t), false);
- ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), 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);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
type = TREE_TYPE (ret);
- low_bound = TREE_PURPOSE (t);
- length = TREE_VALUE (t);
+ low_bound = TREE_OPERAND (t, 1);
+ length = TREE_OPERAND (t, 2);
if ((low_bound && type_dependent_expression_p (low_bound))
|| (length && type_dependent_expression_p (length)))
return NULL_TREE;
@@ -5438,7 +5440,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
tree lb = cp_save_expr (low_bound);
if (lb != low_bound)
{
- TREE_PURPOSE (t) = lb;
+ TREE_OPERAND (t, 1) = lb;
low_bound = lb;
}
}
@@ -5469,14 +5471,14 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
array-section-subscript, the array section could be non-contiguous. */
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
- && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
+ && TREE_CODE (TREE_OPERAND (t, 0)) == OMP_ARRAY_SECTION)
{
/* If any prior dimension has a non-one length, then deem this
array section as non-contiguous. */
- for (tree d = TREE_CHAIN (t); TREE_CODE (d) == TREE_LIST;
- d = TREE_CHAIN (d))
+ for (tree d = TREE_OPERAND (t, 0); TREE_CODE (d) == OMP_ARRAY_SECTION;
+ d = TREE_OPERAND (d, 0))
{
- tree d_length = TREE_VALUE (d);
+ tree d_length = TREE_OPERAND (d, 2);
if (d_length == NULL_TREE || !integer_onep (d_length))
{
error_at (OMP_CLAUSE_LOCATION (c),
@@ -5499,7 +5501,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
tree lb = cp_save_expr (low_bound);
if (lb != low_bound)
{
- TREE_PURPOSE (t) = lb;
+ TREE_OPERAND (t, 1) = lb;
low_bound = lb;
}
/* Temporarily disable -fstrong-eval-order for array reductions.
@@ -5577,10 +5579,12 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
return false;
for (i = num, t = OMP_CLAUSE_DECL (c); i > 0;
- t = TREE_CHAIN (t))
+ t = TREE_OPERAND (t, 0))
{
- tree low_bound = TREE_PURPOSE (t);
- tree length = TREE_VALUE (t);
+ gcc_assert (TREE_CODE (t) == OMP_ARRAY_SECTION);
+
+ tree low_bound = TREE_OPERAND (t, 1);
+ tree length = TREE_OPERAND (t, 2);
i--;
if (low_bound
@@ -6690,8 +6694,8 @@ cp_oacc_check_attachments (tree c)
tree t = OMP_CLAUSE_DECL (c);
tree type;
- while (TREE_CODE (t) == TREE_LIST)
- t = TREE_CHAIN (t);
+ while (TREE_CODE (t) == OMP_ARRAY_SECTION)
+ t = TREE_OPERAND (t, 0);
type = TREE_TYPE (t);
@@ -6798,7 +6802,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_TASK_REDUCTION:
field_ok = ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP);
t = OMP_CLAUSE_DECL (c);
- if (TREE_CODE (t) == TREE_LIST)
+ if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
if (handle_omp_array_sections (c, ort))
{
@@ -6814,10 +6818,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
remove = true;
break;
}
- if (TREE_CODE (t) == TREE_LIST)
+ if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
- while (TREE_CODE (t) == TREE_LIST)
- t = TREE_CHAIN (t);
+ while (TREE_CODE (t) == OMP_ARRAY_SECTION)
+ t = TREE_OPERAND (t, 0);
}
else
{
@@ -7838,7 +7842,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else
last_iterators = NULL_TREE;
- if (TREE_CODE (t) == TREE_LIST)
+ if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
if (handle_omp_array_sections (c, ort))
remove = true;
@@ -7998,7 +8002,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
auto_vec<omp_addr_token *, 10> addr_tokens;
t = OMP_CLAUSE_DECL (c);
- if (TREE_CODE (t) == TREE_LIST)
+ if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
grp_start_p = pc;
grp_sentinel = OMP_CLAUSE_CHAIN (c);
@@ -8008,7 +8012,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else
{
t = OMP_CLAUSE_DECL (c);
- if (TREE_CODE (t) != TREE_LIST
+ if (TREE_CODE (t) != OMP_ARRAY_SECTION
&& !type_dependent_expression_p (t)
&& !omp_mappable_type (TREE_TYPE (t)))
{
@@ -8189,7 +8193,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH))
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
+ || (!openacc && EXPR_P (t))))
break;
if (DECL_P (t))
error_at (OMP_CLAUSE_LOCATION (c),
@@ -8584,15 +8589,15 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_HAS_DEVICE_ADDR:
t = OMP_CLAUSE_DECL (c);
- if (TREE_CODE (t) == TREE_LIST)
+ if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
if (handle_omp_array_sections (c, ort))
remove = true;
else
{
t = OMP_CLAUSE_DECL (c);
- while (TREE_CODE (t) == TREE_LIST)
- t = TREE_CHAIN (t);
+ while (TREE_CODE (t) == OMP_ARRAY_SECTION)
+ t = TREE_OPERAND (t, 0);
while (TREE_CODE (t) == INDIRECT_REF
|| TREE_CODE (t) == ARRAY_REF)
t = TREE_OPERAND (t, 0);
@@ -8963,10 +8968,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (DECL_P (t))
bitmap_clear_bit (&aligned_head, DECL_UID (t));
}
- else if (TREE_CODE (t) == TREE_LIST)
+ else if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
- while (TREE_CODE (t) == TREE_LIST)
- t = TREE_CHAIN (t);
+ while (TREE_CODE (t) == OMP_ARRAY_SECTION)
+ t = TREE_OPERAND (t, 0);
if (DECL_P (t))
bitmap_clear_bit (&aligned_head, DECL_UID (t));
t = OMP_CLAUSE_DECL (c);
@@ -4786,6 +4786,56 @@ build_x_array_ref (location_t loc, tree arg1, tree arg2,
return expr;
}
+/* Build an OpenMP array section reference, creating an exact type for the
+ resulting expression based on the element type and bounds if possible. If
+ we have variable bounds, create an incomplete array type for the result
+ instead. */
+
+tree
+build_omp_array_section (location_t loc, tree array_expr, tree index,
+ tree length)
+{
+ tree idxtype;
+
+ /* If we know the integer bounds, create an index type with exact
+ low/high (or zero/length) bounds. Otherwise, create an incomplete
+ array type. (This mostly only affects diagnostics.) */
+ if (index != NULL_TREE
+ && length != NULL_TREE
+ && TREE_CODE (index) == INTEGER_CST
+ && TREE_CODE (length) == INTEGER_CST)
+ {
+ tree low = fold_convert (sizetype, index);
+ tree high = fold_convert (sizetype, length);
+ high = size_binop (PLUS_EXPR, low, high);
+ high = size_binop (MINUS_EXPR, high, size_one_node);
+ idxtype = build_range_type (sizetype, low, high);
+ }
+ else if ((index == NULL_TREE || integer_zerop (index))
+ && length != NULL_TREE
+ && TREE_CODE (length) == INTEGER_CST)
+ idxtype = build_index_type (length);
+ else
+ idxtype = NULL_TREE;
+
+ tree type = TREE_TYPE (array_expr);
+ gcc_assert (type);
+ type = non_reference (type);
+
+ tree sectype, eltype = TREE_TYPE (type);
+
+ /* It's not an array or pointer type. Just reuse the type of the
+ original expression as the type of the array section (an error will be
+ raised anyway, later). */
+ if (eltype == NULL_TREE)
+ sectype = TREE_TYPE (array_expr);
+ else
+ sectype = build_array_type (eltype, idxtype);
+
+ return build3_loc (loc, OMP_ARRAY_SECTION, sectype, array_expr, index,
+ length);
+}
+
/* Return whether OP is an expression of enum type cast to integer
type. In C++ even unsigned enum types are cast to signed integer
types. We do not want to issue warnings about comparisons between
@@ -17246,6 +17246,9 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
case TREE_LIST:
gcc_unreachable ();
+ case OMP_ARRAY_SECTION:
+ gcc_unreachable ();
+
case COMPOUND_EXPR:
ret = gimplify_compound_expr (expr_p, pre_p, fallback != fb_none);
break;
@@ -20,12 +20,12 @@ foo (void)
;
#pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */
- /* { dg-error "'close' has not been declared" "" { target c++ } .-1 } */
+ /* { dg-error "'close' was not declared in this scope" "" { target c++ } .-1 } */
/* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */
;
#pragma omp target map (always a) /* { dg-error "'always' undeclared" "" { target c } } */
- /* { dg-error "'always' has not been declared" "" { target c++ } .-1 } */
+ /* { dg-error "'always' was not declared in this scope" "" { target c++ } .-1 } */
/* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */
;
new file mode 100644
@@ -0,0 +1,38 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-original" }
+
+int x;
+
+template<int C, int D>
+void foo()
+{
+ int arr1[40];
+#pragma omp target map(arr1[x ? C : D])
+// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR <x != 0 \? 3 : 5>\] \[len: [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR <x != 0 \? 3 : 5>\] - \(long int\) &arr1\]\)} "original" } }
+ { }
+#pragma omp target map(arr1[x ? C : D : D])
+// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR <x != 0 \? 3 : 5>\] \[len: [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR <x != 0 \? 3 : 5>\] - \(long int\) &arr1\]\)} "original" } }
+ { }
+#pragma omp target map(arr1[1 : x ? C : D])
+// { dg-final { scan-tree-dump {map\(tofrom:arr1\[1\] \[len: x != 0 \? [0-9]+ : [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: [0-9]+\]\)} "original" } }
+ { }
+}
+
+int main()
+{
+ int arr1[40];
+#pragma omp target map(arr1[x ? 3 : 5])
+// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR <x != 0 \? 3 : 5>\] \[len: [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR <x != 0 \? 3 : 5>\] - \(long int\) &arr1\]\)} "original" } }
+ { }
+#pragma omp target map(arr1[x ? 3 : 5 : 5])
+// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR <x != 0 \? 3 : 5>\] \[len: [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR <x != 0 \? 3 : 5>\] - \(long int\) &arr1\]\)} "original" } }
+ { }
+#pragma omp target map(arr1[1 : x ? 3 : 5])
+// { dg-final { scan-tree-dump {map\(tofrom:arr1\[1\] [len: x != 0 ? [0-9]+ : [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: [0-9]+\]\)} "original" } }
+ { }
+
+ foo<3, 5> ();
+
+ return 0;
+}
+
new file mode 100644
@@ -0,0 +1,63 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-original" }
+
+int x, y;
+
+class C {
+ int x, y;
+
+public:
+ int foo();
+};
+
+int C::foo()
+{
+ int arr1[40];
+ /* There is a parsing ambiguity here without the space. We don't try to
+ resolve that automatically (though maybe we could, in theory). */
+#pragma omp target map(arr1[::x: ::y])
+// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR <x>\] \[len: \(sizetype\) y \* [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR <x>\] - \(long int\) &arr1\]\)} "original" } }
+ { }
+#pragma omp target map(arr1[::x:])
+// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR <x>\] \[len: \(40 - \(sizetype\) SAVE_EXPR <x>\) \* [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR <x>\] - \(long int\) &arr1\]\)} "original" } }
+ { }
+#pragma omp target map(arr1[: ::y])
+// { dg-final { scan-tree-dump {map\(tofrom:arr1\[0\] \[len: \(sizetype\) y \* [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: 0\]\)} "original" } }
+ { }
+ return ::x + ::y;
+}
+
+template<typename T>
+class Ct {
+ T x, y;
+
+public:
+ void foo();
+};
+
+template<typename T>
+void Ct<T>::foo()
+{
+ int arr1[40];
+#pragma omp target map(arr1[::x: ::y])
+// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR <x>\] \[len: \(sizetype\) y \* [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR <x>\] - \(long int\) &arr1\]\)} "original" } }
+ { }
+#pragma omp target map(arr1[::x:])
+// { dg-final { scan-tree-dump {map\(tofrom:arr1\[SAVE_EXPR <x>\] \[len: \(40 - \(sizetype\) SAVE_EXPR <x>\) \* [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: \(long int\) &arr1\[SAVE_EXPR <x>\] - \(long int\) &arr1\]\)} "original" } }
+ { }
+#pragma omp target map(arr1[: ::y])
+// { dg-final { scan-tree-dump {map\(tofrom:arr1\[0\] \[len: \(sizetype\) y \* [0-9]+\]\) map\(firstprivate:arr1 \[pointer assign, bias: 0\]\)} "original" } }
+ { }
+}
+
+int main()
+{
+ C c;
+ Ct<int> ct;
+
+ c.foo ();
+ ct.foo ();
+
+ return 0;
+}
+
new file mode 100644
@@ -0,0 +1,35 @@
+// { dg-do compile }
+
+int foo (int *ptr);
+
+template<typename T>
+T baz (T *ptr);
+
+template<typename T>
+void bar()
+{
+ T arr[20];
+
+#pragma omp target map(baz(arr[3:5]))
+// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 }
+// { dg-error {expected '\)' before ':' token} "" { target *-*-* } .-2 }
+// { dg-error {expected '\)' before '\]' token} "" { target *-*-* } .-3 }
+// { dg-error {expected '#pragma omp' clause before '\]' token} "" { target *-*-* } .-4 }
+ { }
+}
+
+int main()
+{
+ int arr[20];
+ // Reject array section as function argument.
+#pragma omp target map(foo(arr[3:5]))
+// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 }
+// { dg-error {expected '\)' before ':' token} "" { target *-*-* } .-2 }
+// { dg-error {expected '\)' before '\]' token} "" { target *-*-* } .-3 }
+// { dg-error {expected '#pragma omp' clause before '\]' token} "" { target *-*-* } .-4 }
+ { }
+
+ bar<int> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,35 @@
+// { dg-do compile }
+
+template<int C, int D>
+void foo()
+{
+ int arr1[40];
+#pragma omp target map(arr1[4,C:])
+// { dg-warning "top-level comma expression in array subscript is deprecated" "" { target c++20_only } .-1 }
+ { }
+#pragma omp target map(arr1[4,5:C,7])
+// { dg-warning "top-level comma expression in array subscript is deprecated" "" { target c++20_only } .-1 }
+ { }
+#pragma omp target map(arr1[:8,C,10])
+// { dg-warning "top-level comma expression in array subscript is deprecated" "" { target c++20_only } .-1 }
+ { }
+}
+
+int main()
+{
+ int arr1[40];
+#pragma omp target map(arr1[4,5:])
+// { dg-warning "top-level comma expression in array subscript is deprecated" "" { target c++20_only } .-1 }
+ { }
+#pragma omp target map(arr1[4,5:6,7])
+// { dg-warning "top-level comma expression in array subscript is deprecated" "" { target c++20_only } .-1 }
+ { }
+#pragma omp target map(arr1[:8,9,10])
+// { dg-warning "top-level comma expression in array subscript is deprecated" "" { target c++20_only } .-1 }
+ { }
+
+ foo<6, 9> ();
+
+ return 0;
+}
+
new file mode 100644
@@ -0,0 +1,36 @@
+// { dg-do compile }
+// { dg-additional-options "-std=c++23" }
+
+template<int C, int D>
+void foo()
+{
+ int arr1[40];
+#pragma omp target map(arr1[4,C:])
+// { dg-error "cannot use multidimensional subscript in OpenMP array section" "" { target *-*-* } .-1 }
+ { }
+#pragma omp target map(arr1[4,5:C,7])
+// { dg-error "cannot use multidimensional subscript in OpenMP array section" "" { target *-*-* } .-1 }
+ { }
+#pragma omp target map(arr1[:8,C,10])
+// { dg-error "cannot use multidimensional subscript in OpenMP array section" "" { target *-*-* } .-1 }
+ { }
+}
+
+int main()
+{
+ int arr1[40];
+#pragma omp target map(arr1[4,5:])
+// { dg-error "cannot use multidimensional subscript in OpenMP array section" "" { target *-*-* } .-1 }
+ { }
+#pragma omp target map(arr1[4,5:6,7])
+// { dg-error "cannot use multidimensional subscript in OpenMP array section" "" { target *-*-* } .-1 }
+ { }
+#pragma omp target map(arr1[:8,9,10])
+// { dg-error "cannot use multidimensional subscript in OpenMP array section" "" { target *-*-* } .-1 }
+ { }
+
+ foo<6, 9> ();
+
+ return 0;
+}
+
new file mode 100644
@@ -0,0 +1,33 @@
+// { dg-do compile }
+// { dg-additional-options "-std=c++11" }
+
+template<typename T>
+void foo()
+{
+ T arr[20];
+ // Reject array section in lambda function.
+#pragma omp target map([&](const int x) -> T* { return arr[0:x]; } (5))
+// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 }
+// { dg-error {invalid conversion from 'int' to 'int\*'} "" { target *-*-* } .-2 }
+// { dg-error {expected ';' before ':' token} "" { target *-*-* } .-3 }
+// { dg-error {expected primary-expression before ':' token} "" { target *-*-* } .-4 }
+// { dg-message {sorry, unimplemented: unsupported map expression '<lambda closure object>foo<int>\(\)::<lambda\(int\)>\{arr\}.foo<int>\(\)::<lambda\(int\)>\(5\)'} "" { target *-*-* } .-5 }
+ { }
+}
+
+int main()
+{
+ int arr[20];
+ // Reject array section in lambda function.
+#pragma omp target map([&](const int x) -> int* { return arr[0:x]; } (5))
+// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 }
+// { dg-error {invalid conversion from 'int' to 'int\*'} "" { target *-*-* } .-2 }
+// { dg-error {expected ';' before ':' token} "" { target *-*-* } .-3 }
+// { dg-error {expected primary-expression before ':' token} "" { target *-*-* } .-4 }
+// { dg-message {sorry, unimplemented: unsupported map expression '<lambda closure object>main\(\)::<lambda\(int\)>\{arr\}.main\(\)::<lambda\(int\)>\(5\)'} "" { target *-*-* } .-5 }
+ { }
+
+ foo<int> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,28 @@
+// { dg-do compile }
+
+template<typename T>
+void foo()
+{
+ T arr[20];
+ // Reject array section in statement expression.
+#pragma omp target map( ({ int x = 5; arr[0:x]; }) )
+// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 }
+// { dg-error {expected ';' before ':' token} "" { target *-*-* } .-2 }
+// { dg-message {sorry, unimplemented: unsupported map expression '\(\{\.\.\.\}\)'} "" { target *-*-* } .-3 }
+ { }
+}
+
+int main()
+{
+ int arr[20];
+ // Reject array section in statement expression.
+#pragma omp target map( ({ int x = 5; arr[0:x]; }) )
+// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 }
+// { dg-error {expected ';' before ':' token} "" { target *-*-* } .-2 }
+// { dg-message {sorry, unimplemented: unsupported map expression '\(\{\.\.\.\}\)'} "" { target *-*-* } .-3 }
+ { }
+
+ foo<int> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,50 @@
+// { dg-do compile }
+
+template<typename T>
+struct St {
+ T *ptr;
+};
+
+template<typename T>
+void foo()
+{
+ T arr[20];
+
+ // Reject array section in compound initialiser.
+#pragma omp target map( (struct St<T>) { .ptr = (T *) arr[5:5] } )
+// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 }
+// { dg-error {expected primary-expression before 'struct'} "" { target *-*-* } .-2 }
+// { dg-error {expected '\)' before 'struct'} "" { target *-*-* } .-3 }
+ { }
+
+ // ...and this is unsupported too (probably not useful anyway).
+#pragma omp target map( (struct St<T>) { .ptr = &arr[5] } )
+// { dg-message {sorry, unimplemented: unsupported map expression 'St<int>\{\(\& arr\[5\]\)\}'} "" { target *-*-* } .-1 }
+ { }
+}
+
+struct S {
+ int *ptr;
+};
+
+int main()
+{
+ int arr[20];
+
+ // Reject array section in compound initialiser.
+#pragma omp target map( (struct S) { .ptr = (int *) arr[5:5] } )
+// { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 }
+// { dg-warning {cast to pointer from integer of different size} "" { target *-*-* } .-2 }
+// { dg-error {expected primary-expression before 'struct'} "" { target *-*-* } .-3 }
+// { dg-error {expected '\)' before 'struct'} "" { target *-*-* } .-4 }
+ { }
+
+ // ...and this is unsupported too (probably not useful anyway).
+#pragma omp target map( (struct S) { .ptr = &arr[5] } )
+// { dg-message {sorry, unimplemented: unsupported map expression 'S\{\(\& arr\[5\]\)\}'} "" { target *-*-* } .-1 }
+ { }
+
+ foo<int> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,50 @@
+// { dg-do compile }
+
+int x;
+
+template<typename T>
+void foo()
+{
+ T arr[20];
+ T *ptr;
+ /* "arr[1:10]" looks like it might be an expression of array type, hence
+ able to be indexed (again). This isn't allowed, though. */
+#pragma omp target map(arr[1:10][2])
+// { dg-error {'arr\[1\]' does not have pointer or array type} "" { target *-*-* } .-1 }
+ { }
+#pragma omp target map(arr[1:x][2])
+// { dg-error {'arr\[1\]' does not have pointer or array type} "" { target *-*-* } .-1 }
+ { }
+ // ...and nor is this.
+#pragma omp target map(ptr[1:10][2])
+// { dg-error {'\*\(ptr \+ [0-9]+\)' does not have pointer or array type} "" { target *-*-* } .-1 }
+ { }
+#pragma omp target map(ptr[1:x][2])
+// { dg-error {'\*\(ptr \+ [0-9]+\)' does not have pointer or array type} "" { target *-*-* } .-1 }
+ { }
+}
+
+int main()
+{
+ int arr[20];
+ int *ptr;
+ /* "arr[1:10]" looks like it might be an expression of array type, hence
+ able to be indexed (again). This isn't allowed, though. */
+#pragma omp target map(arr[1:10][2])
+// { dg-error {'arr\[1\]' does not have pointer or array type} "" { target *-*-* } .-1 }
+ { }
+#pragma omp target map(arr[1:x][2])
+// { dg-error {'arr\[1\]' does not have pointer or array type} "" { target *-*-* } .-1 }
+ { }
+ // ...and nor is this.
+#pragma omp target map(ptr[1:10][2])
+// { dg-error {'\*\(ptr \+ [0-9]+\)' does not have pointer or array type} "" { target *-*-* } .-1 }
+ { }
+#pragma omp target map(ptr[1:x][2])
+// { dg-error {'\*\(ptr \+ [0-9]+\)' does not have pointer or array type} "" { target *-*-* } .-1 }
+ { }
+
+ foo<int> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,24 @@
+// { dg-do compile }
+
+bool partly = false;
+
+template<typename T>
+void foo()
+{
+ T arr[20];
+#pragma omp target map(partly ? arr[5:5] : arr)
+// { dg-message {sorry, unimplemented: unsupported map expression '\(partly \? \(\(int\*\)\(\& arr\[5:5\]\)\) : \(\(int\*\)\(\& arr\)\)\)'} "" { target *-*-* } .-1 }
+ { }
+}
+
+int main()
+{
+ int arr[20];
+#pragma omp target map(partly ? arr[5:5] : arr)
+// { dg-message {sorry, unimplemented: unsupported map expression '\(partly \? \(\(int\*\)\(\& arr\[5:5\]\)\) : \(\(int\*\)\(\& arr\)\)\)'} "" { target *-*-* } .-1 }
+ { }
+
+ foo<int> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,36 @@
+// { dg-do compile }
+
+int x;
+
+template<typename T>
+void foo()
+{
+ T arr[20];
+ // Here we know the type of the array section (the upper bound is reported)...
+#pragma omp target map(arr[5:5] * 2)
+// { dg-error {invalid operands of types 'int \[10\]' and 'int'} "" { target *-*-* } .-1 }
+ { }
+ /* ...but here, we have an incomplete array type because of the variable
+ low bound 'x'. */
+#pragma omp target map(arr[x:5] * 2)
+// { dg-error {invalid operands of types 'int \[\]' and 'int'} "" { target *-*-* } .-1 }
+ { }
+}
+
+int main()
+{
+ int arr[20];
+ // Here we know the type of the array section (the upper bound is reported)...
+#pragma omp target map(arr[5:5] * 2)
+// { dg-error {invalid operands of types 'int \[10\]' and 'int'} "" { target *-*-* } .-1 }
+ { }
+ /* ...but here, we have an incomplete array type because of the variable
+ low bound 'x'. */
+#pragma omp target map(arr[x:5] * 2)
+// { dg-error {invalid operands of types 'int \[\]' and 'int'} "" { target *-*-* } .-1 }
+ { }
+
+ foo<int> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,53 @@
+// { dg-do compile }
+
+int x;
+
+template<typename X>
+struct Tt {
+ X arr[20];
+};
+
+template<typename X>
+struct St {
+ X *tvec;
+};
+
+template<typename T>
+void foo()
+{
+ struct St<Tt<T> > *s;
+ // You can't use an array section like this. Make sure sensible errors are
+ // reported.
+#pragma omp target map(s->tvec[3:5].arr[0:20])
+// { dg-error {request for member 'arr' in 's->St<Tt<int> >::tvec\[3:5\]', which is of non-class type 'Tt<int> \[8\]'} "" { target *-*-* } .-1 }
+ { }
+#pragma omp target map(s->tvec[5:x].arr[0:20])
+// { dg-error {invalid use of array with unspecified bounds} "" { target *-*-* } .-1 }
+ { }
+}
+
+struct T {
+ int arr[20];
+};
+
+struct S {
+ struct T *tvec;
+};
+
+int main()
+{
+ struct S *s;
+ // You can't use an array section like this. Make sure sensible errors are
+ // reported.
+#pragma omp target map(s->tvec[3:5].arr[0:20])
+// { dg-error {request for member 'arr' in 's->S::tvec\[3:5\]', which is of non-class type 'T \[8\]'} "" { target *-*-* } .-1 }
+ { }
+#pragma omp target map(s->tvec[5:x].arr[0:20])
+// { dg-error {invalid use of array with unspecified bounds} "" { target *-*-* } .-1 }
+// { dg-error {expected '\)' before 'arr'} "" { target *-*-* } .-2 }
+ { }
+
+ foo<int> ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,39 @@
+// { dg-do compile }
+
+int x;
+
+template<typename T>
+void foo()
+{
+ T arr1[40];
+ T arr2[40];
+#pragma omp target map(arr1[arr2[4:5]:arr2[6:7]])
+// { dg-error {low bound 'arr2\[4:5\]' of array section does not have integral type} "" { target *-*-* } .-1 }
+ { }
+#pragma omp target map(arr1[arr2[:1]:arr2[6:1]])
+// { dg-error {low bound 'arr2\[:1\]' of array section does not have integral type} "" { target *-*-* } .-1 }
+ { }
+#pragma omp target map(arr1[x:arr2[6:1]])
+// { dg-error {length 'arr2\[6:1\]' of array section does not have integral type} "" { target *-*-* } .-1 }
+ { }
+}
+
+int main()
+{
+ int arr1[40];
+ int arr2[40];
+#pragma omp target map(arr1[arr2[4:5]:arr2[6:7]])
+// { dg-error {low bound 'arr2\[4:5\]' of array section does not have integral type} "" { target *-*-* } .-1 }
+ { }
+#pragma omp target map(arr1[arr2[:1]:arr2[6:1]])
+// { dg-error {low bound 'arr2\[:1\]' of array section does not have integral type} "" { target *-*-* } .-1 }
+ { }
+#pragma omp target map(arr1[x:arr2[6:1]])
+// { dg-error {length 'arr2\[6:1\]' of array section does not have integral type} "" { target *-*-* } .-1 }
+ { }
+
+ foo<int> ();
+
+ return 0;
+}
+
new file mode 100644
@@ -0,0 +1,36 @@
+// { dg-do compile }
+
+#include <cstdio>
+#include <cstring>
+#include <cassert>
+
+typedef struct {
+ int arr[100];
+} S;
+
+int main()
+{
+ S *s = new S;
+
+ memset (s->arr, '\0', sizeof s->arr);
+
+#pragma omp target enter data map(to: (*s).arr)
+ /* You can't do this, at least as of OpenMP 5.2. "has_device_addr" takes
+ a "variable list" item type
+ (OpenMP 5.2, "5.4.9 has_device_addr Clause"). */
+#pragma omp target has_device_addr((*s).arr[5:20])
+// { dg-error {expected unqualified-id before '\(' token} "" { target *-*-* } .-1 }
+ {
+ for (int i = 5; i < 25; i++)
+ s->arr[i] = i;
+ }
+
+#pragma omp target exit data map(from: (*s).arr)
+
+ for (int i = 0; i < 100; i++)
+ assert (i >= 5 && i < 25 ? s->arr[i] == i : s->arr[i] == 0);
+
+ delete s;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,37 @@
+#include <cassert>
+
+struct S {
+ int x[10];
+};
+
+S *
+choose (S *a, S *b, int c)
+{
+ if (c < 5)
+ return a;
+ else
+ return b;
+}
+
+int main (int argc, char *argv[])
+{
+ S a, b;
+
+ for (int i = 0; i < 10; i++)
+ a.x[i] = b.x[i] = 0;
+
+ for (int i = 0; i < 10; i++)
+ {
+#pragma omp target map(choose(&a, &b, i)->x[:10])
+/* { dg-message {sorry, unimplemented: unsupported map expression 'choose\(\(& a\), \(& b\), i\)->S::x\[0\]'} "" { target *-*-* } .-1 } */
+ for (int j = 0; j < 10; j++)
+ choose (&a, &b, i)->x[j]++;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (a.x[i] == 5 && b.x[i] == 5);
+
+ return 0;
+}
+
+
new file mode 100644
@@ -0,0 +1,12 @@
+#include <cassert>
+
+int main (int argc, char *argv[])
+{
+ int a = 5, b = 2;
+#pragma omp target map(a += b)
+ /* { dg-message {sorry, unimplemented: unsupported map expression '\(a = \(a \+ b\)\)'} "" { target *-*-* } .-1 } */
+ {
+ a++;
+ }
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,10 @@
+int main (int argc, char *argv[])
+{
+ int a = 5;
+#pragma omp target map(++a)
+ /* { dg-message {sorry, unimplemented: unsupported map expression '\+\+ a'} "" { target *-*-* } .-1 } */
+ {
+ a++;
+ }
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,19 @@
+#include <cassert>
+
+int glob = 10;
+
+int& foo ()
+{
+ return glob;
+}
+
+int main (int argc, char *argv[])
+{
+#pragma omp target map(foo())
+ /* { dg-message {sorry, unimplemented: unsupported map expression 'foo\(\)'} "" { target *-*-* } .-1 } */
+ {
+ foo()++;
+ }
+ assert (glob == 11);
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,37 @@
+#include <cassert>
+
+struct S {
+ int x;
+ int *ptr;
+};
+
+int
+main (int argc, char *argv[])
+{
+ S s;
+ int S::* xp = &S::x;
+ int* S::* ptrp = &S::ptr;
+
+ s.ptr = new int[64];
+
+ s.*xp = 6;
+ for (int i = 0; i < 64; i++)
+ (s.*ptrp)[i] = i;
+
+#pragma omp target map(s.*xp, s.*ptrp, (s.*ptrp)[:64])
+ /* { dg-message {sorry, unimplemented: pointer-to-member mapping '\*\(\*\(\(\(int\*\*\)\(& s\)\) \+ \(\(sizetype\)ptrp\)\)\)' not supported} "" { target *-*-* } .-1 } */
+ /* { dg-message {sorry, unimplemented: pointer-to-member mapping '\*\(\(\(int\*\*\)\(& s\)\) \+ \(\(sizetype\)ptrp\)\)' not supported} "" { target *-*-* } .-2 } */
+ /* { dg-message {sorry, unimplemented: pointer-to-member mapping '\*\(\(\(int\*\)\(& s\)\) \+ \(\(sizetype\)xp\)\)' not supported} "" { target *-*-* } .-3 } */
+#pragma omp teams distribute parallel for
+ for (int i = 0; i < 64; i++)
+ {
+ (s.*xp)++;
+ (s.*ptrp)[i]++;
+ }
+
+ assert (s.*xp == 70);
+ for (int i = 0; i < 64; i++)
+ assert ((s.*ptrp)[i] == i + 1);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,40 @@
+#include <cassert>
+
+struct S {
+ int x;
+ int *ptr;
+};
+
+int
+main (int argc, char *argv[])
+{
+ S *s = new S;
+ int S::* xp = &S::x;
+ int* S::* ptrp = &S::ptr;
+
+ s->ptr = new int[64];
+
+ s->*xp = 4;
+ for (int i = 0; i < 64; i++)
+ (s->*ptrp)[i] = i;
+
+#pragma omp target map(s->*xp, s->*ptrp, (s->*ptrp)[:64])
+ /* { dg-message {sorry, unimplemented: pointer-to-member mapping '\*\(\(\(int\*\*\)s\) \+ \(\(sizetype\)ptrp\)\)' not supported} "" { target *-*-* } .-1 } */
+ /* { dg-message {sorry, unimplemented: pointer-to-member mapping '\*\(\(\(int\*\)s\) \+ \(\(sizetype\)xp\)\)' not supported} "" { target *-*-* } .-2 } */
+ /* { dg-message {sorry, unimplemented: pointer-to-member mapping '\*\(\*\(\(\(int\*\*\)s\) \+ \(\(sizetype\)ptrp\)\)\)' not supported} "" { target *-*-* } .-3 } */
+#pragma omp teams distribute parallel for
+ for (int i = 0; i < 64; i++)
+ {
+ (s->*xp)++;
+ (s->*ptrp)[i]++;
+ }
+
+ assert (s->*xp == 68);
+ for (int i = 0; i < 64; i++)
+ assert ((s->*ptrp)[i] == i + 1);
+
+ delete s->ptr;
+ delete s;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,17 @@
+#include <cassert>
+
+int foo (int x)
+{
+#pragma omp target map(static_cast<int&>(x))
+ /* { dg-message {sorry, unimplemented: unsupported map expression '& x'} "" { target *-*-* } .-1 } */
+ {
+ x += 3;
+ }
+ return x;
+}
+
+int main (int argc, char *argv[])
+{
+ assert (foo (5) == 8);
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,20 @@
+#include <cassert>
+
+int foo (bool yesno)
+{
+ int x = 5, y = 7;
+#pragma omp target map(yesno ? x : y)
+ /* { dg-message {sorry, unimplemented: unsupported map expression '\(yesno \? x : y\)'} "" { target *-*-* } .-1 } */
+ {
+ x += 3;
+ y += 5;
+ }
+ return yesno ? x : y;
+}
+
+int main (int argc, char *argv[])
+{
+ assert (foo (true) == 8);
+ assert (foo (false) == 12);
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,91 @@
+#include <cassert>
+
+typedef int intarr100[100];
+
+class C {
+ int arr[100];
+ int *ptr;
+
+public:
+ C();
+ ~C();
+ void zero ();
+ void do_operation ();
+ void check (int, int);
+ intarr100 &get_arr () { return arr; }
+ int *get_ptr() { return ptr; }
+};
+
+C::C()
+{
+ ptr = new int[100];
+ for (int i = 0; i < 100; i++)
+ arr[i] = 0;
+}
+
+C::~C()
+{
+ delete ptr;
+}
+
+void
+C::zero ()
+{
+ for (int i = 0; i < 100; i++)
+ arr[i] = ptr[i] = 0;
+}
+
+void
+C::do_operation ()
+{
+#pragma omp target map(arr, ptr, ptr[:100])
+#pragma omp teams distribute parallel for
+ for (int i = 0; i < 100; i++)
+ {
+ arr[i] = arr[i] + 3;
+ ptr[i] = ptr[i] + 5;
+ }
+}
+
+void
+C::check (int arrval, int ptrval)
+{
+ for (int i = 0; i < 100; i++)
+ {
+ assert (arr[i] == arrval);
+ assert (ptr[i] == ptrval);
+ }
+}
+
+int
+main (int argc, char *argv[])
+{
+ C c;
+
+ c.zero ();
+ c.do_operation ();
+ c.check (3, 5);
+
+ /* It might sort of make sense to be able to do this, but we don't support
+ it for now. */
+ #pragma omp target map(c.get_arr()[:100])
+ /* { dg-message {sorry, unimplemented: unsupported map expression 'c\.C::get_arr\(\)\[0\]'} "" { target *-*-* } .-1 } */
+ #pragma omp teams distribute parallel for
+ for (int i = 0; i < 100; i++)
+ c.get_arr()[i] += 2;
+
+ c.check (5, 5);
+
+ /* Same for this. */
+ #pragma omp target map(c.get_ptr(), c.get_ptr()[:100])
+ /* { dg-message {sorry, unimplemented: unsupported map expression 'c\.C::get_ptr\(\)'} "" { target *-*-* } .-1 } */
+ /* { dg-message {sorry, unimplemented: unsupported map expression '\* c\.C::get_ptr\(\)'} "" { target *-*-* } .-2 } */
+ #pragma omp teams distribute parallel for
+ for (int i = 0; i < 100; i++)
+ c.get_ptr()[i] += 3;
+
+ c.check (5, 8);
+
+ return 0;
+}
+
@@ -12,7 +12,7 @@ foo (void)
for (int i = 0; i < 16; i++)
;
- #pragma omp target map (S[0:10]) // { dg-error "is not a variable in" }
+ #pragma omp target map (S[0:10]) // { dg-error "expected primary-expression before '\\\[' token" }
;
#pragma omp task depend (inout: S[0:10]) // { dg-error "is not a variable in" }
@@ -2549,6 +2549,20 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags,
}
break;
+ case OMP_ARRAY_SECTION:
+ op0 = TREE_OPERAND (node, 0);
+ if (op_prio (op0) < op_prio (node))
+ pp_left_paren (pp);
+ dump_generic_node (pp, op0, spc, flags, false);
+ if (op_prio (op0) < op_prio (node))
+ pp_right_paren (pp);
+ pp_left_bracket (pp);
+ 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);
+ pp_right_bracket (pp);
+ break;
+
case CONSTRUCTOR:
{
unsigned HOST_WIDE_INT ix;
@@ -1313,6 +1313,9 @@ DEFTREECODE (OMP_ATOMIC_CAPTURE_NEW, "omp_atomic_capture_new", tcc_statement, 2)
/* OpenMP clauses. */
DEFTREECODE (OMP_CLAUSE, "omp_clause", tcc_exceptional, 0)
+/* An OpenMP array section. */
+DEFTREECODE (OMP_ARRAY_SECTION, "omp_array_section", tcc_expression, 3)
+
/* TRANSACTION_EXPR tree code.
Operand 0: BODY: contains body of the transaction. */
DEFTREECODE (TRANSACTION_EXPR, "transaction_expr", tcc_expression, 1)
@@ -11,11 +11,9 @@
#define REF2PTR_DECL_BASE
#define ARRAY_DECL_BASE
-// Needs map clause "lvalue"-parsing support.
-//#define REF2ARRAY_DECL_BASE
+#define REF2ARRAY_DECL_BASE
#define PTR_OFFSET_DECL_BASE
-// Needs map clause "lvalue"-parsing support.
-//#define REF2PTR_OFFSET_DECL_BASE
+#define REF2PTR_OFFSET_DECL_BASE
#define MAP_SECTIONS
@@ -30,25 +28,21 @@
#define ARRAY_DECL_MEMBER_SLICE
#define ARRAY_DECL_MEMBER_SLICE_BASEPTR
-// Needs map clause "lvalue"-parsing support.
-//#define REF2ARRAY_DECL_MEMBER_SLICE
-//#define REF2ARRAY_DECL_MEMBER_SLICE_BASEPTR
+#define REF2ARRAY_DECL_MEMBER_SLICE
+#define REF2ARRAY_DECL_MEMBER_SLICE_BASEPTR
#define PTR_OFFSET_DECL_MEMBER_SLICE
#define PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
-// Needs map clause "lvalue"-parsing support.
-//#define REF2PTR_OFFSET_DECL_MEMBER_SLICE
-//#define REF2PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
+#define REF2PTR_OFFSET_DECL_MEMBER_SLICE
+#define REF2PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
#define PTRARRAY_DECL_MEMBER_SLICE
#define PTRARRAY_DECL_MEMBER_SLICE_BASEPTR
-// Needs map clause "lvalue"-parsing support.
-//#define REF2PTRARRAY_DECL_MEMBER_SLICE
-//#define REF2PTRARRAY_DECL_MEMBER_SLICE_BASEPTR
+#define REF2PTRARRAY_DECL_MEMBER_SLICE
+#define REF2PTRARRAY_DECL_MEMBER_SLICE_BASEPTR
#define PTRPTR_OFFSET_DECL_MEMBER_SLICE
#define PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
-// Needs map clause "lvalue"-parsing support.
-//#define REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE
-//#define REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
+#define REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE
+#define REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
#define NONREF_COMPONENT_BASE
#define NONREF_COMPONENT_MEMBER_SLICE
new file mode 100644
@@ -0,0 +1,3199 @@
+// { dg-do run }
+
+// This is essentially baseptrs-4.C with templates.
+
+#include <cstring>
+#include <cassert>
+
+#define MAP_DECLS
+
+#define NONREF_DECL_BASE
+#define REF_DECL_BASE
+#define PTR_DECL_BASE
+#define REF2PTR_DECL_BASE
+
+#define ARRAY_DECL_BASE
+#define REF2ARRAY_DECL_BASE
+#define PTR_OFFSET_DECL_BASE
+#define REF2PTR_OFFSET_DECL_BASE
+
+#define MAP_SECTIONS
+
+#define NONREF_DECL_MEMBER_SLICE
+#define NONREF_DECL_MEMBER_SLICE_BASEPTR
+#define REF_DECL_MEMBER_SLICE
+#define REF_DECL_MEMBER_SLICE_BASEPTR
+#define PTR_DECL_MEMBER_SLICE
+#define PTR_DECL_MEMBER_SLICE_BASEPTR
+#define REF2PTR_DECL_MEMBER_SLICE
+#define REF2PTR_DECL_MEMBER_SLICE_BASEPTR
+
+#define ARRAY_DECL_MEMBER_SLICE
+#define ARRAY_DECL_MEMBER_SLICE_BASEPTR
+#define REF2ARRAY_DECL_MEMBER_SLICE
+#define REF2ARRAY_DECL_MEMBER_SLICE_BASEPTR
+#define PTR_OFFSET_DECL_MEMBER_SLICE
+#define PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
+#define REF2PTR_OFFSET_DECL_MEMBER_SLICE
+#define REF2PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
+
+#define PTRARRAY_DECL_MEMBER_SLICE
+#define PTRARRAY_DECL_MEMBER_SLICE_BASEPTR
+#define REF2PTRARRAY_DECL_MEMBER_SLICE
+#define REF2PTRARRAY_DECL_MEMBER_SLICE_BASEPTR
+#define PTRPTR_OFFSET_DECL_MEMBER_SLICE
+#define PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
+#define REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE
+#define REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
+
+#define NONREF_COMPONENT_BASE
+#define NONREF_COMPONENT_MEMBER_SLICE
+#define NONREF_COMPONENT_MEMBER_SLICE_BASEPTR
+
+#define REF_COMPONENT_BASE
+#define REF_COMPONENT_MEMBER_SLICE
+#define REF_COMPONENT_MEMBER_SLICE_BASEPTR
+
+#define PTR_COMPONENT_BASE
+#define PTR_COMPONENT_MEMBER_SLICE
+#define PTR_COMPONENT_MEMBER_SLICE_BASEPTR
+
+#define REF2PTR_COMPONENT_BASE
+#define REF2PTR_COMPONENT_MEMBER_SLICE
+#define REF2PTR_COMPONENT_MEMBER_SLICE_BASEPTR
+
+#ifdef MAP_DECLS
+template<int N>
+void
+map_decls (void)
+{
+ int x = 0;
+ int &y = x;
+ int arr[4];
+ int (&arrref)[N] = arr;
+ int *z = &arr[0];
+ int *&t = z;
+
+ memset (arr, 0, sizeof arr);
+
+ #pragma omp target map(x)
+ {
+ x++;
+ }
+
+ #pragma omp target map(y)
+ {
+ y++;
+ }
+
+ assert (x == 2);
+ assert (y == 2);
+
+ /* "A variable that is of type pointer is treated as if it is the base
+ pointer of a zero-length array section that appeared as a list item in a
+ map clause." */
+ #pragma omp target map(z)
+ {
+ z++;
+ }
+
+ /* "A variable that is of type reference to pointer is treated as if it had
+ appeared in a map clause as a zero-length array section."
+
+ The pointer here is *not* associated with a target address, so we're not
+ disallowed from modifying it. */
+ #pragma omp target map(t)
+ {
+ t++;
+ }
+
+ assert (z == &arr[2]);
+ assert (t == &arr[2]);
+
+ #pragma omp target map(arr)
+ {
+ arr[2]++;
+ }
+
+ #pragma omp target map(arrref)
+ {
+ arrref[2]++;
+ }
+
+ assert (arr[2] == 2);
+ assert (arrref[2] == 2);
+}
+#endif
+
+template<typename T>
+struct S {
+ T a;
+ T &b;
+ T *c;
+ T *&d;
+ T e[4];
+ T (&f)[4];
+
+ S(T a1, T &b1, T *c1, T *&d1) :
+ a(a1), b(b1), c(c1), d(d1), f(e)
+ {
+ memset (e, 0, sizeof e);
+ }
+};
+
+#ifdef NONREF_DECL_BASE
+template<typename T>
+void
+nonref_decl_base (void)
+{
+ T a = 0, b = 0, c, *d = &c;
+ S<T> mys(a, b, &c, d);
+
+ #pragma omp target map(mys.a)
+ {
+ mys.a++;
+ }
+
+ #pragma omp target map(mys.b)
+ {
+ mys.b++;
+ }
+
+ assert (mys.a == 1);
+ assert (mys.b == 1);
+
+ #pragma omp target map(mys.c)
+ {
+ mys.c++;
+ }
+
+ #pragma omp target map(mys.d)
+ {
+ mys.d++;
+ }
+
+ assert (mys.c == &c + 1);
+ assert (mys.d == &c + 1);
+
+ #pragma omp target map(mys.e)
+ {
+ mys.e[0]++;
+ }
+
+ #pragma omp target map(mys.f)
+ {
+ mys.f[0]++;
+ }
+
+ assert (mys.e[0] == 2);
+ assert (mys.f[0] == 2);
+}
+#endif
+
+#ifdef REF_DECL_BASE
+template<typename T>
+void
+ref_decl_base (void)
+{
+ T a = 0, b = 0, c, *d = &c;
+ S<T> mys_orig(a, b, &c, d);
+ S<T> &mys = mys_orig;
+
+ #pragma omp target map(mys.a)
+ {
+ mys.a++;
+ }
+
+ #pragma omp target map(mys.b)
+ {
+ mys.b++;
+ }
+
+ assert (mys.a == 1);
+ assert (mys.b == 1);
+
+ #pragma omp target map(mys.c)
+ {
+ mys.c++;
+ }
+
+ #pragma omp target map(mys.d)
+ {
+ mys.d++;
+ }
+
+ assert (mys.c == &c + 1);
+ assert (mys.d == &c + 1);
+
+ #pragma omp target map(mys.e)
+ {
+ mys.e[0]++;
+ }
+
+ #pragma omp target map(mys.f)
+ {
+ mys.f[0]++;
+ }
+
+ assert (mys.e[0] == 2);
+ assert (mys.f[0] == 2);
+}
+#endif
+
+#ifdef PTR_DECL_BASE
+template<typename A>
+void
+ptr_decl_base (void)
+{
+ A a = 0, b = 0, c, *d = &c;
+ S<A> mys_orig(a, b, &c, d);
+ S<A> *mys = &mys_orig;
+
+ #pragma omp target map(mys->a)
+ {
+ mys->a++;
+ }
+
+ #pragma omp target map(mys->b)
+ {
+ mys->b++;
+ }
+
+ assert (mys->a == 1);
+ assert (mys->b == 1);
+
+ #pragma omp target map(mys->c)
+ {
+ mys->c++;
+ }
+
+ #pragma omp target map(mys->d)
+ {
+ mys->d++;
+ }
+
+ assert (mys->c == &c + 1);
+ assert (mys->d == &c + 1);
+
+ #pragma omp target map(mys->e)
+ {
+ mys->e[0]++;
+ }
+
+ #pragma omp target map(mys->f)
+ {
+ mys->f[0]++;
+ }
+
+ assert (mys->e[0] == 2);
+ assert (mys->f[0] == 2);
+}
+#endif
+
+#ifdef REF2PTR_DECL_BASE
+template<typename A>
+void
+ref2ptr_decl_base (void)
+{
+ A a = 0, b = 0, c, *d = &c;
+ S<A> mys_orig(a, b, &c, d);
+ S<A> *mysp = &mys_orig;
+ S<A> *&mys = mysp;
+
+ #pragma omp target map(mys->a)
+ {
+ mys->a++;
+ }
+
+ #pragma omp target map(mys->b)
+ {
+ mys->b++;
+ }
+
+ assert (mys->a == 1);
+ assert (mys->b == 1);
+
+ #pragma omp target map(mys->c)
+ {
+ mys->c++;
+ }
+
+ #pragma omp target map(mys->d)
+ {
+ mys->d++;
+ }
+
+ assert (mys->c == &c + 1);
+ assert (mys->d == &c + 1);
+
+ #pragma omp target map(mys->e)
+ {
+ mys->e[0]++;
+ }
+
+ #pragma omp target map(mys->f)
+ {
+ mys->f[0]++;
+ }
+
+ assert (mys->e[0] == 2);
+ assert (mys->f[0] == 2);
+}
+#endif
+
+#ifdef ARRAY_DECL_BASE
+template<typename A>
+void
+array_decl_base (void)
+{
+ A a = 0, b = 0, c, *d = &c;
+ S<A> mys[4] =
+ {
+ S<A>(a, b, &c, d),
+ S<A>(a, b, &c, d),
+ S<A>(a, b, &c, d),
+ S<A>(a, b, &c, d)
+ };
+
+ #pragma omp target map(mys[2].a)
+ {
+ mys[2].a++;
+ }
+
+ #pragma omp target map(mys[2].b)
+ {
+ mys[2].b++;
+ }
+
+ assert (mys[2].a == 1);
+ assert (mys[2].b == 1);
+
+ #pragma omp target map(mys[2].c)
+ {
+ mys[2].c++;
+ }
+
+ #pragma omp target map(mys[2].d)
+ {
+ mys[2].d++;
+ }
+
+ assert (mys[2].c == &c + 1);
+ assert (mys[2].d == &c + 1);
+
+ #pragma omp target map(mys[2].e)
+ {
+ mys[2].e[0]++;
+ }
+
+ #pragma omp target map(mys[2].f)
+ {
+ mys[2].f[0]++;
+ }
+
+ assert (mys[2].e[0] == 2);
+ assert (mys[2].f[0] == 2);
+}
+#endif
+
+#ifdef REF2ARRAY_DECL_BASE
+template<typename A>
+void
+ref2array_decl_base (void)
+{
+ A a = 0, b = 0, c, *d = &c;
+ S<A> mys_orig[4] =
+ {
+ S<A>(a, b, &c, d),
+ S<A>(a, b, &c, d),
+ S<A>(a, b, &c, d),
+ S<A>(a, b, &c, d)
+ };
+ S<A> (&mys)[4] = mys_orig;
+
+ #pragma omp target map(mys[2].a)
+ {
+ mys[2].a++;
+ }
+
+ #pragma omp target map(mys[2].b)
+ {
+ mys[2].b++;
+ }
+
+ assert (mys[2].a == 1);
+ assert (mys[2].b == 1);
+
+ #pragma omp target map(mys[2].c)
+ {
+ mys[2].c++;
+ }
+
+ #pragma omp target map(mys[2].d)
+ {
+ mys[2].d++;
+ }
+
+ assert (mys[2].c == &c + 1);
+ assert (mys[2].d == &c + 1);
+
+ #pragma omp target map(mys[2].e)
+ {
+ mys[2].e[0]++;
+ }
+
+ #pragma omp target map(mys[2].f)
+ {
+ mys[2].f[0]++;
+ }
+
+ assert (mys[2].e[0] == 2);
+ assert (mys[2].f[0] == 2);
+}
+#endif
+
+#ifdef PTR_OFFSET_DECL_BASE
+template<typename A>
+void
+ptr_offset_decl_base (void)
+{
+ A a = 0, b = 0, c, *d = &c;
+ S<A> mys_orig[4] =
+ {
+ S<A>(a, b, &c, d),
+ S<A>(a, b, &c, d),
+ S<A>(a, b, &c, d),
+ S<A>(a, b, &c, d)
+ };
+ S<A> *mys = &mys_orig[0];
+
+ #pragma omp target map(mys[2].a)
+ {
+ mys[2].a++;
+ }
+
+ #pragma omp target map(mys[2].b)
+ {
+ mys[2].b++;
+ }
+
+ assert (mys[2].a == 1);
+ assert (mys[2].b == 1);
+
+ #pragma omp target map(mys[2].c)
+ {
+ mys[2].c++;
+ }
+
+ #pragma omp target map(mys[2].d)
+ {
+ mys[2].d++;
+ }
+
+ assert (mys[2].c == &c + 1);
+ assert (mys[2].d == &c + 1);
+
+ #pragma omp target map(mys[2].e)
+ {
+ mys[2].e[0]++;
+ }
+
+ #pragma omp target map(mys[2].f)
+ {
+ mys[2].f[0]++;
+ }
+
+ assert (mys[2].e[0] == 2);
+ assert (mys[2].f[0] == 2);
+}
+#endif
+
+#ifdef REF2PTR_OFFSET_DECL_BASE
+template<typename A>
+void
+ref2ptr_offset_decl_base (void)
+{
+ A a = 0, b = 0, c, *d = &c;
+ S<A> mys_orig[4] =
+ {
+ S<A>(a, b, &c, d),
+ S<A>(a, b, &c, d),
+ S<A>(a, b, &c, d),
+ S<A>(a, b, &c, d)
+ };
+ S<A> *mys_ptr = &mys_orig[0];
+ S<A> *&mys = mys_ptr;
+
+ #pragma omp target map(mys[2].a)
+ {
+ mys[2].a++;
+ }
+
+ #pragma omp target map(mys[2].b)
+ {
+ mys[2].b++;
+ }
+
+ assert (mys[2].a == 1);
+ assert (mys[2].b == 1);
+
+ #pragma omp target map(mys[2].c)
+ {
+ mys[2].c++;
+ }
+
+ #pragma omp target map(mys[2].d)
+ {
+ mys[2].d++;
+ }
+
+ assert (mys[2].c == &c + 1);
+ assert (mys[2].d == &c + 1);
+
+ #pragma omp target map(mys[2].e)
+ {
+ mys[2].e[0]++;
+ }
+
+ #pragma omp target map(mys[2].f)
+ {
+ mys[2].f[0]++;
+ }
+
+ assert (mys[2].e[0] == 2);
+ assert (mys[2].f[0] == 2);
+}
+#endif
+
+#ifdef MAP_SECTIONS
+template<typename A, int B>
+void
+map_sections (void)
+{
+ A arr[B];
+ A *ptr;
+ A (&arrref)[B] = arr;
+ A *&ptrref = ptr;
+
+ ptr = new int[B];
+ memset (ptr, 0, sizeof (int) * B);
+ memset (arr, 0, sizeof (int) * B);
+
+ #pragma omp target map(arr[0:B])
+ {
+ arr[2]++;
+ }
+
+ #pragma omp target map(ptr[0:B])
+ {
+ ptr[2]++;
+ }
+
+ #pragma omp target map(arrref[0:B])
+ {
+ arrref[2]++;
+ }
+
+ #pragma omp target map(ptrref[0:B])
+ {
+ ptrref[2]++;
+ }
+
+ assert (arr[2] == 2);
+ assert (ptr[2] == 2);
+
+ delete ptr;
+}
+#endif
+
+template<typename A>
+struct T {
+ A a[10];
+ A (&b)[10];
+ A *c;
+ A *&d;
+
+ T(A (&b1)[10], A *c1, A *&d1) : b(b1), c(c1), d(d1)
+ {
+ memset (a, 0, sizeof a);
+ }
+};
+
+#ifdef NONREF_DECL_MEMBER_SLICE
+template<typename A, int B>
+void
+nonref_decl_member_slice (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt(c, &c[0], d);
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target map(myt.a[0:B])
+ {
+ myt.a[2]++;
+ }
+
+ #pragma omp target map(myt.b[0:B])
+ {
+ myt.b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myt.c)
+
+ #pragma omp target map(myt.c[0:B])
+ {
+ myt.c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt.c)
+
+ #pragma omp target enter data map(to: myt.d)
+
+ #pragma omp target map(myt.d[0:B])
+ {
+ myt.d[2]++;
+ }
+
+ #pragma omp target exit data map(from: myt.d)
+
+ assert (myt.a[2] == 1);
+ assert (myt.b[2] == 3);
+ assert (myt.c[2] == 3);
+ assert (myt.d[2] == 3);
+}
+#endif
+
+#ifdef NONREF_DECL_MEMBER_SLICE_BASEPTR
+template<typename A, int B>
+void
+nonref_decl_member_slice_baseptr (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt(c, &c[0], d);
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target map(to:myt.c) map(myt.c[0:B])
+ {
+ myt.c[2]++;
+ }
+
+ #pragma omp target map(to:myt.d) map(myt.d[0:B])
+ {
+ myt.d[2]++;
+ }
+
+ assert (myt.c[2] == 2);
+ assert (myt.d[2] == 2);
+}
+#endif
+
+#ifdef REF_DECL_MEMBER_SLICE
+template<typename A, int B>
+void
+ref_decl_member_slice (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real(c, &c[0], d);
+ T<A> &myt = myt_real;
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target map(myt.a[0:B])
+ {
+ myt.a[2]++;
+ }
+
+ #pragma omp target map(myt.b[0:B])
+ {
+ myt.b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myt.c)
+
+ #pragma omp target map(myt.c[0:B])
+ {
+ myt.c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt.c)
+
+ #pragma omp target enter data map(to: myt.d)
+
+ #pragma omp target map(myt.d[0:B])
+ {
+ myt.d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt.d)
+
+ assert (myt.a[2] == 1);
+ assert (myt.b[2] == 3);
+ assert (myt.c[2] == 3);
+ assert (myt.d[2] == 3);
+}
+#endif
+
+#ifdef REF_DECL_MEMBER_SLICE_BASEPTR
+template<typename A, int B>
+void
+ref_decl_member_slice_baseptr (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real(c, &c[0], d);
+ T<A> &myt = myt_real;
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target map(to:myt.c) map(myt.c[0:B])
+ {
+ myt.c[2]++;
+ }
+
+ #pragma omp target map(to:myt.d) map(myt.d[0:B])
+ {
+ myt.d[2]++;
+ }
+
+ assert (myt.c[2] == 2);
+ assert (myt.d[2] == 2);
+}
+#endif
+
+#ifdef PTR_DECL_MEMBER_SLICE
+template<typename A, int B>
+void
+ptr_decl_member_slice (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real(c, &c[0], d);
+ T<A> *myt = &myt_real;
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target enter data map(to: myt)
+
+ #pragma omp target map(myt->a[0:B])
+ {
+ myt->a[2]++;
+ }
+
+ #pragma omp target map(myt->b[0:B])
+ {
+ myt->b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myt->c)
+
+ #pragma omp target map(myt->c[0:B])
+ {
+ myt->c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt->c)
+
+ #pragma omp target enter data map(to: myt->d)
+
+ #pragma omp target map(myt->d[0:B])
+ {
+ myt->d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt, myt->d)
+
+ assert (myt->a[2] == 1);
+ assert (myt->b[2] == 3);
+ assert (myt->c[2] == 3);
+ assert (myt->d[2] == 3);
+}
+#endif
+
+#ifdef PTR_DECL_MEMBER_SLICE_BASEPTR
+template<typename A, int B>
+void
+ptr_decl_member_slice_baseptr (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real(c, &c[0], d);
+ T<A> *myt = &myt_real;
+
+ memset (c, 0, sizeof c);
+
+ // These ones have an implicit firstprivate for 'myt'.
+ #pragma omp target map(to:myt->c) map(myt->c[0:B])
+ {
+ myt->c[2]++;
+ }
+
+ #pragma omp target map(to:myt->d) map(myt->d[0:B])
+ {
+ myt->d[2]++;
+ }
+
+ // These ones have an explicit "TO" mapping for 'myt'.
+ #pragma omp target map(to:myt) map(to:myt->c) map(myt->c[0:B])
+ {
+ myt->c[2]++;
+ }
+
+ #pragma omp target map(to:myt) map(to:myt->d) map(myt->d[0:B])
+ {
+ myt->d[2]++;
+ }
+
+ assert (myt->c[2] == 4);
+ assert (myt->d[2] == 4);
+}
+#endif
+
+#ifdef REF2PTR_DECL_MEMBER_SLICE
+template<typename A, int B>
+void
+ref2ptr_decl_member_slice (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real(c, &c[0], d);
+ T<A> *myt_ptr = &myt_real;
+ T<A> *&myt = myt_ptr;
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target enter data map(to: myt)
+
+ #pragma omp target map(myt->a[0:B])
+ {
+ myt->a[2]++;
+ }
+
+ #pragma omp target map(myt->b[0:B])
+ {
+ myt->b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myt->c)
+
+ #pragma omp target map(myt->c[0:B])
+ {
+ myt->c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt->c)
+
+ #pragma omp target enter data map(to: myt->d)
+
+ #pragma omp target map(myt->d[0:B])
+ {
+ myt->d[2]++;
+ }
+
+ #pragma omp target exit data map(from: myt, myt->d)
+
+ assert (myt->a[2] == 1);
+ assert (myt->b[2] == 3);
+ assert (myt->c[2] == 3);
+ assert (myt->d[2] == 3);
+}
+#endif
+
+#ifdef REF2PTR_DECL_MEMBER_SLICE_BASEPTR
+template<typename A, int B>
+void
+ref2ptr_decl_member_slice_baseptr (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real(c, &c[0], d);
+ T<A> *myt_ptr = &myt_real;
+ T<A> *&myt = myt_ptr;
+
+ memset (c, 0, sizeof c);
+
+ // These ones have an implicit firstprivate for 'myt'.
+ #pragma omp target map(to:myt->c) map(myt->c[0:B])
+ {
+ myt->c[2]++;
+ }
+
+ #pragma omp target map(to:myt->d) map(myt->d[0:B])
+ {
+ myt->d[2]++;
+ }
+
+ // These ones have an explicit "TO" mapping for 'myt'.
+ #pragma omp target map(to:myt) map(to:myt->c) map(myt->c[0:B])
+ {
+ myt->c[2]++;
+ }
+
+ #pragma omp target map(to:myt) map(to:myt->d) map(myt->d[0:B])
+ {
+ myt->d[2]++;
+ }
+
+ assert (myt->c[2] == 4);
+ assert (myt->d[2] == 4);
+}
+#endif
+
+#ifdef ARRAY_DECL_MEMBER_SLICE
+template<typename A, int B>
+void
+array_decl_member_slice (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt[4] =
+ {
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d)
+ };
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target map(myt[2].a[0:B])
+ {
+ myt[2].a[2]++;
+ }
+
+ #pragma omp target map(myt[2].b[0:B])
+ {
+ myt[2].b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myt[2].c)
+
+ #pragma omp target map(myt[2].c[0:B])
+ {
+ myt[2].c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt[2].c)
+
+ #pragma omp target enter data map(to: myt[2].d)
+
+ #pragma omp target map(myt[2].d[0:B])
+ {
+ myt[2].d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt[2].d)
+
+ assert (myt[2].a[2] == 1);
+ assert (myt[2].b[2] == 3);
+ assert (myt[2].c[2] == 3);
+ assert (myt[2].d[2] == 3);
+}
+#endif
+
+#ifdef ARRAY_DECL_MEMBER_SLICE_BASEPTR
+template<typename A, int B>
+void
+array_decl_member_slice_baseptr (void)
+{
+ A c[10];
+ A *d = &c[0];
+ T<A> myt[4] =
+ {
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d)
+ };
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target map(to:myt[2].c) map(myt[2].c[0:B])
+ {
+ myt[2].c[2]++;
+ }
+
+ #pragma omp target map(to:myt[2].d) map(myt[2].d[0:B])
+ {
+ myt[2].d[2]++;
+ }
+
+ assert (myt[2].c[2] == 2);
+ assert (myt[2].d[2] == 2);
+}
+#endif
+
+#ifdef REF2ARRAY_DECL_MEMBER_SLICE
+template<typename A, int B>
+void
+ref2array_decl_member_slice (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real[4] =
+ {
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d)
+ };
+ T<A> (&myt)[4] = myt_real;
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target map(myt[2].a[0:B])
+ {
+ myt[2].a[2]++;
+ }
+
+ #pragma omp target map(myt[2].b[0:B])
+ {
+ myt[2].b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myt[2].c)
+
+ #pragma omp target map(myt[2].c[0:B])
+ {
+ myt[2].c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt[2].c)
+
+ #pragma omp target enter data map(to: myt[2].d)
+
+ #pragma omp target map(myt[2].d[0:B])
+ {
+ myt[2].d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt[2].d)
+
+ assert (myt[2].a[2] == 1);
+ assert (myt[2].b[2] == 3);
+ assert (myt[2].c[2] == 3);
+ assert (myt[2].d[2] == 3);
+}
+#endif
+
+#ifdef REF2ARRAY_DECL_MEMBER_SLICE_BASEPTR
+template<typename A, int B>
+void
+ref2array_decl_member_slice_baseptr (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real[4] =
+ {
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d)
+ };
+ T<A> (&myt)[4] = myt_real;
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target map(to:myt[2].c) map(myt[2].c[0:B])
+ {
+ myt[2].c[2]++;
+ }
+
+ #pragma omp target map(to:myt[2].d) map(myt[2].d[0:B])
+ {
+ myt[2].d[2]++;
+ }
+
+ assert (myt[2].c[2] == 2);
+ assert (myt[2].d[2] == 2);
+}
+#endif
+
+#ifdef PTR_OFFSET_DECL_MEMBER_SLICE
+template<typename A, int B>
+void
+ptr_offset_decl_member_slice (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real[4] =
+ {
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d)
+ };
+ T<A> *myt = &myt_real[0];
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target map(myt[2].a[0:B])
+ {
+ myt[2].a[2]++;
+ }
+
+ #pragma omp target map(myt[2].b[0:B])
+ {
+ myt[2].b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myt[2].c)
+
+ #pragma omp target map(myt[2].c[0:B])
+ {
+ myt[2].c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt[2].c)
+
+ #pragma omp target enter data map(to: myt[2].d)
+
+ #pragma omp target map(myt[2].d[0:B])
+ {
+ myt[2].d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt[2].d)
+
+ assert (myt[2].a[2] == 1);
+ assert (myt[2].b[2] == 3);
+ assert (myt[2].c[2] == 3);
+ assert (myt[2].d[2] == 3);
+}
+#endif
+
+#ifdef PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
+template<typename A, int B>
+void
+ptr_offset_decl_member_slice_baseptr (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real[4] =
+ {
+ T (c, &c[0], d),
+ T (c, &c[0], d),
+ T (c, &c[0], d),
+ T (c, &c[0], d)
+ };
+ T<A> *myt = &myt_real[0];
+
+ memset (c, 0, sizeof c);
+
+ /* Implicit 'myt'. */
+ #pragma omp target map(to:myt[2].c) map(myt[2].c[0:B])
+ {
+ myt[2].c[2]++;
+ }
+
+ #pragma omp target map(to:myt[2].d) map(myt[2].d[0:B])
+ {
+ myt[2].d[2]++;
+ }
+
+ /* Explicit 'to'-mapped 'myt'. */
+ #pragma omp target map(to:myt) map(to:myt[2].c) map(myt[2].c[0:B])
+ {
+ myt[2].c[2]++;
+ }
+
+ #pragma omp target map(to:myt) map(to:myt[2].d) map(myt[2].d[0:B])
+ {
+ myt[2].d[2]++;
+ }
+
+ assert (myt[2].c[2] == 4);
+ assert (myt[2].d[2] == 4);
+}
+#endif
+
+#ifdef REF2PTR_OFFSET_DECL_MEMBER_SLICE
+template<typename A, int B>
+void
+ref2ptr_offset_decl_member_slice (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real[4] =
+ {
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d)
+ };
+ T<A> *myt_ptr = &myt_real[0];
+ T<A> *&myt = myt_ptr;
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target map(myt[2].a[0:B])
+ {
+ myt[2].a[2]++;
+ }
+
+ #pragma omp target map(myt[2].b[0:B])
+ {
+ myt[2].b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myt[2].c)
+
+ #pragma omp target map(myt[2].c[0:B])
+ {
+ myt[2].c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt[2].c)
+
+ #pragma omp target enter data map(to: myt[2].d)
+
+ #pragma omp target map(myt[2].d[0:B])
+ {
+ myt[2].d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt[2].d)
+
+ assert (myt[2].a[2] == 1);
+ assert (myt[2].b[2] == 3);
+ assert (myt[2].c[2] == 3);
+ assert (myt[2].d[2] == 3);
+}
+#endif
+
+#ifdef REF2PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
+template<typename A, int B>
+void
+ref2ptr_offset_decl_member_slice_baseptr (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real[4] =
+ {
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d),
+ T<A> (c, &c[0], d)
+ };
+ T<A> *myt_ptr = &myt_real[0];
+ T<A> *&myt = myt_ptr;
+
+ memset (c, 0, sizeof c);
+
+ /* Implicit 'myt'. */
+ #pragma omp target map(to:myt[2].c) map(myt[2].c[0:B])
+ {
+ myt[2].c[2]++;
+ }
+
+ #pragma omp target map(to:myt[2].d) map(myt[2].d[0:B])
+ {
+ myt[2].d[2]++;
+ }
+
+ /* Explicit 'to'-mapped 'myt'. */
+ #pragma omp target map(to:myt) map(to:myt[2].c) map(myt[2].c[0:B])
+ {
+ myt[2].c[2]++;
+ }
+
+ #pragma omp target map(to:myt) map(to:myt[2].d) map(myt[2].d[0:B])
+ {
+ myt[2].d[2]++;
+ }
+
+ assert (myt[2].c[2] == 4);
+ assert (myt[2].d[2] == 4);
+}
+#endif
+
+#ifdef PTRARRAY_DECL_MEMBER_SLICE
+template<typename A, int B>
+void
+ptrarray_decl_member_slice (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real(c, &c[0], d);
+ T<A> *myt[4] =
+ {
+ &myt_real,
+ &myt_real,
+ &myt_real,
+ &myt_real
+ };
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target enter data map(to: myt[2])
+
+ #pragma omp target map(myt[2]->a[0:B])
+ {
+ myt[2]->a[2]++;
+ }
+
+ #pragma omp target map(myt[2]->b[0:B])
+ {
+ myt[2]->b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myt[2]->c)
+
+ #pragma omp target map(myt[2]->c[0:B])
+ {
+ myt[2]->c[2]++;
+ }
+
+ #pragma omp target exit data map(from: myt[2]->c)
+
+ #pragma omp target enter data map(to: myt[2]->d)
+
+ #pragma omp target map(myt[2]->d[0:B])
+ {
+ myt[2]->d[2]++;
+ }
+
+ #pragma omp target exit data map(from: myt[2]->d)
+
+ #pragma omp target exit data map(release: myt[2])
+
+ assert (myt[2]->a[2] == 1);
+ assert (myt[2]->b[2] == 3);
+ assert (myt[2]->c[2] == 3);
+ assert (myt[2]->d[2] == 3);
+}
+#endif
+
+#ifdef PTRARRAY_DECL_MEMBER_SLICE_BASEPTR
+template<typename A, int B>
+void
+ptrarray_decl_member_slice_baseptr (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real(c, &c[0], d);
+ T<A> *myt[4] =
+ {
+ &myt_real,
+ &myt_real,
+ &myt_real,
+ &myt_real
+ };
+
+ memset (c, 0, sizeof c);
+
+ // Implicit 'myt'
+ #pragma omp target map(to: myt[2]->c) map(myt[2]->c[0:B])
+ {
+ myt[2]->c[2]++;
+ }
+
+ #pragma omp target map(to: myt[2]->d) map(myt[2]->d[0:B])
+ {
+ myt[2]->d[2]++;
+ }
+
+ // One element of 'myt'
+ #pragma omp target map(to:myt[2], myt[2]->c) map(myt[2]->c[0:B])
+ {
+ myt[2]->c[2]++;
+ }
+
+ #pragma omp target map(to:myt[2], myt[2]->d) map(myt[2]->d[0:B])
+ {
+ myt[2]->d[2]++;
+ }
+
+ // Explicit map of all of 'myt'
+ #pragma omp target map(to:myt, myt[2]->c) map(myt[2]->c[0:B])
+ {
+ myt[2]->c[2]++;
+ }
+
+ #pragma omp target map(to:myt, myt[2]->d) map(myt[2]->d[0:B])
+ {
+ myt[2]->d[2]++;
+ }
+
+ // Explicit map slice of 'myt'
+ #pragma omp target map(to:myt[1:3], myt[2]->c) map(myt[2]->c[0:B])
+ {
+ myt[2]->c[2]++;
+ }
+
+ #pragma omp target map(to:myt[1:3], myt[2]->d) map(myt[2]->d[0:B])
+ {
+ myt[2]->d[2]++;
+ }
+
+ assert (myt[2]->c[2] == 8);
+ assert (myt[2]->d[2] == 8);
+}
+#endif
+
+#ifdef REF2PTRARRAY_DECL_MEMBER_SLICE
+template<typename A, int B>
+void
+ref2ptrarray_decl_member_slice (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real(c, &c[0], d);
+ T<A> *myt_ptrarr[4] =
+ {
+ &myt_real,
+ &myt_real,
+ &myt_real,
+ &myt_real
+ };
+ T<A> *(&myt)[4] = myt_ptrarr;
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target enter data map(to: myt[2])
+
+ #pragma omp target map(myt[2]->a[0:B])
+ {
+ myt[2]->a[2]++;
+ }
+
+ #pragma omp target map(myt[2]->b[0:B])
+ {
+ myt[2]->b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myt[2]->c)
+
+ #pragma omp target map(myt[2]->c[0:B])
+ {
+ myt[2]->c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt[2]->c)
+
+ #pragma omp target enter data map(to: myt[2]->d)
+
+ #pragma omp target map(myt[2]->d[0:B])
+ {
+ myt[2]->d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt[2]->d)
+
+ #pragma omp target exit data map(release: myt[2])
+
+ assert (myt[2]->a[2] == 1);
+ assert (myt[2]->b[2] == 3);
+ assert (myt[2]->c[2] == 3);
+ assert (myt[2]->d[2] == 3);
+}
+#endif
+
+#ifdef REF2PTRARRAY_DECL_MEMBER_SLICE_BASEPTR
+template<typename A, int B>
+void
+ref2ptrarray_decl_member_slice_baseptr (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real(c, &c[0], d);
+ T<A> *myt_ptrarr[4] =
+ {
+ &myt_real,
+ &myt_real,
+ &myt_real,
+ &myt_real
+ };
+ T<A> *(&myt)[4] = myt_ptrarr;
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target map(to:myt[2], myt[2]->c) map(myt[2]->c[0:B])
+ {
+ myt[2]->c[2]++;
+ }
+
+ #pragma omp target map(to:myt[2], myt[2]->d) map(myt[2]->d[0:B])
+ {
+ myt[2]->d[2]++;
+ }
+
+ #pragma omp target map(to:myt, myt[2]->c) map(myt[2]->c[0:B])
+ {
+ myt[2]->c[2]++;
+ }
+
+ #pragma omp target map(to:myt, myt[2]->d) map(myt[2]->d[0:B])
+ {
+ myt[2]->d[2]++;
+ }
+
+ assert (myt[2]->c[2] == 4);
+ assert (myt[2]->d[2] == 4);
+}
+#endif
+
+#ifdef PTRPTR_OFFSET_DECL_MEMBER_SLICE
+template<typename A, int B>
+void
+ptrptr_offset_decl_member_slice (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real(c, &c[0], d);
+ T<A> *myt_ptrarr[4] =
+ {
+ &myt_real,
+ &myt_real,
+ &myt_real,
+ &myt_real
+ };
+ T<A> **myt = &myt_ptrarr[0];
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target enter data map(to: myt[0:3])
+
+ /* NOTE: For the implicit firstprivate 'myt' to work, the zeroth element of
+ myt[] must be mapped above -- otherwise the zero-length array section
+ lookup fails. */
+ #pragma omp target map(myt[2]->a[0:B])
+ {
+ myt[2]->a[2]++;
+ }
+
+ #pragma omp target map(myt[2]->b[0:B])
+ {
+ myt[2]->b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myt[2]->c)
+
+ #pragma omp target map(myt[2]->c[0:B])
+ {
+ myt[2]->c[2]++;
+ }
+
+ #pragma omp target exit data map(from: myt[2]->c)
+
+ #pragma omp target enter data map(to: myt[2]->d)
+
+ #pragma omp target map(myt[2]->d[0:B])
+ {
+ myt[2]->d[2]++;
+ }
+
+ #pragma omp target exit data map(from: myt[0:3], myt[2]->d)
+
+ assert (myt[2]->a[2] == 1);
+ assert (myt[2]->b[2] == 3);
+ assert (myt[2]->c[2] == 3);
+ assert (myt[2]->d[2] == 3);
+}
+#endif
+
+#ifdef PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
+template<typename A, int B>
+void
+ptrptr_offset_decl_member_slice_baseptr (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real(c, &c[0], d);
+ T<A> *myt_ptrarr[4] =
+ {
+ 0,
+ 0,
+ 0,
+ &myt_real
+ };
+ T<A> **myt = &myt_ptrarr[0];
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target map(to:myt[3], myt[3]->c) map(myt[3]->c[0:B])
+ {
+ myt[3]->c[2]++;
+ }
+
+ #pragma omp target map(to:myt[3], myt[3]->d) map(myt[3]->d[0:B])
+ {
+ myt[3]->d[2]++;
+ }
+
+ #pragma omp target map(to:myt, myt[3], myt[3]->c) map(myt[3]->c[0:B])
+ {
+ myt[3]->c[2]++;
+ }
+
+ #pragma omp target map(to:myt, myt[3], myt[3]->d) map(myt[3]->d[0:B])
+ {
+ myt[3]->d[2]++;
+ }
+
+ assert (myt[3]->c[2] == 4);
+ assert (myt[3]->d[2] == 4);
+}
+#endif
+
+#ifdef REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE
+template<typename A, int B>
+void
+ref2ptrptr_offset_decl_member_slice (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real(c, &c[0], d);
+ T<A> *myt_ptrarr[4] =
+ {
+ 0,
+ 0,
+ &myt_real,
+ 0
+ };
+ T<A> **myt_ptrptr = &myt_ptrarr[0];
+ T<A> **&myt = myt_ptrptr;
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target enter data map(to: myt[0:3])
+
+ #pragma omp target map(myt[2]->a[0:B])
+ {
+ myt[2]->a[2]++;
+ }
+
+ #pragma omp target map(myt[2]->b[0:B])
+ {
+ myt[2]->b[2]++;
+ }
+
+ #pragma omp target enter data map(to:myt[2]->c)
+
+ #pragma omp target map(myt[2]->c[0:B])
+ {
+ myt[2]->c[2]++;
+ }
+
+ #pragma omp target exit data map(release:myt[2]->c)
+
+ #pragma omp target enter data map(to:myt[2]->d)
+
+ #pragma omp target map(myt[2]->d[0:B])
+ {
+ myt[2]->d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myt[0:3], myt[2]->d)
+
+ assert (myt[2]->a[2] == 1);
+ assert (myt[2]->b[2] == 3);
+ assert (myt[2]->c[2] == 3);
+ assert (myt[2]->d[2] == 3);
+}
+#endif
+
+#ifdef REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
+template<typename A, int B>
+void
+ref2ptrptr_offset_decl_member_slice_baseptr (void)
+{
+ A c[B];
+ A *d = &c[0];
+ T<A> myt_real(c, &c[0], d);
+ T<A> *myt_ptrarr[4] =
+ {
+ 0,
+ 0,
+ &myt_real,
+ 0
+ };
+ T<A> **myt_ptrptr = &myt_ptrarr[0];
+ T<A> **&myt = myt_ptrptr;
+
+ memset (c, 0, sizeof c);
+
+ #pragma omp target map(to:myt[2], myt[2]->c) map(myt[2]->c[0:B])
+ {
+ myt[2]->c[2]++;
+ }
+
+ #pragma omp target map(to:myt[2], myt[2]->d) map(myt[2]->d[0:B])
+ {
+ myt[2]->d[2]++;
+ }
+
+ #pragma omp target map(to:myt, myt[2], myt[2]->c) map(myt[2]->c[0:B])
+ {
+ myt[2]->c[2]++;
+ }
+
+ #pragma omp target map(to:myt, myt[2], myt[2]->d) map(myt[2]->d[0:B])
+ {
+ myt[2]->d[2]++;
+ }
+
+ assert (myt[2]->c[2] == 4);
+ assert (myt[2]->d[2] == 4);
+}
+#endif
+
+template<typename A>
+struct U
+{
+ S<A> s1;
+ T<A> t1;
+ S<A> &s2;
+ T<A> &t2;
+ S<A> *s3;
+ T<A> *t3;
+ S<A> *&s4;
+ T<A> *&t4;
+
+ U(S<A> &sptr1, T<A> &tptr1, S<A> &sptr2, T<A> &tptr2,
+ S<A> *sptr3, T<A> *tptr3, S<A> *&sptr4, T<A> *&tptr4)
+ : s1(sptr1), t1(tptr1), s2(sptr2), t2(tptr2), s3(sptr3), t3(tptr3),
+ s4(sptr4), t4(tptr4)
+ {
+ }
+};
+
+#define INIT_S(N) \
+ A a##N = 0, b##N = 0, c##N = 0, d##N = 0; \
+ A *d##N##ptr = &d##N; \
+ S<A> s##N(a##N, b##N, &c##N, d##N##ptr)
+
+#define INIT_T(N) \
+ A arr##N[10]; \
+ A *ptr##N = &arr##N[0]; \
+ T<A> t##N(arr##N, &arr##N[0], ptr##N); \
+ memset (arr##N, 0, sizeof arr##N)
+
+#define INIT_ST \
+ INIT_S(1); \
+ INIT_T(1); \
+ INIT_S(2); \
+ INIT_T(2); \
+ INIT_S(3); \
+ INIT_T(3); \
+ A a4 = 0, b4 = 0, c4 = 0, d4 = 0; \
+ A *d4ptr = &d4; \
+ S<A> *s4 = new S<A>(a4, b4, &c4, d4ptr); \
+ A arr4[10]; \
+ A *ptr4 = &arr4[0]; \
+ T<A> *t4 = new T<A>(arr4, &arr4[0], ptr4); \
+ memset (arr4, 0, sizeof arr4)
+
+#ifdef NONREF_COMPONENT_BASE
+template<typename A>
+void
+nonref_component_base (void)
+{
+ INIT_ST;
+ U<A> myu(s1, t1, s2, t2, &s3, &t3, s4, t4);
+
+ #pragma omp target map(myu.s1.a, myu.s1.b, myu.s1.c, myu.s1.d)
+ {
+ myu.s1.a++;
+ myu.s1.b++;
+ myu.s1.c++;
+ myu.s1.d++;
+ }
+
+ assert (myu.s1.a == 1);
+ assert (myu.s1.b == 1);
+ assert (myu.s1.c == &c1 + 1);
+ assert (myu.s1.d == &d1 + 1);
+
+ #pragma omp target map(myu.s2.a, myu.s2.b, myu.s2.c, myu.s2.d)
+ {
+ myu.s2.a++;
+ myu.s2.b++;
+ myu.s2.c++;
+ myu.s2.d++;
+ }
+
+ assert (myu.s2.a == 1);
+ assert (myu.s2.b == 1);
+ assert (myu.s2.c == &c2 + 1);
+ assert (myu.s2.d == &d2 + 1);
+
+ #pragma omp target map(to:myu.s3) \
+ map(myu.s3->a, myu.s3->b, myu.s3->c, myu.s3->d)
+ {
+ myu.s3->a++;
+ myu.s3->b++;
+ myu.s3->c++;
+ myu.s3->d++;
+ }
+
+ assert (myu.s3->a == 1);
+ assert (myu.s3->b == 1);
+ assert (myu.s3->c == &c3 + 1);
+ assert (myu.s3->d == &d3 + 1);
+
+ #pragma omp target map(to:myu.s4) \
+ map(myu.s4->a, myu.s4->b, myu.s4->c, myu.s4->d)
+ {
+ myu.s4->a++;
+ myu.s4->b++;
+ myu.s4->c++;
+ myu.s4->d++;
+ }
+
+ assert (myu.s4->a == 1);
+ assert (myu.s4->b == 1);
+ assert (myu.s4->c == &c4 + 1);
+ assert (myu.s4->d == &d4 + 1);
+
+ delete s4;
+ delete t4;
+}
+#endif
+
+#ifdef NONREF_COMPONENT_MEMBER_SLICE
+template<typename A>
+void
+nonref_component_member_slice (void)
+{
+ INIT_ST;
+ U<A> myu(s1, t1, s2, t2, &s3, &t3, s4, t4);
+
+ #pragma omp target map(myu.t1.a[2:5])
+ {
+ myu.t1.a[2]++;
+ }
+
+ #pragma omp target map(myu.t1.b[2:5])
+ {
+ myu.t1.b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu.t1.c)
+
+ #pragma omp target map(myu.t1.c[2:5])
+ {
+ myu.t1.c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t1.c)
+
+ #pragma omp target enter data map(to: myu.t1.d)
+
+ #pragma omp target map(myu.t1.d[2:5])
+ {
+ myu.t1.d[2]++;
+ }
+
+ #pragma omp target exit data map(from: myu.t1.d)
+
+ assert (myu.t1.a[2] == 1);
+ assert (myu.t1.b[2] == 3);
+ assert (myu.t1.c[2] == 3);
+ assert (myu.t1.d[2] == 3);
+
+ #pragma omp target map(myu.t2.a[2:5])
+ {
+ myu.t2.a[2]++;
+ }
+
+ #pragma omp target map(myu.t2.b[2:5])
+ {
+ myu.t2.b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu.t2.c)
+
+ #pragma omp target map(myu.t2.c[2:5])
+ {
+ myu.t2.c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t2.c)
+
+ #pragma omp target enter data map(to: myu.t2.d)
+
+ #pragma omp target map(myu.t2.d[2:5])
+ {
+ myu.t2.d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t2.d)
+
+ assert (myu.t2.a[2] == 1);
+ assert (myu.t2.b[2] == 3);
+ assert (myu.t2.c[2] == 3);
+ assert (myu.t2.d[2] == 3);
+
+ #pragma omp target enter data map(to: myu.t3)
+
+ #pragma omp target map(myu.t3->a[2:5])
+ {
+ myu.t3->a[2]++;
+ }
+
+ #pragma omp target map(myu.t3->b[2:5])
+ {
+ myu.t3->b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu.t3->c)
+
+ #pragma omp target map(myu.t3->c[2:5])
+ {
+ myu.t3->c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t3->c)
+
+ #pragma omp target enter data map(to: myu.t3->d)
+
+ #pragma omp target map(myu.t3->d[2:5])
+ {
+ myu.t3->d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t3, myu.t3->d)
+
+ assert (myu.t3->a[2] == 1);
+ assert (myu.t3->b[2] == 3);
+ assert (myu.t3->c[2] == 3);
+ assert (myu.t3->d[2] == 3);
+
+ #pragma omp target enter data map(to: myu.t4)
+
+ #pragma omp target map(myu.t4->a[2:5])
+ {
+ myu.t4->a[2]++;
+ }
+
+ #pragma omp target map(myu.t4->b[2:5])
+ {
+ myu.t4->b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu.t4->c)
+
+ #pragma omp target map(myu.t4->c[2:5])
+ {
+ myu.t4->c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t4->c)
+
+ #pragma omp target enter data map(to: myu.t4->d)
+
+ #pragma omp target map(myu.t4->d[2:5])
+ {
+ myu.t4->d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t4, myu.t4->d)
+
+ assert (myu.t4->a[2] == 1);
+ assert (myu.t4->b[2] == 3);
+ assert (myu.t4->c[2] == 3);
+ assert (myu.t4->d[2] == 3);
+
+ delete s4;
+ delete t4;
+}
+#endif
+
+#ifdef NONREF_COMPONENT_MEMBER_SLICE_BASEPTR
+template<typename A>
+void
+nonref_component_member_slice_baseptr (void)
+{
+ INIT_ST;
+ U<A> myu(s1, t1, s2, t2, &s3, &t3, s4, t4);
+
+ #pragma omp target map(to: myu.t1.c) map(myu.t1.c[2:5])
+ {
+ myu.t1.c[2]++;
+ }
+
+ #pragma omp target map(to: myu.t1.d) map(myu.t1.d[2:5])
+ {
+ myu.t1.d[2]++;
+ }
+
+ assert (myu.t1.c[2] == 2);
+ assert (myu.t1.d[2] == 2);
+
+ #pragma omp target map(to: myu.t2.c) map(myu.t2.c[2:5])
+ {
+ myu.t2.c[2]++;
+ }
+
+ #pragma omp target map(to: myu.t2.d) map(myu.t2.d[2:5])
+ {
+ myu.t2.d[2]++;
+ }
+
+ assert (myu.t2.c[2] == 2);
+ assert (myu.t2.d[2] == 2);
+
+ #pragma omp target map(to: myu.t3, myu.t3->c) map(myu.t3->c[2:5])
+ {
+ myu.t3->c[2]++;
+ }
+
+ #pragma omp target map(to: myu.t3, myu.t3->d) map(myu.t3->d[2:5])
+ {
+ myu.t3->d[2]++;
+ }
+
+ assert (myu.t3->c[2] == 2);
+ assert (myu.t3->d[2] == 2);
+
+ #pragma omp target map(to: myu.t4, myu.t4->c) map(myu.t4->c[2:5])
+ {
+ myu.t4->c[2]++;
+ }
+
+ #pragma omp target map(to: myu.t4, myu.t4->d) map(myu.t4->d[2:5])
+ {
+ myu.t4->d[2]++;
+ }
+
+ assert (myu.t4->c[2] == 2);
+ assert (myu.t4->d[2] == 2);
+
+ delete s4;
+ delete t4;
+}
+#endif
+
+#ifdef REF_COMPONENT_BASE
+template<typename A>
+void
+ref_component_base (void)
+{
+ INIT_ST;
+ U<A> myu_real(s1, t1, s2, t2, &s3, &t3, s4, t4);
+ U<A> &myu = myu_real;
+
+ #pragma omp target map(myu.s1.a, myu.s1.b, myu.s1.c, myu.s1.d)
+ {
+ myu.s1.a++;
+ myu.s1.b++;
+ myu.s1.c++;
+ myu.s1.d++;
+ }
+
+ assert (myu.s1.a == 1);
+ assert (myu.s1.b == 1);
+ assert (myu.s1.c == &c1 + 1);
+ assert (myu.s1.d == &d1 + 1);
+
+ #pragma omp target map(myu.s2.a, myu.s2.b, myu.s2.c, myu.s2.d)
+ {
+ myu.s2.a++;
+ myu.s2.b++;
+ myu.s2.c++;
+ myu.s2.d++;
+ }
+
+ assert (myu.s2.a == 1);
+ assert (myu.s2.b == 1);
+ assert (myu.s2.c == &c2 + 1);
+ assert (myu.s2.d == &d2 + 1);
+
+ #pragma omp target map(to:myu.s3) \
+ map(myu.s3->a, myu.s3->b, myu.s3->c, myu.s3->d)
+ {
+ myu.s3->a++;
+ myu.s3->b++;
+ myu.s3->c++;
+ myu.s3->d++;
+ }
+
+ assert (myu.s3->a == 1);
+ assert (myu.s3->b == 1);
+ assert (myu.s3->c == &c3 + 1);
+ assert (myu.s3->d == &d3 + 1);
+
+ #pragma omp target map(to:myu.s4) \
+ map(myu.s4->a, myu.s4->b, myu.s4->c, myu.s4->d)
+ {
+ myu.s4->a++;
+ myu.s4->b++;
+ myu.s4->c++;
+ myu.s4->d++;
+ }
+
+ assert (myu.s4->a == 1);
+ assert (myu.s4->b == 1);
+ assert (myu.s4->c == &c4 + 1);
+ assert (myu.s4->d == &d4 + 1);
+
+ delete s4;
+ delete t4;
+}
+#endif
+
+#ifdef REF_COMPONENT_MEMBER_SLICE
+template<typename A>
+void
+ref_component_member_slice (void)
+{
+ INIT_ST;
+ U<A> myu_real(s1, t1, s2, t2, &s3, &t3, s4, t4);
+ U<A> &myu = myu_real;
+
+ #pragma omp target map(myu.t1.a[2:5])
+ {
+ myu.t1.a[2]++;
+ }
+
+ #pragma omp target map(myu.t1.b[2:5])
+ {
+ myu.t1.b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu.t1.c)
+
+ #pragma omp target map(myu.t1.c[2:5])
+ {
+ myu.t1.c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t1.c)
+
+ #pragma omp target enter data map(to: myu.t1.d)
+
+ #pragma omp target map(myu.t1.d[2:5])
+ {
+ myu.t1.d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t1.d)
+
+ assert (myu.t1.a[2] == 1);
+ assert (myu.t1.b[2] == 3);
+ assert (myu.t1.c[2] == 3);
+ assert (myu.t1.d[2] == 3);
+
+ #pragma omp target map(myu.t2.a[2:5])
+ {
+ myu.t2.a[2]++;
+ }
+
+ #pragma omp target map(myu.t2.b[2:5])
+ {
+ myu.t2.b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu.t2.c)
+
+ #pragma omp target map(myu.t2.c[2:5])
+ {
+ myu.t2.c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t2.c)
+
+ #pragma omp target enter data map(to: myu.t2.d)
+
+ #pragma omp target map(myu.t2.d[2:5])
+ {
+ myu.t2.d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t2.d)
+
+ assert (myu.t2.a[2] == 1);
+ assert (myu.t2.b[2] == 3);
+ assert (myu.t2.c[2] == 3);
+ assert (myu.t2.d[2] == 3);
+
+ #pragma omp target enter data map(to: myu.t3)
+
+ #pragma omp target map(myu.t3->a[2:5])
+ {
+ myu.t3->a[2]++;
+ }
+
+ #pragma omp target map(myu.t3->b[2:5])
+ {
+ myu.t3->b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu.t3->c)
+
+ #pragma omp target map(myu.t3->c[2:5])
+ {
+ myu.t3->c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t3->c)
+
+ #pragma omp target enter data map(to: myu.t3->d)
+
+ #pragma omp target map(myu.t3->d[2:5])
+ {
+ myu.t3->d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t3, myu.t3->d)
+
+ assert (myu.t3->a[2] == 1);
+ assert (myu.t3->b[2] == 3);
+ assert (myu.t3->c[2] == 3);
+ assert (myu.t3->d[2] == 3);
+
+ #pragma omp target enter data map(to: myu.t4)
+
+ #pragma omp target map(myu.t4->a[2:5])
+ {
+ myu.t4->a[2]++;
+ }
+
+ #pragma omp target map(myu.t4->b[2:5])
+ {
+ myu.t4->b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu.t4->c)
+
+ #pragma omp target map(myu.t4->c[2:5])
+ {
+ myu.t4->c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t4->c)
+
+ #pragma omp target enter data map(to: myu.t4->d)
+
+ #pragma omp target map(myu.t4->d[2:5])
+ {
+ myu.t4->d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu.t4, myu.t4->d)
+
+ assert (myu.t4->a[2] == 1);
+ assert (myu.t4->b[2] == 3);
+ assert (myu.t4->c[2] == 3);
+ assert (myu.t4->d[2] == 3);
+
+ delete s4;
+ delete t4;
+}
+#endif
+
+#ifdef REF_COMPONENT_MEMBER_SLICE_BASEPTR
+template<typename A>
+void
+ref_component_member_slice_baseptr (void)
+{
+ INIT_ST;
+ U<A> myu_real(s1, t1, s2, t2, &s3, &t3, s4, t4);
+ U<A> &myu = myu_real;
+
+ #pragma omp target map(to: myu.t1.c) map(myu.t1.c[2:5])
+ {
+ myu.t1.c[2]++;
+ }
+
+ #pragma omp target map(to: myu.t1.d) map(myu.t1.d[2:5])
+ {
+ myu.t1.d[2]++;
+ }
+
+ assert (myu.t1.c[2] == 2);
+ assert (myu.t1.d[2] == 2);
+
+ #pragma omp target map(to: myu.t2.c) map(myu.t2.c[2:5])
+ {
+ myu.t2.c[2]++;
+ }
+
+ #pragma omp target map(to: myu.t2.d) map(myu.t2.d[2:5])
+ {
+ myu.t2.d[2]++;
+ }
+
+ assert (myu.t2.c[2] == 2);
+ assert (myu.t2.d[2] == 2);
+
+ #pragma omp target map(to: myu.t3, myu.t3->c) map(myu.t3->c[2:5])
+ {
+ myu.t3->c[2]++;
+ }
+
+ #pragma omp target map(to: myu.t3, myu.t3->d) map(myu.t3->d[2:5])
+ {
+ myu.t3->d[2]++;
+ }
+
+ assert (myu.t3->c[2] == 2);
+ assert (myu.t3->d[2] == 2);
+
+ #pragma omp target map(to: myu.t4, myu.t4->c) map(myu.t4->c[2:5])
+ {
+ myu.t4->c[2]++;
+ }
+
+ #pragma omp target map(to: myu.t4, myu.t4->d) map(myu.t4->d[2:5])
+ {
+ myu.t4->d[2]++;
+ }
+
+ assert (myu.t4->c[2] == 2);
+ assert (myu.t4->d[2] == 2);
+
+ delete s4;
+ delete t4;
+}
+#endif
+
+#ifdef PTR_COMPONENT_BASE
+template<typename A>
+void
+ptr_component_base (void)
+{
+ INIT_ST;
+ U<A> *myu = new U<A>(s1, t1, s2, t2, &s3, &t3, s4, t4);
+
+ #pragma omp target map(myu->s1.a, myu->s1.b, myu->s1.c, myu->s1.d)
+ {
+ myu->s1.a++;
+ myu->s1.b++;
+ myu->s1.c++;
+ myu->s1.d++;
+ }
+
+ assert (myu->s1.a == 1);
+ assert (myu->s1.b == 1);
+ assert (myu->s1.c == &c1 + 1);
+ assert (myu->s1.d == &d1 + 1);
+
+ #pragma omp target map(myu->s2.a, myu->s2.b, myu->s2.c, myu->s2.d)
+ {
+ myu->s2.a++;
+ myu->s2.b++;
+ myu->s2.c++;
+ myu->s2.d++;
+ }
+
+ assert (myu->s2.a == 1);
+ assert (myu->s2.b == 1);
+ assert (myu->s2.c == &c2 + 1);
+ assert (myu->s2.d == &d2 + 1);
+
+ #pragma omp target map(to:myu->s3) \
+ map(myu->s3->a, myu->s3->b, myu->s3->c, myu->s3->d)
+ {
+ myu->s3->a++;
+ myu->s3->b++;
+ myu->s3->c++;
+ myu->s3->d++;
+ }
+
+ assert (myu->s3->a == 1);
+ assert (myu->s3->b == 1);
+ assert (myu->s3->c == &c3 + 1);
+ assert (myu->s3->d == &d3 + 1);
+
+ #pragma omp target map(to:myu->s4) \
+ map(myu->s4->a, myu->s4->b, myu->s4->c, myu->s4->d)
+ {
+ myu->s4->a++;
+ myu->s4->b++;
+ myu->s4->c++;
+ myu->s4->d++;
+ }
+
+ assert (myu->s4->a == 1);
+ assert (myu->s4->b == 1);
+ assert (myu->s4->c == &c4 + 1);
+ assert (myu->s4->d == &d4 + 1);
+
+ delete s4;
+ delete t4;
+ delete myu;
+}
+#endif
+
+#ifdef PTR_COMPONENT_MEMBER_SLICE
+template<typename A>
+void
+ptr_component_member_slice (void)
+{
+ INIT_ST;
+ U<A> *myu = new U<A>(s1, t1, s2, t2, &s3, &t3, s4, t4);
+
+ #pragma omp target map(myu->t1.a[2:5])
+ {
+ myu->t1.a[2]++;
+ }
+
+ #pragma omp target map(myu->t1.b[2:5])
+ {
+ myu->t1.b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu->t1.c)
+
+ #pragma omp target map(myu->t1.c[2:5])
+ {
+ myu->t1.c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t1.c)
+
+ #pragma omp target enter data map(to: myu->t1.d)
+
+ #pragma omp target map(myu->t1.d[2:5])
+ {
+ myu->t1.d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t1.d)
+
+ assert (myu->t1.a[2] == 1);
+ assert (myu->t1.b[2] == 3);
+ assert (myu->t1.c[2] == 3);
+ assert (myu->t1.d[2] == 3);
+
+ #pragma omp target map(myu->t2.a[2:5])
+ {
+ myu->t2.a[2]++;
+ }
+
+ #pragma omp target map(myu->t2.b[2:5])
+ {
+ myu->t2.b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu->t2.c)
+
+ #pragma omp target map(myu->t2.c[2:5])
+ {
+ myu->t2.c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t2.c)
+
+ #pragma omp target enter data map(to: myu->t2.d)
+
+ #pragma omp target map(myu->t2.d[2:5])
+ {
+ myu->t2.d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t2.d)
+
+ assert (myu->t2.a[2] == 1);
+ assert (myu->t2.b[2] == 3);
+ assert (myu->t2.c[2] == 3);
+ assert (myu->t2.d[2] == 3);
+
+ #pragma omp target enter data map(to: myu->t3)
+
+ #pragma omp target map(myu->t3->a[2:5])
+ {
+ myu->t3->a[2]++;
+ }
+
+ #pragma omp target map(myu->t3->b[2:5])
+ {
+ myu->t3->b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu->t3->c)
+
+ #pragma omp target map(myu->t3->c[2:5])
+ {
+ myu->t3->c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t3->c)
+
+ #pragma omp target enter data map(to: myu->t3->d)
+
+ #pragma omp target map(myu->t3->d[2:5])
+ {
+ myu->t3->d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t3, myu->t3->d)
+
+ assert (myu->t3->a[2] == 1);
+ assert (myu->t3->b[2] == 3);
+ assert (myu->t3->c[2] == 3);
+ assert (myu->t3->d[2] == 3);
+
+ #pragma omp target enter data map(to: myu->t4)
+
+ #pragma omp target map(myu->t4->a[2:5])
+ {
+ myu->t4->a[2]++;
+ }
+
+ #pragma omp target map(myu->t4->b[2:5])
+ {
+ myu->t4->b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu->t4->c)
+
+ #pragma omp target map(myu->t4->c[2:5])
+ {
+ myu->t4->c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t4->c)
+
+ #pragma omp target enter data map(to: myu->t4->d)
+
+ #pragma omp target map(myu->t4->d[2:5])
+ {
+ myu->t4->d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t4, myu->t4->d)
+
+ assert (myu->t4->a[2] == 1);
+ assert (myu->t4->b[2] == 3);
+ assert (myu->t4->c[2] == 3);
+ assert (myu->t4->d[2] == 3);
+
+ delete s4;
+ delete t4;
+ delete myu;
+}
+#endif
+
+#ifdef PTR_COMPONENT_MEMBER_SLICE_BASEPTR
+template<typename A>
+void
+ptr_component_member_slice_baseptr (void)
+{
+ INIT_ST;
+ U<A> *myu = new U<A>(s1, t1, s2, t2, &s3, &t3, s4, t4);
+
+ /* Implicit firstprivate 'myu'. */
+ #pragma omp target map(to: myu->t1.c) map(myu->t1.c[2:5])
+ {
+ myu->t1.c[2]++;
+ }
+
+ #pragma omp target map(to: myu->t1.d) map(myu->t1.d[2:5])
+ {
+ myu->t1.d[2]++;
+ }
+
+ assert (myu->t1.c[2] == 2);
+ assert (myu->t1.d[2] == 2);
+
+ /* Explicitly-mapped 'myu'. */
+ #pragma omp target map(to: myu, myu->t1.c) map(myu->t1.c[2:5])
+ {
+ myu->t1.c[2]++;
+ }
+
+ #pragma omp target map(to: myu, myu->t1.d) map(myu->t1.d[2:5])
+ {
+ myu->t1.d[2]++;
+ }
+
+ assert (myu->t1.c[2] == 4);
+ assert (myu->t1.d[2] == 4);
+
+ /* Implicit firstprivate 'myu'. */
+ #pragma omp target map(to: myu->t2.c) map(myu->t2.c[2:5])
+ {
+ myu->t2.c[2]++;
+ }
+
+ #pragma omp target map(to: myu->t2.d) map(myu->t2.d[2:5])
+ {
+ myu->t2.d[2]++;
+ }
+
+ assert (myu->t2.c[2] == 2);
+ assert (myu->t2.d[2] == 2);
+
+ /* Explicitly-mapped 'myu'. */
+ #pragma omp target map(to: myu, myu->t2.c) map(myu->t2.c[2:5])
+ {
+ myu->t2.c[2]++;
+ }
+
+ #pragma omp target map(to: myu, myu->t2.d) map(myu->t2.d[2:5])
+ {
+ myu->t2.d[2]++;
+ }
+
+ assert (myu->t2.c[2] == 4);
+ assert (myu->t2.d[2] == 4);
+
+ /* Implicit firstprivate 'myu'. */
+ #pragma omp target map(to: myu->t3, myu->t3->c) map(myu->t3->c[2:5])
+ {
+ myu->t3->c[2]++;
+ }
+
+ #pragma omp target map(to: myu->t3, myu->t3->d) map(myu->t3->d[2:5])
+ {
+ myu->t3->d[2]++;
+ }
+
+ assert (myu->t3->c[2] == 2);
+ assert (myu->t3->d[2] == 2);
+
+ /* Explicitly-mapped 'myu'. */
+ #pragma omp target map(to: myu, myu->t3, myu->t3->c) map(myu->t3->c[2:5])
+ {
+ myu->t3->c[2]++;
+ }
+
+ #pragma omp target map(to: myu, myu->t3, myu->t3->d) map(myu->t3->d[2:5])
+ {
+ myu->t3->d[2]++;
+ }
+
+ assert (myu->t3->c[2] == 4);
+ assert (myu->t3->d[2] == 4);
+
+ /* Implicit firstprivate 'myu'. */
+ #pragma omp target map(to: myu->t4, myu->t4->c) map(myu->t4->c[2:5])
+ {
+ myu->t4->c[2]++;
+ }
+
+ #pragma omp target map(to: myu->t4, myu->t4->d) map(myu->t4->d[2:5])
+ {
+ myu->t4->d[2]++;
+ }
+
+ assert (myu->t4->c[2] == 2);
+ assert (myu->t4->d[2] == 2);
+
+ /* Explicitly-mapped 'myu'. */
+ #pragma omp target map(to: myu, myu->t4, myu->t4->c) map(myu->t4->c[2:5])
+ {
+ myu->t4->c[2]++;
+ }
+
+ #pragma omp target map(to: myu, myu->t4, myu->t4->d) map(myu->t4->d[2:5])
+ {
+ myu->t4->d[2]++;
+ }
+
+ assert (myu->t4->c[2] == 4);
+ assert (myu->t4->d[2] == 4);
+
+ delete s4;
+ delete t4;
+ delete myu;
+}
+#endif
+
+#ifdef REF2PTR_COMPONENT_BASE
+template<typename A>
+void
+ref2ptr_component_base (void)
+{
+ INIT_ST;
+ U<A> *myu_ptr = new U<A>(s1, t1, s2, t2, &s3, &t3, s4, t4);
+ U<A> *&myu = myu_ptr;
+
+ #pragma omp target map(myu->s1.a, myu->s1.b, myu->s1.c, myu->s1.d)
+ {
+ myu->s1.a++;
+ myu->s1.b++;
+ myu->s1.c++;
+ myu->s1.d++;
+ }
+
+ assert (myu->s1.a == 1);
+ assert (myu->s1.b == 1);
+ assert (myu->s1.c == &c1 + 1);
+ assert (myu->s1.d == &d1 + 1);
+
+ #pragma omp target map(myu->s2.a, myu->s2.b, myu->s2.c, myu->s2.d)
+ {
+ myu->s2.a++;
+ myu->s2.b++;
+ myu->s2.c++;
+ myu->s2.d++;
+ }
+
+ assert (myu->s2.a == 1);
+ assert (myu->s2.b == 1);
+ assert (myu->s2.c == &c2 + 1);
+ assert (myu->s2.d == &d2 + 1);
+
+ #pragma omp target map(to:myu->s3) \
+ map(myu->s3->a, myu->s3->b, myu->s3->c, myu->s3->d)
+ {
+ myu->s3->a++;
+ myu->s3->b++;
+ myu->s3->c++;
+ myu->s3->d++;
+ }
+
+ assert (myu->s3->a == 1);
+ assert (myu->s3->b == 1);
+ assert (myu->s3->c == &c3 + 1);
+ assert (myu->s3->d == &d3 + 1);
+
+ #pragma omp target map(to:myu->s4) \
+ map(myu->s4->a, myu->s4->b, myu->s4->c, myu->s4->d)
+ {
+ myu->s4->a++;
+ myu->s4->b++;
+ myu->s4->c++;
+ myu->s4->d++;
+ }
+
+ assert (myu->s4->a == 1);
+ assert (myu->s4->b == 1);
+ assert (myu->s4->c == &c4 + 1);
+ assert (myu->s4->d == &d4 + 1);
+
+ delete s4;
+ delete t4;
+ delete myu_ptr;
+}
+#endif
+
+#ifdef REF2PTR_COMPONENT_MEMBER_SLICE
+template<typename A>
+void
+ref2ptr_component_member_slice (void)
+{
+ INIT_ST;
+ U<A> *myu_ptr = new U<A>(s1, t1, s2, t2, &s3, &t3, s4, t4);
+ U<A> *&myu = myu_ptr;
+
+ #pragma omp target map(myu->t1.a[2:5])
+ {
+ myu->t1.a[2]++;
+ }
+
+ #pragma omp target map(myu->t1.b[2:5])
+ {
+ myu->t1.b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu->t1.c)
+
+ #pragma omp target map(myu->t1.c[2:5])
+ {
+ myu->t1.c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t1.c)
+
+ #pragma omp target enter data map(to: myu->t1.d)
+
+ #pragma omp target map(myu->t1.d[2:5])
+ {
+ myu->t1.d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t1.d)
+
+ assert (myu->t1.a[2] == 1);
+ assert (myu->t1.b[2] == 3);
+ assert (myu->t1.c[2] == 3);
+ assert (myu->t1.d[2] == 3);
+
+ #pragma omp target map(myu->t2.a[2:5])
+ {
+ myu->t2.a[2]++;
+ }
+
+ #pragma omp target map(myu->t2.b[2:5])
+ {
+ myu->t2.b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu->t2.c)
+
+ #pragma omp target map(myu->t2.c[2:5])
+ {
+ myu->t2.c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t2.c)
+
+ #pragma omp target enter data map(to: myu->t2.d)
+
+ #pragma omp target map(myu->t2.d[2:5])
+ {
+ myu->t2.d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t2.d)
+
+ assert (myu->t2.a[2] == 1);
+ assert (myu->t2.b[2] == 3);
+ assert (myu->t2.c[2] == 3);
+ assert (myu->t2.d[2] == 3);
+
+ #pragma omp target enter data map(to: myu->t3)
+
+ #pragma omp target map(myu->t3->a[2:5])
+ {
+ myu->t3->a[2]++;
+ }
+
+ #pragma omp target map(myu->t3->b[2:5])
+ {
+ myu->t3->b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu->t3->c)
+
+ #pragma omp target map(myu->t3->c[2:5])
+ {
+ myu->t3->c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t3->c)
+
+ #pragma omp target enter data map(to: myu->t3->d)
+
+ #pragma omp target map(myu->t3->d[2:5])
+ {
+ myu->t3->d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t3, myu->t3->d)
+
+ assert (myu->t3->a[2] == 1);
+ assert (myu->t3->b[2] == 3);
+ assert (myu->t3->c[2] == 3);
+ assert (myu->t3->d[2] == 3);
+
+ #pragma omp target enter data map(to: myu->t4)
+
+ #pragma omp target map(myu->t4->a[2:5])
+ {
+ myu->t4->a[2]++;
+ }
+
+ #pragma omp target map(myu->t4->b[2:5])
+ {
+ myu->t4->b[2]++;
+ }
+
+ #pragma omp target enter data map(to: myu->t4->c)
+
+ #pragma omp target map(myu->t4->c[2:5])
+ {
+ myu->t4->c[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t4->c)
+
+ #pragma omp target enter data map(to: myu->t4->d)
+
+ #pragma omp target map(myu->t4->d[2:5])
+ {
+ myu->t4->d[2]++;
+ }
+
+ #pragma omp target exit data map(release: myu->t4, myu->t4->d)
+
+ assert (myu->t4->a[2] == 1);
+ assert (myu->t4->b[2] == 3);
+ assert (myu->t4->c[2] == 3);
+ assert (myu->t4->d[2] == 3);
+
+ delete s4;
+ delete t4;
+ delete myu_ptr;
+}
+#endif
+
+#ifdef REF2PTR_COMPONENT_MEMBER_SLICE_BASEPTR
+template<typename A>
+void
+ref2ptr_component_member_slice_baseptr (void)
+{
+ INIT_ST;
+ U<A> *myu_ptr = new U<A>(s1, t1, s2, t2, &s3, &t3, s4, t4);
+ U<A> *&myu = myu_ptr;
+
+ /* Implicit firstprivate 'myu'. */
+ #pragma omp target map(to: myu->t1.c) map(myu->t1.c[2:5])
+ {
+ myu->t1.c[2]++;
+ }
+
+ #pragma omp target map(to: myu->t1.d) map(myu->t1.d[2:5])
+ {
+ myu->t1.d[2]++;
+ }
+
+ assert (myu->t1.c[2] == 2);
+ assert (myu->t1.d[2] == 2);
+
+ /* Explicitly-mapped 'myu'. */
+ #pragma omp target map(to: myu, myu->t1.c) map(myu->t1.c[2:5])
+ {
+ myu->t1.c[2]++;
+ }
+
+ #pragma omp target map(to: myu, myu->t1.d) map(myu->t1.d[2:5])
+ {
+ myu->t1.d[2]++;
+ }
+
+ assert (myu->t1.c[2] == 4);
+ assert (myu->t1.d[2] == 4);
+
+ /* Implicit firstprivate 'myu'. */
+ #pragma omp target map(to: myu->t2.c) map(myu->t2.c[2:5])
+ {
+ myu->t2.c[2]++;
+ }
+
+ #pragma omp target map(to: myu->t2.d) map(myu->t2.d[2:5])
+ {
+ myu->t2.d[2]++;
+ }
+
+ assert (myu->t2.c[2] == 2);
+ assert (myu->t2.d[2] == 2);
+
+ /* Explicitly-mapped 'myu'. */
+ #pragma omp target map(to: myu, myu->t2.c) map(myu->t2.c[2:5])
+ {
+ myu->t2.c[2]++;
+ }
+
+ #pragma omp target map(to: myu, myu->t2.d) map(myu->t2.d[2:5])
+ {
+ myu->t2.d[2]++;
+ }
+
+ assert (myu->t2.c[2] == 4);
+ assert (myu->t2.d[2] == 4);
+
+ /* Implicit firstprivate 'myu'. */
+ #pragma omp target map(to: myu->t3, myu->t3->c) map(myu->t3->c[2:5])
+ {
+ myu->t3->c[2]++;
+ }
+
+ #pragma omp target map(to: myu->t3, myu->t3->d) map(myu->t3->d[2:5])
+ {
+ myu->t3->d[2]++;
+ }
+
+ assert (myu->t3->c[2] == 2);
+ assert (myu->t3->d[2] == 2);
+
+ /* Explicitly-mapped 'myu'. */
+ #pragma omp target map(to: myu, myu->t3, myu->t3->c) map(myu->t3->c[2:5])
+ {
+ myu->t3->c[2]++;
+ }
+
+ #pragma omp target map(to: myu, myu->t3, myu->t3->d) map(myu->t3->d[2:5])
+ {
+ myu->t3->d[2]++;
+ }
+
+ assert (myu->t3->c[2] == 4);
+ assert (myu->t3->d[2] == 4);
+
+ /* Implicit firstprivate 'myu'. */
+ #pragma omp target map(to: myu->t4, myu->t4->c) map(myu->t4->c[2:5])
+ {
+ myu->t4->c[2]++;
+ }
+
+ #pragma omp target map(to: myu->t4, myu->t4->d) map(myu->t4->d[2:5])
+ {
+ myu->t4->d[2]++;
+ }
+
+ assert (myu->t4->c[2] == 2);
+ assert (myu->t4->d[2] == 2);
+
+ /* Explicitly-mapped 'myu'. */
+ #pragma omp target map(to: myu, myu->t4, myu->t4->c) map(myu->t4->c[2:5])
+ {
+ myu->t4->c[2]++;
+ }
+
+ #pragma omp target map(to: myu, myu->t4, myu->t4->d) map(myu->t4->d[2:5])
+ {
+ myu->t4->d[2]++;
+ }
+
+ assert (myu->t4->c[2] == 4);
+ assert (myu->t4->d[2] == 4);
+
+ delete s4;
+ delete t4;
+ delete myu_ptr;
+}
+#endif
+
+int main (int argc, char *argv[])
+{
+#ifdef MAP_DECLS
+ map_decls<4> ();
+#endif
+
+#ifdef NONREF_DECL_BASE
+ nonref_decl_base<int> ();
+#endif
+#ifdef REF_DECL_BASE
+ ref_decl_base<int> ();
+#endif
+#ifdef PTR_DECL_BASE
+ ptr_decl_base<int> ();
+#endif
+#ifdef REF2PTR_DECL_BASE
+ ref2ptr_decl_base<int> ();
+#endif
+
+#ifdef ARRAY_DECL_BASE
+ array_decl_base<int> ();
+#endif
+#ifdef REF2ARRAY_DECL_BASE
+ ref2array_decl_base<int> ();
+#endif
+#ifdef PTR_OFFSET_DECL_BASE
+ ptr_offset_decl_base<int> ();
+#endif
+#ifdef REF2PTR_OFFSET_DECL_BASE
+ ref2ptr_offset_decl_base<int> ();
+#endif
+
+#ifdef MAP_SECTIONS
+ map_sections<int, 10> ();
+#endif
+
+#ifdef NONREF_DECL_MEMBER_SLICE
+ nonref_decl_member_slice<int, 10> ();
+#endif
+#ifdef NONREF_DECL_MEMBER_SLICE_BASEPTR
+ nonref_decl_member_slice_baseptr<int, 10> ();
+#endif
+#ifdef REF_DECL_MEMBER_SLICE
+ ref_decl_member_slice<int, 10> ();
+#endif
+#ifdef REF_DECL_MEMBER_SLICE_BASEPTR
+ ref_decl_member_slice_baseptr<int, 10> ();
+#endif
+#ifdef PTR_DECL_MEMBER_SLICE
+ ptr_decl_member_slice<int, 10> ();
+#endif
+#ifdef PTR_DECL_MEMBER_SLICE_BASEPTR
+ ptr_decl_member_slice_baseptr<int, 10> ();
+#endif
+#ifdef REF2PTR_DECL_MEMBER_SLICE
+ ref2ptr_decl_member_slice<int, 10> ();
+#endif
+#ifdef REF2PTR_DECL_MEMBER_SLICE_BASEPTR
+ ref2ptr_decl_member_slice_baseptr<int, 10> ();
+#endif
+
+#ifdef ARRAY_DECL_MEMBER_SLICE
+ array_decl_member_slice<int, 10> ();
+#endif
+#ifdef ARRAY_DECL_MEMBER_SLICE_BASEPTR
+ array_decl_member_slice_baseptr<int, 10> ();
+#endif
+#ifdef REF2ARRAY_DECL_MEMBER_SLICE
+ ref2array_decl_member_slice<int, 10> ();
+#endif
+#ifdef REF2ARRAY_DECL_MEMBER_SLICE_BASEPTR
+ ref2array_decl_member_slice_baseptr<int, 10> ();
+#endif
+#ifdef PTR_OFFSET_DECL_MEMBER_SLICE
+ ptr_offset_decl_member_slice<int, 10> ();
+#endif
+#ifdef PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
+ ptr_offset_decl_member_slice_baseptr<int, 10> ();
+#endif
+#ifdef REF2PTR_OFFSET_DECL_MEMBER_SLICE
+ ref2ptr_offset_decl_member_slice<int, 10> ();
+#endif
+#ifdef REF2PTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
+ ref2ptr_offset_decl_member_slice_baseptr<int, 10> ();
+#endif
+
+#ifdef PTRARRAY_DECL_MEMBER_SLICE
+ ptrarray_decl_member_slice<int, 10> ();
+#endif
+#ifdef PTRARRAY_DECL_MEMBER_SLICE_BASEPTR
+ ptrarray_decl_member_slice_baseptr<int, 10> ();
+#endif
+#ifdef REF2PTRARRAY_DECL_MEMBER_SLICE
+ ref2ptrarray_decl_member_slice<int, 10> ();
+#endif
+#ifdef REF2PTRARRAY_DECL_MEMBER_SLICE_BASEPTR
+ ref2ptrarray_decl_member_slice_baseptr<int, 10> ();
+#endif
+#ifdef PTRPTR_OFFSET_DECL_MEMBER_SLICE
+ ptrptr_offset_decl_member_slice<int, 10> ();
+#endif
+#ifdef PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
+ ptrptr_offset_decl_member_slice_baseptr<int, 10> ();
+#endif
+#ifdef REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE
+ ref2ptrptr_offset_decl_member_slice<int, 10> ();
+#endif
+#ifdef REF2PTRPTR_OFFSET_DECL_MEMBER_SLICE_BASEPTR
+ ref2ptrptr_offset_decl_member_slice_baseptr<int, 10> ();
+#endif
+
+#ifdef NONREF_COMPONENT_BASE
+ nonref_component_base<int> ();
+#endif
+#ifdef NONREF_COMPONENT_MEMBER_SLICE
+ nonref_component_member_slice<int> ();
+#endif
+#ifdef NONREF_COMPONENT_MEMBER_SLICE_BASEPTR
+ nonref_component_member_slice_baseptr<int> ();
+#endif
+
+#ifdef REF_COMPONENT_BASE
+ ref_component_base<int> ();
+#endif
+#ifdef REF_COMPONENT_MEMBER_SLICE
+ ref_component_member_slice<int> ();
+#endif
+#ifdef REF_COMPONENT_MEMBER_SLICE_BASEPTR
+ ref_component_member_slice_baseptr<int> ();
+#endif
+
+#ifdef PTR_COMPONENT_BASE
+ ptr_component_base<int> ();
+#endif
+#ifdef PTR_COMPONENT_MEMBER_SLICE
+ ptr_component_member_slice<int> ();
+#endif
+#ifdef PTR_COMPONENT_MEMBER_SLICE_BASEPTR
+ ptr_component_member_slice_baseptr<int> ();
+#endif
+
+#ifdef REF2PTR_COMPONENT_BASE
+ ref2ptr_component_base<int> ();
+#endif
+#ifdef REF2PTR_COMPONENT_MEMBER_SLICE
+ ref2ptr_component_member_slice<int> ();
+#endif
+#ifdef REF2PTR_COMPONENT_MEMBER_SLICE_BASEPTR
+ ref2ptr_component_member_slice_baseptr<int> ();
+#endif
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,162 @@
+// { dg-do run }
+// { dg-options "-fopenmp" }
+
+#include <cassert>
+
+struct S
+{
+ int x[10];
+};
+
+struct T
+{
+ struct S *s;
+};
+
+struct U
+{
+ struct T *t;
+};
+
+void
+foo_siblist (void)
+{
+ U *u = new U;
+ u->t = new T;
+ u->t->s = new S;
+ for (int i = 0; i < 10; i++)
+ u->t->s->x[i] = 0;
+#pragma omp target map(u->t, *(u->t), u->t->s, *u->t->s)
+ for (int i = 0; i < 10; i++)
+ u->t->s->x[i] = i * 3;
+ for (int i = 0; i < 10; i++)
+ assert (u->t->s->x[i] == i * 3);
+ delete u->t->s;
+ delete u->t;
+ delete u;
+}
+
+void
+foo (void)
+{
+ U *u = new U;
+ u->t = new T;
+ u->t->s = new S;
+ for (int i = 0; i < 10; i++)
+ u->t->s->x[i] = 0;
+#pragma omp target map(*u, u->t, *(u->t), u->t->s, *u->t->s)
+ for (int i = 0; i < 10; i++)
+ u->t->s->x[i] = i * 3;
+ for (int i = 0; i < 10; i++)
+ assert (u->t->s->x[i] == i * 3);
+ delete u->t->s;
+ delete u->t;
+ delete u;
+}
+
+void
+foo_tofrom (void)
+{
+ U *u = new U;
+ u->t = new T;
+ u->t->s = new S;
+ for (int i = 0; i < 10; i++)
+ u->t->s->x[i] = 0;
+#pragma omp target map(u, *u, u->t, *(u->t), u->t->s, *u->t->s)
+ for (int i = 0; i < 10; i++)
+ u->t->s->x[i] = i * 3;
+ for (int i = 0; i < 10; i++)
+ assert (u->t->s->x[i] == i * 3);
+ delete u->t->s;
+ delete u->t;
+ delete u;
+}
+
+void
+bar (void)
+{
+ U *u = new U;
+ U **up = &u;
+ u->t = new T;
+ u->t->s = new S;
+ for (int i = 0; i < 10; i++)
+ (*up)->t->s->x[i] = 0;
+#pragma omp target map(*up, (*up)->t, *(*up)->t, (*up)->t->s, *(*up)->t->s)
+ for (int i = 0; i < 10; i++)
+ (*up)->t->s->x[i] = i * 3;
+ for (int i = 0; i < 10; i++)
+ assert ((*up)->t->s->x[i] == i * 3);
+ delete u->t->s;
+ delete u->t;
+ delete u;
+}
+
+void
+bar_pp (void)
+{
+ U *u = new U;
+ U **up = &u;
+ u->t = new T;
+ u->t->s = new S;
+ for (int i = 0; i < 10; i++)
+ (*up)->t->s->x[i] = 0;
+#pragma omp target map(*up, **up, (*up)->t, *(*up)->t, (*up)->t->s, *(*up)->t->s)
+ for (int i = 0; i < 10; i++)
+ (*up)->t->s->x[i] = i * 3;
+ for (int i = 0; i < 10; i++)
+ assert ((*up)->t->s->x[i] == i * 3);
+ delete u->t->s;
+ delete u->t;
+ delete u;
+}
+
+void
+bar_tofrom (void)
+{
+ U *u = new U;
+ U **up = &u;
+ u->t = new T;
+ u->t->s = new S;
+ for (int i = 0; i < 10; i++)
+ (*up)->t->s->x[i] = 0;
+#pragma omp target map(*up, up, (*up)->t, *(*up)->t, (*up)->t->s, *(*up)->t->s)
+ for (int i = 0; i < 10; i++)
+ (*up)->t->s->x[i] = i * 3;
+ for (int i = 0; i < 10; i++)
+ assert ((*up)->t->s->x[i] == i * 3);
+ delete u->t->s;
+ delete u->t;
+ delete u;
+}
+
+void
+bar_tofrom_pp (void)
+{
+ U *u = new U;
+ U **up = &u;
+ u->t = new T;
+ u->t->s = new S;
+ for (int i = 0; i < 10; i++)
+ (*up)->t->s->x[i] = 0;
+#pragma omp target map(**up, *up, up, (*up)->t, *(*up)->t, (*up)->t->s, \
+ *(*up)->t->s)
+ for (int i = 0; i < 10; i++)
+ (*up)->t->s->x[i] = i * 3;
+ for (int i = 0; i < 10; i++)
+ assert ((*up)->t->s->x[i] == i * 3);
+ delete u->t->s;
+ delete u->t;
+ delete u;
+}
+
+int main (int argc, char *argv[])
+{
+ foo_siblist ();
+ foo ();
+ foo_tofrom ();
+ bar ();
+ bar_pp ();
+ bar_tofrom ();
+ bar_tofrom_pp ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,93 @@
+// { dg-do run }
+// { dg-options "-fopenmp" }
+
+#include <cassert>
+
+struct S
+{
+ int x[10];
+};
+
+struct T
+{
+ struct S ***s;
+};
+
+struct U
+{
+ struct T **t;
+};
+
+void
+foo (void)
+{
+ U *u = new U;
+ T *real_t = new T;
+ S *real_s = new S;
+ T **t_pp = &real_t;
+ S **s_pp = &real_s;
+ S ***s_ppp = &s_pp;
+ u->t = t_pp;
+ (*u->t)->s = s_ppp;
+ for (int i = 0; i < 10; i++)
+ (**((*u->t)->s))->x[i] = 0;
+#pragma omp target map(u->t, *u->t, (*u->t)->s, *(*u->t)->s, **(*u->t)->s, \
+ (**(*u->t)->s)->x[0:10])
+ for (int i = 0; i < 10; i++)
+ (**((*u->t)->s))->x[i] = i * 3;
+ for (int i = 0; i < 10; i++)
+ assert ((**((*u->t)->s))->x[i] == i * 3);
+ delete real_s;
+ delete real_t;
+ delete u;
+}
+
+template<typename X>
+struct St
+{
+ X x[10];
+};
+
+template<typename X>
+struct Tt
+{
+ X ***s;
+};
+
+template<typename X>
+struct Ut
+{
+ X **t;
+};
+
+template<typename I>
+void
+tfoo (void)
+{
+ Ut<Tt<St<I> > > *u = new Ut<Tt<St<I> > >;
+ Tt<St<I> > *real_t = new Tt<St<int> >;
+ St<I> *real_s = new St<int>;
+ Tt<St<I> > **t_pp = &real_t;
+ St<I> **s_pp = &real_s;
+ St<I> ***s_ppp = &s_pp;
+ u->t = t_pp;
+ (*u->t)->s = s_ppp;
+ for (int i = 0; i < 10; i++)
+ (**((*u->t)->s))->x[i] = 0;
+#pragma omp target map(u->t, *u->t, (*u->t)->s, *(*u->t)->s, **(*u->t)->s, \
+ (**(*u->t)->s)->x[0:10])
+ for (int i = 0; i < 10; i++)
+ (**((*u->t)->s))->x[i] = i * 3;
+ for (int i = 0; i < 10; i++)
+ assert ((**((*u->t)->s))->x[i] == i * 3);
+ delete real_s;
+ delete real_t;
+ delete u;
+}
+
+int main (int argc, char *argv[])
+{
+ foo ();
+ tfoo<int> ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,75 @@
+#include <cstring>
+#include <cassert>
+
+static int lo()
+{
+ return 30;
+}
+
+static int len()
+{
+ return 10;
+}
+
+template<typename T>
+void foo ()
+{
+ T arr[100];
+ T *ptr;
+
+ memset (arr, '\0', sizeof arr);
+
+#pragma omp target enter data map(to: arr[0:100])
+
+ for (int i = 0; i < 100; i++)
+ arr[i] = i;
+
+ ptr = &arr[10];
+
+#pragma omp target update to(*ptr)
+
+ for (int i = lo (); i < lo () + len (); i++)
+ arr[i] = i * 2;
+
+#pragma omp target update to(arr[lo():len()])
+
+#pragma omp target exit data map(from: arr[0:100])
+
+ assert (arr[10] == 10);
+ for (int i = lo (); i < lo () + len (); i++)
+ assert (arr[i] == i * 2);
+}
+
+int
+main ()
+{
+ char arr[100];
+ char *ptr;
+
+ memset (arr, '\0', sizeof arr);
+
+#pragma omp target enter data map(to: arr[0:100])
+
+ for (int i = 0; i < 100; i++)
+ arr[i] = i;
+
+ ptr = &arr[10];
+
+#pragma omp target update to(*ptr)
+
+ for (int i = lo (); i < lo () + len (); i++)
+ arr[i] = i * 2;
+
+#pragma omp target update to(arr[lo():len()])
+
+#pragma omp target exit data map(from: arr[0:100])
+
+ assert (arr[10] == 10);
+ for (int i = lo (); i < lo () + len (); i++)
+ assert (arr[i] == i * 2);
+
+ foo<short> ();
+
+ return 0;
+}
+
new file mode 100644
@@ -0,0 +1,71 @@
+#include <cstring>
+#include <cassert>
+#include <cstdlib>
+
+template<typename T>
+struct t_array_wrapper {
+ T *data;
+ unsigned int length;
+};
+
+template<typename T>
+void foo()
+{
+ struct t_array_wrapper<T> aw;
+
+ aw.data = new T[100];
+ aw.length = 100;
+
+#pragma omp target enter data map(to: aw.data, aw.length) \
+ map(alloc: aw.data[0:aw.length])
+
+#pragma omp target
+ for (int i = 0; i < aw.length; i++)
+ aw.data[i] = i;
+
+#pragma omp target update from(aw.data[:aw.length])
+
+#pragma omp target exit data map(delete: aw.data, aw.length, \
+ aw.data[0:aw.length])
+
+ for (int i = 0; i < aw.length; i++)
+ assert (aw.data[i] == i);
+
+ delete[] aw.data;
+}
+
+struct array_wrapper {
+ int *data;
+ unsigned int length;
+};
+
+int
+main ()
+{
+ struct array_wrapper aw;
+
+ aw.data = new int[100];
+ aw.length = 100;
+
+#pragma omp target enter data map(to: aw.data, aw.length) \
+ map(alloc: aw.data[0:aw.length])
+
+#pragma omp target
+ for (int i = 0; i < aw.length; i++)
+ aw.data[i] = i;
+
+#pragma omp target update from(aw.data[:aw.length])
+
+#pragma omp target exit data map(delete: aw.data, aw.length, \
+ aw.data[0:aw.length])
+
+ for (int i = 0; i < aw.length; i++)
+ assert (aw.data[i] == i);
+
+ delete[] aw.data;
+
+ foo<unsigned long> ();
+
+ return 0;
+}
+
new file mode 100644
@@ -0,0 +1,15 @@
+/* { dg-do run } */
+
+#include <cassert>
+
+int main (int argc, char *argv[])
+{
+ int a = 5, b = 7;
+#pragma omp target map((a, b))
+ {
+ a++;
+ b++;
+ }
+ assert (a == 5 && b == 8);
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,22 @@
+/* { dg-do run } */
+
+#include <cassert>
+
+int foo (int &&x)
+{
+ int y;
+#pragma omp target map(x, y)
+ {
+ x++;
+ y = x;
+ }
+ return y;
+}
+
+int main (int argc, char *argv[])
+{
+ int y = 5;
+ y = foo (y + 3);
+ assert (y == 9);
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,97 @@
+// { dg-do run }
+// { dg-options "-fopenmp" }
+
+#include <cassert>
+
+struct S
+{
+ int x[10];
+};
+
+void
+foo (S *s, int x)
+{
+ S *&r = s;
+ for (int i = 0; i < x; i++)
+ s[0].x[i] = s[1].x[i] = 0;
+ #pragma omp target map (s, x)
+ ;
+ #pragma omp target map (s[0], x)
+ for (int i = 0; i < x; i++)
+ s[0].x[i] = i;
+ #pragma omp target map (s[1], x)
+ for (int i = 0; i < x; i++)
+ s[1].x[i] = i * 2;
+ for (int i = 0; i < x; i++)
+ {
+ assert (s[0].x[i] == i);
+ assert (s[1].x[i] == i * 2);
+ s[0].x[i] = 0;
+ s[1].x[i] = 0;
+ }
+ #pragma omp target map (r, x)
+ ;
+ #pragma omp target map (r[0], x)
+ for (int i = 0; i < x; i++)
+ r[0].x[i] = i;
+ #pragma omp target map (r[1], x)
+ for (int i = 0; i < x; i++)
+ r[1].x[i] = i * 2;
+ for (int i = 0; i < x; i++)
+ {
+ assert (r[0].x[i] == i);
+ assert (r[1].x[i] == i * 2);
+ }
+}
+
+template <int N>
+struct T
+{
+ int x[N];
+};
+
+template <int N>
+void
+bar (T<N> *t, int x)
+{
+ T<N> *&r = t;
+ for (int i = 0; i < x; i++)
+ t[0].x[i] = t[1].x[i] = 0;
+ #pragma omp target map (t, x)
+ ;
+ #pragma omp target map (t[0], x)
+ for (int i = 0; i < x; i++)
+ t[0].x[i] = i;
+ #pragma omp target map (t[1], x)
+ for (int i = 0; i < x; i++)
+ t[1].x[i] = i * 2;
+ for (int i = 0; i < x; i++)
+ {
+ assert (t[0].x[i] == i);
+ assert (t[1].x[i] == i * 2);
+ t[0].x[i] = 0;
+ t[1].x[i] = 0;
+ }
+ #pragma omp target map (r, x)
+ ;
+ #pragma omp target map (r[0], x)
+ for (int i = 0; i < x; i++)
+ r[0].x[i] = i;
+ #pragma omp target map (r[1], x)
+ for (int i = 0; i < x; i++)
+ r[1].x[i] = i * 2;
+ for (int i = 0; i < x; i++)
+ {
+ assert (r[0].x[i] == i);
+ assert (r[1].x[i] == i * 2);
+ }
+}
+
+int main (int argc, char *argv[])
+{
+ S s[2];
+ foo (s, 10);
+ T<10> t[2];
+ bar (t, 10);
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,35 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <string.h>
+#include <assert.h>
+
+#define N 16
+
+struct Z {
+ int *ptr;
+ int arr[N];
+ int c;
+};
+
+int main (int argc, char *argv[])
+{
+ struct Z *myz;
+ myz = (struct Z *) calloc (1, sizeof *myz);
+
+#pragma omp target map(tofrom:myz->arr[0:N], myz->c)
+ {
+ for (int i = 0; i < N; i++)
+ myz->arr[i]++;
+ myz->c++;
+ }
+
+ for (int i = 0; i < N; i++)
+ assert (myz->arr[i] == 1);
+ assert (myz->c == 1);
+
+ free (myz);
+
+ return 0;
+}
+
new file mode 100644
@@ -0,0 +1,65 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <string.h>
+#include <assert.h>
+
+#define N 16
+
+/* NOTE: This test is the same as array-of-struct-2.c, except the fields of
+ this struct are in a different order. */
+
+struct Z {
+ int arr[N];
+ int *ptr;
+ int c;
+};
+
+void
+foo (struct Z *zarr, int len)
+{
+#pragma omp target map(to:zarr, zarr[5].ptr) map(tofrom:zarr[5].ptr[0:len])
+ {
+ for (int i = 0; i < len; i++)
+ zarr[5].ptr[i]++;
+ }
+
+#pragma omp target map(to:zarr) map(tofrom:zarr[4].arr[0:len])
+ {
+ for (int i = 0; i < len; i++)
+ zarr[4].arr[i]++;
+ }
+
+#pragma omp target map (to:zarr[3].ptr) map(tofrom:zarr[3].ptr[0:len])
+ {
+ for (int i = 0; i < len; i++)
+ zarr[3].ptr[i]++;
+ }
+
+#pragma omp target map(tofrom:zarr[2].arr[0:len])
+ {
+ for (int i = 0; i < len; i++)
+ zarr[2].arr[i]++;
+ }
+}
+
+int main (int argc, char *argv[])
+{
+ struct Z zs[10];
+ memset (zs, 0, sizeof zs);
+
+ for (int i = 0; i < 10; i++)
+ zs[i].ptr = (int *) calloc (N, sizeof (int));
+
+ foo (zs, N);
+
+ for (int i = 0; i < N; i++)
+ {
+ assert (zs[2].arr[i] == 1);
+ assert (zs[4].arr[i] == 1);
+ assert (zs[3].ptr[i] == 1);
+ assert (zs[5].ptr[i] == 1);
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,65 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <string.h>
+#include <assert.h>
+
+#define N 16
+
+/* NOTE: This test is the same as array-of-struct-1.c, except the fields of
+ this struct are in a different order. */
+
+struct Z {
+ int *ptr;
+ int arr[N];
+ int c;
+};
+
+void
+foo (struct Z *zarr, int len)
+{
+#pragma omp target map(to:zarr, zarr[5].ptr) map(tofrom:zarr[5].ptr[0:len])
+ {
+ for (int i = 0; i < len; i++)
+ zarr[5].ptr[i]++;
+ }
+
+#pragma omp target map(to:zarr) map(tofrom:zarr[4].arr[0:len])
+ {
+ for (int i = 0; i < len; i++)
+ zarr[4].arr[i]++;
+ }
+
+#pragma omp target map (to:zarr[3].ptr) map(tofrom:zarr[3].ptr[0:len])
+ {
+ for (int i = 0; i < len; i++)
+ zarr[3].ptr[i]++;
+ }
+
+#pragma omp target map(tofrom:zarr[2].arr[0:len])
+ {
+ for (int i = 0; i < len; i++)
+ zarr[2].arr[i]++;
+ }
+}
+
+int main (int argc, char *argv[])
+{
+ struct Z zs[10];
+ memset (zs, 0, sizeof zs);
+
+ for (int i = 0; i < 10; i++)
+ zs[i].ptr = (int *) calloc (N, sizeof (int));
+
+ foo (zs, N);
+
+ for (int i = 0; i < N; i++)
+ {
+ assert (zs[2].arr[i] == 1);
+ assert (zs[4].arr[i] == 1);
+ assert (zs[3].ptr[i] == 1);
+ assert (zs[5].ptr[i] == 1);
+ }
+
+ return 0;
+}