From 694e5e4baebff76320ecbb0c119bc086126c4095 Mon Sep 17 00:00:00 2001 From: Tom de Vries Date: Wed, 2 Dec 2015 13:32:51 +0000 Subject: [PATCH] Handle BUILT_IN_GOACC_PARALLEL in ipa-pta 2015-12-02 Tom de Vries * tree-ssa-structalias.c (find_func_aliases_for_builtin_call) (find_func_clobbers, ipa_pta_execute): Handle BUILT_IN_GOACC_PARALLEL. * 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. From-SVN: r231169 --- gcc/ChangeLog | 5 +++ gcc/testsuite/ChangeLog | 6 +++ .../goacc/kernels-alias-ipa-pta-2.c | 37 +++++++++++++++++++ .../goacc/kernels-alias-ipa-pta-3.c | 36 ++++++++++++++++++ .../goacc/kernels-alias-ipa-pta.c | 23 ++++++++++++ gcc/tree-ssa-structalias.c | 28 ++++++++++++-- .../kernels-alias-ipa-pta-2.c | 27 ++++++++++++++ .../kernels-alias-ipa-pta-3.c | 26 +++++++++++++ .../kernels-alias-ipa-pta.c | 26 +++++++++++++ 9 files changed, 210 insertions(+), 4 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-2.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-3.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 58c472868ed..2c30b300006 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2015-12-02 Tom de Vries + + * tree-ssa-structalias.c (find_func_aliases_for_builtin_call) + (find_func_clobbers, ipa_pta_execute): Handle BUILT_IN_GOACC_PARALLEL. + 2015-12-02 Segher Boessenkool * config/rs6000/rs6000.md (cstore_si_as_di): New expander. diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index f7bb329eed7..05022b3b470 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,9 @@ +2015-12-02 Tom de Vries + + * 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. + 2015-12-02 Richard Biener * gcc.dg/vect/vect-strided-a-u8-i8-gap7-big-array.c: Fix uninitialized diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-2.c new file mode 100644 index 00000000000..f16d698af0d --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-2.c @@ -0,0 +1,37 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fipa-pta -fdump-tree-optimized" } */ + +#ifdef __cplusplus +extern "C" { +#endif +typedef __SIZE_TYPE__ size_t; +void *malloc (size_t); +void free (void *); +#ifdef __cplusplus +} +#endif + +#define N 2 + +void +foo (void) +{ + unsigned int *a = (unsigned int *)malloc (N * sizeof (unsigned int)); + unsigned int *b = (unsigned int *)malloc (N * sizeof (unsigned int)); + unsigned int *c = (unsigned int *)malloc (N * sizeof (unsigned int)); + +#pragma acc kernels pcopyout (a[0:N], b[0:N], c[0:N]) + { + a[0] = 0; + b[0] = 1; + c[0] = a[0]; + } + + free (a); + free (b); + free (c); +} + +/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 2 "optimized" } } */ +/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" } } */ +/* { dg-final { scan-tree-dump-times "(?n)= \\*a" 0 "optimized" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c new file mode 100644 index 00000000000..1eb56eb9e62 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c @@ -0,0 +1,36 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fipa-pta -fdump-tree-optimized" } */ + +#ifdef __cplusplus +extern "C" { +#endif +typedef __SIZE_TYPE__ size_t; +void *malloc (size_t); +void free (void *); +#ifdef __cplusplus +} +#endif + +#define N 2 + +void +foo (void) +{ + unsigned int *a = (unsigned int *)malloc (N * sizeof (unsigned int)); + unsigned int *b = a; + unsigned int *c = (unsigned int *)malloc (N * sizeof (unsigned int)); + +#pragma acc kernels pcopyout (a[0:N], b[0:N], c[0:N]) + { + a[0] = 0; + b[0] = 1; + c[0] = a[0]; + } + + free (a); + free (c); +} + +/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 1 "optimized" } } */ +/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" } } */ +/* { dg-final { scan-tree-dump-times "(?n)= \\*a" 1 "optimized" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta.c new file mode 100644 index 00000000000..969b466e8a8 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta.c @@ -0,0 +1,23 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fipa-pta -fdump-tree-optimized" } */ + +#define N 2 + +void +foo (void) +{ + unsigned int a[N]; + unsigned int b[N]; + unsigned int c[N]; + +#pragma acc kernels pcopyout (a, b, c) + { + a[0] = 0; + b[0] = 1; + c[0] = a[0]; + } +} + +/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 2 "optimized" } } */ +/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" } } */ +/* { dg-final { scan-tree-dump-times "(?n)= \\*_\[0-9\]\\\[0\\\];$" 0 "optimized" } } */ diff --git a/gcc/tree-ssa-structalias.c b/gcc/tree-ssa-structalias.c index 7f4a8ad60e4..060ff3efc1d 100644 --- a/gcc/tree-ssa-structalias.c +++ b/gcc/tree-ssa-structalias.c @@ -4507,15 +4507,32 @@ find_func_aliases_for_builtin_call (struct function *fn, gcall *t) return true; } case BUILT_IN_GOMP_PARALLEL: + case BUILT_IN_GOACC_PARALLEL: { - /* Handle __builtin_GOMP_parallel (fn, data, num_threads, flags) as - fn (data). */ if (in_ipa_mode) { - tree fnarg = gimple_call_arg (t, 0); + unsigned int fnpos, argpos; + switch (DECL_FUNCTION_CODE (fndecl)) + { + case BUILT_IN_GOMP_PARALLEL: + /* __builtin_GOMP_parallel (fn, data, num_threads, flags). */ + fnpos = 0; + argpos = 1; + break; + case BUILT_IN_GOACC_PARALLEL: + /* __builtin_GOACC_parallel (device, fn, mapnum, hostaddrs, + sizes, kinds, ...). */ + fnpos = 1; + argpos = 3; + break; + default: + gcc_unreachable (); + } + + tree fnarg = gimple_call_arg (t, fnpos); gcc_assert (TREE_CODE (fnarg) == ADDR_EXPR); tree fndecl = TREE_OPERAND (fnarg, 0); - tree arg = gimple_call_arg (t, 1); + tree arg = gimple_call_arg (t, argpos); gcc_assert (TREE_CODE (arg) == ADDR_EXPR); varinfo_t fi = get_vi_for_tree (fndecl); @@ -5064,6 +5081,7 @@ find_func_clobbers (struct function *fn, gimple *origt) case BUILT_IN_VA_END: return; case BUILT_IN_GOMP_PARALLEL: + case BUILT_IN_GOACC_PARALLEL: return; /* printf-style functions may have hooks to set pointers to point to somewhere into the generated string. Leave them @@ -7547,6 +7565,8 @@ ipa_pta_execute (void) /* Handle direct calls to functions with body. */ if (gimple_call_builtin_p (stmt, BUILT_IN_GOMP_PARALLEL)) decl = TREE_OPERAND (gimple_call_arg (stmt, 0), 0); + else if (gimple_call_builtin_p (stmt, BUILT_IN_GOACC_PARALLEL)) + decl = TREE_OPERAND (gimple_call_arg (stmt, 1), 0); else decl = gimple_call_fndecl (stmt); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-2.c new file mode 100644 index 00000000000..0f323c82d2c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-2.c @@ -0,0 +1,27 @@ +/* { dg-additional-options "-O2 -fipa-pta" } */ + +#include + +#define N 2 + +int +main (void) +{ + unsigned int *a = (unsigned int *)malloc (N * sizeof (unsigned int)); + unsigned int *b = (unsigned int *)malloc (N * sizeof (unsigned int)); + unsigned int *c = (unsigned int *)malloc (N * sizeof (unsigned int)); + +#pragma acc kernels pcopyout (a[0:N], b[0:N], c[0:N]) + { + a[0] = 0; + b[0] = 1; + c[0] = a[0]; + } + + if (a[0] != 0 || b[0] != 1 || c[0] != 0) + abort (); + + free (a); + free (b); + free (c); +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-3.c new file mode 100644 index 00000000000..654e750c3b4 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-3.c @@ -0,0 +1,26 @@ +/* { dg-additional-options "-O2 -fipa-pta" } */ + +#include + +#define N 2 + +int +main (void) +{ + unsigned int *a = (unsigned int *)malloc (N * sizeof (unsigned int)); + unsigned int *b = a; + unsigned int *c = (unsigned int *)malloc (N * sizeof (unsigned int)); + +#pragma acc kernels pcopyout (a[0:N], b[0:N], c[0:N]) + { + a[0] = 0; + b[0] = 1; + c[0] = a[0]; + } + + if (a[0] != 1 || b[0] != 1 || c[0] != 1) + abort (); + + free (a); + free (c); +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta.c new file mode 100644 index 00000000000..44d4fd25cd2 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta.c @@ -0,0 +1,26 @@ +/* { dg-additional-options "-O2 -fipa-pta" } */ + +#include + +#define N 2 + +int +main (void) +{ + unsigned int a[N]; + unsigned int b[N]; + unsigned int c[N]; + +#pragma acc kernels pcopyout (a, b, c) + { + a[0] = 0; + b[0] = 1; + c[0] = a[0]; + } + + if (a[0] != 0 || b[0] != 1 || c[0] != 0) + abort (); + + return 0; +} + -- 2.30.2