Re: [PATCH, 4/16] Implement -foffload-alias

[Tom de Vries <Tom_deVries at mentor dot com>]

On 13/11/15 12:39, Jakub Jelinek wrote:
> On Fri, Nov 13, 2015 at 12:29:51PM +0100, Richard Biener wrote:
> > > thanks for the explanation. Filed as PR68331 - '[meta-bug] fipa-pta issues'.
> > >
> > > Any feedback on the '#pragma GCC offload-alias=<none|pointer|all>' bit above?
> > > Is that sort of what you had in mind?
> >
> > Yes.  Whether that makes sense is another question of course.  You can
> > annotate memory references with MR_DEPENDENCE_BASE/CLIQUE yourself
> > as well if you know dependences without the users intervention.
>
> I really don't like even the GCC offload-alias, I just don't see anything
> special on the offload code.  Not to mention that the same issue is already
> with other outlined functions, like OpenMP tasks or parallel regions, those
> aren't offloaded, yet they can suffer from worse alias/points-to analysis
> too.

AFAIU there is one aspect that is different for offloaded code: the 
setup of the data on the device.

Consider this example:
...
unsigned int a[N];
unsigned int b[N];
unsigned int c[N];

int
main (void)
{
  ...

#pragma acc kernels copyin (a) copyin (b) copyout (c)
  {
    for (COUNTERTYPE ii = 0; ii < N; ii++)
      c[ii] = a[ii] + b[ii];
  }

  ...
...

At gimple level, we have:
...
#pragma omp target oacc_kernels \
  map(force_from:c [len: 2097152]) \
  map(force_to:b [len: 2097152]) \
  map(force_to:a [len: 2097152])
...

[ The meaning of the force_from/force_to mappings is given in 
include/gomp-constants.h:
...
    /* Allocate.  */
    GOMP_MAP_FORCE_ALLOC = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_ALLOC),
    /* ..., and copy to device.  */
    GOMP_MAP_FORCE_TO = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_TO),
    /* ..., and copy from device.  */
    GOMP_MAP_FORCE_FROM = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_FROM),
    /* ..., and copy to and from device.  */
    GOMP_MAP_FORCE_TOFROM = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_TOFROM),
...  ]

So before calling the offloaded function, a separate alloc is done for 
a, b and c, and the base pointers of the newly allocated objects are 
passed to the offloaded function.

This means we can mark those base pointers as restrict in the offloaded 
function.

Attached proof-of-concept patch implements that.

> We simply have some compiler internal interface between the caller and
> callee of the outlined regions, each interface in between those has
> its own structure type used to communicate the info;
> we can attach attributes on the fields, or some flags to indicate some
> properties interesting from aliasing POV.
> We don't really need to perform
> full IPA-PTA, perhaps it would be enough to a) record somewhere in cgraph
> the relationship in between such callers and callees (for offloading regions
> we already have "omp target entrypoint" attribute on the callee and a
> singler caller), tell LTO if possible not to split those into different
> partitions if easily possible, and then just for these pairs perform
> aliasing/points-to analysis in the caller and the result record using
> cliques/special attributes/whatever to the callee side, so that the callee
> (outlined OpenMP/OpenACC/Cilk+ region) can then improve its alias analysis.

As a start, is the approach of this patch OK?

It will allow us to commit the oacc kernels patch series with the 
ability to parallelize non-trivial testcases, and work on improving the 
alias bit after that.

Thanks,
- Tom



Mark pointers to allocated target vars as restricted, if possible

---
 gcc/omp-low.c | 67 ++++++++++++++++++++++++++++++++++++++++++++++++++++++-----
 1 file changed, 62 insertions(+), 5 deletions(-)

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 268b67b..0ce822d 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1372,7 +1372,8 @@ build_sender_ref (tree var, omp_context *ctx)
 /* Add a new field for VAR inside the structure CTX->SENDER_DECL.  */
 
 static void
-install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
+install_var_field_1 (tree var, bool by_ref, int mask, omp_context *ctx,
+		     bool base_pointers_restrict)
 {
   tree field, type, sfield = NULL_TREE;
   splay_tree_key key = (splay_tree_key) var;
@@ -1396,7 +1397,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);
 
@@ -1460,6 +1465,12 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
     splay_tree_insert (ctx->sfield_map, key, (splay_tree_value) sfield);
 }
 
+static void
+install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
+{
+  install_var_field_1 (var, by_ref, mask, ctx, false);
+}
+
 static tree
 install_var_local (tree var, omp_context *ctx)
 {
@@ -1816,7 +1827,8 @@ fixup_child_record_type (omp_context *ctx)
    specified by CLAUSES.  */
 
 static void
-scan_sharing_clauses (tree clauses, omp_context *ctx)
+scan_sharing_clauses_1 (tree clauses, omp_context *ctx,
+			bool base_pointers_restrict)
 {
   tree c, decl;
   bool scan_array_reductions = false;
@@ -2073,7 +2085,7 @@ 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_1 (decl, true, 3, ctx, base_pointers_restrict);
 		  if (is_gimple_omp_offloaded (ctx->stmt))
 		    install_var_local (decl, ctx);
 		}
@@ -2339,6 +2351,12 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	scan_omp (&OMP_CLAUSE_LINEAR_GIMPLE_SEQ (c), ctx);
 }
 
+static void
+scan_sharing_clauses (tree clauses, omp_context *ctx)
+{
+  scan_sharing_clauses_1 (clauses, ctx, false);
+}
+
 /* Create a new name for omp child function.  Returns an identifier.  If
    IS_CILK_FOR is true then the suffix for the child function is
    "_cilk_for_fn."  */
@@ -3056,13 +3074,52 @@ 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);
+
+      /* If all the clauses force allocation, we can be certain that the objects
+	 on the target are disjoint, and therefore mark the base pointers as
+	 restrict.  */
+      base_pointers_restrict = true;
+      tree c;
+      for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+	{
+	  switch (OMP_CLAUSE_CODE (c))
+	    {
+	    case OMP_CLAUSE_MAP:
+	      switch (OMP_CLAUSE_MAP_KIND (c))
+		{
+		case GOMP_MAP_ALLOC:
+		case GOMP_MAP_FORCE_TO:
+		case GOMP_MAP_FORCE_FROM:
+		case GOMP_MAP_FORCE_TOFROM:
+		  break;
+		default:
+		  base_pointers_restrict = false;
+		  break;
+		}
+	      break;
+
+	    default:
+	      base_pointers_restrict = false;
+	      break;
+	    }
+
+	  if (!base_pointers_restrict)
+	    break;
+	}
+      if (base_pointers_restrict)
+	{
+	  if (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_1 (clauses, ctx, base_pointers_restrict);
   scan_omp (gimple_omp_body_ptr (stmt), ctx);
 
   if (TYPE_FIELDS (ctx->record_type) == NULL)