@@ -9813,7 +9813,10 @@ accumulate_sibling_list (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;
if (!struct_map_to_clause || struct_map_to_clause->get (base) == NULL)
@@ -9862,9 +9865,37 @@ accumulate_sibling_list (enum omp_region_type region_type, enum tree_code code,
tree noind = strip_indirections (base);
- if (!openmp
+ /* TODO: the following two stanzas handling reference-typed struct
+ members (for OpenMP) and nested base pointers (for OpenACC) could
+ probably both be better handled in the frontends. Doing that would
+ avoid this late manipulation of the clause list. */
+
+ 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
@@ -9878,13 +9909,28 @@ accumulate_sibling_list (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;
+ 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) == 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);
@@ -10228,11 +10274,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.deleted = false;
new file mode 100644
@@ -0,0 +1,275 @@
+#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;
+};
+
+struct sb
+{
+ int arr[10];
+};
+
+struct sc
+{
+ sa &a;
+ sb &b;
+ sc (sa &my_a, sb &my_b) : a(my_a), b(my_b) {}
+};
+
+void
+foo ()
+{
+ sa my_a;
+ sb my_b;
+
+ my_a.ptr = (int *) malloc (sizeof (int) * 10);
+ sc 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);
+}
+
+void
+bar ()
+{
+ sa my_a;
+ sb my_b;
+
+ my_a.ptr = (int *) malloc (sizeof (int) * 10);
+ sc my_c(my_a, my_b);
+ sc &my_cref = my_c;
+
+ memset (my_cref.a.ptr, 0, sizeof (int) * 10);
+
+ #pragma omp target map (my_cref.a.ptr, my_cref.a.ptr[:10])
+ {
+ for (int i = 0; i < 10; i++)
+ my_cref.a.ptr[i] = i;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_cref.a.ptr[i] == i);
+
+ memset (my_cref.b.arr, 0, sizeof (int) * 10);
+
+ #pragma omp target map (my_cref.b.arr[:10])
+ {
+ for (int i = 0; i < 10; i++)
+ my_cref.b.arr[i] = i;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_cref.b.arr[i] == i);
+
+ 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;
+ sb *&b;
+ scp (sa *&my_a, sb *&my_b) : a(my_a), b(my_b) {}
+};
+
+void
+foop ()
+{
+ sa *my_a = new sa;
+ sb *my_b = new sb;
+
+ my_a->ptr = new int[10];
+ scp *my_c = new scp(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;
+}
+
+void
+barp ()
+{
+ sa *my_a = new sa;
+ sb *my_b = new sb;
+
+ my_a->ptr = new int[10];
+ scp *my_c = new scp(my_a, my_b);
+ scp *&my_cref = my_c;
+
+ memset (my_cref->a->ptr, 0, sizeof (int) * 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;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_cref->a->ptr[i] == i);
+
+ memset (my_cref->b->arr, 0, sizeof (int) * 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;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_cref->b->arr[i] == i);
+
+ delete my_a->ptr;
+ delete my_a;
+ delete my_b;
+}
+
+int main (int argc, char *argv[])
+{
+ foo0 ();
+ foo ();
+ bar ();
+ foop0 ();
+ foop ();
+ barp ();
+ return 0;
+}