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

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 58c472868ed9c7d1b7d3ead40595cc72c57ca121..2c30b300006a52b6131482e50148c7e46f8fdaa5 100644 (file)
--- 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 f7bb329eed73db95176edbaaebb0d71f322f4d90..05022b3b470c8e4df0ecee5b5059dab392eb7fc6 100644 (file)
--- 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 (file)
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 (file)
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 (file)
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 7f4a8ad60e486e0f41975901ca7e446e26021c32..060ff3efc1d71e7bac339fae917bfe95f7036806 100644 (file)
--- 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 (file)
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 (file)
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 (file)
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;
+}
+