[4/4] OpenACC 2.6 manual deep copy support (attach/detach)
diff mbox series

Message ID 20191006223237.81842-5-julian@codesourcery.com
State New
Headers show
Series
  • OpenACC 2.6 manual deep copy (repost)
Related show

Commit Message

Julian Brown Oct. 6, 2019, 10:32 p.m. UTC
This patch provides support for OpenACC 2.6 manual deep copy, without the
reference-counting changes included in the previous version of the patch:

  https://gcc.gnu.org/ml/gcc-patches/2018-12/msg01084.html

Relative to that version of the patch, I've introduced a
GOMP_MAP_ATTACH_DETACH mapping type (only used internally in GCC,
not exposed to libgomp) to distinguish the OpenACC semantics from the
previously-used GOMP_MAP_ALWAYS_POINTER mapping, and reimplementing the
Fortran derived-type component mapping to fix some bugs, improve code
re-use, prepare for "class" support (a forthcoming patch), and to keep
better separation between OpenACC-specific code and OpenMP.

This version of the patch also includes several new tests.

OK for trunk?

Thanks,

Julian

ChangeLog

	gcc/c-family/
	* c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_ATTACH and
	PRAGMA_OACC_CLAUSE_DETACH.

	gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Add parsing of attach and
	detach clauses.
	(c_parser_omp_variable_list): Add ALLOW_DEREF optional parameter.
	Allow deref (->) in variable lists if true.
	(c_parser_omp_var_list_parens): Add ALLOW_DEREF optional parameter.
	Pass to c_parser_omp_variable_list.
	(c_parser_oacc_data_clause): Support attach and detach clauses.  Update
	call to c_parser_omp_variable_list.
	(c_parser_oacc_all_clauses): Support attach and detach clauses.
	(OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK)
	(OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK): Add
	PRAGMA_OACC_CLAUSE_ATTACH.
	(OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH.
	* c-typeck.c (handle_omp_array_sections_1): Reject subarrays for attach
	and detach.  Support deref.
	(handle_omp_array_sections): Use GOMP_MAP_ATTACH_DETACH instead of
	GOMP_MAP_ALWAYS_POINTER for OpenACC.
	(c_oacc_check_attachments): New function.
	(c_finish_omp_clauses): Check attach/detach arguments for being
	pointers using above.  Support deref.

	gcc/cp/
	* parser.c (cp_parser_omp_clause_name): Support attach and detach
	clauses.
	(cp_parser_omp_var_list_no_open): Add ALLOW_DEREF optional parameter.
	Parse deref if true.
	(cp_parser_omp_var_list): Add ALLOW_DEREF optional parameter.  Pass to
	cp_parser_omp_var_list_no_open.
	(cp_parser_oacc_data_clause): Support attach and detach clauses.
	Update call to cp_parser_omp_var_list_no_open.
	(cp_parser_oacc_all_clauses): Support attach and detach.
	(OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK)
	(OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK): Add
	PRAGMA_OACC_CLAUSE_ATTACH.
	(OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH.
	* semantics.c (handle_omp_array_sections_1): Reject subarrays for
	attach and detach.
	(handle_omp_array_sections): Use GOMP_MAP_ATTACH_DETACH instead of
	GOMP_MAP_ALWAYS_POINTER for OpenACC.
	(cp_oacc_check_attachments): New function.
	(finish_omp_clauses): Use above function.  Allow structure fields and
	class members to appear in OpenACC data clauses.  Support
	GOMP_MAP_ATTACH_DETACH.  Support deref.

	gcc/fortran/
	* gfortran.h (gfc_omp_map_op): Add OMP_MAP_ATTACH, OMP_MAP_DETACH.
	* openmp.c (gfc_match_omp_variable_list): Add allow_derived parameter.
	Parse derived-type member accesses if true.
	(omp_mask2): Add OMP_CLAUSE_ATTACH and OMP_CLAUSE_DETACH.
	(gfc_match_omp_map_clause): Add allow_derived parameter.  Pass to
	gfc_match_omp_variable_list.
	(gfc_match_omp_clauses): Support attach and detach.  Support derived
	types for appropriate OpenACC directives.
	(OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES, OACC_DATA_CLAUSES)
	(OACC_ENTER_DATA_CLAUSES): Add OMP_CLAUSE_ATTACH.
	(OACC_EXIT_DATA_CLAUSES): Add OMP_CLAUSE_DETACH.
	(check_symbol_not_pointer): Don't disallow pointer objects of derived
	type.
	(resolve_oacc_data_clauses): Don't disallow allocatable derived types.
	(resolve_omp_clauses): Perform duplicate checking only for non-derived
	type component accesses (plain variables and arrays or array sections).
	Support component refs.
	* trans-expr.c (gfc_conv_component_ref,
	conv_parent_component_references): Make global.
	(gfc_auto_dereference_var): New function, broken out of...
	(gfc_conv_variable): ...here.  Call above function.
	* trans-openmp.c (gfc_omp_privatize_by_reference): Support component
	refs.
	(gfc_trans_omp_array_section): New function, broken out of...
	(gfc_trans_omp_clauses): ...here.  Support component refs/derived
	types, attach and detach clauses.
	* trans.h (gfc_conv_component_ref, conv_parent_component_references,
	gfc_auto_dereference_var): Add prototypes.

	gcc/
	* gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_HAS_ATTACHMENTS.
	(insert_struct_comp_map): Support derived-type member mappings
	for arrays with descriptors which use GOMP_MAP_TO_PSET.  Support
	GOMP_MAP_ATTACH_DETACH.
	(gimplify_scan_omp_clauses): Tidy up OACC_ENTER_DATA/OACC_EXIT_DATA
	mappings.  Handle attach/detach clauses and component references.
	(gimplify_adjust_omp_clauses_1): Skip adjustments for explicit
	attach/detach clauses.
	(gimplify_omp_target_update): Handle finalize for detach.
	* omp-low.c (lower_omp_target): Support GOMP_MAP_ATTACH,
	GOMP_MAP_DETACH, GOMP_MAP_FORCE_DETACH.
	* tree-pretty-print.c (dump_omp_clause): Likewise, plus
	GOMP_MAP_ATTACH_DETACH.

	include/
	* gomp-constants.h (GOMP_MAP_DEEP_COPY): Define.
	(gomp_map_kind): Add GOMP_MAP_ATTACH, GOMP_MAP_DETACH,
	GOMP_MAP_FORCE_DETACH.

	gcc/testsuite/
	* c-c++-common/goacc/deep-copy-arrayofstruct.c: New test.
	* c-c++-common/goacc/mdc-1.c: New test.
	* c-c++-common/goacc/mdc-2.c: New test.
	* gcc.dg/goacc/mdc.C: New test.
	* gfortran.dg/goacc/data-clauses.f95: New test.
	* gfortran.dg/goacc/derived-types.f90: New test.
	* gfortran.dg/goacc/derived-types-2.f90: New test.
	* gfortran.dg/goacc/enter-exit-data.f95: New test.

	libgomp/
	* libgomp.h (struct target_var_desc): Add do_detach flag.
	(struct splay_tree_key_s): Add attach_count field to union.
	(gomp_attach_pointer, gomp_detach_pointer): Add prototypes.
	* libgomp.map (OACC_2.6): New section. Add acc_attach,
	acc_attach_async, acc_detach, acc_detach_async, acc_detach_finalize,
	acc_detach_finalize_async.
	* oacc-mem.c (acc_attach_async, acc_attach, goacc_detach_internal,
	acc_detach, acc_detach_async, acc_detach_finalize,
	acc_detach_finalize_async): New functions.
	* oacc-parallel.c (find_pointer): Support attach/detach.
	(GOACC_enter_exit_data): Support attach/detach and GOMP_MAP_STRUCT.
	Don't call gomp_acc_insert_pointer.
	* openacc.h (acc_attach, acc_attach_async, acc_detach,
	(acc_detach_async, acc_detach_finalize, acc_detach_finalize_async): Add
	prototypes.
	* target.c (dump_tgt): Support attach_count field.
	(gomp_map_vars_existing): Initialise do_detach field of tgt_var_desc.
	(gomp_attach_pointer, gomp_detach_pointer): New functions.
	(gomp_map_vars_internal): Support attach and detach.
	(gomp_remove_var_internal): Free attach count array if present.
	(gomp_unmap_vars_internal): Support detach.
	(gomp_load_image_to_device): Zero-initialise attach_count field.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c: New test.
	* testsuite/libgomp.oacc-c++/deep-copy-12.C: Mew test.
	* testsuite/libgomp.oacc-c++/deep-copy-13.C: Mew test.
	* testsuite/libgomp.oacc-fortran/deep-copy-1.f90: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-2.f90: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-3.f90: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-4.f90: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-5.f90: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-6.f90: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-7.f90: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-8.f90: New test.
	* testsuite/libgomp.oacc-fortran/derived-type-1.f90: New test.
	* testsuite/libgomp.oacc-fortran/derivedtype-1.f95: New test.
	* testsuite/libgomp.oacc-fortran/derivedtype-2.f95: New test.
	* testsuite/libgomp.oacc-fortran/multidim-slice.f95: New test.
	* testsuite/libgomp.oacc-fortran/update-2.f90: New test.
---
 gcc/c-family/c-pragma.h                       |   2 +
 gcc/c/c-parser.c                              |  52 +++-
 gcc/c/c-typeck.c                              |  76 ++++-
 gcc/cp/parser.c                               |  55 +++-
 gcc/cp/semantics.c                            |  98 +++++-
 gcc/fortran/gfortran.h                        |   2 +
 gcc/fortran/openmp.c                          | 150 ++++++---
 gcc/fortran/trans-expr.c                      | 184 +++++------
 gcc/fortran/trans-openmp.c                    | 290 ++++++++++++++----
 gcc/fortran/trans.h                           |   9 +
 gcc/gimplify.c                                | 224 ++++++++++++--
 gcc/omp-low.c                                 |   3 +
 .../goacc/deep-copy-arrayofstruct.c           |  84 +++++
 gcc/testsuite/c-c++-common/goacc/mdc-1.c      |  55 ++++
 gcc/testsuite/c-c++-common/goacc/mdc-2.c      |  62 ++++
 gcc/testsuite/g++.dg/goacc/mdc.C              |  68 ++++
 .../gfortran.dg/goacc/data-clauses.f95        |  38 +--
 .../gfortran.dg/goacc/derived-types-2.f90     |  14 +
 .../gfortran.dg/goacc/derived-types.f90       |  77 +++++
 .../gfortran.dg/goacc/enter-exit-data.f95     |  24 +-
 gcc/tree-pretty-print.c                       |  18 ++
 include/gomp-constants.h                      |  16 +-
 libgomp/libgomp.h                             |  12 +
 libgomp/libgomp.map                           |  10 +
 libgomp/oacc-mem.c                            |  82 ++++-
 libgomp/oacc-parallel.c                       | 193 +++++++++---
 libgomp/openacc.h                             |   6 +
 libgomp/target.c                              | 192 +++++++++++-
 .../testsuite/libgomp.oacc-c++/deep-copy-12.C |  72 +++++
 .../testsuite/libgomp.oacc-c++/deep-copy-13.C |  72 +++++
 .../libgomp.oacc-c-c++-common/deep-copy-1.c   |  24 ++
 .../libgomp.oacc-c-c++-common/deep-copy-10.c  |  53 ++++
 .../libgomp.oacc-c-c++-common/deep-copy-11.c  |  72 +++++
 .../libgomp.oacc-c-c++-common/deep-copy-14.c  |  63 ++++
 .../libgomp.oacc-c-c++-common/deep-copy-2.c   |  29 ++
 .../libgomp.oacc-c-c++-common/deep-copy-3.c   |  34 ++
 .../libgomp.oacc-c-c++-common/deep-copy-4.c   |  87 ++++++
 .../libgomp.oacc-c-c++-common/deep-copy-5.c   |  81 +++++
 .../libgomp.oacc-c-c++-common/deep-copy-6.c   |  59 ++++
 .../libgomp.oacc-c-c++-common/deep-copy-7.c   |  45 +++
 .../libgomp.oacc-c-c++-common/deep-copy-8.c   |  54 ++++
 .../libgomp.oacc-c-c++-common/deep-copy-9.c   |  53 ++++
 .../libgomp.oacc-fortran/deep-copy-1.f90      |  35 +++
 .../libgomp.oacc-fortran/deep-copy-2.f90      |  33 ++
 .../libgomp.oacc-fortran/deep-copy-3.f90      |  34 ++
 .../libgomp.oacc-fortran/deep-copy-4.f90      |  49 +++
 .../libgomp.oacc-fortran/deep-copy-5.f90      |  57 ++++
 .../libgomp.oacc-fortran/deep-copy-6.f90      |  61 ++++
 .../libgomp.oacc-fortran/deep-copy-7.f90      |  89 ++++++
 .../libgomp.oacc-fortran/deep-copy-8.f90      |  41 +++
 .../libgomp.oacc-fortran/derived-type-1.f90   |  28 ++
 .../libgomp.oacc-fortran/derivedtype-1.f95    |  30 ++
 .../libgomp.oacc-fortran/derivedtype-2.f95    |  41 +++
 .../libgomp.oacc-fortran/multidim-slice.f95   |  50 +++
 .../libgomp.oacc-fortran/update-2.f90         | 284 +++++++++++++++++
 55 files changed, 3389 insertions(+), 337 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/mdc-1.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/mdc-2.c
 create mode 100644 gcc/testsuite/g++.dg/goacc/mdc.C
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/derived-types-2.f90
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/derived-types.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C
 create mode 100644 libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-1.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-4.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-5.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-7.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-8.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/derived-type-1.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/derivedtype-1.f95
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/derivedtype-2.f95
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/multidim-slice.f95
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/update-2.f90

Patch
diff mbox series

diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index e0aa774555a..0bbfd35d261 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -142,11 +142,13 @@  enum pragma_omp_clause {
 
   /* Clauses for OpenACC.  */
   PRAGMA_OACC_CLAUSE_ASYNC,
+  PRAGMA_OACC_CLAUSE_ATTACH,
   PRAGMA_OACC_CLAUSE_AUTO,
   PRAGMA_OACC_CLAUSE_COPY,
   PRAGMA_OACC_CLAUSE_COPYOUT,
   PRAGMA_OACC_CLAUSE_CREATE,
   PRAGMA_OACC_CLAUSE_DELETE,
+  PRAGMA_OACC_CLAUSE_DETACH,
   PRAGMA_OACC_CLAUSE_DEVICEPTR,
   PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
   PRAGMA_OACC_CLAUSE_FINALIZE,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 6957297b3a5..85001394317 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -11745,6 +11745,8 @@  c_parser_omp_clause_name (c_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_ALIGNED;
 	  else if (!strcmp ("async", p))
 	    result = PRAGMA_OACC_CLAUSE_ASYNC;
+	  else if (!strcmp ("attach", p))
+	    result = PRAGMA_OACC_CLAUSE_ATTACH;
 	  break;
 	case 'b':
 	  if (!strcmp ("bind", p))
@@ -11771,6 +11773,8 @@  c_parser_omp_clause_name (c_parser *parser)
 	    result = PRAGMA_OACC_CLAUSE_DELETE;
 	  else if (!strcmp ("depend", p))
 	    result = PRAGMA_OMP_CLAUSE_DEPEND;
+	  else if (!strcmp ("detach", p))
+	    result = PRAGMA_OACC_CLAUSE_DETACH;
 	  else if (!strcmp ("device", p))
 	    result = PRAGMA_OMP_CLAUSE_DEVICE;
 	  else if (!strcmp ("deviceptr", p))
@@ -12014,12 +12018,16 @@  c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list)
    If KIND is nonzero, CLAUSE_LOC is the location of the clause.
 
    If KIND is zero, create a TREE_LIST with the decl in TREE_PURPOSE;
-   return the list created.  */
+   return the list created.
+
+   The optional ALLOW_DEREF argument is true if list items can use the deref
+   (->) operator.  */
 
 static tree
 c_parser_omp_variable_list (c_parser *parser,
 			    location_t clause_loc,
-			    enum omp_clause_code kind, tree list)
+			    enum omp_clause_code kind, tree list,
+			    bool allow_deref = false)
 {
   auto_vec<c_token> tokens;
   unsigned int tokens_avail = 0;
@@ -12146,9 +12154,13 @@  c_parser_omp_variable_list (c_parser *parser,
 	    case OMP_CLAUSE_MAP:
 	    case OMP_CLAUSE_FROM:
 	    case OMP_CLAUSE_TO:
-	      while (c_parser_next_token_is (parser, CPP_DOT))
+	      while (c_parser_next_token_is (parser, CPP_DOT)
+		     || (allow_deref
+			 && c_parser_next_token_is (parser, CPP_DEREF)))
 		{
 		  location_t op_loc = c_parser_peek_token (parser)->location;
+		  if (c_parser_next_token_is (parser, CPP_DEREF))
+		    t = build_simple_mem_ref (t);
 		  c_parser_consume_token (parser);
 		  if (!c_parser_next_token_is (parser, CPP_NAME))
 		    {
@@ -12270,11 +12282,12 @@  c_parser_omp_variable_list (c_parser *parser,
 }
 
 /* Similarly, but expect leading and trailing parenthesis.  This is a very
-   common case for OpenACC and OpenMP clauses.  */
+   common case for OpenACC and OpenMP clauses.  The optional ALLOW_DEREF
+   argument is true if list items can use the deref (->) operator.  */
 
 static tree
 c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
-			      tree list)
+			      tree list, bool allow_deref = false)
 {
   /* The clauses location.  */
   location_t loc = c_parser_peek_token (parser)->location;
@@ -12282,18 +12295,20 @@  c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
   matching_parens parens;
   if (parens.require_open (parser))
     {
-      list = c_parser_omp_variable_list (parser, loc, kind, list);
+      list = c_parser_omp_variable_list (parser, loc, kind, list, allow_deref);
       parens.skip_until_found_close (parser);
     }
   return list;
 }
 
-/* OpenACC 2.0:
+/* OpenACC 2.0+:
+   attach ( variable-list )
    copy ( variable-list )
    copyin ( variable-list )
    copyout ( variable-list )
    create ( variable-list )
    delete ( variable-list )
+   detach ( variable-list )
    present ( variable-list ) */
 
 static tree
@@ -12303,6 +12318,9 @@  c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
   enum gomp_map_kind kind;
   switch (c_kind)
     {
+    case PRAGMA_OACC_CLAUSE_ATTACH:
+      kind = GOMP_MAP_ATTACH;
+      break;
     case PRAGMA_OACC_CLAUSE_COPY:
       kind = GOMP_MAP_TOFROM;
       break;
@@ -12318,6 +12336,9 @@  c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
     case PRAGMA_OACC_CLAUSE_DELETE:
       kind = GOMP_MAP_RELEASE;
       break;
+    case PRAGMA_OACC_CLAUSE_DETACH:
+      kind = GOMP_MAP_DETACH;
+      break;
     case PRAGMA_OACC_CLAUSE_DEVICE:
       kind = GOMP_MAP_FORCE_TO;
       break;
@@ -12337,7 +12358,7 @@  c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
       gcc_unreachable ();
     }
   tree nl, c;
-  nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list);
+  nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true);
 
   for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
     OMP_CLAUSE_SET_MAP_KIND (c, kind);
@@ -15052,6 +15073,10 @@  c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 						 clauses);
 	  c_name = "auto";
 	  break;
+	case PRAGMA_OACC_CLAUSE_ATTACH:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "attach";
+	  break;
 	case PRAGMA_OACC_CLAUSE_COLLAPSE:
 	  clauses = c_parser_omp_clause_collapse (parser, clauses);
 	  c_name = "collapse";
@@ -15080,6 +15105,10 @@  c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  clauses = c_parser_omp_clause_default (parser, clauses, true);
 	  c_name = "default";
 	  break;
+	case PRAGMA_OACC_CLAUSE_DETACH:
+	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "detach";
+	  break;
 	case PRAGMA_OACC_CLAUSE_DEVICE:
 	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
 	  c_name = "device";
@@ -15574,7 +15603,8 @@  c_parser_oacc_cache (location_t loc, c_parser *parser)
 */
 
 #define OACC_DATA_CLAUSE_MASK						\
-	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY)		\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE)		\
@@ -15757,6 +15787,7 @@  c_parser_oacc_declare (c_parser *parser)
 #define OACC_ENTER_DATA_CLAUSE_MASK					\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@@ -15766,6 +15797,7 @@  c_parser_oacc_declare (c_parser *parser)
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) 		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH) 		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FINALIZE) 		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
@@ -15900,6 +15932,7 @@  c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
 
 #define OACC_KERNELS_CLAUSE_MASK					\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
@@ -15915,6 +15948,7 @@  c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
 
 #define OACC_PARALLEL_CLAUSE_MASK					\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index c7339509bd1..0b9b207e653 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12884,7 +12884,6 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	  return error_mark_node;
 	}
       if (TREE_CODE (t) == COMPONENT_REF
-	  && ort == C_ORT_OMP
 	  && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	      || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
 	      || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM))
@@ -12905,6 +12904,15 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		  return error_mark_node;
 		}
 	      t = TREE_OPERAND (t, 0);
+	      if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
+		{
+		  if (maybe_ne (mem_ref_offset (t), 0))
+	            error_at (OMP_CLAUSE_LOCATION (c),
+			      "cannot dereference %qE in %qs clause", t,
+			      omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		  else
+		    t = TREE_OPERAND (t, 0);
+		}
 	    }
 	}
       if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
@@ -12990,7 +12998,18 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
     length = fold_convert (sizetype, length);
   if (low_bound == NULL_TREE)
     low_bound = integer_zero_node;
-
+  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+	  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+    {
+      if (length != integer_one_node)
+	{
+	  error_at (OMP_CLAUSE_LOCATION (c),
+		    "expected single pointer in %qs clause",
+		    c_omp_map_clause_name (c, ort == C_ORT_ACC));
+	  return error_mark_node;
+	}
+    }
   if (length != NULL_TREE)
     {
       if (!integer_nonzerop (length))
@@ -13430,7 +13449,11 @@  handle_omp_array_sections (tree c, enum c_omp_region_type ort)
       if (ort != C_ORT_OMP && ort != C_ORT_ACC)
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
       else if (TREE_CODE (t) == COMPONENT_REF)
-	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+	{
+	  gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
+					       : GOMP_MAP_ALWAYS_POINTER;
+	  OMP_CLAUSE_SET_MAP_KIND (c2, k);
+	}
       else
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
       if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
@@ -13667,6 +13690,35 @@  c_omp_finish_iterators (tree iter)
   return ret;
 }
 
+/* Ensure that pointers are used in OpenACC attach and detach clauses.
+   Return true if an error has been detected.  */
+
+static bool
+c_oacc_check_attachments (tree c)
+{
+  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+    return false;
+
+  /* OpenACC attach / detach clauses must be pointers.  */
+  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+    {
+      tree t = OMP_CLAUSE_DECL (c);
+
+      while (TREE_CODE (t) == TREE_LIST)
+	t = TREE_CHAIN (t);
+
+      if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
+	{
+	  error_at (OMP_CLAUSE_LOCATION (c), "expected pointer in %qs clause",
+		    c_omp_map_clause_name (c, true));
+	  return true;
+	}
+    }
+
+  return false;
+}
+
 /* For all elements of CLAUSES, validate them against their constraints.
    Remove any elements from the list that are invalid.  */
 
@@ -14420,6 +14472,8 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			}
 		    }
 		}
+	      if (c_oacc_check_attachments (c))
+		remove = true;
 	      break;
 	    }
 	  if (t == error_mark_node)
@@ -14427,8 +14481,13 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      remove = true;
 	      break;
 	    }
+	  /* OpenACC attach / detach clauses must be pointers.  */
+	  if (c_oacc_check_attachments (c))
+	    {
+	      remove = true;
+	      break;
+	    }
 	  if (TREE_CODE (t) == COMPONENT_REF
-	      && (ort & C_ORT_OMP)
 	      && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
 	    {
 	      if (DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
@@ -14463,6 +14522,15 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		      break;
 		    }
 		  t = TREE_OPERAND (t, 0);
+		  if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
+	            {
+		      if (maybe_ne (mem_ref_offset (t), 0))
+			error_at (OMP_CLAUSE_LOCATION (c),
+				  "cannot dereference %qE in %qs clause", t,
+				  omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		      else
+			t = TREE_OPERAND (t, 0);
+		    }
 		}
 	      if (remove)
 		break;
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 018a03d4965..d3e71e38983 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -32480,6 +32480,8 @@  cp_parser_omp_clause_name (cp_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_ALIGNED;
 	  else if (!strcmp ("async", p))
 	    result = PRAGMA_OACC_CLAUSE_ASYNC;
+	  else if (!strcmp ("attach", p))
+	    result = PRAGMA_OACC_CLAUSE_ATTACH;
 	  break;
 	case 'b':
 	  if (!strcmp ("bind", p))
@@ -32504,6 +32506,8 @@  cp_parser_omp_clause_name (cp_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_DEFAULTMAP;
 	  else if (!strcmp ("depend", p))
 	    result = PRAGMA_OMP_CLAUSE_DEPEND;
+	  else if (!strcmp ("detach", p))
+	    result = PRAGMA_OACC_CLAUSE_DETACH;
 	  else if (!strcmp ("device", p))
 	    result = PRAGMA_OMP_CLAUSE_DEVICE;
 	  else if (!strcmp ("deviceptr", p))
@@ -32706,11 +32710,15 @@  check_no_duplicate_clause (tree clauses, enum omp_clause_code code,
 
    COLON can be NULL if only closing parenthesis should end the list,
    or pointer to bool which will receive false if the list is terminated
-   by closing parenthesis or true if the list is terminated by colon.  */
+   by closing parenthesis or true if the list is terminated by colon.
+
+   The optional ALLOW_DEREF argument is true if list items can use the deref
+   (->) operator.  */
 
 static tree
 cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
-				tree list, bool *colon)
+				tree list, bool *colon,
+				bool allow_deref = false)
 {
   cp_token *token;
   bool saved_colon_corrects_to_scope_p = parser->colon_corrects_to_scope_p;
@@ -32791,15 +32799,20 @@  cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 	    case OMP_CLAUSE_MAP:
 	    case OMP_CLAUSE_FROM:
 	    case OMP_CLAUSE_TO:
-	      while (cp_lexer_next_token_is (parser->lexer, CPP_DOT))
+	      while (cp_lexer_next_token_is (parser->lexer, CPP_DOT)
+		     || (allow_deref
+			 && cp_lexer_next_token_is (parser->lexer, CPP_DEREF)))
 		{
+		  cpp_ttype ttype
+		    = cp_lexer_next_token_is (parser->lexer, CPP_DOT)
+		      ? CPP_DOT : CPP_DEREF;
 		  location_t loc
 		    = cp_lexer_peek_token (parser->lexer)->location;
 		  cp_id_kind idk = CP_ID_KIND_NONE;
 		  cp_lexer_consume_token (parser->lexer);
 		  decl = convert_from_reference (decl);
 		  decl
-		    = cp_parser_postfix_dot_deref_expression (parser, CPP_DOT,
+		    = cp_parser_postfix_dot_deref_expression (parser, ttype,
 							      decl, false,
 							      &idk, loc);
 		}
@@ -32917,19 +32930,23 @@  cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
    common case for omp clauses.  */
 
 static tree
-cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list)
+cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list,
+			bool allow_deref = false)
 {
   if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
-    return cp_parser_omp_var_list_no_open (parser, kind, list, NULL);
+    return cp_parser_omp_var_list_no_open (parser, kind, list, NULL,
+					   allow_deref);
   return list;
 }
 
-/* OpenACC 2.0:
+/* OpenACC 2.0+:
+   attach ( variable-list )
    copy ( variable-list )
    copyin ( variable-list )
    copyout ( variable-list )
    create ( variable-list )
    delete ( variable-list )
+   detach ( variable-list )
    present ( variable-list ) */
 
 static tree
@@ -32939,6 +32956,9 @@  cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
   enum gomp_map_kind kind;
   switch (c_kind)
     {
+    case PRAGMA_OACC_CLAUSE_ATTACH:
+      kind = GOMP_MAP_ATTACH;
+      break;
     case PRAGMA_OACC_CLAUSE_COPY:
       kind = GOMP_MAP_TOFROM;
       break;
@@ -32954,6 +32974,9 @@  cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
     case PRAGMA_OACC_CLAUSE_DELETE:
       kind = GOMP_MAP_RELEASE;
       break;
+    case PRAGMA_OACC_CLAUSE_DETACH:
+      kind = GOMP_MAP_DETACH;
+      break;
     case PRAGMA_OACC_CLAUSE_DEVICE:
       kind = GOMP_MAP_FORCE_TO;
       break;
@@ -32973,7 +32996,7 @@  cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
       gcc_unreachable ();
     }
   tree nl, c;
-  nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list);
+  nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true);
 
   for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
     OMP_CLAUSE_SET_MAP_KIND (c, kind);
@@ -35451,6 +35474,10 @@  cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 						  clauses);
 	  c_name = "auto";
 	  break;
+	case PRAGMA_OACC_CLAUSE_ATTACH:
+	  clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "attach";
+	  break;
 	case PRAGMA_OACC_CLAUSE_COLLAPSE:
 	  clauses = cp_parser_omp_clause_collapse (parser, clauses, here);
 	  c_name = "collapse";
@@ -35479,6 +35506,10 @@  cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  clauses = cp_parser_omp_clause_default (parser, clauses, here, true);
 	  c_name = "default";
 	  break;
+	case PRAGMA_OACC_CLAUSE_DETACH:
+	  clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+	  c_name = "detach";
+	  break;
 	case PRAGMA_OACC_CLAUSE_DEVICE:
 	  clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
 	  c_name = "device";
@@ -39313,10 +39344,12 @@  cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
      structured-block  */
 
 #define OACC_DATA_CLAUSE_MASK						\
-	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY)		\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) )
@@ -39516,6 +39549,7 @@  cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 
 #define OACC_ENTER_DATA_CLAUSE_MASK					\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE)		\
@@ -39526,6 +39560,7 @@  cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) 		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FINALIZE) 		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
@@ -39629,6 +39664,7 @@  cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
 
 #define OACC_KERNELS_CLAUSE_MASK					\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
@@ -39644,6 +39680,7 @@  cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
 
 #define OACC_PARALLEL_CLAUSE_MASK					\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 5a8c4d237e0..f48bff42fc2 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -4700,7 +4700,6 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	t = TREE_OPERAND (t, 0);
       ret = t;
       if (TREE_CODE (t) == COMPONENT_REF
-	  && ort == C_ORT_OMP
 	  && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	      || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
 	      || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM)
@@ -4724,6 +4723,8 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		  return error_mark_node;
 		}
 	      t = TREE_OPERAND (t, 0);
+	      if (ort == C_ORT_ACC && TREE_CODE (t) == INDIRECT_REF)
+	        t = TREE_OPERAND (t, 0);
 	    }
 	  if (REFERENCE_REF_P (t))
 	    t = TREE_OPERAND (t, 0);
@@ -4823,6 +4824,18 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
   if (low_bound == NULL_TREE)
     low_bound = integer_zero_node;
 
+  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+	  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+    {
+      if (length != integer_one_node)
+	{
+	  error_at (OMP_CLAUSE_LOCATION (c),
+		    "expected single pointer in %qs clause",
+		    c_omp_map_clause_name (c, ort == C_ORT_ACC));
+	  return error_mark_node;
+	}
+    }
   if (length != NULL_TREE)
     {
       if (!integer_nonzerop (length))
@@ -5259,12 +5272,18 @@  handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	  if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC)
 	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
 	  else if (TREE_CODE (t) == COMPONENT_REF)
-	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+	    {
+	      gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
+						   : GOMP_MAP_ALWAYS_POINTER;
+	      OMP_CLAUSE_SET_MAP_KIND (c2, k);
+	    }
 	  else if (REFERENCE_REF_P (t)
 		   && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
 	    {
 	      t = TREE_OPERAND (t, 0);
-	      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+	      gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
+						   : GOMP_MAP_ALWAYS_POINTER;
+	      OMP_CLAUSE_SET_MAP_KIND (c2, k);
 	    }
 	  else
 	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
@@ -6180,6 +6199,41 @@  cp_omp_finish_iterators (tree iter)
   return ret;
 }
 
+/* Ensure that pointers are used in OpenACC attach and detach clauses.
+   Return true if an error has been detected.  */
+
+static bool
+cp_oacc_check_attachments (tree c)
+{
+  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+    return false;
+
+  /* OpenACC attach / detach clauses must be pointers.  */
+  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+    {
+      tree t = OMP_CLAUSE_DECL (c);
+      tree type;
+
+      while (TREE_CODE (t) == TREE_LIST)
+	t = TREE_CHAIN (t);
+
+      type = TREE_TYPE (t);
+
+      if (TREE_CODE (type) == REFERENCE_TYPE)
+	type = TREE_TYPE (type);
+
+      if (TREE_CODE (type) != POINTER_TYPE)
+	{
+	  error_at (OMP_CLAUSE_LOCATION (c), "expected pointer in %qs clause",
+		    c_omp_map_clause_name (c, true));
+	  return true;
+	}
+    }
+
+  return false;
+}
+
 /* For all elements of CLAUSES, validate them vs OpenMP constraints.
    Remove any elements from the list that are invalid.  */
 
@@ -6444,7 +6498,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    t = OMP_CLAUSE_DECL (c);
 	check_dup_generic_t:
 	  if (t == current_class_ptr
-	      && (ort != C_ORT_OMP_DECLARE_SIMD
+	      && ((ort != C_ORT_OMP_DECLARE_SIMD && ort != C_ORT_ACC)
 		  || (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_LINEAR
 		      && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_UNIFORM)))
 	    {
@@ -6514,8 +6568,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	handle_field_decl:
 	  if (!remove
 	      && TREE_CODE (t) == FIELD_DECL
-	      && t == OMP_CLAUSE_DECL (c)
-	      && ort != C_ORT_ACC)
+	      && t == OMP_CLAUSE_DECL (c))
 	    {
 	      OMP_CLAUSE_DECL (c)
 		= omp_privatize_field (t, (OMP_CLAUSE_CODE (c)
@@ -6582,7 +6635,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
 	  else
 	    t = OMP_CLAUSE_DECL (c);
-	  if (t == current_class_ptr)
+	  if (ort != C_ORT_ACC && t == current_class_ptr)
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
 			"%<this%> allowed in OpenMP only in %<declare simd%>"
@@ -7071,7 +7124,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    }
 	  if (t == error_mark_node)
 	    remove = true;
-	  else if (t == current_class_ptr)
+	  else if (ort != C_ORT_ACC && t == current_class_ptr)
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
 			"%<this%> allowed in OpenMP only in %<declare simd%>"
@@ -7202,6 +7255,8 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			}
 		    }
 		}
+	      if (cp_oacc_check_attachments (c))
+		remove = true;
 	      break;
 	    }
 	  if (t == error_mark_node)
@@ -7209,14 +7264,25 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      remove = true;
 	      break;
 	    }
+	  /* OpenACC attach / detach clauses must be pointers.  */
+	  if (cp_oacc_check_attachments (c))
+	    {
+	      remove = true;
+	      break;
+	    }
 	  if (REFERENCE_REF_P (t)
 	      && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
 	    {
 	      t = TREE_OPERAND (t, 0);
 	      OMP_CLAUSE_DECL (c) = t;
 	    }
+	  if (ort == C_ORT_ACC
+	      && TREE_CODE (t) == COMPONENT_REF
+	      && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
+	    t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
 	  if (TREE_CODE (t) == COMPONENT_REF
-	      && (ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
+	      && ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
+		  || ort == C_ORT_ACC)
 	      && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
 	    {
 	      if (type_dependent_expression_p (t))
@@ -7266,7 +7332,8 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		break;
 	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		  && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
-		      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER))
+		      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER
+		      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH))
 		break;
 	      if (DECL_P (t))
 		error_at (OMP_CLAUSE_LOCATION (c),
@@ -7348,7 +7415,9 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      else
 		bitmap_set_bit (&generic_head, DECL_UID (t));
 	    }
-	  else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+	  else if (bitmap_bit_p (&map_head, DECL_UID (t))
+		   && (ort != C_ORT_ACC
+		       || !bitmap_bit_p (&map_field_head, DECL_UID (t))))
 	    {
 	      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 		error_at (OMP_CLAUSE_LOCATION (c),
@@ -7403,7 +7472,12 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
 					      OMP_CLAUSE_MAP);
 		  if (TREE_CODE (t) == COMPONENT_REF)
-		    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+		    {
+		      gomp_map_kind k
+			= (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
+					     : GOMP_MAP_ALWAYS_POINTER;
+		      OMP_CLAUSE_SET_MAP_KIND (c2, k);
+		    }
 		  else
 		    OMP_CLAUSE_SET_MAP_KIND (c2,
 					     GOMP_MAP_FIRSTPRIVATE_REFERENCE);
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index d84d1fa7f7e..45fdfdd4dd3 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1191,10 +1191,12 @@  enum gfc_omp_depend_op
 enum gfc_omp_map_op
 {
   OMP_MAP_ALLOC,
+  OMP_MAP_ATTACH,
   OMP_MAP_TO,
   OMP_MAP_FROM,
   OMP_MAP_TOFROM,
   OMP_MAP_DELETE,
+  OMP_MAP_DETACH,
   OMP_MAP_FORCE_ALLOC,
   OMP_MAP_FORCE_TO,
   OMP_MAP_FORCE_FROM,
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 7df7384c187..f08f77ce940 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -222,7 +222,8 @@  static match
 gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
 			     bool allow_common, bool *end_colon = NULL,
 			     gfc_omp_namelist ***headp = NULL,
-			     bool allow_sections = false)
+			     bool allow_sections = false,
+			     bool allow_derived = false)
 {
   gfc_omp_namelist *head, *tail, *p;
   locus old_loc, cur_loc;
@@ -248,7 +249,8 @@  gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
 	case MATCH_YES:
 	  gfc_expr *expr;
 	  expr = NULL;
-	  if (allow_sections && gfc_peek_ascii_char () == '(')
+	  if ((allow_sections && gfc_peek_ascii_char () == '(')
+	      || (allow_derived && gfc_peek_ascii_char () == '%'))
 	    {
 	      gfc_current_locus = cur_loc;
 	      m = gfc_match_variable (&expr, 0);
@@ -786,7 +788,7 @@  enum omp_mask1
   OMP_MASK1_LAST
 };
 
-/* OpenACC 2.0 specific clauses. */
+/* OpenACC 2.0+ specific clauses. */
 enum omp_mask2
 {
   OMP_CLAUSE_ASYNC,
@@ -812,6 +814,8 @@  enum omp_mask2
   OMP_CLAUSE_TILE,
   OMP_CLAUSE_IF_PRESENT,
   OMP_CLAUSE_FINALIZE,
+  OMP_CLAUSE_ATTACH,
+  OMP_CLAUSE_DETACH,
   /* This must come last.  */
   OMP_MASK2_LAST
 };
@@ -915,10 +919,12 @@  omp_inv_mask::omp_inv_mask (const omp_mask &m) : omp_mask (m)
    mapping.  */
 
 static bool
-gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op)
+gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op,
+			  bool allow_derived = false)
 {
   gfc_omp_namelist **head = NULL;
-  if (gfc_match_omp_variable_list ("", list, false, NULL, &head, true)
+  if (gfc_match_omp_variable_list ("", list, false, NULL, &head, true,
+				   allow_derived)
       == MATCH_YES)
     {
       gfc_omp_namelist *n;
@@ -940,6 +946,14 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 {
   gfc_omp_clauses *c = gfc_get_omp_clauses ();
   locus old_loc;
+  /* Determine whether we're dealing with an OpenACC directive that permits
+     derived type member accesses.  This in particular disallows
+     "!$acc declare" from using such accesses, because it's not clear if/how
+     that should work.  */
+  bool allow_derived = (openacc
+			&& ((mask & OMP_CLAUSE_ATTACH)
+			    || (mask & OMP_CLAUSE_DETACH)
+			    || (mask & OMP_CLAUSE_HOST_SELF)));
 
   gcc_checking_assert (OMP_MASK1_LAST <= 64 && OMP_MASK2_LAST <= 64);
   *cp = NULL;
@@ -1013,6 +1027,11 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	      needs_space = true;
 	      continue;
 	    }
+	  if ((mask & OMP_CLAUSE_ATTACH)
+	      && gfc_match ("attach ( ") == MATCH_YES
+	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+					   OMP_MAP_ATTACH, allow_derived))
+	    continue;
 	  break;
 	case 'c':
 	  if ((mask & OMP_CLAUSE_COLLAPSE)
@@ -1040,7 +1059,7 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	  if ((mask & OMP_CLAUSE_COPY)
 	      && gfc_match ("copy ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_TOFROM))
+					   OMP_MAP_TOFROM, allow_derived))
 	    continue;
 	  if (mask & OMP_CLAUSE_COPYIN)
 	    {
@@ -1048,7 +1067,7 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 		{
 		  if (gfc_match ("copyin ( ") == MATCH_YES
 		      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-						   OMP_MAP_TO))
+						   OMP_MAP_TO, allow_derived))
 		    continue;
 		}
 	      else if (gfc_match_omp_variable_list ("copyin (",
@@ -1059,7 +1078,7 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	  if ((mask & OMP_CLAUSE_COPYOUT)
 	      && gfc_match ("copyout ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_FROM))
+					   OMP_MAP_FROM, allow_derived))
 	    continue;
 	  if ((mask & OMP_CLAUSE_COPYPRIVATE)
 	      && gfc_match_omp_variable_list ("copyprivate (",
@@ -1069,7 +1088,7 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	  if ((mask & OMP_CLAUSE_CREATE)
 	      && gfc_match ("create ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_ALLOC))
+					   OMP_MAP_ALLOC, allow_derived))
 	    continue;
 	  break;
 	case 'd':
@@ -1105,7 +1124,7 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	  if ((mask & OMP_CLAUSE_DELETE)
 	      && gfc_match ("delete ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_RELEASE))
+					   OMP_MAP_RELEASE, allow_derived))
 	    continue;
 	  if ((mask & OMP_CLAUSE_DEPEND)
 	      && gfc_match ("depend ( ") == MATCH_YES)
@@ -1148,6 +1167,11 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	      else
 		gfc_current_locus = old_loc;
 	    }
+	  if ((mask & OMP_CLAUSE_DETACH)
+	      && gfc_match ("detach ( ") == MATCH_YES
+	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+					   OMP_MAP_DETACH, allow_derived))
+	    continue;
 	  if ((mask & OMP_CLAUSE_DEVICE)
 	      && !openacc
 	      && c->device == NULL
@@ -1157,12 +1181,13 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	      && openacc
 	      && gfc_match ("device ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_FORCE_TO))
+					   OMP_MAP_FORCE_TO, allow_derived))
 	    continue;
 	  if ((mask & OMP_CLAUSE_DEVICEPTR)
 	      && gfc_match ("deviceptr ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_FORCE_DEVICEPTR))
+					   OMP_MAP_FORCE_DEVICEPTR,
+					   allow_derived))
 	    continue;
 	  if ((mask & OMP_CLAUSE_DEVICE_RESIDENT)
 	      && gfc_match_omp_variable_list
@@ -1240,7 +1265,7 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	  if ((mask & OMP_CLAUSE_HOST_SELF)
 	      && gfc_match ("host ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_FORCE_FROM))
+					   OMP_MAP_FORCE_FROM, allow_derived))
 	    continue;
 	  break;
 	case 'i':
@@ -1512,47 +1537,48 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	  if ((mask & OMP_CLAUSE_COPY)
 	      && gfc_match ("pcopy ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_TOFROM))
+					   OMP_MAP_TOFROM, allow_derived))
 	    continue;
 	  if ((mask & OMP_CLAUSE_COPYIN)
 	      && gfc_match ("pcopyin ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_TO))
+					   OMP_MAP_TO, allow_derived))
 	    continue;
 	  if ((mask & OMP_CLAUSE_COPYOUT)
 	      && gfc_match ("pcopyout ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_FROM))
+					   OMP_MAP_FROM, allow_derived))
 	    continue;
 	  if ((mask & OMP_CLAUSE_CREATE)
 	      && gfc_match ("pcreate ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_ALLOC))
+					   OMP_MAP_ALLOC, allow_derived))
 	    continue;
 	  if ((mask & OMP_CLAUSE_PRESENT)
 	      && gfc_match ("present ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_FORCE_PRESENT))
+					   OMP_MAP_FORCE_PRESENT,
+					   allow_derived))
 	    continue;
 	  if ((mask & OMP_CLAUSE_COPY)
 	      && gfc_match ("present_or_copy ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_TOFROM))
+					   OMP_MAP_TOFROM, allow_derived))
 	    continue;
 	  if ((mask & OMP_CLAUSE_COPYIN)
 	      && gfc_match ("present_or_copyin ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_TO))
+					   OMP_MAP_TO, allow_derived))
 	    continue;
 	  if ((mask & OMP_CLAUSE_COPYOUT)
 	      && gfc_match ("present_or_copyout ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_FROM))
+					   OMP_MAP_FROM, allow_derived))
 	    continue;
 	  if ((mask & OMP_CLAUSE_CREATE)
 	      && gfc_match ("present_or_create ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_ALLOC))
+					   OMP_MAP_ALLOC, allow_derived))
 	    continue;
 	  if ((mask & OMP_CLAUSE_PRIORITY)
 	      && c->priority == NULL
@@ -1670,8 +1696,8 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 
 	      if (gfc_match_omp_variable_list (" :",
 					       &c->lists[OMP_LIST_REDUCTION],
-					       false, NULL, &head,
-					       openacc) == MATCH_YES)
+					       false, NULL, &head, openacc,
+					       allow_derived) == MATCH_YES)
 		{
 		  gfc_omp_namelist *n;
 		  if (rop == OMP_REDUCTION_NONE)
@@ -1770,7 +1796,7 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	  if ((mask & OMP_CLAUSE_HOST_SELF)
 	      && gfc_match ("self ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-					   OMP_MAP_FORCE_FROM))
+					   OMP_MAP_FORCE_FROM, allow_derived))
 	    continue;
 	  if ((mask & OMP_CLAUSE_SEQ)
 	      && !c->seq
@@ -1945,17 +1971,17 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
    | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT		      \
    | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR	      \
    | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT	      \
-   | OMP_CLAUSE_WAIT)
+   | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH)
 #define OACC_KERNELS_CLAUSES \
   (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS	      \
    | OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
    | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT		      \
    | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT	      \
-   | OMP_CLAUSE_WAIT)
+   | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH)
 #define OACC_DATA_CLAUSES \
   (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR  | OMP_CLAUSE_COPY	      \
    | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE		      \
-   | OMP_CLAUSE_PRESENT)
+   | OMP_CLAUSE_PRESENT | OMP_CLAUSE_ATTACH)
 #define OACC_LOOP_CLAUSES \
   (omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER	      \
    | OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT	      \
@@ -1976,10 +2002,11 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
    | OMP_CLAUSE_DEVICE | OMP_CLAUSE_WAIT | OMP_CLAUSE_IF_PRESENT)
 #define OACC_ENTER_DATA_CLAUSES \
   (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT	      \
-   | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE)
+   | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE | OMP_CLAUSE_ATTACH)
 #define OACC_EXIT_DATA_CLAUSES \
   (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT	      \
-   | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE | OMP_CLAUSE_FINALIZE)
+   | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE | OMP_CLAUSE_FINALIZE	      \
+   | OMP_CLAUSE_DETACH)
 #define OACC_WAIT_CLAUSES \
   omp_mask (OMP_CLAUSE_ASYNC)
 #define OACC_ROUTINE_CLAUSES \
@@ -3812,9 +3839,6 @@  resolve_nonnegative_int_expr (gfc_expr *expr, const char *clause)
 static void
 check_symbol_not_pointer (gfc_symbol *sym, locus loc, const char *name)
 {
-  if (sym->ts.type == BT_DERIVED && sym->attr.pointer)
-    gfc_error ("POINTER object %qs of derived type in %s clause at %L",
-	       sym->name, name, &loc);
   if (sym->ts.type == BT_DERIVED && sym->attr.cray_pointer)
     gfc_error ("Cray pointer object %qs of derived type in %s clause at %L",
 	       sym->name, name, &loc);
@@ -3859,9 +3883,6 @@  check_array_not_assumed (gfc_symbol *sym, locus loc, const char *name)
 static void
 resolve_oacc_data_clauses (gfc_symbol *sym, locus loc, const char *name)
 {
-  if (sym->ts.type == BT_DERIVED && sym->attr.allocatable)
-    gfc_error ("ALLOCATABLE object %qs of derived type in %s clause at %L",
-	       sym->name, name, &loc);
   if ((sym->ts.type == BT_ASSUMED && sym->attr.allocatable)
       || (sym->ts.type == BT_CLASS && CLASS_DATA (sym)
 	  && CLASS_DATA (sym)->attr.allocatable))
@@ -4244,11 +4265,26 @@  resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	&& (list != OMP_LIST_REDUCTION || !openacc))
       for (n = omp_clauses->lists[list]; n; n = n->next)
 	{
-	  if (n->sym->mark)
-	    gfc_error ("Symbol %qs present on multiple clauses at %L",
-		       n->sym->name, &n->where);
-	  else
-	    n->sym->mark = 1;
+	  bool array_only_p = true;
+	  /* Disallow duplicate bare variable references and multiple
+	     subarrays of the same array here, but allow multiple components of
+	     the same (e.g. derived-type) variable.  For the latter, duplicate
+	     components are detected elsewhere.  */
+	  if (openacc && n->expr && n->expr->expr_type == EXPR_VARIABLE)
+	    for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next)
+	      if (ref->type != REF_ARRAY)
+	        {
+		  array_only_p = false;
+		  break;
+		}
+	  if (array_only_p)
+	    {
+	      if (n->sym->mark)
+		gfc_error ("Symbol %qs present on multiple clauses at %L",
+			   n->sym->name, &n->where);
+	      else
+		n->sym->mark = 1;
+	    }
 	}
 
   gcc_assert (OMP_LIST_LASTPRIVATE == OMP_LIST_FIRSTPRIVATE + 1);
@@ -4439,23 +4475,43 @@  resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 				 "are allowed on ORDERED directive at %L",
 				 &n->where);
 		  }
+		gfc_ref *array_ref = NULL;
+		bool resolved = false;
 		if (n->expr)
 		  {
-		    if (!gfc_resolve_expr (n->expr)
+		    array_ref = n->expr->ref;
+		    resolved = gfc_resolve_expr (n->expr);
+
+		    /* Look through component refs to find last array
+		       reference.  */
+		    if (openacc)
+		      while (resolved
+			     && array_ref
+			     && (array_ref->type == REF_COMPONENT
+				 || (array_ref->type == REF_ARRAY
+				     && array_ref->next
+			             && (array_ref->next->type
+					 == REF_COMPONENT))))
+			array_ref = array_ref->next;
+		  }
+		if (array_ref
+		    || (n->expr
+			&& (!resolved || n->expr->expr_type != EXPR_VARIABLE)))
+		  {
+		    if (!resolved
 			|| n->expr->expr_type != EXPR_VARIABLE
-			|| n->expr->ref == NULL
-			|| n->expr->ref->next
-			|| n->expr->ref->type != REF_ARRAY)
+			|| array_ref->next
+			|| array_ref->type != REF_ARRAY)
 		      gfc_error ("%qs in %s clause at %L is not a proper "
 				 "array section", n->sym->name, name,
 				 &n->where);
-		    else if (n->expr->ref->u.ar.codimen)
+		    else if (array_ref->u.ar.codimen)
 		      gfc_error ("Coarrays not supported in %s clause at %L",
 				 name, &n->where);
 		    else
 		      {
 			int i;
-			gfc_array_ref *ar = &n->expr->ref->u.ar;
+			gfc_array_ref *ar = &array_ref->u.ar;
 			for (i = 0; i < ar->dimen; i++)
 			  if (ar->stride[i])
 			    {
diff --git a/gcc/fortran/trans-expr.c b/gcc/fortran/trans-expr.c
index 61db4e39210..b0d99f52b2e 100644
--- a/gcc/fortran/trans-expr.c
+++ b/gcc/fortran/trans-expr.c
@@ -2422,7 +2422,7 @@  gfc_conv_substring (gfc_se * se, gfc_ref * ref, int kind,
 
 /* Convert a derived type component reference.  */
 
-static void
+void
 gfc_conv_component_ref (gfc_se * se, gfc_ref * ref)
 {
   gfc_component *c;
@@ -2512,7 +2512,7 @@  gfc_conv_component_ref (gfc_se * se, gfc_ref * ref)
 
 /* This function deals with component references to components of the
    parent type for derived type extensions.  */
-static void
+void
 conv_parent_component_references (gfc_se * se, gfc_ref * ref)
 {
   gfc_component *c;
@@ -2578,6 +2578,95 @@  conv_inquiry (gfc_se * se, gfc_ref * ref, gfc_expr *expr, gfc_typespec *ts)
   se->expr = res;
 }
 
+/* Transparently dereference VAR if it is a pointer, reference, etc.
+   according to Fortran semantics.  */
+
+tree
+gfc_auto_dereference_var (location_t loc, gfc_symbol *sym, tree var,
+			  bool descriptor_only_p, bool is_classarray)
+{
+  /* Characters are entirely different from other types, they are treated
+     separately.  */
+  if (sym->ts.type == BT_CHARACTER)
+    {
+      /* Dereference character pointer dummy arguments
+	 or results.  */
+      if ((sym->attr.pointer || sym->attr.allocatable)
+	  && (sym->attr.dummy
+	      || sym->attr.function
+	      || sym->attr.result))
+	var = build_fold_indirect_ref_loc (input_location, var);
+    }
+  else if (!sym->attr.value)
+    {
+      /* Dereference temporaries for class array dummy arguments.  */
+      if (sym->attr.dummy && is_classarray
+	  && GFC_ARRAY_TYPE_P (TREE_TYPE (var)))
+	{
+	  if (!descriptor_only_p)
+	    var = GFC_DECL_SAVED_DESCRIPTOR (var);
+
+	  var = build_fold_indirect_ref_loc (input_location, var);
+	}
+
+      /* Dereference non-character scalar dummy arguments.  */
+      if (sym->attr.dummy && !sym->attr.dimension
+	  && !(sym->attr.codimension && sym->attr.allocatable)
+	  && (sym->ts.type != BT_CLASS
+	      || (!CLASS_DATA (sym)->attr.dimension
+		  && !(CLASS_DATA (sym)->attr.codimension
+		       && CLASS_DATA (sym)->attr.allocatable))))
+	var = build_fold_indirect_ref_loc (input_location, var);
+
+      /* Dereference scalar hidden result.  */
+      if (flag_f2c && sym->ts.type == BT_COMPLEX
+	  && (sym->attr.function || sym->attr.result)
+	  && !sym->attr.dimension && !sym->attr.pointer
+	  && !sym->attr.always_explicit)
+	var = build_fold_indirect_ref_loc (input_location, var);
+
+      /* Dereference non-character, non-class pointer variables.
+	 These must be dummies, results, or scalars.  */
+      if (!is_classarray
+	  && (sym->attr.pointer || sym->attr.allocatable
+	      || gfc_is_associate_pointer (sym)
+	      || (sym->as && sym->as->type == AS_ASSUMED_RANK))
+	  && (sym->attr.dummy
+	      || sym->attr.function
+	      || sym->attr.result
+	      || (!sym->attr.dimension
+		  && (!sym->attr.codimension || !sym->attr.allocatable))))
+	var = build_fold_indirect_ref_loc (input_location, var);
+      /* Now treat the class array pointer variables accordingly.  */
+      else if (sym->ts.type == BT_CLASS
+	       && sym->attr.dummy
+	       && (CLASS_DATA (sym)->attr.dimension
+		   || CLASS_DATA (sym)->attr.codimension)
+	       && ((CLASS_DATA (sym)->as
+		    && CLASS_DATA (sym)->as->type == AS_ASSUMED_RANK)
+		   || CLASS_DATA (sym)->attr.allocatable
+		   || CLASS_DATA (sym)->attr.class_pointer))
+	var = build_fold_indirect_ref_loc (input_location, var);
+      /* And the case where a non-dummy, non-result, non-function,
+	 non-allotable and non-pointer classarray is present.  This case was
+	 previously covered by the first if, but with introducing the
+	 condition !is_classarray there, that case has to be covered
+	 explicitly.  */
+      else if (sym->ts.type == BT_CLASS
+	       && !sym->attr.dummy
+	       && !sym->attr.function
+	       && !sym->attr.result
+	       && (CLASS_DATA (sym)->attr.dimension
+		   || CLASS_DATA (sym)->attr.codimension)
+	       && (sym->assoc
+		   || !CLASS_DATA (sym)->attr.allocatable)
+	       && !CLASS_DATA (sym)->attr.class_pointer)
+	var = build_fold_indirect_ref_loc (input_location, var);
+    }
+
+  return var;
+}
+
 /* Return the contents of a variable. Also handles reference/pointer
    variables (all Fortran pointer references are implicit).  */
 
@@ -2684,94 +2773,9 @@  gfc_conv_variable (gfc_se * se, gfc_expr * expr)
 	  return;
 	}
 
-
-      /* Dereference the expression, where needed. Since characters
-	 are entirely different from other types, they are treated
-	 separately.  */
-      if (sym->ts.type == BT_CHARACTER)
-	{
-	  /* Dereference character pointer dummy arguments
-	     or results.  */
-	  if ((sym->attr.pointer || sym->attr.allocatable)
-	      && (sym->attr.dummy
-		  || sym->attr.function
-		  || sym->attr.result))
-	    se->expr = build_fold_indirect_ref_loc (input_location,
-						se->expr);
-
-	}
-      else if (!sym->attr.value)
-	{
-	  /* Dereference temporaries for class array dummy arguments.  */
-	  if (sym->attr.dummy && is_classarray
-	      && GFC_ARRAY_TYPE_P (TREE_TYPE (se->expr)))
-	    {
-	      if (!se->descriptor_only)
-		se->expr = GFC_DECL_SAVED_DESCRIPTOR (se->expr);
-
-	      se->expr = build_fold_indirect_ref_loc (input_location,
-						      se->expr);
-	    }
-
-	  /* Dereference non-character scalar dummy arguments.  */
-	  if (sym->attr.dummy && !sym->attr.dimension
-	      && !(sym->attr.codimension && sym->attr.allocatable)
-	      && (sym->ts.type != BT_CLASS
-		  || (!CLASS_DATA (sym)->attr.dimension
-		      && !(CLASS_DATA (sym)->attr.codimension
-			   && CLASS_DATA (sym)->attr.allocatable))))
-	    se->expr = build_fold_indirect_ref_loc (input_location,
-						se->expr);
-
-          /* Dereference scalar hidden result.  */
-	  if (flag_f2c && sym->ts.type == BT_COMPLEX
-	      && (sym->attr.function || sym->attr.result)
-	      && !sym->attr.dimension && !sym->attr.pointer
-	      && !sym->attr.always_explicit)
-	    se->expr = build_fold_indirect_ref_loc (input_location,
-						se->expr);
-
-	  /* Dereference non-character, non-class pointer variables.
-	     These must be dummies, results, or scalars.  */
-	  if (!is_classarray
-	      && (sym->attr.pointer || sym->attr.allocatable
-		  || gfc_is_associate_pointer (sym)
-		  || (sym->as && sym->as->type == AS_ASSUMED_RANK))
-	      && (sym->attr.dummy
-		  || sym->attr.function
-		  || sym->attr.result
-		  || (!sym->attr.dimension
-		      && (!sym->attr.codimension || !sym->attr.allocatable))))
-	    se->expr = build_fold_indirect_ref_loc (input_location,
-						se->expr);
-	  /* Now treat the class array pointer variables accordingly.  */
-	  else if (sym->ts.type == BT_CLASS
-		   && sym->attr.dummy
-		   && (CLASS_DATA (sym)->attr.dimension
-		       || CLASS_DATA (sym)->attr.codimension)
-		   && ((CLASS_DATA (sym)->as
-			&& CLASS_DATA (sym)->as->type == AS_ASSUMED_RANK)
-		       || CLASS_DATA (sym)->attr.allocatable
-		       || CLASS_DATA (sym)->attr.class_pointer))
-	    se->expr = build_fold_indirect_ref_loc (input_location,
-						se->expr);
-	  /* And the case where a non-dummy, non-result, non-function,
-	     non-allotable and non-pointer classarray is present.  This case was
-	     previously covered by the first if, but with introducing the
-	     condition !is_classarray there, that case has to be covered
-	     explicitly.  */
-	  else if (sym->ts.type == BT_CLASS
-		   && !sym->attr.dummy
-		   && !sym->attr.function
-		   && !sym->attr.result
-		   && (CLASS_DATA (sym)->attr.dimension
-		       || CLASS_DATA (sym)->attr.codimension)
-		   && (sym->assoc
-		       || !CLASS_DATA (sym)->attr.allocatable)
-		   && !CLASS_DATA (sym)->attr.class_pointer)
-	    se->expr = build_fold_indirect_ref_loc (input_location,
-						se->expr);
-	}
+      /* Dereference the expression, where needed.  */
+      se->expr = gfc_auto_dereference_var (input_location, sym, se->expr,
+					   se->descriptor_only, is_classarray);
 
       ref = expr->ref;
     }
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index f83bab4850e..775548fe9af 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -75,6 +75,9 @@  gfc_omp_privatize_by_reference (const_tree decl)
 
   if (TREE_CODE (type) == POINTER_TYPE)
     {
+      while (TREE_CODE (decl) == COMPONENT_REF)
+        decl = TREE_OPERAND (decl, 1);
+
       /* Array POINTER/ALLOCATABLE have aggregate types, all user variables
 	 that have POINTER_TYPE type and aren't scalar pointers, scalar
 	 allocatables, Cray pointees or C pointers are supposed to be
@@ -1836,6 +1839,92 @@  gfc_convert_expr_to_tree (stmtblock_t *block, gfc_expr *expr)
 
 static vec<tree, va_heap, vl_embed> *doacross_steps;
 
+
+/* Translate an array section or array element.  */
+
+static void
+gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
+			     tree decl, bool element, gomp_map_kind ptr_kind,
+			     tree node, tree &node2, tree &node3, tree &node4)
+{
+  gfc_se se;
+  tree ptr, ptr2;
+
+  gfc_init_se (&se, NULL);
+
+  if (element)
+    {
+      gfc_conv_expr_reference (&se, n->expr);
+      gfc_add_block_to_block (block, &se.pre);
+      ptr = se.expr;
+      OMP_CLAUSE_SIZE (node)
+	= TYPE_SIZE_UNIT (TREE_TYPE (ptr));
+    }
+  else
+    {
+      gfc_conv_expr_descriptor (&se, n->expr);
+      ptr = gfc_conv_array_data (se.expr);
+      tree type = TREE_TYPE (se.expr);
+      gfc_add_block_to_block (block, &se.pre);
+      OMP_CLAUSE_SIZE (node) = gfc_full_array_size (block, se.expr,
+						    GFC_TYPE_ARRAY_RANK (type));
+      tree elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+      elemsz = fold_convert (gfc_array_index_type, elemsz);
+      OMP_CLAUSE_SIZE (node) = fold_build2 (MULT_EXPR, gfc_array_index_type,
+					    OMP_CLAUSE_SIZE (node), elemsz);
+    }
+  gfc_add_block_to_block (block, &se.post);
+  ptr = fold_convert (build_pointer_type (char_type_node), ptr);
+  OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
+
+  if (POINTER_TYPE_P (TREE_TYPE (decl))
+      && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl)))
+      && ptr_kind == GOMP_MAP_POINTER)
+    {
+      node4 = build_omp_clause (input_location,
+				OMP_CLAUSE_MAP);
+      OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER);
+      OMP_CLAUSE_DECL (node4) = decl;
+      OMP_CLAUSE_SIZE (node4) = size_int (0);
+      decl = build_fold_indirect_ref (decl);
+    }
+  ptr = fold_convert (sizetype, ptr);
+  if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+    {
+      tree type = TREE_TYPE (decl);
+      ptr2 = gfc_conv_descriptor_data_get (decl);
+      node2 = build_omp_clause (input_location,
+				OMP_CLAUSE_MAP);
+      OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
+      OMP_CLAUSE_DECL (node2) = decl;
+      OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+      node3 = build_omp_clause (input_location,
+				OMP_CLAUSE_MAP);
+      OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
+      OMP_CLAUSE_DECL (node3)
+	= gfc_conv_descriptor_data_get (decl);
+      if (ptr_kind == GOMP_MAP_ATTACH_DETACH)
+        STRIP_NOPS (OMP_CLAUSE_DECL (node3));
+    }
+  else
+    {
+      if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+	ptr2 = build_fold_addr_expr (decl);
+      else
+	{
+	  gcc_assert (POINTER_TYPE_P (TREE_TYPE (decl)));
+	  ptr2 = decl;
+	}
+      node3 = build_omp_clause (input_location,
+				OMP_CLAUSE_MAP);
+      OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
+      OMP_CLAUSE_DECL (node3) = decl;
+    }
+  ptr2 = fold_convert (sizetype, ptr2);
+  OMP_CLAUSE_SIZE (node3)
+    = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
+}
+
 static tree
 gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		       locus where, bool declare_simd = false)
@@ -2161,7 +2250,8 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			  || GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
 			  || GFC_DECL_CRAY_POINTEE (decl)
 			  || GFC_DESCRIPTOR_TYPE_P
-					(TREE_TYPE (TREE_TYPE (decl)))))
+					(TREE_TYPE (TREE_TYPE (decl)))
+			  || n->sym->ts.type == BT_DERIVED))
 		    {
 		      tree orig_decl = decl;
 		      node4 = build_omp_clause (input_location,
@@ -2182,10 +2272,13 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			  decl = build_fold_indirect_ref (decl);
 			}
 		    }
-		  if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+		  if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))
+		      && n->u.map_op != OMP_MAP_ATTACH
+		      && n->u.map_op != OMP_MAP_DETACH)
 		    {
 		      tree type = TREE_TYPE (decl);
 		      tree ptr = gfc_conv_descriptor_data_get (decl);
+
 		      ptr = fold_convert (build_pointer_type (char_type_node),
 					  ptr);
 		      ptr = build_fold_indirect_ref (ptr);
@@ -2251,88 +2344,154 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		  else
 		    OMP_CLAUSE_DECL (node) = decl;
 		}
-	      else
+	      else if (n->expr
+		       && n->expr->expr_type == EXPR_VARIABLE
+		       && n->expr->ref->type == REF_COMPONENT)
 		{
-		  tree ptr, ptr2;
+		  gfc_ref *lastcomp;
+
+		  for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next)
+		    if (ref->type == REF_COMPONENT)
+		      lastcomp = ref;
+
+		  symbol_attribute sym_attr;
+
+		  sym_attr = lastcomp->u.c.component->attr;
+
 		  gfc_init_se (&se, NULL);
-		  if (n->expr->ref->u.ar.type == AR_ELEMENT)
-		    {
-		      gfc_conv_expr_reference (&se, n->expr);
-		      gfc_add_block_to_block (block, &se.pre);
-		      ptr = se.expr;
-		      OMP_CLAUSE_SIZE (node)
-			= TYPE_SIZE_UNIT (TREE_TYPE (ptr));
-		    }
-		  else
+
+		  if (!sym_attr.dimension
+		      && lastcomp->u.c.component->ts.type != BT_DERIVED)
 		    {
-		      gfc_conv_expr_descriptor (&se, n->expr);
-		      ptr = gfc_conv_array_data (se.expr);
-		      tree type = TREE_TYPE (se.expr);
+		      /* Last component is a scalar.  */
+		      gfc_conv_expr (&se, n->expr);
 		      gfc_add_block_to_block (block, &se.pre);
-		      OMP_CLAUSE_SIZE (node)
-			= gfc_full_array_size (block, se.expr,
-					       GFC_TYPE_ARRAY_RANK (type));
-		      tree elemsz
-			= TYPE_SIZE_UNIT (gfc_get_element_type (type));
-		      elemsz = fold_convert (gfc_array_index_type, elemsz);
-		      OMP_CLAUSE_SIZE (node)
-			= fold_build2 (MULT_EXPR, gfc_array_index_type,
-				       OMP_CLAUSE_SIZE (node), elemsz);
+		      OMP_CLAUSE_DECL (node) = se.expr;
+		      gfc_add_block_to_block (block, &se.post);
+		      goto finalize_map_clause;
 		    }
-		  gfc_add_block_to_block (block, &se.post);
-		  ptr = fold_convert (build_pointer_type (char_type_node),
-				      ptr);
-		  OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
 
-		  if (POINTER_TYPE_P (TREE_TYPE (decl))
-		      && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl))))
+		  se.expr
+		    = gfc_auto_dereference_var (input_location, n->sym,
+						decl);
+
+		  for (gfc_ref *ref = n->expr->ref;
+		       ref && ref != lastcomp->next;
+		       ref = ref->next)
 		    {
-		      node4 = build_omp_clause (input_location,
-						OMP_CLAUSE_MAP);
-		      OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER);
-		      OMP_CLAUSE_DECL (node4) = decl;
-		      OMP_CLAUSE_SIZE (node4) = size_int (0);
-		      decl = build_fold_indirect_ref (decl);
+		      if (ref->type == REF_COMPONENT)
+			{
+			  if (ref->u.c.sym->attr.extension)
+		            conv_parent_component_references (&se, ref);
+
+			  gfc_conv_component_ref (&se, ref);
+			}
+		      else
+			sorry ("unhandled derived-type component");
 		    }
-		  ptr = fold_convert (sizetype, ptr);
-		  if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+
+		  tree inner = se.expr;
+
+		  /* Last component is a derived type.  */
+		  if (lastcomp->u.c.component->ts.type == BT_DERIVED)
 		    {
-		      tree type = TREE_TYPE (decl);
-		      ptr2 = gfc_conv_descriptor_data_get (decl);
-		      node2 = build_omp_clause (input_location,
-						OMP_CLAUSE_MAP);
-		      OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
-		      OMP_CLAUSE_DECL (node2) = decl;
-		      OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
-		      node3 = build_omp_clause (input_location,
-						OMP_CLAUSE_MAP);
-		      OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
-		      OMP_CLAUSE_DECL (node3)
-			= gfc_conv_descriptor_data_get (decl);
+		      if (sym_attr.allocatable || sym_attr.pointer)
+		        {
+			  tree data = inner;
+			  tree size = TYPE_SIZE_UNIT (TREE_TYPE (inner));
+
+			  OMP_CLAUSE_DECL (node)
+			    = build_fold_indirect_ref (data);
+			  OMP_CLAUSE_SIZE (node) = size;
+			  node2 = build_omp_clause (input_location,
+						    OMP_CLAUSE_MAP);
+			  OMP_CLAUSE_SET_MAP_KIND (node2,
+						   GOMP_MAP_ATTACH_DETACH);
+			  OMP_CLAUSE_DECL (node2) = data;
+			  OMP_CLAUSE_SIZE (node2) = size_int (0);
+			}
+		      else
+		        {
+		          OMP_CLAUSE_DECL (node) = decl;
+			  OMP_CLAUSE_SIZE (node)
+			    = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+			}
 		    }
-		  else
+		  else if (lastcomp->next
+			   && lastcomp->next->type == REF_ARRAY
+			   && lastcomp->next->u.ar.type == AR_FULL)
 		    {
-		      if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
-			ptr2 = build_fold_addr_expr (decl);
-		      else
+		      /* Just pass the (auto-dereferenced) decl through for
+			 bare attach and detach clauses.  */
+		      if (n->u.map_op == OMP_MAP_ATTACH
+			  || n->u.map_op == OMP_MAP_DETACH)
 			{
-			  gcc_assert (POINTER_TYPE_P (TREE_TYPE (decl)));
-			  ptr2 = decl;
+			  OMP_CLAUSE_DECL (node) = inner;
+			  OMP_CLAUSE_SIZE (node) = size_zero_node;
+			  goto finalize_map_clause;
 			}
-		      node3 = build_omp_clause (input_location,
-						OMP_CLAUSE_MAP);
-		      OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
-		      OMP_CLAUSE_DECL (node3) = decl;
+
+		      if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
+		        {
+			  tree type = TREE_TYPE (inner);
+			  tree ptr = gfc_conv_descriptor_data_get (inner);
+			  ptr = build_fold_indirect_ref (ptr);
+			  OMP_CLAUSE_DECL (node) = ptr;
+			  node2 = build_omp_clause (input_location,
+						    OMP_CLAUSE_MAP);
+			  OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
+			  OMP_CLAUSE_DECL (node2) = inner;
+			  OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+			  node3 = build_omp_clause (input_location,
+						    OMP_CLAUSE_MAP);
+			  OMP_CLAUSE_SET_MAP_KIND (node3,
+						   GOMP_MAP_ATTACH_DETACH);
+			  OMP_CLAUSE_DECL (node3)
+			    = gfc_conv_descriptor_data_get (inner);
+		          STRIP_NOPS (OMP_CLAUSE_DECL (node3));
+			  OMP_CLAUSE_SIZE (node3) = size_int (0);
+			  int rank = GFC_TYPE_ARRAY_RANK (type);
+			  OMP_CLAUSE_SIZE (node)
+			    = gfc_full_array_size (block, inner, rank);
+			  tree elemsz
+			    = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+			  elemsz = fold_convert (gfc_array_index_type, elemsz);
+			  OMP_CLAUSE_SIZE (node)
+			    = fold_build2 (MULT_EXPR, gfc_array_index_type,
+					   OMP_CLAUSE_SIZE (node), elemsz);
+			}
+		      else
+		        OMP_CLAUSE_DECL (node) = inner;
+		    }
+		  else  /* An array element or section.  */
+		    {
+		      bool element
+			= (lastcomp->next
+			   && lastcomp->next->type == REF_ARRAY
+			   && lastcomp->next->u.ar.type == AR_ELEMENT);
+
+		      gfc_trans_omp_array_section (block, n, inner, element,
+						   GOMP_MAP_ATTACH_DETACH,
+						   node, node2, node3, node4);
 		    }
-		  ptr2 = fold_convert (sizetype, ptr2);
-		  OMP_CLAUSE_SIZE (node3)
-		    = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
 		}
+	      else  /* An array element or array section.  */
+	        {
+		  bool element = n->expr->ref->u.ar.type == AR_ELEMENT;
+		  gfc_trans_omp_array_section (block, n, decl, element,
+					       GOMP_MAP_POINTER, node, node2,
+					       node3, node4);
+		}
+
+	      finalize_map_clause:
 	      switch (n->u.map_op)
 		{
 		case OMP_MAP_ALLOC:
 		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
 		  break;
+		case OMP_MAP_ATTACH:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ATTACH);
+		  break;
 		case OMP_MAP_TO:
 		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO);
 		  break;
@@ -2357,6 +2516,9 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		case OMP_MAP_DELETE:
 		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DELETE);
 		  break;
+		case OMP_MAP_DETACH:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DETACH);
+		  break;
 		case OMP_MAP_FORCE_ALLOC:
 		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_ALLOC);
 		  break;
diff --git a/gcc/fortran/trans.h b/gcc/fortran/trans.h
index 405e88dd1c4..b9999d0d2ce 100644
--- a/gcc/fortran/trans.h
+++ b/gcc/fortran/trans.h
@@ -565,6 +565,15 @@  tree gfc_conv_expr_present (gfc_symbol *);
 /* Convert a missing, dummy argument into a null or zero.  */
 void gfc_conv_missing_dummy (gfc_se *, gfc_expr *, gfc_typespec, int);
 
+/* Lowering of component references.  */
+void gfc_conv_component_ref (gfc_se * se, gfc_ref * ref);
+void conv_parent_component_references (gfc_se * se, gfc_ref * ref);
+
+/* Automatically dereference var.  */
+tree gfc_auto_dereference_var (location_t, gfc_symbol *, tree,
+			       bool desc_only = false,
+			       bool is_classarray = false);
+
 /* Generate code to allocate a string temporary.  */
 tree gfc_conv_string_tmp (gfc_se *, tree, tree);
 /* Get the string length variable belonging to an expression.  */
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 5d50713a9f1..421c6cb7894 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -123,6 +123,10 @@  enum gimplify_omp_var_data
   /* Flag for GOVD_REDUCTION: inscan seen in {in,ex}clusive clause.  */
   GOVD_REDUCTION_INSCAN = 0x2000000,
 
+  /* Flag for GOVD_MAP: (struct) vars that have pointer attachments for
+     fields.  */
+  GOVD_MAP_HAS_ATTACHMENTS = 8388608,
+
   GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
 			   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
 			   | GOVD_LOCAL)
@@ -8151,7 +8155,13 @@  insert_struct_comp_map (enum tree_code code, tree c, tree struct_node,
   OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
   OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (c));
   OMP_CLAUSE_CHAIN (c2) = scp ? *scp : prev_node;
-  OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node);
+  if (OMP_CLAUSE_CHAIN (prev_node) != c
+      && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP
+      && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))
+	  == GOMP_MAP_TO_PSET))
+    OMP_CLAUSE_SIZE (c2) = OMP_CLAUSE_SIZE (OMP_CLAUSE_CHAIN (prev_node));
+  else
+    OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node);
   if (struct_node)
     OMP_CLAUSE_CHAIN (struct_node) = c2;
 
@@ -8160,8 +8170,10 @@  insert_struct_comp_map (enum tree_code code, tree c, tree struct_node,
      GOMP_MAP_ALWAYS_POINTER though, i.e. a GOMP_MAP_TO_PSET.  */
   if (OMP_CLAUSE_CHAIN (prev_node) != c
       && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP
-      && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))
-	  == GOMP_MAP_ALWAYS_POINTER))
+      && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))
+	   == GOMP_MAP_ALWAYS_POINTER)
+	  || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))
+	      == GOMP_MAP_ATTACH_DETACH)))
     {
       tree c4 = OMP_CLAUSE_CHAIN (prev_node);
       tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
@@ -8285,6 +8297,7 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
   struct gimplify_omp_ctx *ctx, *outer_ctx;
   tree c;
   hash_map<tree, tree> *struct_map_to_clause = NULL;
+  hash_set<tree> *struct_deref_set = NULL;
   tree *prev_list_p = NULL, *orig_list_p = list_p;
   int handled_depend_iterators = -1;
   int nowait = -1;
@@ -8686,8 +8699,6 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    case OMP_TARGET_DATA:
 	    case OMP_TARGET_ENTER_DATA:
 	    case OMP_TARGET_EXIT_DATA:
-	    case OACC_ENTER_DATA:
-	    case OACC_EXIT_DATA:
 	    case OACC_HOST_DATA:
 	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		  || (OMP_CLAUSE_MAP_KIND (c)
@@ -8696,6 +8707,15 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		   mapped, but not the pointer to it.  */
 		remove = true;
 	      break;
+	    case OACC_ENTER_DATA:
+	    case OACC_EXIT_DATA:
+	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET
+		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+		  || (OMP_CLAUSE_MAP_KIND (c)
+		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+		remove = true;
+	      break;
 	    default:
 	      break;
 	    }
@@ -8758,7 +8778,35 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  pd = &TREE_OPERAND (decl, 0);
 		  decl = TREE_OPERAND (decl, 0);
 		}
-	      if (TREE_CODE (decl) == COMPONENT_REF)
+	      bool indir_p = false;
+	      tree orig_decl = decl;
+	      tree decl_ref = NULL_TREE;
+	      if ((region_type & ORT_ACC) != 0
+		  && TREE_CODE (*pd) == COMPONENT_REF
+		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
+		  && code != OACC_UPDATE)
+		{
+		  while (TREE_CODE (decl) == COMPONENT_REF)
+		    {
+		      decl = TREE_OPERAND (decl, 0);
+		      if ((TREE_CODE (decl) == MEM_REF
+			   && integer_zerop (TREE_OPERAND (decl, 1)))
+			  || INDIRECT_REF_P (decl))
+			{
+			  indir_p = true;
+			  decl = TREE_OPERAND (decl, 0);
+			}
+		      if (TREE_CODE (decl) == INDIRECT_REF
+			  && DECL_P (TREE_OPERAND (decl, 0))
+			  && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+			      == REFERENCE_TYPE))
+			{
+			  decl_ref = decl;
+			  decl = TREE_OPERAND (decl, 0);
+			}
+		    }
+		}
+	      else if (TREE_CODE (decl) == COMPONENT_REF)
 		{
 		  while (TREE_CODE (decl) == COMPONENT_REF)
 		    decl = TREE_OPERAND (decl, 0);
@@ -8768,13 +8816,76 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			  == REFERENCE_TYPE))
 		    decl = TREE_OPERAND (decl, 0);
 		}
+	      if (decl != orig_decl && DECL_P (decl) && indir_p)
+		{
+		  gomp_map_kind k = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
+							     : GOMP_MAP_ATTACH;
+		  /* We have a dereference of a struct member.  Make this an
+		     attach/detach operation, and ensure the base pointer is
+		     mapped as a FIRSTPRIVATE_POINTER.  */
+		  OMP_CLAUSE_SET_MAP_KIND (c, k);
+		  flags = GOVD_MAP | GOVD_SEEN | GOVD_EXPLICIT;
+		  tree next_clause = OMP_CLAUSE_CHAIN (c);
+		  if (k == GOMP_MAP_ATTACH
+		      && code != OACC_ENTER_DATA
+		      && (!next_clause
+			   || (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP)
+			   || (OMP_CLAUSE_MAP_KIND (next_clause)
+			       != GOMP_MAP_POINTER)
+			   || OMP_CLAUSE_DECL (next_clause) != decl)
+		      && (!struct_deref_set
+			  || !struct_deref_set->contains (decl)))
+		    {
+		      if (!struct_deref_set)
+		        struct_deref_set = new hash_set<tree> ();
+		      /* As well as the attach, we also need a
+			 FIRSTPRIVATE_POINTER clause to properly map the
+			 pointer to the struct base.  */
+		      tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+						  OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC);
+		      OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c2)
+			= 1;
+		      tree charptr_zero
+		        = build_int_cst (build_pointer_type (char_type_node),
+					 0);
+		      OMP_CLAUSE_DECL (c2)
+			= build2 (MEM_REF, char_type_node,
+				  decl_ref ? decl_ref : decl, charptr_zero);
+		      OMP_CLAUSE_SIZE (c2) = size_zero_node;
+		      tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+						  OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_SET_MAP_KIND (c3,
+					       GOMP_MAP_FIRSTPRIVATE_POINTER);
+		      OMP_CLAUSE_DECL (c3) = decl;
+		      OMP_CLAUSE_SIZE (c3) = size_zero_node;
+		      tree mapgrp = *prev_list_p;
+		      *prev_list_p = c2;
+		      OMP_CLAUSE_CHAIN (c3) = mapgrp;
+		      OMP_CLAUSE_CHAIN (c2) = c3;
+
+		      struct_deref_set->add (decl);
+		    }
+		  goto do_add_decl;
+		}
+	      /* An "attach/detach" operation on an update directive should
+	         behave as a GOMP_MAP_ALWAYS_POINTER.  Beware that
+		 unlike attach or detach map kinds, GOMP_MAP_ALWAYS_POINTER
+		 depends on the previous mapping.  */
+	      if (code == OACC_UPDATE
+		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+		OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
 	      if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
 		  == GS_ERROR)
 		{
 		  remove = true;
 		  break;
 		}
-	      if (DECL_P (decl))
+	      if (DECL_P (decl)
+		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
+		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
+		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
+		  && code != OACC_UPDATE)
 		{
 		  if (error_operand_p (decl))
 		    {
@@ -8795,7 +8906,8 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		      break;
 		    }
 
-		  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)
+		  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER
+		      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
 		    {
 		      /* Error recovery.  */
 		      if (prev_list_p == NULL)
@@ -8827,20 +8939,46 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		    = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
 		  bool ptr = (OMP_CLAUSE_MAP_KIND (c)
 			      == GOMP_MAP_ALWAYS_POINTER);
+		  bool attach_detach = (OMP_CLAUSE_MAP_KIND (c)
+					== GOMP_MAP_ATTACH_DETACH);
+		  bool attach = OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+			        || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH;
+		  bool has_attachments = false;
+		  /* For OpenACC, pointers in structs should trigger an
+		     attach action.  */
+		  if (attach_detach && (region_type & ORT_ACC) != 0)
+		    {
+		      /* Turning a GOMP_MAP_ALWAYS_POINTER clause into a
+			 GOMP_MAP_ATTACH clause after we have detected a case
+			 that needs a GOMP_MAP_STRUCT mapping added.  */
+		      gomp_map_kind k
+			= (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
+						   : GOMP_MAP_ATTACH;
+		      OMP_CLAUSE_SET_MAP_KIND (c, k);
+		      has_attachments = true;
+		    }
 		  if (n == NULL || (n->value & GOVD_MAP) == 0)
 		    {
 		      tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c),
 						 OMP_CLAUSE_MAP);
-		      OMP_CLAUSE_SET_MAP_KIND (l, GOMP_MAP_STRUCT);
+		      gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT
+					       : GOMP_MAP_STRUCT;
+
+		      OMP_CLAUSE_SET_MAP_KIND (l, k);
 		      if (!base_eq_orig_base)
 			OMP_CLAUSE_DECL (l) = unshare_expr (orig_base);
 		      else
 			OMP_CLAUSE_DECL (l) = decl;
-		      OMP_CLAUSE_SIZE (l) = size_int (1);
+		      OMP_CLAUSE_SIZE (l)
+			= (!attach
+			   ? size_int (1)
+			   : DECL_P (OMP_CLAUSE_DECL (l))
+			   ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))
+			   : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l))));
 		      if (struct_map_to_clause == NULL)
 			struct_map_to_clause = new hash_map<tree, tree>;
 		      struct_map_to_clause->put (decl, l);
-		      if (ptr)
+		      if (ptr || attach_detach)
 			{
 			  insert_struct_comp_map (code, c, l, *prev_list_p,
 						  NULL);
@@ -8866,23 +9004,31 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			  OMP_CLAUSE_CHAIN (l) = c2;
 			}
 		      flags = GOVD_MAP | GOVD_EXPLICIT;
-		      if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)
+		      if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))
+			  || ptr
+			  || attach_detach)
 			flags |= GOVD_SEEN;
+		      if (has_attachments)
+			flags |= GOVD_MAP_HAS_ATTACHMENTS;
 		      goto do_add_decl;
 		    }
-		  else
+		  else if (struct_map_to_clause)
 		    {
 		      tree *osc = struct_map_to_clause->get (decl);
 		      tree *sc = NULL, *scp = NULL;
-		      if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)
+		      if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))
+			  || ptr
+			  || attach_detach)
 			n->value |= GOVD_SEEN;
 		      sc = &OMP_CLAUSE_CHAIN (*osc);
 		      if (*sc != c
 			  && (OMP_CLAUSE_MAP_KIND (*sc)
-			      == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) 
+			      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
 			sc = &OMP_CLAUSE_CHAIN (*sc);
+		      /* Here "prev_list_p" is the end of the inserted
+			 alloc/release nodes after the struct node, OSC.  */
 		      for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc))
-			if (ptr && sc == prev_list_p)
+			if ((ptr || attach_detach) && sc == prev_list_p)
 			  break;
 			else if (TREE_CODE (OMP_CLAUSE_DECL (*sc))
 				 != COMPONENT_REF
@@ -8931,7 +9077,7 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			      }
 			    if (same_decl_offset_lt)
 			      {
-				if (ptr)
+				if (ptr || attach_detach)
 				  scp = sc;
 				else
 				  break;
@@ -8939,10 +9085,11 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			  }
 		      if (remove)
 			break;
-		      OMP_CLAUSE_SIZE (*osc)
-			= size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
-				      size_one_node);
-		      if (ptr)
+		      if (!attach)
+			OMP_CLAUSE_SIZE (*osc)
+			  = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
+					size_one_node);
+		      if (ptr || attach_detach)
 			{
 			  tree cl = insert_struct_comp_map (code, c, NULL,
 							    *prev_list_p, scp);
@@ -8972,11 +9119,18 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		}
 	      if (!remove
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
+		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
+		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
 		  && OMP_CLAUSE_CHAIN (c)
 		  && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP
-		  && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
-		      == GOMP_MAP_ALWAYS_POINTER))
+		  && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+		       == GOMP_MAP_ALWAYS_POINTER)
+		      || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+			  == GOMP_MAP_ATTACH_DETACH)
+		      || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+		          == GOMP_MAP_TO_PSET)))
 		prev_list_p = list_p;
+
 	      break;
 	    }
 	  flags = GOVD_MAP | GOVD_EXPLICIT;
@@ -9500,6 +9654,8 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
   gimplify_omp_ctxp = ctx;
   if (struct_map_to_clause)
     delete struct_map_to_clause;
+  if (struct_deref_set)
+    delete struct_deref_set;
 }
 
 /* Return true if DECL is a candidate for shared to firstprivate
@@ -9647,6 +9803,8 @@  gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
     return 0;
   if ((flags & GOVD_SEEN) == 0)
     return 0;
+  if ((flags & GOVD_MAP_HAS_ATTACHMENTS) != 0)
+    return 0;
   if (flags & GOVD_DEBUG_PRIVATE)
     {
       gcc_assert ((flags & GOVD_DATA_SHARE_CLASS) == GOVD_SHARED);
@@ -12550,8 +12708,9 @@  gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
 	   && omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
 			       OMP_CLAUSE_FINALIZE))
     {
-      /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"
-	 semantics apply to all mappings of this OpenACC directive.  */
+      /* Use GOMP_MAP_DELETE, GOMP_MAP_FORCE_DETACH, and
+	 GOMP_MAP_FORCE_FROM to denote that "finalize" semantics apply
+	 to all mappings of this OpenACC directive.  */
       bool finalize_marked = false;
       for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
 	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
@@ -12565,10 +12724,19 @@  gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
 	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
 	      finalize_marked = true;
 	      break;
+	    case GOMP_MAP_DETACH:
+	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_DETACH);
+	      finalize_marked = true;
+	      break;
+	    case GOMP_MAP_STRUCT:
+	    case GOMP_MAP_FORCE_PRESENT:
+	      /* Skip over an initial struct or force_present mapping.  */
+	      break;
 	    default:
-	      /* Check consistency: libgomp relies on the very first data
-		 mapping clause being marked, so make sure we did that before
-		 any other mapping clauses.  */
+	      /* Check consistency: libgomp relies on the very first
+		 non-struct, non-force-present data mapping clause being
+		 marked, so make sure we did that before any other mapping
+		 clauses.  */
 	      gcc_assert (finalize_marked);
 	      break;
 	    }
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index ca7dfdb83a1..59c6678e2dd 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -11322,6 +11322,9 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case GOMP_MAP_FORCE_DEVICEPTR:
 	  case GOMP_MAP_DEVICE_RESIDENT:
 	  case GOMP_MAP_LINK:
+	  case GOMP_MAP_ATTACH:
+	  case GOMP_MAP_DETACH:
+	  case GOMP_MAP_FORCE_DETACH:
 	    gcc_assert (is_gimple_omp_oacc (stmt));
 	    break;
 	  default:
diff --git a/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c
new file mode 100644
index 00000000000..d411bcfa8e7
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c
@@ -0,0 +1,84 @@ 
+/* { dg-do compile } */
+
+#include <stdlib.h>
+#include <stdio.h>
+
+typedef struct {
+  int *a;
+  int *b;
+  int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  const int S = 32;
+  mystruct *m = (mystruct *) calloc (S, sizeof (*m));
+  int i, j;
+
+  for (i = 0; i < S; i++)
+    {
+      m[i].a = (int *) malloc (N * sizeof (int));
+      m[i].b = (int *) malloc (N * sizeof (int));
+      m[i].c = (int *) malloc (N * sizeof (int));
+    }
+
+  for (j = 0; j < S; j++)
+    for (i = 0; i < N; i++)
+      {
+	m[j].a[i] = 0;
+	m[j].b[i] = 0;
+	m[j].c[i] = 0;
+      }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j, k;
+      for (k = 0; k < S; k++)
+#pragma acc parallel loop copy(m[k].a[0:N]) /* { dg-error "expected .\\\). before .\\\.. token" } */
+        for (j = 0; j < N; j++)
+          m[k].a[j]++;
+
+      for (k = 0; k < S; k++)
+#pragma acc parallel loop copy(m[k].b[0:N], m[k].c[5:N-10]) /* { dg-error "expected .\\\). before .\\\.. token" } */
+	/* { dg-error ".m. appears more than once in data clauses" "" { target c++ } .-1 } */
+	for (j = 0; j < N; j++)
+	  {
+	    m[k].b[j]++;
+	    if (j > 5 && j < N - 5)
+	      m[k].c[j]++;
+	}
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  for (j = 0; j < S; j++)
+    {
+      for (i = 0; i < N; i++)
+	{
+	  if (m[j].a[i] != 99)
+	    abort ();
+	  if (m[j].b[i] != 99)
+	    abort ();
+	  if (i > 5 && i < N-5)
+	    {
+	      if (m[j].c[i] != 99)
+		abort ();
+	    }
+	  else
+	    {
+	      if (m[j].c[i] != 0)
+		abort ();
+	    }
+	}
+
+      free (m[j].a);
+      free (m[j].b);
+      free (m[j].c);
+    }
+  free (m);
+
+  return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
new file mode 100644
index 00000000000..6c6a81ea73a
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
@@ -0,0 +1,55 @@ 
+/* Test OpenACC's support for manual deep copy, including the attach
+   and detach clauses.  */
+
+/* { dg-do compile { target int32 } } */
+/* { dg-additional-options "-fdump-tree-omplower" } */
+
+void
+t1 ()
+{
+  struct foo {
+    int *a, *b, c, d, *e;
+  } s;
+
+  int *a, *z;
+
+#pragma acc enter data copyin(s)
+  {
+#pragma acc data copy(s.a[0:10]) copy(z[0:10])
+    {
+      s.e = z;
+#pragma acc parallel loop attach(s.e)
+      for (int i = 0; i < 10; i++)
+        s.a[i] = s.e[i];
+
+
+      a = s.e;
+#pragma acc enter data attach(a)
+#pragma acc exit data detach(a)
+    }
+
+#pragma acc enter data copyin(a)
+#pragma acc acc enter data attach(s.e)
+#pragma acc exit data detach(s.e)
+
+#pragma acc data attach(s.e)
+    {
+    }
+#pragma acc exit data delete(a)
+
+#pragma acc exit data detach(a) finalize
+#pragma acc exit data detach(s.a) finalize
+  }
+}
+
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 8.. map.tofrom:s .len: 32" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .bias: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .bias: 8.." 1 "omplower" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-2.c b/gcc/testsuite/c-c++-common/goacc/mdc-2.c
new file mode 100644
index 00000000000..fae86671fc9
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/mdc-2.c
@@ -0,0 +1,62 @@ 
+/* Test OpenACC's support for manual deep copy, including the attach
+   and detach clauses.  */
+
+void
+t1 ()
+{
+  struct foo {
+    int *a, *b, c, d, *e;
+  } s;
+
+  int *a, *z, scalar, **y;
+
+#pragma acc enter data copyin(s) detach(z) /* { dg-error ".detach. is not valid for" } */
+  {
+#pragma acc data copy(s.a[0:10]) copy(z[0:10])
+    {
+      s.e = z;
+#pragma acc parallel loop attach(s.e) detach(s.b) /* { dg-error ".detach. is not valid for" } */
+      for (int i = 0; i < 10; i++)
+        s.a[i] = s.e[i];
+
+      a = s.e;
+#pragma acc enter data attach(a) detach(s.c) /* { dg-error ".detach. is not valid for" } */
+#pragma acc exit data detach(a)
+    }
+
+#pragma acc enter data attach(z[:5]) /* { dg-error "expected single pointer in .attach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc exit data detach(z[:5]) /* { dg-error "expected single pointer in .detach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc enter data attach(z[1:]) /* { dg-error "expected single pointer in .attach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc exit data detach(z[1:]) /* { dg-error "expected single pointer in .detach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc enter data attach(z[:]) /* { dg-error "expected single pointer in .attach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc exit data detach(z[:]) /* { dg-error "expected single pointer in .detach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc enter data attach(z[3]) /* { dg-error "expected pointer in .attach. clause" } */
+#pragma acc exit data detach(z[3]) /* { dg-error "expected pointer in .detach. clause" } */
+
+#pragma acc acc enter data attach(s.e)
+#pragma acc exit data detach(s.e) attach(z) /* { dg-error ".attach. is not valid for" } */
+
+#pragma acc data attach(s.e)
+    {
+    }
+#pragma acc exit data delete(a) attach(s.a) /* { dg-error ".attach. is not valid for" } */
+
+#pragma acc enter data attach(scalar) /* { dg-error "expected pointer in .attach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc exit data detach(scalar) /* { dg-error "expected pointer in .detach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc enter data attach(s) /* { dg-error "expected pointer in .attach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc exit data detach(s) /* { dg-error "expected pointer in .detach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+  }
+
+#pragma acc enter data attach(y[10])
+#pragma acc exit data detach(y[10])
+}
diff --git a/gcc/testsuite/g++.dg/goacc/mdc.C b/gcc/testsuite/g++.dg/goacc/mdc.C
new file mode 100644
index 00000000000..b3abab30423
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/mdc.C
@@ -0,0 +1,68 @@ 
+/* Test OpenACC's support for manual deep copy, including the attach
+   and detach clauses.  */
+
+void
+t1 ()
+{
+  struct foo {
+    int *a, *b, c, d, *e;
+  } s;
+
+  struct foo& rs = s;
+  
+  int *a, *z, scalar, **y;
+  int* const &ra = a;
+  int* const &rz = z;
+  int& rscalar = scalar;
+  int** const &ry = y;
+
+#pragma acc enter data copyin(rs) detach(rz) /* { dg-error ".detach. is not valid for" } */
+  {
+#pragma acc data copy(rs.a[0:10]) copy(rz[0:10])
+    {
+      s.e = z;
+#pragma acc parallel loop attach(rs.e) detach(rs.b) /* { dg-error ".detach. is not valid for" } */
+      for (int i = 0; i < 10; i++)
+        s.a[i] = s.e[i];
+
+      a = s.e;
+#pragma acc enter data attach(ra) detach(rs.c) /* { dg-error ".detach. is not valid for" } */
+#pragma acc exit data detach(ra)
+    }
+
+#pragma acc enter data attach(rz[:5]) /* { dg-error "expected single pointer in .attach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc exit data detach(rz[:5]) /* { dg-error "expected single pointer in .detach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc enter data attach(rz[1:]) /* { dg-error "expected single pointer in .attach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc exit data detach(rz[1:]) /* { dg-error "expected single pointer in .detach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc enter data attach(rz[:]) /* { dg-error "expected single pointer in .attach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc exit data detach(rz[:]) /* { dg-error "expected single pointer in .detach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc enter data attach(rz[3]) /* { dg-error "expected pointer in .attach. clause" } */
+#pragma acc exit data detach(rz[3]) /* { dg-error "expected pointer in .detach. clause" } */
+
+#pragma acc acc enter data attach(rs.e)
+#pragma acc exit data detach(rs.e) attach(rz) /* { dg-error ".attach. is not valid for" } */
+
+#pragma acc data attach(rs.e)
+    {
+    }
+#pragma acc exit data delete(ra) attach(rs.a) /* { dg-error ".attach. is not valid for" } */
+
+#pragma acc enter data attach(rscalar) /* { dg-error "expected pointer in .attach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc exit data detach(rscalar) /* { dg-error "expected pointer in .detach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc enter data attach(rs) /* { dg-error "expected pointer in .attach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+#pragma acc exit data detach(rs) /* { dg-error "expected pointer in .detach. clause" } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
+  }
+
+#pragma acc enter data attach(ry[10])
+#pragma acc exit data detach(ry[10])
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
index b94214e8b63..1a4a6719987 100644
--- a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
@@ -39,9 +39,9 @@  contains
   !$acc end data
 
 
-  !$acc parallel copy (tip) ! { dg-error "POINTER" }
+  !$acc parallel copy (tip)
   !$acc end parallel
-  !$acc parallel copy (tia) ! { dg-error "ALLOCATABLE" }
+  !$acc parallel copy (tia)
   !$acc end parallel
   !$acc parallel deviceptr (i) copy (i) ! { dg-error "multiple clauses" }
   !$acc end parallel
@@ -54,9 +54,9 @@  contains
   !$acc end data
 
 
-  !$acc parallel copyin (tip) ! { dg-error "POINTER" }
+  !$acc parallel copyin (tip)
   !$acc end parallel
-  !$acc parallel copyin (tia) ! { dg-error "ALLOCATABLE" }
+  !$acc parallel copyin (tia)
   !$acc end parallel
   !$acc parallel deviceptr (i) copyin (i) ! { dg-error "multiple clauses" }
   !$acc end parallel
@@ -71,9 +71,9 @@  contains
   !$acc end data
 
 
-  !$acc parallel copyout (tip) ! { dg-error "POINTER" }
+  !$acc parallel copyout (tip)
   !$acc end parallel
-  !$acc parallel copyout (tia) ! { dg-error "ALLOCATABLE" }
+  !$acc parallel copyout (tia)
   !$acc end parallel
   !$acc parallel deviceptr (i) copyout (i) ! { dg-error "multiple clauses" }
   !$acc end parallel
@@ -90,9 +90,9 @@  contains
   !$acc end data
 
 
-  !$acc parallel create (tip) ! { dg-error "POINTER" }
+  !$acc parallel create (tip)
   !$acc end parallel
-  !$acc parallel create (tia) ! { dg-error "ALLOCATABLE" }
+  !$acc parallel create (tia)
   !$acc end parallel
   !$acc parallel deviceptr (i) create (i) ! { dg-error "multiple clauses" }
   !$acc end parallel
@@ -111,9 +111,9 @@  contains
   !$acc end data
 
 
-  !$acc parallel present (tip) ! { dg-error "POINTER" }
+  !$acc parallel present (tip)
   !$acc end parallel
-  !$acc parallel present (tia) ! { dg-error "ALLOCATABLE" }
+  !$acc parallel present (tia)
   !$acc end parallel
   !$acc parallel deviceptr (i) present (i) ! { dg-error "multiple clauses" }
   !$acc end parallel
@@ -144,9 +144,9 @@  contains
   !$acc end parallel
 
 
-  !$acc parallel present_or_copy (tip) ! { dg-error "POINTER" }
+  !$acc parallel present_or_copy (tip)
   !$acc end parallel
-  !$acc parallel present_or_copy (tia) ! { dg-error "ALLOCATABLE" }
+  !$acc parallel present_or_copy (tia)
   !$acc end parallel
   !$acc parallel deviceptr (i) present_or_copy (i) ! { dg-error "multiple clauses" }
   !$acc end parallel
@@ -169,9 +169,9 @@  contains
   !$acc end data
 
 
-  !$acc parallel present_or_copyin (tip) ! { dg-error "POINTER" }
+  !$acc parallel present_or_copyin (tip)
   !$acc end parallel
-  !$acc parallel present_or_copyin (tia) ! { dg-error "ALLOCATABLE" }
+  !$acc parallel present_or_copyin (tia)
   !$acc end parallel
   !$acc parallel deviceptr (i) present_or_copyin (i) ! { dg-error "multiple clauses" }
   !$acc end parallel
@@ -196,9 +196,9 @@  contains
   !$acc end data
 
 
-  !$acc parallel present_or_copyout (tip) ! { dg-error "POINTER" }
+  !$acc parallel present_or_copyout (tip)
   !$acc end parallel
-  !$acc parallel present_or_copyout (tia) ! { dg-error "ALLOCATABLE" }
+  !$acc parallel present_or_copyout (tia)
   !$acc end parallel
   !$acc parallel deviceptr (i) present_or_copyout (i) ! { dg-error "multiple clauses" }
   !$acc end parallel
@@ -225,9 +225,9 @@  contains
   !$acc end data
 
 
-  !$acc parallel present_or_create (tip) ! { dg-error "POINTER" }
+  !$acc parallel present_or_create (tip)
   !$acc end parallel
-  !$acc parallel present_or_create (tia) ! { dg-error "ALLOCATABLE" }
+  !$acc parallel present_or_create (tia)
   !$acc end parallel
   !$acc parallel deviceptr (i) present_or_create (i) ! { dg-error "multiple clauses" }
   !$acc end parallel
@@ -256,4 +256,4 @@  contains
   !$acc end data
 
   end subroutine foo
-end module test
\ No newline at end of file
+end module test
diff --git a/gcc/testsuite/gfortran.dg/goacc/derived-types-2.f90 b/gcc/testsuite/gfortran.dg/goacc/derived-types-2.f90
new file mode 100644
index 00000000000..d01583fac89
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/derived-types-2.f90
@@ -0,0 +1,14 @@ 
+module bar
+  type :: type1
+     real(8), pointer, public :: p(:) => null()
+  end type
+  type :: type2
+     class(type1), pointer :: p => null()
+  end type
+end module
+
+subroutine foo (var)
+   use bar
+   type(type2), intent(inout) :: var
+   !$acc enter data create(var%p%p)
+end subroutine
diff --git a/gcc/testsuite/gfortran.dg/goacc/derived-types.f90 b/gcc/testsuite/gfortran.dg/goacc/derived-types.f90
new file mode 100644
index 00000000000..5fb29816c42
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/derived-types.f90
@@ -0,0 +1,77 @@ 
+! Test ACC UPDATE with derived types.
+
+module dt
+  integer, parameter :: n = 10
+  type inner
+     integer :: d(n)
+  end type inner
+  type dtype
+     integer(8) :: a, b, c(n)
+     type(inner) :: in
+  end type dtype
+end module dt
+
+program derived_acc
+  use dt
+  
+  implicit none
+  type(dtype):: var
+  integer i
+  !$acc declare create(var)
+  !$acc declare pcopy(var%a) ! { dg-error "Syntax error in OpenMP" }
+
+  !$acc update host(var)
+  !$acc update host(var%a)
+  !$acc update device(var)
+  !$acc update device(var%a)
+  !$acc update self(var)
+  !$acc update self(var%a)
+  
+  !$acc enter data copyin(var)
+  !$acc enter data copyin(var%a)
+
+  !$acc exit data copyout(var)
+  !$acc exit data copyout(var%a)
+
+  !$acc data copy(var)
+  !$acc end data
+
+  !$acc data copyout(var%a)
+  !$acc end data
+
+  !$acc parallel loop pcopyout(var)
+  do i = 1, 10
+  end do  
+  !$acc end parallel loop
+
+  !$acc parallel loop copyout(var%a)
+  do i = 1, 10
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel pcopy(var)
+  !$acc end parallel
+
+  !$acc parallel pcopy(var%a)
+  do i = 1, 10
+  end do
+  !$acc end parallel
+  
+  !$acc kernels pcopyin(var)
+  !$acc end kernels
+
+  !$acc kernels pcopy(var%a)
+  do i = 1, 10
+  end do
+  !$acc end kernels
+
+  !$acc kernels loop pcopyin(var)
+  do i = 1, 10
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop pcopy(var%a)
+  do i = 1, 10
+  end do
+  !$acc end kernels loop
+end program derived_acc
diff --git a/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95
index a414df8d439..c2a49796318 100644
--- a/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95
@@ -44,14 +44,14 @@  contains
   !$acc enter data wait (i, 1) 
   !$acc enter data wait (a) ! { dg-error "INTEGER" }
   !$acc enter data wait (b(5:6)) ! { dg-error "INTEGER" }
-  !$acc enter data copyin (tip) ! { dg-error "POINTER" }
-  !$acc enter data copyin (tia) ! { dg-error "ALLOCATABLE" }
-  !$acc enter data create (tip) ! { dg-error "POINTER" }
-  !$acc enter data create (tia) ! { dg-error "ALLOCATABLE" }
-  !$acc enter data present_or_copyin (tip) ! { dg-error "POINTER" }
-  !$acc enter data present_or_copyin (tia) ! { dg-error "ALLOCATABLE" }
-  !$acc enter data present_or_create (tip) ! { dg-error "POINTER" }
-  !$acc enter data present_or_create (tia) ! { dg-error "ALLOCATABLE" }
+  !$acc enter data copyin (tip)
+  !$acc enter data copyin (tia)
+  !$acc enter data create (tip)
+  !$acc enter data create (tia)
+  !$acc enter data present_or_copyin (tip)
+  !$acc enter data present_or_copyin (tia)
+  !$acc enter data present_or_create (tip)
+  !$acc enter data present_or_create (tia)
   !$acc enter data copyin (i) create (i) ! { dg-error "multiple clauses" }
   !$acc enter data copyin (i) present_or_copyin (i) ! { dg-error "multiple clauses" }
   !$acc enter data create (i) present_or_copyin (i) ! { dg-error "multiple clauses" }
@@ -79,10 +79,10 @@  contains
   !$acc exit data wait (i, 1) 
   !$acc exit data wait (a) ! { dg-error "INTEGER" }
   !$acc exit data wait (b(5:6)) ! { dg-error "INTEGER" }
-  !$acc exit data copyout (tip) ! { dg-error "POINTER" }
-  !$acc exit data copyout (tia) ! { dg-error "ALLOCATABLE" }
-  !$acc exit data delete (tip) ! { dg-error "POINTER" }
-  !$acc exit data delete (tia) ! { dg-error "ALLOCATABLE" }
+  !$acc exit data copyout (tip)
+  !$acc exit data copyout (tia)
+  !$acc exit data delete (tip)
+  !$acc exit data delete (tia)
   !$acc exit data copyout (i) delete (i) ! { dg-error "multiple clauses" }
   !$acc exit data finalize
   !$acc exit data finalize copyout (i)
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 53b3f55a3e6..df617d47899 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -849,6 +849,18 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	case GOMP_MAP_LINK:
 	  pp_string (pp, "link");
 	  break;
+	case GOMP_MAP_ATTACH:
+	  pp_string (pp, "attach");
+	  break;
+	case GOMP_MAP_DETACH:
+	  pp_string (pp, "detach");
+	  break;
+	case GOMP_MAP_FORCE_DETACH:
+	  pp_string (pp, "force_detach");
+	  break;
+	case GOMP_MAP_ATTACH_DETACH:
+	  pp_string (pp, "attach_detach");
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
@@ -870,6 +882,12 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	    case GOMP_MAP_TO_PSET:
 	      pp_string (pp, " [pointer set, len: ");
 	      break;
+	    case GOMP_MAP_ATTACH:
+	    case GOMP_MAP_DETACH:
+	    case GOMP_MAP_FORCE_DETACH:
+	    case GOMP_MAP_ATTACH_DETACH:
+	      pp_string (pp, " [bias: ");
+	      break;
 	    default:
 	      pp_string (pp, " [len: ");
 	      break;
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 82e9094c934..6aa7691a119 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -40,8 +40,11 @@ 
 #define GOMP_MAP_FLAG_SPECIAL_0		(1 << 2)
 #define GOMP_MAP_FLAG_SPECIAL_1		(1 << 3)
 #define GOMP_MAP_FLAG_SPECIAL_2		(1 << 4)
+#define GOMP_MAP_FLAG_SPECIAL_4		(1 << 6)
 #define GOMP_MAP_FLAG_SPECIAL		(GOMP_MAP_FLAG_SPECIAL_1 \
 					 | GOMP_MAP_FLAG_SPECIAL_0)
+#define GOMP_MAP_DEEP_COPY		(GOMP_MAP_FLAG_SPECIAL_4 \
+					 | GOMP_MAP_FLAG_SPECIAL_2)
 /* Flag to force a specific behavior (or else, trigger a run-time error).  */
 #define GOMP_MAP_FLAG_FORCE		(1 << 7)
 
@@ -127,12 +130,23 @@  enum gomp_map_kind
     /* Decrement usage count and deallocate if zero.  */
     GOMP_MAP_RELEASE =			(GOMP_MAP_FLAG_SPECIAL_2
 					 | GOMP_MAP_DELETE),
+    /* In OpenACC, attach a pointer to a mapped struct field.  */
+    GOMP_MAP_ATTACH =			(GOMP_MAP_DEEP_COPY | 0),
+    /* In OpenACC, detach a pointer to a mapped struct field.  */
+    GOMP_MAP_DETACH =			(GOMP_MAP_DEEP_COPY | 1),
+    /* In OpenACC, detach a pointer to a mapped struct field.  */
+    GOMP_MAP_FORCE_DETACH =		(GOMP_MAP_DEEP_COPY
+					 | GOMP_MAP_FLAG_FORCE | 1),
 
     /* Internal to GCC, not used in libgomp.  */
     /* Do not map, but pointer assign a pointer instead.  */
     GOMP_MAP_FIRSTPRIVATE_POINTER =	(GOMP_MAP_LAST | 1),
     /* Do not map, but pointer assign a reference instead.  */
-    GOMP_MAP_FIRSTPRIVATE_REFERENCE =	(GOMP_MAP_LAST | 2)
+    GOMP_MAP_FIRSTPRIVATE_REFERENCE =	(GOMP_MAP_LAST | 2),
+    /* An attach or detach operation.  Rewritten to the appropriate type during
+       gimplification, depending on directive (i.e. "enter data" or
+       parallel/kernels region vs. "exit data").  */
+    GOMP_MAP_ATTACH_DETACH =		(GOMP_MAP_LAST | 3)
   };
 
 #define GOMP_MAP_COPY_TO_P(X) \
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 6870ee34224..9953fabc586 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -867,6 +867,8 @@  struct target_var_desc {
   bool copy_from;
   /* True if data always should be copied from device to host at the end.  */
   bool always_copy_from;
+  /* True if variable should be detached at end of region.  */
+  bool do_detach;
   /* Relative offset against key host_start.  */
   uintptr_t offset;
   /* Actual length.  */
@@ -943,6 +945,9 @@  struct splay_tree_key_s {
   uintptr_t refcount_chk;
 #endif
   union {
+    /* For a block with attached pointers, the attachment counters for each.
+       Only used for OpenACC.  */
+    uintptr_t *attach_count;
     /* Pointer to the original mapping of "omp declare target link" object.
        Only used for OpenMP.  */
     splay_tree_key link_key;
@@ -1096,6 +1101,13 @@  extern void gomp_copy_dev2host (struct gomp_device_descr *,
 				struct goacc_asyncqueue *, void *, const void *,
 				size_t);
 extern uintptr_t gomp_map_val (struct target_mem_desc *, void **, size_t);
+extern void gomp_attach_pointer (struct gomp_device_descr *,
+				 struct goacc_asyncqueue *, splay_tree,
+				 splay_tree_key, uintptr_t, size_t,
+				 struct gomp_coalesce_buf *);
+extern void gomp_detach_pointer (struct gomp_device_descr *,
+				 struct goacc_asyncqueue *, splay_tree_key,
+				 uintptr_t, bool, struct gomp_coalesce_buf *);
 
 #ifdef RC_CHECKING
 extern void dump_tgt (const char *, struct target_mem_desc *);
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index c79430f8d8d..63276f7d29b 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -484,6 +484,16 @@  OACC_2.5.1 {
 	acc_register_library;
 } OACC_2.5;
 
+OACC_2.6 {
+  global:
+	acc_attach;
+	acc_attach_async;
+	acc_detach;
+	acc_detach_async;
+	acc_detach_finalize;
+	acc_detach_finalize_async;
+} OACC_2.5.1;
+
 GOACC_2.0 {
   global:
 	GOACC_data_end;
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 25084b71a2d..c88fe9732d0 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -875,9 +875,13 @@  gomp_acc_remove_pointer (struct gomp_device_descr *acc_dev, void **hostaddrs,
 	case GOMP_MAP_POINTER:
 	case GOMP_MAP_DELETE:
 	case GOMP_MAP_RELEASE:
+	case GOMP_MAP_DETACH:
+	case GOMP_MAP_FORCE_DETACH:
 	  cur_node.host_start = (uintptr_t) hostaddrs[i];
 	  cur_node.host_end = cur_node.host_start
-			      + (kind == GOMP_MAP_POINTER
+			      + ((kind == GOMP_MAP_DETACH
+				  || kind == GOMP_MAP_FORCE_DETACH
+				  || kind == GOMP_MAP_POINTER)
 				 ? sizeof (void *) : sizes[i]);
 	  n = splay_tree_lookup (&acc_dev->mem_map, &cur_node);
 
@@ -926,3 +930,79 @@  gomp_acc_remove_pointer (struct gomp_device_descr *acc_dev, void **hostaddrs,
 
   gomp_mutex_unlock (&acc_dev->lock);
 }
+
+void
+acc_attach_async (void **hostaddr, int async)
+{
+  struct goacc_thread *thr = goacc_thread ();
+  struct gomp_device_descr *acc_dev = thr->dev;
+  goacc_aq aq = get_goacc_asyncqueue (async);
+
+  struct splay_tree_key_s cur_node;
+  splay_tree_key n;
+
+  if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+    return;
+
+  cur_node.host_start = (uintptr_t) hostaddr;
+  cur_node.host_end = cur_node.host_start + sizeof (void *);
+  n = splay_tree_lookup (&acc_dev->mem_map, &cur_node);
+
+  if (n == NULL)
+    gomp_fatal ("struct not mapped for acc_attach");
+
+  gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr,
+		       0, NULL);
+}
+
+void
+acc_attach (void **hostaddr)
+{
+  acc_attach_async (hostaddr, acc_async_sync);
+}
+
+static void
+goacc_detach_internal (void **hostaddr, int async, bool finalize)
+{
+  struct goacc_thread *thr = goacc_thread ();
+  struct gomp_device_descr *acc_dev = thr->dev;
+  struct splay_tree_key_s cur_node;
+  splay_tree_key n;
+  struct goacc_asyncqueue *aq = get_goacc_asyncqueue (async);
+
+  if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+    return;
+
+  cur_node.host_start = (uintptr_t) hostaddr;
+  cur_node.host_end = cur_node.host_start + sizeof (void *);
+  n = splay_tree_lookup (&acc_dev->mem_map, &cur_node);
+
+  if (n == NULL)
+    gomp_fatal ("struct not mapped for acc_detach");
+
+  gomp_detach_pointer (acc_dev, aq, n, (uintptr_t) hostaddr, finalize, NULL);
+}
+
+void
+acc_detach (void **hostaddr)
+{
+  goacc_detach_internal (hostaddr, acc_async_sync, false);
+}
+
+void
+acc_detach_async (void **hostaddr, int async)
+{
+  goacc_detach_internal (hostaddr, async, false);
+}
+
+void
+acc_detach_finalize (void **hostaddr)
+{
+  goacc_detach_internal (hostaddr, acc_async_sync, true);
+}
+
+void
+acc_detach_finalize_async (void **hostaddr, int async)
+{
+  goacc_detach_internal (hostaddr, async, true);
+}
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 02472c56270..69b9f379476 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -71,7 +71,10 @@  find_pointer (int pos, size_t mapnum, unsigned short *kinds)
       {
 	unsigned char kind1 = kinds[pos + 1] & 0xff;
 	if (kind1 == GOMP_MAP_POINTER
-	    || kind1 == GOMP_MAP_ALWAYS_POINTER)
+	    || kind1 == GOMP_MAP_ALWAYS_POINTER
+	    || kind1 == GOMP_MAP_ATTACH
+	    || kind1 == GOMP_MAP_DETACH
+	    || kind1 == GOMP_MAP_FORCE_DETACH)
 	  return 2;
 	else if (kind1 == GOMP_MAP_TO_PSET)
 	  return 3;
@@ -643,6 +646,10 @@  GOACC_enter_exit_data (int flags_m, size_t mapnum,
   if (mapnum > 0)
     {
       unsigned char kind = kinds[0] & 0xff;
+
+      if (kind == GOMP_MAP_STRUCT || kind == GOMP_MAP_FORCE_PRESENT)
+        kind = kinds[1] & 0xff;
+
       if (kind == GOMP_MAP_DELETE
 	  || kind == GOMP_MAP_FORCE_FROM)
 	finalize = true;
@@ -653,11 +660,14 @@  GOACC_enter_exit_data (int flags_m, size_t mapnum,
     {
       unsigned char kind = kinds[i] & 0xff;
 
-      if (kind == GOMP_MAP_POINTER || kind == GOMP_MAP_TO_PSET)
+      if (kind == GOMP_MAP_POINTER
+	  || kind == GOMP_MAP_TO_PSET
+	  || kind == GOMP_MAP_STRUCT)
 	continue;
 
       if (kind == GOMP_MAP_FORCE_ALLOC
 	  || kind == GOMP_MAP_FORCE_PRESENT
+	  || kind == GOMP_MAP_ATTACH
 	  || kind == GOMP_MAP_FORCE_TO
 	  || kind == GOMP_MAP_TO
 	  || kind == GOMP_MAP_ALLOC)
@@ -668,6 +678,8 @@  GOACC_enter_exit_data (int flags_m, size_t mapnum,
 
       if (kind == GOMP_MAP_RELEASE
 	  || kind == GOMP_MAP_DELETE
+	  || kind == GOMP_MAP_DETACH
+	  || kind == GOMP_MAP_FORCE_DETACH
 	  || kind == GOMP_MAP_FROM
 	  || kind == GOMP_MAP_FORCE_FROM)
 	break;
@@ -777,6 +789,19 @@  GOACC_enter_exit_data (int flags_m, size_t mapnum,
 		case GOMP_MAP_FORCE_TO:
 		  acc_copyin_async (hostaddrs[i], sizes[i], async);
 		  break;
+		case GOMP_MAP_STRUCT:
+		  {
+		    int elems = sizes[i];
+		    goacc_aq aq = get_goacc_asyncqueue (async);
+		    gomp_map_vars_async (acc_dev, aq, elems + 1, &hostaddrs[i],
+					 NULL, &sizes[i], &kinds[i], true,
+					 GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+		    i += elems;
+		  }
+		  break;
+		case GOMP_MAP_ATTACH:
+		case GOMP_MAP_FORCE_PRESENT:
+		  break;
 		default:
 		  gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
 			      kind);
@@ -800,48 +825,140 @@  GOACC_enter_exit_data (int flags_m, size_t mapnum,
 	      i += pointer - 1;
 	    }
 	}
+
+      /* This loop only handles explicit "attach" clauses that are not an
+	 implicit part of a copy{,in,out}, etc. mapping.  */
+      for (i = 0; i < mapnum; i++)
+        {
+	  unsigned char kind = kinds[i] & 0xff;
+
+	  /* Scan for pointers and PSETs.  */
+	  int pointer = find_pointer (i, mapnum, kinds);
+
+	  if (!pointer)
+	    {
+	      if (kind == GOMP_MAP_ATTACH)
+		acc_attach (hostaddrs[i]);
+	      else if (kind == GOMP_MAP_STRUCT)
+	        i += sizes[i];
+	    }
+	  else
+	    i += pointer - 1;
+	}
     }
   else
-    for (i = 0; i < mapnum; ++i)
-      {
-	unsigned char kind = kinds[i] & 0xff;
+    {
+      /* Handle "detach" before copyback/deletion of mapped data.  */
+      for (i = 0; i < mapnum; i++)
+        {
+	  unsigned char kind = kinds[i] & 0xff;
 
-	int pointer = find_pointer (i, mapnum, kinds);
+	  int pointer = find_pointer (i, mapnum, kinds);
 
-	if (!pointer)
-	  {
-	    switch (kind)
-	      {
-	      case GOMP_MAP_RELEASE:
-	      case GOMP_MAP_DELETE:
-		if (acc_is_present (hostaddrs[i], sizes[i]))
+	  if (!pointer)
+	    {
+	      if (kind == GOMP_MAP_DETACH)
+		acc_detach (hostaddrs[i]);
+	      else if (kind == GOMP_MAP_FORCE_DETACH)
+		acc_detach_finalize (hostaddrs[i]);
+	      else if (kind == GOMP_MAP_STRUCT)
+	        i += sizes[i];
+	    }
+	  else
+	    {
+	      unsigned char kind2 = kinds[i + pointer - 1] & 0xff;
+
+	      if (kind2 == GOMP_MAP_DETACH)
+		acc_detach (hostaddrs[i + pointer - 1]);
+	      else if (kind2 == GOMP_MAP_FORCE_DETACH)
+	        acc_detach_finalize (hostaddrs[i + pointer - 1]);
+
+	      i += pointer - 1;
+	    }
+	}
+
+      for (i = 0; i < mapnum; ++i)
+	{
+	  unsigned char kind = kinds[i] & 0xff;
+
+	  int pointer = find_pointer (i, mapnum, kinds);
+
+	  if (!pointer)
+	    {
+	      switch (kind)
+		{
+		case GOMP_MAP_RELEASE:
+		case GOMP_MAP_DELETE:
+		  if (acc_is_present (hostaddrs[i], sizes[i]))
+		    {
+		      if (finalize)
+			acc_delete_finalize_async (hostaddrs[i], sizes[i],
+						   async);
+		      else
+			acc_delete_async (hostaddrs[i], sizes[i], async);
+		    }
+		  break;
+		case GOMP_MAP_DETACH:
+		case GOMP_MAP_FORCE_DETACH:
+		case GOMP_MAP_FORCE_PRESENT:
+		  break;
+		case GOMP_MAP_FROM:
+		case GOMP_MAP_FORCE_FROM:
+		  if (finalize)
+		    acc_copyout_finalize_async (hostaddrs[i], sizes[i], async);
+		  else
+		    acc_copyout_async (hostaddrs[i], sizes[i], async);
+		  break;
+		case GOMP_MAP_STRUCT:
 		  {
-		    if (finalize)
-		      acc_delete_finalize_async (hostaddrs[i], sizes[i], async);
-		    else
-		      acc_delete_async (hostaddrs[i], sizes[i], async);
+		    int elems = sizes[i];
+		    goacc_aq aq = get_goacc_asyncqueue (async);
+		    for (int j = 1; j <= elems; j++)
+		      {
+			struct splay_tree_key_s k;
+			k.host_start = (uintptr_t) hostaddrs[i + j];
+			k.host_end = k.host_start + sizes[i + j];
+			splay_tree_key str;
+			gomp_mutex_lock (&acc_dev->lock);
+			str = splay_tree_lookup (&acc_dev->mem_map, &k);
+			gomp_mutex_unlock (&acc_dev->lock);
+			if (str)
+		          {
+			    assert (str->virtual_refcount
+				    != VREFCOUNT_LINK_KEY);
+			    if (finalize)
+			      {
+				str->refcount -= str->virtual_refcount;
+				str->virtual_refcount = 0;
+			      }
+			    if (str->virtual_refcount > 0)
+			      {
+				str->refcount--;
+				str->virtual_refcount--;
+			      }
+			    else if (str->refcount > 0)
+			      str->refcount--;
+			    if (str->refcount == 0)
+			      gomp_remove_var_async (acc_dev, str, aq);
+			  }
+		      }
+		    i += elems;
 		  }
-		break;
-	      case GOMP_MAP_FROM:
-	      case GOMP_MAP_FORCE_FROM:
-		if (finalize)
-		  acc_copyout_finalize_async (hostaddrs[i], sizes[i], async);
-		else
-		  acc_copyout_async (hostaddrs[i], sizes[i], async);
-		break;
-	      default:
-		gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
-			    kind);
-		break;
-	      }
-	  }
-	else
-	  {
-	    gomp_acc_remove_pointer (acc_dev, &hostaddrs[i], &sizes[i],
-				     &kinds[i], async, finalize, pointer);
-	    i += pointer - 1;
-	  }
-      }
+		  break;
+		default:
+		  gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind "
+			      "0x%.2x", kind);
+		  break;
+		}
+	    }
+	  else
+	    {
+	      gomp_acc_remove_pointer (acc_dev, &hostaddrs[i], &sizes[i],
+				       &kinds[i], async, finalize, pointer);
+	      i += pointer - 1;
+	    }
+	}
+    }
 
  out_prof:
   if (profiling_p)
diff --git a/libgomp/openacc.h b/libgomp/openacc.h
index 1bbe6c90e7f..afe31cfd99b 100644
--- a/libgomp/openacc.h
+++ b/libgomp/openacc.h
@@ -108,12 +108,18 @@  void *acc_hostptr (void *) __GOACC_NOTHROW;
 int acc_is_present (void *, size_t) __GOACC_NOTHROW;
 void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW;
 void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW;
+void acc_attach (void **) __GOACC_NOTHROW;
+void acc_attach_async (void **, int) __GOACC_NOTHROW;
+void acc_detach (void **) __GOACC_NOTHROW;
+void acc_detach_async (void **, int) __GOACC_NOTHROW;
 
 /* Finalize versions of copyout/delete functions, specified in OpenACC 2.5.  */
 void acc_copyout_finalize (void *, size_t) __GOACC_NOTHROW;
 void acc_copyout_finalize_async (void *, size_t, int) __GOACC_NOTHROW;
 void acc_delete_finalize (void *, size_t) __GOACC_NOTHROW;
 void acc_delete_finalize_async (void *, size_t, int) __GOACC_NOTHROW;
+void acc_detach_finalize (void **) __GOACC_NOTHROW;
+void acc_detach_finalize_async (void **, int) __GOACC_NOTHROW;
 
 /* Async functions, specified in OpenACC 2.5.  */
 void acc_copyin_async (void *, size_t, int) __GOACC_NOTHROW;
diff --git a/libgomp/target.c b/libgomp/target.c
index aa7a1df1e46..eb28591ee4f 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -381,8 +381,12 @@  dump_tgt (const char *where, struct target_mem_desc *tgt)
 	    fprintf (stderr, "  key.u.link_key=%p\n",
 		     (void*) tgt->list[i].key->u.link_key);
 	  else
-	    fprintf (stderr, "  key.virtual_refcount=%d\n",
-		     (int) tgt->list[i].key->virtual_refcount);
+	    {
+	      fprintf (stderr, "  key.virtual_refcount=%d\n",
+		       (int) tgt->list[i].key->virtual_refcount);
+	      fprintf (stderr, "  key.u.attach_count=%p\n",
+		       (void*) tgt->list[i].key->u.attach_count);
+	    }
 	}
     }
   fprintf (stderr, "\n");
@@ -537,6 +541,7 @@  gomp_map_vars_existing (struct gomp_device_descr *devicep,
   tgt_var->key = oldn;
   tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind);
   tgt_var->always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind);
+  tgt_var->do_detach = kind == GOMP_MAP_ATTACH;
   tgt_var->offset = newn->host_start - oldn->host_start;
   tgt_var->length = newn->host_end - newn->host_start;
 
@@ -670,6 +675,135 @@  gomp_map_fields_existing (struct target_mem_desc *tgt,
 	      (void *) cur_node.host_end);
 }
 
+void
+gomp_attach_pointer (struct gomp_device_descr *devicep,
+		     struct goacc_asyncqueue *aq, splay_tree mem_map,
+		     splay_tree_key n, uintptr_t attach_to, size_t bias,
+		     struct gomp_coalesce_buf *cbufp)
+{
+  struct splay_tree_key_s s;
+  size_t size, idx;
+
+  if (n == NULL)
+    {
+      gomp_mutex_unlock (&devicep->lock);
+      gomp_fatal ("enclosing struct not mapped for attach");
+    }
+
+  size = (n->host_end - n->host_start + sizeof (void *) - 1) / sizeof (void *);
+  /* We might have a pointer in a packed struct: however we cannot have more
+     than one such pointer in each pointer-sized portion of the struct, so
+     this is safe.  */
+  idx = (attach_to - n->host_start) / sizeof (void *);
+
+  assert (n->virtual_refcount != VREFCOUNT_LINK_KEY);
+
+  if (!n->u.attach_count)
+    n->u.attach_count
+      = gomp_malloc_cleared (sizeof (*n->u.attach_count) * size);
+
+  if (n->u.attach_count[idx] < UINTPTR_MAX)
+    n->u.attach_count[idx]++;
+  else
+    {
+      gomp_mutex_unlock (&devicep->lock);
+      gomp_fatal ("attach count overflow");
+    }
+
+  if (n->u.attach_count[idx] == 1)
+    {
+      uintptr_t devptr = n->tgt->tgt_start + n->tgt_offset + attach_to
+			 - n->host_start;
+      uintptr_t target = (uintptr_t) *(void **) attach_to;
+      splay_tree_key tn;
+      uintptr_t data;
+
+      if ((void *) target == NULL)
+	{
+	  gomp_mutex_unlock (&devicep->lock);
+	  gomp_fatal ("attempt to attach null pointer");
+	}
+
+      s.host_start = target + bias;
+      s.host_end = s.host_start + 1;
+      tn = splay_tree_lookup (mem_map, &s);
+
+      if (!tn)
+	{
+	  gomp_mutex_unlock (&devicep->lock);
+	  gomp_fatal ("pointer target not mapped for attach");
+	}
+
+      data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start;
+
+      gomp_debug (1,
+		  "%s: attaching host %p, target %p (struct base %p) to %p\n",
+		  __FUNCTION__, (void *) attach_to, (void *) devptr,
+		  (void *) (n->tgt->tgt_start + n->tgt_offset), (void *) data);
+
+      gomp_copy_host2dev (devicep, aq, (void *) devptr, (void *) &data,
+			  sizeof (void *), cbufp);
+    }
+  else
+    gomp_debug (1, "%s: attach count for %p -> %u\n", __FUNCTION__,
+		(void *) attach_to, (int) n->u.attach_count[idx]);
+}
+
+void
+gomp_detach_pointer (struct gomp_device_descr *devicep,
+		     struct goacc_asyncqueue *aq, splay_tree_key n,
+		     uintptr_t detach_from, bool finalize,
+		     struct gomp_coalesce_buf *cbufp)
+{
+  size_t idx;
+
+  if (n == NULL)
+    {
+      gomp_mutex_unlock (&devicep->lock);
+      gomp_fatal ("enclosing struct not mapped for detach");
+    }
+
+  idx = (detach_from - n->host_start) / sizeof (void *);
+
+  assert (n->virtual_refcount != VREFCOUNT_LINK_KEY);
+
+  if (!n->u.attach_count)
+    {
+      gomp_mutex_unlock (&devicep->lock);
+      gomp_fatal ("no attachment counters for struct");
+    }
+
+  if (finalize)
+    n->u.attach_count[idx] = 1;
+
+  if (n->u.attach_count[idx] == 0)
+    {
+      gomp_mutex_unlock (&devicep->lock);
+      gomp_fatal ("attach count underflow");
+    }
+  else
+    n->u.attach_count[idx]--;
+
+  if (n->u.attach_count[idx] == 0)
+    {
+      uintptr_t devptr = n->tgt->tgt_start + n->tgt_offset + detach_from
+			 - n->host_start;
+      uintptr_t target = (uintptr_t) *(void **) detach_from;
+
+      gomp_debug (1,
+		  "%s: detaching host %p, target %p (struct base %p) to %p\n",
+		  __FUNCTION__, (void *) detach_from, (void *) devptr,
+		  (void *) (n->tgt->tgt_start + n->tgt_offset),
+		  (void *) target);
+
+      gomp_copy_host2dev (devicep, aq, (void *) devptr, (void *) &target,
+			  sizeof (void *), cbufp);
+    }
+  else
+    gomp_debug (1, "%s: attach count for %p -> %u\n", __FUNCTION__,
+		(void *) detach_from, (int) n->u.attach_count[idx]);
+}
+
 uintptr_t
 gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
 {
@@ -822,8 +956,15 @@  gomp_map_vars_internal (struct gomp_device_descr *devicep,
 	  has_firstprivate = true;
 	  continue;
 	}
+      else if ((kind & typemask) == GOMP_MAP_ATTACH)
+	{
+	  tgt->list[i].key = NULL;
+	  has_firstprivate = true;
+	  continue;
+	}
       cur_node.host_start = (uintptr_t) hostaddrs[i];
-      if (!GOMP_MAP_POINTER_P (kind & typemask))
+      if (!GOMP_MAP_POINTER_P (kind & typemask)
+          && (kind & typemask) != GOMP_MAP_ATTACH)
 	cur_node.host_end = cur_node.host_start + sizes[i];
       else
 	cur_node.host_end = cur_node.host_start + sizeof (void *);
@@ -1047,6 +1188,32 @@  gomp_map_vars_internal (struct gomp_device_descr *devicep,
 		cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
 				      + cur_node.host_start - n->host_start;
 		continue;
+	      case GOMP_MAP_ATTACH:
+		{
+		  cur_node.host_start = (uintptr_t) hostaddrs[i];
+		  cur_node.host_end = cur_node.host_start + sizeof (void *);
+		  splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+		  if (n != NULL)
+		    {
+		      tgt->list[i].key = n;
+		      tgt->list[i].offset = cur_node.host_start - n->host_start;
+		      tgt->list[i].length = n->host_end - n->host_start;
+		      tgt->list[i].copy_from = false;
+		      tgt->list[i].always_copy_from = false;
+		      tgt->list[i].do_detach
+		        = (pragma_kind != GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+		      n->refcount++;
+		    }
+		  else
+		    {
+		      gomp_mutex_unlock (&devicep->lock);
+		      gomp_fatal ("outer struct not mapped for attach");
+		    }
+		  gomp_attach_pointer (devicep, aq, mem_map, n,
+				       (uintptr_t) hostaddrs[i], sizes[i],
+				       cbufp);
+		  continue;
+		}
 	      default:
 		break;
 	      }
@@ -1090,10 +1257,12 @@  gomp_map_vars_internal (struct gomp_device_descr *devicep,
 		tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask);
 		tgt->list[i].always_copy_from
 		  = GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
+		tgt->list[i].do_detach = false;
 		tgt->list[i].offset = 0;
 		tgt->list[i].length = k->host_end - k->host_start;
 		k->refcount = 1;
 		k->virtual_refcount = 0;
+		k->u.attach_count = NULL;
 		tgt->refcount++;
 		array->left = NULL;
 		array->right = NULL;
@@ -1144,6 +1313,7 @@  gomp_map_vars_internal (struct gomp_device_descr *devicep,
 			  tgt->list[j].key = k;
 			  tgt->list[j].copy_from = false;
 			  tgt->list[j].always_copy_from = false;
+			  tgt->list[j].do_detach = false;
 			  if (k->refcount != REFCOUNT_INFINITY)
 			    k->refcount++;
 			  gomp_map_pointer (tgt, aq,
@@ -1319,6 +1489,8 @@  gomp_remove_var_internal (struct gomp_device_descr *devicep, splay_tree_key k,
       if (k->u.link_key)
 	splay_tree_insert (&devicep->mem_map, (splay_tree_node) k->u.link_key);
     }
+  else if (k->u.attach_count)
+    free (k->u.attach_count);
   if (aq)
     devicep->openacc.async.queue_callback_func (aq, gomp_unref_tgt_void,
 						(void *) k->tgt);
@@ -1372,6 +1544,18 @@  gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom,
     }
 
   size_t i;
+
+  /* We must perform detachments before any copies back to the host.  */
+  for (i = 0; i < tgt->list_count; i++)
+    {
+      splay_tree_key k = tgt->list[i].key;
+
+      if (k != NULL && tgt->list[i].do_detach)
+	gomp_detach_pointer (devicep, aq, k, tgt->list[i].key->host_start
+					     + tgt->list[i].offset,
+			     k->refcount == 1, NULL);
+    }
+
   for (i = 0; i < tgt->list_count; i++)
     {
       splay_tree_key k = tgt->list[i].key;
@@ -1545,6 +1729,7 @@  gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
       k->tgt_offset = target_table[i].start;
       k->refcount = REFCOUNT_INFINITY;
       k->virtual_refcount = 0;
+      k->u.attach_count = NULL;
       array->left = NULL;
       array->right = NULL;
       splay_tree_insert (&devicep->mem_map, array);
@@ -1577,6 +1762,7 @@  gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
       k->tgt_offset = target_var->start;
       k->refcount = target_size & link_bit ? REFCOUNT_LINK : REFCOUNT_INFINITY;
       k->virtual_refcount = 0;
+      k->u.attach_count = NULL;
       array->left = NULL;
       array->right = NULL;
       splay_tree_insert (&devicep->mem_map, array);
diff --git a/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C
new file mode 100644
index 00000000000..771876afd2d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C
@@ -0,0 +1,72 @@ 
+#include <stdlib.h>
+
+/* Test attach/detach with dereferences of reference to pointer to struct.  */
+
+typedef struct {
+  int *a;
+  int *b;
+  int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct *m = (mystruct *) malloc (sizeof (*m));
+  mystruct *&mref = m;
+  int i;
+
+  mref->a = (int *) malloc (N * sizeof (int));
+  m->b = (int *) malloc (N * sizeof (int));
+  m->c = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      mref->a[i] = 0;
+      m->b[i] = 0;
+      m->c[i] = 0;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+#pragma acc parallel loop copy(mref->a[0:N])
+      for (j = 0; j < N; j++)
+        mref->a[j]++;
+#pragma acc parallel loop copy(mref->b[0:N], m->c[5:N-10])
+      for (j = 0; j < N; j++)
+	{
+	  mref->b[j]++;
+	  if (j > 5 && j < N - 5)
+	    m->c[j]++;
+	}
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 99)
+        abort ();
+      if (m->b[i] != 99)
+        abort ();
+      if (i > 5 && i < N-5)
+	{
+	  if (m->c[i] != 99)
+	    abort ();
+	}
+      else
+	{
+	  if (m->c[i] != 0)
+	    abort ();
+	}
+    }
+
+  free (m->a);
+  free (m->b);
+  free (m->c);
+  free (m);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C
new file mode 100644
index 00000000000..98cf450c61d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C
@@ -0,0 +1,72 @@ 
+#include <stdlib.h>
+
+/* Test array slice with reference to pointer.  */
+
+typedef struct {
+  int *a;
+  int *b;
+  int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct *m = (mystruct *) malloc (sizeof (*m));
+  int i;
+
+  m->a = (int *) malloc (N * sizeof (int));
+  m->b = (int *) malloc (N * sizeof (int));
+  m->c = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      m->a[i] = 0;
+      m->b[i] = 0;
+      m->c[i] = 0;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+      int *&ptr = m->a;
+#pragma acc parallel loop copy(ptr[0:N])
+      for (j = 0; j < N; j++)
+        ptr[j]++;
+#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10])
+      for (j = 0; j < N; j++)
+	{
+	  m->b[j]++;
+	  if (j > 5 && j < N - 5)
+	    m->c[j]++;
+	}
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 99)
+        abort ();
+      if (m->b[i] != 99)
+        abort ();
+      if (i > 5 && i < N-5)
+	{
+	  if (m->c[i] != 99)
+	    abort ();
+	}
+      else
+	{
+	  if (m->c[i] != 0)
+	    abort ();
+	}
+    }
+
+  free (m->a);
+  free (m->b);
+  free (m->c);
+  free (m);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c
new file mode 100644
index 00000000000..d8d7067e452
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c
@@ -0,0 +1,24 @@ 
+#include <stdlib.h>
+#include <assert.h>
+
+struct dc
+{
+  int a;
+  int *b;
+};
+
+int
+main ()
+{
+  int n = 100, i;
+  struct dc v = { .a = 3, .b = (int *) malloc (sizeof (int) * n) };
+
+#pragma acc parallel loop copy(v.a, v.b[:n])
+  for (i = 0; i < n; i++)
+    v.b[i] = v.a;
+
+  for (i = 0; i < 10; i++)
+    assert (v.b[i] == v.a);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c
new file mode 100644
index 00000000000..37cde4ef059
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c
@@ -0,0 +1,53 @@ 
+#include <stdlib.h>
+
+/* Test asyncronous attach and detach operation.  */
+
+typedef struct {
+  int *a;
+  int *b;
+} mystruct;
+
+int
+main (int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct m;
+  int i;
+
+  m.a = (int *) malloc (N * sizeof (int));
+  m.b = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      m.a[i] = 0;
+      m.b[i] = 0;
+    }
+
+#pragma acc enter data copyin(m)
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+#pragma acc parallel loop copy(m.a[0:N]) async(i % 2)
+      for (j = 0; j < N; j++)
+        m.a[j]++;
+#pragma acc parallel loop copy(m.b[0:N]) async((i + 1) % 2)
+      for (j = 0; j < N; j++)
+        m.b[j]++;
+    }
+
+#pragma acc exit data copyout(m) wait(0, 1)
+
+  for (i = 0; i < N; i++)
+    {
+      if (m.a[i] != 99)
+        abort ();
+      if (m.b[i] != 99)
+        abort ();
+    }
+
+  free (m.a);
+  free (m.b);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c
new file mode 100644
index 00000000000..ed8ddcc54fa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c
@@ -0,0 +1,72 @@ 
+#include <stdlib.h>
+
+/* Test multiple struct dereferences on one directive, and slices starting at
+   non-zero.  */
+
+typedef struct {
+  int *a;
+  int *b;
+  int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct *m = (mystruct *) malloc (sizeof (*m));
+  int i;
+
+  m->a = (int *) malloc (N * sizeof (int));
+  m->b = (int *) malloc (N * sizeof (int));
+  m->c = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      m->a[i] = 0;
+      m->b[i] = 0;
+      m->c[i] = 0;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+#pragma acc parallel loop copy(m->a[0:N])
+      for (j = 0; j < N; j++)
+        m->a[j]++;
+#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10])
+      for (j = 0; j < N; j++)
+	{
+	  m->b[j]++;
+	  if (j > 5 && j < N - 5)
+	    m->c[j]++;
+	}
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 99)
+        abort ();
+      if (m->b[i] != 99)
+        abort ();
+      if (i > 5 && i < N-5)
+	{
+	  if (m->c[i] != 99)
+	    abort ();
+	}
+      else
+	{
+	  if (m->c[i] != 0)
+	    abort ();
+	}
+    }
+
+  free (m->a);
+  free (m->b);
+  free (m->c);
+  free (m);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c
new file mode 100644
index 00000000000..04c70aed7b5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c
@@ -0,0 +1,63 @@ 
+#include <openacc.h>
+#include <stdlib.h>
+
+/* Test attach/detach operation with chained dereferences.  */
+
+typedef struct mystruct {
+  int *a;
+  struct mystruct *next;
+} mystruct;
+
+int
+main (int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct *m = (mystruct *) malloc (sizeof (*m));
+  int i;
+
+  m->a = (int *) malloc (N * sizeof (int));
+  m->next = (mystruct *) malloc (sizeof (*m));
+  m->next->a = (int *) malloc (N * sizeof (int));
+  m->next->next = NULL;
+
+  for (i = 0; i < N; i++)
+    {
+      m->a[i] = 0;
+      m->next->a[i] = 0;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+  acc_copyin (m->next, sizeof (*m));
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+      acc_copyin (m->next->a, N * sizeof (int));
+      acc_attach ((void **) &m->next);
+      /* This will attach only the innermost pointer, i.e. "a[0:N]".  That's
+	 why we have to attach the "m->next" pointer manually above.  */
+#pragma acc parallel loop copy(m->next->a[0:N])
+      for (j = 0; j < N; j++)
+        m->next->a[j]++;
+      acc_detach ((void **) &m->next);
+      acc_copyout (m->next->a, N * sizeof (int));
+    }
+
+  acc_copyout (m->next, sizeof (*m));
+#pragma acc exit data copyout(m[0:1])
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 0)
+        abort ();
+      if (m->next->a[i] != 99)
+        abort ();
+    }
+
+  free (m->next->a);
+  free (m->next);
+  free (m->a);
+  free (m);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c
new file mode 100644
index 00000000000..7e26e9aa8b9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c
@@ -0,0 +1,29 @@ 
+#include <assert.h>
+#include <stdlib.h>
+
+int
+main(int argc, char* argv[])
+{
+  struct foo {
+    int *a, *b, c, d, *e;
+  } s;
+
+  s.a = (int *) malloc (16 * sizeof (int));
+  s.b = (int *) malloc (16 * sizeof (int));
+  s.e = (int *) malloc (16 * sizeof (int));
+
+  #pragma acc data copy(s)
+  {
+    #pragma acc data copy(s.a[0:10])
+    {
+      #pragma acc parallel loop attach(s.a)
+      for (int i = 0; i < 10; i++)
+	s.a[i] = i;
+    }
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (s.a[i] == i);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c
new file mode 100644
index 00000000000..cec764bd3e7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c
@@ -0,0 +1,34 @@ 
+#include <assert.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+int
+main ()
+{
+  int n = 100, i;
+  int *a = (int *) malloc (sizeof (int) * n);
+  int *b;
+
+  for (i = 0; i < n; i++)
+    a[i] = i+1;
+
+#pragma acc enter data copyin(a[:n]) create(b)
+
+  b = a;
+  acc_attach ((void **)&b);
+
+#pragma acc parallel loop present (b[:n])
+  for (i = 0; i < n; i++)
+    b[i] = i+1;
+
+  acc_detach ((void **)&b);
+
+#pragma acc exit data copyout(a[:n], b)
+
+  for (i = 0; i < 10; i++)
+    assert (a[i] == b[i]);
+
+  free (a);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c
new file mode 100644
index 00000000000..8874ca0a504
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c
@@ -0,0 +1,87 @@ 
+#include <assert.h>
+#include <stdlib.h>
+
+#define LIST_LENGTH 10
+
+struct node
+{
+  struct node *next;
+  int val;
+};
+
+int
+sum_nodes (struct node *head)
+{
+  int i = 0, sum = 0;
+
+#pragma acc parallel reduction(+:sum) present(head[:1])
+  {
+    for (; head != NULL; head = head->next)
+      sum += head->val;
+  }
+
+  return sum;
+}
+
+void
+insert (struct node *head, int val)
+{
+  struct node *n = (struct node *) malloc (sizeof (struct node));
+
+  if (head->next)
+    {
+#pragma acc exit data detach(head->next)
+    }
+
+  n->val = val;
+  n->next = head->next;
+  head->next = n;
+
+#pragma acc enter data copyin(n[:1])
+#pragma acc enter data attach(head->next)
+  if (n->next)
+    {
+#pragma acc enter data attach(n->next)
+    }
+}
+
+void
+destroy (struct node *head)
+{
+  while (head->next != NULL)
+    {
+#pragma acc exit data detach(head->next)
+      struct node * n = head->next;
+      head->next = n->next;
+      if (n->next)
+	{
+#pragma acc exit data detach(n->next)
+	}
+#pragma acc exit data delete (n[:1])
+      if (head->next)
+	{
+#pragma acc enter data attach(head->next)
+	}
+      free (n);
+    }
+}
+
+int
+main ()
+{
+  struct node list = { .next = NULL, .val = 0 };
+  int i;
+
+#pragma acc enter data copyin(list)
+
+  for (i = 0; i < LIST_LENGTH; i++)
+    insert (&list, i + 1);
+
+  assert (sum_nodes (&list) == (LIST_LENGTH * LIST_LENGTH + LIST_LENGTH) / 2);
+
+  destroy (&list);
+
+#pragma acc exit data delete(list)
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c
new file mode 100644
index 00000000000..89cafbb62ab
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c
@@ -0,0 +1,81 @@ 
+#include <assert.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+struct node
+{
+  struct node *next;
+  int val;
+};
+
+int
+sum_nodes (struct node *head)
+{
+  int i = 0, sum = 0;
+
+#pragma acc parallel reduction(+:sum) present(head[:1])
+  {
+    for (; head != NULL; head = head->next)
+      sum += head->val;
+  }
+
+  return sum;
+}
+
+void
+insert (struct node *head, int val)
+{
+  struct node *n = (struct node *) malloc (sizeof (struct node));
+
+  if (head->next)
+    acc_detach ((void **) &head->next);
+
+  n->val = val;
+  n->next = head->next;
+  head->next = n;
+
+  acc_copyin (n, sizeof (struct node));
+  acc_attach((void **) &head->next);
+
+  if (n->next)
+    acc_attach ((void **) &n->next);
+}
+
+void
+destroy (struct node *head)
+{
+  while (head->next != NULL)
+    {
+      acc_detach ((void **) &head->next);
+      struct node * n = head->next;
+      head->next = n->next;
+      if (n->next)
+	acc_detach ((void **) &n->next);
+
+      acc_delete (n, sizeof (struct node));
+      if (head->next)
+	acc_attach((void **) &head->next);
+
+      free (n);
+    }
+}
+
+int
+main ()
+{
+  struct node list = { .next = NULL, .val = 0 };
+  int i;
+
+  acc_copyin (&list, sizeof (struct node));
+
+  for (i = 0; i < 10; i++)
+    insert (&list, 2);
+
+  assert (sum_nodes (&list) == 10 * 2);
+
+  destroy (&list);
+
+  acc_delete (&list, sizeof (struct node));
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c
new file mode 100644
index 00000000000..81c1c5e1093
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c
@@ -0,0 +1,59 @@ 
+/* { dg-do run { target { ! openacc_host_selected } } } */
+
+#include <stdlib.h>
+#include <assert.h>
+#include <openacc.h>
+
+struct dc
+{
+  int a;
+  int **b;
+};
+
+int
+main ()
+{
+  int n = 100, i, j, k;
+  struct dc v = { .a = 3 };
+
+  v.b = (int **) malloc (sizeof (int *) * n);
+  for (i = 0; i < n; i++)
+    v.b[i] = (int *) malloc (sizeof (int) * n);
+
+  for (k = 0; k < 16; k++)
+    {
+#pragma acc data copy(v)
+      {
+#pragma acc data copy(v.b[:n])
+	{
+	  for (i = 0; i < n; i++)
+	    {
+	      acc_copyin (v.b[i], sizeof (int) * n);
+	      acc_attach ((void **) &v.b[i]);
+	    }
+
+#pragma acc parallel loop
+	  for (i = 0; i < n; i++)
+	    for (j = 0; j < n; j++)
+	      v.b[i][j] = v.a + i + j;
+
+	  for (i = 0; i < n; i++)
+	    {
+	      acc_detach ((void **) &v.b[i]);
+	      acc_copyout (v.b[i], sizeof (int) * n);
+	    }
+	}
+      }
+
+      for (i = 0; i < n; i++)
+	for (j = 0; j < n; j++)
+	  assert (v.b[i][j] == v.a + i + j);
+
+      assert (!acc_is_present (&v, sizeof (v)));
+      assert (!acc_is_present (v.b, sizeof (int *) * n));
+      for (i = 0; i < n; i++)
+        assert (!acc_is_present (v.b[i], sizeof (int) * n));
+    }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c
new file mode 100644
index 00000000000..a59047af520
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c
@@ -0,0 +1,45 @@ 
+/* { dg-do run { target { ! openacc_host_selected } } } */
+
+#include <stdlib.h>
+#include <assert.h>
+#include <openacc.h>
+
+struct dc
+{
+  int a;
+  int *b;
+};
+
+int
+main ()
+{
+  int n = 100, i, j, k;
+  struct dc v = { .a = 3 };
+
+  v.b = (int *) malloc (sizeof (int) * n);
+
+  for (k = 0; k < 16; k++)
+    {
+      /* Here, we do not explicitly copy the enclosing structure, but work
+	 with fields directly.  Make sure attachment counters and reference
+	 counters work properly in that case.  */
+#pragma acc enter data copyin(v.a, v.b[0:n])
+#pragma acc enter data pcopyin(v.b[0:n])
+#pragma acc enter data pcopyin(v.b[0:n])
+
+#pragma acc parallel loop present(v.a, v.b)
+      for (i = 0; i < n; i++)
+	v.b[i] = v.a + i;
+
+#pragma acc exit data copyout(v.b[:n]) finalize
+#pragma acc exit data delete(v.a)
+
+      for (i = 0; i < n; i++)
+	assert (v.b[i] == v.a + i);
+
+      assert (!acc_is_present (&v, sizeof (v)));
+      assert (!acc_is_present (v.b, sizeof (int *) * n));
+    }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c
new file mode 100644
index 00000000000..0ca5990b377
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c
@@ -0,0 +1,54 @@ 
+/* { dg-do run { target { ! openacc_host_selected } } } */
+
+#include <stdlib.h>
+#include <assert.h>
+#include <openacc.h>
+
+struct dc
+{
+  int a;
+  int *b;
+  int *c;
+  int *d;
+};
+
+int
+main ()
+{
+  int n = 100, i, j, k;
+  struct dc v = { .a = 3 };
+
+  v.b = (int *) malloc (sizeof (int) * n);
+  v.c = (int *) malloc (sizeof (int) * n);
+  v.d = (int *) malloc (sizeof (int) * n);
+
+#pragma acc enter data copyin(v)
+
+  for (k = 0; k < 16; k++)
+    {
+#pragma acc enter data copyin(v.a, v.b[:n], v.c[:n], v.d[:n])
+
+#pragma acc parallel loop
+      for (i = 0; i < n; i++)
+	v.b[i] = v.a + i;
+
+#pragma acc exit data copyout(v.b[:n])
+#pragma acc exit data copyout(v.c[:n])
+#pragma acc exit data copyout(v.d[:n])
+#pragma acc exit data copyout(v.a)
+
+      for (i = 0; i < n; i++)
+	assert (v.b[i] == v.a + i);
+
+      assert (acc_is_present (&v, sizeof (v)));
+      assert (!acc_is_present (v.b, sizeof (int *) * n));
+      assert (!acc_is_present (v.c, sizeof (int *) * n));
+      assert (!acc_is_present (v.d, sizeof (int *) * n));
+    }
+
+#pragma acc exit data copyout(v)
+
+  assert (!acc_is_present (&v, sizeof (v)));
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c
new file mode 100644
index 00000000000..28535c9e281
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c
@@ -0,0 +1,53 @@ 
+#include <stdlib.h>
+
+typedef struct {
+  int *a;
+  int *b;
+} mystruct;
+
+int
+main (int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct *m = (mystruct *) malloc (sizeof (*m));
+  int i;
+
+  m->a = (int *) malloc (N * sizeof (int));
+  m->b = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      m->a[i] = 0;
+      m->b[i] = 0;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+      int *ptr = m->a;
+#pragma acc parallel loop copy(m->a[0:N])
+      for (j = 0; j < N; j++)
+        m->a[j]++;
+#pragma acc parallel loop copy(m->b[0:N]) 
+      for (j = 0; j < N; j++)
+	m->b[j]++;
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 99)
+        abort ();
+      if (m->b[i] != 99)
+        abort ();
+    }
+
+  free (m->a);
+  free (m->b);
+  free (m);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-1.f90
new file mode 100644
index 00000000000..c4cea11b571
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-1.f90
@@ -0,0 +1,35 @@ 
+! { dg-do run }
+
+! Test of attach/detach with "acc data".
+
+program dtype
+  implicit none
+  integer, parameter :: n = 512
+  type mytype
+    integer, allocatable :: a(:)
+  end type mytype
+  integer i
+
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+
+!$acc data copy(var)
+!$acc data copy(var%a)
+
+!$acc parallel loop
+  do i = 1,n
+    var%a(i) = i
+  end do
+!$acc end parallel loop
+
+!$acc end data
+!$acc end data
+
+  do i = 1,n
+    if (i .ne. var%a(i)) stop 1
+  end do
+
+  deallocate(var%a)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90
new file mode 100644
index 00000000000..35936617b87
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90
@@ -0,0 +1,33 @@ 
+! { dg-do run }
+
+! Test of attach/detach with "acc data", two clauses at once.
+
+program dtype
+  implicit none
+  integer, parameter :: n = 512
+  type mytype
+    integer, allocatable :: a(:)
+  end type mytype
+  integer i
+
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+
+!$acc data copy(var) copy(var%a)
+
+!$acc parallel loop
+  do i = 1,n
+    var%a(i) = i
+  end do
+!$acc end parallel loop
+
+!$acc end data
+
+  do i = 1,n
+    if (i .ne. var%a(i)) stop 1
+  end do
+
+  deallocate(var%a)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90
new file mode 100644
index 00000000000..667d944fecb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90
@@ -0,0 +1,34 @@ 
+! { dg-do run }
+
+! Test of attach/detach with "acc parallel".
+
+program dtype
+  implicit none
+  integer, parameter :: n = 512
+  type mytype
+    integer, allocatable :: a(:)
+    integer, allocatable :: b(:)
+  end type mytype
+  integer i
+
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+  allocate(var%b(1:n))
+
+!$acc parallel loop copy(var) copy(var%a(1:n)) copy(var%b(1:n))
+  do i = 1,n
+    var%a(i) = i
+    var%b(i) = i
+  end do
+!$acc end parallel loop
+
+  do i = 1,n
+    if (i .ne. var%a(i)) stop 1
+    if (i .ne. var%b(i)) stop 2
+  end do
+
+  deallocate(var%a)
+  deallocate(var%b)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-4.f90
new file mode 100644
index 00000000000..6949e120c9f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-4.f90
@@ -0,0 +1,49 @@ 
+! { dg-do run }
+
+! Test of attach/detach with "acc enter/exit data".
+
+program dtype
+  implicit none
+  integer, parameter :: n = 512
+  type mytype
+    integer, allocatable :: a(:)
+    integer, allocatable :: b(:)
+  end type mytype
+  integer, allocatable :: r(:)
+  integer i
+
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+  allocate(var%b(1:n))
+  allocate(r(1:n))
+
+!$acc enter data copyin(var)
+
+!$acc enter data copyin(var%a, var%b, r)
+
+!$acc parallel loop
+  do i = 1,n
+    var%a(i) = i
+    var%b(i) = i * 2
+    r(i) = i * 3
+  end do
+!$acc end parallel loop
+
+!$acc exit data copyout(var%a)
+!$acc exit data copyout(var%b)
+!$acc exit data copyout(r)
+
+  do i = 1,n
+    if (i .ne. var%a(i)) stop 1
+    if (i * 2 .ne. var%b(i)) stop 2
+    if (i * 3 .ne. r(i)) stop 3
+  end do
+
+!$acc exit data delete(var)
+
+  deallocate(var%a)
+  deallocate(var%b)
+  deallocate(r)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-5.f90
new file mode 100644
index 00000000000..6843cf1d0fa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-5.f90
@@ -0,0 +1,57 @@ 
+! { dg-do run }
+
+! Test of attach/detach, "enter data" inside "data", and subarray.
+
+program dtype
+  implicit none
+  integer, parameter :: n = 512
+  type mytype
+    integer, allocatable :: a(:)
+    integer, allocatable :: b(:)
+  end type mytype
+  integer i
+
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+  allocate(var%b(1:n))
+
+!$acc data copy(var)
+
+  do i = 1, n
+    var%a(i) = 0
+    var%b(i) = 0
+  end do
+
+!$acc enter data copyin(var%a(5:n - 5), var%b(5:n - 5))
+
+!$acc parallel loop
+  do i = 5,n - 5
+    var%a(i) = i
+    var%b(i) = i * 2
+  end do
+!$acc end parallel loop
+
+!$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5))
+
+!$acc end data
+
+  do i = 1,4
+    if (var%a(i) .ne. 0) stop 1
+    if (var%b(i) .ne. 0) stop 2
+  end do
+
+  do i = 5,n - 5
+    if (i .ne. var%a(i)) stop 3
+    if (i * 2 .ne. var%b(i)) stop 4
+  end do
+
+  do i = n - 4,n
+    if (var%a(i) .ne. 0) stop 5
+    if (var%b(i) .ne. 0) stop 6
+  end do
+
+  deallocate(var%a)
+  deallocate(var%b)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
new file mode 100644
index 00000000000..12910d0d655
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
@@ -0,0 +1,61 @@ 
+! { dg-do run }
+
+! Test of attachment counters and finalize.
+
+program dtype
+  implicit none
+  integer, parameter :: n = 512
+  type mytype
+    integer, allocatable :: a(:)
+    integer, allocatable :: b(:)
+  end type mytype
+  integer i
+
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+  allocate(var%b(1:n))
+
+!$acc data copy(var)
+
+  do i = 1, n
+    var%a(i) = 0
+    var%b(i) = 0
+  end do
+
+!$acc enter data copyin(var%a(5:n - 5), var%b(5:n - 5))
+
+  do i = 1,20
+    !$acc enter data attach(var%a)
+  end do
+
+!$acc parallel loop
+  do i = 5,n - 5
+    var%a(i) = i
+    var%b(i) = i * 2
+  end do
+!$acc end parallel loop
+
+!$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5)) finalize
+
+!$acc end data
+
+  do i = 1,4
+    if (var%a(i) .ne. 0) stop 1
+    if (var%b(i) .ne. 0) stop 2
+  end do
+
+  do i = 5,n - 5
+    if (i .ne. var%a(i)) stop 3
+    if (i * 2 .ne. var%b(i)) stop 4
+  end do
+
+  do i = n - 4,n
+    if (var%a(i) .ne. 0) stop 5
+    if (var%b(i) .ne. 0) stop 6
+  end do
+
+  deallocate(var%a)
+  deallocate(var%b)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-7.f90
new file mode 100644
index 00000000000..ab44f0a73b9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-7.f90
@@ -0,0 +1,89 @@ 
+! { dg-do run }
+
+! Test of attach/detach with scalar elements and nested derived types.
+
+program dtype
+  implicit none
+  integer, parameter :: n = 512
+  type subtype
+    integer :: g, h
+    integer, allocatable :: q(:)
+  end type subtype
+  type mytype
+    integer, allocatable :: a(:)
+    integer, allocatable :: c, d
+    integer, allocatable :: b(:)
+    integer :: f
+    type(subtype) :: s
+  end type mytype
+  integer i
+
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+  allocate(var%b(1:n))
+  allocate(var%c)
+  allocate(var%d)
+  allocate(var%s%q(1:n))
+
+  var%c = 16
+  var%d = 20
+  var%f = 7
+  var%s%g = 21
+  var%s%h = 38
+
+!$acc enter data copyin(var)
+
+  do i = 1, n
+    var%a(i) = 0
+    var%b(i) = 0
+    var%s%q(i) = 0
+  end do
+
+!$acc data copy(var%a(5:n - 5), var%b(5:n - 5), var%c, var%d) &
+!$acc & copy(var%s%q)
+
+!$acc parallel loop default(none) present(var)
+  do i = 5,n - 5
+    var%a(i) = i
+    var%b(i) = i * 2
+    var%s%q(i) = i * 3
+    var%s%g = 100
+    var%s%h = 101
+  end do
+!$acc end parallel loop
+
+!$acc end data
+
+!$acc exit data copyout(var)
+
+  do i = 1,4
+    if (var%a(i) .ne. 0) stop 1
+    if (var%b(i) .ne. 0) stop 2
+    if (var%s%q(i) .ne. 0) stop 3
+  end do
+
+  do i = 5,n - 5
+    if (i .ne. var%a(i)) stop 4
+    if (i * 2 .ne. var%b(i)) stop 5
+    if (i * 3 .ne. var%s%q(i)) stop 6
+  end do
+
+  do i = n - 4,n
+    if (var%a(i) .ne. 0) stop 7
+    if (var%b(i) .ne. 0) stop 8
+    if (var%s%q(i) .ne. 0) stop 9
+  end do
+
+  if (var%c .ne. 16) stop 10
+  if (var%d .ne. 20) stop 11
+  if (var%s%g .ne. 100 .or. var%s%h .ne. 101) stop 12
+  if (var%f .ne. 7) stop 13
+
+  deallocate(var%a)
+  deallocate(var%b)
+  deallocate(var%c)
+  deallocate(var%d)
+  deallocate(var%s%q)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-8.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-8.f90
new file mode 100644
index 00000000000..d142763ae59
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-8.f90
@@ -0,0 +1,41 @@ 
+! { dg-do run }
+
+! Test of explicit attach/detach clauses and attachment counters. There are no
+! acc_attach/acc_detach API routines in Fortran.
+
+program dtype
+  use openacc
+  implicit none
+  integer, parameter :: n = 512
+  type mytype
+    integer, allocatable :: a(:)
+  end type mytype
+  integer i
+
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+
+  call acc_copyin(var)
+  call acc_copyin(var%a)
+
+  !$acc enter data attach(var%a)
+
+!$acc parallel loop attach(var%a)
+  do i = 1,n
+    var%a(i) = i
+  end do
+!$acc end parallel loop
+
+  !$acc exit data detach(var%a)
+
+  call acc_copyout(var%a)
+  call acc_copyout(var)
+
+  do i = 1,n
+    if (i .ne. var%a(i)) stop 1
+  end do
+
+  deallocate(var%a)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/derived-type-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/derived-type-1.f90
new file mode 100644
index 00000000000..eb7812d541e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/derived-type-1.f90
@@ -0,0 +1,28 @@ 
+! Test derived types with subarrays
+
+! { dg-do run }
+
+  implicit none
+  type dtype
+     integer :: a, b, c
+  end type dtype
+  integer, parameter :: n = 100
+  integer i
+  type (dtype), dimension(n) :: d
+
+  !$acc data copy(d(1:n))
+  !$acc parallel loop
+  do i = 1, n
+     d(i)%a = i
+     d(i)%b = i-1
+     d(i)%c = i+1
+  end do
+  !$acc end data
+
+  do i = 1, n
+     if (d(i)%a /= i) stop 1
+     if (d(i)%b /= i-1) stop 2
+     if (d(i)%c /= i+1) stop 3
+  end do
+end program
+
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-1.f95 b/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-1.f95
new file mode 100644
index 00000000000..75ce48ddca2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-1.f95
@@ -0,0 +1,30 @@ 
+! { dg-do run }
+
+program main
+  implicit none
+
+  type mytype
+    integer :: a, b, c
+  end type mytype
+
+  type(mytype) :: myvar
+  integer :: i
+
+  myvar%a = 0
+  myvar%b = 0
+  myvar%c = 0
+
+!$acc enter data copyin(myvar)
+
+!$acc parallel present(myvar)
+  myvar%a = 1
+  myvar%b = 2
+  myvar%c = 3
+!$acc end parallel
+
+!$acc exit data copyout(myvar)
+
+  if (myvar%a .ne. 1) stop 1
+  if (myvar%b .ne. 2) stop 2
+  if (myvar%c .ne. 3) stop 3
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-2.f95 b/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-2.f95
new file mode 100644
index 00000000000..3088b832957
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-2.f95
@@ -0,0 +1,41 @@ 
+! { dg-do run }
+
+program main
+  implicit none
+
+  type tnest
+    integer :: ia, ib, ic
+  end type tnest
+
+  type mytype
+    type(tnest) :: nest
+    integer :: a, b, c
+  end type mytype
+
+  type(mytype) :: myvar
+  integer :: i
+
+  myvar%a = 0
+  myvar%b = 0
+  myvar%c = 0
+  myvar%nest%ia = 0
+  myvar%nest%ib = 0
+  myvar%nest%ic = 0
+
+!$acc enter data copyin(myvar%nest)
+
+!$acc parallel present(myvar%nest)
+  myvar%nest%ia = 4
+  myvar%nest%ib = 5
+  myvar%nest%ic = 6
+!$acc end parallel
+
+!$acc exit data copyout(myvar%nest)
+
+  if (myvar%a .ne. 0) stop 1
+  if (myvar%b .ne. 0) stop 2
+  if (myvar%c .ne. 0) stop 3
+  if (myvar%nest%ia .ne. 4) stop 4
+  if (myvar%nest%ib .ne. 5) stop 5
+  if (myvar%nest%ic .ne. 6) stop 6
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/multidim-slice.f95 b/libgomp/testsuite/libgomp.oacc-fortran/multidim-slice.f95
new file mode 100644
index 00000000000..a9b40eeab4c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/multidim-slice.f95
@@ -0,0 +1,50 @@ 
+! { dg-do run }
+
+program main
+  implicit none
+  real, allocatable :: myarr(:,:,:,:,:)
+  integer i, j, k, l, m
+
+  allocate(myarr(1:10,1:10,1:10,1:10,1:10))
+
+  do i=1,10
+    do j=1,10
+      do k=1,10
+        do l=1,10
+          do m=1,10
+            myarr(m,l,k,j,i) = i+j+k+l+m
+          end do
+        end do
+      end do
+    end do
+  end do
+
+  do i=1,10
+    !$acc data copy(myarr(:,:,:,:,i))
+    !$acc parallel loop collapse(4) present(myarr(:,:,:,:,i))
+    do j=1,10
+      do k=1,10
+        do l=1,10
+          do m=1,10
+            myarr(m,l,k,j,i) = myarr(m,l,k,j,i) + 1
+          end do
+        end do
+      end do
+    end do
+    !$acc end parallel loop
+    !$acc end data
+  end do
+
+  do i=1,10
+    do j=1,10
+      do k=1,10
+        do l=1,10
+          do m=1,10
+            if (myarr(m,l,k,j,i) .ne. i+j+k+l+m+1) stop 1
+          end do
+        end do
+      end do
+    end do
+  end do
+
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90
new file mode 100644
index 00000000000..c3c8a07868f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90
@@ -0,0 +1,284 @@ 
+! Test ACC UPDATE with derived types.
+
+! { dg-do run }
+
+module dt
+  integer, parameter :: n = 10
+  type inner
+     integer :: d(n)
+  end type inner
+  type mytype
+     integer(8) :: a, b, c(n)
+     type(inner) :: in
+  end type mytype
+end module dt
+
+program derived_acc
+  use dt
+
+  implicit none
+  integer i, res
+  type(mytype) :: var
+
+  var%a = 0
+  var%b = 1
+  var%c(:) = 10
+  var%in%d(:) = 100
+
+  var%c(:) = 10
+
+  !$acc enter data copyin(var)
+
+  !$acc parallel loop present(var)
+  do i = 1, 1
+     var%a = var%b
+  end do
+  !$acc end parallel loop
+
+  !$acc update host(var%a)
+
+  if (var%a /= var%b) stop 1
+
+  var%b = 100
+
+  !$acc update device(var%b)
+
+  !$acc parallel loop present(var)
+  do i = 1, 1
+     var%a = var%b
+  end do
+  !$acc end parallel loop
+
+  !$acc update host(var%a)
+
+  if (var%a /= var%b) stop 2
+
+  !$acc parallel loop present (var)
+  do i = 1, n
+     var%c(i) = i
+  end do
+  !$acc end parallel loop
+
+  !$acc update host(var%c)
+
+  var%a = -1
+
+  do i = 1, n
+     if (var%c(i) /= i) stop 3
+     var%c(i) = var%a
+  end do
+
+  !$acc update device(var%a)
+  !$acc update device(var%c)
+
+  res = 0
+
+  !$acc parallel loop present(var) reduction(+:res)
+  do i = 1, n
+     if (var%c(i) /= var%a) res = res + 1
+  end do
+
+  if (res /= 0) stop 4
+
+  var%c(:) = 0
+
+  !$acc update device(var%c)
+
+  !$acc parallel loop present(var)
+  do i = 5, 5
+     var%c(i) = 1
+  end do
+  !$acc end parallel loop
+
+  !$acc update host(var%c(5))
+
+  do i = 1, n
+     if (i /= 5 .and. var%c(i) /= 0) stop 5
+     if (i == 5 .and. var%c(i) /= 1) stop 6
+  end do
+
+  !$acc parallel loop present(var)
+  do i = 1, n
+     var%in%d = var%a
+  end do
+  !$acc end parallel loop
+
+  !$acc update host(var%in%d)
+
+  do i = 1, n
+     if (var%in%d(i) /= var%a) stop 7
+  end do
+
+  var%c(:) = 0
+
+  !$acc update device(var%c)
+
+  var%c(:) = -1
+
+  !$acc parallel loop present(var)
+  do i = n/2, n
+     var%c(i) = i
+  end do
+  !$acc end parallel loop
+
+  !$acc update host(var%c(n/2:n))
+
+  do i = 1,n
+     if (i < n/2 .and. var%c(i) /= -1) stop 8
+     if (i >= n/2 .and. var%c(i) /= i) stop 9
+  end do
+
+  var%in%d(:) = 0
+  !$acc update device(var%in%d)
+
+  !$acc parallel loop present(var)
+  do i = 5, 5
+     var%in%d(i) = 1
+  end do
+  !$acc end parallel loop
+
+  !$acc update host(var%in%d(5))
+
+  do i = 1, n
+     if (i /= 5 .and. var%in%d(i) /= 0) stop 10
+     if (i == 5 .and. var%in%d(i) /= 1) stop 11
+  end do
+
+  !$acc exit data delete(var)
+
+  call derived_acc_subroutine(var)
+end program derived_acc
+
+subroutine derived_acc_subroutine(var)
+  use dt
+
+  implicit none
+  integer i, res
+  type(mytype) :: var
+
+  var%a = 0
+  var%b = 1
+  var%c(:) = 10
+  var%in%d(:) = 100
+
+  var%c(:) = 10
+
+  !$acc enter data copyin(var)
+
+  !$acc parallel loop present(var)
+  do i = 1, 1
+     var%a = var%b
+  end do
+  !$acc end parallel loop
+
+  !$acc update host(var%a)
+
+  if (var%a /= var%b) stop 12
+
+  var%b = 100
+
+  !$acc update device(var%b)
+
+  !$acc parallel loop present(var)
+  do i = 1, 1
+     var%a = var%b
+  end do
+  !$acc end parallel loop
+
+  !$acc update host(var%a)
+
+  if (var%a /= var%b) stop 13
+
+  !$acc parallel loop present (var)
+  do i = 1, n
+     var%c(i) = i
+  end do
+  !$acc end parallel loop
+
+  !$acc update host(var%c)
+
+  var%a = -1
+
+  do i = 1, n
+     if (var%c(i) /= i) stop 14
+     var%c(i) = var%a
+  end do
+
+  !$acc update device(var%a)
+  !$acc update device(var%c)
+
+  res = 0
+
+  !$acc parallel loop present(var) reduction(+:res)
+  do i = 1, n
+     if (var%c(i) /= var%a) res = res + 1
+  end do
+
+  if (res /= 0) stop 15
+
+  var%c(:) = 0
+
+  !$acc update device(var%c)
+
+  !$acc parallel loop present(var)
+  do i = 5, 5
+     var%c(i) = 1
+  end do
+  !$acc end parallel loop
+
+  !$acc update host(var%c(5))
+
+  do i = 1, n
+     if (i /= 5 .and. var%c(i) /= 0) stop 16
+     if (i == 5 .and. var%c(i) /= 1) stop 17
+  end do
+
+  !$acc parallel loop present(var)
+  do i = 1, n
+     var%in%d = var%a
+  end do
+  !$acc end parallel loop
+
+  !$acc update host(var%in%d)
+
+  do i = 1, n
+     if (var%in%d(i) /= var%a) stop 18
+  end do
+
+  var%c(:) = 0
+
+  !$acc update device(var%c)
+
+  var%c(:) = -1
+
+  !$acc parallel loop present(var)
+  do i = n/2, n
+     var%c(i) = i
+  end do
+  !$acc end parallel loop
+
+  !$acc update host(var%c(n/2:n))
+
+  do i = 1,n
+     if (i < n/2 .and. var%c(i) /= -1) stop 19
+     if (i >= n/2 .and. var%c(i) /= i) stop 20
+  end do
+
+  var%in%d(:) = 0
+  !$acc update device(var%in%d)
+
+  !$acc parallel loop present(var)
+  do i = 5, 5
+     var%in%d(i) = 1
+  end do
+  !$acc end parallel loop
+
+  !$acc update host(var%in%d(5))
+
+  do i = 1, n
+     if (i /= 5 .and. var%in%d(i) /= 0) stop 21
+     if (i == 5 .and. var%in%d(i) /= 1) stop 22
+  end do
+
+  !$acc exit data delete(var)
+end subroutine derived_acc_subroutine