diff mbox

[gomp4] declare directive [3/5]

Message ID 5575AEEB.6020902@codesourcery.com
State New
Headers show

Commit Message

James Norris June 8, 2015, 3:04 p.m. UTC

Comments

Thomas Schwinge June 17, 2015, 9:59 a.m. UTC | #1
Hi Jim!

I had mentioned that the Fortran front end changes cause regressions in a
few libgomp execution tests, if configured for Intel MIC (emulation)
offloading.  I have now located *where* this is coming from, but would
you please work on figuring out *why*?

Fortunately, you'll be able to work on the problem even without Intel MIC
(emulation) offloading configured: to reproduce, just look at the
difference in -fdump-tree-original without and with your patch applied.
You'll notice that clauses are getting "lost" from OpenMP target update
directives; for example, for
libgomp/testsuite/libgomp.fortran/declare-target-1.f90 I see:

    --- GOOD/declare-target-2.f90.003t.original     2015-06-16 18:16:07.472763339 +0200
    +++ ./declare-target-2.f90.003t.original        2015-06-16 19:28:22.706845250 +0200
    @@ -3,14 +3,14 @@
       extern integer(kind=4) var_x;
     
       var_x = 10;
    -  #pragma omp target update to(var_x)
    +  #pragma omp target update
       #pragma omp target
         {
           {
             var_x = var_x * 2;
           }
         }
    -  #pragma omp target update from(var_x)
    +  #pragma omp target update
       if (var_x != 20)
         {
           _gfortran_abort ();

(This is the only test case that I looked at, so far.)

I tracked this down to:

On Mon, 8 Jun 2015 10:04:11 -0500, James Norris <jnorris@codesourcery.com> wrote:
> --- a/gcc/fortran/trans-decl.c
> +++ b/gcc/fortran/trans-decl.c

> +static void
> +find_module_oacc_declare_clauses (gfc_symbol *sym)
> +{
> +  if (sym->attr.use_assoc)
> +    {
> +      gfc_omp_map_op map_op;
> +
> +      sym->attr.referenced = sym->attr.oacc_declare_create
> +			     | sym->attr.oacc_declare_copyin
> +			     | sym->attr.oacc_declare_deviceptr
> +			     | sym->attr.oacc_declare_device_resident;
> +
> +      if (sym->attr.oacc_declare_create)
> +	map_op = OMP_MAP_FORCE_ALLOC;
> +
> +      if (sym->attr.oacc_declare_copyin)
> +	map_op = OMP_MAP_FORCE_TO;
> +
> +      if (sym->attr.oacc_declare_deviceptr)
> +	map_op = OMP_MAP_FORCE_DEVICEPTR;
> +
> +      if (sym->attr.oacc_declare_device_resident)
> +	map_op = OMP_MAP_DEVICE_RESIDENT;
> +
> +      if (sym->attr.referenced)
> +	add_clause (sym, map_op);
> +    }
> +}

... this function apparently doing "something inappropriate".  It gets
(unconditionally) called from:

> +finish_oacc_declare (gfc_namespace *ns, enum sym_flavor flavor)
>  {
> [...]
> +  gfc_traverse_ns (ns, find_module_oacc_declare_clauses);

... here, which in turn gets (unconditionally) called from:

> @@ -5946,8 +6237,7 @@ gfc_generate_function_code (gfc_namespace * ns)
>      add_argument_checking (&body, sym);
>  
>    /* Generate !$ACC DECLARE directive. */
> -  if (ns->oacc_declare)
> -    insert_oacc_declare (ns);
> +  finish_oacc_declare (ns, sym->attr.flavor);
>  
>    tmp = gfc_trans_code (ns->code);
>    gfc_add_expr_to_block (&body, tmp);

... here, and:

> --- a/gcc/fortran/trans-stmt.c
> +++ b/gcc/fortran/trans-stmt.c
> @@ -1588,8 +1588,7 @@ gfc_trans_block_construct (gfc_code* code)
>    code->exit_label = exit_label;
>  
>    /* Generate !$ACC DECLARE directive. */
> -  if (ns->oacc_declare)
> -    insert_oacc_declare (ns);
> +  finish_oacc_declare (ns, FL_UNKNOWN);
>  
>    gfc_add_expr_to_block (&body, gfc_trans_code (ns->code));
>    gfc_add_expr_to_block (&body, build1_v (LABEL_EXPR, exit_label));

... here.

Is that sufficient information for you to reproduce the problem?

As soon as you have a patch to bring back the lost clauses in the
-fdump-tree-original, I'll be happy to test it in my Intel MIC (emulated)
offloading build.


Grüße,
 Thomas
diff mbox

Patch

diff --git a/gcc/fortran/f95-lang.c b/gcc/fortran/f95-lang.c
index 5003581..a889342 100644
--- a/gcc/fortran/f95-lang.c
+++ b/gcc/fortran/f95-lang.c
@@ -119,6 +119,8 @@  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, 0, true,  false, false,
     gfc_handle_omp_declare_target_attribute, false },
   { NULL,		  0, 0, false, false, false, NULL, false }
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index e73c269..a90b0f8 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -861,6 +861,13 @@  typedef struct
   /* Mentioned in OMP DECLARE TARGET.  */
   unsigned omp_declare_target:1;
 
+  /* Mentioned in OACC DECLARE.  */
+  unsigned oacc_declare_create:1;
+  unsigned oacc_declare_copyin:1;
+  unsigned oacc_declare_deviceptr:1;
+  unsigned oacc_declare_device_resident:1;
+  unsigned oacc_declare_link:1;
+
   /* This is an OpenACC acclerator function.  */
   unsigned oacc_function:1;
 
@@ -1132,6 +1139,8 @@  typedef enum
   OMP_MAP_FORCE_TOFROM,
   OMP_MAP_FORCE_PRESENT,
   OMP_MAP_FORCE_DEVICEPTR,
+  OMP_MAP_DEVICE_RESIDENT,
+  OMP_MAP_LINK,
   OMP_MAP_FORCE_TO_GANGLOCAL
 }
 gfc_omp_map_op;
@@ -1174,6 +1183,7 @@  enum
   OMP_LIST_FROM,
   OMP_LIST_REDUCTION,
   OMP_LIST_DEVICE_RESIDENT,
+  OMP_LIST_LINK,
   OMP_LIST_USE_DEVICE,
   OMP_LIST_CACHE,
   OMP_LIST_NUM
@@ -1269,6 +1279,7 @@  typedef struct gfc_oacc_declare
 {
   struct gfc_oacc_declare *next;
   locus where;
+  bool module_var;
   gfc_omp_clauses *clauses;
 }
 gfc_oacc_declare;
@@ -3276,6 +3287,6 @@  void gfc_convert_mpz_to_signed (mpz_t, int);
 
 /* trans-decl.c */
 
-void insert_oacc_declare (gfc_namespace *);
+void finish_oacc_declare (gfc_namespace *, enum sym_flavor);
 
 #endif /* GCC_GFORTRAN_H  */
diff --git a/gcc/fortran/module.c b/gcc/fortran/module.c
index 1abfc46..c174902 100644
--- a/gcc/fortran/module.c
+++ b/gcc/fortran/module.c
@@ -1894,7 +1894,9 @@  typedef enum
   AB_IS_CLASS, AB_PROCEDURE, AB_PROC_POINTER, AB_ASYNCHRONOUS, AB_CODIMENSION,
   AB_COARRAY_COMP, AB_VTYPE, AB_VTAB, AB_CONTIGUOUS, AB_CLASS_POINTER,
   AB_IMPLICIT_PURE, AB_ARTIFICIAL, AB_UNLIMITED_POLY, AB_OMP_DECLARE_TARGET,
-  AB_ARRAY_OUTER_DEPENDENCY
+  AB_ARRAY_OUTER_DEPENDENCY, AB_OACC_DECLARE_CREATE, AB_OACC_DECLARE_COPYIN,
+  AB_OACC_DECLARE_DEVICEPTR, AB_OACC_DECLARE_DEVICE_RESIDENT,
+  AB_OACC_DECLARE_LINK
 }
 ab_attribute;
 
@@ -1951,6 +1953,11 @@  static const mstring attr_bits[] =
     minit ("UNLIMITED_POLY", AB_UNLIMITED_POLY),
     minit ("OMP_DECLARE_TARGET", AB_OMP_DECLARE_TARGET),
     minit ("ARRAY_OUTER_DEPENDENCY", AB_ARRAY_OUTER_DEPENDENCY),
+    minit ("OACC_DECLARE_CREATE", AB_OACC_DECLARE_CREATE),
+    minit ("OACC_DECLARE_COPYIN", AB_OACC_DECLARE_COPYIN),
+    minit ("OACC_DECLARE_DEVICEPTR", AB_OACC_DECLARE_DEVICEPTR),
+    minit ("OACC_DECLARE_DEVICE_RESIDENT", AB_OACC_DECLARE_DEVICE_RESIDENT),
+    minit ("OACC_DECLARE_LINK", AB_OACC_DECLARE_LINK),
     minit (NULL, -1)
 };
 
@@ -2133,6 +2140,16 @@  mio_symbol_attribute (symbol_attribute *attr)
 	MIO_NAME (ab_attribute) (AB_OMP_DECLARE_TARGET, attr_bits);
       if (attr->array_outer_dependency)
 	MIO_NAME (ab_attribute) (AB_ARRAY_OUTER_DEPENDENCY, attr_bits);
+      if (attr->oacc_declare_create)
+	MIO_NAME (ab_attribute) (AB_OACC_DECLARE_CREATE, attr_bits);
+      if (attr->oacc_declare_copyin)
+	MIO_NAME (ab_attribute) (AB_OACC_DECLARE_COPYIN, attr_bits);
+      if (attr->oacc_declare_deviceptr)
+	MIO_NAME (ab_attribute) (AB_OACC_DECLARE_DEVICEPTR, attr_bits);
+      if (attr->oacc_declare_device_resident)
+	MIO_NAME (ab_attribute) (AB_OACC_DECLARE_DEVICE_RESIDENT, attr_bits);
+      if (attr->oacc_declare_link)
+	MIO_NAME (ab_attribute) (AB_OACC_DECLARE_LINK, attr_bits);
 
       mio_rparen ();
 
@@ -2302,6 +2319,21 @@  mio_symbol_attribute (symbol_attribute *attr)
 	    case AB_ARRAY_OUTER_DEPENDENCY:
 	      attr->array_outer_dependency =1;
 	      break;
+	    case AB_OACC_DECLARE_CREATE:
+	      attr->oacc_declare_create = 1;
+	      break;
+	    case AB_OACC_DECLARE_COPYIN:
+	      attr->oacc_declare_copyin = 1;
+	      break;
+	    case AB_OACC_DECLARE_DEVICEPTR:
+	      attr->oacc_declare_deviceptr = 1;
+	      break;
+	    case AB_OACC_DECLARE_DEVICE_RESIDENT:
+	      attr->oacc_declare_device_resident = 1;
+	      break;
+	    case AB_OACC_DECLARE_LINK:
+	      attr->oacc_declare_link = 1;
+	      break;
 	    }
 	}
     }
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index fc16d8c..46bf865 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -475,6 +475,7 @@  match_oacc_clause_gang (gfc_omp_clauses *cp)
 #define OMP_CLAUSE_BIND			((uint64_t) 1 << 58)
 #define OMP_CLAUSE_NOHOST		((uint64_t) 1 << 59)
 #define OMP_CLAUSE_DEVICE_TYPE		((uint64_t) 1 << 60)
+#define OMP_CLAUSE_LINK			((uint64_t) 1 << 61)
 
 /* Helper function for OpenACC and OpenMP clauses involving memory
    mapping.  */
@@ -749,6 +750,12 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
 					  true)
 	     == MATCH_YES)
 	continue;
+      if ((mask & OMP_CLAUSE_LINK)
+	  && gfc_match_omp_variable_list ("link (",
+					  &c->lists[OMP_LIST_LINK],
+					  true)
+	     == MATCH_YES)
+	continue;
       if ((mask & OMP_CLAUSE_OACC_DEVICE)
 	  && gfc_match ("device ( ") == MATCH_YES
 	  && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
@@ -1352,7 +1359,7 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
    | OMP_CLAUSE_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_DEVICE_RESIDENT    \
    | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY                          \
    | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT             \
-   | OMP_CLAUSE_PRESENT_OR_CREATE)
+   | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_LINK)
 #define OACC_UPDATE_CLAUSES \
   (OMP_CLAUSE_IF | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST \
    | OMP_CLAUSE_OACC_DEVICE | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE)
@@ -1501,12 +1508,18 @@  gfc_match_oacc_declare (void)
   gfc_omp_namelist *n;
   gfc_namespace *ns = gfc_current_ns;
   gfc_oacc_declare *new_oc, *oc;
-  locus where = gfc_current_locus;
+  bool module_var = false;
 
   if (gfc_match_omp_clauses (&c, OACC_DECLARE_CLAUSES, 0, false, false, true)
       != MATCH_YES)
     return MATCH_ERROR;
 
+  for (n = c->lists[OMP_LIST_DEVICE_RESIDENT]; n != NULL; n = n->next)
+    n->sym->attr.oacc_declare_device_resident = 1;
+
+  for (n = c->lists[OMP_LIST_LINK]; n != NULL; n = n->next)
+    n->sym->attr.oacc_declare_link = 1;
+
   for (n = c->lists[OMP_LIST_MAP]; n != NULL; n = n->next)
     {
       gfc_symbol *s = n->sym;
@@ -1520,6 +1533,14 @@  gfc_match_oacc_declare (void)
 			 "$!ACC DECLARE at %C");
 	      return MATCH_ERROR;
 	    }
+
+	  module_var = true;
+	}
+
+      if (ns->proc_name->attr.oacc_function)
+	{
+	  gfc_error ("Invalid declare in routine with " "$!ACC DECLARE at %C");
+	  return MATCH_ERROR;
 	}
 
       if (s->attr.in_common)
@@ -1543,12 +1564,31 @@  gfc_match_oacc_declare (void)
 		     "$!ACC DECLARE at %C");
 	  return MATCH_ERROR;
 	}
+
+      switch (n->u.map_op)
+	{
+	  case OMP_MAP_FORCE_ALLOC:
+	    s->attr.oacc_declare_create = 1;
+	    break;
+
+	  case OMP_MAP_FORCE_TO:
+	    s->attr.oacc_declare_copyin = 1;
+	    break;
+
+	  case OMP_MAP_FORCE_DEVICEPTR:
+	    s->attr.oacc_declare_deviceptr = 1;
+	    break;
+
+	  default:
+	    break;
+	}
     }
 
   new_oc = gfc_get_oacc_declare ();
   new_oc->next = ns->oacc_declare;
-  new_oc->where = where;
+  new_oc->module_var = module_var;
   new_oc->clauses = c;
+  new_oc->where = gfc_current_locus;
 
   for (oc = new_oc; oc; oc = oc->next)
     {
@@ -4961,6 +5001,33 @@  gfc_resolve_oacc_declare (gfc_namespace *ns)
 			 n->sym->name, &loc);
 	}
     }
+
+  for (oc = ns->oacc_declare; oc; oc = oc->next)
+    {
+      for (list = OMP_LIST_LINK; list <= OMP_LIST_LINK; list++)
+	for (n = oc->clauses->lists[list]; n; n = n->next)
+	  n->sym->mark = 0;
+    }
+
+  for (oc = ns->oacc_declare; oc; oc = oc->next)
+    {
+      for (list = OMP_LIST_LINK; list <= OMP_LIST_LINK; list++)
+	for (n = oc->clauses->lists[list]; n; n = n->next)
+	  {
+	    if (n->sym->mark)
+	      gfc_error ("Symbol %qs present on multiple clauses at %L",
+			 n->sym->name, &loc);
+	    else
+	      n->sym->mark = 1;
+	  }
+    }
+
+  for (oc = ns->oacc_declare; oc; oc = oc->next)
+    {
+      for (list = OMP_LIST_LINK; list <= OMP_LIST_LINK; list++)
+	for (n = oc->clauses->lists[list]; n; n = n->next)
+	  n->sym->mark = 0;
+    }
 }
 
 
diff --git a/gcc/fortran/symbol.c b/gcc/fortran/symbol.c
index b18608b..1ecc16d 100644
--- a/gcc/fortran/symbol.c
+++ b/gcc/fortran/symbol.c
@@ -375,6 +375,11 @@  check_conflict (symbol_attribute *attr, const char *name, locus *where)
     *contiguous = "CONTIGUOUS", *generic = "GENERIC";
   static const char *threadprivate = "THREADPRIVATE";
   static const char *omp_declare_target = "OMP DECLARE TARGET";
+  static const char *oacc_declare_copyin = "OACC DECLARE COPYIN";
+  static const char *oacc_declare_create = "OACC DECLARE CREATE";
+  static const char *oacc_declare_deviceptr = "OACC DECLARE DEVICEPTR";
+  static const char *oacc_declare_device_resident =
+						"OACC DECLARE DEVICE_RESIDENT";
 
   const char *a1, *a2;
   int standard;
@@ -506,6 +511,10 @@  check_conflict (symbol_attribute *attr, const char *name, locus *where)
   conf (in_equivalence, allocatable);
   conf (in_equivalence, threadprivate);
   conf (in_equivalence, omp_declare_target);
+  conf (in_equivalence, oacc_declare_create);
+  conf (in_equivalence, oacc_declare_copyin);
+  conf (in_equivalence, oacc_declare_deviceptr);
+  conf (in_equivalence, oacc_declare_device_resident);
 
   conf (dummy, result);
   conf (entry, result);
@@ -555,6 +564,10 @@  check_conflict (symbol_attribute *attr, const char *name, locus *where)
   conf (cray_pointee, in_equivalence);
   conf (cray_pointee, threadprivate);
   conf (cray_pointee, omp_declare_target);
+  conf (cray_pointee, oacc_declare_create);
+  conf (cray_pointee, oacc_declare_copyin);
+  conf (cray_pointee, oacc_declare_deviceptr);
+  conf (cray_pointee, oacc_declare_device_resident);
 
   conf (data, dummy);
   conf (data, function);
@@ -609,6 +622,10 @@  check_conflict (symbol_attribute *attr, const char *name, locus *where)
   conf (proc_pointer, abstract)
 
   conf (entry, omp_declare_target)
+  conf (entry, oacc_declare_create)
+  conf (entry, oacc_declare_copyin)
+  conf (entry, oacc_declare_deviceptr)
+  conf (entry, oacc_declare_device_resident)
 
   a1 = gfc_code2string (flavors, attr->flavor);
 
@@ -646,6 +663,10 @@  check_conflict (symbol_attribute *attr, const char *name, locus *where)
       conf2 (subroutine);
       conf2 (threadprivate);
       conf2 (omp_declare_target);
+      conf2 (oacc_declare_create);
+      conf2 (oacc_declare_copyin);
+      conf2 (oacc_declare_deviceptr);
+      conf2 (oacc_declare_device_resident);
 
       if (attr->access == ACCESS_PUBLIC || attr->access == ACCESS_PRIVATE)
 	{
@@ -728,6 +749,10 @@  check_conflict (symbol_attribute *attr, const char *name, locus *where)
       conf2 (threadprivate);
       conf2 (result);
       conf2 (omp_declare_target);
+      conf2 (oacc_declare_create);
+      conf2 (oacc_declare_copyin);
+      conf2 (oacc_declare_deviceptr);
+      conf2 (oacc_declare_device_resident);
 
       if (attr->intent != INTENT_UNKNOWN)
 	{
@@ -1239,6 +1264,62 @@  gfc_add_omp_declare_target (symbol_attribute *attr, const char *name,
 
 
 bool
+gfc_add_oacc_declare_create (symbol_attribute *attr, const char *name, locus *where)
+{
+  if (check_used (attr, name, where))
+    return false;
+
+  if (attr->oacc_declare_create)
+    return true;
+
+  attr->oacc_declare_create = 1;
+  return check_conflict (attr, name, where);
+}
+
+
+bool
+gfc_add_oacc_declare_copyin (symbol_attribute *attr, const char *name, locus *where)
+{
+  if (check_used (attr, name, where))
+    return false;
+
+  if (attr->oacc_declare_copyin)
+    return true;
+
+  attr->oacc_declare_copyin = 1;
+  return check_conflict (attr, name, where);
+}
+
+
+bool
+gfc_add_oacc_declare_deviceptr (symbol_attribute *attr, const char *name, locus *where)
+{
+  if (check_used (attr, name, where))
+    return false;
+
+  if (attr->oacc_declare_deviceptr)
+    return true;
+
+  attr->oacc_declare_deviceptr = 1;
+  return check_conflict (attr, name, where);
+}
+
+
+bool
+gfc_add_oacc_declare_device_resident (symbol_attribute *attr, const char *name, locus *where)
+{
+  if (check_used (attr, name, where))
+    return false;
+
+  if (attr->oacc_declare_device_resident)
+    return true;
+
+  attr->oacc_declare_device_resident = 1;
+  return check_conflict (attr, name, where);
+}
+
+
+bool
 gfc_add_target (symbol_attribute *attr, locus *where)
 {
 
@@ -1796,6 +1877,18 @@  gfc_copy_attr (symbol_attribute *dest, symbol_attribute *src, locus *where)
   if (src->omp_declare_target
       && !gfc_add_omp_declare_target (dest, NULL, where))
     goto fail;
+  if (src->oacc_declare_create
+      && !gfc_add_oacc_declare_create (dest, NULL, where))
+    goto fail;
+  if (src->oacc_declare_copyin
+      && !gfc_add_oacc_declare_copyin (dest, NULL, where))
+    goto fail;
+  if (src->oacc_declare_deviceptr
+      && !gfc_add_oacc_declare_deviceptr (dest, NULL, where))
+    goto fail;
+  if (src->oacc_declare_device_resident
+      && !gfc_add_oacc_declare_device_resident (dest, NULL, where))
+    goto fail;
   if (src->target && !gfc_add_target (dest, where))
     goto fail;
   if (src->dummy && !gfc_add_dummy (dest, NULL, where))
diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c
index 6cdc472..77fdc8b 100644
--- a/gcc/fortran/trans-decl.c
+++ b/gcc/fortran/trans-decl.c
@@ -1309,6 +1309,16 @@  add_attributes_to_decl (symbol_attribute sym_attr, tree list)
     list = tree_cons (get_identifier ("omp declare target"),
 		      NULL_TREE, list);
 
+  if (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);
+    }
+
   if (sym_attr.oacc_function)
     list = tree_cons (get_identifier ("oacc function"),
 		      NULL_TREE, list);
@@ -5754,14 +5764,49 @@  is_ieee_module_used (gfc_namespace *ns)
 }
 
 
+static struct oacc_return
+{
+  gfc_code *code;
+  struct oacc_return *next;
+} *oacc_returns;
+
+
+static void
+find_oacc_return (gfc_code *code)
+{
+  if (code->next)
+    {
+      if (code->next->op == EXEC_RETURN)
+	{
+	  struct oacc_return *r;
+
+	  r = XCNEW (struct oacc_return);
+	  r->code = code;
+	  r->next = NULL;
+
+	  if (oacc_returns)
+	    r->next = oacc_returns;
+
+	  oacc_returns = r;
+	}
+      else
+	{
+	  find_oacc_return (code->next);
+	}
+    }
+
+  if (code->block)
+    find_oacc_return (code->block);
+
+  return;
+}
+
+
 static gfc_code *
 find_end (gfc_code *code)
 {
   gcc_assert (code);
 
-  if (code->op == EXEC_END_PROCEDURE)
-    return code;
-
   if (code->next)
     {
       if (code->next->op == EXEC_END_PROCEDURE)
@@ -5774,38 +5819,284 @@  find_end (gfc_code *code)
 }
 
 
+static gfc_omp_clauses *module_oacc_clauses;
+
+
+static void
+add_clause (gfc_symbol *sym, gfc_omp_map_op map_op)
+{
+  gfc_omp_namelist *n;
+
+  n = gfc_get_omp_namelist ();
+  n->sym = sym;
+  n->u.map_op = map_op;
+
+  if (!module_oacc_clauses)
+    module_oacc_clauses = gfc_get_omp_clauses ();
+
+  if (module_oacc_clauses->lists[OMP_LIST_MAP])
+    n->next = module_oacc_clauses->lists[OMP_LIST_MAP];
+
+  module_oacc_clauses->lists[OMP_LIST_MAP] = n;
+}
+
+
+static void
+find_module_oacc_declare_clauses (gfc_symbol *sym)
+{
+  if (sym->attr.use_assoc)
+    {
+      gfc_omp_map_op map_op;
+
+      sym->attr.referenced = sym->attr.oacc_declare_create
+			     | sym->attr.oacc_declare_copyin
+			     | sym->attr.oacc_declare_deviceptr
+			     | sym->attr.oacc_declare_device_resident;
+
+      if (sym->attr.oacc_declare_create)
+	map_op = OMP_MAP_FORCE_ALLOC;
+
+      if (sym->attr.oacc_declare_copyin)
+	map_op = OMP_MAP_FORCE_TO;
+
+      if (sym->attr.oacc_declare_deviceptr)
+	map_op = OMP_MAP_FORCE_DEVICEPTR;
+
+      if (sym->attr.oacc_declare_device_resident)
+	map_op = OMP_MAP_DEVICE_RESIDENT;
+
+      if (sym->attr.referenced)
+	add_clause (sym, map_op);
+    }
+}
+
+
 void
-insert_oacc_declare (gfc_namespace *ns)
+finish_oacc_declare (gfc_namespace *ns, enum sym_flavor flavor)
 {
-  gfc_code *code;
+  gfc_code *code, *end_c, *code2;
+  gfc_oacc_declare *oc;
+  gfc_omp_clauses *omp_clauses = NULL, *ret_clauses = NULL;
+  gfc_omp_namelist *n;
+  locus where = gfc_current_locus;
+
+  gfc_traverse_ns (ns, find_module_oacc_declare_clauses);
+
+  if (module_oacc_clauses && flavor == FL_PROGRAM)
+    {
+      gfc_oacc_declare *new_oc;
+
+      new_oc = gfc_get_oacc_declare ();
+      new_oc->next = ns->oacc_declare;
+      new_oc->clauses = module_oacc_clauses;
+
+      ns->oacc_declare = new_oc;
+      module_oacc_clauses = NULL;
+    }
+
+  if (!ns->oacc_declare)
+    return;
+
+  for (oc = ns->oacc_declare; oc; oc = oc->next)
+    {
+      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 = ns->oacc_declare->where;
+  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;
+
+	  new_n = gfc_get_omp_namelist ();
+	  new_n->sym = n->sym;
+	  new_n->u.map_op = new_op;
+
+	  if (!ret_clauses)
+	    ret_clauses = gfc_get_omp_clauses ();
+
+	  if (ret_clauses->lists[OMP_LIST_MAP])
+	    new_n->next = ret_clauses->lists[OMP_LIST_MAP];
+
+	  ret_clauses->lists[OMP_LIST_MAP] = new_n;
+	  ret = false;
+	}
+    }
 
-  code->ext.oacc_declare = ns->oacc_declare;
+  if (!ret_clauses)
+    {
+      code->next = ns->code;
+      ns->code = code;
+      return;
+    }
 
-  code->block = XCNEW (gfc_code);
-  code->block->op = EXEC_OACC_DECLARE;
-  code->block->loc = ns->oacc_declare->where;
+  code2 = XCNEW (gfc_code);
+  code2->op = EXEC_OACC_DECLARE;
+  code2->loc = where;
+  code2->ext.omp_clauses = ret_clauses;
 
   if (ns->code)
     {
-      gfc_code *c;
+      find_oacc_return (ns->code);
+
+      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;
+
+	      last = ns->code;
+
+	      while (last->next)
+		last = last->next;
+
+	      last->next = code2;
+	      code->next = ns->code;
+	    }
+	}
+    }
+  else
+    {
+      code->next = code2;
+    }
+
+  while (oacc_returns)
+    {
+      struct oacc_return *r;
+
+      r = oacc_returns;
 
-      c = find_end (ns->code);
-      if (c)
+      ret_clauses = gfc_get_omp_clauses ();
+
+      for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
 	{
-	  code->next = c->next;
-	  c->next = NULL;
+	  if (n->u.map_op == OMP_MAP_FORCE_ALLOC
+	      || n->u.map_op == OMP_MAP_FORCE_TO)
+	    {
+	      gfc_omp_namelist *new_n;
+
+	      new_n = gfc_get_omp_namelist ();
+	      new_n->sym = n->sym;
+	      new_n->u.map_op = OMP_MAP_FORCE_DEALLOC;
+
+	      if (ret_clauses->lists[OMP_LIST_MAP])
+		new_n->next = ret_clauses->lists[OMP_LIST_MAP];
+
+	      ret_clauses->lists[OMP_LIST_MAP] = new_n;
+	    }
 	}
 
-      code->block->next = ns->code;
-      code->block->ext.oacc_declare = NULL;
+      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;
-  ns->oacc_declare = NULL;
+    ns->code = code;
 }
 
 
@@ -5946,8 +6237,7 @@  gfc_generate_function_code (gfc_namespace * ns)
     add_argument_checking (&body, sym);
 
   /* Generate !$ACC DECLARE directive. */
-  if (ns->oacc_declare)
-    insert_oacc_declare (ns);
+  finish_oacc_declare (ns, sym->attr.flavor);
 
   tmp = gfc_trans_code (ns->code);
   gfc_add_expr_to_block (&body, tmp);
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 1aa33c0..f73e366 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1784,12 +1784,12 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	case OMP_LIST_USE_DEVICE:
 	  clause_code = OMP_CLAUSE_USE_DEVICE;
 	  goto add_clause;
-	case OMP_LIST_DEVICE_RESIDENT:
-	  clause_code = OMP_CLAUSE_DEVICE_RESIDENT;
-	  goto add_clause;
 	case OMP_LIST_CACHE:
 	  clause_code = OMP_CLAUSE__CACHE_;
 	  goto add_clause;
+	case OMP_LIST_DEVICE_RESIDENT:
+	case OMP_LIST_LINK:
+	  continue;
 
 	add_clause:
 	  omp_clauses
@@ -1937,6 +1937,9 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	      if (!n->sym->attr.referenced)
 		continue;
 
+	      if (n->sym->attr.use_assoc && n->sym->attr.oacc_declare_link)
+		continue;
+
 	      tree node = build_omp_clause (input_location, OMP_CLAUSE_MAP);
 	      tree node2 = NULL_TREE;
 	      tree node3 = NULL_TREE;
@@ -2160,6 +2163,9 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		case OMP_MAP_FORCE_TO_GANGLOCAL:
 		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TO_GANGLOCAL);
 		  break;
+		case OMP_MAP_DEVICE_RESIDENT:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DEVICE_RESIDENT);
+		  break;
 		default:
 		  gcc_unreachable ();
 		}
@@ -4391,25 +4397,18 @@  tree
 gfc_trans_oacc_declare (gfc_code *code)
 {
   stmtblock_t block;
-  struct gfc_oacc_declare *d;
-  tree stmt, clauses = NULL_TREE;
+  tree stmt, oacc_clauses;
+  enum tree_code construct_code;
 
   gfc_start_block (&block);
 
-  for (d = code->ext.oacc_declare; d; d = d->next)
-    {
-      tree t;
-
-      t = gfc_trans_omp_clauses (&block, d->clauses, d->clauses->loc);
+  construct_code = OACC_DECLARE;
 
-      if (clauses)
-	OMP_CLAUSE_CHAIN (clauses) = t;
-      else
-	clauses = t;
-    }
-
-  stmt = gfc_trans_omp_code (code->block->next, true);
-  stmt = build2_loc (input_location, OACC_DATA, void_type_node, stmt, clauses);
+  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);
   gfc_add_expr_to_block (&block, stmt);
   return gfc_finish_block (&block);
 }
diff --git a/gcc/fortran/trans-stmt.c b/gcc/fortran/trans-stmt.c
index c6be9ad..352b383 100644
--- a/gcc/fortran/trans-stmt.c
+++ b/gcc/fortran/trans-stmt.c
@@ -1588,8 +1588,7 @@  gfc_trans_block_construct (gfc_code* code)
   code->exit_label = exit_label;
 
   /* Generate !$ACC DECLARE directive. */
-  if (ns->oacc_declare)
-    insert_oacc_declare (ns);
+  finish_oacc_declare (ns, FL_UNKNOWN);
 
   gfc_add_expr_to_block (&body, gfc_trans_code (ns->code));
   gfc_add_expr_to_block (&body, build1_v (LABEL_EXPR, exit_label));
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index 067882f..cc11d11 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -147,6 +147,7 @@  DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I4_INT, BT_VOID, BT_VOLATILE_PTR, BT_I4, BT
 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_INT_INT_INT_INT, BT_INT, BT_INT, BT_INT, 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)