From: Jason Merrill <jason@redhat.com>
To: Julian Brown <julian@codesourcery.com>, gcc-patches@gcc.gnu.org
Cc: Thomas Schwinge <thomas@codesourcery.com>,
Jakub Jelinek <jakub@redhat.com>,
Tobias Burnus <tobias@codesourcery.com>,
Catherine_Moore@mentor.com, fortran@gcc.gnu.org
Subject: Re: [PATCH 09/13] OpenACC 2.6 deep copy: C and C++ front-end parts
Date: Thu, 26 Dec 2019 19:04:00 -0000 [thread overview]
Message-ID: <21c2d1d0-de2f-8752-8683-d756afefe218@redhat.com> (raw)
In-Reply-To: <5dfedf23eedbf91e58142a2bc853922c8502cba4.1576648001.git.julian@codesourcery.com>
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])
> +}
>
next prev parent reply other threads:[~2019-12-26 19:00 UTC|newest]
Thread overview: 81+ messages / expand[flat|nested] mbox.gz Atom feed top
2018-11-10 17:11 [PATCH 0/3] OpenACC 2.6 manual deep copy support (attach/detach) Julian Brown
2018-11-10 17:11 ` [PATCH 2/3] Factor out duplicate code in gimplify_scan_omp_clauses Julian Brown
2018-12-18 14:16 ` Julian Brown
2018-12-18 14:50 ` Jakub Jelinek
2018-11-10 17:11 ` [PATCH 1/3] Host-to-device transfer coalescing & magic offset value self-documentation Julian Brown
2018-12-21 10:56 ` libgomp/target.c magic constants self-documentation Thomas Schwinge
2019-05-29 14:48 ` Thomas Schwinge
2018-11-10 17:12 ` [PATCH 3/3] OpenACC 2.6 manual deep copy support (attach/detach) Julian Brown
2018-11-11 17:04 ` Bernhard Reutner-Fischer
2018-11-30 11:41 ` [PATCH] " Julian Brown
2018-12-03 17:03 ` Julian Brown
2018-12-07 13:50 ` Jakub Jelinek
2018-12-10 19:42 ` Julian Brown
2018-12-13 10:57 ` Jakub Jelinek
2018-12-14 19:00 ` Julian Brown
2018-12-18 12:25 ` Jakub Jelinek
2018-12-22 13:37 ` Thomas Schwinge
2019-10-18 17:20 ` Thomas Schwinge
2019-11-06 18:44 ` Julian Brown
2019-11-22 23:54 ` Julian Brown
2019-11-25 10:53 ` Tobias Burnus
2019-11-26 2:54 ` Julian Brown
2019-12-17 12:16 ` Thomas Schwinge
2019-12-17 17:28 ` [WIP] OpenACC 'acc_attach*', 'acc_detach*' runtime library routines (was: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)) Thomas Schwinge
2019-12-18 6:03 ` [PATCH 00/13] OpenACC 2.6 manual deep copy support Julian Brown
2019-12-18 6:03 ` [PATCH 02/13] OpenACC reference count overhaul Julian Brown
2020-05-19 15:42 ` Thomas Schwinge
2020-06-04 18:13 ` [OpenACC] Use 'tgt' returned from 'gomp_map_vars' (was: [PATCH 02/13] OpenACC reference count overhaul) Thomas Schwinge
2020-05-19 15:49 ` [PATCH 02/13] OpenACC reference count overhaul Thomas Schwinge
2020-05-19 15:58 ` Thomas Schwinge
2020-06-25 11:03 ` Thomas Schwinge
2020-07-03 15:29 ` Thomas Schwinge
2019-12-18 6:03 ` [PATCH 03/13] OpenACC reference count consistency checking Julian Brown
2019-12-18 6:03 ` [PATCH 01/13] Use aux struct in libgomp for infrequently-used/API-specific data Julian Brown
2019-12-18 6:04 ` [PATCH 05/13] Factor out duplicate code in gimplify_scan_omp_clauses Julian Brown
2019-12-18 6:04 ` [PATCH 04/13] Use gomp_map_val for OpenACC host-to-device address translation Julian Brown
2019-12-18 6:04 ` [PATCH 09/13] OpenACC 2.6 deep copy: C and C++ front-end parts Julian Brown
2019-12-24 5:05 ` Thomas Schwinge
2019-12-26 19:04 ` Jason Merrill [this message]
2021-06-10 11:03 ` Thomas Schwinge
2019-12-18 6:04 ` [PATCH 06/13] OpenACC 2.6 deep copy: attach/detach API routines Julian Brown
2019-12-18 6:04 ` [PATCH 08/13] OpenACC 2.6 deep copy: middle-end parts Julian Brown
2019-12-21 21:51 ` Thomas Schwinge
2019-12-18 6:05 ` [PATCH 12/13] OpenACC 2.6 deep copy: Fortran execution tests Julian Brown
2019-12-18 6:05 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Julian Brown
2019-12-21 23:37 ` Thomas Schwinge
2020-01-03 12:26 ` Julian Brown
2020-05-20 9:37 ` Thomas Schwinge
2020-06-05 16:23 ` [OpenACC 'exit data'] Simplify 'GOMP_MAP_STRUCT' handling (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts) Thomas Schwinge
2020-06-05 16:36 ` [OpenACC 'exit data'] Strip 'GOMP_MAP_STRUCT' mappings " Thomas Schwinge
2020-05-20 14:52 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Thomas Schwinge
2020-05-20 19:11 ` Julian Brown
2020-06-04 18:35 ` [OpenACC] Repair/restore 'is_tgt_unmapped' checking (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts) Thomas Schwinge
2020-06-04 18:53 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Thomas Schwinge
2020-06-05 10:39 ` Thomas Schwinge
2020-06-05 20:28 ` Julian Brown
2020-06-05 11:17 ` Thomas Schwinge
2020-06-05 20:31 ` Julian Brown
2020-06-09 10:41 ` OpenACC 'attach'/'detach' has no business affecting user-visible reference counting (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts) Thomas Schwinge
2020-06-09 12:23 ` Julian Brown
2020-06-18 18:21 ` Julian Brown
2020-07-16 8:35 ` OpenACC 'attach'/'detach' has no business affecting user-visible reference counting Thomas Schwinge
2020-06-26 9:20 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Thomas Schwinge
2020-07-16 9:35 ` Thomas Schwinge
2020-07-16 21:21 ` Julian Brown
2020-07-17 9:12 ` Thomas Schwinge
2020-06-30 15:58 ` Thomas Schwinge
2019-12-18 6:05 ` [PATCH 13/13] Fortran polymorphic class-type support for OpenACC Julian Brown
2019-12-18 6:05 ` [PATCH 11/13] OpenACC 2.6 deep copy: C and C++ execution tests Julian Brown
2020-06-04 18:43 ` Fix 'sizeof' usage in 'libgomp.oacc-c-c++-common/deep-copy-{7, 8}.c' (was: [PATCH 11/13] OpenACC 2.6 deep copy: C and C++ execution tests) Thomas Schwinge
2023-10-31 14:00 ` Add OpenACC 'acc_map_data' variant to 'libgomp.oacc-c-c++-common/deep-copy-8.c' " Thomas Schwinge
2019-12-18 7:20 ` [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts Julian Brown
2019-12-18 23:30 ` Tobias Burnus
2019-12-20 12:25 ` [committed] Improve is-coindexed check for OpenACC/OpenMP (was: [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts) Tobias Burnus
2019-12-20 13:25 ` [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts Tobias Burnus
2019-12-20 10:08 ` [patch,committed] Fix testsuite-fallout of OpenACC deep-copy patch (was: [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts) Tobias Burnus
2019-12-18 18:24 ` [PATCH 00/13] OpenACC 2.6 manual deep copy support Thomas Schwinge
2019-12-20 1:21 ` Julian Brown
2019-12-20 14:36 ` OpenACC regression and development pace Thomas Koenig
2020-06-04 18:07 ` [OpenACC] XFAIL behavior of over-eager 'finalize' clause (was: [PATCH 00/13] OpenACC 2.6 manual deep copy support) Thomas Schwinge
2019-12-17 16:53 ` In 'libgomp/target.c', 'struct splay_tree_key_s', use 'struct splay_tree_aux' for infrequently-used or API-specific data (was: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)) Thomas Schwinge
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=21c2d1d0-de2f-8752-8683-d756afefe218@redhat.com \
--to=jason@redhat.com \
--cc=Catherine_Moore@mentor.com \
--cc=fortran@gcc.gnu.org \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
--cc=julian@codesourcery.com \
--cc=thomas@codesourcery.com \
--cc=tobias@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).