Hi Chung-Lin!
On 2021-05-11T19:28:04+0800, Chung-Lin Tang <clt...@codesourcery.com> wrote:
> 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:
> #pragma omp target enter data 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:
> #pragma omp target enter data map(to:*_1 [len: 400]) 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:
> #pragma omp target enter data map(to: s.ptr, s.ptr[:100])
>
> This adjustment in behavior required a number of small adjustments here and
> there in gimplify, including
> to accomodate map sequences for C++ references.
I'm a bit confused by that -- this mandates the bulk of the testsuite
changes that you've included, and these seem a step backwards in terms of
user experience, but then, I have no state on the exact OpenMP
specification requirements, so you certainly may be right on that. (And
also, as Julian mentioned, how this relates to OpenACC semantics, which I
also haven't considered in detail -- but I note you didn't adjust any
OpenACC testcases for that, so I suppose that's really conditionalized to
OpenMP only.)
> There is also a small Fortran front-end patch involved (hence CCing Tobias).
> The new gimplify processing changed behavior in handling
> GOMP_MAP_ALWAYS_POINTER maps such that
> the libgomp.fortran/struct-elem-map-1.f90 regressed. It appeared that the
> Fortran FE was generating
> a GOMP_MAP_ALWAYS_POINTER for array types, which didn't seem quite correct,
> and the pre-patch behavior
> was removing this map anyways. I have a small change in
> trans-openmp.c:gfc_trans_omp_array_section
> to not generate the map in this case, and so far no bad test results.
Makes sense to argue that one separately, with testcases, for the master
branch submission?
> (2) The second part (though kind of related to the first above) are fixes in
> libgomp/target.c
> to not overwrite attached pointers when handling device<->host copies, mainly
> for the "always" case.
> This behavior is also noted in the 5.0 spec, but not yet properly coded
> before.
Likewise, if that makes sense?
> (3) The third is 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, so despite in the long
> term the entire map clause syntax parsing is probably going to be revamped,
> we're still adding this in
> for now. These changes are enabled for both OpenACC and OpenMP.
Likewise, if that makes sense? ;-)
> Tested on x86_64-linux with nvptx offloading with no regressions.
I'm seeing a regression with
'libgomp.oacc-c-c++-common/noncontig_array-1.c' execution testing, both C
and C++, for '-O2' (but not '-O0'), and only for about half of the
invocations. But it seems to reliable reproduce in GDB:
Thread 1 "a.out" received signal SIGSEGV, Segmentation fault.
gomp_decrement_refcount (do_remove=<synthetic pointer>, do_copy=<synthetic
pointer>, delete_p=false, refcount_set=0x0, k=0xc4d450) at
[...]/source-gcc/libgomp/target.c:468
468 uintptr_t orig_refcount = *refcount_ptr;
(gdb) bt
#0 gomp_decrement_refcount (do_remove=<synthetic pointer>,
do_copy=<synthetic pointer>, delete_p=false, refcount_set=0x0, k=0xc4d450) at
[...]/source-gcc/libgomp/target.c:468
#1 gomp_unmap_vars_internal (aq=0x0, aq@entry=0x8223c0, refcount_set=0x0,
do_copyfrom=<optimized out>, do_copyfrom@entry=true, tgt=tgt@entry=0xc696a0) at
[...]/source-gcc/libgomp/target.c:2065
#2 goacc_unmap_vars (tgt=tgt@entry=0xc696a0,
do_copyfrom=do_copyfrom@entry=true, aq=aq@entry=0x0) at
[...]/source-gcc/libgomp/target.c:2118
#3 0x00007ffff7daa41c in GOACC_parallel_keyed (flags_m=flags_m@entry=-1,
fn=fn@entry=0x400ae0 <test3._omp_fn.0>, mapnum=mapnum@entry=2,
hostaddrs=hostaddrs@entry=0x7fffffffd7a0, sizes=sizes@entry=0x604500
<omp_data_sizes.40>, kinds=kinds@entry=0x6044f0 <omp_data_kinds.41>) at
[...]/source-gcc/libgomp/oacc-parallel.c:639
#4 0x0000000000400f11 in test3 () at
source-gcc/libgomp/testsuite/libgomp.oacc-c/../libgomp.oacc-c-c++-common/noncontig_array-1.c:75
#5 0x00000000004008f3 in main () at
source-gcc/libgomp/testsuite/libgomp.oacc-c/../libgomp.oacc-c-c++-common/noncontig_array-1.c:101
(gdb) print refcount_ptr
$1 = (uintptr_t *) 0x100000000
(gdb) list 457,468
457 uintptr_t *refcount_ptr = &k->refcount;
458
459 if (REFCOUNT_STRUCTELEM_FIRST_P (k->refcount))
460 refcount_ptr = &k->structelem_refcount;
461 else if (REFCOUNT_STRUCTELEM_P (k->refcount))
462 refcount_ptr = k->structelem_refcount_ptr;
[...]
468 uintptr_t orig_refcount = *refcount_ptr;
(gdb) print &k->refcount
$2 = (uintptr_t *) 0xc4d470
(gdb) print &k->structelem_refcount
$3 = (uintptr_t *) 0xc4d478
(gdb) print k->structelem_refcount_ptr
$4 = (uintptr_t *) 0x100000000
Grüße
Thomas
> Pushed to devel/omp/gcc-10, will
> send mainline version of patch later.
>
> Chung-Lin
>
> 2021-05-11 Chung-Lin Tang <clt...@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 0a6aee439f6..ecc3e12cf78 100644
> --- a/gcc/c/c-parser.c
> +++ b/gcc/c/c-parser.c
> @@ -12893,6 +12893,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,
> @@ -12906,6 +12915,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)
> {
> @@ -13025,6 +13035,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)))
> @@ -13051,10 +13062,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))
> @@ -13065,9 +13080,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 `:'. */
> @@ -13096,8 +13115,35 @@ c_parser_omp_variable_list (c_parser *parser,
> break;
> }
>
> - t = tree_cons (low_bound, length, t);
> + dims.safe_push (omp_dim (low_bound, length, loc, no_colon));
> + }
> +
> + if (t != error_mark_node)
> + {
> + if ((kind == OMP_CLAUSE_MAP
> + || kind == OMP_CLAUSE_FROM
> + || kind == OMP_CLAUSE_TO)
> + && !array_section_p
> + && (c_parser_next_token_is (parser, CPP_DOT)
> + || (allow_deref
> + && c_parser_next_token_is (parser,
> + CPP_DEREF))))
> + {
> + for (unsigned i = 0; i < dims.length (); i++)
> + {
> + gcc_assert (dims[i].length == integer_one_node);
> + t = build_array_ref (dims[i].loc,
> + t, dims[i].low_bound);
> + }
> + goto start_component_ref;
> + }
> + else
> + {
> + for (unsigned i = 0; i < dims.length (); i++)
> + t = tree_cons (dims[i].low_bound, dims[i].length, t);
> + }
> }
> +
> if (kind == OMP_CLAUSE_DEPEND
> && t != error_mark_node
> && parser->tokens_avail != 2)
> @@ -15892,7 +15938,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:
> @@ -15901,7 +15948,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 7c887a80ce9..c8bcbdd4473 100644
> --- a/gcc/c/c-typeck.c
> +++ b/gcc/c/c-typeck.c
> @@ -12896,6 +12896,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
> @@ -12917,12 +12929,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))
> @@ -13204,20 +13220,30 @@ 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
> && 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 @@ 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 (bitmap_bit_p (&map_field_head, DECL_UID (t)))
> break;
> if (bitmap_bit_p (&map_head, DECL_UID (t)))
> @@ -14573,15 +14606,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_)
> {
> @@ -14617,7 +14668,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),
> @@ -14626,6 +14678,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;
> @@ -14690,7 +14751,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 9fc2a9b05eb..27aef0dd245 100644
> --- a/gcc/cp/parser.c
> +++ b/gcc/cp/parser.c
> @@ -34219,12 +34219,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)
> @@ -34306,6 +34317,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)))
> @@ -34328,20 +34340,30 @@ 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))
> - 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 @@ cp_parser_omp_var_list_no_open (cp_parser *parser,
> enum omp_clause_code kind,
> }
> 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 @@ 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;
> @@ -37472,11 +37520,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 3e290767d5c..57d5df337b0 100644
> --- a/gcc/cp/semantics.c
> +++ b/gcc/cp/semantics.c
> @@ -4762,6 +4762,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
> @@ -4786,12 +4798,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);
> @@ -5085,20 +5101,30 @@ 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
> && 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 @@ 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
> @@ -5424,8 +5469,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))))
> {
> @@ -7412,15 +7459,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 (bitmap_bit_p (&map_field_head, DECL_UID (t)))
> break;
> if (bitmap_bit_p (&map_head, DECL_UID (t)))
> @@ -7481,16 +7535,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
> @@ -7527,6 +7599,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;
> @@ -7627,7 +7718,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),
> @@ -7675,8 +7767,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:
> @@ -7705,7 +7802,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 e3df4bbf84e..d3667031ca9 100644
> --- a/gcc/fortran/trans-openmp.c
> +++ b/gcc/fortran/trans-openmp.c
> @@ -2242,6 +2242,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 ba071e8ae68..e51f0dd7787 100644
> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -8331,7 +8331,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;
> @@ -8378,10 +8378,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;
>
> @@ -8390,6 +8391,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)
> @@ -8403,12 +8405,17 @@ 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)
> + 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 @@ 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;
> + }
> + }
> + }
> + }
> + }
> }
>
> /* Scan the OMP clauses in *LIST_P, installing mappings into a new
> @@ -8516,6 +8609,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;
> @@ -9092,6 +9186,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
> @@ -9110,6 +9205,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);
> }
> @@ -9156,7 +9252,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> ();
> @@ -9200,7 +9298,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
> @@ -9235,7 +9334,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)
> @@ -9248,13 +9355,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)
> @@ -9286,6 +9395,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)
> @@ -9325,9 +9460,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
> @@ -9412,9 +9552,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)
> @@ -9502,16 +9644,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);
> }
> @@ -10139,6 +10286,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 ecda2efe34c..500631e0151 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)
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank
Thürauf
