diff options
| -rw-r--r-- | gcc/ChangeLog | 8 | ||||
| -rw-r--r-- | gcc/omp-low.c | 93 | ||||
| -rw-r--r-- | gcc/testsuite/ChangeLog | 11 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c | 27 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c | 20 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c | 22 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c | 19 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/kernels-alias-6.c | 23 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/kernels-alias-7.c | 25 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c | 22 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/kernels-alias.c | 29 |
11 files changed, 292 insertions, 7 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 990aab8..0e4a3dd 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,11 @@ +2015-12-02 Tom de Vries <tom@codesourcery.com> + + * omp-low.c (install_var_field, scan_sharing_clauses): Add and handle + parameter base_pointers_restrict. + (omp_target_base_pointers_restrict_p): New function. + (scan_omp_target): Call scan_sharing_clauses with base_pointers_restrict + arg. + 2015-12-02 Nathan Sidwell <nathan@acm.org> * config/nvptx/nvptx-protos.h (nvptx_output_mov_insn): Declare. diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 15cc839..d1d1e3c 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -1366,10 +1366,12 @@ build_sender_ref (tree var, omp_context *ctx) return build_sender_ref ((splay_tree_key) var, ctx); } -/* Add a new field for VAR inside the structure CTX->SENDER_DECL. */ +/* Add a new field for VAR inside the structure CTX->SENDER_DECL. If + BASE_POINTERS_RESTRICT, declare the field with restrict. */ static void -install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) +install_var_field (tree var, bool by_ref, int mask, omp_context *ctx, + bool base_pointers_restrict = false) { tree field, type, sfield = NULL_TREE; splay_tree_key key = (splay_tree_key) var; @@ -1393,7 +1395,11 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) type = build_pointer_type (build_pointer_type (type)); } else if (by_ref) - type = build_pointer_type (type); + { + type = build_pointer_type (type); + if (base_pointers_restrict) + type = build_qualified_type (type, TYPE_QUAL_RESTRICT); + } else if ((mask & 3) == 1 && is_reference (var)) type = TREE_TYPE (type); @@ -1810,10 +1816,12 @@ fixup_child_record_type (omp_context *ctx) } /* Instantiate decls as necessary in CTX to satisfy the data sharing - specified by CLAUSES. */ + specified by CLAUSES. If BASE_POINTERS_RESTRICT, install var field with + restrict. */ static void -scan_sharing_clauses (tree clauses, omp_context *ctx) +scan_sharing_clauses (tree clauses, omp_context *ctx, + bool base_pointers_restrict = false) { tree c, decl; bool scan_array_reductions = false; @@ -2075,7 +2083,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) install_var_field (decl, true, 7, ctx); else - install_var_field (decl, true, 3, ctx); + install_var_field (decl, true, 3, ctx, + base_pointers_restrict); if (is_gimple_omp_offloaded (ctx->stmt)) install_var_local (decl, ctx); } @@ -3036,6 +3045,68 @@ scan_omp_single (gomp_single *stmt, omp_context *outer_ctx) layout_type (ctx->record_type); } +/* Return true if the CLAUSES of an omp target guarantee that the base pointers + used in the corresponding offloaded function are restrict. */ + +static bool +omp_target_base_pointers_restrict_p (tree clauses) +{ + /* The analysis relies on the GOMP_MAP_FORCE_* mapping kinds, which are only + used by OpenACC. */ + if (flag_openacc == 0) + return false; + + /* I. Basic example: + + void foo (void) + { + unsigned int a[2], b[2]; + + #pragma acc kernels \ + copyout (a) \ + copyout (b) + { + a[0] = 0; + b[0] = 1; + } + } + + After gimplification, we have: + + #pragma omp target oacc_kernels \ + map(force_from:a [len: 8]) \ + map(force_from:b [len: 8]) + { + a[0] = 0; + b[0] = 1; + } + + Because both mappings have the force prefix, we know that they will be + allocated when calling the corresponding offloaded function, which means we + can mark the base pointers for a and b in the offloaded function as + restrict. */ + + tree c; + for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) + return false; + + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TOFROM: + break; + default: + return false; + } + } + + return true; +} + /* Scan a GIMPLE_OMP_TARGET. */ static void @@ -3057,13 +3128,21 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx) DECL_NAMELESS (name) = 1; TYPE_NAME (ctx->record_type) = name; TYPE_ARTIFICIAL (ctx->record_type) = 1; + + bool base_pointers_restrict = false; if (offloaded) { create_omp_child_function (ctx, false); gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn); + + base_pointers_restrict = omp_target_base_pointers_restrict_p (clauses); + if (base_pointers_restrict + && dump_file && (dump_flags & TDF_DETAILS)) + fprintf (dump_file, + "Base pointers in offloaded function are restrict\n"); } - scan_sharing_clauses (clauses, ctx); + scan_sharing_clauses (clauses, ctx, base_pointers_restrict); scan_omp (gimple_omp_body_ptr (stmt), ctx); if (TYPE_FIELDS (ctx->record_type) == NULL) diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 05022b3..bb523db 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,5 +1,16 @@ 2015-12-02 Tom de Vries <tom@codesourcery.com> + * c-c++-common/goacc/kernels-alias-2.c: New test. + * c-c++-common/goacc/kernels-alias-3.c: New test. + * c-c++-common/goacc/kernels-alias-4.c: New test. + * c-c++-common/goacc/kernels-alias-5.c: New test. + * c-c++-common/goacc/kernels-alias-6.c: New test. + * c-c++-common/goacc/kernels-alias-7.c: New test. + * c-c++-common/goacc/kernels-alias-8.c: New test. + * c-c++-common/goacc/kernels-alias.c: New test. + +2015-12-02 Tom de Vries <tom@codesourcery.com> + * c-c++-common/goacc/kernels-alias-ipa-pta-2.c: New test. * c-c++-common/goacc/kernels-alias-ipa-pta-3.c: New test. * c-c++-common/goacc/kernels-alias-ipa-pta.c: New test. diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c new file mode 100644 index 0000000..d437c47 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c @@ -0,0 +1,27 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +void +foo (void) +{ + unsigned int a; + unsigned int b; + unsigned int c; + unsigned int d; + +#pragma acc kernels copyin (a) create (b) copyout (c) copy (d) + { + a = 0; + b = 0; + c = 0; + d = 0; + } +} + +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } */ + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c new file mode 100644 index 0000000..0eda7e1 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c @@ -0,0 +1,20 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +void +foo (void) +{ + unsigned int a; + unsigned int *p = &a; + +#pragma acc kernels pcopyin (a, p[0:1]) + { + a = 0; + *p = 1; + } +} + +/* Only the omp_data_i related loads should be annotated with cliques. */ +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */ + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c new file mode 100644 index 0000000..037901f --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c @@ -0,0 +1,22 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +#define N 2 + +void +foo (void) +{ + unsigned int a[N]; + unsigned int *p = &a[0]; + +#pragma acc kernels pcopyin (a, p[0:2]) + { + a[0] = 0; + *p = 1; + } +} + +/* Only the omp_data_i related loads should be annotated with cliques. */ +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */ + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c new file mode 100644 index 0000000..69cd3fb --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c @@ -0,0 +1,19 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +void +foo (int *a) +{ + int *p = a; + +#pragma acc kernels pcopyin (a[0:1], p[0:1]) + { + *a = 0; + *p = 1; + } +} + +/* Only the omp_data_i related loads should be annotated with cliques. */ +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */ + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-6.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-6.c new file mode 100644 index 0000000..6ebce15 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-6.c @@ -0,0 +1,23 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +typedef __SIZE_TYPE__ size_t; +extern void *acc_copyin (void *, size_t); + +void +foo (void) +{ + int a = 0; + int *p = (int *)acc_copyin (&a, sizeof (a)); + +#pragma acc kernels deviceptr (p) pcopy(a) + { + a = 0; + *p = 1; + } +} + +/* Only the omp_data_i related loads should be annotated with cliques. */ +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */ + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-7.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-7.c new file mode 100644 index 0000000..40eb235 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-7.c @@ -0,0 +1,25 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +typedef __SIZE_TYPE__ size_t; +extern void *acc_copyin (void *, size_t); + +#define N 2 + +void +foo (void) +{ + int a[N]; + int *p = (int *)acc_copyin (&a[0], sizeof (a)); + +#pragma acc kernels deviceptr (p) pcopy(a) + { + a[0] = 0; + *p = 1; + } +} + +/* Only the omp_data_i related loads should be annotated with cliques. */ +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */ + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c new file mode 100644 index 0000000..0b93e35 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c @@ -0,0 +1,22 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +typedef __SIZE_TYPE__ size_t; +extern void *acc_copyin (void *, size_t); + +void +foo (int *a, size_t n) +{ + int *p = (int *)acc_copyin (&a, n); + +#pragma acc kernels deviceptr (p) pcopy(a[0:n]) + { + a = 0; + *p = 1; + } +} + +/* Only the omp_data_i related loads should be annotated with cliques. */ +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */ + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias.c new file mode 100644 index 0000000..25821ab2 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias.c @@ -0,0 +1,29 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-ealias-all" } */ + +#define N 2 + +void +foo (void) +{ + unsigned int a[N]; + unsigned int b[N]; + unsigned int c[N]; + unsigned int d[N]; + +#pragma acc kernels copyin (a) create (b) copyout (c) copy (d) + { + a[0] = 0; + b[0] = 0; + c[0] = 0; + d[0] = 0; + } +} + +/* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } */ + |