This patch has been broken out of the "OpenACC 2.6 manual deep copy support" patch, last posted here:
https://gcc.gnu.org/ml/gcc-patches/2019-11/msg02376.html This part contains the C and C++ changes to parse attach and detach clauses and struct member accesses via "." or "->" on other data-movement clauses (copyin, copyout, etc.). Tested alongside other patches in this series with offloading to NVPTX. OK? Thanks, Julian ChangeLog gcc/c-family/ * c-common.h (c_omp_map_clause_name): Add prototype. * c-omp.c (c_omp_map_clause_name): New function. * c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_ATTACH and PRAGMA_OACC_CLAUSE_DETACH. gcc/c/ * c-parser.c (c_parser_omp_clause_name): Add parsing of attach and detach clauses. (c_parser_omp_variable_list): Add ALLOW_DEREF optional parameter. Allow deref (->) in variable lists if true. (c_parser_omp_var_list_parens): Add ALLOW_DEREF optional parameter. Pass to c_parser_omp_variable_list. (c_parser_oacc_data_clause): Support attach and detach clauses. Update call to c_parser_omp_variable_list. (c_parser_oacc_all_clauses): Support attach and detach clauses. (OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK, OACC_SERIAL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_ATTACH. (OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH. * c-typeck.c (handle_omp_array_sections_1): Reject subarrays for attach and detach. Support deref. (handle_omp_array_sections): Use GOMP_MAP_ATTACH_DETACH instead of GOMP_MAP_ALWAYS_POINTER for OpenACC. (c_oacc_check_attachments): New function. (c_finish_omp_clauses): Check attach/detach arguments for being pointers using above. Support deref. gcc/cp/ * parser.c (cp_parser_omp_clause_name): Support attach and detach clauses. (cp_parser_omp_var_list_no_open): Add ALLOW_DEREF optional parameter. Parse deref if true. (cp_parser_omp_var_list): Add ALLOW_DEREF optional parameter. Pass to cp_parser_omp_var_list_no_open. (cp_parser_oacc_data_clause): Support attach and detach clauses. Update call to cp_parser_omp_var_list_no_open. (cp_parser_oacc_all_clauses): Support attach and detach. (OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK, OACC_SERIAL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_ATTACH. (OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH. * semantics.c (handle_omp_array_sections_1): Reject subarrays for attach and detach. (handle_omp_array_sections): Use GOMP_MAP_ATTACH_DETACH instead of GOMP_MAP_ALWAYS_POINTER for OpenACC. (cp_oacc_check_attachments): New function. (finish_omp_clauses): Use above function. Allow structure fields and class members to appear in OpenACC data clauses. Support GOMP_MAP_ATTACH_DETACH. Support deref. gcc/testsuite/ * c-c++-common/goacc/deep-copy-arrayofstruct.c: New test. * c-c++-common/goacc/mdc-1.c: New test. * c-c++-common/goacc/mdc-2.c: New test. * gcc.dg/goacc/mdc.C: New test. --- gcc/c-family/c-common.h | 1 + gcc/c-family/c-omp.c | 33 +++++++ gcc/c-family/c-pragma.h | 2 + gcc/c/c-parser.c | 53 ++++++++-- gcc/c/c-typeck.c | 76 +++++++++++++- gcc/cp/parser.c | 56 +++++++++-- gcc/cp/semantics.c | 98 ++++++++++++++++--- .../goacc/deep-copy-arrayofstruct.c | 84 ++++++++++++++++ gcc/testsuite/c-c++-common/goacc/mdc-1.c | 55 +++++++++++ gcc/testsuite/c-c++-common/goacc/mdc-2.c | 62 ++++++++++++ gcc/testsuite/g++.dg/goacc/mdc.C | 68 +++++++++++++ 11 files changed, 554 insertions(+), 34 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c create mode 100644 gcc/testsuite/c-c++-common/goacc/mdc-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/mdc-2.c create mode 100644 gcc/testsuite/g++.dg/goacc/mdc.C 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]) +} -- 2.23.0
