@@ -12970,6 +12970,15 @@ c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list)
The optional ALLOW_DEREF argument is true if list items can use the deref
(->) operator. */
+struct omp_dim
+{
+ tree low_bound, length;
+ location_t loc;
+ bool no_colon;
+ omp_dim (tree lb, tree len, location_t lo, bool nc)
+ : low_bound (lb), length (len), loc (lo), no_colon (nc) {}
+};
+
static tree
c_parser_omp_variable_list (c_parser *parser,
location_t clause_loc,
@@ -12982,6 +12991,7 @@ c_parser_omp_variable_list (c_parser *parser,
while (1)
{
+ auto_vec<omp_dim> dims;
bool array_section_p = false;
if (kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY)
{
@@ -13101,6 +13111,7 @@ c_parser_omp_variable_list (c_parser *parser,
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
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)))
@@ -13128,9 +13139,13 @@ c_parser_omp_variable_list (c_parser *parser,
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_IN_REDUCTION:
case OMP_CLAUSE_TASK_REDUCTION:
+ array_section_p = false;
+ dims.truncate (0);
while (c_parser_next_token_is (parser, CPP_OPEN_SQUARE))
{
+ location_t loc = UNKNOWN_LOCATION;
tree low_bound = NULL_TREE, length = NULL_TREE;
+ bool no_colon = false;
c_parser_consume_token (parser);
if (!c_parser_next_token_is (parser, CPP_COLON))
@@ -13141,9 +13156,13 @@ c_parser_omp_variable_list (c_parser *parser,
expr = convert_lvalue_to_rvalue (expr_loc, expr,
false, true);
low_bound = expr.value;
+ loc = expr_loc;
}
if (c_parser_next_token_is (parser, CPP_CLOSE_SQUARE))
- length = integer_one_node;
+ {
+ length = integer_one_node;
+ no_colon = true;
+ }
else
{
/* Look for `:'. */
@@ -13172,8 +13191,35 @@ c_parser_omp_variable_list (c_parser *parser,
break;
}
- t = tree_cons (low_bound, length, t);
+ dims.safe_push (omp_dim (low_bound, length, loc, no_colon));
+ }
+
+ if (t != error_mark_node)
+ {
+ if ((kind == OMP_CLAUSE_MAP
+ || kind == OMP_CLAUSE_FROM
+ || 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))))
+ {
+ for (unsigned i = 0; i < dims.length (); i++)
+ {
+ gcc_assert (dims[i].length == integer_one_node);
+ t = build_array_ref (dims[i].loc,
+ t, dims[i].low_bound);
+ }
+ goto start_component_ref;
+ }
+ else
+ {
+ for (unsigned i = 0; i < dims.length (); i++)
+ t = tree_cons (dims[i].low_bound, dims[i].length, t);
+ }
}
+
if ((kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY)
&& t != error_mark_node
&& parser->tokens_avail != 2)
@@ -16420,7 +16466,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);
+ return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO, list, true);
}
/* OpenMP 4.0:
@@ -16429,7 +16475,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);
+ return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_FROM, list, true);
}
/* OpenMP 4.0:
@@ -13217,6 +13217,18 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
+ while (TREE_CODE (t) == INDIRECT_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ STRIP_NOPS (t);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
+ }
+ while (TREE_CODE (t) == COMPOUND_EXPR)
+ {
+ t = TREE_OPERAND (t, 1);
+ STRIP_NOPS (t);
+ }
if (TREE_CODE (t) == COMPONENT_REF
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
@@ -13238,10 +13250,14 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
- if (TREE_CODE (t) == MEM_REF)
+ while (TREE_CODE (t) == MEM_REF
+ || TREE_CODE (t) == INDIRECT_REF
+ || TREE_CODE (t) == ARRAY_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
}
if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
{
@@ -13530,15 +13546,25 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
/* If there is a pointer type anywhere but in the very first
- array-section-subscript, the array section can't be contiguous. */
+ 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)
{
- error_at (OMP_CLAUSE_LOCATION (c),
- "array section is not contiguous in %qs clause",
- omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
- return error_mark_node;
+ /* If 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))
+ {
+ tree d_length = TREE_VALUE (d);
+ if (d_length == NULL_TREE || !integer_onep (d_length))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "array section is not contiguous in %qs clause",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ return error_mark_node;
+ }
+ }
}
}
else
@@ -14888,13 +14914,20 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
{
- while (TREE_CODE (t) == COMPONENT_REF)
- t = TREE_OPERAND (t, 0);
- if (TREE_CODE (t) == MEM_REF)
+ do
{
t = TREE_OPERAND (t, 0);
- STRIP_NOPS (t);
+ if (TREE_CODE (t) == MEM_REF
+ || TREE_CODE (t) == INDIRECT_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ STRIP_NOPS (t);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
+ }
}
+ while (TREE_CODE (t) == COMPONENT_REF);
+
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IMPLICIT (c)
&& (bitmap_bit_p (&map_head, DECL_UID (t))
@@ -14961,14 +14994,32 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bias) to zero here, so it is not set erroneously to the pointer
size later on in gimplify.c. */
OMP_CLAUSE_SIZE (c) = size_zero_node;
+ while (TREE_CODE (t) == INDIRECT_REF
+ || TREE_CODE (t) == ARRAY_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ STRIP_NOPS (t);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
+ }
+ while (TREE_CODE (t) == COMPOUND_EXPR)
+ {
+ t = TREE_OPERAND (t, 1);
+ STRIP_NOPS (t);
+ }
indir_component_ref_p = false;
if (TREE_CODE (t) == COMPONENT_REF
- && TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF)
+ && (TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF
+ || TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF
+ || TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF))
{
t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
indir_component_ref_p = true;
STRIP_NOPS (t);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
}
+
if (TREE_CODE (t) == COMPONENT_REF
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
{
@@ -15004,7 +15055,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
}
t = TREE_OPERAND (t, 0);
- if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
+ if (TREE_CODE (t) == MEM_REF)
{
if (maybe_ne (mem_ref_offset (t), 0))
error_at (OMP_CLAUSE_LOCATION (c),
@@ -15013,6 +15064,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else
t = TREE_OPERAND (t, 0);
}
+ while (TREE_CODE (t) == MEM_REF
+ || TREE_CODE (t) == INDIRECT_REF
+ || TREE_CODE (t) == ARRAY_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ STRIP_NOPS (t);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
+ }
}
if (remove)
break;
@@ -15084,7 +15144,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
"%qD appears more than once in data clauses", t);
remove = true;
}
- else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ else if (bitmap_bit_p (&map_head, DECL_UID (t))
+ && !bitmap_bit_p (&map_field_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
error_at (OMP_CLAUSE_LOCATION (c),
@@ -36323,11 +36323,22 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code,
The optional ALLOW_DEREF argument is true if list items can use the deref
(->) operator. */
+struct omp_dim
+{
+ tree low_bound, length;
+ location_t loc;
+ bool no_colon;
+ omp_dim (tree lb, tree len, location_t lo, bool nc)
+ : low_bound (lb), length (len), loc (lo), no_colon (nc) {}
+};
+
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)
{
+ auto_vec<omp_dim> dims;
+ bool array_section_p;
cp_token *token;
bool saved_colon_corrects_to_scope_p = parser->colon_corrects_to_scope_p;
if (colon)
@@ -36408,6 +36419,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
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)))
@@ -36431,14 +36443,19 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_IN_REDUCTION:
case OMP_CLAUSE_TASK_REDUCTION:
+ array_section_p = false;
+ dims.truncate (0);
while (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE))
{
+ location_t loc = UNKNOWN_LOCATION;
tree low_bound = NULL_TREE, length = NULL_TREE;
+ bool no_colon = false;
parser->colon_corrects_to_scope_p = false;
cp_lexer_consume_token (parser->lexer);
if (!cp_lexer_next_token_is (parser->lexer, CPP_COLON))
{
+ loc = cp_lexer_peek_token (parser->lexer)->location;
low_bound = cp_parser_expression (parser);
/* Later handling is not prepared to see through these. */
gcc_checking_assert (!location_wrapper_p (low_bound));
@@ -36447,7 +36464,10 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
parser->colon_corrects_to_scope_p
= saved_colon_corrects_to_scope_p;
if (cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_SQUARE))
- length = integer_one_node;
+ {
+ length = integer_one_node;
+ no_colon = true;
+ }
else
{
/* Look for `:'. */
@@ -36460,6 +36480,8 @@ 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_commit_to_tentative_parse (parser);
+ else
+ array_section_p = true;
if (!cp_lexer_next_token_is (parser->lexer,
CPP_CLOSE_SQUARE))
{
@@ -36478,8 +36500,32 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
goto skip_comma;
}
- decl = tree_cons (low_bound, length, decl);
+ dims.safe_push (omp_dim (low_bound, length, loc, no_colon));
+ }
+
+ if ((kind == OMP_CLAUSE_MAP
+ || kind == OMP_CLAUSE_FROM
+ || 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))))
+ {
+ for (unsigned i = 0; i < dims.length (); i++)
+ {
+ gcc_assert (dims[i].length == integer_one_node);
+ decl = build_array_ref (dims[i].loc,
+ decl, dims[i].low_bound);
+ }
+ goto start_component_ref;
}
+ else
+ {
+ for (unsigned i = 0; i < dims.length (); i++)
+ decl = tree_cons (dims[i].low_bound, dims[i].length, decl);
+ }
+
break;
default:
break;
@@ -39981,11 +40027,13 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO_DECLARE,
clauses);
else
- clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, clauses);
+ clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, clauses,
+ true);
c_name = "to";
break;
case PRAGMA_OMP_CLAUSE_FROM:
- clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, clauses);
+ clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, clauses,
+ true);
c_name = "from";
break;
case PRAGMA_OMP_CLAUSE_UNIFORM:
@@ -5030,6 +5030,18 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
t = TREE_OPERAND (t, 0);
ret = t;
+ while (TREE_CODE (t) == INDIRECT_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ STRIP_NOPS (t);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
+ }
+ while (TREE_CODE (t) == COMPOUND_EXPR)
+ {
+ t = TREE_OPERAND (t, 1);
+ STRIP_NOPS (t);
+ }
if (TREE_CODE (t) == COMPONENT_REF
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
@@ -5054,10 +5066,14 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
- if (TREE_CODE (t) == INDIRECT_REF)
+ while (TREE_CODE (t) == MEM_REF
+ || TREE_CODE (t) == INDIRECT_REF
+ || TREE_CODE (t) == ARRAY_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
}
}
if (REFERENCE_REF_P (t))
@@ -5341,15 +5357,25 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
/* If there is a pointer type anywhere but in the very first
- array-section-subscript, the array section can't be contiguous. */
+ 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)
{
- error_at (OMP_CLAUSE_LOCATION (c),
- "array section is not contiguous in %qs clause",
- omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
- return error_mark_node;
+ /* If 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))
+ {
+ tree d_length = TREE_VALUE (d);
+ if (d_length == NULL_TREE || !integer_onep (d_length))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "array section is not contiguous in %qs clause",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ return error_mark_node;
+ }
+ }
}
}
else
@@ -5619,16 +5645,37 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
default:
break;
}
+ bool reference_always_pointer = true;
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
if (TREE_CODE (t) == COMPONENT_REF)
- OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
+ {
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
+
+ if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
+ && TYPE_REF_P (TREE_TYPE (t)))
+ {
+ if (TREE_CODE (TREE_TYPE (TREE_TYPE (t))) == ARRAY_TYPE)
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ else
+ t = convert_from_reference (t);
+
+ reference_always_pointer = false;
+ }
+ }
else if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
{
- t = TREE_OPERAND (t, 0);
- gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
- : GOMP_MAP_ALWAYS_POINTER;
+ gomp_map_kind k;
+ if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
+ && TREE_CODE (TREE_TYPE (t)) == POINTER_TYPE)
+ k = GOMP_MAP_ATTACH_DETACH;
+ else
+ {
+ t = TREE_OPERAND (t, 0);
+ k = (ort == C_ORT_ACC
+ ? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER);
+ }
OMP_CLAUSE_SET_MAP_KIND (c2, k);
}
else
@@ -5652,8 +5699,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
OMP_CLAUSE_SIZE (c2) = t;
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = c2;
+
ptr = OMP_CLAUSE_DECL (c2);
- if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ if (reference_always_pointer
+ && OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& TYPE_REF_P (TREE_TYPE (ptr))
&& INDIRECT_TYPE_P (TREE_TYPE (TREE_TYPE (ptr))))
{
@@ -7855,15 +7904,22 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
{
- while (TREE_CODE (t) == COMPONENT_REF)
- t = TREE_OPERAND (t, 0);
- if (REFERENCE_REF_P (t))
- t = TREE_OPERAND (t, 0);
- if (TREE_CODE (t) == INDIRECT_REF)
+ do
{
t = TREE_OPERAND (t, 0);
- STRIP_NOPS (t);
+ if (REFERENCE_REF_P (t))
+ t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == MEM_REF
+ || TREE_CODE (t) == INDIRECT_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ STRIP_NOPS (t);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
+ }
}
+ while (TREE_CODE (t) == COMPONENT_REF);
+
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IMPLICIT (c)
&& (bitmap_bit_p (&map_head, DECL_UID (t))
@@ -7934,15 +7990,33 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
{
t = TREE_OPERAND (t, 0);
- OMP_CLAUSE_DECL (c) = t;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH)
+ OMP_CLAUSE_DECL (c) = t;
+ }
+ while (TREE_CODE (t) == INDIRECT_REF
+ || TREE_CODE (t) == ARRAY_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ STRIP_NOPS (t);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
+ }
+ while (TREE_CODE (t) == COMPOUND_EXPR)
+ {
+ t = TREE_OPERAND (t, 1);
+ STRIP_NOPS (t);
}
indir_component_ref_p = false;
if (TREE_CODE (t) == COMPONENT_REF
- && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
+ && (TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF
+ || TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF))
{
t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
indir_component_ref_p = true;
STRIP_NOPS (t);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
}
if (TREE_CODE (t) == COMPONENT_REF
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
@@ -7977,6 +8051,24 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
}
t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == MEM_REF)
+ {
+ if (maybe_ne (mem_ref_offset (t), 0))
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "cannot dereference %qE in %qs clause", t,
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ else
+ t = TREE_OPERAND (t, 0);
+ }
+ while (TREE_CODE (t) == MEM_REF
+ || TREE_CODE (t) == INDIRECT_REF
+ || TREE_CODE (t) == ARRAY_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ STRIP_NOPS (t);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
+ }
}
if (remove)
break;
@@ -8074,7 +8166,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
"%qD appears more than once in data clauses", t);
remove = true;
}
- else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ else if (bitmap_bit_p (&map_head, DECL_UID (t))
+ && !bitmap_bit_p (&map_field_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
error_at (OMP_CLAUSE_LOCATION (c),
@@ -8121,8 +8214,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else
{
bitmap_set_bit (&map_head, DECL_UID (t));
- if (t != OMP_CLAUSE_DECL (c)
- && TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
+
+ tree decl = OMP_CLAUSE_DECL (c);
+ if (t != decl
+ && (TREE_CODE (decl) == COMPONENT_REF
+ || (INDIRECT_REF_P (decl)
+ && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+ && TYPE_REF_P (TREE_TYPE (TREE_OPERAND (decl, 0))))))
bitmap_set_bit (&map_field_head, DECL_UID (t));
}
handle_map_references:
@@ -8151,7 +8249,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
if (TREE_CODE (t) == COMPONENT_REF)
- OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
else
OMP_CLAUSE_SET_MAP_KIND (c2,
GOMP_MAP_FIRSTPRIVATE_REFERENCE);
@@ -2460,6 +2460,9 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
TREE_TYPE (TREE_TYPE (decl)),
decl, offset, NULL_TREE, NULL_TREE);
OMP_CLAUSE_DECL (node) = offset;
+
+ if (ptr_kind == GOMP_MAP_ALWAYS_POINTER)
+ return;
}
else
{
@@ -8666,7 +8666,7 @@ insert_struct_comp_map (enum tree_code code, tree c, tree struct_node,
static tree
extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
- poly_offset_int *poffsetp)
+ poly_offset_int *poffsetp, tree *offsetp)
{
tree offset;
poly_int64 bitsize, bitpos;
@@ -8713,10 +8713,11 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
&& TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == REFERENCE_TYPE)
base = TREE_OPERAND (base, 0);
- gcc_assert (offset == NULL_TREE || poly_int_tree_p (offset));
-
- if (offset)
- poffset = wi::to_poly_offset (offset);
+ if (offset && poly_int_tree_p (offset))
+ {
+ poffset = wi::to_poly_offset (offset);
+ offset = NULL_TREE;
+ }
else
poffset = 0;
@@ -8725,6 +8726,7 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
*bitposp = bitpos;
*poffsetp = poffset;
+ *offsetp = offset;
/* Set *BASE_REF if BASE was a dereferenced reference variable. */
if (base_ref && orig_base != base)
@@ -8738,12 +8740,22 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
static bool
is_or_contains_p (tree expr, tree base_ptr)
{
- while (expr != base_ptr)
- if (TREE_CODE (base_ptr) == COMPONENT_REF)
- base_ptr = TREE_OPERAND (base_ptr, 0);
- else
- break;
- return expr == base_ptr;
+ if ((TREE_CODE (expr) == INDIRECT_REF && TREE_CODE (base_ptr) == MEM_REF)
+ || (TREE_CODE (expr) == MEM_REF && TREE_CODE (base_ptr) == INDIRECT_REF))
+ return operand_equal_p (TREE_OPERAND (expr, 0),
+ TREE_OPERAND (base_ptr, 0));
+ while (!operand_equal_p (expr, base_ptr))
+ {
+ if (TREE_CODE (base_ptr) == COMPOUND_EXPR)
+ base_ptr = TREE_OPERAND (base_ptr, 1);
+ if (TREE_CODE (base_ptr) == COMPONENT_REF
+ || TREE_CODE (base_ptr) == POINTER_PLUS_EXPR
+ || TREE_CODE (base_ptr) == SAVE_EXPR)
+ base_ptr = TREE_OPERAND (base_ptr, 0);
+ else
+ break;
+ }
+ return operand_equal_p (expr, base_ptr);
}
/* Implement OpenMP 5.x map ordering rules for target directives. There are
@@ -8823,21 +8835,107 @@ omp_target_reorder_clauses (tree *list_p)
tree base_ptr = TREE_OPERAND (decl, 0);
STRIP_TYPE_NOPS (base_ptr);
for (unsigned int j = i + 1; j < atf.length (); j++)
- {
- tree *cp2 = atf[j];
- tree decl2 = OMP_CLAUSE_DECL (*cp2);
- if (is_or_contains_p (decl2, base_ptr))
- {
- /* Move *cp2 to before *cp. */
- tree c = *cp2;
- *cp2 = OMP_CLAUSE_CHAIN (c);
- OMP_CLAUSE_CHAIN (c) = *cp;
- *cp = c;
- atf[j] = NULL;
+ if (atf[j])
+ {
+ tree *cp2 = atf[j];
+ tree decl2 = OMP_CLAUSE_DECL (*cp2);
+
+ decl2 = OMP_CLAUSE_DECL (*cp2);
+ if (is_or_contains_p (decl2, base_ptr))
+ {
+ /* Move *cp2 to before *cp. */
+ tree c = *cp2;
+ *cp2 = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = *cp;
+ *cp = c;
+
+ if (*cp2 != NULL_TREE
+ && OMP_CLAUSE_CODE (*cp2) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (*cp2) == GOMP_MAP_ALWAYS_POINTER)
+ {
+ tree c2 = *cp2;
+ *cp2 = OMP_CLAUSE_CHAIN (c2);
+ OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = c2;
+ }
+
+ atf[j] = NULL;
}
- }
+ }
}
}
+
+ /* For attach_detach map clauses, if there is another map that maps the
+ attached/detached pointer, make sure that map is ordered before the
+ attach_detach. */
+ atf.truncate (0);
+ for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
+ if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
+ {
+ /* Collect alloc, to, from, to/from clauses, and
+ always_pointer/attach_detach clauses. */
+ gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
+ if (k == GOMP_MAP_ALLOC
+ || k == GOMP_MAP_TO
+ || k == GOMP_MAP_FROM
+ || k == GOMP_MAP_TOFROM
+ || k == GOMP_MAP_ALWAYS_TO
+ || k == GOMP_MAP_ALWAYS_FROM
+ || k == GOMP_MAP_ALWAYS_TOFROM
+ || k == GOMP_MAP_ATTACH_DETACH
+ || k == GOMP_MAP_ALWAYS_POINTER)
+ atf.safe_push (cp);
+ }
+
+ for (unsigned int i = 0; i < atf.length (); i++)
+ if (atf[i])
+ {
+ tree *cp = atf[i];
+ tree ptr = OMP_CLAUSE_DECL (*cp);
+ STRIP_TYPE_NOPS (ptr);
+ if (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH)
+ for (unsigned int j = i + 1; j < atf.length (); j++)
+ {
+ tree *cp2 = atf[j];
+ tree decl2 = OMP_CLAUSE_DECL (*cp2);
+ if (OMP_CLAUSE_MAP_KIND (*cp2) != GOMP_MAP_ATTACH_DETACH
+ && OMP_CLAUSE_MAP_KIND (*cp2) != GOMP_MAP_ALWAYS_POINTER
+ && is_or_contains_p (decl2, ptr))
+ {
+ /* Move *cp2 to before *cp. */
+ tree c = *cp2;
+ *cp2 = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = *cp;
+ *cp = c;
+ atf[j] = NULL;
+
+ /* If decl2 is of the form '*decl2_opnd0', and followed by an
+ ALWAYS_POINTER or ATTACH_DETACH of 'decl2_opnd0', move the
+ pointer operation along with *cp2. This can happen for C++
+ reference sequences. */
+ if (j + 1 < atf.length ()
+ && (TREE_CODE (decl2) == INDIRECT_REF
+ || TREE_CODE (decl2) == MEM_REF))
+ {
+ tree *cp3 = atf[j + 1];
+ tree decl3 = OMP_CLAUSE_DECL (*cp3);
+ tree decl2_opnd0 = TREE_OPERAND (decl2, 0);
+ if ((OMP_CLAUSE_MAP_KIND (*cp3) == GOMP_MAP_ALWAYS_POINTER
+ || OMP_CLAUSE_MAP_KIND (*cp3) == GOMP_MAP_ATTACH_DETACH)
+ && operand_equal_p (decl3, decl2_opnd0))
+ {
+ /* Also move *cp3 to before *cp. */
+ c = *cp3;
+ *cp2 = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = *cp;
+ *cp = c;
+ atf[j + 1] = NULL;
+ j += 1;
+ }
+ }
+ }
+ }
+ }
}
/* DECL is supposed to have lastprivate semantics in the outer contexts
@@ -8929,6 +9027,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
struct gimplify_omp_ctx *ctx, *outer_ctx;
tree c;
hash_map<tree_operand_hash, tree> *struct_map_to_clause = NULL;
+ hash_map<tree_operand_hash, tree *> *struct_seen_clause = NULL;
hash_set<tree> *struct_deref_set = NULL;
tree *prev_list_p = NULL, *orig_list_p = list_p;
int handled_depend_iterators = -1;
@@ -9404,6 +9503,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
bool indir_p = false;
bool component_ref_p = false;
+ tree indir_base = NULL_TREE;
tree orig_decl = decl;
tree decl_ref = NULL_TREE;
if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
@@ -9422,6 +9522,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
== POINTER_TYPE))
{
indir_p = true;
+ indir_base = decl;
decl = TREE_OPERAND (decl, 0);
STRIP_NOPS (decl);
}
@@ -9468,7 +9569,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
!= GOMP_MAP_POINTER)
|| OMP_CLAUSE_DECL (next_clause) != decl)
&& (!struct_deref_set
- || !struct_deref_set->contains (decl)))
+ || !struct_deref_set->contains (decl))
+ && (!struct_map_to_clause
+ || !struct_map_to_clause->get (indir_base)))
{
if (!struct_deref_set)
struct_deref_set = new hash_set<tree> ();
@@ -9512,7 +9615,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
if ((DECL_P (decl)
|| (component_ref_p
&& (INDIRECT_REF_P (decl)
- || TREE_CODE (decl) == MEM_REF)))
+ || TREE_CODE (decl) == MEM_REF
+ || TREE_CODE (decl) == ARRAY_REF)))
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
@@ -9547,7 +9651,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
remove = true;
break;
}
- if (OMP_CLAUSE_CHAIN (*prev_list_p) != c)
+
+ /* The below prev_list_p based error recovery code is
+ currently no longer valid for OpenMP. */
+ if (code != OMP_TARGET
+ && code != OMP_TARGET_DATA
+ && code != OMP_TARGET_UPDATE
+ && code != OMP_TARGET_ENTER_DATA
+ && code != OMP_TARGET_EXIT_DATA
+ && OMP_CLAUSE_CHAIN (*prev_list_p) != c)
{
tree ch = OMP_CLAUSE_CHAIN (*prev_list_p);
if (ch == NULL_TREE || OMP_CLAUSE_CHAIN (ch) != c)
@@ -9560,13 +9672,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
poly_offset_int offset1;
poly_int64 bitpos1;
+ tree tree_offset1;
tree base_ref;
tree base
= extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ref,
- &bitpos1, &offset1);
+ &bitpos1, &offset1,
+ &tree_offset1);
- gcc_assert (base == decl);
+ bool do_map_struct = (base == decl && !tree_offset1);
splay_tree_node n
= (DECL_P (decl)
@@ -9598,6 +9712,32 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
OMP_CLAUSE_SET_MAP_KIND (c, k);
has_attachments = true;
}
+
+ /* We currently don't handle non-constant offset accesses wrt to
+ GOMP_MAP_STRUCT elements. */
+ if (!do_map_struct)
+ goto skip_map_struct;
+
+ /* Nor for attach_detach for OpenMP. */
+ if ((code == OMP_TARGET
+ || code == OMP_TARGET_DATA
+ || code == OMP_TARGET_UPDATE
+ || code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA)
+ && attach_detach)
+ {
+ if (DECL_P (decl))
+ {
+ if (struct_seen_clause == NULL)
+ struct_seen_clause
+ = new hash_map<tree_operand_hash, tree *>;
+ if (!struct_seen_clause->get (decl))
+ struct_seen_clause->put (decl, list_p);
+ }
+
+ goto skip_map_struct;
+ }
+
if ((DECL_P (decl)
&& (n == NULL || (n->value & GOVD_MAP) == 0))
|| (!DECL_P (decl)
@@ -9637,9 +9777,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
struct_map_to_clause->put (decl, l);
if (ptr || attach_detach)
{
- insert_struct_comp_map (code, c, l, *prev_list_p,
+ tree **sc = (struct_seen_clause
+ ? struct_seen_clause->get (decl)
+ : NULL);
+ tree *insert_node_pos = sc ? *sc : prev_list_p;
+
+ insert_struct_comp_map (code, c, l, *insert_node_pos,
NULL);
- *prev_list_p = l;
+ *insert_node_pos = l;
prev_list_p = NULL;
}
else
@@ -9725,9 +9870,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
tree sc_decl = OMP_CLAUSE_DECL (*sc);
poly_offset_int offsetn;
poly_int64 bitposn;
+ tree tree_offsetn;
tree base
= extract_base_bit_offset (sc_decl, NULL,
- &bitposn, &offsetn);
+ &bitposn, &offsetn,
+ &tree_offsetn);
if (base != decl)
break;
if (scp)
@@ -9815,16 +9962,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
continue;
}
}
+ skip_map_struct:
+ ;
}
else if ((code == OACC_ENTER_DATA
|| code == OACC_EXIT_DATA
|| code == OACC_DATA
|| code == OACC_PARALLEL
|| code == OACC_KERNELS
- || code == OACC_SERIAL)
+ || code == OACC_SERIAL
+ || code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA)
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
{
- gomp_map_kind k = (code == OACC_EXIT_DATA
+ gomp_map_kind k = ((code == OACC_EXIT_DATA
+ || code == OMP_TARGET_EXIT_DATA)
? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
OMP_CLAUSE_SET_MAP_KIND (c, k);
}
@@ -10656,6 +10808,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
ctx->clauses = *orig_list_p;
gimplify_omp_ctxp = ctx;
+ if (struct_seen_clause)
+ delete struct_seen_clause;
if (struct_map_to_clause)
delete struct_map_to_clause;
if (struct_deref_set)
@@ -37,13 +37,12 @@ int main(int argc, char* argv[])
{
int j, k;
for (k = 0; k < S; k++)
-#pragma acc parallel loop copy(m[k].a[0:N]) /* { dg-error "expected .\\\). before .\\\.. token" } */
+#pragma acc parallel loop copy(m[k].a[0:N])
for (j = 0; j < N; j++)
m[k].a[j]++;
for (k = 0; k < S; k++)
-#pragma acc parallel loop copy(m[k].b[0:N], m[k].c[5:N-10]) /* { dg-error "expected .\\\). before .\\\.. token" } */
- /* { dg-error ".m. appears more than once in data clauses" "" { target c++ } .-1 } */
+#pragma acc parallel loop copy(m[k].b[0:N], m[k].c[5:N-10])
for (j = 0; j < N; j++)
{
m[k].b[j]++;
new file mode 100644
@@ -0,0 +1,24 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fopenmp -fdump-tree-gimple" } */
+
+struct bar
+{
+ int num_vectors;
+ double *vectors;
+};
+
+struct foo
+{
+ int num_vectors;
+ struct bar *bars;
+ double **vectors;
+};
+
+void func (struct foo *f, int n, int m)
+{
+ #pragma omp target enter data map (to: f->vectors[m][:n])
+ #pragma omp target enter data map (to: f->bars[n].vectors[:m])
+ #pragma omp target enter data map (to: f->bars[n].vectors[:f->bars[n].num_vectors])
+}
+
+/* { dg-final { scan-tree-dump-times "map\\(to:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:\[^-\]+->vectors \\\[bias: \[^\]\]+\\\]\\)" 3 "gimple" } } */
new file mode 100644
@@ -0,0 +1,52 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+#include <stdlib.h>
+
+#define N 10
+
+struct S
+{
+ int a, b;
+ int *ptr;
+ int c, d;
+};
+
+int
+main (void)
+{
+ struct S a;
+ a.ptr = (int *) malloc (sizeof (int) * N);
+
+ for (int i = 0; i < N; i++)
+ a.ptr[i] = 0;
+
+ #pragma omp target enter data map(to: a.ptr, a.ptr[:N])
+
+ #pragma omp target
+ for (int i = 0; i < N; i++)
+ a.ptr[i] += 1;
+
+ #pragma omp target update from(a.ptr[:N])
+
+ for (int i = 0; i < N; i++)
+ if (a.ptr[i] != 1)
+ abort ();
+
+ #pragma omp target map(a.ptr[:N])
+ for (int i = 0; i < N; i++)
+ a.ptr[i] += 1;
+
+ #pragma omp target update from(a.ptr[:N])
+
+ for (int i = 0; i < N; i++)
+ if (a.ptr[i] != 2)
+ abort ();
+
+ #pragma omp target exit data map(from:a.ptr, a.ptr[:N])
+
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:a \[len: [0-9]+\]\[implicit\]\)} "gimple" } } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:a \[len: [0-9]+\]\[implicit\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(attach:a\.ptr \[bias: 0\]\)} "gimple" } } */
@@ -581,11 +581,30 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
address/length adjustment is a TODO. */
assert (!implicit_subset);
- gomp_copy_host2dev (devicep, aq,
- (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
- + newn->host_start - oldn->host_start),
- (void *) newn->host_start,
- newn->host_end - newn->host_start, false, cbuf);
+ if (oldn->aux && oldn->aux->attach_count)
+ {
+ /* We have to be careful not to overwrite still attached pointers
+ during the copyback to host. */
+ uintptr_t addr = newn->host_start;
+ while (addr < newn->host_end)
+ {
+ size_t i = (addr - oldn->host_start) / sizeof (void *);
+ if (oldn->aux->attach_count[i] == 0)
+ gomp_copy_host2dev (devicep, aq,
+ (void *) (oldn->tgt->tgt_start
+ + oldn->tgt_offset
+ + addr - oldn->host_start),
+ (void *) addr,
+ sizeof (void *), false, cbuf);
+ addr += sizeof (void *);
+ }
+ }
+ else
+ gomp_copy_host2dev (devicep, aq,
+ (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
+ + newn->host_start - oldn->host_start),
+ (void *) newn->host_start,
+ newn->host_end - newn->host_start, false, cbuf);
}
gomp_increment_refcount (oldn, refcount_set);
@@ -2011,17 +2030,45 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
(void *) n->host_end);
}
-
- void *hostaddr = (void *) cur_node.host_start;
- void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset
- + cur_node.host_start - n->host_start);
- size_t size = cur_node.host_end - cur_node.host_start;
-
- if (GOMP_MAP_COPY_TO_P (kind & typemask))
- gomp_copy_host2dev (devicep, NULL, devaddr, hostaddr, size,
- false, NULL);
- if (GOMP_MAP_COPY_FROM_P (kind & typemask))
- gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size);
+ if (n->aux && n->aux->attach_count)
+ {
+ uintptr_t addr = cur_node.host_start;
+ while (addr < cur_node.host_end)
+ {
+ /* We have to be careful not to overwrite still attached
+ pointers during host<->device updates. */
+ size_t i = (addr - cur_node.host_start) / sizeof (void *);
+ if (n->aux->attach_count[i] == 0)
+ {
+ void *devaddr = (void *) (n->tgt->tgt_start
+ + n->tgt_offset
+ + addr - n->host_start);
+ if (GOMP_MAP_COPY_TO_P (kind & typemask))
+ gomp_copy_host2dev (devicep, NULL,
+ devaddr, (void *) addr,
+ sizeof (void *), false, NULL);
+ if (GOMP_MAP_COPY_FROM_P (kind & typemask))
+ gomp_copy_dev2host (devicep, NULL,
+ (void *) addr, devaddr,
+ sizeof (void *));
+ }
+ addr += sizeof (void *);
+ }
+ }
+ else
+ {
+ void *hostaddr = (void *) cur_node.host_start;
+ void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset
+ + cur_node.host_start
+ - n->host_start);
+ size_t size = cur_node.host_end - cur_node.host_start;
+
+ if (GOMP_MAP_COPY_TO_P (kind & typemask))
+ gomp_copy_host2dev (devicep, NULL, devaddr, hostaddr, size,
+ false, NULL);
+ if (GOMP_MAP_COPY_FROM_P (kind & typemask))
+ gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size);
+ }
}
}
gomp_mutex_unlock (&devicep->lock);
@@ -2915,11 +2962,31 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
if ((kind == GOMP_MAP_FROM && do_copy)
|| kind == GOMP_MAP_ALWAYS_FROM)
- gomp_copy_dev2host (devicep, NULL, (void *) cur_node.host_start,
- (void *) (k->tgt->tgt_start + k->tgt_offset
- + cur_node.host_start
- - k->host_start),
- cur_node.host_end - cur_node.host_start);
+ {
+ if (k->aux && k->aux->attach_count)
+ {
+ /* We have to be careful not to overwrite still attached
+ pointers during the copyback to host. */
+ uintptr_t addr = k->host_start;
+ while (addr < k->host_end)
+ {
+ size_t i = (addr - k->host_start) / sizeof (void *);
+ if (k->aux->attach_count[i] == 0)
+ gomp_copy_dev2host (devicep, NULL, (void *) addr,
+ (void *) (k->tgt->tgt_start
+ + k->tgt_offset
+ + addr - k->host_start),
+ sizeof (void *));
+ addr += sizeof (void *);
+ }
+ }
+ else
+ gomp_copy_dev2host (devicep, NULL, (void *) cur_node.host_start,
+ (void *) (k->tgt->tgt_start + k->tgt_offset
+ + cur_node.host_start
+ - k->host_start),
+ cur_node.host_end - cur_node.host_start);
+ }
/* Structure elements lists are removed altogether at once, which
may cause immediate deallocation of the target_mem_desc, causing
@@ -23,9 +23,11 @@ foo ()
e = c + 18;
D s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e };
int err = 0;
- #pragma omp target map (to:s.v.b[0:z + 7], s.template u[z + 1:z + 4]) \
- map (tofrom:s.s[3:3], s. template v. template d[z + 1:z + 3]) \
- map (from: s.w[z:4], s.x[1:3], err) private (i)
+ #pragma omp target map (to: s.v.b, s.v.b[0:z + 7]) \
+ map (s.template u, s.template u[z + 1:z + 4]) \
+ map (tofrom: s.s, s.s[3:3]) \
+ map (tofrom: s. template v. template d[z + 1:z + 3])\
+ map (from: s.w, s.w[z:4], s.x, s.x[1:3], err) private (i)
{
err = 0;
for (i = 0; i < 7; i++)
@@ -80,9 +82,9 @@ main ()
e = c + 18;
S s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e };
int err = 0;
- #pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \
- map (tofrom:s.s[3:3], s.v.d[z + 1:z + 3]) \
- map (from: s.w[z:4], s.x[1:3], err) private (i)
+ #pragma omp target map (to: s.v.b, s.v.b[0:z + 7], s.u, s.u[z + 1:z + 4]) \
+ map (tofrom: s.s, s.s[3:3], s.v.d[z + 1:z + 3]) \
+ map (from: s.w, s.w[z:4], s.x, s.x[1:3], err) private (i)
{
err = 0;
for (i = 0; i < 7; i++)
@@ -53,7 +53,7 @@ main ()
int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
S s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
int *v = u + 4;
- #pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3])
+ #pragma omp target enter data map (to: s.s, s.u, s.u[0:5]) map (alloc: s.v[1:3])
s.s++;
u[3]++;
s.v[1]++;
@@ -14,7 +14,7 @@ foo (S s)
d = id;
int err;
- #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+ #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err)
{
err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@@ -48,7 +48,7 @@ foo (S s)
|| omp_target_is_present (&s.h, d)
|| omp_target_is_present (&s.h[2], d)))
abort ();
- #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
{
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
@@ -61,8 +61,8 @@ foo (S s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
- #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
- #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+ #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
{
err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
@@ -73,7 +73,7 @@ foo (S s)
s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
}
- #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
}
if (sep
&& (omp_target_is_present (&s.a, d)
@@ -97,7 +97,7 @@ foo (S s)
s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
- #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
@@ -109,8 +109,8 @@ foo (S s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
- #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
- #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+ #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
{
err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
@@ -121,7 +121,7 @@ foo (S s)
s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
}
- #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
@@ -133,7 +133,7 @@ foo (S s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
- #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (sep
&& (omp_target_is_present (&s.a, d)
|| omp_target_is_present (s.b, d)
@@ -16,7 +16,7 @@ foo (S<C, I, L, UC, SH> s)
d = id;
int err;
- #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+ #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err)
{
err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@@ -50,7 +50,7 @@ foo (S<C, I, L, UC, SH> s)
|| omp_target_is_present (&s.h, d)
|| omp_target_is_present (&s.h[2], d)))
abort ();
- #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
{
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
@@ -63,8 +63,8 @@ foo (S<C, I, L, UC, SH> s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
- #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
- #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+ #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
{
err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
@@ -75,7 +75,7 @@ foo (S<C, I, L, UC, SH> s)
s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
}
- #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
}
if (sep
&& (omp_target_is_present (&s.a, d)
@@ -99,7 +99,7 @@ foo (S<C, I, L, UC, SH> s)
s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
- #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
@@ -111,8 +111,8 @@ foo (S<C, I, L, UC, SH> s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
- #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
- #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+ #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
{
err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
@@ -123,7 +123,7 @@ foo (S<C, I, L, UC, SH> s)
s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
}
- #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
@@ -135,7 +135,7 @@ foo (S<C, I, L, UC, SH> s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
- #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (sep
&& (omp_target_is_present (&s.a, d)
|| omp_target_is_present (s.b, d)
@@ -16,7 +16,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
d = id;
int err;
- #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+ #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err)
{
err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@@ -50,7 +50,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
|| omp_target_is_present (&s.h, d)
|| omp_target_is_present (&s.h[2], d)))
abort ();
- #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
{
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
@@ -63,8 +63,8 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
- #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
- #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+ #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
{
err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
@@ -75,7 +75,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
}
- #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
}
if (sep
&& (omp_target_is_present (&s.a, d)
@@ -99,7 +99,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
- #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
@@ -111,8 +111,8 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
- #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
- #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+ #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
{
err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
@@ -123,7 +123,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
}
- #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
@@ -135,7 +135,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
- #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (sep
&& (omp_target_is_present (&s.a, d)
|| omp_target_is_present (s.b, d)
@@ -7,7 +7,7 @@ void
foo (S s)
{
int err;
- #pragma omp target map (s.x[0:N], s.y[0:N]) map (s.t.t[16:3]) map (from: err)
+ #pragma omp target map (s.x[0:N], s.y, s.y[0:N]) map (s.t.t[16:3]) map (from: err)
{
err = s.x[2] != 28 || s.y[2] != 37 || s.t.t[17] != 81;
s.x[2]++;
@@ -38,7 +38,7 @@ void
foo2 (S &s)
{
int err;
- #pragma omp target map (s.x[N:10], s.y[N:10]) map (from: err) map (s.t.t[N+16:N+3])
+ #pragma omp target map (s.x[N:10], s.y, s.y[N:10]) map (from: err) map (s.t.t[N+16:N+3])
{
err = s.x[2] != 30 || s.y[2] != 38 || s.t.t[17] != 81;
s.x[2]++;
@@ -69,7 +69,7 @@ void
foo3 (U s)
{
int err;
- #pragma omp target map (s.x[0:10], s.y[0:10]) map (from: err) map (s.t.t[16:3])
+ #pragma omp target map (s.x[0:10], s.y, s.y[0:10]) map (from: err) map (s.t.t[16:3])
{
err = s.x[2] != 32 || s.y[2] != 39 || s.t.t[17] != 82;
s.x[2]++;
@@ -100,7 +100,7 @@ void
foo4 (U &s)
{
int err;
- #pragma omp target map (s.x[0:10], s.y[0:10]) map (from: err) map (s.t.t[16:3])
+ #pragma omp target map (s.x[0:10], s.y, s.y[0:10]) map (from: err) map (s.t.t[16:3])
{
err = s.x[2] != 34 || s.y[2] != 40 || s.t.t[17] != 82;
s.x[2]++;
@@ -16,13 +16,13 @@ main (void)
s->data[i] = 0;
#pragma omp target enter data map(to: s)
- #pragma omp target enter data map(to: s->data[:SZ])
+ #pragma omp target enter data map(to: s->data, s->data[:SZ])
#pragma omp target
{
for (int i = 0; i < SZ; i++)
s->data[i] = i;
}
- #pragma omp target exit data map(from: s->data[:SZ])
+ #pragma omp target exit data map(from: s->data, s->data[:SZ])
#pragma omp target exit data map(from: s)
for (int i = 0; i < SZ; i++)
new file mode 100644
@@ -0,0 +1,46 @@
+#include <stdlib.h>
+
+#define N 10
+
+struct S
+{
+ int a, b;
+ int *ptr;
+ int c, d;
+};
+
+int
+main (void)
+{
+ struct S a;
+ a.ptr = (int *) malloc (sizeof (int) * N);
+
+ for (int i = 0; i < N; i++)
+ a.ptr[i] = 0;
+
+ #pragma omp target enter data map(to: a.ptr, a.ptr[:N])
+
+ #pragma omp target
+ for (int i = 0; i < N; i++)
+ a.ptr[i] += 1;
+
+ #pragma omp target update from(a.ptr[:N])
+
+ for (int i = 0; i < N; i++)
+ if (a.ptr[i] != 1)
+ abort ();
+
+ #pragma omp target map(a.ptr[:N])
+ for (int i = 0; i < N; i++)
+ a.ptr[i] += 1;
+
+ #pragma omp target update from(a.ptr[:N])
+
+ for (int i = 0; i < N; i++)
+ if (a.ptr[i] != 2)
+ abort ();
+
+ #pragma omp target exit data map(from:a.ptr, a.ptr[:N])
+
+ return 0;
+}
@@ -8,7 +8,7 @@ main ()
int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
struct S s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
int *v = u + 4;
- #pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3])
+ #pragma omp target enter data map (to: s.s, s.u, s.u[0:5]) map (alloc: s.v[1:3])
s.s++;
u[3]++;
s.v[1]++;
@@ -14,7 +14,7 @@ foo (struct S s)
d = id;
int err;
- #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3]) map(to: sep) map(from: err)
+ #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(to: sep) map(from: err)
{
err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@@ -35,7 +35,7 @@ foo (struct S s)
|| omp_target_is_present (s.d, d)
|| omp_target_is_present (&s.d[-2], d)))
abort ();
- #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3])
+ #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
{
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
@@ -43,15 +43,15 @@ foo (struct S s)
|| !omp_target_is_present (s.d, d)
|| !omp_target_is_present (&s.d[-2], d))
abort ();
- #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3])
- #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err)
+ #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(from: err)
{
err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
s.a = 17; s.b[0] = 18; s.b[1] = 19;
s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
}
- #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3])
+ #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
}
if (sep
&& (omp_target_is_present (&s.a, d)
@@ -66,29 +66,29 @@ foo (struct S s)
if (err) abort ();
s.a = 33; s.b[0] = 34; s.b[1] = 35;
s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
- #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3])
+ #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
|| !omp_target_is_present (s.d, d)
|| !omp_target_is_present (&s.d[-2], d))
abort ();
- #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3])
- #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err)
+ #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(from: err)
{
err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
s.a = 49; s.b[0] = 48; s.b[1] = 47;
s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
}
- #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3])
+ #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
|| !omp_target_is_present (s.d, d)
|| !omp_target_is_present (&s.d[-2], d))
abort ();
- #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3])
+ #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
if (sep
&& (omp_target_is_present (&s.a, d)
|| omp_target_is_present (s.b, d)