OpenACC: "declare create" fixes wrt. "allocatable" variables

This patch fixes a case revealed by the previous patch where a synthetic
"acc data" region created for a "declare create" variable could interact
strangely with lexical inheritance behaviour.  In fact, it doesn't seem
right to create the "acc data" region for allocatable variables at all
-- doing so means that a data region is likely to be created for an
unallocated variable.

The fix is not to add such variables to the synthetic "acc data" region
at all, and defer to the code that performs "enter data"/"exit data"
for them when allocated/deallocated on the host instead. Then, "declare
create" variables are implicitly turned into "present" clauses on in-scope
offload regions.

gcc/fortran/
	* trans-openmp.cc (gfc_omp_finish_clause): Handle "declare create" for
	scalar allocatable variables.
	(gfc_trans_omp_clauses): Don't include allocatable vars in synthetic
	"acc data" region created for "declare create" variables.  Mark such
	variables with the "oacc declare create" attribute instead.  Don't
	create ALWAYS_POINTER mapping for target-to-host updates of declare
	create variables.
	(gfc_trans_oacc_declare): Handle empty clause list.

gcc/
	* gimplify.cc (gimplify_adjust_omp_clauses_1): Handle "oacc declare
	create" attribute.

gcc/testsuite/
	* c-c++-common/goacc/readonly-1.c: Adjust patterns.

libgomp/
	* testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90:
	Remove xfails.
	* testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90:
	Remove xfails.
	* testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90:
	Remove xfails.
	* testsuite/libgomp.oacc-fortran/declare-create-1.f90: New test.
	* testsuite/libgomp.oacc-fortran/declare-create-2.f90: New test.
	* testsuite/libgomp.oacc-fortran/declare-create-3.f90: New test.

Co-Authored-By: Paul-Antoine Arras <parras@baylibre.com>
Co-Authored-By: Sandra Loosemore <sandra@baylibre.com>

3 files changed, 52 insertions, 9 deletions

diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 36a3bcd..c28d937 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -1627,7 +1627,16 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p, bool openacc)
       orig_decl = decl;
 
       c4 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
-      OMP_CLAUSE_SET_MAP_KIND (c4, GOMP_MAP_POINTER);
+      if (openacc
+	  && GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
+	  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_PRESENT)
+	/* This allows "declare create" to work for scalar allocatables.  The
+	   resulting mapping nodes are:
+	     force_present(*var) firstprivate_pointer(var)
+	   which is the same as an explicit "present" clause gives.  */
+	OMP_CLAUSE_SET_MAP_KIND (c4, GOMP_MAP_FIRSTPRIVATE_POINTER);
+      else
+	OMP_CLAUSE_SET_MAP_KIND (c4, GOMP_MAP_POINTER);
       OMP_CLAUSE_DECL (c4) = decl;
       OMP_CLAUSE_SIZE (c4) = size_int (0);
       decl = build_fold_indirect_ref (decl);
@@ -4025,6 +4034,29 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	      if (!n->sym->attr.referenced)
 		continue;
 
+	      /* We do not want to include allocatable vars in a synthetic
+		 "acc data" region created for "!$acc declare create" vars.
+		 Such variables are handled by augmenting allocate/deallocate
+		 statements elsewhere (with
+		 "acc enter data declare_allocate(...)", etc.).  */
+	      if (op == EXEC_OACC_DECLARE
+		  && n->u.map.op == OMP_MAP_ALLOC
+		  && n->sym->attr.allocatable
+		  && n->sym->attr.oacc_declare_create)
+		{
+		  tree tree_var = gfc_get_symbol_decl (n->sym);
+		  if (!lookup_attribute ("oacc declare create",
+					 DECL_ATTRIBUTES (tree_var)))
+		    DECL_ATTRIBUTES (tree_var)
+		      = tree_cons (get_identifier ("oacc declare create"),
+				   NULL_TREE, DECL_ATTRIBUTES (tree_var));
+		  /* We might need to turn what would normally be a
+		     "firstprivate" mapping into a "present" mapping.  For the
+		     latter, we need the decl to be addressable.  */
+		  TREE_ADDRESSABLE (tree_var) = 1;
+		  continue;
+		}
+
 	      location_t map_loc = gfc_get_location (&n->where);
 	      bool always_modifier = false;
 	      tree node = build_omp_clause (map_loc, OMP_CLAUSE_MAP);
@@ -4255,7 +4287,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
 			    gmk = GOMP_MAP_RELEASE;
 			  else if (GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
-				   && n->sym->attr.oacc_declare_create)
+				   && n->sym->attr.oacc_declare_create
+				   && n->u.map.op != OMP_MAP_FORCE_FROM)
 			    {
 			      if (clauses->update_allocatable)
 				gmk = GOMP_MAP_ALWAYS_POINTER;
@@ -9363,10 +9396,12 @@ gfc_trans_oacc_declare (gfc_code *code)
   gfc_start_block (&block);
 
   oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.oacc_declare->clauses,
-					code->loc, false, true);
+					code->loc, false, true,
+					EXEC_OACC_DECLARE);
   stmt = gfc_trans_omp_code (code->block->next, true);
-  stmt = build2_loc (input_location, construct_code, void_type_node, stmt,
-		     oacc_clauses);
+  if (oacc_clauses)
+    stmt = build2_loc (input_location, construct_code, void_type_node, stmt,
+		       oacc_clauses);
   gfc_add_expr_to_block (&block, stmt);
 
   return gfc_finish_block (&block);
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 31e4b4d..00fb34f 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -14647,6 +14647,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
 		g->have_offload = true;
 	    }
 	}
+      if (lookup_attribute ("oacc declare create", DECL_ATTRIBUTES (decl)))
+	flags |= GOVD_MAP_FORCE_PRESENT;
     }
   else if (flags & GOVD_SHARED)
     {
@@ -14686,6 +14688,12 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
 		 "%<target%> construct", decl);
 	  return 0;
 	}
+      if (lookup_attribute ("oacc declare create", DECL_ATTRIBUTES (decl)))
+	{
+	  code = OMP_CLAUSE_MAP;
+	  flags &= ~GOVD_FIRSTPRIVATE;
+	  flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
+	}
     }
   else if (flags & GOVD_LASTPRIVATE)
     code = OMP_CLAUSE_LASTPRIVATE;
diff --git a/gcc/testsuite/c-c++-common/goacc/readonly-1.c b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
index 300464c..bd4d229 100644
--- a/gcc/testsuite/c-c++-common/goacc/readonly-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
@@ -51,14 +51,14 @@ int main (void)
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
 
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
-/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
 
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\);$" 4 "original" } } */
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\);$" 4 "original" } } */