Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - middle end

[Cesar Philippidis <cesar at codesourcery dot com>]

This patch implements the OpenACC 2.5 data clause semantics in the
middle end.

Is it OK for trunk?

Cesar

2018-06-19  Chung-Lin Tang <cltang@codesourcery.com>
	    Thomas Schwinge <thomas@codesourcery.com>
	    Cesar Philippidis  <cesar@codesourcery.com>

	gcc/c-family/
	* c-pragma.h (enum pragma_omp_clause): Add
	PRAGMA_OACC_CLAUSE_{FINALIZE,IF_PRESENT}. Remove
	PRAGMA_OACC_CLAUSE_PRESENT_OR_{COPY,COPYIN,COPYOUT,CREATE}.

	gcc/
	* gimplify.c (gimplify_scan_omp_clauses): Add support for
	OMP_CLAUSE_{IF_PRESENT,FINALIZE}.
	(gimplify_adjust_omp_clauses): Likewise.
	(gimplify_oacc_declare_1): Add support for GOMP_MAP_RELEASE, remove
	support for GOMP_MAP_FORCE_{ALLOC,TO,FROM,TOFROM}.
	(gimplify_omp_target_update): Update handling of acc update and
	enter/exit data.
	* omp-low.c (install_var_field): Remove unused parameter
	base_pointers_restrict.
	(scan_sharing_clauses): Remove base_pointers_restrict parameter.
	Update call to install_var_field. Handle OMP_CLAUSE_{IF_PRESENT,
	FINALIZE}
	(omp_target_base_pointers_restrict_p): Delete.
	(scan_omp_target): Update call to scan_sharing_clauses.
	* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_{IF_PRESENT,
	FINALIZE}.
	* tree-nested.c (convert_nonlocal_omp_clauses): Handle
	OMP_CLAUSE_{IF_PRESENT,FINALIZE}.
	(convert_local_omp_clauses): Likewise.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	* tree.c (omp_clause_num_ops): Add entries for 	OMP_CLAUSE_{IF_PRESENT,
	FINALIZE}.
	(omp_clause_code_name): Likewise.


>From f79b0e6f0d796dc18ef1faf20b9fad0b7feeaa94 Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <cesar@codesourcery.com>
Date: Tue, 19 Jun 2018 09:32:20 -0700
Subject: [PATCH 6/7] gcc middle end

---
 gcc/c-family/c-pragma.h |  6 +--
 gcc/gimplify.c          | 67 ++++++++++++++++++++++-------
 gcc/omp-low.c           | 93 +++++------------------------------------
 gcc/tree-core.h         |  8 +++-
 gcc/tree-nested.c       |  4 ++
 gcc/tree-pretty-print.c |  6 +++
 gcc/tree.c              |  8 +++-
 7 files changed, 88 insertions(+), 104 deletions(-)

diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index c70380c211b..b322547b11a 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -138,16 +138,13 @@ enum pragma_omp_clause {
   PRAGMA_OACC_CLAUSE_DELETE,
   PRAGMA_OACC_CLAUSE_DEVICEPTR,
   PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
+  PRAGMA_OACC_CLAUSE_FINALIZE,
   PRAGMA_OACC_CLAUSE_GANG,
   PRAGMA_OACC_CLAUSE_HOST,
   PRAGMA_OACC_CLAUSE_INDEPENDENT,
   PRAGMA_OACC_CLAUSE_NUM_GANGS,
   PRAGMA_OACC_CLAUSE_NUM_WORKERS,
   PRAGMA_OACC_CLAUSE_PRESENT,
-  PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY,
-  PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN,
-  PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT,
-  PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE,
   PRAGMA_OACC_CLAUSE_SELF,
   PRAGMA_OACC_CLAUSE_SEQ,
   PRAGMA_OACC_CLAUSE_TILE,
@@ -156,6 +153,7 @@ enum pragma_omp_clause {
   PRAGMA_OACC_CLAUSE_VECTOR_LENGTH,
   PRAGMA_OACC_CLAUSE_WAIT,
   PRAGMA_OACC_CLAUSE_WORKER,
+  PRAGMA_OACC_CLAUSE_IF_PRESENT,
   PRAGMA_OACC_CLAUSE_COLLAPSE = PRAGMA_OMP_CLAUSE_COLLAPSE,
   PRAGMA_OACC_CLAUSE_COPYIN = PRAGMA_OMP_CLAUSE_COPYIN,
   PRAGMA_OACC_CLAUSE_DEVICE = PRAGMA_OMP_CLAUSE_DEVICE,
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 1523a27e828..97543ed5f70 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8524,6 +8524,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	case OMP_CLAUSE_NOGROUP:
 	case OMP_CLAUSE_THREADS:
 	case OMP_CLAUSE_SIMD:
+	case OMP_CLAUSE_IF_PRESENT:
+	case OMP_CLAUSE_FINALIZE:
 	  break;
 
 	case OMP_CLAUSE_DEFAULTMAP:
@@ -9305,6 +9307,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
 	case OMP_CLAUSE_TILE:
+	case OMP_CLAUSE_IF_PRESENT:
+	case OMP_CLAUSE_FINALIZE:
 	  break;
 
 	default:
@@ -9361,21 +9365,7 @@ gimplify_oacc_declare_1 (tree clause)
   switch (kind)
     {
       case GOMP_MAP_ALLOC:
-      case GOMP_MAP_FORCE_ALLOC:
-      case GOMP_MAP_FORCE_TO:
-	new_op = GOMP_MAP_DELETE;
-	ret = true;
-	break;
-
-      case GOMP_MAP_FORCE_FROM:
-	OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC);
-	new_op = GOMP_MAP_FORCE_FROM;
-	ret = true;
-	break;
-
-      case GOMP_MAP_FORCE_TOFROM:
-	OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_TO);
-	new_op = GOMP_MAP_FORCE_FROM;
+	new_op = GOMP_MAP_RELEASE;
 	ret = true;
 	break;
 
@@ -10817,6 +10807,53 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
 			     ort, TREE_CODE (expr));
   gimplify_adjust_omp_clauses (pre_p, NULL, &OMP_STANDALONE_CLAUSES (expr),
 			       TREE_CODE (expr));
+  if (TREE_CODE (expr) == OACC_UPDATE
+      && omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
+			  OMP_CLAUSE_IF_PRESENT))
+    {
+      /* The runtime uses GOMP_MAP_{TO,FROM} to denote the if_present
+	 clause.  */
+      for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
+	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+	  switch (OMP_CLAUSE_MAP_KIND (c))
+	    {
+	    case GOMP_MAP_FORCE_TO:
+	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+	      break;
+	    case GOMP_MAP_FORCE_FROM:
+	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FROM);
+	      break;
+	    default:
+	      break;
+	    }
+    }
+  else if (TREE_CODE (expr) == OACC_EXIT_DATA
+	   && omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
+			       OMP_CLAUSE_FINALIZE))
+    {
+      /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"
+	 semantics apply to all mappings of this OpenACC directive.  */
+      bool finalize_marked = false;
+      for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
+	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+	  switch (OMP_CLAUSE_MAP_KIND (c))
+	    {
+	    case GOMP_MAP_FROM:
+	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_FROM);
+	      finalize_marked = true;
+	      break;
+	    case GOMP_MAP_RELEASE:
+	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
+	      finalize_marked = true;
+	      break;
+	    default:
+	      /* Check consistency: libgomp relies on the very first data
+		 mapping clause being marked, so make sure we did that before
+		 any other mapping clauses.  */
+	      gcc_assert (finalize_marked);
+	      break;
+	    }
+    }
   stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
 
   gimplify_seq_add_stmt (pre_p, stmt);
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index ba6c705cf8b..c591231d8f1 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -642,8 +642,7 @@ 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,
-		   bool base_pointers_restrict = false)
+install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
 {
   tree field, type, sfield = NULL_TREE;
   splay_tree_key key = (splay_tree_key) var;
@@ -674,11 +673,7 @@ 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);
-      if (base_pointers_restrict)
-	type = build_qualified_type (type, TYPE_QUAL_RESTRICT);
-    }
+    type = build_pointer_type (type);
   else if ((mask & 3) == 1 && omp_is_reference (var))
     type = TREE_TYPE (type);
 
@@ -992,12 +987,10 @@ fixup_child_record_type (omp_context *ctx)
 }
 
 /* Instantiate decls as necessary in CTX to satisfy the data sharing
-   specified by CLAUSES.  If BASE_POINTERS_RESTRICT, install var field with
-   restrict.  */
+   specified by CLAUSES.  */
 
 static void
-scan_sharing_clauses (tree clauses, omp_context *ctx,
-		      bool base_pointers_restrict = false)
+scan_sharing_clauses (tree clauses, omp_context *ctx)
 {
   tree c, decl;
   bool scan_array_reductions = false;
@@ -1256,8 +1249,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,
-				       base_pointers_restrict);
+		    install_var_field (decl, true, 3, ctx);
 		  if (is_gimple_omp_offloaded (ctx->stmt)
 		      && !OMP_CLAUSE_MAP_IN_REDUCTION (c))
 		    install_var_local (decl, ctx);
@@ -1328,6 +1320,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE__SIMT_:
 	case OMP_CLAUSE_DEFAULT:
+	case OMP_CLAUSE_IF_PRESENT:
+	case OMP_CLAUSE_FINALIZE:
 	  break;
 
 	case OMP_CLAUSE_ALIGNED:
@@ -1499,6 +1493,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE__GRIDDIM_:
 	case OMP_CLAUSE__SIMT_:
+	case OMP_CLAUSE_IF_PRESENT:
+	case OMP_CLAUSE_FINALIZE:
 	  break;
 
 	case OMP_CLAUSE__CACHE_:
@@ -2266,68 +2262,6 @@ 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
@@ -2349,20 +2283,13 @@ 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, base_pointers_restrict);
+  scan_sharing_clauses (clauses, ctx);
   scan_omp (gimple_omp_body_ptr (stmt), ctx);
 
   if (TYPE_FIELDS (ctx->record_type) == NULL)
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 2bebb22a7e9..4a04e9e8b26 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -454,7 +454,13 @@ enum omp_clause_code {
 
   /* OpenMP internal-only clause to specify grid dimensions of a gridified
      kernel.  */
-  OMP_CLAUSE__GRIDDIM_
+  OMP_CLAUSE__GRIDDIM_,
+
+  /* OpenACC clause: if_present.  */
+  OMP_CLAUSE_IF_PRESENT,
+
+  /* OpenACC clause: finalize.  */
+  OMP_CLAUSE_FINALIZE
 };
 
 #undef DEFTREESTRUCT
diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c
index b335d6b0afe..257ceae6f2d 100644
--- a/gcc/tree-nested.c
+++ b/gcc/tree-nested.c
@@ -1333,6 +1333,8 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_SEQ:
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_AUTO:
+	case OMP_CLAUSE_IF_PRESENT:
+	case OMP_CLAUSE_FINALIZE:
 	  break;
 
 	  /* The following clause belongs to the OpenACC cache directive, which
@@ -2022,6 +2024,8 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_SEQ:
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_AUTO:
+	case OMP_CLAUSE_IF_PRESENT:
+	case OMP_CLAUSE_FINALIZE:
 	  break;
 
 	  /* The following clause belongs to the OpenACC cache directive, which
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 63ec823c0ba..e65c40a41a3 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -1045,6 +1045,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 			 false);
       pp_right_paren (pp);
       break;
+    case OMP_CLAUSE_IF_PRESENT:
+      pp_string (pp, "if_present");
+      break;
+    case OMP_CLAUSE_FINALIZE:
+      pp_string (pp, "finalize");
+      break;
 
     default:
       /* Should never happen.  */
diff --git a/gcc/tree.c b/gcc/tree.c
index 889d88c50b4..cf0b5f6dad3 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -343,6 +343,8 @@ unsigned const char omp_clause_num_ops[] =
   1, /* OMP_CLAUSE_VECTOR_LENGTH  */
   3, /* OMP_CLAUSE_TILE  */
   2, /* OMP_CLAUSE__GRIDDIM_  */
+  0, /* OMP_CLAUSE_IF_PRESENT */
+  0, /* OMP_CLAUSE_FINALIZE */
 };
 
 const char * const omp_clause_code_name[] =
@@ -413,7 +415,9 @@ const char * const omp_clause_code_name[] =
   "num_workers",
   "vector_length",
   "tile",
-  "_griddim_"
+  "_griddim_",
+  "if_present",
+  "finalize",
 };
 
 
@@ -11579,6 +11583,8 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
 	case OMP_CLAUSE_SEQ:
 	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE__SIMT_:
+	case OMP_CLAUSE_IF_PRESENT:
+	case OMP_CLAUSE_FINALIZE:
 	  WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp));
 
 	case OMP_CLAUSE_LASTPRIVATE:
-- 
2.17.1