Reinstate kernels-restrict behaviour

This patch contains a small fix for upstream churn relative to the last version
posted.

2018-09-20  Cesar Philippidis  <cesar@codesourcery.com>
	    Julian Brown  <julian@codesourcery.com>

	* omp-low.c (install_var_field): New base_pointer_restrict
	argument.
	(scan_sharing_clauses): Update call to install_var_field.
	(omp_target_base_pointers_restrict_p): New function.
	(scan_omp_target): Update call to install_var_field.

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index d1a9bcef5235b734b7eb3e9d36e84f6d06161c6c..6c0729986fa16c760489d3326bfc8c93d061ffd3 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,12 @@
+2018-09-20  Cesar Philippidis  <cesar@codesourcery.com>
+           Julian Brown  <julian@codesourcery.com>
+
+       * omp-low.c (install_var_field): New base_pointer_restrict
+       argument.
+       (scan_sharing_clauses): Update call to install_var_field.
+       (omp_target_base_pointers_restrict_p): New function.
+       (scan_omp_target): Update call to install_var_field.
+
 2018-10-30  Cesar Philippidis  <cesar@codesourcery.com>
 
        * config/nvptx/nvptx.cc (nvptx_propagate_unified): New.

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 70a6dc3c3995ee6669f2331cde732909817c9e0c..edc22e25e91db458b72a892a281f1bf31ae4b121 100644 (file)
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -777,7 +777,8 @@ build_sender_ref (tree var, omp_context *ctx)
    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;
@@ -816,7 +817,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 & (32 | 3)) == 1
           && omp_privatize_by_reference (var))
     type = TREE_TYPE (type);
@@ -1256,10 +1261,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;
@@ -1796,7 +1803,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)
                      && !(is_gimple_omp_oacc (ctx->stmt)
                           && OMP_CLAUSE_MAP_IN_REDUCTION (c)))
@@ -3197,6 +3205,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;
+}
+
 /* Reorder clauses so that non-contiguous array map clauses are placed at the very
    front of the chain.  */
 
@@ -3269,13 +3339,19 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
 
   tree clauses = gimple_omp_target_clauses (stmt);
 
+  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)