Mark pointers to allocated target vars as restricted, if possible

2015-12-02  Tom de Vries  <tom@codesourcery.com>

	* omp-low.c (install_var_field, scan_sharing_clauses): Add and handle
	parameter base_pointers_restrict.
	(omp_target_base_pointers_restrict_p): New function.
	(scan_omp_target): Call scan_sharing_clauses with base_pointers_restrict
	arg.

	* c-c++-common/goacc/kernels-alias-2.c: New test.
	* c-c++-common/goacc/kernels-alias-3.c: New test.
	* c-c++-common/goacc/kernels-alias-4.c: New test.
	* c-c++-common/goacc/kernels-alias-5.c: New test.
	* c-c++-common/goacc/kernels-alias-6.c: New test.
	* c-c++-common/goacc/kernels-alias-7.c: New test.
	* c-c++-common/goacc/kernels-alias-8.c: New test.
	* c-c++-common/goacc/kernels-alias.c: New test.

From-SVN: r231182

gcc/ChangeLog

@@ -1,3 +1,11 @@
+2015-12-02  Tom de Vries  <tom@codesourcery.com>
+
+	* omp-low.c (install_var_field, scan_sharing_clauses): Add and handle
+	parameter base_pointers_restrict.
+	(omp_target_base_pointers_restrict_p): New function.
+	(scan_omp_target): Call scan_sharing_clauses with base_pointers_restrict
+	arg.
+
 2015-12-02  Nathan Sidwell  <nathan@acm.org>
 
 	* config/nvptx/nvptx-protos.h (nvptx_output_mov_insn): Declare.

gcc/omp-low.c

@@ -1366,10 +1366,12 @@ build_sender_ref (tree var, omp_context *ctx)
   return build_sender_ref ((splay_tree_key) var, ctx);
 }
 
-/* Add a new field for VAR inside the structure CTX->SENDER_DECL.  */
+/* Add a new field for VAR inside the structure CTX->SENDER_DECL.  If
+   BASE_POINTERS_RESTRICT, declare the field with restrict.  */
 
 static void
-install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
+install_var_field (tree var, bool by_ref, int mask, omp_context *ctx,
+		   bool base_pointers_restrict = false)
 {
   tree field, type, sfield = NULL_TREE;
   splay_tree_key key = (splay_tree_key) var;
@@ -1393,7 +1395,11 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
       type = build_pointer_type (build_pointer_type (type));
     }
   else if (by_ref)
-    type = build_pointer_type (type);
+    {
+      type = build_pointer_type (type);
+      if (base_pointers_restrict)
+	type = build_qualified_type (type, TYPE_QUAL_RESTRICT);
+    }
   else if ((mask & 3) == 1 && is_reference (var))
     type = TREE_TYPE (type);
 
@@ -1810,10 +1816,12 @@ fixup_child_record_type (omp_context *ctx)
 }
 
 /* Instantiate decls as necessary in CTX to satisfy the data sharing
-   specified by CLAUSES.  */
+   specified by CLAUSES.  If BASE_POINTERS_RESTRICT, install var field with
+   restrict.  */
 
 static void
-scan_sharing_clauses (tree clauses, omp_context *ctx)
+scan_sharing_clauses (tree clauses, omp_context *ctx,
+		      bool base_pointers_restrict = false)
 {
   tree c, decl;
   bool scan_array_reductions = false;
@@ -2075,7 +2083,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 		      && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
 		    install_var_field (decl, true, 7, ctx);
 		  else
-		    install_var_field (decl, true, 3, ctx);
+		    install_var_field (decl, true, 3, ctx,
+				       base_pointers_restrict);
 		  if (is_gimple_omp_offloaded (ctx->stmt))
 		    install_var_local (decl, ctx);
 		}
@@ -3036,6 +3045,68 @@ scan_omp_single (gomp_single *stmt, omp_context *outer_ctx)
     layout_type (ctx->record_type);
 }
 
+/* Return true if the CLAUSES of an omp target guarantee that the base pointers
+   used in the corresponding offloaded function are restrict.  */
+
+static bool
+omp_target_base_pointers_restrict_p (tree clauses)
+{
+  /* The analysis relies on the GOMP_MAP_FORCE_* mapping kinds, which are only
+     used by OpenACC.  */
+  if (flag_openacc == 0)
+    return false;
+
+  /* I.  Basic example:
+
+       void foo (void)
+       {
+	 unsigned int a[2], b[2];
+
+	 #pragma acc kernels \
+	   copyout (a) \
+	   copyout (b)
+	 {
+	   a[0] = 0;
+	   b[0] = 1;
+	 }
+       }
+
+     After gimplification, we have:
+
+       #pragma omp target oacc_kernels \
+	 map(force_from:a [len: 8]) \
+	 map(force_from:b [len: 8])
+       {
+	 a[0] = 0;
+	 b[0] = 1;
+       }
+
+     Because both mappings have the force prefix, we know that they will be
+     allocated when calling the corresponding offloaded function, which means we
+     can mark the base pointers for a and b in the offloaded function as
+     restrict.  */
+
+  tree c;
+  for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+    {
+      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+	return false;
+
+      switch (OMP_CLAUSE_MAP_KIND (c))
+	{
+	case GOMP_MAP_FORCE_ALLOC:
+	case GOMP_MAP_FORCE_TO:
+	case GOMP_MAP_FORCE_FROM:
+	case GOMP_MAP_FORCE_TOFROM:
+	  break;
+	default:
+	  return false;
+	}
+    }
+
+  return true;
+}
+
 /* Scan a GIMPLE_OMP_TARGET.  */
 
 static void
@@ -3057,13 +3128,21 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
   DECL_NAMELESS (name) = 1;
   TYPE_NAME (ctx->record_type) = name;
   TYPE_ARTIFICIAL (ctx->record_type) = 1;
+
+  bool base_pointers_restrict = false;
   if (offloaded)
     {
       create_omp_child_function (ctx, false);
       gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
+
+      base_pointers_restrict = omp_target_base_pointers_restrict_p (clauses);
+      if (base_pointers_restrict
+	  && dump_file && (dump_flags & TDF_DETAILS))
+	fprintf (dump_file,
+		 "Base pointers in offloaded function are restrict\n");
     }
 
-  scan_sharing_clauses (clauses, ctx);
+  scan_sharing_clauses (clauses, ctx, base_pointers_restrict);
   scan_omp (gimple_omp_body_ptr (stmt), ctx);
 
   if (TYPE_FIELDS (ctx->record_type) == NULL)

gcc/testsuite/ChangeLog

@@ -1,3 +1,14 @@
+2015-12-02  Tom de Vries  <tom@codesourcery.com>
+
+	* c-c++-common/goacc/kernels-alias-2.c: New test.
+	* c-c++-common/goacc/kernels-alias-3.c: New test.
+	* c-c++-common/goacc/kernels-alias-4.c: New test.
+	* c-c++-common/goacc/kernels-alias-5.c: New test.
+	* c-c++-common/goacc/kernels-alias-6.c: New test.
+	* c-c++-common/goacc/kernels-alias-7.c: New test.
+	* c-c++-common/goacc/kernels-alias-8.c: New test.
+	* c-c++-common/goacc/kernels-alias.c: New test.
+
 2015-12-02  Tom de Vries  <tom@codesourcery.com>
 
 	* c-c++-common/goacc/kernels-alias-ipa-pta-2.c: New test.

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

@@ -0,0 +1,27 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fdump-tree-ealias-all" } */
+
+void
+foo (void)
+{
+  unsigned int a;
+  unsigned int b;
+  unsigned int c;
+  unsigned int d;
+
+#pragma acc kernels copyin (a) create (b) copyout (c) copy (d)
+  {
+    a = 0;
+    b = 0;
+    c = 0;
+    d = 0;
+  }
+}
+
+/* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } */
+

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

@@ -0,0 +1,20 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fdump-tree-ealias-all" } */
+
+void
+foo (void)
+{
+  unsigned int a;
+  unsigned int *p = &a;
+
+#pragma acc kernels pcopyin (a, p[0:1])
+  {
+    a = 0;
+    *p = 1;
+  }
+}
+
+/* Only the omp_data_i related loads should be annotated with cliques.  */
+/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */
+

gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c

@@ -0,0 +1,22 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fdump-tree-ealias-all" } */
+
+#define N 2
+
+void
+foo (void)
+{
+  unsigned int a[N];
+  unsigned int *p = &a[0];
+
+#pragma acc kernels pcopyin (a, p[0:2])
+  {
+    a[0] = 0;
+    *p = 1;
+  }
+}
+
+/* Only the omp_data_i related loads should be annotated with cliques.  */
+/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */
+

gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c

@@ -0,0 +1,19 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fdump-tree-ealias-all" } */
+
+void
+foo (int *a)
+{
+  int *p = a;
+
+#pragma acc kernels pcopyin (a[0:1], p[0:1])
+  {
+    *a = 0;
+    *p = 1;
+  }
+}
+
+/* Only the omp_data_i related loads should be annotated with cliques.  */
+/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */
+

gcc/testsuite/c-c++-common/goacc/kernels-alias-6.c

@@ -0,0 +1,23 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fdump-tree-ealias-all" } */
+
+typedef __SIZE_TYPE__ size_t;
+extern void *acc_copyin (void *, size_t);
+
+void
+foo (void)
+{
+  int a = 0;
+  int *p = (int *)acc_copyin (&a, sizeof (a));
+
+#pragma acc kernels deviceptr (p) pcopy(a)
+  {
+    a = 0;
+    *p = 1;
+  }
+}
+
+/* Only the omp_data_i related loads should be annotated with cliques.  */
+/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */
+

gcc/testsuite/c-c++-common/goacc/kernels-alias-7.c

@@ -0,0 +1,25 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fdump-tree-ealias-all" } */
+
+typedef __SIZE_TYPE__ size_t;
+extern void *acc_copyin (void *, size_t);
+
+#define N 2
+
+void
+foo (void)
+{
+  int a[N];
+  int *p = (int *)acc_copyin (&a[0], sizeof (a));
+
+#pragma acc kernels deviceptr (p) pcopy(a)
+  {
+    a[0] = 0;
+    *p = 1;
+  }
+}
+
+/* Only the omp_data_i related loads should be annotated with cliques.  */
+/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */
+

gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c

@@ -0,0 +1,22 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fdump-tree-ealias-all" } */
+
+typedef __SIZE_TYPE__ size_t;
+extern void *acc_copyin (void *, size_t);
+
+void
+foo (int *a, size_t n)
+{
+  int *p = (int *)acc_copyin (&a, n);
+
+#pragma acc kernels deviceptr (p) pcopy(a[0:n])
+  {
+    a = 0;
+    *p = 1;
+  }
+}
+
+/* Only the omp_data_i related loads should be annotated with cliques.  */
+/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 2 "ealias" } } */
+

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

@@ -0,0 +1,29 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fdump-tree-ealias-all" } */
+
+#define N 2
+
+void
+foo (void)
+{
+  unsigned int a[N];
+  unsigned int b[N];
+  unsigned int c[N];
+  unsigned int d[N];
+
+#pragma acc kernels copyin (a) create (b) copyout (c) copy (d)
+  {
+    a[0] = 0;
+    b[0] = 0;
+    c[0] = 0;
+    d[0] = 0;
+  }
+}
+
+/* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } */
+