diff options
author | Tom de Vries <tom@codesourcery.com> | 2015-12-02 13:32:51 +0000 |
---|---|---|
committer | Tom de Vries <vries@gcc.gnu.org> | 2015-12-02 13:32:51 +0000 |
commit | 694e5e4baebff76320ecbb0c119bc086126c4095 (patch) | |
tree | 6b472c95b371572ac65f8248f47d8aac5477666b /gcc | |
parent | 7dbf36f7cb76a7e2bb8a8658f608ee659fe07cb9 (diff) | |
download | gcc-694e5e4baebff76320ecbb0c119bc086126c4095.zip gcc-694e5e4baebff76320ecbb0c119bc086126c4095.tar.gz gcc-694e5e4baebff76320ecbb0c119bc086126c4095.tar.bz2 |
Handle BUILT_IN_GOACC_PARALLEL in ipa-pta
2015-12-02 Tom de Vries <tom@codesourcery.com>
* 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
Diffstat (limited to 'gcc')
-rw-r--r-- | gcc/ChangeLog | 5 | ||||
-rw-r--r-- | gcc/testsuite/ChangeLog | 6 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-2.c | 37 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c | 36 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta.c | 23 | ||||
-rw-r--r-- | gcc/tree-ssa-structalias.c | 28 |
6 files changed, 131 insertions, 4 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 58c4728..2c30b30 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2015-12-02 Tom de Vries <tom@codesourcery.com> + + * 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 <segher@kernel.crashing.org> * config/rs6000/rs6000.md (cstore_si_as_di): New expander. diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index f7bb329..05022b3 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,9 @@ +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. + 2015-12-02 Richard Biener <rguenther@suse.de> * 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 0000000..f16d698 --- /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 0000000..1eb56eb --- /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 0000000..969b466 --- /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 7f4a8ad..060ff3e 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); |