Re: [patch,openacc] Generate sequential loop for OpenACC loop directive inside kernels

[Cesar Philippidis <cesar at codesourcery dot com>]

On 09/20/2018 10:14 AM, Cesar Philippidis wrote:
> As Chung-Lin noted here
> <https://gcc.gnu.org/ml/gcc-patches/2015-06/msg01079.html>:
> 
>   This patch adjusts omp-low.c:expand_omp_for_generic() to expand to a
>   "sequential" loop form (without the OMP runtime calls), used for loop
>   directives inside OpenACC kernels constructs. Tom mentions that this
>   allows the kernels parallelization to work when '#pragma acc loop'
>   makes the front-ends create OMP_FOR, which the loop analysis phases
>   don't understand.
> 
> I bootstrapped and regtested it on x86_64 Linux with nvptx offloading.
> Is this patch OK for trunk?

I forgot to mention how that patch depends on the
omp_target_base_pointers_restrict_p functionality from omp lowering that
I removed back in June when I added support for the OpenACC 2.5 data
clause semantics. It turned out that I was too aggressive when I was
removing unused code. That's because, at least initially, there was no
test cases that exercised that functionality in trunk until Chung-Lin's
kernels patch goes in.

Anyway, this patch is specifically required to get
kernels-acc-loop-reduction.c working.

Is this OK for trunk? I bootstrapped and regression tested it on x86_64
Linux with nvptx offloading.

Thanks,
Cesar

[OpenACC] Reintroduce omp_target_base_pointers_restrict_p

It turns out that existing acc kernels instructure based on parloops
will benefit if the variables used in OpenACC data clauses maintained
the restrict pointer qualifier. This code is present in GCC 8, but I
removed it back in June when I committed a patch to update the
behavior of the data clauses match the semantics in OpenACC 2.5.

Is this patch OK for trunk? A forthcoming acc kernels patch depends on
it.

2018-XX-YY  Cesar Philippidis  <cesar@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.
---
 gcc/omp-low.c | 89 +++++++++++++++++++++++++++++++++++++++++++++++----
 1 file changed, 83 insertions(+), 6 deletions(-)

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 24685fd012c..a59c15ae5fd 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -642,7 +642,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;
@@ -673,7 +674,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);
 
@@ -987,10 +992,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;
@@ -1252,7 +1259,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);
@@ -2265,6 +2273,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
@@ -2286,13 +2356,20 @@ 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");
     }
 
-  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)
-- 
2.17.1