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.

(cherry picked from openacc-gcc-9-branch commit
99a4260fe264624136eda666363464eba48bdc3e)

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 3c516287c92e82a6354858b669ee5996ee3d3abf..f23ad4ed2202994cf21f3741b5f46dc24493d15c 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.c (nvptx_propagate_unified): New.

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index a17e54aa7935b57a068077cc3c40422e6a89c390..ca37ddfb66e4527446b768147f35bbc69956c1ea 100644 (file)
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -668,7 +668,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;
@@ -699,7 +700,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 && omp_is_reference (var))
     type = TREE_TYPE (type);
 
@@ -1155,10 +1160,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;
@@ -1505,7 +1512,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)
                      && !OMP_CLAUSE_MAP_IN_REDUCTION (c))
                    install_var_local (decl, ctx);
@@ -2612,6 +2620,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 dynamic array map clauses are placed at the very
    front of the chain.  */
 
@@ -2664,6 +2734,7 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
   omp_context *ctx;
   tree name;
   bool offloaded = is_gimple_omp_offloaded (stmt);
+  tree clauses = gimple_omp_target_clauses (stmt);
 
   ctx = new_omp_context (stmt, outer_ctx);
   ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
@@ -2676,21 +2747,28 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
   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");
     }
 
   /* If OpenACC construct, put dynamic array clauses (if any) in front of
      clause chain. The runtime can then test the first to see if the
      additional map processing for them is required.  */
   if (is_gimple_omp_oacc (stmt))
-    reorder_dynamic_array_clauses (gimple_omp_target_clauses_ptr (stmt));
+    {
+      reorder_dynamic_array_clauses (gimple_omp_target_clauses_ptr (stmt));
+      clauses = gimple_omp_target_clauses (stmt);
+    }
 
-  tree clauses = gimple_omp_target_clauses (stmt);
-  
-  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)