[PATCH 09/13] OpenACC 2.6 deep copy: C and C++ front-end parts
Jason Merrill
jason@redhat.com
Thu Dec 26 19:00:00 GMT 2019
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])
> +}
>
More information about the Fortran
mailing list