+2019-01-31 Julian Brown <julian@codesourcery.com>
+
+ * gimplify.c (gimplify_scan_omp_clauses): Handle array sections on
+ dereferenced struct members.
+
2018-12-14 Julian Brown <julian@codesourcery.com>
* gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_HAS_ATTACHMENTS.
+2019-01-31 Julian Brown <julian@codesourcery.com>
+
+ * c-typeck.c (handle_omp_array_sections_1): Handle chained dereferences.
+ (c_finish_omp_clauses): Likewise.
+
2018-12-14 Julian Brown <julian@codesourcery.com>
* c-parser.c (c_parser_omp_clause_name): Add parsing of attach and
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
- }
- if (ort == C_ORT_ACC && 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,
- c_omp_map_clause_name (c, true));
- else
- t = TREE_OPERAND (t, 0);
+ if (ort == C_ORT_ACC && 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);
+ }
}
}
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
break;
}
t = TREE_OPERAND (t, 0);
+ if (ort == C_ORT_ACC && 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);
+ }
}
if (remove)
break;
- if (ort == C_ORT_ACC && 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,
- c_omp_map_clause_name (c, true));
- else
- t = TREE_OPERAND (t, 0);
- }
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+2019-01-31 Julian Brown <julian@codesourcery.com>
+
+ * semantics.c (handle_omp_array_sections_1): Handle array section on
+ dereferenced struct member.
+ (finish_omp_clauses): Don't error on multiple dereferenced struct
+ elements with the same base.
+
2018-12-14 Julian Brown <julian@codesourcery.com>
* parser.c (cp_parser_omp_clause_name): Support attach and detach
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
+ if (ort == C_ORT_ACC && TREE_CODE (t) == INDIRECT_REF)
+ t = TREE_OPERAND (t, 0);
}
if (REFERENCE_REF_P (t))
t = TREE_OPERAND (t, 0);
else
bitmap_set_bit (&generic_head, DECL_UID (t));
}
- else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ else if (bitmap_bit_p (&map_head, DECL_UID (t))
+ && (ort != C_ORT_ACC
+ || !bitmap_bit_p (&map_field_head, DECL_UID (t))))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error_at (OMP_CLAUSE_LOCATION (c),
struct gimplify_omp_ctx *ctx, *outer_ctx;
tree c;
hash_map<tree, tree> *struct_map_to_clause = NULL;
+ hash_set<tree> *struct_deref_set = NULL;
tree *prev_list_p = NULL;
int handled_depend_iterators = -1;
int nowait = -1;
pd = &TREE_OPERAND (decl, 0);
decl = TREE_OPERAND (decl, 0);
}
- if (TREE_CODE (decl) == COMPONENT_REF)
+ bool indir_p = false;
+ tree orig_decl = decl;
+ tree decl_ref = NULL_TREE;
+ if ((region_type & ORT_ACC) != 0
+ && TREE_CODE (*pd) == COMPONENT_REF
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER
+ && code != OACC_UPDATE)
+ {
+ while (TREE_CODE (decl) == COMPONENT_REF)
+ {
+ decl = TREE_OPERAND (decl, 0);
+ if ((TREE_CODE (decl) == MEM_REF
+ && integer_zerop (TREE_OPERAND (decl, 1)))
+ || INDIRECT_REF_P (decl))
+ {
+ indir_p = true;
+ decl = TREE_OPERAND (decl, 0);
+ }
+ if (TREE_CODE (decl) == INDIRECT_REF
+ && DECL_P (TREE_OPERAND (decl, 0))
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == REFERENCE_TYPE))
+ {
+ decl_ref = decl;
+ decl = TREE_OPERAND (decl, 0);
+ }
+ }
+ }
+ else if (TREE_CODE (decl) == COMPONENT_REF)
{
while (TREE_CODE (decl) == COMPONENT_REF)
decl = TREE_OPERAND (decl, 0);
== REFERENCE_TYPE))
decl = TREE_OPERAND (decl, 0);
}
+ if (decl != orig_decl && DECL_P (decl) && indir_p)
+ {
+ gomp_map_kind k = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
+ : GOMP_MAP_ATTACH;
+ /* We have a dereference of a struct member. Make this an
+ attach/detach operation, and ensure the base pointer is
+ mapped as a FIRSTPRIVATE_POINTER. */
+ OMP_CLAUSE_SET_MAP_KIND (c, k);
+ flags = GOVD_MAP | GOVD_SEEN | GOVD_EXPLICIT;
+ if (k == GOMP_MAP_ATTACH
+ && (!struct_deref_set
+ || !struct_deref_set->contains (decl)))
+ {
+ if (!struct_deref_set)
+ struct_deref_set = new hash_set<tree> ();
+ /* As well as the attach, we also need a
+ FIRSTPRIVATE_POINTER clause to properly map the
+ pointer to the struct base. */
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC);
+ OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c2)
+ = 1;
+ tree charptr_zero
+ = build_int_cst (build_pointer_type (char_type_node),
+ 0);
+ OMP_CLAUSE_DECL (c2)
+ = build2 (MEM_REF, char_type_node,
+ decl_ref ? decl_ref : decl, charptr_zero);
+ OMP_CLAUSE_SIZE (c2) = size_zero_node;
+ tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c3,
+ GOMP_MAP_FIRSTPRIVATE_POINTER);
+ OMP_CLAUSE_DECL (c3) = decl;
+ OMP_CLAUSE_SIZE (c3) = size_zero_node;
+ tree mapgrp = *prev_list_p;
+ *prev_list_p = c2;
+ OMP_CLAUSE_CHAIN (c3) = mapgrp;
+ OMP_CLAUSE_CHAIN (c2) = c3;
+
+ struct_deref_set->add (decl);
+ }
+ goto do_add_decl;
+ }
if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
== GS_ERROR)
{
gimplify_omp_ctxp = ctx;
if (struct_map_to_clause)
delete struct_map_to_clause;
+ if (struct_deref_set)
+ delete struct_deref_set;
}
/* Return true if DECL is a candidate for shared to firstprivate
+2019-01-31 Julian Brown <julian@codesourcery.com>
+
+ * c-c++-common/goacc/deep-copy-arrayofstruct.c: New test.
+
2018-12-14 Julian Brown <julian@codesourcery.com>
* c-c++-common/goacc/mdc-1.c: New test.
--- /dev/null
+/* { dg-do compile } */
+
+#include <stdlib.h>
+#include <stdio.h>
+
+typedef struct {
+ int *a;
+ int *b;
+ int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+ const int N = 1024;
+ const int S = 32;
+ mystruct *m = (mystruct *) calloc (S, sizeof (*m));
+ int i, j;
+
+ for (i = 0; i < S; i++)
+ {
+ m[i].a = (int *) malloc (N * sizeof (int));
+ m[i].b = (int *) malloc (N * sizeof (int));
+ m[i].c = (int *) malloc (N * sizeof (int));
+ }
+
+ for (j = 0; j < S; j++)
+ for (i = 0; i < N; i++)
+ {
+ m[j].a[i] = 0;
+ m[j].b[i] = 0;
+ m[j].c[i] = 0;
+ }
+
+#pragma acc enter data copyin(m[0:1])
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j, k;
+ for (k = 0; k < S; k++)
+#pragma acc parallel loop copy(m[k].a[0:N]) /* { dg-error "expected .\\\). before .\\\.. token" } */
+ 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 } */
+ for (j = 0; j < N; j++)
+ {
+ m[k].b[j]++;
+ if (j > 5 && j < N - 5)
+ m[k].c[j]++;
+ }
+ }
+
+#pragma acc exit data copyout(m[0:1])
+
+ for (j = 0; j < S; j++)
+ {
+ for (i = 0; i < N; i++)
+ {
+ if (m[j].a[i] != 99)
+ abort ();
+ if (m[j].b[i] != 99)
+ abort ();
+ if (i > 5 && i < N-5)
+ {
+ if (m[j].c[i] != 99)
+ abort ();
+ }
+ else
+ {
+ if (m[j].c[i] != 0)
+ abort ();
+ }
+ }
+
+ free (m[j].a);
+ free (m[j].b);
+ free (m[j].c);
+ }
+ free (m);
+
+ return 0;
+}
--- /dev/null
+module bar
+ type :: type1
+ real(8), pointer, public :: p(:) => null()
+ end type
+ type :: type2
+ class(type1), pointer :: p => null()
+ end type
+end module
+
+subroutine foo (var)
+ use bar
+ type(type2), intent(inout) :: var
+ !$acc enter data create(var%p%p)
+end subroutine
+2019-01-31 Julian Brown <julian@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c++/deep-copy-12.C: New test.
+ * testsuite/libgomp.oacc-c++/deep-copy-13.C: New test.
+ * testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c: New test.
+
2018-11-28 Julian Brown <julian@codesourcery.com>
* libgomp.h (RC_CHECKING): New macro, disabled by default, guarding all
--- /dev/null
+#include <stdlib.h>
+
+/* Test attach/detach with dereferences of reference to pointer to struct. */
+
+typedef struct {
+ int *a;
+ int *b;
+ int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+ const int N = 1024;
+ mystruct *m = (mystruct *) malloc (sizeof (*m));
+ mystruct *&mref = m;
+ int i;
+
+ mref->a = (int *) malloc (N * sizeof (int));
+ m->b = (int *) malloc (N * sizeof (int));
+ m->c = (int *) malloc (N * sizeof (int));
+
+ for (i = 0; i < N; i++)
+ {
+ mref->a[i] = 0;
+ m->b[i] = 0;
+ m->c[i] = 0;
+ }
+
+#pragma acc enter data copyin(m[0:1])
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j;
+#pragma acc parallel loop copy(mref->a[0:N])
+ for (j = 0; j < N; j++)
+ mref->a[j]++;
+#pragma acc parallel loop copy(mref->b[0:N], m->c[5:N-10])
+ for (j = 0; j < N; j++)
+ {
+ mref->b[j]++;
+ if (j > 5 && j < N - 5)
+ m->c[j]++;
+ }
+ }
+
+#pragma acc exit data copyout(m[0:1])
+
+ for (i = 0; i < N; i++)
+ {
+ if (m->a[i] != 99)
+ abort ();
+ if (m->b[i] != 99)
+ abort ();
+ if (i > 5 && i < N-5)
+ {
+ if (m->c[i] != 99)
+ abort ();
+ }
+ else
+ {
+ if (m->c[i] != 0)
+ abort ();
+ }
+ }
+
+ free (m->a);
+ free (m->b);
+ free (m->c);
+ free (m);
+
+ return 0;
+}
--- /dev/null
+#include <stdlib.h>
+
+/* Test array slice with reference to pointer. */
+
+typedef struct {
+ int *a;
+ int *b;
+ int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+ const int N = 1024;
+ mystruct *m = (mystruct *) malloc (sizeof (*m));
+ int i;
+
+ m->a = (int *) malloc (N * sizeof (int));
+ m->b = (int *) malloc (N * sizeof (int));
+ m->c = (int *) malloc (N * sizeof (int));
+
+ for (i = 0; i < N; i++)
+ {
+ m->a[i] = 0;
+ m->b[i] = 0;
+ m->c[i] = 0;
+ }
+
+#pragma acc enter data copyin(m[0:1])
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j;
+ int *&ptr = m->a;
+#pragma acc parallel loop copy(ptr[0:N])
+ for (j = 0; j < N; j++)
+ ptr[j]++;
+#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10])
+ for (j = 0; j < N; j++)
+ {
+ m->b[j]++;
+ if (j > 5 && j < N - 5)
+ m->c[j]++;
+ }
+ }
+
+#pragma acc exit data copyout(m[0:1])
+
+ for (i = 0; i < N; i++)
+ {
+ if (m->a[i] != 99)
+ abort ();
+ if (m->b[i] != 99)
+ abort ();
+ if (i > 5 && i < N-5)
+ {
+ if (m->c[i] != 99)
+ abort ();
+ }
+ else
+ {
+ if (m->c[i] != 0)
+ abort ();
+ }
+ }
+
+ free (m->a);
+ free (m->b);
+ free (m->c);
+ free (m);
+
+ return 0;
+}
--- /dev/null
+#include <stdlib.h>
+
+/* Test asyncronous attach and detach operation. */
+
+typedef struct {
+ int *a;
+ int *b;
+} mystruct;
+
+int
+main (int argc, char* argv[])
+{
+ const int N = 1024;
+ mystruct m;
+ int i;
+
+ m.a = (int *) malloc (N * sizeof (int));
+ m.b = (int *) malloc (N * sizeof (int));
+
+ for (i = 0; i < N; i++)
+ {
+ m.a[i] = 0;
+ m.b[i] = 0;
+ }
+
+#pragma acc enter data copyin(m)
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j;
+#pragma acc parallel loop copy(m.a[0:N]) async(i % 2)
+ for (j = 0; j < N; j++)
+ m.a[j]++;
+#pragma acc parallel loop copy(m.b[0:N]) async((i + 1) % 2)
+ for (j = 0; j < N; j++)
+ m.b[j]++;
+ }
+
+#pragma acc exit data copyout(m) wait(0, 1)
+
+ for (i = 0; i < N; i++)
+ {
+ if (m.a[i] != 99)
+ abort ();
+ if (m.b[i] != 99)
+ abort ();
+ }
+
+ free (m.a);
+ free (m.b);
+
+ return 0;
+}
--- /dev/null
+#include <stdlib.h>
+
+/* Test multiple struct dereferences on one directive, and slices starting at
+ non-zero. */
+
+typedef struct {
+ int *a;
+ int *b;
+ int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+ const int N = 1024;
+ mystruct *m = (mystruct *) malloc (sizeof (*m));
+ int i;
+
+ m->a = (int *) malloc (N * sizeof (int));
+ m->b = (int *) malloc (N * sizeof (int));
+ m->c = (int *) malloc (N * sizeof (int));
+
+ for (i = 0; i < N; i++)
+ {
+ m->a[i] = 0;
+ m->b[i] = 0;
+ m->c[i] = 0;
+ }
+
+#pragma acc enter data copyin(m[0:1])
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j;
+#pragma acc parallel loop copy(m->a[0:N])
+ for (j = 0; j < N; j++)
+ m->a[j]++;
+#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10])
+ for (j = 0; j < N; j++)
+ {
+ m->b[j]++;
+ if (j > 5 && j < N - 5)
+ m->c[j]++;
+ }
+ }
+
+#pragma acc exit data copyout(m[0:1])
+
+ for (i = 0; i < N; i++)
+ {
+ if (m->a[i] != 99)
+ abort ();
+ if (m->b[i] != 99)
+ abort ();
+ if (i > 5 && i < N-5)
+ {
+ if (m->c[i] != 99)
+ abort ();
+ }
+ else
+ {
+ if (m->c[i] != 0)
+ abort ();
+ }
+ }
+
+ free (m->a);
+ free (m->b);
+ free (m->c);
+ free (m);
+
+ return 0;
+}
--- /dev/null
+#include <openacc.h>
+#include <stdlib.h>
+
+/* Test attach/detach operation with chained dereferences. */
+
+typedef struct mystruct {
+ int *a;
+ struct mystruct *next;
+} mystruct;
+
+int
+main (int argc, char* argv[])
+{
+ const int N = 1024;
+ mystruct *m = (mystruct *) malloc (sizeof (*m));
+ int i;
+
+ m->a = (int *) malloc (N * sizeof (int));
+ m->next = (mystruct *) malloc (sizeof (*m));
+ m->next->a = (int *) malloc (N * sizeof (int));
+ m->next->next = NULL;
+
+ for (i = 0; i < N; i++)
+ {
+ m->a[i] = 0;
+ m->next->a[i] = 0;
+ }
+
+#pragma acc enter data copyin(m[0:1])
+ acc_copyin (m->next, sizeof (*m));
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j;
+ acc_copyin (m->next->a, N * sizeof (int));
+ acc_attach ((void **) &m->next);
+ /* This will attach only the innermost pointer, i.e. "a[0:N]". That's
+ why we have to attach the "m->next" pointer manually above. */
+#pragma acc parallel loop copy(m->next->a[0:N])
+ for (j = 0; j < N; j++)
+ m->next->a[j]++;
+ acc_detach ((void **) &m->next);
+ acc_copyout (m->next->a, N * sizeof (int));
+ }
+
+ acc_copyout (m->next, sizeof (*m));
+#pragma acc exit data copyout(m[0:1])
+
+ for (i = 0; i < N; i++)
+ {
+ if (m->a[i] != 0)
+ abort ();
+ if (m->next->a[i] != 99)
+ abort ();
+ }
+
+ free (m->next->a);
+ free (m->next);
+ free (m->a);
+ free (m);
+
+ return 0;
+}
--- /dev/null
+#include <stdlib.h>
+
+typedef struct {
+ int *a;
+ int *b;
+} mystruct;
+
+int
+main (int argc, char* argv[])
+{
+ const int N = 1024;
+ mystruct *m = (mystruct *) malloc (sizeof (*m));
+ int i;
+
+ m->a = (int *) malloc (N * sizeof (int));
+ m->b = (int *) malloc (N * sizeof (int));
+
+ for (i = 0; i < N; i++)
+ {
+ m->a[i] = 0;
+ m->b[i] = 0;
+ }
+
+#pragma acc enter data copyin(m[0:1])
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j;
+ int *ptr = m->a;
+#pragma acc parallel loop copy(m->a[0:N])
+ for (j = 0; j < N; j++)
+ m->a[j]++;
+#pragma acc parallel loop copy(m->b[0:N])
+ for (j = 0; j < N; j++)
+ m->b[j]++;
+ }
+
+#pragma acc exit data copyout(m[0:1])
+
+ for (i = 0; i < N; i++)
+ {
+ if (m->a[i] != 99)
+ abort ();
+ if (m->b[i] != 99)
+ abort ();
+ }
+
+ free (m->a);
+ free (m->b);
+ free (m);
+
+ return 0;
+}
projects / thirdparty / gcc.git / commitdiff
