[og9] OpenACC assumed-size arrays with non-lexical data mappings

Message ID 20190710015251.118ea4cb@squid.athome
State New
Headers show
Series
  • [og9] OpenACC assumed-size arrays with non-lexical data mappings
Related show

Commit Message

Julian Brown July 10, 2019, 12:52 a.m.
Hi,

This patch provides support for implicit mapping of assumed-sized
arrays for OpenACC, in cases where those arrays have previously been
mapped using non-lexical data mappings (e.g. "#pragma acc enter data").

Previously posted here:

https://gcc.gnu.org/ml/gcc-patches/2016-08/msg02090.html

and then revised:

https://gcc.gnu.org/ml/gcc-patches/2017-07/msg00069.html

It's not clear if this is required behaviour for OpenACC, but at least
one test program we are using relies on the semantics introduced by this
patch.

Tested with offloading to nvptx. I will apply to the
openacc-gcc-9-branch shortly.

Julian

ChangeLog

2019-07-10  Cesar Philippidis  <cesar@codesourcery.com>
            Thomas Schwinge  <thomas@codesourcery.com>
            Julian Brown  <julian@codesourcery.com>

	gcc/
	* gimplify.c (gimplify_adjust_omp_clauses_1): Raise error for
	assumed-size arrays in map clauses for Fortran/OpenMP.
	* omp-low.c (lower_omp_target): Set the size of assumed-size
	Fortran arrays to one to allow use of data already mapped on
	the offload device.

	gcc/fortran/
	* trans-openmp.c (gfc_omp_finish_clause): Change clauses mapping
	assumed-size arrays to use the GOMP_MAP_FORCE_PRESENT map type.

Patch

From 2c5a7e445ebadc920730c732279732d2f9b40598 Mon Sep 17 00:00:00 2001
From: Julian Brown <julian@codesourcery.com>
Date: Thu, 4 Jul 2019 18:14:41 -0700
Subject: [PATCH 2/3] Assumed-size arrays with non-lexical data mappings

	gcc/fortran/
	* trans-openmp.c (gfc_omp_finish_clause): Change clauses mapping
	assumed-size arrays to use the GOMP_MAP_FORCE_PRESENT map type.
	* gimplify.c (gimplify_adjust_omp_clauses_1): Raise error for
	assumed-size arrays in map clauses for Fortran/OpenMP.
	* omp-low.c (lower_omp_target): Set the size of assumed-size Fortran
	arrays to one to allow use of data already mapped on the offload device.
---
 gcc/fortran/ChangeLog.openacc |  9 +++++++++
 gcc/fortran/trans-openmp.c    | 22 +++++++++++++---------
 gcc/gimplify.c                | 14 ++++++++++++++
 gcc/omp-low.c                 |  5 +++++
 4 files changed, 41 insertions(+), 9 deletions(-)

diff --git a/gcc/fortran/ChangeLog.openacc b/gcc/fortran/ChangeLog.openacc
index c44a5ebdb3b..beba7d94ad2 100644
--- a/gcc/fortran/ChangeLog.openacc
+++ b/gcc/fortran/ChangeLog.openacc
@@ -1,3 +1,12 @@ 
+2019-07-10  Julian Brown  <julian@codesourcery.com>
+
+	* trans-openmp.c (gfc_omp_finish_clause): Change clauses mapping
+	assumed-size arrays to use the GOMP_MAP_FORCE_PRESENT map type.
+	* gimplify.c (gimplify_adjust_omp_clauses_1): Raise error for
+	assumed-size arrays in map clauses for Fortran/OpenMP.
+	* omp-low.c (lower_omp_target): Set the size of assumed-size Fortran
+	arrays to one to allow use of data already mapped on the offload device.
+
 2019-07-10  Julian Brown  <julian@codesourcery.com>
 
 	* openmp.c (resolve_oacc_data_clauses): Allow polymorphic allocatable
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index d5ae0b717df..db009130c85 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1137,10 +1137,18 @@  gfc_omp_finish_clause (tree c, gimple_seq *pre_p)
   tree decl = OMP_CLAUSE_DECL (c);
 
   /* Assumed-size arrays can't be mapped implicitly, they have to be mapped
-     explicitly using array sections.  An exception is if the array is
-     mapped explicitly in an enclosing data construct for OpenACC, in which
-     case we see GOMP_MAP_FORCE_PRESENT here and do not need to raise an
-     error.  */
+     explicitly using array sections.  For OpenACC this restriction is lifted
+     if the array has already been mapped:
+
+       - Using a lexically-enclosing data region: in that case we see the
+         GOMP_MAP_FORCE_PRESENT mapping kind here.
+
+       - Using a non-lexical data mapping ("acc enter data").
+
+     In the latter case we change the mapping type to GOMP_MAP_FORCE_PRESENT.
+     This raises an error for OpenMP in our the caller
+     (gimplify.c:gimplify_adjust_omp_clauses_1).  OpenACC will raise a runtime
+     error if the assumed-size array is not mapped.  */
   if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_PRESENT
       && TREE_CODE (decl) == PARM_DECL
       && GFC_ARRAY_TYPE_P (TREE_TYPE (decl))
@@ -1148,11 +1156,7 @@  gfc_omp_finish_clause (tree c, gimple_seq *pre_p)
       && GFC_TYPE_ARRAY_UBOUND (TREE_TYPE (decl),
 				GFC_TYPE_ARRAY_RANK (TREE_TYPE (decl)) - 1)
 	 == NULL)
-    {
-      error_at (OMP_CLAUSE_LOCATION (c),
-		"implicit mapping of assumed size array %qD", decl);
-      return;
-    }
+    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_PRESENT);
 
   tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE;
   if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 60e04ff8353..58142c9eb90 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10088,7 +10088,21 @@  gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
   *list_p = clause;
   struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
   gimplify_omp_ctxp = ctx->outer_context;
+  gomp_map_kind kind = (code == OMP_CLAUSE_MAP) ? OMP_CLAUSE_MAP_KIND (clause)
+						: (gomp_map_kind) GOMP_MAP_LAST;
   lang_hooks.decls.omp_finish_clause (clause, pre_p);
+  /* Allow OpenACC to have implicit assumed-size arrays via FORCE_PRESENT,
+     which should work as long as the array has previously been mapped
+     explicitly on the target (e.g. by "enter data").  Raise an error for
+     OpenMP.  */
+  if (lang_GNU_Fortran ()
+      && code == OMP_CLAUSE_MAP
+      && (ctx->region_type & ORT_ACC) == 0
+      && kind == GOMP_MAP_TOFROM
+      && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_FORCE_PRESENT)
+    error_at (OMP_CLAUSE_LOCATION (clause),
+	      "implicit mapping of assumed size array %qD",
+	      OMP_CLAUSE_DECL (clause));
   if (gimplify_omp_ctxp)
     for (; clause != chain; clause = OMP_CLAUSE_CHAIN (clause))
       if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 249cc1aac53..97c00217d9f 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -10453,6 +10453,11 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	      s = OMP_CLAUSE_SIZE (c);
 	    if (s == NULL_TREE)
 	      s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
+	    /* Fortran assumed-size arrays have zero size because the type is
+	       incomplete.  Set the size to one to allow the runtime to remap
+	       any existing data that is already present on the accelerator.  */
+	    if (s == NULL_TREE && is_gimple_omp_oacc (ctx->stmt))
+	      s = integer_one_node;
 	    s = fold_convert (size_type_node, s);
 	    decl_args = append_decl_arg (ovar, decl_args, ctx);
 	    purpose = size_int (map_idx++);
-- 
2.22.0