@@ -10486,7 +10486,7 @@ c_parser_postfix_expression_after_primary (c_parser *parser,
struct c_expr expr)
{
struct c_expr orig_expr;
- tree ident, idx;
+ tree ident, idx, len;
location_t sizeof_arg_loc[3], comp_loc;
tree sizeof_arg[3];
unsigned int literal_zero_mask;
@@ -10505,15 +10505,44 @@ c_parser_postfix_expression_after_primary (c_parser *parser,
case CPP_OPEN_SQUARE:
/* Array reference. */
c_parser_consume_token (parser);
- idx = c_parser_expression (parser).value;
- 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;
+ 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;
+
+ c_parser_skip_until_found (parser, CPP_CLOSE_SQUARE,
+ "expected %<]%>");
+
+ /* NOTE: We are reusing using the type of the whole array as the
+ type of the array section here, which isn't necessarily
+ entirely correct. Might need revisiting. */
+ start = expr.get_start ();
+ finish = parser->tokens_buf[0].location;
+ expr.value = build3_loc (op_loc, OMP_ARRAY_SECTION,
+ TREE_TYPE (expr.value), expr.value,
+ idx, len);
+ set_c_expr_source_range (&expr, start, finish);
+ expr.original_code = ERROR_MARK;
+ expr.original_type = NULL;
+ }
+ else
+ {
+ 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;
+ }
break;
case CPP_OPEN_PAREN:
/* Function call. */
@@ -13009,7 +13038,7 @@ static tree
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;
@@ -13019,6 +13048,8 @@ c_parser_omp_variable_list (c_parser *parser,
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)
@@ -13087,8 +13118,84 @@ c_parser_omp_variable_list (c_parser *parser,
parser->tokens = tokens.address ();
parser->tokens_avail = tokens.length ();
}
+ else if (map_lvalue && kind == OMP_CLAUSE_MAP)
+ {
+ 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;
- tree t = NULL_TREE;
+ 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 (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 = tree_cons (dims[i].low_bound, dims[i].length, decl);
+ }
+ else if (TREE_CODE (decl) == INDIRECT_REF)
+ {
+ /* Turn *foo into the representation previously used for
+ foo[0]. */
+ decl = TREE_OPERAND (decl, 0);
+ STRIP_NOPS (decl);
+
+ decl = tree_cons (integer_zero_node, integer_one_node, decl);
+ }
+ else if (TREE_CODE (decl) == ARRAY_REF)
+ {
+ tree idx = TREE_OPERAND (decl, 1);
+
+ decl = TREE_OPERAND (decl, 0);
+ STRIP_NOPS (decl);
+
+ decl = tree_cons (idx, integer_one_node, decl);
+ }
+ 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)
@@ -13139,8 +13246,7 @@ c_parser_omp_variable_list (c_parser *parser,
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;
if (c_parser_next_token_is (parser, CPP_DEREF))
@@ -13227,9 +13333,7 @@ c_parser_omp_variable_list (c_parser *parser,
|| 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++)
{
@@ -13289,6 +13393,8 @@ c_parser_omp_variable_list (c_parser *parser,
parser->tokens = &parser->tokens_buf[0];
parser->tokens_avail = tokens_avail;
}
+
+ next_item:
if (c_parser_next_token_is_not (parser, CPP_COMMA))
break;
@@ -13305,7 +13411,7 @@ c_parser_omp_variable_list (c_parser *parser,
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;
@@ -13313,7 +13419,7 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
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;
@@ -13382,7 +13488,7 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
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);
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_SET_MAP_KIND (c, kind);
@@ -16490,7 +16596,7 @@ c_parser_omp_clause_device_type (c_parser *parser, tree list)
static tree
c_parser_omp_clause_to (c_parser *parser, tree list)
{
- return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO, list, true);
+ return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO, list);
}
/* OpenMP 4.0:
@@ -16499,7 +16605,7 @@ c_parser_omp_clause_to (c_parser *parser, tree list)
static tree
c_parser_omp_clause_from (c_parser *parser, tree list)
{
- return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_FROM, list, true);
+ return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_FROM, list);
}
/* OpenMP 4.0:
@@ -673,6 +673,7 @@ extern int in_alignof;
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;
@@ -75,6 +75,9 @@ int in_typeof;
/* 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;
@@ -2020,6 +2023,13 @@ mark_exp_read (tree exp)
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;
}
@@ -13233,7 +13243,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
t = t_insp.get_deref_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),
@@ -14901,6 +14913,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (!DECL_P (rt))
break;
+ if (!t_insp.map_supported_p ())
+ {
+ sorry_at (OMP_CLAUSE_LOCATION (c),
+ "unsupported map expression %qE",
+ OMP_CLAUSE_DECL (c));
+ remove = true;
+ break;
+ }
+
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IMPLICIT (c)
&& (bitmap_bit_p (&map_head, DECL_UID (rt))
@@ -14951,6 +14972,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
if (t == error_mark_node)
{
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "unmappable expression in %qs clause",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
break;
}
@@ -14979,6 +15003,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else
t = t_insp.get_deref_origin ();
+ if (!t_insp.map_supported_p ())
+ {
+ sorry_at (OMP_CLAUSE_LOCATION (c),
+ "unsupported map expression %qE",
+ OMP_CLAUSE_DECL (c));
+ remove = true;
+ break;
+ }
+
if (t == error_mark_node)
{
remove = true;
@@ -14999,6 +15032,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
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)]);
@@ -39,7 +39,8 @@ foo (int g[3][10], int h[4][8], int i[2][10], int j[][9],
;
#pragma omp target map(alloc: s1) /* { dg-error "'s1' does not have a mappable type in 'map' clause" } */
;
- #pragma omp target map(alloc: s2) /* { dg-error "'s2' does not have a mappable type in 'map' clause" } */
+ #pragma omp target map(alloc: s2) /* { dg-error "'s2' does not have a mappable type in 'map' clause" "" { target c++ } } */
+ /* { dg-error "unmappable expression in 'map' clause" "" { target c } .-1 } */
;
#pragma omp target map(to: a[:][:]) /* { dg-error "array type length expression must be specified" } */
bar (&a[0][0]); /* { dg-error "referenced in target region does not have a mappable type" } */
@@ -22,11 +22,13 @@ foo (void)
#pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */
/* { dg-error "'close' was not declared in this scope" "" { target c++ } .-1 } */
/* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */
+ /* { dg-error "unmappable expression in 'map' clause" "" { target c } .-3 } */
;
#pragma omp target map (always a) /* { dg-error "'always' undeclared" "" { target c } } */
/* { dg-error "'always' was not declared in this scope" "" { target c++ } .-1 } */
/* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */
+ /* { dg-error "unmappable expression in 'map' clause" "" { target c } .-3 } */
;
#pragma omp target map (close to:a)
new file mode 100644
@@ -0,0 +1,50 @@
+// { 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;
+}
new file mode 100644
@@ -0,0 +1,16 @@
+#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;
+}