| Message ID | 5384ea0b-6e89-6009-b486-e1b5540ca330@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | [OG10,OpenMP,5.0,committed] Remove array section base-pointer mapping semantics, and other front-end adjustments. | expand |
On Tue, 11 May 2021 19:28:04 +0800 Chung-Lin Tang <cltang@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) Oh, that's not going to play nicely (eventually?) with the patch series I just posted... we probably need to clarify what the intention is for OpenACC, but IIUC "user expectation" (i.e. existing code) expects the base-pointer mapping to happen. Julian
Hi Chung-Lin! On 2021-05-11T19:28:04+0800, Chung-Lin Tang <cltang@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 <cltang@codesourcery.com> > > gcc/c/ChangeLog: > > * c-parser.c (struct omp_dim): New struct type for use inside > c_parser_omp_variable_list. > (c_parser_omp_variable_list): Allow multiple levels of array and > component accesses in array section base-pointer expression. > (c_parser_omp_clause_to): Set 'allow_deref' to true in call to > c_parser_omp_var_list_parens. > (c_parser_omp_clause_from): Likewise. > * c-typeck.c (handle_omp_array_sections_1): Extend allowed range > of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and > POINTER_PLUS_EXPR. > (c_finish_omp_clauses): Extend allowed ranged of expressions > involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR. > > gcc/cp/ChangeLog: > > * parser.c (struct omp_dim): New struct type for use inside > cp_parser_omp_var_list_no_open. > (cp_parser_omp_var_list_no_open): Allow multiple levels of array and > component accesses in array section base-pointer expression. > (cp_parser_omp_all_clauses): Set 'allow_deref' to true in call to > cp_parser_omp_var_list for to/from clauses. > * semantics.c (handle_omp_array_sections_1): Extend allowed range > of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and > POINTER_PLUS_EXPR. > (handle_omp_array_sections): Adjust pointer map generation of > references. > (finish_omp_clauses): Extend allowed ranged of expressions > involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR. > > gcc/fortran/ChangeLog: > > * trans-openmp.c (gfc_trans_omp_array_section): Do not generate > GOMP_MAP_ALWAYS_POINTER map for main array maps of ARRAY_TYPE type. > > > gcc/ChangeLog: > > * gimplify.c (extract_base_bit_offset): Add 'tree *offsetp' parameter, > accomodate case where 'offset' return of get_inner_reference is > non-NULL. > (is_or_contains_p): Further robustify conditions. > (omp_target_reorder_clauses): In alloc/to/from sorting phase, also > move following GOMP_MAP_ALWAYS_POINTER maps along. Add new sorting > phase where we make sure pointers with an attach/detach map are ordered > correctly. > (gimplify_scan_omp_clauses): Add modifications to avoid creating > GOMP_MAP_STRUCT and associated alloc map for attach/detach maps. > > gcc/testsuite/ChangeLog: > > * c-c++-common/goacc/deep-copy-arrayofstruct.c: Adjust testcase. > * c-c++-common/gomp/target-enter-data-1.c: New testcase. > > libgomp/ChangeLog: > > * target.c (gomp_map_vars_existing): Make sure attached pointer is > not overwritten during cross-host/device copying. > (gomp_update): Likewise. > (gomp_exit_data): Likewise. > > * testsuite/libgomp.c++/target-11.C: Adjust testcase. > * testsuite/libgomp.c++/target-12.C: Likewise. > * testsuite/libgomp.c++/target-15.C: Likewise. > * testsuite/libgomp.c++/target-16.C: Likewise. > * testsuite/libgomp.c++/target-17.C: Likewise. > * testsuite/libgomp.c++/target-21.C: Likewise. > * testsuite/libgomp.c++/target-23.C: Likewise. > * testsuite/libgomp.c/target-23.c: Likewise. > * testsuite/libgomp.c/target-29.c: Likewise. > > diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c > index 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
On 2021/5/11 11:15 , Thomas Schwinge wrote: > Hi Chung-Lin! > > On 2021-05-11T19:28:04+0800, Chung-Lin Tang <cltang@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.) It is indeed a bit awkward to use, but that's what the omp-lang list seemed to decide. This change is OpenMP only. I took care to only handle OpenMP constructs like this in the middle-end, of course this does not preclude some mistake in adjusting the shared code paths... > >> 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? Maybe. although this part was needed to solve a regression caused by the above changes. >> (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? Some of the separation of base-pointer/array-section in map clauses seemed to step on this bug (e.g. if one mechanically updates "s.ptr[:N]" into "s.ptr, s.ptr[:N]", and a target-update overwrites the base-pointer) So it's arguably separate, but also can cause some testsuite chaos if not included together. > >> (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? ;-) Yeah, this might be separated :P >> 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 Weird, I did not see this in my own testing. (OTOH, I do have another bug fix for this structelem refcounting, about to submit) I'll check again if I can reproduce this, and see if I can fix this. Thanks, Chung-Lin > Grüße > Thomas > > >> Pushed to devel/omp/gcc-10, will >> send mainline version of patch later. >> >> Chung-Lin >> >> 2021-05-11 Chung-Lin Tang <cltang@codesourcery.com> >> >> gcc/c/ChangeLog: >> >> * c-parser.c (struct omp_dim): New struct type for use inside >> c_parser_omp_variable_list. >> (c_parser_omp_variable_list): Allow multiple levels of array and >> component accesses in array section base-pointer expression. >> (c_parser_omp_clause_to): Set 'allow_deref' to true in call to >> c_parser_omp_var_list_parens. >> (c_parser_omp_clause_from): Likewise. >> * c-typeck.c (handle_omp_array_sections_1): Extend allowed range >> of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and >> POINTER_PLUS_EXPR. >> (c_finish_omp_clauses): Extend allowed ranged of expressions >> involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR. >> >> gcc/cp/ChangeLog: >> >> * parser.c (struct omp_dim): New struct type for use inside >> cp_parser_omp_var_list_no_open. >> (cp_parser_omp_var_list_no_open): Allow multiple levels of array and >> component accesses in array section base-pointer expression. >> (cp_parser_omp_all_clauses): Set 'allow_deref' to true in call to >> cp_parser_omp_var_list for to/from clauses. >> * semantics.c (handle_omp_array_sections_1): Extend allowed range >> of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and >> POINTER_PLUS_EXPR. >> (handle_omp_array_sections): Adjust pointer map generation of >> references. >> (finish_omp_clauses): Extend allowed ranged of expressions >> involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR. >> >> gcc/fortran/ChangeLog: >> >> * trans-openmp.c (gfc_trans_omp_array_section): Do not generate >> GOMP_MAP_ALWAYS_POINTER map for main array maps of ARRAY_TYPE type. >> >> >> gcc/ChangeLog: >> >> * gimplify.c (extract_base_bit_offset): Add 'tree *offsetp' parameter, >> accomodate case where 'offset' return of get_inner_reference is >> non-NULL. >> (is_or_contains_p): Further robustify conditions. >> (omp_target_reorder_clauses): In alloc/to/from sorting phase, also >> move following GOMP_MAP_ALWAYS_POINTER maps along. Add new sorting >> phase where we make sure pointers with an attach/detach map are ordered >> correctly. >> (gimplify_scan_omp_clauses): Add modifications to avoid creating >> GOMP_MAP_STRUCT and associated alloc map for attach/detach maps. >> >> gcc/testsuite/ChangeLog: >> >> * c-c++-common/goacc/deep-copy-arrayofstruct.c: Adjust testcase. >> * c-c++-common/gomp/target-enter-data-1.c: New testcase. >> >> libgomp/ChangeLog: >> >> * target.c (gomp_map_vars_existing): Make sure attached pointer is >> not overwritten during cross-host/device copying. >> (gomp_update): Likewise. >> (gomp_exit_data): Likewise. >> >> * testsuite/libgomp.c++/target-11.C: Adjust testcase. >> * testsuite/libgomp.c++/target-12.C: Likewise. >> * testsuite/libgomp.c++/target-15.C: Likewise. >> * testsuite/libgomp.c++/target-16.C: Likewise. >> * testsuite/libgomp.c++/target-17.C: Likewise. >> * testsuite/libgomp.c++/target-21.C: Likewise. >> * testsuite/libgomp.c++/target-23.C: Likewise. >> * testsuite/libgomp.c/target-23.c: Likewise. >> * testsuite/libgomp.c/target-29.c: Likewise. >> >> diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c >> index 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)
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)
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. 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. (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. (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. Tested on x86_64-linux with nvptx offloading with no regressions. Pushed to devel/omp/gcc-10, will send mainline version of patch later. Chung-Lin 2021-05-11 Chung-Lin Tang <cltang@codesourcery.com> gcc/c/ChangeLog: * c-parser.c (struct omp_dim): New struct type for use inside c_parser_omp_variable_list. (c_parser_omp_variable_list): Allow multiple levels of array and component accesses in array section base-pointer expression. (c_parser_omp_clause_to): Set 'allow_deref' to true in call to c_parser_omp_var_list_parens. (c_parser_omp_clause_from): Likewise. * c-typeck.c (handle_omp_array_sections_1): Extend allowed range of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR. (c_finish_omp_clauses): Extend allowed ranged of expressions involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR. gcc/cp/ChangeLog: * parser.c (struct omp_dim): New struct type for use inside cp_parser_omp_var_list_no_open. (cp_parser_omp_var_list_no_open): Allow multiple levels of array and component accesses in array section base-pointer expression. (cp_parser_omp_all_clauses): Set 'allow_deref' to true in call to cp_parser_omp_var_list for to/from clauses. * semantics.c (handle_omp_array_sections_1): Extend allowed range of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR. (handle_omp_array_sections): Adjust pointer map generation of references. (finish_omp_clauses): Extend allowed ranged of expressions involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR. gcc/fortran/ChangeLog: * trans-openmp.c (gfc_trans_omp_array_section): Do not generate GOMP_MAP_ALWAYS_POINTER map for main array maps of ARRAY_TYPE type. gcc/ChangeLog: * gimplify.c (extract_base_bit_offset): Add 'tree *offsetp' parameter, accomodate case where 'offset' return of get_inner_reference is non-NULL. (is_or_contains_p): Further robustify conditions. (omp_target_reorder_clauses): In alloc/to/from sorting phase, also move following GOMP_MAP_ALWAYS_POINTER maps along. Add new sorting phase where we make sure pointers with an attach/detach map are ordered correctly. (gimplify_scan_omp_clauses): Add modifications to avoid creating GOMP_MAP_STRUCT and associated alloc map for attach/detach maps. gcc/testsuite/ChangeLog: * c-c++-common/goacc/deep-copy-arrayofstruct.c: Adjust testcase. * c-c++-common/gomp/target-enter-data-1.c: New testcase. libgomp/ChangeLog: * target.c (gomp_map_vars_existing): Make sure attached pointer is not overwritten during cross-host/device copying. (gomp_update): Likewise. (gomp_exit_data): Likewise. * testsuite/libgomp.c++/target-11.C: Adjust testcase. * testsuite/libgomp.c++/target-12.C: Likewise. * testsuite/libgomp.c++/target-15.C: Likewise. * testsuite/libgomp.c++/target-16.C: Likewise. * testsuite/libgomp.c++/target-17.C: Likewise. * testsuite/libgomp.c++/target-21.C: Likewise. * testsuite/libgomp.c++/target-23.C: Likewise. * testsuite/libgomp.c/target-23.c: Likewise. * testsuite/libgomp.c/target-29.c: Likewise.