From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 71C673858C2C; Thu, 14 Dec 2023 14:26:47 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 71C673858C2C Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 71C673858C2C Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=68.232.137.252 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1702564011; cv=none; b=IpI/A7n2munAN3vM/umsZQbFAoCDaisoiB6UlF7klizNeZwJRzRN/wULJfUKxH7A5d1gmFF6e7ATif5jloRXgCrH0/Z1CR/+hhFyAcWau5sl/T7RVjcdTroTQUG+MyxyKkJbXwLWE6Lmq+9ov/220voQPlw1OVCqMWXkjul8sHI= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1702564011; c=relaxed/simple; bh=xmRyvr1KRv3Zgn/UU12XQ2VirWVSlt8OoNZYc1USWTg=; h=Message-ID:Date:MIME-Version:Subject:To:From; b=jslY6zP1ljp08rBBkAoeMbbcrDT7e2VmXvQpMvp/nZA4jaoxvVNGpE6eOp3PfWDxtTOiHifSyWTmJ2ToIExSsAyC1axSjhVk1wJFbtBCc2KYa3ehZgLO/4vWtD2Wy/5ROQyMNxWbsFjOUsXOr3/QsqiWycPM5+KfkDoWfbGC65U= ARC-Authentication-Results: i=1; server2.sourceware.org X-CSE-ConnectionGUID: 3SLGY1LiTgipVuVx77xX0Q== X-CSE-MsgGUID: oUtZplsARyG3afKjK1nTJQ== X-IronPort-AV: E=Sophos;i="6.04,275,1695715200"; d="scan'208";a="25258057" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 14 Dec 2023 06:26:43 -0800 IronPort-SDR: av+dvhx/D1GdFS4HbzwiaktBe3sdxQ39+tVbMwYoM7BMOqA7gNF5QTiuLM23FRflw50oJBXshV Z1xbllQbFpTtVbwHK4f57B1Bp++gzsRAKbN5FxZt17yYscOBqv+fzBE1jdHL+LULkYfAD2ttdg pG0IJaqBxbej3xl1glFa2uoZu/0c0B6km2eXhVYv6LXlBEy4GDvsOlsNESqxLNr45D6RJCKUiA zSnQWy+BClvMN6rTmUaKD+prG7aOUXu+FjrKjM0S7jzICini7hy+vJ/zWG6hynu/mHt8TYLxpw Dp8= Message-ID: Date: Thu, 14 Dec 2023 15:26:38 +0100 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird Subject: Re: [PATCH v7 4/5] OpenMP/OpenACC: Unordered/non-constant component offset runtime diagnostic Content-Language: en-US To: Julian Brown , CC: , References: <81839b2435cb8b4ae46c09f2ff240eb9f679d389.1692398074.git.julian@codesourcery.com> From: Tobias Burnus In-Reply-To: <81839b2435cb8b4ae46c09f2ff240eb9f679d389.1692398074.git.julian@codesourcery.com> Content-Type: text/plain; charset="UTF-8"; format=flowed Content-Transfer-Encoding: quoted-printable X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-11.3 required=5.0 tests=BAYES_00,GIT_PATCH_0,HEADER_FROM_DIFFERENT_DOMAINS,KAM_DMARC_STATUS,KAM_SHORT,SPF_HELO_PASS,SPF_PASS,TXREP,T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: On 19.08.23 00:47, Julian Brown wrote: > This patch adds support for non-constant component offsets in "map" > clauses for OpenMP (and the equivalants for OpenACC), which are not able > to be sorted into order at compile time. Normally struct accesses in > such clauses are gathered together and sorted into increasing address > order after a "GOMP_MAP_STRUCT" node: if we have variable indices, > that is no longer possible. > > This version of the patch scales back the previously-posted version to > merely add a diagnostic for incorrect usage of component accesses with > variably-indexed arrays of structs: the only permitted variant is where > we have multiple indices that are the same, but we could not prove so > at compile time. Rather than silently producing the wrong result for > cases where the indices are in fact different, we error out (e.g., > "map(dtarr(i)%arrptr, dtarr(j)%arrptr(4:8))", for different i/j). > > For now, multiple *constant* array indices are still supported (see > map-arrayofstruct-1.c). That could perhaps be addressed with a follow-up > patch, if necessary. > > This version of the patch renumbers the GOMP_MAP_STRUCT_UNORD kind to > avoid clashing with the OpenACC "non-contiguous" dynamic array support > (though that is not yet applied to mainline). LGTM with: - inclusion of your follow-up fix for shared-memory systems (see email of August 21) - adding a comment to map-arrayofstruct-1.c indicating that this usage is an extension, violating a restriction (be a bit more explicit that just that) See https://gcc.gnu.org/pipermail/gcc-patches/2022-October/603126.html for a quote of the specification or (same wording, newer spec) in TR12 under "Restrictions to the map clause are as follows:" in "6.8.3 map Clause" [218+219:36-37+1-3] Thanks, Tobias > 2023-08-18 Julian Brown > > gcc/ > * gimplify.cc (extract_base_bit_offset): Add VARIABLE_OFFSET parame= ter. > (omp_get_attachment, omp_group_last, omp_group_base, > omp_directive_maps_explicitly): Add GOMP_MAP_STRUCT_UNORD support. > (omp_accumulate_sibling_list): Update calls to extract_base_bit_off= set. > Support GOMP_MAP_STRUCT_UNORD. > (omp_build_struct_sibling_lists, gimplify_scan_omp_clauses, > gimplify_adjust_omp_clauses, gimplify_omp_target_update): Add > GOMP_MAP_STRUCT_UNORD support. > * omp-low.cc (lower_omp_target): Add GOMP_MAP_STRUCT_UNORD support. > * tree-pretty-print.cc (dump_omp_clause): Likewise. > > include/ > * gomp-constants.h (gomp_map_kind): Add GOMP_MAP_STRUCT_UNORD. > > libgomp/ > * oacc-mem.c (find_group_last, goacc_enter_data_internal, > goacc_exit_data_internal, GOACC_enter_exit_data): Add > GOMP_MAP_STRUCT_UNORD support. > * target.c (gomp_map_vars_internal): Add GOMP_MAP_STRUCT_UNORD supp= ort. > Detect incorrect use of variable indexing of arrays of structs. > (GOMP_target_enter_exit_data, gomp_target_task_fn): Add > GOMP_MAP_STRUCT_UNORD support. > * testsuite/libgomp.c-c++-common/map-arrayofstruct-1.c: New test. > * testsuite/libgomp.c-c++-common/map-arrayofstruct-2.c: New test. > * testsuite/libgomp.c-c++-common/map-arrayofstruct-3.c: New test. > * testsuite/libgomp.fortran/map-subarray-5.f90: New test. > --- > gcc/gimplify.cc | 110 ++++++++++++++---- > gcc/omp-low.cc | 1 + > gcc/tree-pretty-print.cc | 3 + > include/gomp-constants.h | 6 + > libgomp/oacc-mem.c | 6 +- > libgomp/target.c | 60 +++++++++- > .../map-arrayofstruct-1.c | 38 ++++++ > .../map-arrayofstruct-2.c | 58 +++++++++ > .../map-arrayofstruct-3.c | 68 +++++++++++ > .../libgomp.fortran/map-subarray-5.f90 | 54 +++++++++ > 10 files changed, 377 insertions(+), 27 deletions(-) > create mode 100644 libgomp/testsuite/libgomp.c-c++-common/map-arrayofst= ruct-1.c > create mode 100644 libgomp/testsuite/libgomp.c-c++-common/map-arrayofst= ruct-2.c > create mode 100644 libgomp/testsuite/libgomp.c-c++-common/map-arrayofst= ruct-3.c > create mode 100644 libgomp/testsuite/libgomp.fortran/map-subarray-5.f90 > > diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc > index fad4308a0eb4..e682583054b0 100644 > --- a/gcc/gimplify.cc > +++ b/gcc/gimplify.cc > @@ -8965,7 +8965,8 @@ build_omp_struct_comp_nodes (enum tree_code code, t= ree grp_start, tree grp_end, > > static tree > extract_base_bit_offset (tree base, poly_int64 *bitposp, > - poly_offset_int *poffsetp) > + poly_offset_int *poffsetp, > + bool *variable_offset) > { > tree offset; > poly_int64 bitsize, bitpos; > @@ -8983,10 +8984,13 @@ extract_base_bit_offset (tree base, poly_int64 *b= itposp, > if (offset && poly_int_tree_p (offset)) > { > poffset =3D wi::to_poly_offset (offset); > - offset =3D NULL_TREE; > + *variable_offset =3D false; > } > else > - poffset =3D 0; > + { > + poffset =3D 0; > + *variable_offset =3D (offset !=3D NULL_TREE); > + } > > if (maybe_ne (bitpos, 0)) > poffset +=3D bits_to_bytes_round_down (bitpos); > @@ -9166,6 +9170,7 @@ omp_get_attachment (omp_mapping_group *grp) > return error_mark_node; > > case GOMP_MAP_STRUCT: > + case GOMP_MAP_STRUCT_UNORD: > case GOMP_MAP_FORCE_DEVICEPTR: > case GOMP_MAP_DEVICE_RESIDENT: > case GOMP_MAP_LINK: > @@ -9271,6 +9276,7 @@ omp_group_last (tree *start_p) > break; > > case GOMP_MAP_STRUCT: > + case GOMP_MAP_STRUCT_UNORD: > { > unsigned HOST_WIDE_INT num_mappings > =3D tree_to_uhwi (OMP_CLAUSE_SIZE (c)); > @@ -9437,6 +9443,7 @@ omp_group_base (omp_mapping_group *grp, unsigned in= t *chained, > return error_mark_node; > > case GOMP_MAP_STRUCT: > + case GOMP_MAP_STRUCT_UNORD: > { > unsigned HOST_WIDE_INT num_mappings > =3D tree_to_uhwi (OMP_CLAUSE_SIZE (node)); > @@ -10079,7 +10086,8 @@ omp_directive_maps_explicitly (hash_map /* We might be called during omp_build_struct_sibling_lists, when > GOMP_MAP_STRUCT might have been inserted at the start of the group= . > Skip over that, and also possibly the node after it. */ > - if (OMP_CLAUSE_MAP_KIND (grp_first) =3D=3D GOMP_MAP_STRUCT) > + if (OMP_CLAUSE_MAP_KIND (grp_first) =3D=3D GOMP_MAP_STRUCT > + || OMP_CLAUSE_MAP_KIND (grp_first) =3D=3D GOMP_MAP_STRUCT_UNORD) > { > grp_first =3D OMP_CLAUSE_CHAIN (grp_first); > if (OMP_CLAUSE_MAP_KIND (grp_first) =3D=3D GOMP_MAP_FIRSTPRIVATE_= POINTER > @@ -10816,7 +10824,9 @@ omp_accumulate_sibling_list (enum omp_region_type= region_type, > } > } > > - tree base =3D extract_base_bit_offset (ocd, &cbitpos, &coffset); > + bool variable_offset; > + tree base > + =3D extract_base_bit_offset (ocd, &cbitpos, &coffset, &variable_offs= et); > > int base_token; > for (base_token =3D addr_tokens.length () - 1; base_token >=3D 0; bas= e_token--) > @@ -10850,14 +10860,20 @@ omp_accumulate_sibling_list (enum omp_region_ty= pe region_type, > > if (!struct_map_to_clause || struct_map_to_clause->get (base) =3D=3D = NULL) > { > - tree l =3D build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), OMP_CL= AUSE_MAP); > - > - OMP_CLAUSE_SET_MAP_KIND (l, GOMP_MAP_STRUCT); > - OMP_CLAUSE_DECL (l) =3D unshare_expr (base); > - OMP_CLAUSE_SIZE (l) =3D size_int (1); > + enum gomp_map_kind str_kind =3D GOMP_MAP_STRUCT; > > if (struct_map_to_clause =3D=3D NULL) > struct_map_to_clause =3D new hash_map; > + > + if (variable_offset) > + str_kind =3D GOMP_MAP_STRUCT_UNORD; > + > + tree l =3D build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), OMP_CL= AUSE_MAP); > + > + OMP_CLAUSE_SET_MAP_KIND (l, str_kind); > + OMP_CLAUSE_DECL (l) =3D unshare_expr (base); > + OMP_CLAUSE_SIZE (l) =3D size_int (1); > + > struct_map_to_clause->put (base, l); > > /* On first iterating through the clause list, we insert the stru= ct node > @@ -11097,6 +11113,11 @@ omp_accumulate_sibling_list (enum omp_region_typ= e region_type, > { > tree *osc =3D struct_map_to_clause->get (base); > tree *sc =3D NULL, *scp =3D NULL; > + bool unordered =3D false; > + > + if (osc && OMP_CLAUSE_MAP_KIND (*osc) =3D=3D GOMP_MAP_STRUCT_UNORD= ) > + unordered =3D true; > + > unsigned HOST_WIDE_INT i, elems =3D tree_to_uhwi (OMP_CLAUSE_SIZE= (*osc)); > sc =3D &OMP_CLAUSE_CHAIN (*osc); > /* The struct mapping might be immediately followed by a > @@ -11137,12 +11158,20 @@ omp_accumulate_sibling_list (enum omp_region_ty= pe region_type, > =3D=3D REFERENCE_TYPE)) > sc_decl =3D TREE_OPERAND (sc_decl, 0); > > - tree base2 =3D extract_base_bit_offset (sc_decl, &bitpos, &offs= et); > + bool variable_offset2; > + tree base2 =3D extract_base_bit_offset (sc_decl, &bitpos, &offs= et, > + &variable_offset2); > if (!base2 || !operand_equal_p (base2, base, 0)) > break; > if (scp) > continue; > - if ((region_type & ORT_ACC) !=3D 0) > + if (variable_offset2) > + { > + OMP_CLAUSE_SET_MAP_KIND (*osc, GOMP_MAP_STRUCT_UNORD); > + unordered =3D true; > + break; > + } > + else if ((region_type & ORT_ACC) !=3D 0) > { > /* For OpenACC, allow (ignore) duplicate struct accesses in > the middle of a mapping clause, e.g. "mystruct->foo" in: > @@ -11174,6 +11203,15 @@ omp_accumulate_sibling_list (enum omp_region_typ= e region_type, > } > } > > + /* If this is an unordered struct, just insert the new element at = the > + end of the list. */ > + if (unordered) > + { > + for (; i < elems; i++) > + sc =3D &OMP_CLAUSE_CHAIN (*sc); > + scp =3D NULL; > + } > + > OMP_CLAUSE_SIZE (*osc) > =3D size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc), size_one_node); > > @@ -11565,14 +11603,42 @@ omp_build_struct_sibling_lists (enum tree_code = code, > > /* This is the first sorted node in the struct sibling list. Use i= t > to recalculate the correct bias to use. > - (&first_node - attach_decl). */ > - tree first_node =3D OMP_CLAUSE_DECL (OMP_CLAUSE_CHAIN (attach)); > - first_node =3D build_fold_addr_expr (first_node); > - first_node =3D fold_convert (ptrdiff_type_node, first_node); > + (&first_node - attach_decl). > + For GOMP_MAP_STRUCT_UNORD, we need e.g. the > + min(min(min(first,second),third),fourth) element, because the > + elements aren't in any particular order. */ > + tree lowest_addr; > + if (OMP_CLAUSE_MAP_KIND (struct_node) =3D=3D GOMP_MAP_STRUCT_UNORD) > + { > + tree first_node =3D OMP_CLAUSE_CHAIN (attach); > + unsigned HOST_WIDE_INT num_mappings > + =3D tree_to_uhwi (OMP_CLAUSE_SIZE (struct_node)); > + lowest_addr =3D OMP_CLAUSE_DECL (first_node); > + lowest_addr =3D build_fold_addr_expr (lowest_addr); > + lowest_addr =3D fold_convert (pointer_sized_int_node, lowest_ad= dr); > + tree next_node =3D OMP_CLAUSE_CHAIN (first_node); > + while (num_mappings > 1) > + { > + tree tmp =3D OMP_CLAUSE_DECL (next_node); > + tmp =3D build_fold_addr_expr (tmp); > + tmp =3D fold_convert (pointer_sized_int_node, tmp); > + lowest_addr =3D fold_build2 (MIN_EXPR, pointer_sized_int_no= de, > + lowest_addr, tmp); > + next_node =3D OMP_CLAUSE_CHAIN (next_node); > + num_mappings--; > + } > + lowest_addr =3D fold_convert (ptrdiff_type_node, lowest_addr); > + } > + else > + { > + tree first_node =3D OMP_CLAUSE_DECL (OMP_CLAUSE_CHAIN (attach))= ; > + first_node =3D build_fold_addr_expr (first_node); > + lowest_addr =3D fold_convert (ptrdiff_type_node, first_node); > + } > tree attach_decl =3D OMP_CLAUSE_DECL (attach); > attach_decl =3D fold_convert (ptrdiff_type_node, attach_decl); > OMP_CLAUSE_SIZE (attach) > - =3D fold_build2 (MINUS_EXPR, ptrdiff_type_node, first_node, > + =3D fold_build2 (MINUS_EXPR, ptrdiff_type_node, lowest_addr, > attach_decl); > > /* Remove GOMP_MAP_ATTACH node from after struct node. */ > @@ -12129,7 +12195,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_s= eq *pre_p, > GOVD_FIRSTPRIVATE | GOVD_SEEN); > } > > - if (OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_STRUCT > + if ((OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_STRUCT > + || OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_STRUCT_UNORD) > && (addr_tokens[0]->type =3D=3D STRUCTURE_BASE > || addr_tokens[0]->type =3D=3D ARRAY_BASE) > && addr_tokens[0]->u.structure_base_kind =3D=3D BASE_DECL) > @@ -13718,7 +13785,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, g= imple_seq body, tree *list_p, > } > } > } > - if (OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_STRUCT > + if ((OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_STRUCT > + || OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_STRUCT_UNORD) > && (code =3D=3D OMP_TARGET_EXIT_DATA || code =3D=3D OACC_EXIT= _DATA)) > { > remove =3D true; > @@ -13762,7 +13830,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, g= imple_seq body, tree *list_p, > in target block and none of the mapping has always modifie= r, > remove all the struct element mappings, which immediately > follow the GOMP_MAP_STRUCT map clause. */ > - if (OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_STRUCT) > + if (OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_STRUCT > + || OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_STRUCT_UNORD) > { > HOST_WIDE_INT cnt =3D tree_to_shwi (OMP_CLAUSE_SIZE (c)); > while (cnt--) > @@ -16541,6 +16610,7 @@ gimplify_omp_target_update (tree *expr_p, gimple_= seq *pre_p) > have_clause =3D false; > break; > case GOMP_MAP_STRUCT: > + case GOMP_MAP_STRUCT_UNORD: > have_clause =3D false; > break; > default: > diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc > index 2b2aa7f13146..3e2c984f8815 100644 > --- a/gcc/omp-low.cc > +++ b/gcc/omp-low.cc > @@ -12811,6 +12811,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, om= p_context *ctx) > case GOMP_MAP_FIRSTPRIVATE_POINTER: > case GOMP_MAP_FIRSTPRIVATE_REFERENCE: > case GOMP_MAP_STRUCT: > + case GOMP_MAP_STRUCT_UNORD: > case GOMP_MAP_ALWAYS_POINTER: > case GOMP_MAP_ATTACH: > case GOMP_MAP_DETACH: > diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc > index 25d191b10fd7..0c1d6722c5ca 100644 > --- a/gcc/tree-pretty-print.cc > +++ b/gcc/tree-pretty-print.cc > @@ -967,6 +967,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int= spc, dump_flags_t flags) > case GOMP_MAP_STRUCT: > pp_string (pp, "struct"); > break; > + case GOMP_MAP_STRUCT_UNORD: > + pp_string (pp, "struct_unord"); > + break; > case GOMP_MAP_ALWAYS_POINTER: > pp_string (pp, "always_pointer"); > break; > diff --git a/include/gomp-constants.h b/include/gomp-constants.h > index 8d4e8e813031..20c722665680 100644 > --- a/include/gomp-constants.h > +++ b/include/gomp-constants.h > @@ -153,6 +153,12 @@ enum gomp_map_kind > (address of the last adjacent entry plus its size). */ > GOMP_MAP_STRUCT =3D (GOMP_MAP_FLAG_SPECIAL_2 > | GOMP_MAP_FLAG_SPECIAL | 0), > + /* As above, but followed by an unordered list of adjacent entries. > + At present, this is used only to diagnose incorrect usage of vari= able > + indices into arrays of structs. */ > + GOMP_MAP_STRUCT_UNORD =3D (GOMP_MAP_FLAG_SPECIAL_4 > + | GOMP_MAP_FLAG_SPECIAL_2 > + | GOMP_MAP_FLAG_SPECIAL | 0), > /* On a location of a pointer/reference that is assumed to be alrea= dy mapped > earlier, store the translated address of the preceeding mapping. > No refcount is bumped by this, and the store is done uncondition= ally. */ > diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c > index fe6327407693..79d6e32c0b4e 100644 > --- a/libgomp/oacc-mem.c > +++ b/libgomp/oacc-mem.c > @@ -1028,6 +1028,7 @@ find_group_last (int pos, size_t mapnum, size_t *si= zes, unsigned short *kinds) > break; > > case GOMP_MAP_STRUCT: > + case GOMP_MAP_STRUCT_UNORD: > pos +=3D sizes[pos]; > break; > > @@ -1088,6 +1089,7 @@ goacc_enter_data_internal (struct gomp_device_descr= *acc_dev, size_t mapnum, > switch (kinds[i] & 0xff) > { > case GOMP_MAP_STRUCT: > + case GOMP_MAP_STRUCT_UNORD: > { > size =3D (uintptr_t) hostaddrs[group_last] + sizes[group_last] > - (uintptr_t) hostaddrs[i]; > @@ -1334,6 +1336,7 @@ goacc_exit_data_internal (struct gomp_device_descr = *acc_dev, size_t mapnum, > break; > > case GOMP_MAP_STRUCT: > + case GOMP_MAP_STRUCT_UNORD: > /* Skip the 'GOMP_MAP_STRUCT' itself, and use the regular process= ing > for all its entries. This special handling exists for GCC 10.= 1 > compatibility; afterwards, we're not generating these no-op > @@ -1472,7 +1475,8 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, = void **hostaddrs, > > if (kind =3D=3D GOMP_MAP_POINTER > || kind =3D=3D GOMP_MAP_TO_PSET > - || kind =3D=3D GOMP_MAP_STRUCT) > + || kind =3D=3D GOMP_MAP_STRUCT > + || kind =3D=3D GOMP_MAP_STRUCT_UNORD) > continue; > > if (kind =3D=3D GOMP_MAP_FORCE_ALLOC > diff --git a/libgomp/target.c b/libgomp/target.c > index b9137e703045..a94fcea154f9 100644 > --- a/libgomp/target.c > +++ b/libgomp/target.c > @@ -1083,7 +1083,8 @@ gomp_map_vars_internal (struct gomp_device_descr *d= evicep, > tgt->list[i].offset =3D 0; > continue; > } > - else if ((kind & typemask) =3D=3D GOMP_MAP_STRUCT) > + else if ((kind & typemask) =3D=3D GOMP_MAP_STRUCT > + || (kind & typemask) =3D=3D GOMP_MAP_STRUCT_UNORD) > { > size_t first =3D i + 1; > size_t last =3D i + sizes[i]; > @@ -1467,6 +1468,20 @@ gomp_map_vars_internal (struct gomp_device_descr *= devicep, > tgt->list[i].offset =3D OFFSET_INLINED; > } > continue; > + case GOMP_MAP_STRUCT_UNORD: > + if (sizes[i] > 1) > + { > + void *first =3D hostaddrs[i + 1]; > + for (size_t j =3D i + 1; j < i + sizes[i]; j++) > + if (hostaddrs[j + 1] !=3D first) > + { > + gomp_mutex_unlock (&devicep->lock); > + gomp_fatal ("Mapped array elements must be the " > + "same (%p vs %p)", first, > + hostaddrs[j + 1]); > + } > + } > + /* Fallthrough. */ > case GOMP_MAP_STRUCT: > first =3D i + 1; > last =3D i + sizes[i]; > @@ -1585,9 +1600,40 @@ gomp_map_vars_internal (struct gomp_device_descr *= devicep, > k->host_end =3D k->host_start + sizeof (void *); > splay_tree_key n =3D splay_tree_lookup (mem_map, k); > if (n && n->refcount !=3D REFCOUNT_LINK) > - gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i], > - kind & typemask, false, implicit, cbu= fp, > - refcount_set); > + { > + if (field_tgt_clear !=3D FIELD_TGT_EMPTY) > + { > + /* For this condition to be true, there must be a > + duplicate struct element mapping. This can happen w= ith > + GOMP_MAP_STRUCT_UNORD mappings, for example. */ > + tgt->list[i].key =3D n; > + if (openmp_p) > + { > + assert ((n->refcount & REFCOUNT_STRUCTELEM) !=3D 0)= ; > + assert (field_tgt_structelem_first !=3D NULL); > + > + if (i =3D=3D field_tgt_clear) > + { > + n->refcount |=3D REFCOUNT_STRUCTELEM_FLAG_LAST; > + field_tgt_structelem_first =3D NULL; > + } > + } > + if (i =3D=3D field_tgt_clear) > + field_tgt_clear =3D FIELD_TGT_EMPTY; > + gomp_increment_refcount (n, refcount_set); > + tgt->list[i].copy_from > + =3D GOMP_MAP_COPY_FROM_P (kind & typemask); > + tgt->list[i].always_copy_from > + =3D GOMP_MAP_ALWAYS_FROM_P (kind & typemask); > + tgt->list[i].is_attach =3D false; > + tgt->list[i].offset =3D 0; > + tgt->list[i].length =3D k->host_end - k->host_start; > + } > + else > + gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i], > + kind & typemask, false, implicit, > + cbufp, refcount_set); > + } > else > { > k->aux =3D NULL; > @@ -4154,7 +4200,8 @@ GOMP_target_enter_exit_data (int device, size_t map= num, void **hostaddrs, > size_t i, j; > if ((flags & GOMP_TARGET_FLAG_EXIT_DATA) =3D=3D 0) > for (i =3D 0; i < mapnum; i++) > - if ((kinds[i] & 0xff) =3D=3D GOMP_MAP_STRUCT) > + if ((kinds[i] & 0xff) =3D=3D GOMP_MAP_STRUCT > + || (kinds[i] & 0xff) =3D=3D GOMP_MAP_STRUCT_UNORD) > { > gomp_map_vars (devicep, sizes[i] + 1, &hostaddrs[i], NULL, &sizes= [i], > &kinds[i], true, &refcount_set, > @@ -4252,7 +4299,8 @@ gomp_target_task_fn (void *data) > htab_t refcount_set =3D htab_create (ttask->mapnum); > if ((ttask->flags & GOMP_TARGET_FLAG_EXIT_DATA) =3D=3D 0) > for (i =3D 0; i < ttask->mapnum; i++) > - if ((ttask->kinds[i] & 0xff) =3D=3D GOMP_MAP_STRUCT) > + if ((ttask->kinds[i] & 0xff) =3D=3D GOMP_MAP_STRUCT > + || (ttask->kinds[i] & 0xff) =3D=3D GOMP_MAP_STRUCT_UNORD) > { > gomp_map_vars (devicep, ttask->sizes[i] + 1, &ttask->hostaddr= s[i], > NULL, &ttask->sizes[i], &ttask->kinds[i], true= , > diff --git a/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-1.c= b/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-1.c > new file mode 100644 > index 000000000000..b0994c0a7bb4 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-1.c > @@ -0,0 +1,38 @@ > +#include > +#include > + > +struct st { > + int *p; > +}; > + > +int main (void) > +{ > + struct st s[2]; > + s[0].p =3D (int *) calloc (5, sizeof (int)); > + s[1].p =3D (int *) calloc (5, sizeof (int)); > + > +#pragma omp target map(s[0].p, s[1].p, s[0].p[0:2], s[1].p[1:3]) > + { > + s[0].p[0] =3D 5; > + s[1].p[1] =3D 7; > + } > + > +#pragma omp target map(s, s[0].p[0:2], s[1].p[1:3]) > + { > + s[0].p[0]++; > + s[1].p[1]++; > + } > + > +#pragma omp target map(s[0:2], s[0].p[0:2], s[1].p[1:3]) > + { > + s[0].p[0]++; > + s[1].p[1]++; > + } > + > + assert (s[0].p[0] =3D=3D 7); > + assert (s[1].p[1] =3D=3D 9); > + > + free (s[0].p); > + free (s[1].p); > + return 0; > +} > diff --git a/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-2.c= b/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-2.c > new file mode 100644 > index 000000000000..81f7efc27c98 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-2.c > @@ -0,0 +1,58 @@ > +#include > +#include > + > +struct st { > + int *p; > +}; > + > +int main (void) > +{ > + struct st s[10]; > + > + for (int i =3D 0; i < 10; i++) > + s[i].p =3D (int *) calloc (5, sizeof (int)); > + > + for (int i =3D 0; i < 10; i++) > + for (int j =3D 0; j < 10; j++) > + for (int k =3D 0; k < 10; k++) > + { > + if (i =3D=3D j || j =3D=3D k || i =3D=3D k) > + continue; > + > +#pragma omp target map(s[i].p, s[j].p, s[k].p, s[i].p[0:2], s[j].p[1:3],= \ > + s[k].p[2]) > + { > + s[i].p[0]++; > + s[j].p[1]++; > + s[k].p[2]++; > + } > + > +#pragma omp target map(s, s[i].p[0:2], s[j].p[1:3], s[k].p[2]) > + { > + s[i].p[0]++; > + s[j].p[1]++; > + s[k].p[2]++; > + } > + > +#pragma omp target map(s[0:10], s[i].p[0:2], s[j].p[1:3], s[k].p[2]) > + { > + s[i].p[0]++; > + s[j].p[1]++; > + s[k].p[2]++; > + } > + } > + > + for (int i =3D 0; i < 10; i++) > + { > + assert (s[i].p[0] =3D=3D 216); > + assert (s[i].p[1] =3D=3D 216); > + assert (s[i].p[2] =3D=3D 216); > + free (s[i].p); > + } > + > + return 0; > +} > + > +/* { dg-output "(\n|\r|\r\n)" } */ > +/* { dg-output "libgomp: Mapped array elements must be the same .*(\n|\r= |\r\n)+" } */ > +/* { dg-shouldfail "" { offload_device_nonshared_as } } */ > diff --git a/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-3.c= b/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-3.c > new file mode 100644 > index 000000000000..639a0d2bc1e3 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-3.c > @@ -0,0 +1,68 @@ > +#include > +#include > + > +struct st { > + int *p; > +}; > + > +struct tt { > + struct st a[10]; > +}; > + > +struct ut { > + struct tt *t; > +}; > + > +int main (void) > +{ > + struct tt *t =3D (struct tt *) malloc (sizeof *t); > + struct ut *u =3D (struct ut *) malloc (sizeof *u); > + > + for (int i =3D 0; i < 10; i++) > + t->a[i].p =3D (int *) calloc (5, sizeof (int)); > + > + u->t =3D t; > + > + for (int i =3D 0; i < 10; i++) > + for (int j =3D 0; j < 10; j++) > + for (int k =3D 0; k < 10; k++) > + { > + if (i =3D=3D j || j =3D=3D k || i =3D=3D k) > + continue; > + > + /* This one can use "firstprivate" for T... */ > +#pragma omp target map(t->a[i].p, t->a[j].p, t->a[k].p, \ > + t->a[i].p[0:2], t->a[j].p[1:3], t->a[k].p[2]) > + { > + t->a[i].p[0]++; > + t->a[j].p[1]++; > + t->a[k].p[2]++; > + } > + > + /* ...but this one must use attach/detach for T. */ > +#pragma omp target map(u->t, u->t->a[i].p, u->t->a[j].p, u->t->a[k].p, \ > + u->t->a[i].p[0:2], u->t->a[j].p[1:3], u->t->a[k].p[2= ]) > + { > + u->t->a[i].p[0]++; > + u->t->a[j].p[1]++; > + u->t->a[k].p[2]++; > + } > + } > + > + for (int i =3D 0; i < 10; i++) > + { > + assert (t->a[i].p[0] =3D=3D 144); > + assert (t->a[i].p[1] =3D=3D 144); > + assert (t->a[i].p[2] =3D=3D 144); > + free (t->a[i].p); > + } > + > + free (u); > + free (t); > + > + return 0; > +} > + > +/* { dg-output "(\n|\r|\r\n)" } */ > +/* { dg-output "libgomp: Mapped array elements must be the same .*(\n|\r= |\r\n)+" } */ > +/* { dg-shouldfail "" { offload_device_nonshared_as } } */ > diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-5.f90 b/libgo= mp/testsuite/libgomp.fortran/map-subarray-5.f90 > new file mode 100644 > index 000000000000..e7cdf11e6108 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.fortran/map-subarray-5.f90 > @@ -0,0 +1,54 @@ > +! { dg-do run } > + > +type t > + integer, pointer :: p(:) > +end type t > + > +type(t) :: var(3) > +integer :: i, j > + > +allocate (var(1)%p, source=3D[1,2,3,5]) > +allocate (var(2)%p, source=3D[2,3,5]) > +allocate (var(3)%p(1:3)) > + > +var(3)%p =3D 0 > + > +do i =3D 1, 3 > + do j =3D 1, 3 > +!$omp target map(var(i)%p, var(j)%p) > + var(i)%p(1) =3D 5 > + var(j)%p(2) =3D 7 > +!$omp end target > + > + if (i.ne.j) then > +!$omp target map(var(i)%p(1:3), var(i)%p, var(j)%p) > + var(i)%p(1) =3D var(i)%p(1) + 1 > + var(j)%p(2) =3D var(j)%p(2) + 1 > +!$omp end target > + > +!$omp target map(var(i)%p, var(j)%p, var(j)%p(1:3)) > + var(i)%p(1) =3D var(i)%p(1) + 1 > + var(j)%p(2) =3D var(j)%p(2) + 1 > +!$omp end target > + > +!$omp target map(var(i)%p, var(i)%p(1:3), var(j)%p, var(j)%p(2)) > + var(i)%p(1) =3D var(i)%p(1) + 1 > + var(j)%p(2) =3D var(j)%p(2) + 1 > +!$omp end target > + end if > + > + if (i.eq.j) then > + if (var(i)%p(1).ne.5) stop 1 > + if (var(j)%p(2).ne.7) stop 2 > + else > + if (var(i)%p(1).ne.8) stop 3 > + if (var(j)%p(2).ne.10) stop 4 > + end if > + end do > +end do > + > +end > + > +! { dg-output "(\n|\r|\r\n)" } > +! { dg-output "libgomp: Mapped array elements must be the same .*(\n|\r|= \r\n)+" } > +! { dg-shouldfail "" { offload_device_nonshared_as } } ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstra=C3=9Fe 201= , 80634 M=C3=BCnchen; Gesellschaft mit beschr=C3=A4nkter Haftung; Gesch=C3= =A4ftsf=C3=BChrer: Thomas Heurung, Frank Th=C3=BCrauf; Sitz der Gesellschaf= t: M=C3=BCnchen; Registergericht M=C3=BCnchen, HRB 106955