Hi!

On Mon, 23 Nov 2015 13:37:20 +0100, I wrote:
> A few things I noticed when working on merging your trunk r230275 into
> gomp-4_0-branch.  [...]

I have now committed the merge of your trunk r230275 into
gomp-4_0-branch, in r230787.  This was non-trivial: the implementation of
OpenACC declare on gomp-4_0-branch has been rather different from what
got committed to trunk, so aside from resolving all the textual merge
conflicts, I also had to fix that up, clean up things which were only
relevant/needed for earlier revisions of OpenACC declare support on
gomp-4_0-branch: "oacc declare" function attribute removed,
BUILT_IN_GOACC_STATIC/GOACC_register_static builtin and
BT_FN_VOID_PTR_INT_UINT function type removed, goacc_allocate_static and
goacc_deallocate_static functions removed in libgomp, OACC_DECLARE tree
code restored to trunk variant, OACC_DECLARE_RETURN_CLAUSES removed, stmt
member of struct gimplify_omp_ctx removed,
gcc/varpool.c:varpool_node::get_create changes reverted.  I also had to
XFAIL a few Fortran test cases, because on gomp-4_0-branch the Fortran
front end has not yet been updated to match the changes in the generic
middle end and libgomp parts of the OpenACC declare implementation, that
are now done differently on trunk.  I expect to remove these XFAILs once
(soon) merging into gomp-4_0-branch the respective trunk commit
containing these Fortran front end changes.  So, I now effectively
applied the following patch to gomp-4_0-branch; please verify that's
alright:

commit d60891e4995b13368b305db2db1d3babebb018c9
Merge: 4002b8b 2fc5e98
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Tue Nov 24 08:10:01 2015 +0000

    svn merge -r 230274:230275 svn+ssh://gcc.gnu.org/svn/gcc/trunk
    
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@230787 
138bc75d-0d04-0410-961f-82ee72b054a4

 gcc/ChangeLog                                      |  23 +++
 gcc/builtin-types.def                              |   1 -
 gcc/c-family/ChangeLog                             |   8 +
 gcc/c-family/c-common.c                            |   1 -
 gcc/c-family/c-pragma.h                            |   4 +-
 gcc/c/ChangeLog                                    |  12 ++
 gcc/c/c-parser.c                                   | 227 +++++----------------
 gcc/cp/ChangeLog                                   |  14 ++
 gcc/cp/parser.c                                    | 224 +++++---------------
 gcc/cp/pt.c                                        |   3 -
 gcc/fortran/f95-lang.c                             |   2 -
 gcc/fortran/gfortran.h                             |   2 +-
 gcc/fortran/trans-decl.c                           |  19 +-
 gcc/fortran/trans-openmp.c                         |   6 +-
 gcc/fortran/types.def                              |   1 -
 gcc/gimplify.c                                     | 205 +++++++++++++------
 gcc/omp-builtins.def                               |   2 -
 gcc/testsuite/ChangeLog                            |   6 +
 gcc/testsuite/c-c++-common/goacc/declare-2.c       |  15 +-
 gcc/tree-pretty-print.c                            |  12 +-
 gcc/tree.def                                       |   9 +-
 gcc/tree.h                                         |   2 -
 gcc/varpool.c                                      |  52 +----
 include/ChangeLog                                  |   6 +
 include/gomp-constants.h                           |   3 +-
 libgomp/ChangeLog                                  |  11 +
 libgomp/libgomp.map                                |   3 +-
 libgomp/oacc-init.c                                |   9 +-
 libgomp/oacc-int.h                                 |   3 -
 libgomp/oacc-parallel.c                            |  88 +-------
 libgomp/testsuite/libgomp.oacc-c++/declare-1.C     |  14 +-
 .../libgomp.oacc-c-c++-common/declare-1.c          |   6 +-
 .../libgomp.oacc-c-c++-common/declare-5.c          |   4 +-
 .../testsuite/libgomp.oacc-fortran/declare-1.f90   |   1 +
 .../testsuite/libgomp.oacc-fortran/declare-3.f90   |   1 +
 35 files changed, 405 insertions(+), 594 deletions(-)

[diff --git gcc/ChangeLog gcc/ChangeLog]
diff --git gcc/builtin-types.def gcc/builtin-types.def
index 7744f01..c68fb19 100644
--- gcc/builtin-types.def
+++ gcc/builtin-types.def
@@ -450,7 +450,6 @@ DEF_FUNCTION_TYPE_3 (BT_FN_BOOL_ULONG_ULONG_ULONGPTR, 
BT_BOOL, BT_ULONG,
                     BT_ULONG, BT_PTR_ULONG)
 DEF_FUNCTION_TYPE_3 (BT_FN_BOOL_ULONGLONG_ULONGLONG_ULONGLONGPTR, BT_BOOL,
                     BT_ULONGLONG, BT_ULONGLONG, BT_PTR_ULONGLONG)
-DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_INT_UINT, BT_VOID, BT_PTR, BT_INT, BT_UINT)
 
 DEF_FUNCTION_TYPE_4 (BT_FN_SIZE_CONST_PTR_SIZE_SIZE_FILEPTR,
                     BT_SIZE, BT_CONST_PTR, BT_SIZE, BT_SIZE, BT_FILEPTR)
[diff --git gcc/c-family/ChangeLog gcc/c-family/ChangeLog]
diff --git gcc/c-family/c-common.c gcc/c-family/c-common.c
index cbe2f01..eff919c 100644
--- gcc/c-family/c-common.c
+++ gcc/c-family/c-common.c
@@ -831,7 +831,6 @@ const struct attribute_spec c_common_attribute_table[] =
                              handle_bnd_legacy, false },
   { "bnd_instrument",         0, 0, true, false, false,
                              handle_bnd_instrument, false },
-  { "oacc declare",           0, -1, true,  false, false, NULL, false },
   { NULL,                     0, 0, false, false, false, NULL, false }
 };
 
diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h
index 7049b79..aaebc50 100644
--- gcc/c-family/c-pragma.h
+++ gcc/c-family/c-pragma.h
@@ -181,9 +181,9 @@ enum pragma_omp_clause {
   PRAGMA_OACC_CLAUSE_DEFAULT = PRAGMA_OMP_CLAUSE_DEFAULT,
   PRAGMA_OACC_CLAUSE_FIRSTPRIVATE = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE,
   PRAGMA_OACC_CLAUSE_IF = PRAGMA_OMP_CLAUSE_IF,
-  PRAGMA_OACC_CLAUSE_LINK = PRAGMA_OMP_CLAUSE_LINK,
   PRAGMA_OACC_CLAUSE_PRIVATE = PRAGMA_OMP_CLAUSE_PRIVATE,
-  PRAGMA_OACC_CLAUSE_REDUCTION = PRAGMA_OMP_CLAUSE_REDUCTION
+  PRAGMA_OACC_CLAUSE_REDUCTION = PRAGMA_OMP_CLAUSE_REDUCTION,
+  PRAGMA_OACC_CLAUSE_LINK = PRAGMA_OMP_CLAUSE_LINK
 };
 
 extern struct cpp_reader* parse_in;
[diff --git gcc/c/ChangeLog gcc/c/ChangeLog]
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 3f66e6a..fc88e3e 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -1228,6 +1228,7 @@ static vec<tree, va_gc> *c_parser_expr_list (c_parser *, 
bool, bool,
                                             vec<tree, va_gc> **, location_t *,
                                             tree *, vec<location_t> *,
                                             unsigned int * = NULL);
+static void c_parser_oacc_declare (c_parser *);
 static void c_parser_oacc_enter_exit_data (c_parser *, bool);
 static void c_parser_oacc_update (c_parser *);
 static void c_parser_omp_construct (c_parser *);
@@ -1248,7 +1249,6 @@ static bool c_parser_omp_target (c_parser *, enum 
pragma_context);
 static void c_parser_omp_end_declare_target (c_parser *);
 static void c_parser_omp_declare (c_parser *, enum pragma_context);
 static bool c_parser_omp_ordered (c_parser *, enum pragma_context);
-static void c_parser_oacc_declare (c_parser *parser);
 static void c_parser_oacc_routine (c_parser *parser, enum pragma_context);
 
 /* These Objective-C parser functions are only ever called when
@@ -10054,13 +10054,13 @@ c_parser_omp_clause_name (c_parser *parser, bool 
consume_token = true)
            result = PRAGMA_OMP_CLAUSE_DEPEND;
          else if (!strcmp ("device", p))
            result = PRAGMA_OMP_CLAUSE_DEVICE;
+         else if (!strcmp ("deviceptr", p))
+           result = PRAGMA_OACC_CLAUSE_DEVICEPTR;
          else if (!strcmp ("device_resident", p))
            result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT;
          else if (!strcmp ("device_type", p)
                   || !strcmp ("dtype", p))
            result = PRAGMA_OACC_CLAUSE_DEVICE_TYPE;
-         else if (!strcmp ("deviceptr", p))
-           result = PRAGMA_OACC_CLAUSE_DEVICEPTR;
          else if (!strcmp ("dist_schedule", p))
            result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
          break;
@@ -12914,6 +12914,10 @@ c_parser_oacc_all_clauses (c_parser *parser, 
omp_clause_mask mask,
          clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
          c_name = "device";
          break;
+       case PRAGMA_OACC_CLAUSE_DEVICEPTR:
+         clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses);
+         c_name = "deviceptr";
+         break;
        case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT:
          clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
          c_name = "device_resident";
@@ -12924,10 +12928,6 @@ c_parser_oacc_all_clauses (c_parser *parser, 
omp_clause_mask mask,
          c_name = "device_type";
          seen_dtype = true;
          break;
-       case PRAGMA_OACC_CLAUSE_DEVICEPTR:
-         clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses);
-         c_name = "deviceptr";
-         break;
        case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
          clauses = c_parser_omp_clause_firstprivate (parser, clauses);
          c_name = "firstprivate";
@@ -13419,8 +13419,6 @@ c_parser_oacc_data (location_t loc, c_parser *parser)
    # pragma acc declare oacc-data-clause[optseq] new-line
 */
 
-static int oacc_dcl_idx = 0;
-
 #define OACC_DECLARE_CLAUSE_MASK                                       \
        ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY)                \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)              \
@@ -13439,7 +13437,7 @@ static void
 c_parser_oacc_declare (c_parser *parser)
 {
   location_t pragma_loc = c_parser_peek_token (parser)->location;
-  tree c, clauses, ret_clauses, stmt, t;
+  tree clauses, stmt, t, decl;
 
   bool error = false;
 
@@ -13457,11 +13455,10 @@ c_parser_oacc_declare (c_parser *parser)
   for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
     {
       location_t loc = OMP_CLAUSE_LOCATION (t);
-      tree decl = OMP_CLAUSE_DECL (t);
-      tree devres = NULL_TREE;
+      decl = OMP_CLAUSE_DECL (t);
       if (!DECL_P (decl))
        {
-         error_at (loc, "subarray in %<#pragma acc declare%>");
+         error_at (loc, "array section in %<#pragma acc declare%>");
          error = true;
          continue;
        }
@@ -13471,10 +13468,7 @@ c_parser_oacc_declare (c_parser *parser)
        case GOMP_MAP_FORCE_ALLOC:
        case GOMP_MAP_FORCE_TO:
        case GOMP_MAP_FORCE_DEVICEPTR:
-         break;
-
        case GOMP_MAP_DEVICE_RESIDENT:
-         devres = t;
          break;
 
        case GOMP_MAP_POINTER:
@@ -13483,7 +13477,9 @@ c_parser_oacc_declare (c_parser *parser)
          break;
 
        case GOMP_MAP_LINK:
-         if (!global_bindings_p () && !DECL_EXTERNAL (decl))
+         if (!global_bindings_p ()
+             && (TREE_STATIC (decl)
+              || !DECL_EXTERNAL (decl)))
            {
              error_at (loc,
                        "%qD must be a global variable in"
@@ -13520,173 +13516,58 @@ c_parser_oacc_declare (c_parser *parser)
          break;
        }
 
-      tree decl_for_attr = decl;
-      tree prev_attr = lookup_attribute ("oacc declare",
-                                        DECL_ATTRIBUTES (decl));
-      if (prev_attr)
+      if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))
+         || lookup_attribute ("omp declare target link",
+                              DECL_ATTRIBUTES (decl)))
        {
-         tree p = TREE_VALUE (prev_attr);
-         tree cl = TREE_VALUE (p);
+         error_at (loc, "variable %qD used more than once with "
+                   "%<#pragma acc declare%>", decl);
+         error = true;
+         continue;
+       }
 
-         if (!devres && OMP_CLAUSE_MAP_KIND (cl) != GOMP_MAP_DEVICE_RESIDENT)
+      if (!error)
+       {
+         tree id;
+
+         if (OMP_CLAUSE_MAP_KIND (t) == GOMP_MAP_LINK)
+           id = get_identifier ("omp declare target link");
+         else
+           id = get_identifier ("omp declare target");
+
+         DECL_ATTRIBUTES (decl)
+                          = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl));
+
+         if (global_bindings_p ())
            {
-             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;
+             symtab_node *node = symtab_node::get (decl);
+             if (node != NULL)
+               {
+                 node->offloadable = 1;
+#ifdef ENABLE_OFFLOADING
+                 g->have_offload = true;
+                 if (is_a <varpool_node *> (node))
+                   {
+                     vec_safe_push (offload_vars, decl);
+                     node->force_output = 1;
+                   }
+#endif
+               }
            }
        }
-
-      if (!error && global_bindings_p ())
-       {
-         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);
-       }
     }
 
-  if (error)
+  if (error || global_bindings_p ())
     return;
 
-  ret_clauses = NULL_TREE;
+  stmt = make_node (OACC_DECLARE);
+  TREE_TYPE (stmt) = void_type_node;
+  OACC_DECLARE_CLAUSES (stmt) = clauses;
+  SET_EXPR_LOCATION (stmt, pragma_loc);
 
-  for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
-    {
-      bool ret = false;
-      HOST_WIDE_INT kind, new_op;
+  add_stmt (stmt);
 
-      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];
-      struct c_declarator *target;
-      tree stmt, attrs;
-      c_arg_info *arg_info = build_arg_info ();
-      struct c_declarator *declarator;
-      struct c_declspecs *specs;
-      struct c_typespec spec;
-      location_t loc = UNKNOWN_LOCATION;
-      tree f, t, fnbody, call_fn;
-
-      sprintf (buf, "__openacc_c_constructor__%d", oacc_dcl_idx++);
-      target = build_id_declarator (get_identifier (buf));
-      arg_info->types = void_list_node;
-      declarator = build_function_declarator (arg_info, target);
-
-      specs = build_null_declspecs ();
-      spec.kind = ctsk_resword;
-      spec.spec = get_identifier ("void");
-      spec.expr = NULL_TREE;
-      spec.expr_const_operands = true;
-
-      declspecs_add_type (pragma_loc, specs, spec);
-      finish_declspecs (specs);
-
-      attrs = tree_cons (get_identifier ("constructor") , NULL_TREE, 
NULL_TREE);
-      start_function (specs, declarator, attrs);
-      store_parm_decls ();
-      f = c_begin_compound_stmt (true);
-      TREE_USED (current_function_decl) = 1;
-      call_fn = builtin_decl_explicit (BUILT_IN_GOACC_STATIC);
-      TREE_SIDE_EFFECTS (call_fn) = 1;
-
-      for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
-       {
-         tree d, a1, a2, a3;
-         vec<tree, va_gc> *args;
-         vec_alloc (args, 3);
-
-         d = OMP_CLAUSE_DECL (t);
-
-         a1 = build_unary_op (loc, ADDR_EXPR, d, 0);
-         a2 = DECL_SIZE_UNIT (d);
-         a3 = build_int_cst (unsigned_type_node, OMP_CLAUSE_MAP_KIND (t));
-
-         args->quick_push (a1);
-         args->quick_push (a2);
-         args->quick_push (a3);
-
-         stmt = build_function_call_vec (loc, vNULL, call_fn, args, NULL);
-         add_stmt (stmt);
-       }
-
-       fnbody = c_end_compound_stmt (loc, f, true);
-       add_stmt (fnbody);
-
-      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);
-    }
+  return;
 }
 
 /* OpenACC 2.0:
[diff --git gcc/cp/ChangeLog gcc/cp/ChangeLog]
diff --git gcc/cp/parser.c gcc/cp/parser.c
index 1e15e79..3158905 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -29133,13 +29133,13 @@ cp_parser_omp_clause_name (cp_parser *parser, bool 
consume_token = true)
            result = PRAGMA_OMP_CLAUSE_DEPEND;
          else if (!strcmp ("device", p))
            result = PRAGMA_OMP_CLAUSE_DEVICE;
+         else if (!strcmp ("deviceptr", p))
+           result = PRAGMA_OACC_CLAUSE_DEVICEPTR;
          else if (!strcmp ("device_resident", p))
            result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT;
          else if (!strcmp ("device_type", p)
                   || !strcmp ("dtype", p))
            result = PRAGMA_OACC_CLAUSE_DEVICE_TYPE;
-         else if (!strcmp ("deviceptr", p))
-           result = PRAGMA_OACC_CLAUSE_DEVICEPTR;
          else if (!strcmp ("dist_schedule", p))
            result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
          break;
@@ -31696,6 +31696,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, 
omp_clause_mask mask,
          clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
          c_name = "device";
          break;
+       case PRAGMA_OACC_CLAUSE_DEVICEPTR:
+         clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses);
+         c_name = "deviceptr";
+         break;
        case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT:
          clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
          c_name = "device_resident";
@@ -31706,10 +31710,6 @@ cp_parser_oacc_all_clauses (cp_parser *parser, 
omp_clause_mask mask,
          c_name = "device_type";
          seen_dtype = true;
          break;
-       case PRAGMA_OACC_CLAUSE_DEVICEPTR:
-         clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses);
-         c_name = "deviceptr";
-         break;
        case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
          clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FIRSTPRIVATE,
                                            clauses);
@@ -34707,8 +34707,6 @@ cp_parser_oacc_data (cp_parser *parser, cp_token 
*pragma_tok)
    # pragma acc declare oacc-data-clause[optseq] new-line
 */
 
-static int oacc_dcl_idx = 0;
-
 #define OACC_DECLARE_CLAUSE_MASK                                       \
        ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY)                \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)              \
@@ -34726,14 +34724,12 @@ static int oacc_dcl_idx = 0;
 static tree
 cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 {
-  tree c, clauses, ret_clauses, stmt, t;
+  tree clauses, stmt;
   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,
@@ -34745,10 +34741,9 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token 
*pragma_tok)
     {
       location_t loc = OMP_CLAUSE_LOCATION (t);
       tree decl = OMP_CLAUSE_DECL (t);
-      tree devres = NULL_TREE;
       if (!DECL_P (decl))
        {
-         error_at (loc, "subarray in %<#pragma acc declare%>");
+         error_at (loc, "array section in %<#pragma acc declare%>");
          error = true;
          continue;
        }
@@ -34758,10 +34753,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token 
*pragma_tok)
        case GOMP_MAP_FORCE_ALLOC:
        case GOMP_MAP_FORCE_TO:
        case GOMP_MAP_FORCE_DEVICEPTR:
-         break;
-
        case GOMP_MAP_DEVICE_RESIDENT:
-         devres = t;
          break;
 
        case GOMP_MAP_POINTER:
@@ -34770,7 +34762,9 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token 
*pragma_tok)
          break;
 
        case GOMP_MAP_LINK:
-         if (!global_bindings_p () && !DECL_EXTERNAL (decl))
+         if (!global_bindings_p ()
+             && (TREE_STATIC (decl)
+              || !DECL_EXTERNAL (decl)))
            {
              error_at (loc,
                        "%qD must be a global variable in"
@@ -34807,173 +34801,55 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token 
*pragma_tok)
          break;
        }
 
-      tree decl_for_attr = decl;
-      tree prev_attr = lookup_attribute ("oacc declare",
-                                            DECL_ATTRIBUTES (decl));
-      if (prev_attr)
+      if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))
+         || lookup_attribute ("omp declare target link",
+                              DECL_ATTRIBUTES (decl)))
        {
-         tree p = TREE_VALUE (prev_attr);
-         tree cl = TREE_VALUE (p);
+         error_at (loc, "variable %qD used more than once with "
+                   "%<#pragma acc declare%>", decl);
+         error = true;
+         continue;
+       }
 
-         if (!devres && OMP_CLAUSE_MAP_KIND (cl) != GOMP_MAP_DEVICE_RESIDENT)
+      if (!error)
+       {
+         tree id;
+
+         if (OMP_CLAUSE_MAP_KIND (t) == GOMP_MAP_LINK)
+           id = get_identifier ("omp declare target link");
+         else
+           id = get_identifier ("omp declare target");
+
+         DECL_ATTRIBUTES (decl)
+                          = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl));
+         if (global_bindings_p ())
            {
-             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;
+             symtab_node *node = symtab_node::get (decl);
+             if (node != NULL)
+               {
+                 node->offloadable = 1;
+#ifdef ENABLE_OFFLOADING
+                 g->have_offload = true;
+                 if (is_a <varpool_node *> (node))
+                   {
+                     vec_safe_push (offload_vars, decl);
+                     node->force_output = 1;
+                   }
+#endif
+               }
            }
        }
-
-      if (!error && global_bindings_p ())
-       {
-         tree attr = tree_cons (NULL_TREE, t, NULL_TREE);
-         tree attrs = tree_cons (get_identifier ("oacc declare"),
-                                 attr, NULL_TREE);
-         decl_attributes (&decl_for_attr, attrs, 0);
-       }
     }
 
-  if (error)
+  if (error || global_bindings_p ())
     return NULL_TREE;
 
-  ret_clauses = NULL_TREE;
+  stmt = make_node (OACC_DECLARE);
+  TREE_TYPE (stmt) = void_type_node;
+  OACC_DECLARE_CLAUSES (stmt) = clauses;
+  SET_EXPR_LOCATION (stmt, pragma_tok->location);
 
-  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];
-      cp_decl_specifier_seq decl_specifiers;
-      cp_declarator *declarator;
-      tree attrs, parms;
-      tree f, t, call_fn, stmt;
-      location_t loc = UNKNOWN_LOCATION;
-      void *p;
-
-      p = obstack_alloc (&declarator_obstack, 0);
-      clear_decl_specs (&decl_specifiers);
-      decl_specifiers.type = void_type_node;
-      sprintf (buf, "__openacc_cp_constructor__%d", oacc_dcl_idx++);
-
-      declarator = make_id_declarator (NULL_TREE, get_identifier (buf),
-                                      sfk_none);
-      parms = void_list_node;
-      declarator = make_call_declarator (declarator, parms,
-                                        TYPE_UNQUALIFIED,
-                                        VIRT_SPEC_UNSPECIFIED,
-                                        REF_QUAL_NONE,
-                                        /*tx_qualifier=*/NULL_TREE,
-                                        /*exception_specification=*/NULL_TREE,
-                                        /*late_return_type=*/NULL_TREE,
-                                        /*requires_clause=*/NULL_TREE);
-      attrs = tree_cons (get_identifier ("constructor") , NULL_TREE, 
NULL_TREE);
-      start_function (&decl_specifiers, declarator, attrs);
-      f = begin_compound_stmt (0);
-      call_fn = builtin_decl_explicit (BUILT_IN_GOACC_STATIC);
-      TREE_SIDE_EFFECTS (call_fn) = 1;
-
-      for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
-       {
-         tree d, a1, a2, a3;
-         vec<tree, va_gc> *args;
-         vec_alloc (args, 3);
-
-         d = OMP_CLAUSE_DECL (t);
-
-         a1 = build_unary_op (loc, ADDR_EXPR, d, 0);
-         a2 = DECL_SIZE_UNIT (d);
-         a3 = build_int_cst (unsigned_type_node, OMP_CLAUSE_MAP_KIND (t));
-
-         args->quick_push (a1);
-         args->quick_push (a2);
-         args->quick_push (a3);
-
-         stmt = build_function_call_vec (loc, vNULL, call_fn, args, NULL);
-         finish_expr_stmt (stmt);
-       }
-
-      finish_compound_stmt (f);
-      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);
-    }
+  add_stmt (stmt);
 
   return NULL_TREE;
 }
diff --git gcc/cp/pt.c gcc/cp/pt.c
index cfaf9b2..f70a444 100644
--- gcc/cp/pt.c
+++ gcc/cp/pt.c
@@ -15416,9 +15416,6 @@ tsubst_expr (tree t, tree args, tsubst_flags_t 
complain, tree in_decl,
       tmp = tsubst_omp_clauses (OACC_DECLARE_CLAUSES (t), false, false,
                                args, complain, in_decl);
       OACC_DECLARE_CLAUSES (t) = tmp;
-      tmp = tsubst_omp_clauses (OACC_DECLARE_RETURN_CLAUSES (t), false, false,
-                               args, complain, in_decl);
-      OACC_DECLARE_RETURN_CLAUSES (t) = tmp;
       add_stmt (t);
       break;
 
diff --git gcc/fortran/f95-lang.c gcc/fortran/f95-lang.c
index 7038fa3..0958102 100644
--- gcc/fortran/f95-lang.c
+++ gcc/fortran/f95-lang.c
@@ -94,8 +94,6 @@ static const struct attribute_spec gfc_attribute_table[] =
        affects_type_identity } */
   { "omp declare target", 0, 0, true,  false, false,
     gfc_handle_omp_declare_target_attribute, false },
-  { "oacc declare", 0, 0, true,  false, false,
-    gfc_handle_omp_declare_target_attribute, false },
   { "oacc function", 0, -1, true,  false, false,
     gfc_handle_omp_declare_target_attribute, false },
   { NULL,                0, 0, false, false, false, NULL, false }
diff --git gcc/fortran/gfortran.h gcc/fortran/gfortran.h
index 5e7db13..c8401cf 100644
--- gcc/fortran/gfortran.h
+++ gcc/fortran/gfortran.h
@@ -1250,7 +1250,7 @@ gfc_omp_clauses;
 
 #define gfc_get_omp_clauses() XCNEW (gfc_omp_clauses)
 
-/* Node in the linked list used for storing !$oacc declare constructs.  */
+/* Node in the linked list used for storing OpenACC declare constructs.  */
 
 typedef struct gfc_oacc_declare
 {
diff --git gcc/fortran/trans-decl.c gcc/fortran/trans-decl.c
index 976a518..058fd17 100644
--- gcc/fortran/trans-decl.c
+++ gcc/fortran/trans-decl.c
@@ -1301,19 +1301,16 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree 
list)
        list = chainon (list, attr);
       }
 
-  if (sym_attr.omp_declare_target)
-    list = tree_cons (get_identifier ("omp declare target"),
-                     NULL_TREE, list);
-
-  if (sym_attr.oacc_declare_create
+  if (sym_attr.omp_declare_target
+      || sym_attr.oacc_declare_create
       || sym_attr.oacc_declare_copyin
       || sym_attr.oacc_declare_deviceptr
-      || sym_attr.oacc_declare_device_resident
-      || sym_attr.oacc_declare_link)
-    {
-      list = tree_cons (get_identifier ("oacc declare"),
-                       NULL_TREE, list);
-    }
+      || sym_attr.oacc_declare_device_resident)
+    list = tree_cons (get_identifier ("omp declare target"),
+                     NULL_TREE, list);
+  if (sym_attr.oacc_declare_link)
+    list = tree_cons (get_identifier ("omp declare target link"),
+                     NULL_TREE, list);
 
   if (sym_attr.oacc_function)
     {
diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c
index c95c2e6..87ecc5a 100644
--- gcc/fortran/trans-openmp.c
+++ gcc/fortran/trans-openmp.c
@@ -4672,7 +4672,7 @@ tree
 gfc_trans_oacc_declare (gfc_code *code)
 {
   stmtblock_t block;
-  tree stmt, c1, c2;
+  tree stmt, c1;
   enum tree_code construct_code;
 
   gfc_start_block (&block);
@@ -4683,10 +4683,12 @@ gfc_trans_oacc_declare (gfc_code *code)
   c1 = gfc_trans_omp_clauses (&block, code->ext.oacc_declare->clauses,
                              code->loc);
 
+#if 0 /* TODO */
   c2 = gfc_trans_omp_clauses (&block, code->ext.oacc_declare->return_clauses,
                              code->loc);
+#endif
 
-  stmt = build2_loc (input_location, construct_code, void_type_node, c1, c2);
+  stmt = build1_loc (input_location, construct_code, void_type_node, c1);
   gfc_add_expr_to_block (&block, stmt);
 
   return gfc_finish_block (&block);
diff --git gcc/fortran/types.def gcc/fortran/types.def
index 7def2a8..a37e856 100644
--- gcc/fortran/types.def
+++ gcc/fortran/types.def
@@ -145,7 +145,6 @@ DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I2_INT, BT_VOID, 
BT_VOLATILE_PTR, BT_I2, BT
 DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I4_INT, BT_VOID, BT_VOLATILE_PTR, BT_I4, 
BT_INT)
 DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I8_INT, BT_VOID, BT_VOLATILE_PTR, BT_I8, 
BT_INT)
 DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I16_INT, BT_VOID, BT_VOLATILE_PTR, 
BT_I16, BT_INT)
-DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_INT_UINT, BT_VOID, BT_PTR, BT_INT, BT_UINT)
 
 DEF_FUNCTION_TYPE_4 (BT_FN_VOID_OMPFN_PTR_UINT_UINT,
                      BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, BT_UINT)
diff --git gcc/gimplify.c gcc/gimplify.c
index b91d69e..aea8bff 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -177,7 +177,6 @@ struct gimplify_omp_ctx
   bool target_map_scalars_firstprivate;
   bool target_map_pointers_as_0len_arrays;
   bool target_firstprivatize_array_bases;
-  gomp_target *stmt;
 };
 
 struct privatize_reduction
@@ -190,6 +189,7 @@ static struct gimplify_omp_ctx *gimplify_omp_ctxp;
 
 /* Forward declaration.  */
 static enum gimplify_status gimplify_compound_expr (tree *, gimple_seq *, 
bool);
+static hash_map<tree, tree> *oacc_declare_returns;
 
 /* Shorter alias name for the above function for use in gimplify.c
    only.  */
@@ -1092,6 +1092,7 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p)
   gimple_seq body, cleanup;
   gcall *stack_save;
   location_t start_locus = 0, end_locus = 0;
+  tree ret_clauses = NULL;
 
   tree temp = voidify_wrapper_expr (bind_expr, NULL);
 
@@ -1193,9 +1194,39 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p)
          clobber_stmt = gimple_build_assign (t, clobber);
          gimple_set_location (clobber_stmt, end_locus);
          gimplify_seq_add_stmt (&cleanup, clobber_stmt);
+
+         if (flag_openacc && oacc_declare_returns != NULL)
+           {
+             tree *c = oacc_declare_returns->get (t);
+             if (c != NULL)
+               {
+                 if (ret_clauses)
+                   OMP_CLAUSE_CHAIN (*c) = ret_clauses;
+
+                 ret_clauses = *c;
+
+                 oacc_declare_returns->remove (t);
+
+                 if (oacc_declare_returns->elements () == 0)
+                   {
+                     delete oacc_declare_returns;
+                     oacc_declare_returns = NULL;
+                   }
+               }
+           }
        }
     }
 
+  if (ret_clauses)
+    {
+      gomp_target *stmt;
+      gimple_stmt_iterator si = gsi_start (cleanup);
+
+      stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE,
+                                     ret_clauses);
+      gsi_insert_seq_before_without_update (&si, stmt, GSI_NEW_STMT);
+    }
+
   if (cleanup)
     {
       gtry *gs;
@@ -5823,6 +5854,26 @@ omp_notice_threadprivate_variable (struct 
gimplify_omp_ctx *ctx, tree decl,
   return false;
 }
 
+/* Return true if global var DECL is device resident.  */
+
+static bool
+device_resident_p (tree decl)
+{
+  tree attr = lookup_attribute ("oacc declare target", DECL_ATTRIBUTES (decl));
+
+  if (!attr)
+    return false;
+
+  for (tree t = TREE_VALUE (attr); t; t = TREE_PURPOSE (t))
+    {
+      tree c = TREE_VALUE (t);
+      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DEVICE_RESIDENT)
+       return true;
+    }
+
+  return false;
+}
+
 /* Determine outer default flags for DECL mentioned in an OMP region
    but not declared in an enclosing clause.
 
@@ -5914,25 +5965,6 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree 
decl,
   return flags;
 }
 
-/* Return true if global var DECL is device resident.  */
-
-static bool
-device_resident_p (tree decl)
-{
-  tree attr = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (decl));
-
-  if (!attr)
-    return false;
-  
-  for (tree t = TREE_VALUE (attr); t; t = TREE_PURPOSE (t))
-    {
-      tree c = TREE_VALUE (t);
-      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DEVICE_RESIDENT)
-       return true;
-    }
-
-  return false;
-}
 
 /* Determine outer default flags for DECL mentioned in an OACC region
    but not declared in an enclosing clause.  */
@@ -5941,10 +5973,15 @@ static unsigned
 oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
 {
   const char *rkind;
-  bool on_device = is_global_var (decl) && device_resident_p (decl);
+  bool on_device = false;
 
-  if (on_device)
-    flags |= GOVD_MAP_TO_ONLY;
+  if ((ctx->region_type & (ORT_ACC_PARALLEL | ORT_ACC_KERNELS)) != 0
+      && is_global_var (decl)
+      && device_resident_p (decl))
+    {
+      on_device = true;
+      flags |= GOVD_MAP_TO_ONLY;
+    }
 
   switch (ctx->region_type)
     {
@@ -7925,6 +7962,76 @@ gimplify_oacc_cache (tree *expr_p, gimple_seq *pre_p)
   *expr_p = NULL_TREE;
 }
 
+/* Helper function of gimplify_oacc_declare.  The helper's purpose is to,
+   if required, translate 'kind' in CLAUSE into an 'entry' kind and 'exit'
+   kind.  The entry kind will replace the one in CLAUSE, while the exit
+   kind will be used in a new omp_clause and returned to the caller.  */
+
+static tree
+gimplify_oacc_declare_1 (tree clause)
+{
+  HOST_WIDE_INT kind, new_op;
+  bool ret = false;
+  tree c = NULL;
+
+  kind = OMP_CLAUSE_MAP_KIND (clause);
+
+  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 (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;
+       ret = true;
+       break;
+
+      case GOMP_MAP_FROM:
+       OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC);
+       new_op = GOMP_MAP_FROM;
+       ret = true;
+       break;
+
+      case GOMP_MAP_TOFROM:
+       OMP_CLAUSE_SET_MAP_KIND (clause, 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)
+    {
+      c = build_omp_clause (OMP_CLAUSE_LOCATION (clause), OMP_CLAUSE_MAP);
+      OMP_CLAUSE_SET_MAP_KIND (c, new_op);
+      OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clause);
+    }
+
+  return c;
+}
+
 /* Gimplify OACC_DECLARE.  */
 
 static void
@@ -7936,22 +8043,30 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
 
   clauses = OACC_DECLARE_CLAUSES (expr);
 
-  gimplify_scan_omp_clauses (&clauses, pre_p, ORT_ACC_DATA, OACC_DECLARE);
-
-  gimplify_omp_ctxp->stmt = NULL;
+  gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, OACC_DECLARE);
 
   for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
     {
-      tree attrs, decl = OMP_CLAUSE_DECL (t);
+      tree decl = OMP_CLAUSE_DECL (t);
 
       if (TREE_CODE (decl) == MEM_REF)
        continue;
 
+      if (TREE_CODE (decl) == VAR_DECL
+         && !is_global_var (decl)
+         && DECL_CONTEXT (decl) == current_function_decl)
+       {
+         tree c = gimplify_oacc_declare_1 (t);
+         if (c)
+           {
+             if (oacc_declare_returns == NULL)
+               oacc_declare_returns = new hash_map<tree, tree>;
+
+             oacc_declare_returns->put (decl, c);
+           }
+       }
+
       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,
@@ -7959,23 +8074,6 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
 
   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_ACC_DATA, OACC_DECLARE);
-
-      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;
 }
 
@@ -11215,17 +11313,6 @@ gimplify_body (tree fndecl, bool do_parms)
       gimplify_seq_add_stmt (&seq, outer_stmt);
     }
 
-  if (flag_openacc)
-    while (gimplify_omp_ctxp)
-      {
-       if (gimplify_omp_ctxp->stmt)
-         gimplify_seq_add_stmt (&seq, gimplify_omp_ctxp->stmt);
-
-       struct gimplify_omp_ctx *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 gcc/omp-builtins.def gcc/omp-builtins.def
index 30cf24d..63e5e6e 100644
--- gcc/omp-builtins.def
+++ gcc/omp-builtins.def
@@ -355,7 +355,5 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA,
                  BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, ATTR_NOTHROW_LIST)
 DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS, "GOMP_teams",
                  BT_FN_VOID_UINT_UINT, ATTR_NOTHROW_LIST)
-DEF_GOACC_BUILTIN (BUILT_IN_GOACC_STATIC, "GOACC_register_static",
-                  BT_FN_VOID_PTR_INT_UINT, ATTR_NOTHROW_LIST)
 DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DECLARE, "GOACC_declare",
                   BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
[diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog]
diff --git gcc/testsuite/c-c++-common/goacc/declare-2.c 
gcc/testsuite/c-c++-common/goacc/declare-2.c
index 7979f0c..d24cb22 100644
--- gcc/testsuite/c-c++-common/goacc/declare-2.c
+++ gcc/testsuite/c-c++-common/goacc/declare-2.c
@@ -7,13 +7,13 @@
 /* { dg-error "no valid clauses" "second error" { target *-*-* } 6 } */
 
 int v0[10];
-#pragma acc declare create(v0[1:3]) /* { dg-error "subarray" } */
+#pragma acc declare create(v0[1:3]) /* { dg-error "array section" } */
 
 int v1;
 #pragma acc declare create(v1, v1) /* { dg-error "more than once" } */
 
 int v2;
-#pragma acc declare create(v2) /* { dg-message "previous directive" } */
+#pragma acc declare create(v2)
 #pragma acc declare copyin(v2) /* { dg-error "more than once" } */
 
 int v3;
@@ -37,6 +37,17 @@ int v8;
 int v9;
 #pragma acc declare present_or_create(v9) /* { dg-error "at file scope" } */
 
+int va10;
+#pragma acc declare create (va10)
+#pragma acc declare link (va10) /* { dg-error "more than once" } */
+
+int va11;
+#pragma acc declare link (va11)
+#pragma acc declare link (va11) /* { dg-error "more than once" } */
+
+int va12;
+#pragma acc declare create (va12) link (va12) /* { dg-error "more than once" } 
*/
+
 void
 f (void)
 {
diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c
index 6e82cf4..1f6502a 100644
--- gcc/tree-pretty-print.c
+++ gcc/tree-pretty-print.c
@@ -630,12 +630,6 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, 
int flags)
        case GOMP_MAP_FORCE_DEVICEPTR:
          pp_string (pp, "force_deviceptr");
          break;
-       case GOMP_MAP_DEVICE_RESIDENT:
-         pp_string (pp, "device_resident");
-         break;
-       case GOMP_MAP_LINK:
-         pp_string (pp, "link");
-         break;
        case GOMP_MAP_ALWAYS_TO:
          pp_string (pp, "always,to");
          break;
@@ -660,6 +654,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, 
int flags)
        case GOMP_MAP_ALWAYS_POINTER:
          pp_string (pp, "always_pointer");
          break;
+       case GOMP_MAP_DEVICE_RESIDENT:
+         pp_string (pp, "device_resident");
+         break;
+       case GOMP_MAP_LINK:
+         pp_string (pp, "link");
+         break;
        default:
          gcc_unreachable ();
        }
diff --git gcc/tree.def gcc/tree.def
index 1c0d42c..41064a8 100644
--- gcc/tree.def
+++ gcc/tree.def
@@ -1064,11 +1064,6 @@ 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.  */
@@ -1190,6 +1185,10 @@ DEFTREECODE (OMP_TASKGROUP, "omp_taskgroup", 
tcc_statement, 1)
        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 gcc/tree.h gcc/tree.h
index 2b03de8..f04a2b1 100644
--- gcc/tree.h
+++ gcc/tree.h
@@ -1253,8 +1253,6 @@ 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 gcc/varpool.c gcc/varpool.c
index 3f3d6b2..478f365 100644
--- gcc/varpool.c
+++ gcc/varpool.c
@@ -137,42 +137,6 @@ varpool_node::create_empty (void)
   return node;
 }   
 
-static void
-make_offloadable_1 (varpool_node *node, tree decl ATTRIBUTE_UNUSED)
-{
-  node->offloadable = 1;
-#ifdef ENABLE_OFFLOADING
-  g->have_offload = true;
-  if (!in_lto_p)
-    vec_safe_push (offload_vars, decl);
-  node->force_output = 1;
-#endif
-}
-
-void
-make_offloadable (varpool_node *node, tree decl)
-{
-  tree attrs;
-
-  if (node->offloadable)
-    return;
-
-  if (flag_openmp)
-    {
-      make_offloadable_1 (node, decl);
-      return;
-    }
-
-  attrs = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (decl));
-  if (attrs)
-    {
-      make_offloadable_1 (node, decl);
-
-      DECL_ATTRIBUTES (decl)
-         = remove_attribute ("oacc declare", DECL_ATTRIBUTES (decl));
-    }
-}
-
 /* Return varpool node assigned to DECL.  Create new one when needed.  */
 varpool_node *
 varpool_node::get_create (tree decl)
@@ -180,18 +144,22 @@ varpool_node::get_create (tree decl)
   varpool_node *node = varpool_node::get (decl);
   gcc_checking_assert (TREE_CODE (decl) == VAR_DECL);
   if (node)
-    {
-      if (flag_openacc && !DECL_EXTERNAL (decl))
-       make_offloadable (node, decl);
-      return node;
-    }
+    return node;
 
   node = varpool_node::create_empty ();
   node->decl = decl;
 
   if ((flag_openacc || flag_openmp) && !DECL_EXTERNAL (decl)
       && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)))
-    make_offloadable (node, decl);
+    {
+      node->offloadable = 1;
+#ifdef ENABLE_OFFLOADING
+      g->have_offload = true;
+      if (!in_lto_p)
+       vec_safe_push (offload_vars, decl);
+      node->force_output = 1;
+#endif
+    }
 
   node->register_symbol ();
   return node;
[diff --git include/ChangeLog include/ChangeLog]
diff --git include/gomp-constants.h include/gomp-constants.h
index d2d632e..dffd631 100644
--- include/gomp-constants.h
+++ include/gomp-constants.h
@@ -71,11 +71,12 @@ enum gomp_map_kind
     /* Is a device pointer.  OMP_CLAUSE_SIZE for these is unused; is implicitly
        POINTER_SIZE_UNITS.  */
     GOMP_MAP_FORCE_DEVICEPTR =         (GOMP_MAP_FLAG_SPECIAL_1 | 0),
+    /* Do not map, copy bits for firstprivate instead.  */
     /* OpenACC device_resident.  */
     GOMP_MAP_DEVICE_RESIDENT =         (GOMP_MAP_FLAG_SPECIAL_1 | 1),
     /* OpenACC link.  */
     GOMP_MAP_LINK =                    (GOMP_MAP_FLAG_SPECIAL_1 | 2),
-    /* Do not map, copy bits for firstprivate instead.  */
+    /* Allocate.  */
     GOMP_MAP_FIRSTPRIVATE =            (GOMP_MAP_FLAG_SPECIAL | 0),
     /* Similarly, but store the value in the pointer rather than
        pointed by the pointer.  */
[diff --git libgomp/ChangeLog libgomp/ChangeLog]
diff --git libgomp/libgomp.map libgomp/libgomp.map
index 7a2faf2..98a86cd 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -382,7 +382,6 @@ GOACC_2.0 {
   global:
        GOACC_data_end;
        GOACC_data_start;
-       GOACC_declare;
        GOACC_enter_exit_data;
        GOACC_parallel;
        GOACC_update;
@@ -393,13 +392,13 @@ GOACC_2.0 {
 
 GOACC_2.0.1 {
   global:
+       GOACC_declare;
        GOACC_parallel_keyed;
 } GOACC_2.0;
 
 GOACC_2.0.GOMP_4_BRANCH {
   global:
        GOACC_deviceptr;
-       GOACC_register_static;
        GOMP_set_offload_targets;
 } GOACC_2.0.1;
 
diff --git libgomp/oacc-init.c libgomp/oacc-init.c
index 5e8525f..d6fad51 100644
--- libgomp/oacc-init.c
+++ libgomp/oacc-init.c
@@ -256,8 +256,6 @@ acc_shutdown_1 (acc_device_t d)
   /* Get the base device for this device type.  */
   base_dev = resolve_device (d, true);
 
-  goacc_deallocate_static (d);
-
   ndevs = base_dev->get_num_devices_func ();
 
   /* Unload all the devices of this type that have been opened.  */
@@ -443,9 +441,7 @@ goacc_attach_host_thread_to_device (int ord)
 void
 acc_init (acc_device_t d)
 {
-  bool init = !cached_base_dev;
-
-  if (init)
+  if (!cached_base_dev)
     gomp_init_targets_once ();
 
   gomp_mutex_lock (&acc_device_lock);
@@ -453,9 +449,6 @@ acc_init (acc_device_t d)
   cached_base_dev = acc_init_1 (d);
 
   gomp_mutex_unlock (&acc_device_lock);
-
-  if (init)
-    goacc_allocate_static (d);
   
   goacc_attach_host_thread_to_device (-1);
 }
diff --git libgomp/oacc-int.h libgomp/oacc-int.h
index 560f68b..f11e216 100644
--- libgomp/oacc-int.h
+++ libgomp/oacc-int.h
@@ -99,9 +99,6 @@ void goacc_restore_bind (void);
 void goacc_lazy_initialize (void);
 void goacc_host_init (void);
 
-void goacc_allocate_static (acc_device_t);
-void goacc_deallocate_static (acc_device_t);
-
 #ifdef HAVE_ATTRIBUTE_VISIBILITY
 # pragma GCC visibility pop
 #endif
diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c
index 8dfa622..d66e343 100644
--- libgomp/oacc-parallel.c
+++ libgomp/oacc-parallel.c
@@ -57,68 +57,6 @@ find_pointer (int pos, size_t mapnum, unsigned short *kinds)
   return 0;
 }
 
-static struct oacc_static
-{
-  void *addr;
-  size_t size;
-  unsigned short mask;
-  bool free;
-  struct oacc_static *next;
-} *oacc_statics;
-
-static bool alloc_done = false;
-
-void
-goacc_allocate_static (acc_device_t d)
-{
-  struct oacc_static *s;
-
-  if (alloc_done)
-    assert (0);
-
-  for (s = oacc_statics; s; s = s->next)
-    {
-      void *d;
-
-      switch (s->mask)
-       {
-       case GOMP_MAP_FORCE_ALLOC:
-         break;
-
-       case GOMP_MAP_FORCE_TO:
-         d = acc_deviceptr (s->addr);
-         acc_memcpy_to_device (d, s->addr, s->size);
-         break;
-
-       case GOMP_MAP_FORCE_DEVICEPTR:
-       case GOMP_MAP_DEVICE_RESIDENT:
-       case GOMP_MAP_LINK:
-         break;
-
-       default:
-         assert (0);
-         break;
-       }
-    }
-
-  alloc_done = true;
-}
-
-void
-goacc_deallocate_static (acc_device_t d)
-{
-  struct oacc_static *s;
-  unsigned short mask = GOMP_MAP_FORCE_DEALLOC;
-
-  if (!alloc_done)
-    return;
-
-  for (s = oacc_statics; s; s = s->next)
-    GOACC_enter_exit_data (d, 1, &s->addr, &s->size, &mask, 0, 0);
-
-  alloc_done = false;
-}
-
 static void goacc_wait (int async, int num_waits, va_list *ap);
 
 
@@ -630,21 +568,6 @@ GOACC_get_thread_num (int gang, int worker, int vector)
 }
 
 void
-GOACC_register_static (void *addr, int size, unsigned int mask)
-{
-  struct oacc_static *s;
-
-  s = (struct oacc_static *) malloc (sizeof (struct oacc_static));
-  s->addr = addr;
-  s->size = (size_t) size;
-  s->mask = mask;
-  s->free = false;
-  s->next = oacc_statics;
-
-  oacc_statics = s;
-}
-
-void
 GOACC_declare (int device, size_t mapnum,
               void **hostaddrs, size_t *sizes, unsigned short *kinds)
 {
@@ -673,10 +596,8 @@ GOACC_declare (int device, size_t mapnum,
 
          case GOMP_MAP_ALLOC:
            if (!acc_is_present (hostaddrs[i], sizes[i]))
-             {
-               GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
-                                      &kinds[i], 0, 0);
-             }
+             GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
+                                    &kinds[i], 0, 0);
            break;
 
          case GOMP_MAP_TO:
@@ -688,12 +609,13 @@ GOACC_declare (int device, size_t mapnum,
          case GOMP_MAP_FROM:
            kinds[i] = GOMP_MAP_FORCE_FROM;
            GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
-                                      &kinds[i], 0, 0);
+                                  &kinds[i], 0, 0);
            break;
 
          case GOMP_MAP_FORCE_PRESENT:
            if (!acc_is_present (hostaddrs[i], sizes[i]))
-             gomp_fatal ("[%p,%zd] is not mapped", hostaddrs[i], sizes[i]);
+             gomp_fatal ("[%p,%ld] is not mapped", hostaddrs[i],
+                         (unsigned long) sizes[i]);
            break;
 
          default:
diff --git libgomp/testsuite/libgomp.oacc-c++/declare-1.C 
libgomp/testsuite/libgomp.oacc-c++/declare-1.C
index e82a8e5..0286955 100644
--- libgomp/testsuite/libgomp.oacc-c++/declare-1.C
+++ libgomp/testsuite/libgomp.oacc-c++/declare-1.C
@@ -1,17 +1,20 @@
 /* { dg-do run { target openacc_nvidia_accel_selected } } */
 
+#include <stdlib.h>
+
 template<class T>
 T foo()
 {
-  T a = 0;
+  T a, b;
   #pragma acc declare create (a)
 
-  #pragma acc parallel
+  #pragma acc parallel copyout (b)
   {
     a = 5;
+    b = a;
   }
 
-  return a;
+  return b;
 }
 
 int
@@ -21,5 +24,8 @@ main (void)
 
   rc = foo<int>();
 
-  return rc;
+  if (rc != 5)
+    abort ();
+
+  return 0;
 }
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c 
libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c
index 8fbec4d..c63a68d 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c
@@ -71,8 +71,8 @@ main (int argc, char **argv)
   {
     for (i = 0; i < N; i++)
       {
-        b[i] = a[i];
-        a[i] = b[i];
+       b[i] = a[i];
+       a[i] = b[i];
       }
   }
 
@@ -86,7 +86,7 @@ main (int argc, char **argv)
   {
     for (i = 0; i < N; i++)
       {
-        e[i] = a[i] + d[i];
+       e[i] = a[i] + d[i];
        a[i] = e[i];
       }
   }
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c 
libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c
index 1e2f6ce..38c5de0 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c
@@ -1,11 +1,13 @@
 /* { dg-do run { target openacc_nvidia_accel_selected } } */
 
+#include <stdio.h>
+
 int
 main (int argc, char **argv)
 {
   int a[8] __attribute__((unused));
 
-  __builtin_printf ("CheCKpOInT\n");
+  fprintf (stderr, "CheCKpOInT\n");
 #pragma acc declare present (a)
 }
 
diff --git libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90 
libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
index 18dd1bb..dffbedd 100644
--- libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
+++ libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
@@ -1,4 +1,5 @@
 ! { dg-do run  { target openacc_nvidia_accel_selected } }
+! { dg-xfail-run-if "TODO" { *-*-* } }
 
 module vars
   integer z
diff --git libgomp/testsuite/libgomp.oacc-fortran/declare-3.f90 
libgomp/testsuite/libgomp.oacc-fortran/declare-3.f90
index 79fc011..1d19bfd 100644
--- libgomp/testsuite/libgomp.oacc-fortran/declare-3.f90
+++ libgomp/testsuite/libgomp.oacc-fortran/declare-3.f90
@@ -1,4 +1,5 @@
 ! { dg-do run  { target openacc_nvidia_accel_selected } }
+! { dg-xfail-if "TODO" { *-*-* } }
 
 module globalvars
   real b


Grüße
 Thomas

Reply via email to