[09/13] OpenACC 2.6 deep copy: C and C++ front-end parts

This patch has been broken out of the "OpenACC 2.6 manual deep copy
support" patch, last posted here:

  https://gcc.gnu.org/ml/gcc-patches/2019-11/msg02376.html

This part contains the C and C++ changes to parse attach and detach
clauses and struct member accesses via "." or "->" on other data-movement
clauses (copyin, copyout, etc.).

Tested alongside other patches in this series with offloading to
NVPTX. OK?

Thanks,

Julian

ChangeLog

	gcc/c-family/
	* c-common.h (c_omp_map_clause_name): Add prototype.
	* c-omp.c (c_omp_map_clause_name): New function.
	* 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,
	OACC_SERIAL_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,
	OACC_SERIAL_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/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.
---
 gcc/c-family/c-common.h                       |  1 +
 gcc/c-family/c-omp.c                          | 33 +++++++
 gcc/c-family/c-pragma.h                       |  2 +
 gcc/c/c-parser.c                              | 53 ++++++++--
 gcc/c/c-typeck.c                              | 76 +++++++++++++-
 gcc/cp/parser.c                               | 56 +++++++++--
 gcc/cp/semantics.c                            | 98 ++++++++++++++++---
 .../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 +++++++++++++
 11 files changed, 554 insertions(+), 34 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

Hi!

On 2019-12-17T22:03:49-0800, Julian Brown <julian@codesourcery.com> wrote:
> --- /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;

These data types...

> +#pragma acc enter data copyin(s)
> +[...]

..., and these uses...

> +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
> +[...]

..., and these tree dump scanning directives don't match up: a lot of
FAILs for anything that doesn't use 64-bit pointers, such as x86_64
GNU/Linux's '-m32' multilib.  This will need further tweaking to enable
tree dump scanning for all configurations, but for now, see attached
"Restrict 'c-c++-common/goacc/mdc-1.c' to LP64, LLP64"; committed to
trunk in r279720.


Grüße
 Thomas

On 12/18/19 1:03 AM, Julian Brown wrote:
> This patch has been broken out of the "OpenACC 2.6 manual deep copy
> support" patch, last posted here:
> 
>    https://gcc.gnu.org/ml/gcc-patches/2019-11/msg02376.html
> 
> This part contains the C and C++ changes to parse attach and detach
> clauses and struct member accesses via "." or "->" on other data-movement
> clauses (copyin, copyout, etc.).
> 
> Tested alongside other patches in this series with offloading to
> NVPTX. OK?

The C++ changes look fine to me if they make sense to Jakub.

> Thanks,
> 
> Julian
> 
> ChangeLog
> 
> 	gcc/c-family/
> 	* c-common.h (c_omp_map_clause_name): Add prototype.
> 	* c-omp.c (c_omp_map_clause_name): New function.
> 	* 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,
> 	OACC_SERIAL_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,
> 	OACC_SERIAL_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/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.
> ---
>   gcc/c-family/c-common.h                       |  1 +
>   gcc/c-family/c-omp.c                          | 33 +++++++
>   gcc/c-family/c-pragma.h                       |  2 +
>   gcc/c/c-parser.c                              | 53 ++++++++--
>   gcc/c/c-typeck.c                              | 76 +++++++++++++-
>   gcc/cp/parser.c                               | 56 +++++++++--
>   gcc/cp/semantics.c                            | 98 ++++++++++++++++---
>   .../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 +++++++++++++
>   11 files changed, 554 insertions(+), 34 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
> 
> diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
> index 2bcb54f66b9..2d89451b693 100644
> --- a/gcc/c-family/c-common.h
> +++ b/gcc/c-family/c-common.h
> @@ -1205,6 +1205,7 @@ extern bool c_omp_predefined_variable (tree);
>   extern enum omp_clause_default_kind c_omp_predetermined_sharing (tree);
>   extern tree c_omp_check_context_selector (location_t, tree);
>   extern void c_omp_mark_declare_variant (location_t, tree, tree);
> +extern const char *c_omp_map_clause_name (tree, bool);
>   
>   /* Return next tree in the chain for chain_next walking of tree nodes.  */
>   static inline tree
> diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
> index a4be2d68b9a..04f2c0b0682 100644
> --- a/gcc/c-family/c-omp.c
> +++ b/gcc/c-family/c-omp.c
> @@ -2259,3 +2259,36 @@ c_omp_mark_declare_variant (location_t loc, tree variant, tree construct)
>       error_at (loc, "%qD used as a variant with incompatible %<construct%> "
>   		   "selector sets", variant);
>   }
> +
> +/* For OpenACC, the OMP_CLAUSE_MAP_KIND of an OMP_CLAUSE_MAP is used internally
> +   to distinguish clauses as seen by the user.  Return the "friendly" clause
> +   name for error messages etc., where possible.  See also
> +   c/c-parser.c:c_parser_oacc_data_clause and
> +   cp/parser.c:cp_parser_oacc_data_clause.  */
> +
> +const char *
> +c_omp_map_clause_name (tree clause, bool oacc)
> +{
> +  if (oacc && OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP)
> +    switch (OMP_CLAUSE_MAP_KIND (clause))
> +    {
> +    case GOMP_MAP_FORCE_ALLOC:
> +    case GOMP_MAP_ALLOC: return "create";
> +    case GOMP_MAP_FORCE_TO:
> +    case GOMP_MAP_TO: return "copyin";
> +    case GOMP_MAP_FORCE_FROM:
> +    case GOMP_MAP_FROM: return "copyout";
> +    case GOMP_MAP_FORCE_TOFROM:
> +    case GOMP_MAP_TOFROM: return "copy";
> +    case GOMP_MAP_RELEASE: return "delete";
> +    case GOMP_MAP_FORCE_PRESENT: return "present";
> +    case GOMP_MAP_ATTACH: return "attach";
> +    case GOMP_MAP_FORCE_DETACH:
> +    case GOMP_MAP_DETACH: return "detach";
> +    case GOMP_MAP_DEVICE_RESIDENT: return "device_resident";
> +    case GOMP_MAP_LINK: return "link";
> +    case GOMP_MAP_FORCE_DEVICEPTR: return "deviceptr";
> +    default: break;
> +    }
> +  return omp_clause_code_name[OMP_CLAUSE_CODE (clause)];
> +}
> diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
> index bfe681bb430..8a04e611bc7 100644
> --- a/gcc/c-family/c-pragma.h
> +++ b/gcc/c-family/c-pragma.h
> @@ -143,11 +143,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 bfe56998996..3839636f6ef 100644
> --- a/gcc/c/c-parser.c
> +++ b/gcc/c/c-parser.c
> @@ -12564,6 +12564,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))
> @@ -12590,6 +12592,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))
> @@ -12833,12 +12837,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;
> @@ -12965,9 +12973,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))
>   		    {
> @@ -13089,11 +13101,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;
> @@ -13101,18 +13114,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
> @@ -13122,6 +13137,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;
> @@ -13137,6 +13155,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;
> @@ -13156,7 +13177,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);
> @@ -15871,6 +15892,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";
> @@ -15899,6 +15924,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";
> @@ -16409,7 +16438,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)		\
> @@ -16592,6 +16622,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) )
> @@ -16601,6 +16632,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) )
>   
> @@ -16740,6 +16772,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)		\
> @@ -16755,6 +16788,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)		\
> @@ -16773,6 +16807,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
>   
>   #define OACC_SERIAL_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 36aedc063d2..db03b3c97d4 100644
> --- a/gcc/c/c-typeck.c
> +++ b/gcc/c/c-typeck.c
> @@ -12897,7 +12897,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))
> @@ -12918,6 +12917,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)
> @@ -13003,7 +13011,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))
> @@ -13443,7 +13462,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
> @@ -13680,6 +13703,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.  */
>   
> @@ -14433,6 +14485,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)
> @@ -14440,8 +14494,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)))
> @@ -14476,6 +14535,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 16d1359c47d..c7aa071088d 100644
> --- a/gcc/cp/parser.c
> +++ b/gcc/cp/parser.c
> @@ -33124,6 +33124,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))
> @@ -33148,6 +33150,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))
> @@ -33350,11 +33354,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;
> @@ -33435,15 +33443,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);
>   		}
> @@ -33561,19 +33574,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
> @@ -33583,6 +33600,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;
> @@ -33598,6 +33618,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;
> @@ -33617,7 +33640,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);
> @@ -36095,6 +36118,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";
> @@ -36123,6 +36150,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";
> @@ -39971,10 +40002,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) )
> @@ -40174,6 +40207,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)		\
> @@ -40184,6 +40218,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) )
>   
> @@ -40291,6 +40326,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)		\
> @@ -40306,6 +40342,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)		\
> @@ -40324,6 +40361,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
>   
>   #define OACC_SERIAL_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 42611682549..dec22494cd9 100644
> --- a/gcc/cp/semantics.c
> +++ b/gcc/cp/semantics.c
> @@ -4740,7 +4740,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)
> @@ -4764,6 +4763,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);
> @@ -4863,6 +4864,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))
> @@ -5310,12 +5323,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);
> @@ -6238,6 +6257,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.  */
>   
> @@ -6502,7 +6556,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)))
>   	    {
> @@ -6572,8 +6626,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)
> @@ -6640,7 +6693,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%>"
> @@ -7129,7 +7182,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%>"
> @@ -7261,6 +7314,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)
> @@ -7268,14 +7323,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))
> @@ -7325,7 +7391,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),
> @@ -7407,7 +7474,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),
> @@ -7462,7 +7531,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/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])
> +}
>

Hi!

While working on something else...  ;-)

On 2019-12-17T22:03:49-0800, Julian Brown <julian@codesourcery.com> wrote:
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
> @@ -0,0 +1,55 @@
> +[...]
> +#pragma acc acc enter data attach(s.e)
> +[...]

... I noticed 'acc acc'.

> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/mdc-2.c

Similar.

> --- /dev/null
> +++ b/gcc/testsuite/g++.dg/goacc/mdc.C

Similar.

Pushed "Fix '#pragma acc acc [...]' typos" to master branch in
commit 05c4dabb71476ddea8d409fd41f1e97d62d0b5f4, see attached.


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf