Remove array section base-pointer mapping semantics, and other front-end adjustments.
This patch largely implements three pieces of functionality:
(1) Per discussion and clarification on the omp-lang mailing list, standards
conforming behavior for mapping array sections should *NOT* also map the
base-pointer. This patch adjusts OpenMP map clause behavior to do this.
(2) Fixes in libgomp/target.c to not overwrite attached pointers when handling
device<->host copying.
(3) Changes to the C/C++ front-ends to extend the allowed component access syntax
in map clauses.
2021-05-11 Chung-Lin Tang <cltang@codesourcery.com>
gcc/c/ChangeLog:
* c-parser.c (struct omp_dim): New struct type for use inside
c_parser_omp_variable_list.
(c_parser_omp_variable_list): Allow multiple levels of array and
component accesses in array section base-pointer expression.
(c_parser_omp_clause_to): Set 'allow_deref' to true in call to
c_parser_omp_var_list_parens.
(c_parser_omp_clause_from): Likewise.
* c-typeck.c (handle_omp_array_sections_1): Extend allowed range
of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and
POINTER_PLUS_EXPR.
(c_finish_omp_clauses): Extend allowed ranged of expressions
involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR.
gcc/cp/ChangeLog:
* parser.c (struct omp_dim): New struct type for use inside
cp_parser_omp_var_list_no_open.
(cp_parser_omp_var_list_no_open): Allow multiple levels of array and
component accesses in array section base-pointer expression.
(cp_parser_omp_all_clauses): Set 'allow_deref' to true in call to
cp_parser_omp_var_list for to/from clauses.
* semantics.c (handle_omp_array_sections_1): Extend allowed range
of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and
POINTER_PLUS_EXPR.
(handle_omp_array_sections): Adjust pointer map generation of
references.
(finish_omp_clauses): Extend allowed ranged of expressions
involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR.
gcc/fortran/ChangeLog:
* trans-openmp.c (gfc_trans_omp_array_section): Do not generate
GOMP_MAP_ALWAYS_POINTER map for main array maps of ARRAY_TYPE type.
gcc/ChangeLog:
* gimplify.c (extract_base_bit_offset): Add 'tree *offsetp' parameter,
accomodate case where 'offset' return of get_inner_reference is
non-NULL.
(is_or_contains_p): Further robustify conditions.
(omp_target_reorder_clauses): In alloc/to/from sorting phase, also
move following GOMP_MAP_ALWAYS_POINTER maps along. Add new sorting
phase where we make sure pointers with an attach/detach map are ordered
correctly.
(gimplify_scan_omp_clauses): Add modifications to avoid creating
GOMP_MAP_STRUCT and associated alloc map for attach/detach maps.
gcc/testsuite/ChangeLog:
* c-c++-common/goacc/deep-copy-arrayofstruct.c: Adjust testcase.
* c-c++-common/gomp/target-enter-data-1.c: New testcase.
libgomp/ChangeLog:
* target.c (gomp_map_vars_existing): Make sure attached pointer is
not overwritten during cross-host/device copying.
(gomp_update): Likewise.
(gomp_exit_data): Likewise.
* testsuite/libgomp.c++/target-11.C: Adjust testcase.
* testsuite/libgomp.c++/target-12.C: Likewise.
* testsuite/libgomp.c++/target-15.C: Likewise.
* testsuite/libgomp.c++/target-16.C: Likewise.
* testsuite/libgomp.c++/target-17.C: Likewise.
* testsuite/libgomp.c++/target-21.C: Likewise.
* testsuite/libgomp.c++/target-23.C: Likewise.
* testsuite/libgomp.c/target-23.c: Likewise.
* testsuite/libgomp.c/target-29.c: Likewise.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 0a6aee43..ecc3e12 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -12893,6 +12893,15 @@
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,
@@ -12906,6 +12915,7 @@
while (1)
{
+ auto_vec<omp_dim> dims;
bool array_section_p = false;
if (kind == OMP_CLAUSE_DEPEND)
{
@@ -13025,6 +13035,7 @@
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)))
@@ -13051,10 +13062,14 @@
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_IN_REDUCTION:
case OMP_CLAUSE_TASK_REDUCTION:
+ array_section_p = false;
+ dims.truncate (0);
while ((ort != C_ORT_ACC || kind != OMP_CLAUSE_REDUCTION)
&& 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))
@@ -13065,9 +13080,13 @@
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 `:'. */
@@ -13096,8 +13115,35 @@
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
&& t != error_mark_node
&& parser->tokens_avail != 2)
@@ -15892,7 +15938,8 @@
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, C_ORT_OMP,
+ true);
}
/* OpenMP 4.0:
@@ -15901,7 +15948,8 @@
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, C_ORT_OMP,
+ true);
}
/* OpenMP 4.0:
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 7c887a8..c8bcbdd 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12896,6 +12896,18 @@
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
@@ -12917,12 +12929,16 @@
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
- if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
- && TREE_CODE (t) == MEM_REF)
- {
- t = TREE_OPERAND (t, 0);
- STRIP_NOPS (t);
- }
+ if (ort == C_ORT_ACC || ort == C_ORT_OMP)
+ 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)
{
if (maybe_ne (mem_ref_offset (t), 0))
@@ -13204,20 +13220,30 @@
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.
- Note that OpenACC does accept these kinds of non-contiguous pointer
- based arrays. */
+ array-section-subscript, the array section could be non-contiguous. */
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
&& TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
{
if (ort == C_ORT_ACC)
+ /* Note that OpenACC does accept these kinds of non-contiguous
+ pointer based arrays. */
non_contiguous = true;
else
{
- error_at (OMP_CLAUSE_LOCATION (c),
- "array section is not contiguous in %qs clause",
- omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
- return error_mark_node;
+ /* 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;
+ }
+ }
}
}
}
@@ -14510,13 +14536,20 @@
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 (bitmap_bit_p (&map_field_head, DECL_UID (t)))
break;
if (bitmap_bit_p (&map_head, DECL_UID (t)))
@@ -14573,15 +14606,33 @@
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 ((ort == C_ORT_ACC || ort == C_ORT_OMP)
&& 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_)
{
@@ -14617,7 +14668,8 @@
break;
}
t = TREE_OPERAND (t, 0);
- if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
+ if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+ && TREE_CODE (t) == MEM_REF)
{
if (maybe_ne (mem_ref_offset (t), 0))
error_at (OMP_CLAUSE_LOCATION (c),
@@ -14626,6 +14678,15 @@
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;
@@ -14690,7 +14751,8 @@
"%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),
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 9fc2a9b..27aef0d 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -34219,12 +34219,23 @@
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,
enum c_omp_region_type ort = C_ORT_OMP,
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)
@@ -34306,6 +34317,7 @@
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)))
@@ -34328,20 +34340,30 @@
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_IN_REDUCTION:
case OMP_CLAUSE_TASK_REDUCTION:
+ array_section_p = false;
+ dims.truncate (0);
while ((ort != C_ORT_ACC || kind != OMP_CLAUSE_REDUCTION)
&& 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))
- low_bound = cp_parser_expression (parser);
+ {
+ loc = cp_lexer_peek_token (parser->lexer)->location;
+ low_bound = cp_parser_expression (parser);
+ }
if (!colon)
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 `:'. */
@@ -34354,6 +34376,8 @@
}
if (kind == OMP_CLAUSE_DEPEND)
cp_parser_commit_to_tentative_parse (parser);
+ else
+ array_section_p = true;
if (!cp_lexer_next_token_is (parser->lexer,
CPP_CLOSE_SQUARE))
length = cp_parser_expression (parser);
@@ -34368,8 +34392,32 @@
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;
@@ -37472,11 +37520,13 @@
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,
+ C_ORT_OMP, 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,
+ C_ORT_OMP, true);
c_name = "from";
break;
case PRAGMA_OMP_CLAUSE_UNIFORM:
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 3e29076..57d5df3 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -4762,6 +4762,18 @@
&& 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
@@ -4786,12 +4798,16 @@
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
- if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
- && TREE_CODE (t) == INDIRECT_REF)
- {
- t = TREE_OPERAND (t, 0);
- STRIP_NOPS (t);
- }
+ if (ort == C_ORT_ACC || ort == C_ORT_OMP)
+ 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))
t = TREE_OPERAND (t, 0);
@@ -5085,20 +5101,30 @@
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.
- Note that OpenACC does accept these kinds of non-contiguous pointer
- based arrays. */
+ array-section-subscript, the array section could be non-contiguous. */
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
&& TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
{
if (ort == C_ORT_ACC)
+ /* Note that OpenACC does accept these kinds of non-contiguous
+ pointer based arrays. */
non_contiguous = true;
else
{
- error_at (OMP_CLAUSE_LOCATION (c),
- "array section is not contiguous in %qs clause",
- omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
- return error_mark_node;
+ /* 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;
+ }
+ }
}
}
}
@@ -5390,18 +5416,37 @@
default:
break;
}
+ bool reference_always_pointer = true;
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
else 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 && 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 && 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
@@ -5424,8 +5469,10 @@
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))))
{
@@ -7412,15 +7459,22 @@
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 (bitmap_bit_p (&map_field_head, DECL_UID (t)))
break;
if (bitmap_bit_p (&map_head, DECL_UID (t)))
@@ -7481,16 +7535,34 @@
&& 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 ((ort == C_ORT_ACC || ort == C_ORT_OMP)
&& 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
&& ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
@@ -7527,6 +7599,25 @@
break;
}
t = TREE_OPERAND (t, 0);
+ if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+ && 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;
@@ -7627,7 +7718,8 @@
"%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),
@@ -7675,8 +7767,13 @@
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:
@@ -7705,7 +7802,7 @@
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);
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index e3df4bb..d366703 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2242,6 +2242,9 @@
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
{
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index ba071e8..e51f0dd 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8331,7 +8331,7 @@
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;
@@ -8378,10 +8378,11 @@
&& 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;
@@ -8390,6 +8391,7 @@
*bitposp = bitpos;
*poffsetp = poffset;
+ *offsetp = offset;
/* Set *BASE_REF if BASE was a dereferenced reference variable. */
if (base_ref && orig_base != base)
@@ -8403,12 +8405,17 @@
static bool
is_or_contains_p (tree expr, tree base_ptr)
{
- while (expr != base_ptr)
- if (TREE_CODE (base_ptr) == COMPONENT_REF)
+ 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) == COMPONENT_REF
+ || TREE_CODE (base_ptr) == POINTER_PLUS_EXPR)
base_ptr = TREE_OPERAND (base_ptr, 0);
else
break;
- return expr == base_ptr;
+ return operand_equal_p (expr, base_ptr);
}
/* Implement OpenMP 5.x map ordering rules for target directives. There are
@@ -8488,21 +8495,107 @@
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;
+ }
+ }
+ }
+ }
+ }
}
/* Scan the OMP clauses in *LIST_P, installing mappings into a new
@@ -8516,6 +8609,7 @@
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;
@@ -9092,6 +9186,7 @@
}
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
@@ -9110,6 +9205,7 @@
== POINTER_TYPE))
{
indir_p = true;
+ indir_base = decl;
decl = TREE_OPERAND (decl, 0);
STRIP_NOPS (decl);
}
@@ -9156,7 +9252,9 @@
!= 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> ();
@@ -9200,7 +9298,8 @@
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
@@ -9235,7 +9334,15 @@
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)
@@ -9248,13 +9355,15 @@
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)
@@ -9286,6 +9395,32 @@
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)
@@ -9325,9 +9460,14 @@
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
@@ -9412,9 +9552,11 @@
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)
@@ -9502,16 +9644,21 @@
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);
}
@@ -10139,6 +10286,8 @@
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)
diff --git a/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c
index d411bcf..4247607 100644
--- a/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c
+++ b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c
@@ -37,13 +37,12 @@
{
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]++;
diff --git a/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c b/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c
new file mode 100644
index 0000000..ce766d2
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c
@@ -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" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c
new file mode 100644
index 0000000..3aa1a8f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c
@@ -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" } } */
diff --git a/libgomp/target.c b/libgomp/target.c
index ecda2ef..500631e 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -552,11 +552,30 @@
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);
@@ -2142,16 +2161,46 @@
}
- 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);
@@ -3025,11 +3074,31 @@
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
diff --git a/libgomp/testsuite/libgomp.c++/target-11.C b/libgomp/testsuite/libgomp.c++/target-11.C
index fe99603..87c2980 100644
--- a/libgomp/testsuite/libgomp.c++/target-11.C
+++ b/libgomp/testsuite/libgomp.c++/target-11.C
@@ -23,9 +23,11 @@
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 @@
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++)
diff --git a/libgomp/testsuite/libgomp.c++/target-12.C b/libgomp/testsuite/libgomp.c++/target-12.C
index 3b4ed57..480e479 100644
--- a/libgomp/testsuite/libgomp.c++/target-12.C
+++ b/libgomp/testsuite/libgomp.c++/target-12.C
@@ -53,7 +53,7 @@
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]++;
diff --git a/libgomp/testsuite/libgomp.c++/target-15.C b/libgomp/testsuite/libgomp.c++/target-15.C
index 4b320c3..53626b2 100644
--- a/libgomp/testsuite/libgomp.c++/target-15.C
+++ b/libgomp/testsuite/libgomp.c++/target-15.C
@@ -14,7 +14,7 @@
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 @@
|| 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 @@
|| !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 @@
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 @@
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 @@
|| !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 @@
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 @@
|| !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)
diff --git a/libgomp/testsuite/libgomp.c++/target-16.C b/libgomp/testsuite/libgomp.c++/target-16.C
index cd102d9..b8be7cc 100644
--- a/libgomp/testsuite/libgomp.c++/target-16.C
+++ b/libgomp/testsuite/libgomp.c++/target-16.C
@@ -16,7 +16,7 @@
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 @@
|| 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 @@
|| !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 @@
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 @@
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 @@
|| !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 @@
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 @@
|| !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)
diff --git a/libgomp/testsuite/libgomp.c++/target-17.C b/libgomp/testsuite/libgomp.c++/target-17.C
index d81ff19..f97476a 100644
--- a/libgomp/testsuite/libgomp.c++/target-17.C
+++ b/libgomp/testsuite/libgomp.c++/target-17.C
@@ -16,7 +16,7 @@
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 @@
|| 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 @@
|| !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 @@
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 @@
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 @@
|| !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 @@
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 @@
|| !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)
diff --git a/libgomp/testsuite/libgomp.c++/target-21.C b/libgomp/testsuite/libgomp.c++/target-21.C
index 21a2f29..da17b57 100644
--- a/libgomp/testsuite/libgomp.c++/target-21.C
+++ b/libgomp/testsuite/libgomp.c++/target-21.C
@@ -7,7 +7,7 @@
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 @@
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 @@
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 @@
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]++;
diff --git a/libgomp/testsuite/libgomp.c++/target-23.C b/libgomp/testsuite/libgomp.c++/target-23.C
index d4f9ff3..63d3436 100644
--- a/libgomp/testsuite/libgomp.c++/target-23.C
+++ b/libgomp/testsuite/libgomp.c++/target-23.C
@@ -16,13 +16,13 @@
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++)
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c
new file mode 100644
index 0000000..974a978
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c
@@ -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;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-23.c b/libgomp/testsuite/libgomp.c/target-23.c
index fb1532a..d56b13a 100644
--- a/libgomp/testsuite/libgomp.c/target-23.c
+++ b/libgomp/testsuite/libgomp.c/target-23.c
@@ -8,7 +8,7 @@
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]++;
diff --git a/libgomp/testsuite/libgomp.c/target-29.c b/libgomp/testsuite/libgomp.c/target-29.c
index e5095a1..4a28664 100644
--- a/libgomp/testsuite/libgomp.c/target-29.c
+++ b/libgomp/testsuite/libgomp.c/target-29.c
@@ -14,7 +14,7 @@
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 @@
|| 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 @@
|| !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 @@
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)
