[ was: Re: [PATCH, PR46032] Handle BUILT_IN_GOMP_PARALLEL in ipa-pta ]
On 30/11/15 17:36, Tom de Vries wrote:
On 30/11/15 14:24, Richard Biener wrote:
On Mon, 30 Nov 2015, Tom de Vries wrote:
On 30/11/15 10:16, Richard Biener wrote:
On Mon, 30 Nov 2015, Tom de Vries wrote:
Hi,
this patch fixes PR46032.
It handles a call:
...
__builtin_GOMP_parallel (fn, data, num_threads, flags)
...
as:
...
fn (data)
...
in ipa-pta.
This improves ipa-pta alias analysis in the parallelized function
fn,
This follow-up patch does the same for BUILT_IN_GOACC_PARALLEL.
Bootstrapped and reg-tested on x86_64.
OK for stage3 trunk?
Thanks,
- Tom
Handle BUILT_IN_GOACC_PARALLEL in ipa-pta
2015-12-01 Tom de Vries <t...@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.
---
.../c-c++-common/goacc/kernels-alias-ipa-pta-2.c | 37 ++++++++++++++++++++++
.../c-c++-common/goacc/kernels-alias-ipa-pta-3.c | 36 +++++++++++++++++++++
.../c-c++-common/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 +++++++++++++++
7 files changed, 199 insertions(+), 4 deletions(-)
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);
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 0000000..0f323c8
--- /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 <stdlib.h>
+
+#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 0000000..654e750
--- /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 <stdlib.h>
+
+#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 0000000..44d4fd2
--- /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 <stdlib.h>
+
+#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;
+}
+
