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

gcc/ChangeLog

+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.

gcc/testsuite/ChangeLog

+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

gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-2.c

+/* { 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" } } */

gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c

+/* { 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" } } */

gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta.c

+/* { 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" } } */

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);
 

libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-2.c

+/* { 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);
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-3.c

+/* { 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);
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta.c

+/* { 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;
+}
+