public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-11] OpenMP 5.0: Remove array section base-pointer mapping semantics, and other front-end adjustments
@ 2021-05-31 9:03 Chung-Lin Tang
0 siblings, 0 replies; only message in thread
From: Chung-Lin Tang @ 2021-05-31 9:03 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:b63222b04ae06eba3cf1285ca469ad3d097e8e73
commit b63222b04ae06eba3cf1285ca469ad3d097e8e73
Author: Chung-Lin Tang <cltang@codesourcery.com>
Date: Thu May 27 22:57:09 2021 +0800
OpenMP 5.0: 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, i.e for this code:
struct S { int *ptr; ... };
struct S s;
#pragma omp target enter data map(to: s.ptr[:100])
Currently we generate after gimplify:
map(struct:s [len: 1]) map(alloc:s.ptr [len: 8]) \
map(to:*_1 [len: 400]) map(attach:s.ptr [bias: 0])
which is deemed incorrect. After this patch, the gimplify results are now
adjusted to:
map(attach:s.ptr [bias: 0])
(the attach operation is still generated, and if s.ptr is already mapped prior,
attachment will happen)
The correct way of achieving the base-pointer-also-mapped behavior would be
to use:
This adjustment in behavior required a number of small adjustments here and
there in gimplify, including to accomodate map sequences for C++ references.
(2) Fixes in libgomp/target.c to not overwrite attached pointers when handling
device<->host copies, mainly for the "always" case.
(3) A set of changes to the C/C++ front-ends to extend the allowed component
access syntax in map clauses. This is actually mainly an effort to allow
SPEC HPC to compile. These changes are enabled for both OpenACC and OpenMP.
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.
* c-c++-common/gomp/target-implicit-map-2.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.
* testsuite/libgomp.c-c++-common/target-implicit-map-2.c: New testcase.
Diff:
---
gcc/c/c-parser.c | 55 +++++-
gcc/c/c-typeck.c | 103 ++++++++--
gcc/cp/parser.c | 56 +++++-
gcc/cp/semantics.c | 155 ++++++++++++---
gcc/fortran/trans-openmp.c | 3 +
gcc/gimplify.c | 220 +++++++++++++++++----
.../c-c++-common/goacc/deep-copy-arrayofstruct.c | 5 +-
.../c-c++-common/gomp/target-enter-data-1.c | 24 +++
.../c-c++-common/gomp/target-implicit-map-2.c | 52 +++++
libgomp/target.c | 107 ++++++++--
libgomp/testsuite/libgomp.c++/target-11.C | 14 +-
libgomp/testsuite/libgomp.c++/target-12.C | 2 +-
libgomp/testsuite/libgomp.c++/target-15.C | 20 +-
libgomp/testsuite/libgomp.c++/target-16.C | 20 +-
libgomp/testsuite/libgomp.c++/target-17.C | 20 +-
libgomp/testsuite/libgomp.c++/target-21.C | 8 +-
libgomp/testsuite/libgomp.c++/target-23.C | 4 +-
.../libgomp.c-c++-common/target-implicit-map-2.c | 46 +++++
libgomp/testsuite/libgomp.c/target-23.c | 2 +-
libgomp/testsuite/libgomp.c/target-29.c | 20 +-
20 files changed, 771 insertions(+), 165 deletions(-)
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 0fc4ddb61e0..a99e3bf4bae 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -12897,6 +12897,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,
@@ -12910,6 +12919,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)
{
@@ -13029,6 +13039,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)))
@@ -13056,10 +13067,14 @@ 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 ((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))
@@ -13070,9 +13085,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 `:'. */
@@ -13101,7 +13120,33 @@ 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
@@ -16023,7 +16068,8 @@ 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,
+ C_ORT_OMP, true);
}
/* OpenMP 4.0:
@@ -16032,7 +16078,8 @@ 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,
+ C_ORT_OMP, true);
}
/* OpenMP 4.0:
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index f5aa8af21bd..a84e14f087a 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13073,6 +13073,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
@@ -13094,12 +13106,16 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
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))
@@ -13388,21 +13404,31 @@ 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.
- 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
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
&& 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;
+ }
+ }
}
}
}
@@ -14771,13 +14797,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))
@@ -14788,6 +14821,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
remove = true;
break;
}
+
if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
break;
if (bitmap_bit_p (&map_head, DECL_UID (t)))
@@ -14844,15 +14878,33 @@ 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 ((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_)
{
@@ -14888,7 +14940,8 @@ 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 ((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),
@@ -14897,6 +14950,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;
@@ -14967,7 +15029,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),
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 916a96482d7..9f11e5c15fe 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -35325,12 +35325,23 @@ 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,
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)
@@ -35411,6 +35422,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)))
@@ -35434,15 +35446,20 @@ 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 ((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))
{
+ 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));
@@ -35451,7 +35468,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 `:'. */
@@ -35464,6 +35484,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))
{
@@ -35482,8 +35504,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;
@@ -38747,11 +38793,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,
+ 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 3abc70e447c..0f2c5621b82 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -4940,6 +4940,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
@@ -4964,12 +4976,16 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
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);
@@ -5268,7 +5284,7 @@ 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.
Note that OpenACC does accept these kinds of non-contiguous pointer
based arrays. */
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
@@ -5276,13 +5292,25 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
&& 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;
+ }
+ }
}
}
}
@@ -5576,18 +5604,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 ((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
@@ -5611,8 +5658,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))))
{
@@ -7727,15 +7776,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))
@@ -7806,16 +7862,34 @@ 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 ((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
@@ -7852,6 +7926,25 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
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;
@@ -7959,7 +8052,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),
@@ -8010,8 +8104,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:
@@ -8040,7 +8139,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);
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 4d91f8023e1..94d22f31cfb 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2432,6 +2432,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
{
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 32e5b8ee5f2..ff8c5159d0c 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8518,7 +8518,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;
@@ -8565,10 +8565,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;
@@ -8577,6 +8578,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)
@@ -8590,12 +8592,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
@@ -8675,21 +8687,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
@@ -8781,6 +8879,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;
@@ -9309,6 +9408,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
@@ -9327,6 +9427,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);
}
@@ -9373,7 +9474,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> ();
@@ -9417,7 +9520,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
@@ -9452,7 +9556,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)
@@ -9465,13 +9577,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)
@@ -9503,6 +9617,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)
@@ -9542,9 +9682,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
@@ -9629,9 +9774,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)
@@ -9719,16 +9866,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);
}
@@ -10384,6 +10536,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)
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 d411bcfa8e7..4247607b61c 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 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]++;
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 00000000000..ce766d29e2d
--- /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 00000000000..3aa1a8fc55e
--- /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 7a8fa87bd85..d373fed8956 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -552,11 +552,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);
@@ -2142,16 +2161,46 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
}
- 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 @@ 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
diff --git a/libgomp/testsuite/libgomp.c++/target-11.C b/libgomp/testsuite/libgomp.c++/target-11.C
index fe99603351d..87c2980b4b5 100644
--- a/libgomp/testsuite/libgomp.c++/target-11.C
+++ b/libgomp/testsuite/libgomp.c++/target-11.C
@@ -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++)
diff --git a/libgomp/testsuite/libgomp.c++/target-12.C b/libgomp/testsuite/libgomp.c++/target-12.C
index 3b4ed57df68..480e479c262 100644
--- a/libgomp/testsuite/libgomp.c++/target-12.C
+++ b/libgomp/testsuite/libgomp.c++/target-12.C
@@ -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]++;
diff --git a/libgomp/testsuite/libgomp.c++/target-15.C b/libgomp/testsuite/libgomp.c++/target-15.C
index 4b320c31229..53626b2547e 100644
--- a/libgomp/testsuite/libgomp.c++/target-15.C
+++ b/libgomp/testsuite/libgomp.c++/target-15.C
@@ -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)
diff --git a/libgomp/testsuite/libgomp.c++/target-16.C b/libgomp/testsuite/libgomp.c++/target-16.C
index cd102d90594..b8be7cc922f 100644
--- a/libgomp/testsuite/libgomp.c++/target-16.C
+++ b/libgomp/testsuite/libgomp.c++/target-16.C
@@ -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)
diff --git a/libgomp/testsuite/libgomp.c++/target-17.C b/libgomp/testsuite/libgomp.c++/target-17.C
index d81ff19a411..f97476aafc4 100644
--- a/libgomp/testsuite/libgomp.c++/target-17.C
+++ b/libgomp/testsuite/libgomp.c++/target-17.C
@@ -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)
diff --git a/libgomp/testsuite/libgomp.c++/target-21.C b/libgomp/testsuite/libgomp.c++/target-21.C
index 21a2f299bbb..da17b5745de 100644
--- a/libgomp/testsuite/libgomp.c++/target-21.C
+++ b/libgomp/testsuite/libgomp.c++/target-21.C
@@ -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]++;
diff --git a/libgomp/testsuite/libgomp.c++/target-23.C b/libgomp/testsuite/libgomp.c++/target-23.C
index d4f9ff3e983..63d343624b0 100644
--- a/libgomp/testsuite/libgomp.c++/target-23.C
+++ b/libgomp/testsuite/libgomp.c++/target-23.C
@@ -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++)
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 00000000000..974a9786c3f
--- /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 fb1532a07b2..d56b13acf82 100644
--- a/libgomp/testsuite/libgomp.c/target-23.c
+++ b/libgomp/testsuite/libgomp.c/target-23.c
@@ -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]++;
diff --git a/libgomp/testsuite/libgomp.c/target-29.c b/libgomp/testsuite/libgomp.c/target-29.c
index e5095a1b6b8..4a286649811 100644
--- a/libgomp/testsuite/libgomp.c/target-29.c
+++ b/libgomp/testsuite/libgomp.c/target-29.c
@@ -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)
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-05-31 9:03 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-31 9:03 [gcc/devel/omp/gcc-11] OpenMP 5.0: Remove array section base-pointer mapping semantics, and other front-end adjustments Chung-Lin Tang
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).