diff mbox

[gomp4] declare directive

Message ID 55CB9109.6030202@codesourcery.com
State New
Headers show

Commit Message

James Norris Aug. 12, 2015, 6:31 p.m. UTC 
Hi,

The attached patch is a revision of the functionality required to
handle the declare directive. The motivation for the change was that
the original code was deemed to fragile and difficult to maintain.
The new code is now smaller, in terms of line count, and hopefully,
easier to understand.  Given the significant amount of changes, I've
included a commentary on entire functions, rather than comments
about the specific changes. Hopefully, this will help in the
review process.

     gimple.c : gimplify_oacc_declare ()

         The first part of this function deals with handling
         the entry clauses. After scanning the clauses, we
         spin through the list and while adding the appropriate
         ones to the context variable list. Also if any
         'oacc declare' attributes are around, toss them.

         The next section deals with the return clauses.
         Again scan them, but after scanning them toss the
         context that was created as a result of the scan.
         We don't need the context as the context that was
         created with the entry clauses will suffice. After
         creation of the gimple statement, save it away so
         that it can be used at exit time.

     gimple.c : gimplify_body ()

         This is the place where the 'exit' clauses are
         used and inserted via a gimple statement. This
         is also the place where the omp context is
         'unwound'.

     fortran/trans-decl.c: finish_oacc_declare ()

         This function is called to handle the declare
         directives. The information that was contained
         in the directives during parsing of the variable
         section is hanging off the namespace (ns).

         The primary function here is to create the two
         sets of derived clauses and associate them with
         a newly created statement, i.e, gfc_code. This
         new statement is then inserted at the beginning
         of the statement chain associated with 'ns'.

     c/c-parser.c c_parser_oacc_declare ()

         The initial section of the function is doing
         syntax checking. If no errors are found and
         the variable is 'global', then an attribute
         is added to the variable. This is followed by
         the creation of the derived clauses.

         In the final section, when a global is at hand,
         a constructor is created. The constructor will
         pass on information to the runtime. On the other
         hand if the variable is local, then a statement
         will be inserted and created that will contain
         both sets of derived clauses.

     cp/parser.c: cp_parser_oacc_declare ()

         The initial section of the function is doing
         syntax checking. If no errors are found and
         the variable is 'global', then an attribute
         is added to the variable. This is followed by
         the creation of the derived clauses.

         In the final section, when a global is at hand,
         a constructor is created. The constructor will
         pass on information to the runtime. On the other
         hand if the variable is local, then a statement
         will be inserted and created that will contain
         both sets of derived clauses.

I'll commit this to gomp-4_0-branch sometime during the week of 17 August.

Thanks!
Jim
diff mbox

Patch

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index b452235..3f3e8c0 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -1452,319 +1452,6 @@  c_parser_external_declaration (c_parser *parser)
     }
 }
 
-static tree
-check_oacc_vars_1 (tree *tp, int *, void *l)
-{
-  if (TREE_CODE (*tp) == VAR_DECL && TREE_PUBLIC (*tp))
-    {
-      location_t loc = DECL_SOURCE_LOCATION (*tp);
-      tree attrs;
-      attrs = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (*tp));
-      if (attrs)
-	{
-	  tree t;
-
-	  for (t = TREE_VALUE (attrs); t; t = TREE_CHAIN (t))
-	    {
-	      loc = EXPR_LOCATION ((tree) l);
-
-	      if (OMP_CLAUSE_MAP_KIND (TREE_VALUE (t)) == GOMP_MAP_LINK)
-		{
-		  error_at (loc, "%<link%> clause cannot be used with %qE",
-			    *tp);
-		  break;
-		}
-	    }
-	}
-      else
-	error_at (loc, "no %<#pragma acc declare%> for %qE", *tp);
-    }
-  return NULL_TREE;
-}
-
-static tree
-check_oacc_vars (tree *tp, int *, void *)
-{
-  if (TREE_CODE (*tp) == STATEMENT_LIST)
-    {
-      tree_stmt_iterator i;
-
-      for (i = tsi_start (*tp); !tsi_end_p (i); tsi_next (&i))
-	{
-	  tree t = tsi_stmt (i);
-	  walk_tree_without_duplicates (&t, check_oacc_vars_1, t);
-	}
-    }
-
-  return NULL_TREE;
-}
-
-static struct oacc_return
-{
-  tree_stmt_iterator iter;
-  tree stmt;
-  int op;
-  struct oacc_return *next;
-} *oacc_returns;
-
-static tree
-find_oacc_return (tree *tp, int *, void *)
-{
-  if (TREE_CODE (*tp) == STATEMENT_LIST)
-    {
-      tree_stmt_iterator i;
-
-      for (i = tsi_start (*tp); !tsi_end_p (i); tsi_next (&i))
-	{
-	  tree t;
-	  struct oacc_return *r;
-
-	  t = tsi_stmt (i);
-
-	  if (TREE_CODE (t) == RETURN_EXPR)
-	    {
-	      r = XNEW (struct oacc_return);
-	      r->iter = i;
-	      r->stmt = NULL_TREE;
-	      r->op = 1;
-	      r->next = NULL;
-
-	      if (oacc_returns)
-		r->next = oacc_returns;
-
-	      oacc_returns = r;
-	    }
-	  else if (TREE_CODE (t) == COND_EXPR)
-	    {
-	       bool op1, op2;
-	       tree op;
-
-	       op1 = op2 = false;
-
-	       op = TREE_OPERAND (t, 1);
-	       op1 = (op && TREE_CODE (op) == RETURN_EXPR);
-
-	       op = TREE_OPERAND (t, 2);
-	       op2 = (op && TREE_CODE (op) == RETURN_EXPR);
-
-	       if (op1 || op2)
-		{
-		  r = XNEW (struct oacc_return);
-		  r->stmt = t;
-		  r->op = op1 ? 1 : 2;
-		  r->next = NULL;
-
-		  if (oacc_returns)
-		    r->next = oacc_returns;
-
-		  oacc_returns = r;
-		}
-	    }
-	}
-    }
-
-  return NULL_TREE;
-}
-
-static void
-finish_oacc_declare (tree fnbody, tree decls)
-{
-  tree t, stmt, body, c, ret_clauses, clauses;
-  location_t loc;
-  tree_stmt_iterator i;
-  tree fndecl = current_function_decl;
-
-  if (lookup_attribute ("oacc function", DECL_ATTRIBUTES (fndecl)))
-    {
-      if (lookup_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl)))
-	{
-	  location_t loc = DECL_SOURCE_LOCATION (fndecl);
-	  error_at (loc, "%<#pragma acc declare%> not allowed in %qE", fndecl);
-	}
-
-      walk_tree_without_duplicates (&fnbody, check_oacc_vars, NULL);
-      return;
-    }
-
-  if (!decls)
-    return;
-
-  body = BIND_EXPR_BODY (fnbody);
-
-  if (TREE_CODE (body) != STATEMENT_LIST)
-    {
-      tree list;
-
-      list = alloc_stmt_list ();
-      append_to_statement_list (body, &list);
-      BIND_EXPR_BODY (fnbody) = list;
-      body = list;
-    }
-
-  walk_tree_without_duplicates (&body, find_oacc_return, NULL);
-
-  clauses = NULL_TREE;
-
-  for (t = decls; t; t = TREE_CHAIN (t))
-    {
-      c = TREE_VALUE (TREE_VALUE (t));
-
-      if (clauses)
-	OMP_CLAUSE_CHAIN (c) = clauses;
-      else
-	loc = OMP_CLAUSE_LOCATION (c);
-
-      clauses = c;
-    }
-
-  ret_clauses = NULL_TREE;
-
-  for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
-    {
-      bool ret = false;
-      HOST_WIDE_INT kind, new_op;
-
-      kind = OMP_CLAUSE_MAP_KIND (c);
-
-      switch (kind)
-	{
-	  case GOMP_MAP_ALLOC:
-	  case GOMP_MAP_FORCE_ALLOC:
-	  case GOMP_MAP_FORCE_TO:
-	    new_op = GOMP_MAP_FORCE_DEALLOC;
-	    ret = true;
-	    break;
-
-	  case GOMP_MAP_FORCE_FROM:
-	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
-	    new_op = GOMP_MAP_FORCE_FROM;
-	    ret = true;
-	    break;
-
-	  case GOMP_MAP_FORCE_TOFROM:
-	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_TO);
-	    new_op = GOMP_MAP_FORCE_FROM;
-	    ret = true;
-	    break;
-
-	  case GOMP_MAP_FROM:
-	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
-	    new_op = GOMP_MAP_FROM;
-	    ret = true;
-	    break;
-
-	  case GOMP_MAP_TOFROM:
-	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
-	    new_op = GOMP_MAP_FROM;
-	    ret = true;
-	    break;
-
-	  case GOMP_MAP_DEVICE_RESIDENT:
-	  case GOMP_MAP_FORCE_DEVICEPTR:
-	  case GOMP_MAP_FORCE_PRESENT:
-	  case GOMP_MAP_LINK:
-	  case GOMP_MAP_POINTER:
-	  case GOMP_MAP_TO:
-	    break;
-
-	  default:
-	    gcc_unreachable ();
-	    break;
-	}
-
-      if (ret)
-	{
-	  t = copy_node (c);
-
-	  OMP_CLAUSE_SET_MAP_KIND (t, new_op);
-
-	  if (ret_clauses)
-	    OMP_CLAUSE_CHAIN (t) = ret_clauses;
-
-	  ret_clauses = t;
-	}
-    }
-
-  if (clauses)
-    {
-      bool found = false;
-
-      stmt = make_node (OACC_DECLARE);
-      TREE_TYPE (stmt) = void_type_node;
-      OACC_DECLARE_CLAUSES (stmt) = clauses;
-      SET_EXPR_LOCATION (stmt, loc);
-
-      c = OMP_CLAUSE_DECL (TREE_VALUE (TREE_VALUE (decls)));
-
-      for (i = tsi_start (body); !tsi_end_p (i); tsi_next (&i))
-	{
-	  tree it;
-
-	  it = tsi_stmt (i);
-
-	  if ((TREE_CODE (it) == DECL_EXPR) && (DECL_EXPR_DECL (it) == c))
-	    {
-	      tsi_link_after (&i, stmt, TSI_CONTINUE_LINKING);
-	      found = true;
-	      break;
-	    }
-	}
-
-	if (!found)
-	  {
-	    i = tsi_start (body);
-	    tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
-	  }
-    }
-
-    while (oacc_returns)
-      {
-	struct oacc_return *r;
-
-	stmt = make_node (OACC_DECLARE);
-	TREE_TYPE (stmt) = void_type_node;
-	OACC_DECLARE_CLAUSES (stmt) = ret_clauses;
-	SET_EXPR_LOCATION (stmt, loc);
-
-	r = oacc_returns;
-	if (r->stmt)
-	  {
-	    tree l;
-
-	    l = alloc_stmt_list ();
-	    append_to_statement_list (stmt, &l);
-	    stmt = TREE_OPERAND (r->stmt, r->op);
-	    append_to_statement_list (stmt, &l);
-	    TREE_OPERAND (r->stmt, r->op) = l;
-	  }
-	else
-	  tsi_link_before (&r->iter, stmt, TSI_CONTINUE_LINKING);
-
-	oacc_returns = r->next;
-	free (r);
-     }
-
-  for (i = tsi_start (body); !tsi_end_p (i); tsi_next (&i))
-    {
-      if (tsi_end_p (i))
-	break;
-    }
-
-  if (ret_clauses)
-    {
-      stmt = make_node (OACC_DECLARE);
-      TREE_TYPE (stmt) = void_type_node;
-      OACC_DECLARE_CLAUSES (stmt) = ret_clauses;
-      SET_EXPR_LOCATION (stmt, loc);
-
-      tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
-    }
-
-  DECL_ATTRIBUTES (fndecl)
-	  = remove_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl));
-}
-
-
 static void c_finish_omp_declare_simd (c_parser *, tree, tree, vec<c_token>);
 static void c_finish_oacc_routine (c_parser *, tree, tree, bool);
 
@@ -2312,9 +1999,6 @@  c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
       fnbody = c_parser_compound_statement (parser);
       if (flag_cilkplus && contains_array_notation_expr (fnbody))
 	fnbody = expand_array_notation_exprs (fnbody);
-      tree decls = lookup_attribute ("oacc declare",
-				     DECL_ATTRIBUTES (current_function_decl));
-      finish_oacc_declare (fnbody, decls);
       if (nested)
 	{
 	  tree decl = current_function_decl;
@@ -12770,7 +12454,8 @@  static void
 c_parser_oacc_declare (c_parser *parser)
 {
   location_t pragma_loc = c_parser_peek_token (parser)->location;
-  tree clauses;
+  tree c, clauses, ret_clauses, stmt, t;
+
   bool error = false;
 
   c_parser_consume_pragma (parser);
@@ -12783,7 +12468,8 @@  c_parser_oacc_declare (c_parser *parser)
 		"no valid clauses specified in %<#pragma acc declare%>");
       return;
     }
-  for (tree t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
+
+  for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
     {
       location_t loc = OMP_CLAUSE_LOCATION (t);
       tree decl = OMP_CLAUSE_DECL (t);
@@ -12849,60 +12535,27 @@  c_parser_oacc_declare (c_parser *parser)
 	  break;
 	}
 
-      /* Store the clause in an attribute on the variable, at file
-	 scope, or the function, at block scope.  */
-      tree decl_for_attr;
-      if (global_bindings_p ())
+      tree decl_for_attr = decl;
+      tree prev_attr = lookup_attribute ("oacc declare",
+					 DECL_ATTRIBUTES (decl));
+      if (prev_attr)
 	{
-	  decl_for_attr = decl;
-	  tree prev_attr = lookup_attribute ("oacc declare",
-					     DECL_ATTRIBUTES (decl));
-	  if (prev_attr)
-	    {
-	      tree p = TREE_VALUE (prev_attr);
-	      tree cl = TREE_VALUE (p);
+	  tree p = TREE_VALUE (prev_attr);
+	  tree cl = TREE_VALUE (p);
 
-	      if (!devres
-		  && OMP_CLAUSE_MAP_KIND (cl) != GOMP_MAP_DEVICE_RESIDENT)
-		{
-		  error_at (loc,
-			    "variable %qD used more than once with "
-			    "%<#pragma acc declare%>", decl);
-		  inform (OMP_CLAUSE_LOCATION (cl),
-			  "previous directive was here");
-		  error = true;
-		  continue;
-		}
-	    }
-	}
-      else
-	{
-	  decl_for_attr = current_function_decl;
-	  tree prev_attr = lookup_attribute ("oacc declare",
-					     DECL_ATTRIBUTES (decl_for_attr));
-	  for (;
-	       prev_attr;
-	       prev_attr = lookup_attribute ("oacc declare",
-					     TREE_CHAIN (prev_attr)))
+	  if (!devres && OMP_CLAUSE_MAP_KIND (cl) != GOMP_MAP_DEVICE_RESIDENT)
 	    {
-	      tree p = TREE_VALUE (prev_attr);
-	      tree cl = TREE_VALUE (p);
-	      if (OMP_CLAUSE_DECL (cl) == decl)
-		{
-		  error_at (loc,
-			    "variable %qD used more than once with "
-			    "%<#pragma acc declare%>", decl);
-		  inform (OMP_CLAUSE_LOCATION (cl),
-			  "previous directive was here");
-		  error = true;
-		  break;
-		}
+	      error_at (loc, "variable %qD used more than once with "
+			     "%<#pragma acc declare%>", decl);
+	      inform (OMP_CLAUSE_LOCATION (cl), "previous directive was here");
+	      error = true;
+	      continue;
 	    }
 	}
 
-      if (!error)
+      if (!error && global_bindings_p ())
 	{
-	  tree attr = tree_cons (NULL_TREE, t, NULL_TREE);
+	  tree attr = tree_cons (NULL_TREE, clauses, NULL_TREE);
 	  tree attrs = tree_cons (get_identifier ("oacc declare"),
 				  attr, NULL_TREE);
 	  decl_attributes (&decl_for_attr, attrs, 0);
@@ -12912,6 +12565,74 @@  c_parser_oacc_declare (c_parser *parser)
   if (error)
     return;
 
+  ret_clauses = NULL_TREE;
+
+  for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+    {
+      bool ret = false;
+      HOST_WIDE_INT kind, new_op;
+
+      kind = OMP_CLAUSE_MAP_KIND (c);
+
+      switch (kind)
+	{
+	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_FORCE_ALLOC:
+	  case GOMP_MAP_FORCE_TO:
+	    new_op = GOMP_MAP_FORCE_DEALLOC;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_FORCE_FROM:
+	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
+	    new_op = GOMP_MAP_FORCE_FROM;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_FORCE_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_TO);
+	    new_op = GOMP_MAP_FORCE_FROM;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_FROM:
+	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
+	    new_op = GOMP_MAP_FROM;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+	    new_op = GOMP_MAP_FROM;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_DEVICE_RESIDENT:
+	  case GOMP_MAP_FORCE_DEVICEPTR:
+	  case GOMP_MAP_FORCE_PRESENT:
+	  case GOMP_MAP_LINK:
+	  case GOMP_MAP_POINTER:
+	  case GOMP_MAP_TO:
+	    break;
+
+	  default:
+	    gcc_unreachable ();
+	    break;
+	}
+
+      if (ret)
+	{
+	  t = copy_node (c);
+
+	  OMP_CLAUSE_SET_MAP_KIND (t, new_op);
+
+	  if (ret_clauses)
+	    OMP_CLAUSE_CHAIN (t) = ret_clauses;
+
+	  ret_clauses = t;
+	}
+    }
+
   if (global_bindings_p ())
     {
       char buf[128];
@@ -12971,6 +12692,16 @@  c_parser_oacc_declare (c_parser *parser)
 
       finish_function ();
     }
+  else
+    {
+      stmt = make_node (OACC_DECLARE);
+      TREE_TYPE (stmt) = void_type_node;
+      OACC_DECLARE_CLAUSES (stmt) = clauses;
+      OACC_DECLARE_RETURN_CLAUSES (stmt) = ret_clauses;
+      SET_EXPR_LOCATION (stmt, pragma_loc);
+
+      add_stmt (stmt);
+    }
 }
 
 /* OpenACC 2.0:
diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c
index 1d80ef2..069db46 100644
--- a/gcc/cp/decl.c
+++ b/gcc/cp/decl.c
@@ -14114,332 +14114,6 @@  maybe_save_function_definition (tree fun)
     register_constexpr_fundef (fun, DECL_SAVED_TREE (fun));
 }
 
-static tree
-check_oacc_vars_1 (tree *tp, int *, void *l)
-{
-  if (TREE_CODE (*tp) == VAR_DECL && TREE_PUBLIC (*tp))
-    {
-      location_t loc = DECL_SOURCE_LOCATION (*tp);
-      tree attrs;
-      attrs = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (*tp));
-      if (attrs)
-	{
-	  tree t;
-
-	  for (t = TREE_VALUE (attrs); t; t = TREE_CHAIN (t))
-	    {
-	      loc = EXPR_LOCATION ((tree) l);
-
-	      if (OMP_CLAUSE_MAP_KIND (TREE_VALUE (t)) == GOMP_MAP_LINK)
-		{
-		  error_at (loc, "%<link%> clause cannot be used with %qE",
-			    *tp);
-		  break;
-		}
-	    }
-	}
-      else
-	error_at (loc, "no %<#pragma acc declare%> for %qE", *tp);
-    }
-  return NULL_TREE;
-}
-
-static tree
-check_oacc_vars (tree *tp, int *, void *)
-{
-  if (TREE_CODE (*tp) == STATEMENT_LIST)
-    {
-      tree_stmt_iterator i;
-
-      for (i = tsi_start (*tp); !tsi_end_p (i); tsi_next (&i))
-	{
-	  tree t = tsi_stmt (i);
-	  walk_tree_without_duplicates (&t, check_oacc_vars_1, t);
-	}
-    }
-
-  return NULL_TREE;
-}
-
-static struct oacc_return
-{
-  tree_stmt_iterator iter;
-  tree stmt;
-  int op;
-  struct oacc_return *next;
-} *oacc_returns;
-
-static tree
-find_oacc_return (tree *tp, int *, void *)
-{
-  if (TREE_CODE (*tp) == STATEMENT_LIST)
-    {
-      tree_stmt_iterator i;
-
-      for (i = tsi_start (*tp); !tsi_end_p (i); tsi_next (&i))
-	{
-	  tree t;
-	  struct oacc_return *r;
-
-	  t = tsi_stmt (i);
-
-	  if (TREE_CODE (t) == RETURN_EXPR)
-	    {
-	      r = XNEW (struct oacc_return);
-	      r->iter = i;
-	      r->stmt = NULL_TREE;
-	      r->op = 1;
-	      r->next = NULL;
-
-	      if (oacc_returns)
-		r->next = oacc_returns;
-
-	      oacc_returns = r;
-	    }
-	  else if (TREE_CODE (t) == IF_STMT)
-	    {
-	       bool op1, op2;
-	       tree op;
-
-	       op1 = op2 = false;
-
-	       op = TREE_OPERAND (t, 1);
-	       op1 = (op && TREE_CODE (op) == RETURN_EXPR);
-
-	       op = TREE_OPERAND (t, 2);
-	       op2 = (op && TREE_CODE (op) == RETURN_EXPR);
-
-	       if (op1 || op2)
-		{
-		  r = XNEW (struct oacc_return);
-		  r->stmt = t;
-		  r->op = op1 ? 1 : 2;
-		  r->next = NULL;
-
-		  if (oacc_returns)
-		    r->next = oacc_returns;
-
-		  oacc_returns = r;
-		}
-	    }
-	}
-    }
-
-  return NULL_TREE;
-}
-
-static void
-finish_oacc_declare (tree fndecl)
-{
-  tree t, stmt, list, c, ret_clauses, clauses, decls;
-  location_t loc;
-  tree_stmt_iterator i;
-
-  if (DECL_USE_TEMPLATE (fndecl))
-    return;
-
-  list = cur_stmt_list;
-
-  decls = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl));
-
-  if (lookup_attribute ("oacc function", DECL_ATTRIBUTES (fndecl)))
-    {
-      if (decls)
-	{
-	  location_t loc = DECL_SOURCE_LOCATION (fndecl);
-	  error_at (loc, "%<#pragma acc declare%> not allowed in %qE", fndecl);
-	}
-
-      walk_tree_without_duplicates (&list, check_oacc_vars, NULL);
-      return;
-    }
-
-  if (!decls)
-    return;
-
-  walk_tree_without_duplicates (&list, find_oacc_return, NULL);
-
-  clauses = NULL_TREE;
-
-  for (t = decls; t; t = TREE_CHAIN (t))
-    {
-      c = TREE_VALUE (TREE_VALUE (t));
-
-      if (clauses)
-	OMP_CLAUSE_CHAIN (c) = clauses;
-      else
-	loc = OMP_CLAUSE_LOCATION (c);
-
-      clauses = c;
-    }
-
-  ret_clauses = NULL_TREE;
-
-  for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
-    {
-      bool ret = false;
-      HOST_WIDE_INT kind, new_op;
-
-      kind = OMP_CLAUSE_MAP_KIND (c);
-
-      switch (kind)
-	{
-	  case GOMP_MAP_ALLOC:
-	  case GOMP_MAP_FORCE_ALLOC:
-	  case GOMP_MAP_FORCE_TO:
-	    new_op = GOMP_MAP_FORCE_DEALLOC;
-	    ret = true;
-	    break;
-
-	  case GOMP_MAP_FORCE_FROM:
-	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
-	    new_op = GOMP_MAP_FORCE_FROM;
-	    ret = true;
-	    break;
-
-	  case GOMP_MAP_FORCE_TOFROM:
-	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_TO);
-	    new_op = GOMP_MAP_FORCE_FROM;
-	    ret = true;
-	    break;
-
-	  case GOMP_MAP_FROM:
-	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
-	    new_op = GOMP_MAP_FROM;
-	    ret = true;
-	    break;
-
-	  case GOMP_MAP_TOFROM:
-	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
-	    new_op = GOMP_MAP_FROM;
-	    ret = true;
-	    break;
-
-	  case GOMP_MAP_DEVICE_RESIDENT:
-	  case GOMP_MAP_FORCE_DEVICEPTR:
-	  case GOMP_MAP_FORCE_PRESENT:
-	  case GOMP_MAP_POINTER:
-	  case GOMP_MAP_TO:
-	    break;
-
-	  case GOMP_MAP_LINK:
-	    continue;
-
-	  default:
-	    gcc_unreachable ();
-	    break;
-	}
-
-      if (ret)
-	{
-	  t = copy_node (c);
-
-	  OMP_CLAUSE_SET_MAP_KIND (t, new_op);
-
-	  if (ret_clauses)
-	    OMP_CLAUSE_CHAIN (t) = ret_clauses;
-
-	  ret_clauses = t;
-	}
-    }
-
-   i = tsi_start (list);
-   if (!tsi_end_p (i))
-     {
-	t = tsi_stmt (i);
-	if (TREE_CODE (t) == BIND_EXPR)
-	  {
-	    list = BIND_EXPR_BODY (t);
-	    if (TREE_CODE (list) != STATEMENT_LIST)
-	      {
-		stmt = list;
-		list = alloc_stmt_list ();
-		BIND_EXPR_BODY (t) = list;
-		i = tsi_start (list);
-		tsi_link_after (&i, stmt, TSI_CONTINUE_LINKING);
-	      }
-	  }
-      }
-
-  if (clauses)
-    {
-      bool found = false;
-
-      stmt = make_node (OACC_DECLARE);
-      TREE_TYPE (stmt) = void_type_node;
-      OMP_STANDALONE_CLAUSES (stmt) = clauses;
-      SET_EXPR_LOCATION (stmt, loc);
-
-      c = OMP_CLAUSE_DECL (TREE_VALUE (TREE_VALUE (decls)));
-
-      for (i = tsi_start (list); !tsi_end_p (i); tsi_next (&i))
-	{
-	  tree it;
-
-	  it = tsi_stmt (i);
-
-	  if ((TREE_CODE (it) == DECL_EXPR) && (DECL_EXPR_DECL (it) == c))
-	    {
-	      tsi_link_after (&i, stmt, TSI_CONTINUE_LINKING);
-	      found = true;
-	      break;
-	    }
-	}
-
-      if (!found)
-	{
-	  i = tsi_start (list);
-	  tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
-	}
-    }
-
-    while (oacc_returns)
-      {
-	struct oacc_return *r;
-
-	stmt = make_node (OACC_DECLARE);
-	TREE_TYPE (stmt) = void_type_node;
-	OMP_STANDALONE_CLAUSES (stmt) = ret_clauses;
-	SET_EXPR_LOCATION (stmt, loc);
-
-	r = oacc_returns;
-	if (r->stmt)
-	  {
-	    tree l;
-
-	    l = alloc_stmt_list ();
-	    append_to_statement_list (stmt, &l);
-	    stmt = TREE_OPERAND (r->stmt, r->op);
-	    append_to_statement_list (stmt, &l);
-	    TREE_OPERAND (r->stmt, r->op) = l;
-	  }
-	else
-	  tsi_link_before (&r->iter, stmt, TSI_CONTINUE_LINKING);
-
-	oacc_returns = r->next;
-	free (r);
-     }
-
-  if (ret_clauses)
-    {
-      for (i = tsi_start (list); !tsi_end_p (i); tsi_next (&i))
-	{
-	  if (tsi_end_p (i))
-	    break;
-	}
-
-      stmt = make_node (OACC_DECLARE);
-      TREE_TYPE (stmt) = void_type_node;
-      OMP_STANDALONE_CLAUSES (stmt) = ret_clauses;
-      SET_EXPR_LOCATION (stmt, loc);
-
-      tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
-    }
-
-  DECL_ATTRIBUTES (fndecl)
-	  = remove_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl));
-}
-
 /* Finish up a function declaration and compile that function
    all the way to assembler language output.  The free the storage
    for the function definition.
@@ -14468,8 +14142,6 @@  finish_function (int flags)
   gcc_assert (!defer_mark_used_calls);
   defer_mark_used_calls = true;
 
-  finish_oacc_declare (fndecl);
-
   record_key_method_defined (fndecl);
 
   fntype = TREE_TYPE (fndecl);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 85236bf..14e7f8e 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -32179,12 +32179,14 @@  static int oacc_dcl_idx = 0;
 static tree
 cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 {
-  tree clauses;
+  tree c, clauses, ret_clauses, stmt, t;
   bool error = false;
 
+
   clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK,
 					"#pragma acc declare", pragma_tok);
 
+
   if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
     {
       error_at (pragma_tok->location,
@@ -32258,58 +32260,26 @@  cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 	  break;
 	}
 
-      /* Store the clause in an attribute on the variable, at file
-	 scope, or the function, at block scope.  */
-      tree decl_for_attr;
-      if (global_bindings_p ())
-	{
-	  decl_for_attr = decl;
-	  tree prev_attr = lookup_attribute ("oacc declare",
+      tree decl_for_attr = decl;
+      tree prev_attr = lookup_attribute ("oacc declare",
 					     DECL_ATTRIBUTES (decl));
-	  if (prev_attr)
-	    {
-	      tree p = TREE_VALUE (prev_attr);
-	      tree cl = TREE_VALUE (p);
-
-	      if (!devres
-		  && OMP_CLAUSE_MAP_KIND (cl) != GOMP_MAP_DEVICE_RESIDENT)
-		{
-		  error_at (loc,
-			    "variable %qD used more than once with "
-			    "%<#pragma acc declare%>", decl);
-		  inform (OMP_CLAUSE_LOCATION (TREE_VALUE (p)),
-			  "previous directive was here");
-		  error = true;
-		  continue;
-		}
-	    }
-	}
-      else
+      if (prev_attr)
 	{
-	  decl_for_attr = current_function_decl;
-	  tree prev_attr = lookup_attribute ("oacc declare",
-					     DECL_ATTRIBUTES (decl_for_attr));
-	  for (;
-	       prev_attr;
-	       prev_attr = lookup_attribute ("oacc declare",
-					     TREE_CHAIN (prev_attr)))
+	  tree p = TREE_VALUE (prev_attr);
+	  tree cl = TREE_VALUE (p);
+
+	  if (!devres && OMP_CLAUSE_MAP_KIND (cl) != GOMP_MAP_DEVICE_RESIDENT)
 	    {
-	      tree p = TREE_VALUE (prev_attr);
-	      tree cl = TREE_VALUE (p);
-	      if (OMP_CLAUSE_DECL (cl) == decl)
-		{
-		  error_at (loc,
-			    "variable %qD used more than once with "
-			    "%<#pragma acc declare%>", decl);
-		  inform (OMP_CLAUSE_LOCATION (cl),
-			  "previous directive was here");
-		  error = true;
-		  break;
-		}
+	      error_at (loc, "variable %qD used more than once with "
+			"%<#pragma acc declare%>", decl);
+	      inform (OMP_CLAUSE_LOCATION (TREE_VALUE (p)),
+		      "previous directive was here");
+	      error = true;
+	      continue;
 	    }
 	}
 
-      if (!error)
+      if (!error && global_bindings_p ())
 	{
 	  tree attr = tree_cons (NULL_TREE, t, NULL_TREE);
 	  tree attrs = tree_cons (get_identifier ("oacc declare"),
@@ -32321,6 +32291,76 @@  cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
   if (error)
     return NULL_TREE;
 
+  ret_clauses = NULL_TREE;
+
+  for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+    {
+      bool ret = false;
+      HOST_WIDE_INT kind, new_op;
+
+      kind = OMP_CLAUSE_MAP_KIND (c);
+
+      switch (kind)
+	{
+	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_FORCE_ALLOC:
+	  case GOMP_MAP_FORCE_TO:
+	    new_op = GOMP_MAP_FORCE_DEALLOC;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_FORCE_FROM:
+	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
+	    new_op = GOMP_MAP_FORCE_FROM;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_FORCE_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_TO);
+	    new_op = GOMP_MAP_FORCE_FROM;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_FROM:
+	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
+	    new_op = GOMP_MAP_FROM;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+	    new_op = GOMP_MAP_FROM;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_DEVICE_RESIDENT:
+	  case GOMP_MAP_FORCE_DEVICEPTR:
+	  case GOMP_MAP_FORCE_PRESENT:
+	  case GOMP_MAP_POINTER:
+	  case GOMP_MAP_TO:
+	    break;
+
+	  case GOMP_MAP_LINK:
+	    continue;
+
+	  default:
+	    gcc_unreachable ();
+	    break;
+	}
+
+      if (ret)
+	{
+	  t = copy_node (c);
+
+	  OMP_CLAUSE_SET_MAP_KIND (t, new_op);
+
+	  if (ret_clauses)
+	    OMP_CLAUSE_CHAIN (t) = ret_clauses;
+
+	  ret_clauses = t;
+	}
+    }
+
   if (global_bindings_p ())
     {
       char buf[128];
@@ -32375,6 +32415,16 @@  cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
       expand_or_defer_fn (finish_function (0));
       obstack_free (&declarator_obstack, p);
     }
+  else
+    {
+      stmt = make_node (OACC_DECLARE);
+      TREE_TYPE (stmt) = void_type_node;
+      OACC_DECLARE_CLAUSES (stmt) = clauses;
+      OACC_DECLARE_RETURN_CLAUSES (stmt) = ret_clauses;
+      SET_EXPR_LOCATION (stmt, pragma_tok->location);
+
+      add_stmt (stmt);
+    }
 
   return NULL_TREE;
 }
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 056b2c1..8ace93c 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -13907,6 +13907,7 @@  tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
 	       integral_constant_expression_p)
 
   tree stmt, tmp;
+tree s;
   tree r;
   location_t loc;
 
@@ -14396,8 +14397,18 @@  tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
       add_stmt (t);
       break;
 
-    case OMP_TARGET_UPDATE:
     case OACC_DECLARE:
+      t = copy_node (t);
+      tmp = tsubst_omp_clauses (OACC_DECLARE_CLAUSES (t), false,
+				args, complain, in_decl);
+      OACC_DECLARE_CLAUSES (t) = tmp;
+      tmp = tsubst_omp_clauses (OACC_DECLARE_RETURN_CLAUSES (t), false,
+				args, complain, in_decl);
+      OACC_DECLARE_RETURN_CLAUSES (t) = tmp;
+      add_stmt (t);
+      break;
+
+    case OMP_TARGET_UPDATE:
     case OACC_ENTER_DATA:
     case OACC_EXIT_DATA:
     case OACC_UPDATE:
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 1e87d6c..9d76db7 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1287,6 +1287,7 @@  typedef struct gfc_oacc_declare
   locus where;
   bool module_var;
   gfc_omp_clauses *clauses;
+  gfc_omp_clauses *return_clauses;
 }
 gfc_oacc_declare;
 #define gfc_get_oacc_declare() XCNEW (gfc_oacc_declare)
diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c
index bc54067..eee5340 100644
--- a/gcc/fortran/trans-decl.c
+++ b/gcc/fortran/trans-decl.c
@@ -5864,8 +5864,7 @@  void
 finish_oacc_declare (gfc_namespace *ns, enum sym_flavor flavor)
 {
   gfc_code *code, *end_c, *code2;
-  gfc_oacc_declare *oc;
-  gfc_omp_clauses *omp_clauses = NULL, *ret_clauses = NULL;
+  gfc_oacc_declare *oc, *new_oc;
   gfc_omp_namelist *n;
   locus where = gfc_current_locus;
 
@@ -5888,204 +5887,109 @@  finish_oacc_declare (gfc_namespace *ns, enum sym_flavor flavor)
 
   for (oc = ns->oacc_declare; oc; oc = oc->next)
     {
+      gfc_omp_clauses *omp_clauses, *ret_clauses;
+
       if (oc->module_var)
 	continue;
 
       if (oc->clauses)
 	{
-	  if (omp_clauses)
-	    {
-	      gfc_omp_namelist *p;
-
-	      for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
-		p = n;
-
-	      p->next = oc->clauses->lists[OMP_LIST_MAP];
-	    }
-	  else
-	    {
-	      omp_clauses = oc->clauses;
-	    }
-	}
-    }
-
-  while (ns->oacc_declare)
-    {
-      oc = ns->oacc_declare;
-      ns->oacc_declare = oc->next;
-      free (oc);
-    }
-
-  code = XCNEW (gfc_code);
-  code->op = EXEC_OACC_DECLARE;
-  code->loc = where;
-  code->ext.omp_clauses = omp_clauses;
-
-  for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
-    {
-      bool ret = false;
-      gfc_omp_map_op new_op;
-
-      switch (n->u.map_op)
-	{
-	case OMP_MAP_ALLOC:
-	case OMP_MAP_FORCE_ALLOC:
-	  new_op = OMP_MAP_FORCE_DEALLOC;
-	  ret = true;
-	  break;
-
-	case OMP_MAP_DEVICE_RESIDENT:
-	  n->u.map_op = OMP_MAP_FORCE_ALLOC;
-	  new_op = OMP_MAP_FORCE_DEALLOC;
-	  ret = true;
-	  break;
-
-	case OMP_MAP_FORCE_FROM:
-	  n->u.map_op = OMP_MAP_FORCE_ALLOC;
-	  new_op = OMP_MAP_FORCE_FROM;
-	  ret = true;
-	  break;
-
-	case OMP_MAP_FORCE_TO:
-	  new_op = OMP_MAP_FORCE_DEALLOC;
-	  ret = true;
-	  break;
-
-	case OMP_MAP_FORCE_TOFROM:
-	  n->u.map_op = OMP_MAP_FORCE_TO;
-	  new_op = OMP_MAP_FORCE_FROM;
-	  ret = true;
-	  break;
-
-	case OMP_MAP_FROM:
-	  n->u.map_op = OMP_MAP_FORCE_ALLOC;
-	  new_op = OMP_MAP_FROM;
-	  ret = true;
-	  break;
-
-	case OMP_MAP_FORCE_DEVICEPTR:
-	case OMP_MAP_FORCE_PRESENT:
-	case OMP_MAP_LINK:
-	case OMP_MAP_TO:
-	  break;
-
-	case OMP_MAP_TOFROM:
-	  n->u.map_op = OMP_MAP_TO;
-	  new_op = OMP_MAP_FROM;
-	  ret = true;
-	  break;
-
-	default:
-	  gcc_unreachable ();
-	  break;
-	}
-
-      if (ret)
-	{
-	  gfc_omp_namelist *new_n;
+	   code = XCNEW (gfc_code);
+	   code->op = EXEC_OACC_DECLARE;
+	   code->loc = where;
 
-	  new_n = gfc_get_omp_namelist ();
-	  new_n->sym = n->sym;
-	  new_n->u.map_op = new_op;
+	   ret_clauses = NULL;
+	   omp_clauses = oc->clauses;
 
-	  if (!ret_clauses)
-	    ret_clauses = gfc_get_omp_clauses ();
+	   for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
+	     {
+		bool ret = false;
+		gfc_omp_map_op new_op;
 
-	  if (ret_clauses->lists[OMP_LIST_MAP])
-	    new_n->next = ret_clauses->lists[OMP_LIST_MAP];
+		switch (n->u.map_op)
+		  {
+		    case OMP_MAP_ALLOC:
+		    case OMP_MAP_FORCE_ALLOC:
+		      new_op = OMP_MAP_FORCE_DEALLOC;
+		      ret = true;
+		      break;
 
-	  ret_clauses->lists[OMP_LIST_MAP] = new_n;
-	  ret = false;
-	}
-    }
+		    case OMP_MAP_DEVICE_RESIDENT:
+		      n->u.map_op = OMP_MAP_FORCE_ALLOC;
+		      new_op = OMP_MAP_FORCE_DEALLOC;
+		      ret = true;
+		      break;
 
-  if (!ret_clauses)
-    {
-      code->next = ns->code;
-      ns->code = code;
-      return;
-    }
+		    case OMP_MAP_FORCE_FROM:
+		      n->u.map_op = OMP_MAP_FORCE_ALLOC;
+		      new_op = OMP_MAP_FORCE_FROM;
+		      ret = true;
+		      break;
 
-  code2 = XCNEW (gfc_code);
-  code2->op = EXEC_OACC_DECLARE;
-  code2->loc = where;
-  code2->ext.omp_clauses = ret_clauses;
+		    case OMP_MAP_FORCE_TO:
+		      new_op = OMP_MAP_FORCE_DEALLOC;
+		      ret = true;
+		      break;
 
-  if (ns->code)
-    {
-      find_oacc_return (ns->code);
+		    case OMP_MAP_FORCE_TOFROM:
+		      n->u.map_op = OMP_MAP_FORCE_TO;
+		      new_op = OMP_MAP_FORCE_FROM;
+		      ret = true;
+		      break;
 
-      if (ns->code->op == EXEC_END_PROCEDURE)
-	{
-	  code2->next = ns->code;
-	  code->next = code2;
-	}
-      else
-	{
-	  end_c = find_end (ns->code);
-	  if (end_c)
-	    {
-	      code2->next = end_c->next;
-	      end_c->next = code2;
-	      code->next = ns->code;
-	    }
-	  else
-	    {
-	      gfc_code *last;
+		    case OMP_MAP_FROM:
+		      n->u.map_op = OMP_MAP_FORCE_ALLOC;
+		      new_op = OMP_MAP_FROM;
+		      ret = true;
+		      break;
 
-	      last = ns->code;
+		    case OMP_MAP_FORCE_DEVICEPTR:
+		    case OMP_MAP_FORCE_PRESENT:
+		    case OMP_MAP_LINK:
+		    case OMP_MAP_TO:
+		      break;
 
-	      while (last->next)
-		last = last->next;
+		    case OMP_MAP_TOFROM:
+		      n->u.map_op = OMP_MAP_TO;
+		      new_op = OMP_MAP_FROM;
+		      ret = true;
+		      break;
 
-	      last->next = code2;
-	      code->next = ns->code;
-	    }
-	}
-    }
-  else
-    {
-      code->next = code2;
-    }
+		    default:
+		      gcc_unreachable ();
+		      break;
+		  }
 
-  while (oacc_returns)
-    {
-      struct oacc_return *r;
+		if (ret)
+		  {
+		    gfc_omp_namelist *new_n;
 
-      r = oacc_returns;
+		    new_n = gfc_get_omp_namelist ();
+		    new_n->sym = n->sym;
+		    new_n->u.map_op = new_op;
 
-      ret_clauses = gfc_get_omp_clauses ();
+		    if (!ret_clauses)
+		      ret_clauses = gfc_get_omp_clauses ();
 
-      for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
-	{
-	  if (n->u.map_op == OMP_MAP_FORCE_ALLOC
-	      || n->u.map_op == OMP_MAP_FORCE_TO)
-	    {
-	      gfc_omp_namelist *new_n;
+		    if (ret_clauses->lists[OMP_LIST_MAP])
+		      new_n->next = ret_clauses->lists[OMP_LIST_MAP];
 
-	      new_n = gfc_get_omp_namelist ();
-	      new_n->sym = n->sym;
-	      new_n->u.map_op = OMP_MAP_FORCE_DEALLOC;
+		    ret_clauses->lists[OMP_LIST_MAP] = new_n;
+		    ret = false;
+		  }
+	     }
 
-	      if (ret_clauses->lists[OMP_LIST_MAP])
-		new_n->next = ret_clauses->lists[OMP_LIST_MAP];
+	   code->ext.oacc_declare = gfc_get_oacc_declare ();
+	   code->ext.oacc_declare->clauses = omp_clauses;
+	   code->ext.oacc_declare->return_clauses = ret_clauses;
 
-	      ret_clauses->lists[OMP_LIST_MAP] = new_n;
-	    }
+	   if (ns->code)
+	     code->next = ns->code;
+	   ns->code = code;
 	}
-
-      code2 = XCNEW (gfc_code);
-      code2->op = EXEC_OACC_DECLARE;
-      code2->loc = where;
-      code2->ext.omp_clauses = ret_clauses;
-      code2->next = r->code->next;
-      r->code->next = code2;
-
-      oacc_returns = r->next;
-      free (r);
     }
 
-    ns->code = code;
+  return;
 }
 
 
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index a6b8928..cd76f2a 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -4448,7 +4448,7 @@  tree
 gfc_trans_oacc_declare (gfc_code *code)
 {
   stmtblock_t block;
-  tree stmt, oacc_clauses;
+  tree stmt, c1, c2;
   enum tree_code construct_code;
 
   gfc_start_block (&block);
@@ -4456,11 +4456,15 @@  gfc_trans_oacc_declare (gfc_code *code)
   construct_code = OACC_DECLARE;
 
   gfc_start_block (&block);
-  oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
-					code->loc);
-  stmt = build1_loc (input_location, construct_code, void_type_node,
-		     oacc_clauses);
+  c1 = gfc_trans_omp_clauses (&block, code->ext.oacc_declare->clauses,
+			      code->loc);
+
+  c2 = gfc_trans_omp_clauses (&block, code->ext.oacc_declare->return_clauses,
+			      code->loc);
+
+  stmt = build2_loc (input_location, construct_code, void_type_node, c1, c2);
   gfc_add_expr_to_block (&block, stmt);
+
   return gfc_finish_block (&block);
 }
 
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index e130af9..b6d4a42 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -132,7 +132,8 @@  enum acc_region_kind
 {
   ARK_GENERAL,  /* Default used for data, etc. regions.  */
   ARK_PARALLEL, /* Parallel construct.  */
-  ARK_KERNELS   /* Kernels construct.  */
+  ARK_KERNELS,  /* Kernels construct.  */
+  ARK_DECLARE   /* Declare directive.  */
 };
 
 /* Gimplify hashtable helper.  */
@@ -176,6 +177,7 @@  struct gimplify_omp_ctx
   enum acc_region_kind acc_region_kind;
   bool combined_loop;
   bool distribute;
+  gomp_target *stmt;
 };
 
 static struct gimplify_ctx *gimplify_ctxp;
@@ -7105,6 +7107,61 @@  gimplify_oacc_cache (tree *expr_p, gimple_seq *pre_p)
   *expr_p = NULL_TREE;
 }
 
+/* Gimplify OACC_DECLARE.  */
+
+static void
+gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
+{
+  tree expr = *expr_p;
+  gomp_target *stmt;
+  tree clauses, t;
+
+  clauses = OACC_DECLARE_CLAUSES (expr);
+
+  gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, ORK_OACC);
+
+  gimplify_omp_ctxp->acc_region_kind = ARK_DECLARE;
+  gimplify_omp_ctxp->stmt = NULL;
+
+  for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
+    {
+      tree attrs, decl = OMP_CLAUSE_DECL (t);
+
+      if (TREE_CODE (decl) == MEM_REF)
+	continue;
+
+      omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN);
+
+      attrs = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (decl));
+      if (attrs)
+	DECL_ATTRIBUTES (decl) = remove_attribute ("oacc declare", attrs);
+    }
+
+  stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE,
+				  clauses);
+
+  gimplify_seq_add_stmt (pre_p, stmt);
+
+  clauses = OACC_DECLARE_RETURN_CLAUSES (expr);
+
+  if (clauses)
+    {
+      struct gimplify_omp_ctx *c;
+
+      gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, ORK_OACC);
+
+      c = gimplify_omp_ctxp;
+      gimplify_omp_ctxp = c->outer_context;
+      delete_omp_context (c);
+
+      stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE,
+				      clauses);
+      gimplify_omp_ctxp->stmt = stmt;
+    }
+
+  *expr_p = NULL_TREE;
+}
+
 static tree
 gimplify_oacc_host_data_1 (tree *tp, int *walk_subtrees, void *data ATTRIBUTE_UNUSED)
 {
@@ -7933,10 +7990,6 @@  gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
 
   switch (TREE_CODE (expr))
     {
-    case OACC_DECLARE:
-      kind = GF_OMP_TARGET_KIND_OACC_DECLARE;
-      ork = ORK_OACC;
-      break;
     case OACC_ENTER_DATA:
       kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
       ork = ORK_OACC;
@@ -8914,6 +8967,11 @@  gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	case OACC_HOST_DATA:
 	  ret = gimplify_oacc_host_data (expr_p, pre_p);
 	  break;
+
+	case OACC_DECLARE:
+	  gimplify_oacc_declare (expr_p, pre_p);
+	  ret = GS_ALL_DONE;
+	  break;
 	  
 	case OACC_KERNELS:
 	case OACC_PARALLEL:
@@ -8927,7 +8985,6 @@  gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	  ret = GS_ALL_DONE;
 	  break;
 
-	case OACC_DECLARE:
 	case OACC_ENTER_DATA:
 	case OACC_EXIT_DATA:
 	case OACC_UPDATE:
@@ -9568,6 +9625,25 @@  gimplify_body (tree fndecl, bool do_parms)
       gimplify_seq_add_stmt (&seq, outer_stmt);
     }
 
+   if (flag_openacc && gimplify_omp_ctxp)
+    {
+      while (gimplify_omp_ctxp)
+	{
+	  struct gimplify_omp_ctx *c;
+
+	  if (gimplify_omp_ctxp->acc_region_kind == ARK_DECLARE
+	      && gimplify_omp_ctxp->stmt)
+	    {
+	      gimplify_seq_add_stmt (&seq, gimplify_omp_ctxp->stmt);
+	      gimplify_omp_ctxp->stmt = NULL;
+	    }
+
+	  c = gimplify_omp_ctxp;
+	  gimplify_omp_ctxp = c->outer_context;
+	  delete_omp_context (c);
+	}
+    }
+
   /* The body must contain exactly one statement, a GIMPLE_BIND.  If this is
      not the case, wrap everything in a GIMPLE_BIND to make it so.  */
   if (gimple_code (outer_stmt) == GIMPLE_BIND
diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-1.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-1.f95
index 7ea58ef..3129f04 100644
--- a/gcc/testsuite/gfortran.dg/goacc/declare-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/declare-1.f95
@@ -16,4 +16,3 @@  contains
   end function foo
 end program test
 ! { dg-final { scan-tree-dump-times "pragma acc declare map\\(force_to:i\\)" 2 "original" } }
-! { dg-final { scan-tree-dump-times "pragma acc declare map\\(force_from:i\\)" 2 "original" } }
diff --git a/gcc/tree.def b/gcc/tree.def
index 56580af..9ea537e 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1056,6 +1056,11 @@  DEFTREECODE (OACC_KERNELS, "oacc_kernels", tcc_statement, 2)
 
 DEFTREECODE (OACC_DATA, "oacc_data", tcc_statement, 2)
 
+/* OpenACC - #pragma acc declare [clause1 ... clauseN]
+   Operand 0: OACC_DECLARE_CLAUSES: List of clauses.
+   Operand 1: OACC_DECLARE_RETURN_CLAUSES: List of clauses for returns.  */
+DEFTREECODE (OACC_DECLARE, "oacc_declare", tcc_statement, 2)
+
 /* OpenACC - #pragma acc host_data [clause1 ... clauseN]
    Operand 0: OACC_HOST_DATA_BODY: Host_data construct body.
    Operand 1: OACC_HOST_DATA_CLAUSES: List of clauses.  */
@@ -1166,10 +1171,6 @@  DEFTREECODE (OMP_CRITICAL, "omp_critical", tcc_statement, 2)
 	OMP_CLAUSE__CACHE_ clauses).  */
 DEFTREECODE (OACC_CACHE, "oacc_cache", tcc_statement, 1)
 
-/* OpenACC - #pragma acc declare [clause1 ... clauseN]
-   Operand 0: OACC_DECLARE_CLAUSES: List of clauses.  */
-DEFTREECODE (OACC_DECLARE, "oacc_declare", tcc_statement, 1)
-
 /* OpenACC - #pragma acc enter data [clause1 ... clauseN]
    Operand 0: OACC_ENTER_DATA_CLAUSES: List of clauses.  */
 DEFTREECODE (OACC_ENTER_DATA, "oacc_enter_data", tcc_statement, 1)
diff --git a/gcc/tree.h b/gcc/tree.h
index 29bce01..5b2e267 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1233,6 +1233,8 @@  extern void protected_set_expr_location (tree, location_t);
 
 #define OACC_DECLARE_CLAUSES(NODE) \
   TREE_OPERAND (OACC_DECLARE_CHECK (NODE), 0)
+#define OACC_DECLARE_RETURN_CLAUSES(NODE) \
+  TREE_OPERAND (OACC_DECLARE_CHECK (NODE), 1)
 
 #define OACC_ENTER_DATA_CLAUSES(NODE) \
   TREE_OPERAND (OACC_ENTER_DATA_CHECK (NODE), 0)
diff --git a/libgomp/testsuite/libgomp.oacc-c++/declare-1.C b/libgomp/testsuite/libgomp.oacc-c++/declare-1.C
index 6618b10..e82a8e5 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/declare-1.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/declare-1.C
@@ -1,3 +1,4 @@ 
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
 
 template<class T>
 T foo()
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c
index 0d75f49..3dfde71 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c
@@ -6,6 +6,8 @@  int
 main (int argc, char **argv)
 {
   float a, b;
+  float c;
+#pragma acc declare create (c)
 
   a = 2.0;
   b = 0.0;
@@ -60,5 +62,28 @@  main (int argc, char **argv)
   if (a != 5.0)
     abort ();
 
+#pragma acc parallel default (none) copy (a)
+  {
+    c = a;
+    a = 1.0;
+    a = a + c;
+  }
+
+  if (a != 6.0)
+    abort ();
+
+#pragma acc data copy (a)
+  {
+#pragma acc parallel default (none)
+    {
+      c = a;
+      a = 1.0;
+      a = a + c;
+    }
+  }
+
+  if (a != 7.0)
+    abort ();
+
   return 0;
 }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
index 4d58e70..18dd1bb 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
@@ -104,13 +104,11 @@  subroutine subr2 (a, b, c)
 
 end subroutine
 
-subroutine subr1 (a, b, c)
+subroutine subr1 (a)
   integer, parameter :: N = 8
   integer :: i
   integer :: a(N)
   !$acc declare present (a)
-  integer :: b(N)
-  integer :: c(N)
 
   i = 0
 
@@ -144,7 +142,7 @@  subroutine subr0 (a, b, c, d)
   call test (b, .false.)
   call test (c, .false.)
 
-  call subr1 (a, b, c)
+  call subr1 (a)
 
   call test (a, .true.)
   call test (b, .false.)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
index f82316e..1059089 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
@@ -3,6 +3,8 @@ 
 program main
   implicit none
   real a, b
+  real c
+  !$acc declare create (c)
 
   a = 2.0
   b = 0.0
@@ -31,4 +33,22 @@  program main
 
   if (a .ne. 5.0) call abort
 
+  !$acc parallel default (none) copy (a)
+    c = a
+    a = 1.0
+    a = a + c
+  !$acc end parallel
+
+  if (a .ne. 6.0) call abort
+
+  !$acc data copy (a)
+  !$acc parallel default (none)
+    c = a
+    a = 1.0
+    a = a + c
+  !$acc end parallel
+  !$acc end data
+
+  if (a .ne. 7.0) call abort
+
 end program main