pp_c_right_bracket (this);
break;
+ case OMP_ARRAY_SECTION:
+ postfix_expression (TREE_OPERAND (e, 0));
+ pp_c_left_bracket (this);
+ if (TREE_OPERAND (e, 1))
+ expression (TREE_OPERAND (e, 1));
+ pp_colon (this);
+ if (TREE_OPERAND (e, 2))
+ expression (TREE_OPERAND (e, 2));
+ pp_c_right_bracket (this);
+ break;
+
case CALL_EXPR:
{
call_expr_arg_iterator iter;
case POSTINCREMENT_EXPR:
case POSTDECREMENT_EXPR:
case ARRAY_REF:
+ case OMP_ARRAY_SECTION:
case CALL_EXPR:
case COMPONENT_REF:
case BIT_FIELD_REF:
location_t brace_loc = c_parser_peek_token (parser)->location;
gcc_obstack_init (&braced_init_obstack);
gcc_assert (c_parser_next_token_is (parser, CPP_OPEN_BRACE));
+ bool save_c_omp_array_section_p = c_omp_array_section_p;
+ c_omp_array_section_p = false;
matching_braces braces;
braces.consume_open (parser);
if (nested_p)
break;
}
}
+ c_omp_array_section_p = save_c_omp_array_section_p;
c_token *next_tok = c_parser_peek_token (parser);
if (next_tok->type != CPP_CLOSE_BRACE)
{
{
struct c_expr cond, exp1, exp2, ret;
location_t start, cond_loc, colon_loc;
+ bool save_c_omp_array_section_p = c_omp_array_section_p;
gcc_assert (!after || c_dialect_objc ());
if (c_parser_next_token_is_not (parser, CPP_QUERY))
return cond;
+ c_omp_array_section_p = false;
if (cond.value != error_mark_node)
start = cond.get_start ();
else
ret.set_error ();
ret.original_code = ERROR_MARK;
ret.original_type = NULL;
+ c_omp_array_section_p = save_c_omp_array_section_p;
return ret;
}
{
}
set_c_expr_source_range (&ret, start, exp2.get_finish ());
ret.m_decimal = 0;
+ c_omp_array_section_p = save_c_omp_array_section_p;
return ret;
}
/* A statement expression. */
tree stmt;
location_t brace_loc;
+ bool save_c_omp_array_section_p = c_omp_array_section_p;
c_parser_consume_token (parser);
brace_loc = c_parser_peek_token (parser)->location;
c_parser_consume_token (parser);
expr.set_error ();
break;
}
+ c_omp_array_section_p = false;
stmt = c_begin_stmt_expr ();
c_parser_compound_statement_nostart (parser);
location_t close_loc = c_parser_peek_token (parser)->location;
expr.value = c_finish_stmt_expr (brace_loc, stmt);
set_c_expr_source_range (&expr, loc, close_loc);
mark_exp_read (expr.value);
+ c_omp_array_section_p = save_c_omp_array_section_p;
}
else
{
struct c_expr expr)
{
struct c_expr orig_expr;
- tree ident, idx;
+ tree ident, idx, len;
location_t sizeof_arg_loc[6], comp_loc;
tree sizeof_arg[6];
unsigned int literal_zero_mask;
case CPP_OPEN_SQUARE:
/* Array reference. */
c_parser_consume_token (parser);
- idx = c_parser_expression (parser).value;
+ idx = len = NULL_TREE;
+ if (!c_omp_array_section_p
+ || c_parser_next_token_is_not (parser, CPP_COLON))
+ idx = c_parser_expression (parser).value;
+
+ if (c_omp_array_section_p
+ && c_parser_next_token_is (parser, CPP_COLON))
+ {
+ c_parser_consume_token (parser);
+ if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE))
+ len = c_parser_expression (parser).value;
+
+ expr.value = build_omp_array_section (op_loc, expr.value, idx,
+ len);
+ }
+ else
+ expr.value = build_array_ref (op_loc, expr.value, idx);
+
c_parser_skip_until_found (parser, CPP_CLOSE_SQUARE,
"expected %<]%>");
+
start = expr.get_start ();
finish = parser->tokens_buf[0].location;
- expr.value = build_array_ref (op_loc, expr.value, idx);
set_c_expr_source_range (&expr, start, finish);
expr.original_code = ERROR_MARK;
expr.original_type = NULL;
vec<tree, va_gc> *orig_types;
struct c_expr expr;
unsigned int idx = 0;
+ bool save_c_omp_array_section_p = c_omp_array_section_p;
+ c_omp_array_section_p = false;
ret = make_tree_vector ();
if (p_orig_types == NULL)
}
if (orig_types)
*p_orig_types = orig_types;
+ c_omp_array_section_p = save_c_omp_array_section_p;
return ret;
}
\f
c_parser_omp_variable_list (c_parser *parser,
location_t clause_loc,
enum omp_clause_code kind, tree list,
- bool allow_deref = false)
+ bool map_lvalue = false)
{
auto_vec<omp_dim> dims;
bool array_section_p;
while (1)
{
+ tree t = NULL_TREE;
+
if (kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY)
{
if (c_parser_next_token_is_not (parser, CPP_NAME)
parser->tokens = tokens.address ();
parser->tokens_avail = tokens.length ();
}
+ else if (map_lvalue
+ && (kind == OMP_CLAUSE_MAP
+ || kind == OMP_CLAUSE_TO
+ || kind == OMP_CLAUSE_FROM))
+ {
+ location_t loc = c_parser_peek_token (parser)->location;
+ bool save_c_omp_array_section_p = c_omp_array_section_p;
+ c_omp_array_section_p = true;
+ c_expr expr = c_parser_expr_no_commas (parser, NULL);
+ if (expr.value != error_mark_node)
+ mark_exp_read (expr.value);
+ c_omp_array_section_p = save_c_omp_array_section_p;
+ tree decl = expr.value;
+
+ /* 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
+ c-typeck.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. (See also equivalent code in cp/parser.cc,
+ cp/semantics.cc). */
+ 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);
+ }
- tree t = NULL_TREE;
+ while (TREE_CODE (decl) == ARRAY_REF
+ || TREE_CODE (decl) == INDIRECT_REF
+ || TREE_CODE (decl) == COMPOUND_EXPR)
+ {
+ 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);
+ }
+ }
+
+ for (int i = dims.length () - 1; i >= 0; i--)
+ decl = build_omp_array_section (loc, decl, dims[i].low_bound,
+ dims[i].length);
+ }
+ else if (TREE_CODE (decl) == INDIRECT_REF)
+ {
+ /* Turn indirection of a pointer "*foo" into "foo[0:1]". */
+ decl = TREE_OPERAND (decl, 0);
+ STRIP_NOPS (decl);
+
+ decl = build_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 = build_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);
+
+ tree u = build_omp_clause (clause_loc, kind);
+ OMP_CLAUSE_DECL (u) = decl;
+ OMP_CLAUSE_CHAIN (u) = list;
+ list = u;
+
+ goto next_item;
+ }
if (c_parser_next_token_is (parser, CPP_NAME)
&& c_parser_peek_token (parser)->id_kind == C_ID_ID)
case OMP_CLAUSE_TO:
start_component_ref:
while (c_parser_next_token_is (parser, CPP_DOT)
- || (allow_deref
- && c_parser_next_token_is (parser, CPP_DEREF)))
+ || c_parser_next_token_is (parser, CPP_DEREF))
{
location_t op_loc = c_parser_peek_token (parser)->location;
location_t arrow_loc = UNKNOWN_LOCATION;
|| kind == OMP_CLAUSE_TO)
&& !array_section_p
&& (c_parser_next_token_is (parser, CPP_DOT)
- || (allow_deref
- && c_parser_next_token_is (parser,
- CPP_DEREF))))
+ || c_parser_next_token_is (parser, CPP_DEREF)))
{
for (unsigned i = 0; i < dims.length (); i++)
{
}
else
for (unsigned i = 0; i < dims.length (); i++)
- t = tree_cons (dims[i].low_bound, dims[i].length, t);
+ t = build_omp_array_section (clause_loc, t,
+ dims[i].low_bound,
+ dims[i].length);
}
if ((kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY)
parser->tokens = saved_tokens;
parser->tokens_avail = tokens_avail;
}
+
+ next_item:
if (c_parser_next_token_is_not (parser, CPP_COMMA))
break;
static tree
c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
- tree list, bool allow_deref = false)
+ tree list, bool map_lvalue = false)
{
/* The clauses location. */
location_t loc = c_parser_peek_token (parser)->location;
matching_parens parens;
if (parens.require_open (parser))
{
- list = c_parser_omp_variable_list (parser, loc, kind, list, allow_deref);
+ list = c_parser_omp_variable_list (parser, loc, kind, list, map_lvalue);
parens.skip_until_found_close (parser);
}
return list;
gcc_unreachable ();
}
tree nl, c;
- nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true);
+ nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, false);
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_SET_MAP_KIND (c, kind);
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
{
tree d = OMP_CLAUSE_DECL (c), type;
- if (TREE_CODE (d) != TREE_LIST)
+ if (TREE_CODE (d) != OMP_ARRAY_SECTION)
type = TREE_TYPE (d);
else
{
int cnt = 0;
tree t;
- for (t = d; TREE_CODE (t) == TREE_LIST; t = TREE_CHAIN (t))
+ for (t = d;
+ TREE_CODE (t) == OMP_ARRAY_SECTION;
+ t = TREE_OPERAND (t, 0))
cnt++;
type = TREE_TYPE (t);
while (cnt > 0)
extern int in_sizeof;
extern int in_typeof;
extern bool c_in_omp_for;
+extern bool c_omp_array_section_p;
extern tree c_last_sizeof_arg;
extern location_t c_last_sizeof_loc;
extern tree build_component_ref (location_t, tree, tree, location_t,
location_t);
extern tree build_array_ref (location_t, tree, tree);
+extern tree build_omp_array_section (location_t, tree, tree, tree);
extern tree build_external_ref (location_t, tree, bool, tree *);
extern void pop_maybe_used (bool);
extern struct c_expr c_expr_sizeof_expr (location_t, struct c_expr);
/* True when parsing OpenMP loop expressions. */
bool c_in_omp_for;
+/* True when parsing OpenMP map clause. */
+bool c_omp_array_section_p;
+
/* The argument of last parsed sizeof expression, only to be tested
if expr.original_code == SIZEOF_EXPR. */
tree c_last_sizeof_arg;
case C_MAYBE_CONST_EXPR:
mark_exp_read (TREE_OPERAND (exp, 1));
break;
+ case OMP_ARRAY_SECTION:
+ mark_exp_read (TREE_OPERAND (exp, 0));
+ if (TREE_OPERAND (exp, 1))
+ mark_exp_read (TREE_OPERAND (exp, 1));
+ if (TREE_OPERAND (exp, 2))
+ mark_exp_read (TREE_OPERAND (exp, 2));
+ break;
default:
break;
}
return ret;
}
}
+
+/* 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, tree index, tree length)
+{
+ tree type = TREE_TYPE (array);
+ gcc_assert (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 || error_operand_p (eltype))
+ sectype = TREE_TYPE (array);
+ else
+ {
+ tree idxtype = NULL_TREE;
+
+ if (index != NULL_TREE
+ && length != NULL_TREE
+ && INTEGRAL_TYPE_P (TREE_TYPE (index))
+ && INTEGRAL_TYPE_P (TREE_TYPE (length)))
+ {
+ 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
+ && INTEGRAL_TYPE_P (TREE_TYPE (length)))
+ idxtype = build_index_type (length);
+
+ gcc_assert (!error_operand_p (idxtype));
+
+ sectype = build_array_type (eltype, idxtype);
+ }
+
+ return build3_loc (loc, OMP_ARRAY_SECTION, sectype, array, index, length);
+}
+
\f
/* Build an external reference to identifier ID. FUN indicates
whether this will be used for a function call. LOC is the source
return error_mark_node;
}
- if (TREE_TYPE (ref) == error_mark_node)
+ /* For an OpenMP map clause, we can get better diagnostics for decls with
+ unmappable types if we return the decl with an error_mark_node type,
+ rather than returning error_mark_node for the decl itself. */
+ if (TREE_TYPE (ref) == error_mark_node
+ && !c_omp_array_section_p)
return error_mark_node;
if (TREE_UNAVAILABLE (ref))
{
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;
t = ai.unconverted_ref_origin ();
if (t == error_mark_node)
return error_mark_node;
- if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
+ if (!VAR_P (t)
+ && (ort == C_ORT_ACC || !EXPR_P (t))
+ && TREE_CODE (t) != PARM_DECL)
{
if (DECL_P (t))
error_at (OMP_CLAUSE_LOCATION (c),
return ret;
}
- ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
+ 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 == error_mark_node || length == error_mark_node)
return error_mark_node;
tree lb = save_expr (low_bound);
if (lb != low_bound)
{
- TREE_PURPOSE (t) = lb;
+ TREE_OPERAND (t, 1) = lb;
low_bound = lb;
}
}
array-section-subscript, the array section could be non-contiguous. */
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
- && 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),
tree lb = save_expr (low_bound);
if (lb != low_bound)
{
- TREE_PURPOSE (t) = lb;
+ TREE_OPERAND (t, 1) = lb;
low_bound = lb;
}
ret = build_array_ref (OMP_CLAUSE_LOCATION (c), ret, low_bound);
maybe_zero_len = true;
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);
+ tree low_bound = TREE_OPERAND (t, 1);
+ tree length = TREE_OPERAND (t, 2);
i--;
if (low_bound
{
tree 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);
if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
{
case OMP_CLAUSE_TASK_REDUCTION:
need_implicitly_determined = true;
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))
{
}
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;
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);
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
+ if (ort != C_ORT_ACC && EXPR_P (t))
+ break;
+
error_at (OMP_CLAUSE_LOCATION (c),
"%qE is not a variable in %qs clause", t,
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
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;
--- /dev/null
+/* { dg-do compile } */
+
+int foo (int *ptr);
+
+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 {passing argument 1 of 'foo' makes pointer from integer without a cast} "" { target *-*-* } .-2 } */
+/* { dg-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-3 } */
+ { }
+
+ return 0;
+}
--- /dev/null
+/* { dg-do compile } */
+
+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-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-2 } */
+ { }
+
+ return 0;
+}
--- /dev/null
+/* { dg-do compile } */
+
+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-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-3 } */
+ { }
+
+ /* ...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} "" { target *-*-* } .-1 } */
+ { }
+
+ return 0;
+}
--- /dev/null
+/* { dg-do compile } */
+
+int x;
+
+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 } */
+ { }
+
+ return 0;
+}
--- /dev/null
+/* { dg-do compile } */
+
+int partly = 0;
+
+int main()
+{
+ int arr[20];
+#pragma omp target map(partly ? arr[5:5] : arr)
+/* { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 } */
+/* { dg-error {pointer/integer type mismatch in conditional expression} "" { target *-*-* } .-2 } */
+/* { dg-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-3 } */
+ { }
+
+ return 0;
+}
--- /dev/null
+/* { dg-do compile } */
+
+int x;
+
+int main()
+{
+ int arr[20];
+#pragma omp target map(arr[5:5] * 2)
+/* { dg-error {invalid operands to binary \*} "" { target *-*-* } .-1 } */
+ { }
+#pragma omp target map(arr[x:5] * 2)
+/* { dg-error {invalid operands to binary \*} "" { target *-*-* } .-1 } */
+ { }
+
+ return 0;
+}
--- /dev/null
+/* { dg-do compile } */
+
+int x;
+
+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 {'\(struct T \*\)&s->tvec\[3:5\]' is a pointer; did you mean to use '->'\?} "" { target *-*-* } .-1 } */
+ { }
+#pragma omp target map(s->tvec[5:x].arr[0:20])
+/* { dg-error {'\(struct T \*\)&s->tvec\[5:x\]' is a pointer; did you mean to use '->'\?} "" { target *-*-* } .-1 } */
+ { }
+
+ return 0;
+}
--- /dev/null
+/* { dg-do compile } */
+
+int x;
+
+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 } */
+ { }
+
+ return 0;
+}
+
@item Discontiguous array section with @code{target update} construct
@tab N @tab
@item C/C++'s lvalue expressions in @code{to}, @code{from}
- and @code{map} clauses @tab N @tab
+ and @code{map} clauses @tab Y @tab
@item C/C++'s lvalue expressions in @code{depend} clauses @tab Y @tab
@item Nested @code{declare target} directive @tab Y @tab
@item Combined @code{master} constructs @tab Y @tab
--- /dev/null
+// { dg-do run }
+// { dg-options "-fopenmp" }
+
+#include <assert.h>
+#include <stdlib.h>
+
+typedef struct
+{
+ int x[10];
+} S;
+
+typedef struct
+{
+ S ***s;
+} T;
+
+typedef struct
+{
+ T **t;
+} U;
+
+void
+foo (void)
+{
+ U *u = (U *) malloc (sizeof (U));
+ T *real_t = (T *) malloc (sizeof (T));
+ S *real_s = (S *) malloc (sizeof (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);
+ free (real_s);
+ free (real_t);
+ free (u);
+}
+
+int main (int argc, char *argv[])
+{
+ foo ();
+ return 0;
+}
--- /dev/null
+#include <assert.h>
+
+int main (int argc, char *argv[])
+{
+ int y = 0;
+ int *x = &y;
+
+#pragma omp target map(*x)
+ {
+ (*x)++;
+ }
+
+ assert (y == 1);
+
+ return 0;
+}
projects / thirdparty / gcc.git / commitdiff
