This patch fixes the baseptrs-3.C test case introduced in the patch:
https://gcc.gnu.org/pipermail/gcc-patches/2021-October/580729.html
The problematic case concerns OpenMP mapping clauses containing struct
members of reference type, e.g. "mystruct.myref.myptr[:N]". To be able
to access the array slice through the reference in the middle, we need
to perform an attach action for that reference, since it is represented
internally as a pointer.
I don't think the spec allows for this case explicitly. The closest
clause is (OpenMP 5.0, "2.19.7.1 map Clause"):
"If the type of a list item is a reference to a type T then the
reference in the device data environment is initialized to refer to
the object in the device data environment that corresponds to the
object referenced by the list item. If mapping occurs, it occurs as
though the object were mapped through a pointer with an array section
of type T and length one."
The patch as is allows the mapping to work with just
"mystruct.myref.myptr[:N]", without an explicit "mystruct.myref"
mapping also (because, would that refer to the hidden pointer used by
the reference, or the automatically-dereferenced data itself?). An
attach/detach operation is thus synthesised for the reference.
Tested with offloading to NVPTX and bootstrapped. OK (pending
previously-posted series?).
Julian Brown
2021-10-11 Julian Brown <julian@codesourcery.com>
gcc/cp/
* semantics.c (finish_omp_clauses): Handle reference-typed members.
gcc/
* gimplify.c (build_struct_group): Arrange for attach/detach nodes to
be created for reference-typed struct members for OpenMP. Only create
firstprivate_pointer/firstprivate_reference nodes for innermost struct
accesses, those with an optionally-indirected DECL_P base.
(omp_build_struct_sibling_lists): Handle two-element chain for inner
struct component returned from build_struct_group.
libgomp/
* testsuite/libgomp.c++/baseptrs-3.C: Remove XFAILs and extend test.
---
gcc/cp/semantics.c | 4 +
gcc/gimplify.c | 56 +++++++++--
libgomp/testsuite/libgomp.c++/baseptrs-3.C | 109 +++++++++++++++++++--
3 files changed, 154 insertions(+), 15 deletions(-)
@@ -7862,6 +7862,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
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);
}
}
while (TREE_CODE (t) == COMPONENT_REF);
@@ -7961,6 +7963,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
{
t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
indir_component_ref_p = true;
+ if (REFERENCE_REF_P (t))
+ t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
@@ -10249,7 +10249,10 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code,
/* FIXME: If we're not mapping the base pointer in some other clause on this
directive, I think we want to create ALLOC/RELEASE here -- i.e. not
early-exit. */
- if (openmp && attach_detach)
+ if (openmp
+ && attach_detach
+ && !(TREE_CODE (TREE_TYPE (ocd)) == REFERENCE_TYPE
+ && TREE_CODE (TREE_TYPE (TREE_TYPE (ocd))) != POINTER_TYPE))
return NULL;
#ifdef NOISY_SIBLING_LISTS
@@ -10317,9 +10320,32 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code,
tree noind = strip_indirections (base);
- if (!openmp
+ if (openmp
+ && TREE_CODE (TREE_TYPE (noind)) == REFERENCE_TYPE
&& (region_type & ORT_TARGET)
&& TREE_CODE (noind) == COMPONENT_REF)
+ {
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_TO);
+ OMP_CLAUSE_DECL (c2) = unshare_expr (base);
+ OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (TREE_TYPE (noind));
+
+ tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH_DETACH);
+ OMP_CLAUSE_DECL (c3) = unshare_expr (noind);
+ OMP_CLAUSE_SIZE (c3) = size_zero_node;
+
+ OMP_CLAUSE_CHAIN (c2) = c3;
+ OMP_CLAUSE_CHAIN (c3) = NULL_TREE;
+
+ *inner = c2;
+ return NULL;
+ }
+ else if (!openmp
+ && (region_type & ORT_TARGET)
+ && TREE_CODE (noind) == COMPONENT_REF)
{
/* The base for this component access is a struct component access
itself. Insert a node to be processed on the next iteration of
@@ -10333,13 +10359,30 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code,
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FORCE_PRESENT);
OMP_CLAUSE_DECL (c2) = unshare_expr (noind);
OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (TREE_TYPE (noind));
+ OMP_CLAUSE_CHAIN (c2) = NULL_TREE;
*inner = c2;
return NULL;
}
- tree sdecl = strip_components_and_deref (base);
+ tree sdecl = base;
+ /* There are too many places we need to do things like this. */
+ if (TREE_CODE (sdecl) == INDIRECT_REF
+ || TREE_CODE (sdecl) == MEM_REF)
+ {
+ sdecl = TREE_OPERAND (sdecl, 0);
+ if (TREE_CODE (sdecl) == INDIRECT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (sdecl, 0)))
+ == REFERENCE_TYPE))
+ sdecl = TREE_OPERAND (sdecl, 0);
+ }
- if (POINTER_TYPE_P (TREE_TYPE (sdecl)) && (region_type & ORT_TARGET))
+ while (TREE_CODE (sdecl) == COMPONENT_REF
+ || TREE_CODE (sdecl) == POINTER_PLUS_EXPR)
+ sdecl = TREE_OPERAND (sdecl, 0);
+
+ if (DECL_P (sdecl)
+ && POINTER_TYPE_P (TREE_TYPE (sdecl))
+ && (region_type & ORT_TARGET))
{
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end),
OMP_CLAUSE_MAP);
@@ -10740,11 +10783,10 @@ omp_build_struct_sibling_lists (enum tree_code code,
else
*tail = inner;
- OMP_CLAUSE_CHAIN (inner) = NULL_TREE;
-
omp_mapping_group newgrp;
newgrp.grp_start = new_next ? new_next : tail;
- newgrp.grp_end = inner;
+ newgrp.grp_end = (OMP_CLAUSE_CHAIN (inner)
+ ? OMP_CLAUSE_CHAIN (inner) : inner);
newgrp.mark = UNVISITED;
newgrp.sibling = NULL;
newgrp.next = NULL;
@@ -1,9 +1,58 @@
-/* { dg-xfail-if "fails to parse correctly" { *-*-* } } */
-
#include <cstdlib>
#include <cstring>
#include <cassert>
+struct sa0
+{
+ int *ptr;
+};
+
+struct sb0
+{
+ int arr[10];
+};
+
+struct sc0
+{
+ sa0 a;
+ sb0 b;
+ sc0 (sa0 &my_a, sb0 &my_b) : a(my_a), b(my_b) {}
+};
+
+void
+foo0 ()
+{
+ sa0 my_a;
+ sb0 my_b;
+
+ my_a.ptr = (int *) malloc (sizeof (int) * 10);
+ sc0 my_c(my_a, my_b);
+
+ memset (my_c.a.ptr, 0, sizeof (int) * 10);
+
+ #pragma omp target map (my_c.a.ptr, my_c.a.ptr[:10])
+ {
+ for (int i = 0; i < 10; i++)
+ my_c.a.ptr[i] = i;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_c.a.ptr[i] == i);
+
+ memset (my_c.b.arr, 0, sizeof (int) * 10);
+
+ #pragma omp target map (my_c.b.arr[:10])
+ {
+ for (int i = 0; i < 10; i++)
+ my_c.b.arr[i] = i;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_c.b.arr[i] == i);
+
+ free (my_a.ptr);
+}
+
struct sa
{
int *ptr;
@@ -90,6 +139,49 @@ bar ()
free (my_a.ptr);
}
+struct scp0
+{
+ sa *a;
+ sb *b;
+ scp0 (sa *my_a, sb *my_b) : a(my_a), b(my_b) {}
+};
+
+void
+foop0 ()
+{
+ sa *my_a = new sa;
+ sb *my_b = new sb;
+
+ my_a->ptr = new int[10];
+ scp0 *my_c = new scp0(my_a, my_b);
+
+ memset (my_c->a->ptr, 0, sizeof (int) * 10);
+
+ #pragma omp target map (my_c->a, my_c->a[:1], my_c->a->ptr, my_c->a->ptr[:10])
+ {
+ for (int i = 0; i < 10; i++)
+ my_c->a->ptr[i] = i;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_c->a->ptr[i] == i);
+
+ memset (my_c->b->arr, 0, sizeof (int) * 10);
+
+ #pragma omp target map (my_c->b, my_c->b[:1], my_c->b->arr[:10])
+ {
+ for (int i = 0; i < 10; i++)
+ my_c->b->arr[i] = i;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_c->b->arr[i] == i);
+
+ delete[] my_a->ptr;
+ delete my_a;
+ delete my_b;
+}
+
struct scp
{
sa *&a;
@@ -108,7 +200,7 @@ foop ()
memset (my_c->a->ptr, 0, sizeof (int) * 10);
- #pragma omp target map (my_c->a->ptr, my_c->a->ptr[:10])
+ #pragma omp target map (my_c->a, my_c->a[:1], my_c->a->ptr, my_c->a->ptr[:10])
{
for (int i = 0; i < 10; i++)
my_c->a->ptr[i] = i;
@@ -119,8 +211,7 @@ foop ()
memset (my_c->b->arr, 0, sizeof (int) * 10);
-/* FIXME: This currently ICEs. */
-/* #pragma omp target map (my_c->b->arr[:10]) */
+ #pragma omp target map (my_c->b, my_c->b[:1], my_c->b->arr[:10])
{
for (int i = 0; i < 10; i++)
my_c->b->arr[i] = i;
@@ -146,7 +237,8 @@ barp ()
memset (my_cref->a->ptr, 0, sizeof (int) * 10);
- #pragma omp target map (my_cref->a->ptr, my_cref->a->ptr[:10])
+ #pragma omp target map (my_cref->a, my_cref->a[:1], my_cref->a->ptr, \
+ my_cref->a->ptr[:10])
{
for (int i = 0; i < 10; i++)
my_cref->a->ptr[i] = i;
@@ -157,8 +249,7 @@ barp ()
memset (my_cref->b->arr, 0, sizeof (int) * 10);
-/* FIXME: This currently ICEs. */
-/* #pragma omp target map (my_cref->b->arr[:10]) */
+ #pragma omp target map (my_cref->b, my_cref->b[:1], my_cref->b->arr[:10])
{
for (int i = 0; i < 10; i++)
my_cref->b->arr[i] = i;
@@ -174,8 +265,10 @@ barp ()
int main (int argc, char *argv[])
{
+ foo0 ();
foo ();
bar ();
+ foop0 ();
foop ();
barp ();
return 0;