Fortran "declare create"/allocate support for OpenACC

This patch incorporates these commits from OG14 branch:
65be1389eeda9b3b97f6587721215c3f31bd7f98
9d43e819d88f97c7ade7f8c95c35ea3464ea7771
f2cf2b994c4d8c871fad5502ffb9aaee9ea4f4e0
2770ce41615557e595065ce0c5db71e9f3d82b0a
a29e58f4b314862a72730119f85e9125879abf0b
ffd990543f805ed448aaa355d190f37103f8f1f0

gcc/ChangeLog
	* gimplify.cc (omp_group_base): Handle GOMP_MAP_DECLARE_ALLOCATE
	and GOMP_MAP_DECLARE_DEALLOCATE.
	(gimplify_adjust_omp_clauses): Likewise.
	* omp-low.cc (scan_sharing_clauses): Update handling of OpenACC declare
	create, declare copyin and declare deviceptr to have local lifetimes.
	(convert_to_firstprivate_int): Handle pointer types.
	(convert_from_firstprivate_int): Likewise.  Create local storage for
	the values being pointed to.  Add new orig_type argument.  Use
	VIEW_CONVERT also for vectors.
	(lower_omp_target): Handle GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}.
	Add orig_type argument to convert_from_firstprivate_int call.
	Allow pointer types with GOMP_MAP_FIRSTPRIVATE_INT.  Don't privatize
	firstprivate VLAs.
	* tree-pretty-print.cc (dump_omp_clause): Handle
	GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}.

gcc/fortran/ChangeLog
	* gfortran.h (enum gfc_omp_map_op): Add OMP_MAP_DECLARE_ALLOCATE,
	OMP_MAP_DECLARE_DEALLOCATE.
	(gfc_omp_clauses): Add update_allocatable.
	* trans-array.cc (gfc_array_allocate): Call
	gfc_trans_oacc_declare_allocate for decls that have oacc_declare_create
	attribute set.
	* trans-decl.cc (find_module_oacc_declare_clauses): Relax
	oacc_declare_create to OMP_MAP_ALLOC, and oacc_declare_copyin to
	OMP_MAP_TO, in order to	match OpenACC 2.5 semantics.
	* trans-openmp.cc (gfc_omp_check_optional_argument): Handle non-decl
	case.
	(gfc_trans_omp_clauses): Use GOMP_MAP_ALWAYS_POINTER (for update
	directive) or GOMP_MAP_FIRSTPRIVATE_POINTER (otherwise) for
	allocatable scalar decls.  Handle OMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}
	clauses.
	(gfc_trans_oacc_executable_directive): Use GOMP_MAP_ALWAYS_POINTER
	for allocatable scalar data clauses inside acc update directives.
	(gfc_trans_oacc_declare_allocate): New function.
	* trans-stmt.cc (gfc_trans_allocate): Call
	gfc_trans_oacc_declare_allocate for decls with oacc_declare_create
	attribute set.
	(gfc_trans_deallocate): Likewise.
	* trans.h (gfc_trans_oacc_declare_allocate): Declare.

gcc/testsuite/ChangeLog
	* gfortran.dg/goacc/declare-allocatable-1.f90: New test.
	* gfortran.dg/goacc/declare-3.f95: Adjust expected dump output.

include/ChangeLog
	* gomp-constants.h (enum gomp_map_kind): Define
	GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE} and GOMP_MAP_FLAG_SPECIAL_4.

libgomp/ChangeLog
	* libgomp.h (gomp_acc_declare_allocate): Remove prototype.
	* oacc-mem.c (gomp_acc_declare_allocate): New function.
	(find_group_last): Handle GOMP_MAP_DECLARE_ALLOCATE and
	GOMP_MAP_DECLARE_DEALLOCATE groupings.
	(goacc_enter_data_internal): Fix kind check for
	GOMP_MAP_DECLARE_ALLOCATE. Pass new pointer argument to
	gomp_acc_declare_allocate.   Unlock mutex before calling
	gomp_acc_declare_allocate and relock it afterwards.
	(goacc_exit_data_internal): Unlock device mutex around
	gomp_acc_declare_allocate call. Pass new pointer argument. Handle
	group pointer mapping for deallocate.
	* testsuite/libgomp.oacc-fortran/allocatable-scalar.f90: New test.
	* testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90:
	Adjust.
	* testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/declare-allocatable-2.f90: New test.
	* testsuite/libgomp.oacc-fortran/declare-allocatable-3.f90: New test.
	* testsuite/libgomp.oacc-fortran/declare-allocatable-4.f90: New test.
	* testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90:
	Adjust.
	* testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1.f90:
	New test.

Co-Authored-By: Julian Brown  <julian@codesourcery.com>
Co-Authored-By: Kwok Cheung Yeung  <kcy@codesourcery.com>
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Paul-Antoine Arras <parras@baylibre.com>

11 files changed, 183 insertions, 29 deletions

diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 5ef7037..5b18bca 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1348,7 +1348,9 @@ enum gfc_omp_map_op
   OMP_MAP_PRESENT_TOFROM,
   OMP_MAP_ALWAYS_PRESENT_TO,
   OMP_MAP_ALWAYS_PRESENT_FROM,
-  OMP_MAP_ALWAYS_PRESENT_TOFROM
+  OMP_MAP_ALWAYS_PRESENT_TOFROM,
+  OMP_MAP_DECLARE_ALLOCATE,
+  OMP_MAP_DECLARE_DEALLOCATE
 };
 
 enum gfc_omp_defaultmap
@@ -1675,7 +1677,7 @@ typedef struct gfc_omp_clauses
   unsigned async:1, gang:1, worker:1, vector:1, seq:1, independent:1;
   unsigned par_auto:1, gang_static:1;
   unsigned if_present:1, finalize:1;
-  unsigned nohost:1;
+  unsigned nohost:1, update_allocatable:1;
   locus loc;
 }
 gfc_omp_clauses;
diff --git a/gcc/fortran/trans-array.cc b/gcc/fortran/trans-array.cc
index 9606131..92254fe 100644
--- a/gcc/fortran/trans-array.cc
+++ b/gcc/fortran/trans-array.cc
@@ -6504,6 +6504,7 @@ gfc_array_allocate (gfc_se * se, gfc_expr * expr, tree status, tree errmsg,
   bool allocatable, coarray, dimension, alloc_w_e3_arr_spec = false,
       non_ulimate_coarray_ptr_comp;
   tree omp_cond = NULL_TREE, omp_alt_alloc = NULL_TREE;
+  bool oacc_declare = false;
 
   ref = expr->ref;
 
@@ -6518,6 +6519,7 @@ gfc_array_allocate (gfc_se * se, gfc_expr * expr, tree status, tree errmsg,
       allocatable = expr->symtree->n.sym->attr.allocatable;
       dimension = expr->symtree->n.sym->attr.dimension;
       non_ulimate_coarray_ptr_comp = false;
+      oacc_declare = expr->symtree->n.sym->attr.oacc_declare_create;
     }
   else
     {
@@ -6755,6 +6757,8 @@ gfc_array_allocate (gfc_se * se, gfc_expr * expr, tree status, tree errmsg,
       gfc_conv_descriptor_offset_set (&set_descriptor_block, se->expr, offset);
       tmp = fold_convert (gfc_array_index_type, element_size);
       gfc_conv_descriptor_span_set (&set_descriptor_block, se->expr, tmp);
+      if (oacc_declare)
+	gfc_trans_oacc_declare_allocate (&set_descriptor_block, expr, true);
     }
 
   set_descriptor = gfc_finish_block (&set_descriptor_block);
diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc
index 4f2ea76..7ec3a1f 100644
--- a/gcc/fortran/trans-decl.cc
+++ b/gcc/fortran/trans-decl.cc
@@ -6941,10 +6941,10 @@ find_module_oacc_declare_clauses (gfc_symbol *sym)
       gfc_omp_map_op map_op;
 
       if (sym->attr.oacc_declare_create)
-	map_op = OMP_MAP_FORCE_ALLOC;
+	map_op = OMP_MAP_ALLOC;
 
       if (sym->attr.oacc_declare_copyin)
-	map_op = OMP_MAP_FORCE_TO;
+	map_op = OMP_MAP_TO;
 
       if (sym->attr.oacc_declare_deviceptr)
 	map_op = OMP_MAP_FORCE_DEVICEPTR;
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 9767f20..c76c98b 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -102,6 +102,10 @@ gfc_omp_check_optional_argument (tree decl, bool for_present_check)
   if (!for_present_check)
     return gfc_omp_is_optional_argument (decl) ? decl : NULL_TREE;
 
+  if (!DECL_P (decl))
+    return fold_build2_loc (input_location, NE_EXPR, boolean_type_node,
+			    decl, null_pointer_node);
+
   if (!DECL_LANG_SPECIFIC (decl))
     return NULL_TREE;
 
@@ -4117,6 +4121,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		case OMP_MAP_FORCE_DEVICEPTR:
 		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_DEVICEPTR);
 		  break;
+		case OMP_MAP_DECLARE_ALLOCATE:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DECLARE_ALLOCATE);
+		  break;
+		case OMP_MAP_DECLARE_DEALLOCATE:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DECLARE_DEALLOCATE);
+		  break;
 		default:
 		  gcc_unreachable ();
 		}
@@ -4239,6 +4249,14 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			    gmk = GOMP_MAP_DELETE;
 			  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)
+			    {
+			      if (clauses->update_allocatable)
+				gmk = GOMP_MAP_ALWAYS_POINTER;
+			      else
+				gmk = GOMP_MAP_FIRSTPRIVATE_POINTER;
+			    }
 			  tree size;
 			  if (gmk == GOMP_MAP_RELEASE || gmk == GOMP_MAP_DELETE)
 			    size = TYPE_SIZE_UNIT (TREE_TYPE (decl));
@@ -5978,12 +5996,14 @@ gfc_trans_oacc_executable_directive (gfc_code *code)
 {
   stmtblock_t block;
   tree stmt, oacc_clauses;
+  gfc_omp_clauses *clauses = code->ext.omp_clauses;
   enum tree_code construct_code;
 
   switch (code->op)
     {
       case EXEC_OACC_UPDATE:
 	construct_code = OACC_UPDATE;
+	clauses->update_allocatable = 1;
 	break;
       case EXEC_OACC_ENTER_DATA:
 	construct_code = OACC_ENTER_DATA;
@@ -5999,8 +6019,8 @@ gfc_trans_oacc_executable_directive (gfc_code *code)
     }
 
   gfc_start_block (&block);
-  oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
-					code->loc, false, true, code->op);
+  oacc_clauses = gfc_trans_omp_clauses (&block, clauses, code->loc,
+					false, true, code->op);
   stmt = build1_loc (input_location, construct_code, void_type_node,
 		     oacc_clauses);
   gfc_add_expr_to_block (&block, stmt);
@@ -9324,6 +9344,41 @@ gfc_trans_oacc_declare (gfc_code *code)
   return gfc_finish_block (&block);
 }
 
+/* Create an OpenACC enter or exit data construct for an OpenACC declared
+   variable that has been allocated or deallocated.  */
+
+tree
+gfc_trans_oacc_declare_allocate (stmtblock_t *block, gfc_expr *expr,
+				 bool allocate)
+{
+  gfc_omp_clauses *clauses = gfc_get_omp_clauses ();
+  gfc_omp_namelist *p = gfc_get_omp_namelist ();
+  tree oacc_clauses, stmt;
+  enum tree_code construct_code;
+
+  p->sym = expr->symtree->n.sym;
+  p->where = expr->where;
+
+  if (allocate)
+    {
+      p->u.map.op = OMP_MAP_DECLARE_ALLOCATE;
+      construct_code = OACC_ENTER_DATA;
+    }
+  else
+    {
+      p->u.map.op = OMP_MAP_DECLARE_DEALLOCATE;
+      construct_code = OACC_EXIT_DATA;
+    }
+  clauses->lists[OMP_LIST_MAP] = p;
+
+  oacc_clauses = gfc_trans_omp_clauses (block, clauses, expr->where);
+  stmt = build1_loc (input_location, construct_code, void_type_node,
+		     oacc_clauses);
+  gfc_add_expr_to_block (block, stmt);
+
+  return stmt;
+}
+
 tree
 gfc_trans_oacc_directive (gfc_code *code)
 {
diff --git a/gcc/fortran/trans-stmt.cc b/gcc/fortran/trans-stmt.cc
index 37f8aca..2dfecff 100644
--- a/gcc/fortran/trans-stmt.cc
+++ b/gcc/fortran/trans-stmt.cc
@@ -7330,6 +7330,10 @@ gfc_trans_allocate (gfc_code * code, gfc_omp_namelist *omp_allocate)
 	  else
 	    gfc_allocate_using_malloc (&se.pre, se.expr, memsz, stat,
 				      omp_cond, omp_alt_alloc, succ_add_expr);
+
+	  /* Allocate memory for OpenACC declared variables.  */
+	  if (expr->symtree->n.sym->attr.oacc_declare_create)
+	    gfc_trans_oacc_declare_allocate (&se.pre, expr, true);
 	}
       else
 	{
@@ -7876,6 +7880,10 @@ gfc_trans_deallocate (gfc_code *code)
 
 	  if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (se.expr)))
 	    {
+	      if (!is_coarray
+		  && expr->symtree->n.sym->attr.oacc_declare_create)
+		gfc_trans_oacc_declare_allocate (&se.pre, expr, false);
+
 	      gfc_coarray_deregtype caf_dtype;
 
 	      if (is_coarray)
@@ -7929,6 +7937,10 @@ gfc_trans_deallocate (gfc_code *code)
 	}
       else
 	{
+	  /* Deallocate memory for OpenACC declared variables.  */
+	  if (expr->symtree->n.sym->attr.oacc_declare_create)
+	    gfc_trans_oacc_declare_allocate (&se.pre, expr, false);
+
 	  tmp = gfc_deallocate_scalar_with_status (se.expr, pstat, label_finish,
 						   false, al->expr,
 						   al->expr->ts, NULL_TREE,
diff --git a/gcc/fortran/trans.h b/gcc/fortran/trans.h
index ae7be9f..f4a562f 100644
--- a/gcc/fortran/trans.h
+++ b/gcc/fortran/trans.h
@@ -851,6 +851,7 @@ bool gfc_omp_private_debug_clause (tree, bool);
 bool gfc_omp_private_outer_ref (tree);
 struct gimplify_omp_ctx;
 void gfc_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *, tree);
+tree gfc_trans_oacc_declare_allocate (stmtblock_t *, gfc_expr *, bool);
 
 /* In trans-intrinsic.cc.  */
 void gfc_conv_intrinsic_mvbits (gfc_se *, gfc_actual_arglist *,
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index bd45816..17e5ca4 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -10755,6 +10755,8 @@ omp_group_base (omp_mapping_group *grp, unsigned int *chained,
     case GOMP_MAP_FIRSTPRIVATE_INT:
     case GOMP_MAP_USE_DEVICE_PTR:
     case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+    case GOMP_MAP_DECLARE_ALLOCATE:
+    case GOMP_MAP_DECLARE_DEALLOCATE:
       return NULL_TREE;
 
     case GOMP_MAP_FIRSTPRIVATE_POINTER:
@@ -14737,7 +14739,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 			     enum tree_code code)
 {
   struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
-  tree *orig_list_p = list_p;
+  tree *prev_list_p = NULL, *orig_list_p = list_p;
   tree c, decl;
   bool has_inscan_reductions = false;
 
@@ -15048,9 +15050,15 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	    case OMP_TARGET_DATA:
 	    case OMP_TARGET_ENTER_DATA:
 	    case OMP_TARGET_EXIT_DATA:
-	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
-		  || (OMP_CLAUSE_MAP_KIND (c)
-		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+	      if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+		   || (OMP_CLAUSE_MAP_KIND (c)
+		       == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+		  && !(prev_list_p
+		       && OMP_CLAUSE_CODE (*prev_list_p) == OMP_CLAUSE_MAP
+		       && ((OMP_CLAUSE_MAP_KIND (*prev_list_p)
+			    == GOMP_MAP_DECLARE_ALLOCATE)
+			   || (OMP_CLAUSE_MAP_KIND (*prev_list_p)
+			       == GOMP_MAP_DECLARE_DEALLOCATE))))
 		/* For target {,enter ,exit }data only the array slice is
 		   mapped, but not the pointer to it.  */
 		remove = true;
@@ -15392,6 +15400,19 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 		      == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)))
 	    move_attach = true;
 
+	  if (!remove && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
+	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
+	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
+	      && OMP_CLAUSE_CHAIN (c)
+	      && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP
+	      && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+		   == GOMP_MAP_ALWAYS_POINTER)
+		  || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+		      == GOMP_MAP_ATTACH_DETACH)
+		  || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+		      == GOMP_MAP_TO_PSET)))
+	    prev_list_p = list_p;
+
 	  break;
 
 	case OMP_CLAUSE_TO:
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 0188251..0d6f947 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1708,7 +1708,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
 	      && varpool_node::get_create (decl)->offloadable
 	      && !lookup_attribute ("omp declare target link",
-				    DECL_ATTRIBUTES (decl)))
+				    DECL_ATTRIBUTES (decl))
+	      && !is_gimple_omp_oacc (ctx->stmt))
 	    break;
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER)
@@ -12875,7 +12876,7 @@ convert_to_firstprivate_int (tree var, gimple_seq *gs)
 {
   tree type = TREE_TYPE (var), new_type = NULL_TREE;
 
-  if (omp_privatize_by_reference (var))
+  if (omp_privatize_by_reference (var) || POINTER_TYPE_P (type))
     {
       type = TREE_TYPE (type);
       tree tmp = create_tmp_var (type);
@@ -12900,7 +12901,8 @@ convert_to_firstprivate_int (tree var, gimple_seq *gs)
 /* Like convert_to_firstprivate_int, but restore the original type.  */
 
 static tree
-convert_from_firstprivate_int (tree var, bool is_ref, gimple_seq *gs)
+convert_from_firstprivate_int (tree var, tree orig_type, bool is_ref,
+			       gimple_seq *gs)
 {
   tree type = TREE_TYPE (var);
   tree new_type = NULL_TREE;
@@ -12909,7 +12911,32 @@ convert_from_firstprivate_int (tree var, bool is_ref, gimple_seq *gs)
   gcc_assert (TREE_CODE (var) == MEM_REF);
   var = TREE_OPERAND (var, 0);
 
-  if (INTEGRAL_TYPE_P (var) || POINTER_TYPE_P (type))
+  if (is_ref || POINTER_TYPE_P (orig_type))
+    {
+      tree_code code = NOP_EXPR;
+
+      if (TREE_CODE (type) == REAL_TYPE || TREE_CODE (type) == COMPLEX_TYPE
+	  || VECTOR_TYPE_P (type))
+	code = VIEW_CONVERT_EXPR;
+
+      if (code == VIEW_CONVERT_EXPR
+	  && TYPE_SIZE (type) != TYPE_SIZE (orig_type))
+	{
+	  tree ptype = build_pointer_type (type);
+	  var = fold_build1 (code, ptype, build_fold_addr_expr (var));
+	  var = build_simple_mem_ref (var);
+	}
+      else
+	var = fold_build1 (code, type, var);
+
+      tree inst = create_tmp_var (type);
+      gimplify_assign (inst, var, gs);
+      var = build_fold_addr_expr (inst);
+
+      return var;
+    }
+
+  if (INTEGRAL_TYPE_P (var))
     return fold_convert (type, var);
 
   gcc_assert (tree_to_uhwi (TYPE_SIZE (type)) <= POINTER_SIZE);
@@ -12920,16 +12947,8 @@ convert_from_firstprivate_int (tree var, bool is_ref, gimple_seq *gs)
   tmp = create_tmp_var (new_type);
   var = fold_convert (new_type, var);
   gimplify_assign (tmp, var, gs);
-  var = fold_build1 (VIEW_CONVERT_EXPR, type, tmp);
-
-  if (is_ref)
-    {
-      tmp = create_tmp_var (build_pointer_type (type));
-      gimplify_assign (tmp, build_fold_addr_expr (var), gs);
-      var = tmp;
-    }
 
-  return var;
+  return fold_build1 (VIEW_CONVERT_EXPR, type, tmp);
 }
 
 /* Lower the GIMPLE_OMP_TARGET in the current statement
@@ -13087,6 +13106,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case GOMP_MAP_NONCONTIG_ARRAY_ALLOC:
 	  case GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC:
 	  case GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT:
+	  case GOMP_MAP_DECLARE_ALLOCATE:
+	  case GOMP_MAP_DECLARE_DEALLOCATE:
 	  case GOMP_MAP_LINK:
 	  case GOMP_MAP_FORCE_DETACH:
 	    gcc_assert (is_gimple_omp_oacc (stmt));
@@ -13191,7 +13212,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		&& !maybe_lookup_field_in_outer_ctx (var, ctx))
 	      {
 		gcc_assert (is_gimple_omp_oacc (ctx->stmt));
-		x = convert_from_firstprivate_int (x,
+		x = convert_from_firstprivate_int (x, TREE_TYPE (new_var),
 						   omp_privatize_by_reference (var),
 						   &fplist);
 		gimplify_assign (new_var, x, &fplist);
@@ -13209,13 +13230,20 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		gcc_assert (is_gimple_omp_oacc (ctx->stmt));
 		if (omp_privatize_by_reference (new_var)
 		    && (TREE_CODE (var_type) != POINTER_TYPE
-		        || DECL_BY_REFERENCE (var)))
+			|| DECL_BY_REFERENCE (var))
+		    /* Accelerators may not have alloca, so it's not
+		       possible to privatize local storage for those
+		       objects.  */
+		    && TREE_CONSTANT (TYPE_SIZE (TREE_TYPE (var_type))))
 		  {
 		    /* Create a local object to hold the instance
 		       value.  */
 		    const char *id = IDENTIFIER_POINTER (DECL_NAME (new_var));
 		    tree inst = create_tmp_var (TREE_TYPE (var_type), id);
-		    gimplify_assign (inst, fold_indirect_ref (x), &fplist);
+		    if (TREE_CODE (var_type) == POINTER_TYPE)
+		      gimplify_assign (inst, x, &fplist);
+		    else
+		      gimplify_assign (inst, fold_indirect_ref (x), &fplist);
 		    x = build_fold_addr_expr (inst);
 		  }
 		gimplify_assign (new_var, x, &fplist);
@@ -13659,9 +13687,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
 		  {
 		    gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
+		    tree new_var = lookup_decl (var, ctx);
 		    tree type = TREE_TYPE (var);
 		    tree inner_type
-		      = omp_privatize_by_reference (var)
+		      = omp_privatize_by_reference (new_var)
 			? TREE_TYPE (type) : type;
 		    if ((FLOAT_TYPE_P (inner_type)
 			 || ANY_INTEGRAL_TYPE_P (inner_type))
diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-3.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-3.f95
index 9127cba..c94f515 100644
--- a/gcc/testsuite/gfortran.dg/goacc/declare-3.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/declare-3.f95
@@ -38,8 +38,7 @@ program test
   use mod_b
   use mod_d
   use mod_e
-
-  ! { dg-final { scan-tree-dump {(?n)#pragma acc data map\(force_alloc:d\) map\(force_to:b\) map\(force_alloc:a\)$} original } }
+  ! { dg-final { scan-tree-dump {(?n)#pragma acc data map\(force_alloc:d\) map\(to:b\) map\(alloc:a\)$} original } }
 end program test
 
 ! { dg-final { scan-tree-dump-times {#pragma acc data} 1 original } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-allocatable-1.f90 b/gcc/testsuite/gfortran.dg/goacc/declare-allocatable-1.f90
new file mode 100644
index 0000000..5349e0d
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/declare-allocatable-1.f90
@@ -0,0 +1,25 @@
+! Verify that OpenACC declared allocatable arrays have implicit
+! OpenACC enter and exit pragmas at the time of allocation and
+! deallocation.
+
+! { dg-additional-options "-fdump-tree-original" }
+
+program allocate
+  implicit none
+  integer, allocatable :: a(:), b
+  integer, parameter :: n = 100
+  integer i
+  !$acc declare create(a,b)
+
+  allocate (a(n), b)
+
+  !$acc parallel loop copyout(a, b)
+  do i = 1, n
+     a(i) = b
+  end do
+
+  deallocate (a, b)
+end program allocate
+
+! { dg-final { scan-tree-dump-times "pragma acc enter data map.declare_allocate" 2 "original" } }
+! { dg-final { scan-tree-dump-times "pragma acc exit data map.declare_deallocate" 2 "original" } }
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 662236f..9525733 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1086,6 +1086,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	case GOMP_MAP_LINK:
 	  pp_string (pp, "link");
 	  break;
+	case GOMP_MAP_DECLARE_ALLOCATE:
+	  pp_string (pp, "declare_allocate");
+	  break;
+	case GOMP_MAP_DECLARE_DEALLOCATE:
+	  pp_string (pp, "declare_deallocate");
+	  break;
 	case GOMP_MAP_ATTACH:
 	  pp_string (pp, "attach");
 	  break;