@@ -1249,6 +1249,7 @@ extern tree c_omp_check_context_selector (location_t, tree);
extern void c_omp_mark_declare_variant (location_t, tree, tree);
extern const char *c_omp_map_clause_name (tree, bool);
extern void c_omp_adjust_map_clauses (tree, bool);
+extern tree c_omp_decompose_attachable_address (tree t, tree *virtbase);
enum c_omp_directive_kind {
C_OMP_DIR_STANDALONE,
@@ -3119,6 +3119,48 @@ c_omp_adjust_map_clauses (tree clauses, bool is_target)
}
}
+tree
+c_omp_decompose_attachable_address (tree t, tree *virtbase)
+{
+ *virtbase = t;
+
+ /* It's already a pointer. Just use that. */
+ if (POINTER_TYPE_P (TREE_TYPE (t)))
+ return NULL_TREE;
+
+ /* Otherwise, look for a base pointer deeper within the expression. */
+
+ while (TREE_CODE (t) == COMPONENT_REF
+ && (TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF
+ || TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF))
+ {
+ t = TREE_OPERAND (t, 0);
+ while (TREE_CODE (t) == ARRAY_REF)
+ t = TREE_OPERAND (t, 0);
+ }
+
+
+ *virtbase = t;
+
+ if (TREE_CODE (t) != COMPONENT_REF)
+ return NULL_TREE;
+
+ t = TREE_OPERAND (t, 0);
+
+ tree attach_pt = NULL_TREE;
+
+ if ((TREE_CODE (t) == INDIRECT_REF
+ || TREE_CODE (t) == MEM_REF)
+ && TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == POINTER_TYPE)
+ {
+ attach_pt = TREE_OPERAND (t, 0);
+ if (TREE_CODE (attach_pt) == POINTER_PLUS_EXPR)
+ attach_pt = TREE_OPERAND (attach_pt, 0);
+ }
+
+ return attach_pt;
+}
+
static const struct c_omp_directive omp_directives[] = {
/* Keep this alphabetically sorted by the first word. Non-null second/third
if any should precede null ones. */
@@ -13696,9 +13696,15 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
if (size)
size = c_fully_fold (size, false, NULL);
OMP_CLAUSE_SIZE (c) = size;
+ tree virtbase = t;
+ tree attach_pt
+ = ((ort != C_ORT_ACC)
+ ? c_omp_decompose_attachable_address (t, &virtbase)
+ : NULL_TREE);
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (TREE_CODE (t) == COMPONENT_REF
- && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE
+ && !attach_pt))
return false;
gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR);
switch (OMP_CLAUSE_MAP_KIND (c))
@@ -13731,10 +13737,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& !c_mark_addressable (t))
return false;
- OMP_CLAUSE_DECL (c2) = t;
+ OMP_CLAUSE_DECL (c2) = attach_pt ? attach_pt : t;
t = build_fold_addr_expr (first);
t = fold_convert_loc (OMP_CLAUSE_LOCATION (c), ptrdiff_type_node, t);
- tree ptr = OMP_CLAUSE_DECL (c2);
+ tree ptr = virtbase;
if (!POINTER_TYPE_P (TREE_TYPE (ptr)))
ptr = build_fold_addr_expr (ptr);
t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
@@ -5608,9 +5608,16 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
OMP_CLAUSE_SIZE (c) = size;
if (TREE_CODE (t) == FIELD_DECL)
t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
+
+ tree virtbase = t;
+ tree attach_pt
+ = ((ort != C_ORT_ACC)
+ ? c_omp_decompose_attachable_address (t, &virtbase)
+ : NULL_TREE);
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (TREE_CODE (t) == COMPONENT_REF
- && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE
+ && !attach_pt))
return false;
switch (OMP_CLAUSE_MAP_KIND (c))
{
@@ -5670,12 +5677,11 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& !cxx_mark_addressable (t))
return false;
- OMP_CLAUSE_DECL (c2) = t;
+ OMP_CLAUSE_DECL (c2) = attach_pt ? attach_pt : t;
t = build_fold_addr_expr (first);
t = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
ptrdiff_type_node, t);
- tree ptr = OMP_CLAUSE_DECL (c2);
- ptr = convert_from_reference (ptr);
+ tree ptr = convert_from_reference (virtbase);
if (!INDIRECT_TYPE_P (TREE_TYPE (ptr)))
ptr = build_fold_addr_expr (ptr);
t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
new file mode 100644
@@ -0,0 +1,182 @@
+/* { dg-xfail-if "fails to parse correctly" { *-*-* } } */
+
+#include <cstdlib>
+#include <cstring>
+#include <cassert>
+
+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 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->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);
+
+/* FIXME: This currently ICEs. */
+/* #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);
+
+ 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->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);
+
+/* FIXME: This currently ICEs. */
+/* #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);
+
+ delete my_a->ptr;
+ delete my_a;
+ delete my_b;
+}
+
+int main (int argc, char *argv[])
+{
+ foo ();
+ bar ();
+ foop ();
+ barp ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,50 @@
+#include <stdlib.h>
+#include <string.h>
+#include <assert.h>
+#include <stdio.h>
+
+#define N 32
+
+typedef struct {
+ int x2[10][N];
+} x1type;
+
+typedef struct {
+ x1type x1[10];
+} p2type;
+
+typedef struct {
+ p2type *p2;
+} p1type;
+
+typedef struct {
+ p1type *p1;
+} x0type;
+
+typedef struct {
+ x0type x0[10];
+} p0type;
+
+int main(int argc, char *argv[])
+{
+ p0type *p0;
+ int k1 = 0, k2 = 0, k3 = 0, n = N;
+
+ p0 = (p0type *) malloc (sizeof *p0);
+ p0->x0[0].p1 = (p1type *) malloc (sizeof *p0->x0[0].p1);
+ p0->x0[0].p1->p2 = (p2type *) malloc (sizeof *p0->x0[0].p1->p2);
+ memset (p0->x0[0].p1->p2, 0, sizeof *p0->x0[0].p1->p2);
+
+#pragma omp target map(tofrom: p0->x0[k1].p1->p2[k2].x1[k3].x2[4][0:n]) \
+ map(to: p0->x0[k1].p1, p0->x0[k1].p1->p2) \
+ map(to: p0->x0[k1].p1[0])
+ {
+ for (int i = 0; i < n; i++)
+ p0->x0[k1].p1->p2[k2].x1[k3].x2[4][i] = i;
+ }
+
+ for (int i = 0; i < n; i++)
+ assert (i == p0->x0[k1].p1->p2[k2].x1[k3].x2[4][i]);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,70 @@
+#include <stdlib.h>
+#include <string.h>
+#include <assert.h>
+
+#define N 32
+
+typedef struct {
+ int arr[N];
+ int *ptr;
+} sc;
+
+typedef struct {
+ sc *c;
+} sb;
+
+typedef struct {
+ sb *b;
+ sc *c;
+} sa;
+
+int main (int argc, char *argv[])
+{
+ sa *p;
+
+ p = (sa *) malloc (sizeof *p);
+ p->b = (sb *) malloc (sizeof *p->b);
+ p->b->c = (sc *) malloc (sizeof *p->b->c);
+ p->c = (sc *) malloc (sizeof *p->c);
+ p->b->c->ptr = (int *) malloc (N * sizeof (int));
+ p->c->ptr = (int *) malloc (N * sizeof (int));
+
+ for (int i = 0; i < N; i++)
+ {
+ p->b->c->ptr[i] = 0;
+ p->c->ptr[i] = 0;
+ p->b->c->arr[i] = 0;
+ p->c->arr[i] = 0;
+ }
+
+#pragma omp target map(to: p->b, p->b[0], p->c, p->c[0], p->b->c, p->b->c[0]) \
+ map(to: p->b->c->ptr, p->c->ptr) \
+ map(tofrom: p->b->c->ptr[:N], p->c->ptr[:N])
+ {
+ for (int i = 0; i < N; i++)
+ {
+ p->b->c->ptr[i] = i;
+ p->c->ptr[i] = i * 2;
+ }
+ }
+
+#pragma omp target map(to: p->b, p->b[0], p->b->c, p->c) \
+ map(tofrom: p->c[0], p->b->c[0])
+ {
+ for (int i = 0; i < N; i++)
+ {
+ p->b->c->arr[i] = i * 3;
+ p->c->arr[i] = i * 4;
+ }
+ }
+
+ for (int i = 0; i < N; i++)
+ {
+ assert (p->b->c->ptr[i] == i);
+ assert (p->c->ptr[i] == i * 2);
+ assert (p->b->c->arr[i] == i * 3);
+ assert (p->c->arr[i] == i * 4);
+ }
+
+ return 0;
+}