diff options
77 files changed, 1420 insertions, 619 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 7b399f1..37fc7da 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,31 @@ +2018-06-20 Chung-Lin Tang <cltang@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + Cesar Philippidis <cesar@codesourcery.com> + + * gimplify.c (gimplify_scan_omp_clauses): Add support for + OMP_CLAUSE_{IF_PRESENT,FINALIZE}. + (gimplify_adjust_omp_clauses): Likewise. + (gimplify_oacc_declare_1): Add support for GOMP_MAP_RELEASE, remove + support for GOMP_MAP_FORCE_{ALLOC,TO,FROM,TOFROM}. + (gimplify_omp_target_update): Update handling of acc update and + enter/exit data. + * omp-low.c (install_var_field): Remove unused parameter + base_pointers_restrict. + (scan_sharing_clauses): Remove base_pointers_restrict parameter. + Update call to install_var_field. Handle OMP_CLAUSE_{IF_PRESENT, + FINALIZE} + (omp_target_base_pointers_restrict_p): Delete. + (scan_omp_target): Update call to scan_sharing_clauses. + * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_{IF_PRESENT, + FINALIZE}. + * tree-nested.c (convert_nonlocal_omp_clauses): Handle + OMP_CLAUSE_{IF_PRESENT,FINALIZE}. + (convert_local_omp_clauses): Likewise. + * tree-pretty-print.c (dump_omp_clause): Likewise. + * tree.c (omp_clause_num_ops): Add entries for OMP_CLAUSE_{IF_PRESENT, + FINALIZE}. + (omp_clause_code_name): Likewise. + 2018-06-20 Jakub Jelinek <jakub@redhat.com> PR debug/86194 diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog index 6b974da..fc7f7de 100644 --- a/gcc/c-family/ChangeLog +++ b/gcc/c-family/ChangeLog @@ -1,3 +1,11 @@ +2018-06-20 Chung-Lin Tang <cltang@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + Cesar Philippidis <cesar@codesourcery.com> + + * c-pragma.h (enum pragma_omp_clause): Add + PRAGMA_OACC_CLAUSE_{FINALIZE,IF_PRESENT}. Remove + PRAGMA_OACC_CLAUSE_PRESENT_OR_{COPY,COPYIN,COPYOUT,CREATE}. + 2018-06-20 Jakub Jelinek <jakub@redhat.com> PR c++/86210 diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index c70380c..b322547b 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -138,16 +138,13 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_DELETE, PRAGMA_OACC_CLAUSE_DEVICEPTR, PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT, + PRAGMA_OACC_CLAUSE_FINALIZE, PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_HOST, PRAGMA_OACC_CLAUSE_INDEPENDENT, PRAGMA_OACC_CLAUSE_NUM_GANGS, PRAGMA_OACC_CLAUSE_NUM_WORKERS, PRAGMA_OACC_CLAUSE_PRESENT, - PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY, - PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN, - PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT, - PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE, PRAGMA_OACC_CLAUSE_SELF, PRAGMA_OACC_CLAUSE_SEQ, PRAGMA_OACC_CLAUSE_TILE, @@ -156,6 +153,7 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_VECTOR_LENGTH, PRAGMA_OACC_CLAUSE_WAIT, PRAGMA_OACC_CLAUSE_WORKER, + PRAGMA_OACC_CLAUSE_IF_PRESENT, PRAGMA_OACC_CLAUSE_COLLAPSE = PRAGMA_OMP_CLAUSE_COLLAPSE, PRAGMA_OACC_CLAUSE_COPYIN = PRAGMA_OMP_CLAUSE_COPYIN, PRAGMA_OACC_CLAUSE_DEVICE = PRAGMA_OMP_CLAUSE_DEVICE, diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index f1ff69e8..1ba1173 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,26 @@ +2018-06-20 Chung-Lin Tang <cltang@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + Cesar Philippidis <cesar@codesourcery.com> + + * c-parser.c (c_parser_omp_clause_name): Add support for finalize + and if_present. Make present_or_{copy,copyin,copyout,create} aliases + to their non-present_or_* counterparts. Make 'self' an alias to + PRAGMA_OACC_CLAUSE_HOST. + (c_parser_oacc_data_clause): Update GOMP mappings for + PRAGMA_OACC_CLAUSE_{COPY,COPYIN,COPYOUT,CREATE,DELETE}. Remove + PRAGMA_OACC_CLAUSE_{SELF,PRESENT_OR_*}. + (c_parser_oacc_all_clauses): Handle finalize and if_present clauses. + Remove support for present_or_* clauses. + (OACC_KERNELS_CLAUSE_MASK): Remove PRESENT_OR_* clauses. + (OACC_PARALLEL_CLAUSE_MASK): Likewise. + (OACC_DECLARE_CLAUSE_MASK): Likewise. + (OACC_DATA_CLAUSE_MASK): Likewise. + (OACC_ENTER_DATA_CLAUSE_MASK): Remove PRESENT_OR_* clauses. + (OACC_EXIT_DATA_CLAUSE_MASK): Add FINALIZE clause. + (OACC_UPDATE_CLAUSE_MASK): Remove SELF, add IF_PRESENT. + (c_parser_oacc_declare): Remove PRESENT_OR_* clauses. + * c-typeck.c (c_finish_omp_clauses): Handle IF_PRESENT and FINALIZE. + 2018-06-16 Kugan Vivekanandarajah <kuganv@linaro.org> * c-typeck.c (build_unary_op): Handle ABSU_EXPR; diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 6b41a61..7a92628 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -11259,6 +11259,8 @@ c_parser_omp_clause_name (c_parser *parser) case 'f': if (!strcmp ("final", p)) result = PRAGMA_OMP_CLAUSE_FINAL; + else if (!strcmp ("finalize", p)) + result = PRAGMA_OACC_CLAUSE_FINALIZE; else if (!strcmp ("firstprivate", p)) result = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE; else if (!strcmp ("from", p)) @@ -11277,7 +11279,9 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OACC_CLAUSE_HOST; break; case 'i': - if (!strcmp ("inbranch", p)) + if (!strcmp ("if_present", p)) + result = PRAGMA_OACC_CLAUSE_IF_PRESENT; + else if (!strcmp ("inbranch", p)) result = PRAGMA_OMP_CLAUSE_INBRANCH; else if (!strcmp ("independent", p)) result = PRAGMA_OACC_CLAUSE_INDEPENDENT; @@ -11325,18 +11329,20 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_PARALLEL; else if (!strcmp ("present", p)) result = PRAGMA_OACC_CLAUSE_PRESENT; + /* As of OpenACC 2.5, these are now aliases of the non-present_or + clauses. */ else if (!strcmp ("present_or_copy", p) || !strcmp ("pcopy", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY; + result = PRAGMA_OACC_CLAUSE_COPY; else if (!strcmp ("present_or_copyin", p) || !strcmp ("pcopyin", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN; + result = PRAGMA_OACC_CLAUSE_COPYIN; else if (!strcmp ("present_or_copyout", p) || !strcmp ("pcopyout", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT; + result = PRAGMA_OACC_CLAUSE_COPYOUT; else if (!strcmp ("present_or_create", p) || !strcmp ("pcreate", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE; + result = PRAGMA_OACC_CLAUSE_CREATE; else if (!strcmp ("priority", p)) result = PRAGMA_OMP_CLAUSE_PRIORITY; else if (!strcmp ("private", p)) @@ -11355,6 +11361,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_SCHEDULE; else if (!strcmp ("sections", p)) result = PRAGMA_OMP_CLAUSE_SECTIONS; + else if (!strcmp ("self", p)) /* "self" is a synonym for "host". */ + result = PRAGMA_OACC_CLAUSE_HOST; else if (!strcmp ("seq", p)) result = PRAGMA_OACC_CLAUSE_SEQ; else if (!strcmp ("shared", p)) @@ -11363,8 +11371,6 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_SIMD; else if (!strcmp ("simdlen", p)) result = PRAGMA_OMP_CLAUSE_SIMDLEN; - else if (!strcmp ("self", p)) - result = PRAGMA_OACC_CLAUSE_SELF; break; case 't': if (!strcmp ("taskgroup", p)) @@ -11646,15 +11652,7 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, copyout ( variable-list ) create ( variable-list ) delete ( variable-list ) - present ( variable-list ) - present_or_copy ( variable-list ) - pcopy ( variable-list ) - present_or_copyin ( variable-list ) - pcopyin ( variable-list ) - present_or_copyout ( variable-list ) - pcopyout ( variable-list ) - present_or_create ( variable-list ) - pcreate ( variable-list ) */ + present ( variable-list ) */ static tree c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, @@ -11664,19 +11662,19 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, switch (c_kind) { case PRAGMA_OACC_CLAUSE_COPY: - kind = GOMP_MAP_FORCE_TOFROM; + kind = GOMP_MAP_TOFROM; break; case PRAGMA_OACC_CLAUSE_COPYIN: - kind = GOMP_MAP_FORCE_TO; + kind = GOMP_MAP_TO; break; case PRAGMA_OACC_CLAUSE_COPYOUT: - kind = GOMP_MAP_FORCE_FROM; + kind = GOMP_MAP_FROM; break; case PRAGMA_OACC_CLAUSE_CREATE: - kind = GOMP_MAP_FORCE_ALLOC; + kind = GOMP_MAP_ALLOC; break; case PRAGMA_OACC_CLAUSE_DELETE: - kind = GOMP_MAP_DELETE; + kind = GOMP_MAP_RELEASE; break; case PRAGMA_OACC_CLAUSE_DEVICE: kind = GOMP_MAP_FORCE_TO; @@ -11685,7 +11683,6 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, kind = GOMP_MAP_DEVICE_RESIDENT; break; case PRAGMA_OACC_CLAUSE_HOST: - case PRAGMA_OACC_CLAUSE_SELF: kind = GOMP_MAP_FORCE_FROM; break; case PRAGMA_OACC_CLAUSE_LINK: @@ -11694,18 +11691,6 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY: - kind = GOMP_MAP_TOFROM; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN: - kind = GOMP_MAP_TO; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT: - kind = GOMP_MAP_FROM; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE: - kind = GOMP_MAP_ALLOC; - break; default: gcc_unreachable (); } @@ -12597,8 +12582,9 @@ c_parser_oacc_shape_clause (c_parser *parser, omp_clause_code kind, return list; } -/* OpenACC: +/* OpenACC 2.5: auto + finalize independent nohost seq */ @@ -13955,6 +13941,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "device_resident"; break; + case PRAGMA_OACC_CLAUSE_FINALIZE: + clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_FINALIZE, + clauses); + c_name = "finalize"; + break; case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE: clauses = c_parser_omp_clause_firstprivate (parser, clauses); c_name = "firstprivate"; @@ -13972,6 +13963,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_omp_clause_if (parser, clauses, false); c_name = "if"; break; + case PRAGMA_OACC_CLAUSE_IF_PRESENT: + clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_IF_PRESENT, + clauses); + c_name = "if_present"; + break; case PRAGMA_OACC_CLAUSE_INDEPENDENT: clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_INDEPENDENT, clauses); @@ -13997,22 +13993,6 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "present"; break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY: - clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_copy"; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN: - clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_copyin"; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT: - clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_copyout"; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE: - clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_create"; - break; case PRAGMA_OACC_CLAUSE_PRIVATE: clauses = c_parser_omp_clause_private (parser, clauses); c_name = "private"; @@ -14021,10 +14001,6 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_omp_clause_reduction (parser, clauses); c_name = "reduction"; break; - case PRAGMA_OACC_CLAUSE_SELF: - clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "self"; - break; case PRAGMA_OACC_CLAUSE_SEQ: clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ, clauses); @@ -14417,11 +14393,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) ) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)) static tree c_parser_oacc_data (location_t loc, c_parser *parser, bool *if_p) @@ -14451,11 +14423,7 @@ c_parser_oacc_data (location_t loc, c_parser *parser, bool *if_p) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) ) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)) static void c_parser_oacc_declare (c_parser *parser) @@ -14490,8 +14458,8 @@ c_parser_oacc_declare (c_parser *parser) switch (OMP_CLAUSE_MAP_KIND (t)) { case GOMP_MAP_FIRSTPRIVATE_POINTER: - case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_TO: + case GOMP_MAP_ALLOC: + case GOMP_MAP_TO: case GOMP_MAP_FORCE_DEVICEPTR: case GOMP_MAP_DEVICE_RESIDENT: break; @@ -14604,8 +14572,6 @@ c_parser_oacc_declare (c_parser *parser) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) #define OACC_EXIT_DATA_CLAUSE_MASK \ @@ -14613,6 +14579,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_FINALIZE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) static void @@ -14756,10 +14723,6 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) @@ -14777,10 +14740,6 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) @@ -15008,7 +14967,7 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_HOST) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF_PRESENT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) static void diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index aa70b23..90ae306 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13897,6 +13897,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: case OMP_CLAUSE_TILE: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: pc = &OMP_CLAUSE_CHAIN (c); continue; diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index 72cd388..56cecfa 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,27 @@ +2018-06-20 Chung-Lin Tang <cltang@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + Cesar Philippidis <cesar@codesourcery.com> + + * parser.c (cp_parser_omp_clause_name): Add support for finalize + and if_present. Make present_or_{copy,copyin,copyout,create} aliases + to their non-present_or_* counterparts. Make 'self' an alias to + PRAGMA_OACC_CLAUSE_HOST. + (cp_parser_oacc_data_clause): Update GOMP mappings for + PRAGMA_OACC_CLAUSE_{COPY,COPYIN,COPYOUT,CREATE,DELETE}. Remove + PRAGMA_OACC_CLAUSE_{SELF,PRESENT_OR_*}. + (cp_parser_oacc_all_clauses): Handle finalize and if_present clauses. + Remove support for present_or_* clauses. + (OACC_KERNELS_CLAUSE_MASK): Remove PRESENT_OR_* clauses. + (OACC_PARALLEL_CLAUSE_MASK): Likewise. + (OACC_DECLARE_CLAUSE_MASK): Likewise. + (OACC_DATA_CLAUSE_MASK): Likewise. + (OACC_ENTER_DATA_CLAUSE_MASK): Remove PRESENT_OR_* clauses. + (OACC_EXIT_DATA_CLAUSE_MASK): Add FINALIZE clause. + (OACC_UPDATE_CLAUSE_MASK): Remove SELF, add IF_PRESENT. + (cp_parser_oacc_declare): Remove PRESENT_OR_* clauses. + * pt.c (tsubst_omp_clauses): Handle IF_PRESENT and FINALIZE. + * semantics.c (finish_omp_clauses): Handle IF_PRESENT and FINALIZE. + 2018-06-20 Marek Polacek <polacek@redhat.com> PR c++/86240 diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index b618485..154729c 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -31375,6 +31375,8 @@ cp_parser_omp_clause_name (cp_parser *parser) case 'f': if (!strcmp ("final", p)) result = PRAGMA_OMP_CLAUSE_FINAL; + else if (!strcmp ("finalize", p)) + result = PRAGMA_OACC_CLAUSE_FINALIZE; else if (!strcmp ("firstprivate", p)) result = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE; else if (!strcmp ("from", p)) @@ -31393,7 +31395,9 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OACC_CLAUSE_HOST; break; case 'i': - if (!strcmp ("inbranch", p)) + if (!strcmp ("if_present", p)) + result = PRAGMA_OACC_CLAUSE_IF_PRESENT; + else if (!strcmp ("inbranch", p)) result = PRAGMA_OMP_CLAUSE_INBRANCH; else if (!strcmp ("independent", p)) result = PRAGMA_OACC_CLAUSE_INDEPENDENT; @@ -31443,16 +31447,16 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OACC_CLAUSE_PRESENT; else if (!strcmp ("present_or_copy", p) || !strcmp ("pcopy", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY; + result = PRAGMA_OACC_CLAUSE_COPY; else if (!strcmp ("present_or_copyin", p) || !strcmp ("pcopyin", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN; + result = PRAGMA_OACC_CLAUSE_COPYIN; else if (!strcmp ("present_or_copyout", p) || !strcmp ("pcopyout", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT; + result = PRAGMA_OACC_CLAUSE_COPYOUT; else if (!strcmp ("present_or_create", p) || !strcmp ("pcreate", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE; + result = PRAGMA_OACC_CLAUSE_CREATE; else if (!strcmp ("priority", p)) result = PRAGMA_OMP_CLAUSE_PRIORITY; else if (!strcmp ("proc_bind", p)) @@ -31469,8 +31473,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_SCHEDULE; else if (!strcmp ("sections", p)) result = PRAGMA_OMP_CLAUSE_SECTIONS; - else if (!strcmp ("self", p)) - result = PRAGMA_OACC_CLAUSE_SELF; + else if (!strcmp ("self", p)) /* "self" is a synonym for "host". */ + result = PRAGMA_OACC_CLAUSE_HOST; else if (!strcmp ("seq", p)) result = PRAGMA_OACC_CLAUSE_SEQ; else if (!strcmp ("shared", p)) @@ -31730,15 +31734,7 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list) copyout ( variable-list ) create ( variable-list ) delete ( variable-list ) - present ( variable-list ) - present_or_copy ( variable-list ) - pcopy ( variable-list ) - present_or_copyin ( variable-list ) - pcopyin ( variable-list ) - present_or_copyout ( variable-list ) - pcopyout ( variable-list ) - present_or_create ( variable-list ) - pcreate ( variable-list ) */ + present ( variable-list ) */ static tree cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, @@ -31748,19 +31744,19 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, switch (c_kind) { case PRAGMA_OACC_CLAUSE_COPY: - kind = GOMP_MAP_FORCE_TOFROM; + kind = GOMP_MAP_TOFROM; break; case PRAGMA_OACC_CLAUSE_COPYIN: - kind = GOMP_MAP_FORCE_TO; + kind = GOMP_MAP_TO; break; case PRAGMA_OACC_CLAUSE_COPYOUT: - kind = GOMP_MAP_FORCE_FROM; + kind = GOMP_MAP_FROM; break; case PRAGMA_OACC_CLAUSE_CREATE: - kind = GOMP_MAP_FORCE_ALLOC; + kind = GOMP_MAP_ALLOC; break; case PRAGMA_OACC_CLAUSE_DELETE: - kind = GOMP_MAP_DELETE; + kind = GOMP_MAP_RELEASE; break; case PRAGMA_OACC_CLAUSE_DEVICE: kind = GOMP_MAP_FORCE_TO; @@ -31769,7 +31765,6 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, kind = GOMP_MAP_DEVICE_RESIDENT; break; case PRAGMA_OACC_CLAUSE_HOST: - case PRAGMA_OACC_CLAUSE_SELF: kind = GOMP_MAP_FORCE_FROM; break; case PRAGMA_OACC_CLAUSE_LINK: @@ -31778,18 +31773,6 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY: - kind = GOMP_MAP_TOFROM; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN: - kind = GOMP_MAP_TO; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT: - kind = GOMP_MAP_FROM; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE: - kind = GOMP_MAP_ALLOC; - break; default: gcc_unreachable (); } @@ -31828,8 +31811,9 @@ cp_parser_oacc_data_clause_deviceptr (cp_parser *parser, tree list) return list; } -/* OpenACC 2.0: +/* OpenACC 2.5: auto + finalize independent nohost seq */ @@ -33794,6 +33778,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "device_resident"; break; + case PRAGMA_OACC_CLAUSE_FINALIZE: + clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_FINALIZE, + clauses, here); + c_name = "finalize"; + break; case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE: clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FIRSTPRIVATE, clauses); @@ -33812,6 +33801,12 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_omp_clause_if (parser, clauses, here, false); c_name = "if"; break; + case PRAGMA_OACC_CLAUSE_IF_PRESENT: + clauses = cp_parser_oacc_simple_clause (parser, + OMP_CLAUSE_IF_PRESENT, + clauses, here); + c_name = "if_present"; + break; case PRAGMA_OACC_CLAUSE_INDEPENDENT: clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_INDEPENDENT, @@ -33838,22 +33833,6 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "present"; break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY: - clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_copy"; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN: - clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_copyin"; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT: - clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_copyout"; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE: - clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_create"; - break; case PRAGMA_OACC_CLAUSE_PRIVATE: clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_PRIVATE, clauses); @@ -33863,10 +33842,6 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_omp_clause_reduction (parser, clauses); c_name = "reduction"; break; - case PRAGMA_OACC_CLAUSE_SELF: - clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "self"; - break; case PRAGMA_OACC_CLAUSE_SEQ: clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ, clauses, here); @@ -36802,11 +36777,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) ) static tree cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) @@ -36861,11 +36832,7 @@ cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) ) static tree cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) @@ -36898,8 +36865,8 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) switch (OMP_CLAUSE_MAP_KIND (t)) { case GOMP_MAP_FIRSTPRIVATE_POINTER: - case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_TO: + case GOMP_MAP_ALLOC: + case GOMP_MAP_TO: case GOMP_MAP_FORCE_DEVICEPTR: case GOMP_MAP_DEVICE_RESIDENT: break; @@ -37010,8 +36977,6 @@ 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_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) #define OACC_EXIT_DATA_CLAUSE_MASK \ @@ -37019,6 +36984,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_FINALIZE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) static tree @@ -37131,10 +37097,6 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) @@ -37151,10 +37113,6 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ @@ -37215,7 +37173,7 @@ cp_parser_oacc_kernels_parallel (cp_parser *parser, cp_token *pragma_tok, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_HOST) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF_PRESENT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT)) static tree diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index be42b20..c5433dc 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -16107,6 +16107,8 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort, case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: break; default: gcc_unreachable (); diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 38b7b66..bad712e 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -7091,6 +7091,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_AUTO: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: break; case OMP_CLAUSE_TILE: diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog index 0ea7a51..99a3114 100644 --- a/gcc/fortran/ChangeLog +++ b/gcc/fortran/ChangeLog @@ -1,3 +1,25 @@ +2018-06-20 Chung-Lin Tang <cltang@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + Cesar Philippidis <cesar@codesourcery.com> + + * gfortran.h (gfc_omp_clauses): Add unsigned if_present, finalize + bitfields. + * openmp.c (enum omp_mask2): Remove OMP_CLAUSE_PRESENT_OR_*. Add + OMP_CLAUSE_{IF_PRESENT,FINALIZE}. + (gfc_match_omp_clauses): Update handling of copy, copyin, copyout, + create, deviceptr, present_of_*. Add support for finalize and + if_present. + (OACC_PARALLEL_CLAUSES): Remove PRESENT_OR_* clauses. + (OACC_KERNELS_CLAUSES): Likewise. + (OACC_DATA_CLAUSES): Likewise. + (OACC_DECLARE_CLAUSES): Likewise. + (OACC_UPDATE_CLAUSES): Add IF_PRESENT clause. + (OACC_ENTER_DATA_CLAUSES): Remove PRESENT_OR_* clauses. + (OACC_EXIT_DATA_CLAUSES): Add FINALIZE clause. + (gfc_match_oacc_declare): Update to OpenACC 2.5 semantics. + * trans-openmp.c (gfc_trans_omp_clauses): Add support for IF_PRESENT + and FINALIZE. + 2018-06-18 Eric Botcazou <ebotcazou@adacore.com> * trans-decl.c (gfc_get_fake_result_decl): Revert latest change. diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 1d98d25..0b89f8d 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1344,6 +1344,7 @@ typedef struct gfc_omp_clauses gfc_expr_list *tile_list; unsigned async:1, gang:1, worker:1, vector:1, seq:1, independent:1; unsigned wait:1, par_auto:1, gang_static:1; + unsigned if_present:1, finalize:1; locus loc; } diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 97d6e78..94a7f7e 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -796,10 +796,6 @@ enum omp_mask2 OMP_CLAUSE_COPYOUT, OMP_CLAUSE_CREATE, OMP_CLAUSE_PRESENT, - OMP_CLAUSE_PRESENT_OR_COPY, - OMP_CLAUSE_PRESENT_OR_COPYIN, - OMP_CLAUSE_PRESENT_OR_COPYOUT, - OMP_CLAUSE_PRESENT_OR_CREATE, OMP_CLAUSE_DEVICEPTR, OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER, @@ -813,6 +809,8 @@ enum omp_mask2 OMP_CLAUSE_DELETE, OMP_CLAUSE_AUTO, OMP_CLAUSE_TILE, + OMP_CLAUSE_IF_PRESENT, + OMP_CLAUSE_FINALIZE, /* This must come last. */ OMP_MASK2_LAST }; @@ -1041,7 +1039,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_COPY) && gfc_match ("copy ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_TOFROM)) + OMP_MAP_TOFROM)) continue; if (mask & OMP_CLAUSE_COPYIN) { @@ -1049,7 +1047,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, { if (gfc_match ("copyin ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_TO)) + OMP_MAP_TO)) continue; } else if (gfc_match_omp_variable_list ("copyin (", @@ -1060,7 +1058,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_COPYOUT) && gfc_match ("copyout ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_FROM)) + OMP_MAP_FROM)) continue; if ((mask & OMP_CLAUSE_COPYPRIVATE) && gfc_match_omp_variable_list ("copyprivate (", @@ -1070,7 +1068,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_CREATE) && gfc_match ("create ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_ALLOC)) + OMP_MAP_ALLOC)) continue; break; case 'd': @@ -1106,7 +1104,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if ((mask & OMP_CLAUSE_DELETE) && gfc_match ("delete ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_DELETE)) + OMP_MAP_RELEASE)) continue; if ((mask & OMP_CLAUSE_DEPEND) && gfc_match ("depend ( ") == MATCH_YES) @@ -1161,19 +1159,10 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, OMP_MAP_FORCE_TO)) continue; if ((mask & OMP_CLAUSE_DEVICEPTR) - && gfc_match ("deviceptr ( ") == MATCH_YES) - { - gfc_omp_namelist **list = &c->lists[OMP_LIST_MAP]; - gfc_omp_namelist **head = NULL; - if (gfc_match_omp_variable_list ("", list, true, NULL, - &head, false) == MATCH_YES) - { - gfc_omp_namelist *n; - for (n = *head; n; n = n->next) - n->u.map_op = OMP_MAP_FORCE_DEVICEPTR; - continue; - } - } + && gfc_match ("deviceptr ( ") == MATCH_YES + && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], + OMP_MAP_FORCE_DEVICEPTR)) + continue; if ((mask & OMP_CLAUSE_DEVICE_RESIDENT) && gfc_match_omp_variable_list ("device_resident (", @@ -1202,6 +1191,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, && c->final_expr == NULL && gfc_match ("final ( %e )", &c->final_expr) == MATCH_YES) continue; + if ((mask & OMP_CLAUSE_FINALIZE) + && !c->finalize + && gfc_match ("finalize") == MATCH_YES) + { + c->finalize = true; + needs_space = true; + continue; + } if ((mask & OMP_CLAUSE_FIRSTPRIVATE) && gfc_match_omp_variable_list ("firstprivate (", &c->lists[OMP_LIST_FIRSTPRIVATE], @@ -1274,6 +1271,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, } gfc_current_locus = old_loc; } + if ((mask & OMP_CLAUSE_IF_PRESENT) + && !c->if_present + && gfc_match ("if_present") == MATCH_YES) + { + c->if_present = true; + needs_space = true; + continue; + } if ((mask & OMP_CLAUSE_INBRANCH) && !c->inbranch && !c->notinbranch @@ -1503,22 +1508,22 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, } break; case 'p': - if ((mask & OMP_CLAUSE_PRESENT_OR_COPY) + if ((mask & OMP_CLAUSE_COPY) && gfc_match ("pcopy ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_TOFROM)) continue; - if ((mask & OMP_CLAUSE_PRESENT_OR_COPYIN) + if ((mask & OMP_CLAUSE_COPYIN) && gfc_match ("pcopyin ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_TO)) continue; - if ((mask & OMP_CLAUSE_PRESENT_OR_COPYOUT) + if ((mask & OMP_CLAUSE_COPYOUT) && gfc_match ("pcopyout ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_FROM)) continue; - if ((mask & OMP_CLAUSE_PRESENT_OR_CREATE) + if ((mask & OMP_CLAUSE_CREATE) && gfc_match ("pcreate ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_ALLOC)) @@ -1528,22 +1533,22 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_FORCE_PRESENT)) continue; - if ((mask & OMP_CLAUSE_PRESENT_OR_COPY) + if ((mask & OMP_CLAUSE_COPY) && gfc_match ("present_or_copy ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_TOFROM)) continue; - if ((mask & OMP_CLAUSE_PRESENT_OR_COPYIN) + if ((mask & OMP_CLAUSE_COPYIN) && gfc_match ("present_or_copyin ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_TO)) continue; - if ((mask & OMP_CLAUSE_PRESENT_OR_COPYOUT) + if ((mask & OMP_CLAUSE_COPYOUT) && gfc_match ("present_or_copyout ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_FROM)) continue; - if ((mask & OMP_CLAUSE_PRESENT_OR_CREATE) + if ((mask & OMP_CLAUSE_CREATE) && gfc_match ("present_or_create ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_ALLOC)) @@ -1925,23 +1930,19 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \ | OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_REDUCTION \ | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ - | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \ - | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \ - | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE \ - | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT) + | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \ + | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \ + | OMP_CLAUSE_WAIT) #define OACC_KERNELS_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \ | OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \ | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ - | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \ - | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \ - | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT) + | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \ + | OMP_CLAUSE_WAIT) #define OACC_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \ | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \ - | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \ - | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \ - | OMP_CLAUSE_PRESENT_OR_CREATE) + | OMP_CLAUSE_PRESENT) #define OACC_LOOP_CLAUSES \ (omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER \ | OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT \ @@ -1955,19 +1956,17 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, #define OACC_DECLARE_CLAUSES \ (omp_mask (OMP_CLAUSE_COPY) | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ | OMP_CLAUSE_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_DEVICE_RESIDENT \ - | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \ - | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \ - | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_LINK) + | OMP_CLAUSE_PRESENT \ + | OMP_CLAUSE_LINK) #define OACC_UPDATE_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST_SELF \ - | OMP_CLAUSE_DEVICE | OMP_CLAUSE_WAIT) + | OMP_CLAUSE_DEVICE | OMP_CLAUSE_WAIT | OMP_CLAUSE_IF_PRESENT) #define OACC_ENTER_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \ - | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT_OR_COPYIN \ - | OMP_CLAUSE_PRESENT_OR_CREATE) + | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE) #define OACC_EXIT_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \ - | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE) + | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE | OMP_CLAUSE_FINALIZE) #define OACC_WAIT_CLAUSES \ omp_mask (OMP_CLAUSE_ASYNC) #define OACC_ROUTINE_CLAUSES \ @@ -2061,8 +2060,7 @@ gfc_match_oacc_declare (void) if (s->ns->proc_name && s->ns->proc_name->attr.proc == PROC_MODULE) { - if (n->u.map_op != OMP_MAP_FORCE_ALLOC - && n->u.map_op != OMP_MAP_FORCE_TO) + if (n->u.map_op != OMP_MAP_ALLOC && n->u.map_op != OMP_MAP_TO) { gfc_error ("Invalid clause in module with !$ACC DECLARE at %L", &where); @@ -2072,6 +2070,13 @@ gfc_match_oacc_declare (void) module_var = true; } + if (ns->proc_name->attr.oacc_function) + { + gfc_error ("Invalid declare in routine with $!ACC DECLARE at %L", + &where); + return MATCH_ERROR; + } + if (s->attr.use_assoc) { gfc_error ("Variable is USE-associated with !$ACC DECLARE at %L", @@ -2090,10 +2095,12 @@ gfc_match_oacc_declare (void) switch (n->u.map_op) { case OMP_MAP_FORCE_ALLOC: + case OMP_MAP_ALLOC: s->attr.oacc_declare_create = 1; break; case OMP_MAP_FORCE_TO: + case OMP_MAP_TO: s->attr.oacc_declare_copyin = 1; break; diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 795175d..f038f4c 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2895,6 +2895,16 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, c = build_omp_clause (where.lb->location, OMP_CLAUSE_AUTO); omp_clauses = gfc_trans_add_clause (c, omp_clauses); } + if (clauses->if_present) + { + c = build_omp_clause (where.lb->location, OMP_CLAUSE_IF_PRESENT); + omp_clauses = gfc_trans_add_clause (c, omp_clauses); + } + if (clauses->finalize) + { + c = build_omp_clause (where.lb->location, OMP_CLAUSE_FINALIZE); + omp_clauses = gfc_trans_add_clause (c, omp_clauses); + } if (clauses->independent) { c = build_omp_clause (where.lb->location, OMP_CLAUSE_INDEPENDENT); diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 1523a27..97543ed 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8524,6 +8524,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_NOGROUP: case OMP_CLAUSE_THREADS: case OMP_CLAUSE_SIMD: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: break; case OMP_CLAUSE_DEFAULTMAP: @@ -9305,6 +9307,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: case OMP_CLAUSE_TILE: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: break; default: @@ -9361,21 +9365,7 @@ gimplify_oacc_declare_1 (tree clause) switch (kind) { case GOMP_MAP_ALLOC: - case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_TO: - new_op = GOMP_MAP_DELETE; - ret = true; - break; - - case GOMP_MAP_FORCE_FROM: - OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC); - new_op = GOMP_MAP_FORCE_FROM; - ret = true; - break; - - case GOMP_MAP_FORCE_TOFROM: - OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_TO); - new_op = GOMP_MAP_FORCE_FROM; + new_op = GOMP_MAP_RELEASE; ret = true; break; @@ -10817,6 +10807,53 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p) ort, TREE_CODE (expr)); gimplify_adjust_omp_clauses (pre_p, NULL, &OMP_STANDALONE_CLAUSES (expr), TREE_CODE (expr)); + if (TREE_CODE (expr) == OACC_UPDATE + && omp_find_clause (OMP_STANDALONE_CLAUSES (expr), + OMP_CLAUSE_IF_PRESENT)) + { + /* The runtime uses GOMP_MAP_{TO,FROM} to denote the if_present + clause. */ + for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_FORCE_TO: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO); + break; + case GOMP_MAP_FORCE_FROM: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FROM); + break; + default: + break; + } + } + else if (TREE_CODE (expr) == OACC_EXIT_DATA + && omp_find_clause (OMP_STANDALONE_CLAUSES (expr), + OMP_CLAUSE_FINALIZE)) + { + /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize" + semantics apply to all mappings of this OpenACC directive. */ + bool finalize_marked = false; + for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_FROM: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_FROM); + finalize_marked = true; + break; + case GOMP_MAP_RELEASE: + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE); + finalize_marked = true; + break; + default: + /* Check consistency: libgomp relies on the very first data + mapping clause being marked, so make sure we did that before + any other mapping clauses. */ + gcc_assert (finalize_marked); + break; + } + } stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr)); gimplify_seq_add_stmt (pre_p, stmt); diff --git a/gcc/omp-low.c b/gcc/omp-low.c index ba6c705..c591231 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -642,8 +642,7 @@ build_sender_ref (tree var, omp_context *ctx) BASE_POINTERS_RESTRICT, declare the field with restrict. */ static void -install_var_field (tree var, bool by_ref, int mask, omp_context *ctx, - bool base_pointers_restrict = false) +install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) { tree field, type, sfield = NULL_TREE; splay_tree_key key = (splay_tree_key) var; @@ -674,11 +673,7 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx, type = build_pointer_type (build_pointer_type (type)); } else if (by_ref) - { - type = build_pointer_type (type); - if (base_pointers_restrict) - type = build_qualified_type (type, TYPE_QUAL_RESTRICT); - } + type = build_pointer_type (type); else if ((mask & 3) == 1 && omp_is_reference (var)) type = TREE_TYPE (type); @@ -992,12 +987,10 @@ fixup_child_record_type (omp_context *ctx) } /* Instantiate decls as necessary in CTX to satisfy the data sharing - specified by CLAUSES. If BASE_POINTERS_RESTRICT, install var field with - restrict. */ + specified by CLAUSES. */ static void -scan_sharing_clauses (tree clauses, omp_context *ctx, - bool base_pointers_restrict = false) +scan_sharing_clauses (tree clauses, omp_context *ctx) { tree c, decl; bool scan_array_reductions = false; @@ -1256,8 +1249,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) install_var_field (decl, true, 7, ctx); else - install_var_field (decl, true, 3, ctx, - base_pointers_restrict); + install_var_field (decl, true, 3, ctx); if (is_gimple_omp_offloaded (ctx->stmt) && !OMP_CLAUSE_MAP_IN_REDUCTION (c)) install_var_local (decl, ctx); @@ -1328,6 +1320,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, case OMP_CLAUSE_TILE: case OMP_CLAUSE__SIMT_: case OMP_CLAUSE_DEFAULT: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: break; case OMP_CLAUSE_ALIGNED: @@ -1499,6 +1493,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, case OMP_CLAUSE_TILE: case OMP_CLAUSE__GRIDDIM_: case OMP_CLAUSE__SIMT_: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: break; case OMP_CLAUSE__CACHE_: @@ -2266,68 +2262,6 @@ scan_omp_single (gomp_single *stmt, omp_context *outer_ctx) layout_type (ctx->record_type); } -/* Return true if the CLAUSES of an omp target guarantee that the base pointers - used in the corresponding offloaded function are restrict. */ - -static bool -omp_target_base_pointers_restrict_p (tree clauses) -{ - /* The analysis relies on the GOMP_MAP_FORCE_* mapping kinds, which are only - used by OpenACC. */ - if (flag_openacc == 0) - return false; - - /* I. Basic example: - - void foo (void) - { - unsigned int a[2], b[2]; - - #pragma acc kernels \ - copyout (a) \ - copyout (b) - { - a[0] = 0; - b[0] = 1; - } - } - - After gimplification, we have: - - #pragma omp target oacc_kernels \ - map(force_from:a [len: 8]) \ - map(force_from:b [len: 8]) - { - a[0] = 0; - b[0] = 1; - } - - Because both mappings have the force prefix, we know that they will be - allocated when calling the corresponding offloaded function, which means we - can mark the base pointers for a and b in the offloaded function as - restrict. */ - - tree c; - for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) - { - if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) - return false; - - switch (OMP_CLAUSE_MAP_KIND (c)) - { - case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_TO: - case GOMP_MAP_FORCE_FROM: - case GOMP_MAP_FORCE_TOFROM: - break; - default: - return false; - } - } - - return true; -} - /* Scan a GIMPLE_OMP_TARGET. */ static void @@ -2349,20 +2283,13 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx) TYPE_NAME (ctx->record_type) = name; TYPE_ARTIFICIAL (ctx->record_type) = 1; - bool base_pointers_restrict = false; if (offloaded) { create_omp_child_function (ctx, false); gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn); - - base_pointers_restrict = omp_target_base_pointers_restrict_p (clauses); - if (base_pointers_restrict - && dump_file && (dump_flags & TDF_DETAILS)) - fprintf (dump_file, - "Base pointers in offloaded function are restrict\n"); } - scan_sharing_clauses (clauses, ctx, base_pointers_restrict); + scan_sharing_clauses (clauses, ctx); scan_omp (gimple_omp_body_ptr (stmt), ctx); if (TYPE_FIELDS (ctx->record_type) == NULL) diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index c30af52..1c06223 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,36 @@ +2018-06-20 Chung-Lin Tang <cltang@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + Cesar Philippidis <cesar@codesourcery.com> + + * c-c++-common/goacc/declare-1.c: Update test case to utilize OpenACC + 2.5 data clause semantics. + * c-c++-common/goacc/declare-2.c: Likewise. + * c-c++-common/goacc/default-4.c: Likewise. + * c-c++-common/goacc/finalize-1.c: New test. + * c-c++-common/goacc/kernels-alias-2.c: Update test case to utilize + OpenACC 2.5 data clause semantics. + * c-c++-common/goacc/kernels-alias.c: Likewise. + * c-c++-common/goacc/routine-5.c: Likewise. + * c-c++-common/goacc/update-if_present-1.c: New test. + * c-c++-common/goacc/update-if_present-2.c: New test. + * g++.dg/goacc/template.C: Update test case to utilize OpenACC + 2.5 data clause semantics. + * gfortran.dg/goacc/combined-directives.f90: Likewise. + * gfortran.dg/goacc/data-tree.f95: Likewise. + * gfortran.dg/goacc/declare-2.f95: Likewise. + * gfortran.dg/goacc/default-4.f: Likewise. + * gfortran.dg/goacc/enter-exit-data.f95: Likewise. + * gfortran.dg/goacc/finalize-1.f: New test. + * gfortran.dg/goacc/kernels-alias-2.f95: Update test case to utilize + OpenACC 2.5 data clause semantics. + * gfortran.dg/goacc/kernels-alias.f95: Likewise. + * gfortran.dg/goacc/kernels-tree.f95: Likewise. + * gfortran.dg/goacc/nested-function-1.f90: Likewise. + * gfortran.dg/goacc/parallel-tree.f95: Likewise. + * gfortran.dg/goacc/reduction-promotions.f90: Likewise. + * gfortran.dg/goacc/update-if_present-1.f90: New test. + * gfortran.dg/goacc/update-if_present-2.f90: New test. + 2018-06-20 Jakub Jelinek <jakub@redhat.com> PR c++/86210 diff --git a/gcc/testsuite/c-c++-common/goacc/declare-1.c b/gcc/testsuite/c-c++-common/goacc/declare-1.c index b036c63..35b1ccd 100644 --- a/gcc/testsuite/c-c++-common/goacc/declare-1.c +++ b/gcc/testsuite/c-c++-common/goacc/declare-1.c @@ -19,6 +19,12 @@ int v4; int v5, v6, v7, v8; #pragma acc declare create(v5, v6) copyin(v7, v8) +int v9; +#pragma acc declare present_or_copyin(v9) + +int v10; +#pragma acc declare present_or_create(v10) + void f (void) { @@ -49,6 +55,12 @@ f (void) extern int ve4; #pragma acc declare link(ve4) + extern int ve5; +#pragma acc declare present_or_copyin(ve5) + + extern int ve6; +#pragma acc declare present_or_create(ve6) + int va5; #pragma acc declare copy(va5) diff --git a/gcc/testsuite/c-c++-common/goacc/declare-2.c b/gcc/testsuite/c-c++-common/goacc/declare-2.c index e41a0f5..33b8245 100644 --- a/gcc/testsuite/c-c++-common/goacc/declare-2.c +++ b/gcc/testsuite/c-c++-common/goacc/declare-2.c @@ -29,13 +29,7 @@ int v6; #pragma acc declare present_or_copy(v6) /* { dg-error "at file scope" } */ int v7; -#pragma acc declare present_or_copyin(v7) /* { dg-error "at file scope" } */ - -int v8; -#pragma acc declare present_or_copyout(v8) /* { dg-error "at file scope" } */ - -int v9; -#pragma acc declare present_or_create(v9) /* { dg-error "at file scope" } */ +#pragma acc declare present_or_copyout(v7) /* { dg-error "at file scope" } */ int va10; #pragma acc declare create (va10) @@ -67,13 +61,7 @@ f (void) #pragma acc declare present_or_copy(ve3) /* { dg-error "invalid use of" } */ extern int ve4; -#pragma acc declare present_or_copyin(ve4) /* { dg-error "invalid use of" } */ - - extern int ve5; -#pragma acc declare present_or_copyout(ve5) /* { dg-error "invalid use of" } */ - - extern int ve6; -#pragma acc declare present_or_create(ve6) /* { dg-error "invalid use of" } */ +#pragma acc declare present_or_copyout(ve4) /* { dg-error "invalid use of" } */ -#pragma acc declare present (v9) /* { dg-error "invalid use of" } */ +#pragma acc declare present (v2) /* { dg-error "invalid use of" } */ } diff --git a/gcc/testsuite/c-c++-common/goacc/default-4.c b/gcc/testsuite/c-c++-common/goacc/default-4.c index dfa79bb..867175d 100644 --- a/gcc/testsuite/c-c++-common/goacc/default-4.c +++ b/gcc/testsuite/c-c++-common/goacc/default-4.c @@ -8,7 +8,7 @@ void f1 () float f1_b[2]; #pragma acc data copyin (f1_a) copyout (f1_b) - /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f1_b \[^\\)\]+\\) map\\(force_to:f1_a" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f1_b \[^\\)\]+\\) map\\(to:f1_a" 1 "gimple" } } */ { #pragma acc kernels /* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } } */ @@ -29,7 +29,7 @@ void f2 () float f2_b[2]; #pragma acc data copyin (f2_a) copyout (f2_b) - /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f2_b \[^\\)\]+\\) map\\(force_to:f2_a" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f2_b \[^\\)\]+\\) map\\(to:f2_a" 1 "gimple" } } */ { #pragma acc kernels default (none) /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } } */ @@ -50,7 +50,7 @@ void f3 () float f3_b[2]; #pragma acc data copyin (f3_a) copyout (f3_b) - /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f3_b \[^\\)\]+\\) map\\(force_to:f3_a" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f3_b \[^\\)\]+\\) map\\(to:f3_a" 1 "gimple" } } */ { #pragma acc kernels default (present) /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/finalize-1.c b/gcc/testsuite/c-c++-common/goacc/finalize-1.c new file mode 100644 index 0000000..9482029 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/finalize-1.c @@ -0,0 +1,28 @@ +/* Test valid usage and processing of the finalize clause. */ + +/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */ + +extern int del_r; +extern float del_f[3]; +extern double cpo_r[8]; +extern long cpo_f; + +void f () +{ +#pragma acc exit data delete (del_r) +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } } + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ + +#pragma acc exit data finalize delete (del_f) +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } } + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } */ + +#pragma acc exit data copyout (cpo_r) +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } } + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ + +#pragma acc exit data copyout (cpo_f) finalize +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:cpo_f\\);$" 1 "original" } } + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ +} + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c index d437c47..7576a64 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c @@ -18,10 +18,12 @@ foo (void) } } +/* The xfails occur due to the OpenACC 2.5 data semantics. */ + /* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } } */ /* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias.c index 25821ab2..e8ff018 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-alias.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias.c @@ -20,10 +20,12 @@ foo (void) } } +/* The xfails occur due to the OpenACC 2.5 data semantics. */ + /* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } } */ /* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c index b967a74..b759db3 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-5.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c @@ -4,11 +4,11 @@ struct PC { -#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */ }; void PC1( /* { dg-bogus "variable or field .PC1. declared void" "TODO" { xfail c++ } } */ -#pragma acc routine +#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" "" { target c } .-1 } { dg-error ".#pragma. is not allowed here" "" { target c++ } .-2 } */ ) /* { dg-bogus "expected declaration specifiers or .\\.\\.\\.. before .\\). token" "TODO" { xfail c } } */ @@ -18,26 +18,26 @@ void PC1( /* { dg-bogus "variable or field .PC1. declared void" "TODO" { xfail c void PC2() { if (0) -#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */ ; } void PC3() { -#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */ } /* "( name )" syntax. */ #pragma acc routine ( /* { dg-error "expected (function name|unqualified-id) before end of line" } */ -#pragma acc routine () /* { dg-error "expected (function name|unqualified-id) before .\\). token" } */ -#pragma acc routine (+) /* { dg-error "expected (function name|unqualified-id) before .\\+. token" } */ -#pragma acc routine (?) /* { dg-error "expected (function name|unqualified-id) before .\\?. token" } */ -#pragma acc routine (:) /* { dg-error "expected (function name|unqualified-id) before .:. token" } */ -#pragma acc routine (4) /* { dg-error "expected (function name|unqualified-id) before numeric constant" } */ +#pragma acc routine () seq /* { dg-error "expected (function name|unqualified-id) before .\\). token" } */ +#pragma acc routine (+) seq /* { dg-error "expected (function name|unqualified-id) before .\\+. token" } */ +#pragma acc routine (?) seq /* { dg-error "expected (function name|unqualified-id) before .\\?. token" } */ +#pragma acc routine (:) seq /* { dg-error "expected (function name|unqualified-id) before .:. token" } */ +#pragma acc routine (4) seq /* { dg-error "expected (function name|unqualified-id) before numeric constant" } */ #pragma acc routine ('4') /* { dg-error "expected (function name|unqualified-id) before .4." } */ -#pragma acc routine ("4") /* { dg-error "expected (function name|unqualified-id) before string constant" } */ +#pragma acc routine ("4") seq /* { dg-error "expected (function name|unqualified-id) before string constant" } */ extern void R1(void); extern void R2(void); #pragma acc routine (R1, R2, R3) worker /* { dg-error "expected .\\). before .,. token" } */ @@ -49,84 +49,84 @@ extern void R2(void); /* "#pragma acc routine" not immediately followed by (a single) function declaration or definition. */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int a; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ void fn1 (void), fn1b (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int b, fn2 (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int b_, fn2_ (void), B_; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ int fn3 (void), b2; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ typedef struct c c; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ struct d {} d; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ void fn1_2 (void), fn1b_2 (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ #pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int b_2, fn2_2 (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int b_2_, fn2_2_ (void), B_2_; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ int fn3_2 (void), b2_2; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ typedef struct c_2 c_2; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ struct d_2 {} d_2; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq int fn4 (void); int fn5a (void); int fn5b (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine (fn5a) -#pragma acc routine (fn5b) +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine (fn5a) seq +#pragma acc routine (fn5b) seq int fn5 (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine (fn6a) /* { dg-error ".fn6a. has not been declared" } */ -#pragma acc routine (fn6b) /* { dg-error ".fn6b. has not been declared" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine (fn6a) seq /* { dg-error ".fn6a. has not been declared" } */ +#pragma acc routine (fn6b) seq /* { dg-error ".fn6b. has not been declared" } */ int fn6 (void); #ifdef __cplusplus -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */ namespace f {} namespace g {} -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */ using namespace g; -#pragma acc routine (g) /* { dg-error ".g. does not refer to a function" "" { target c++ } } */ +#pragma acc routine (g) seq /* { dg-error ".g. does not refer to a function" "" { target c++ } } */ #endif /* __cplusplus */ -#pragma acc routine (a) /* { dg-error ".a. does not refer to a function" } */ +#pragma acc routine (a) seq /* { dg-error ".a. does not refer to a function" } */ -#pragma acc routine (c) /* { dg-error ".c. does not refer to a function" } */ +#pragma acc routine (c) seq /* { dg-error ".c. does not refer to a function" } */ /* Static assert. */ @@ -143,66 +143,24 @@ static_assert(0, ""); /* { dg-error "static assertion failed" "" { target c++11 #endif void f_static_assert(); /* Check that we already recognized "f_static_assert" as an OpenACC routine. */ -#pragma acc routine (f_static_assert) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*f_static_assert" "TODO" { xfail *-*-* } } */ +#pragma acc routine (f_static_assert) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*f_static_assert" "TODO" { xfail *-*-* } } */ /* __extension__ usage. */ -#pragma acc routine +#pragma acc routine seq __extension__ extern void ex1(); #pragma acc routine (ex1) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex1" } */ -#pragma acc routine +#pragma acc routine seq __extension__ __extension__ __extension__ __extension__ __extension__ void ex2() { } #pragma acc routine (ex2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex2" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ __extension__ int ex3; -#pragma acc routine (ex3) /* { dg-error ".ex3. does not refer to a function" } */ - - -/* "#pragma acc routine" already applied. */ - -extern void fungsi_1(); -#pragma acc routine(fungsi_1) gang -#pragma acc routine(fungsi_1) gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */ -#pragma acc routine(fungsi_1) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */ -#pragma acc routine(fungsi_1) vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */ - -#pragma acc routine seq -extern void fungsi_2(); -#pragma acc routine(fungsi_2) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */ -#pragma acc routine(fungsi_2) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */ -#pragma acc routine(fungsi_2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */ - -#pragma acc routine vector -extern void fungsi_3(); -#pragma acc routine vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_3." } */ -void fungsi_3() -{ -} - -extern void fungsi_4(); -#pragma acc routine (fungsi_4) worker -#pragma acc routine gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_4." } */ -void fungsi_4() -{ -} - -#pragma acc routine gang -void fungsi_5() -{ -} -#pragma acc routine (fungsi_5) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_5." } */ - -#pragma acc routine seq -void fungsi_6() -{ -} -#pragma acc routine seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_6." } */ -extern void fungsi_6(); +#pragma acc routine (ex3) seq /* { dg-error ".ex3. does not refer to a function" } */ /* "#pragma acc routine" must be applied before. */ @@ -214,11 +172,11 @@ void Foo () Bar (); } -#pragma acc routine (Bar) // { dg-error ".#pragma acc routine. must be applied before use" } +#pragma acc routine (Bar) seq // { dg-error ".#pragma acc routine. must be applied before use" } #pragma acc routine (Foo) gang // { dg-error ".#pragma acc routine. must be applied before definition" } -#pragma acc routine (Baz) // { dg-error "not been declared" } +#pragma acc routine (Baz) seq // { dg-error "not been declared" } /* OpenACC declare. */ @@ -227,7 +185,7 @@ int vb1; /* { dg-error "directive for use" } */ extern int vb2; /* { dg-error "directive for use" } */ static int vb3; /* { dg-error "directive for use" } */ -#pragma acc routine +#pragma acc routine seq int func1 (int a) { @@ -238,7 +196,7 @@ func1 (int a) return vb3; } -#pragma acc routine +#pragma acc routine seq int func2 (int a) { @@ -256,7 +214,7 @@ extern int vb6; /* { dg-error "clause used in" } */ static int vb7; /* { dg-error "clause used in" } */ #pragma acc declare link (vb7) -#pragma acc routine +#pragma acc routine seq int func3 (int a) { @@ -273,7 +231,7 @@ extern int vb9; static int vb10; #pragma acc declare create (vb10) -#pragma acc routine +#pragma acc routine seq int func4 (int a) { @@ -291,7 +249,7 @@ extern int vb12; extern int vb13; #pragma acc declare device_resident (vb13) -#pragma acc routine +#pragma acc routine seq int func5 (int a) { @@ -302,7 +260,7 @@ func5 (int a) return vb13; } -#pragma acc routine +#pragma acc routine seq int func6 (int a) { diff --git a/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c b/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c new file mode 100644 index 0000000..c34a0e4 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c @@ -0,0 +1,28 @@ +/* Test valid usages of the if_present clause. */ + +/* { dg-additional-options "-fdump-tree-omplower" } */ + +void +t () +{ + int a, b, c[10]; + +#pragma acc update self(a) if_present +#pragma acc update device(b) async if_present +#pragma acc update host(c[1:3]) wait(4) if_present +#pragma acc update self(c) device(b) host (a) async(10) if (a == 5) if_present + +#pragma acc update self(a) +#pragma acc update device(b) async +#pragma acc update host(c[1:3]) wait(4) +#pragma acc update self(c) device(b) host (a) async(10) if (a == 5) +} + +/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present map.from:a .len: 4.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present async.-1. map.to:b .len: 4.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present wait.4. map.from:c.1. .len: 12.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present if.... async.10. map.from:a .len: 4.. map.to:b .len: 4.. map.from:c .len: 40.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_update map.force_from:a .len: 4.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_update async.-1. map.force_to:b .len: 4.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_update wait.4. map.force_from:c.1. .len: 12.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_update if.... async.10. map.force_from:a .len: 4.. map.force_to:b .len: 4.. map.force_from:c .len: 40.." 1 "omplower" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c b/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c new file mode 100644 index 0000000..974f1b8 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c @@ -0,0 +1,42 @@ +/* Test invalid usages of the if_present clause. */ + +#pragma acc routine gang if_present /* { dg-error "'if_present' is not valid" } */ +void +t1 () +{ + int a, b, c[10]; + +#pragma acc enter data copyin(a) if_present /* { dg-error "'if_present' is not valid" } */ +#pragma acc exit data copyout(a) if_present /* { dg-error "'if_present' is not valid" } */ + +#pragma acc data copy(a) if_present /* { dg-error "'if_present' is not valid" } */ + { + } + +#pragma acc declare create(c) if_present /* { dg-error "'if_present' is not valid" } */ + +#pragma acc init if_present +#pragma acc shutdown if_present +} + +void +t2 () +{ + int a, b, c[10]; + +#pragma acc update self(a) +#pragma acc parallel +#pragma acc loop if_present /* { dg-error "'if_present' is not valid" } */ + for (b = 1; b < 10; b++) + ; +#pragma acc end parallel + +#pragma acc kernels loop if_present /* { dg-error "'if_present' is not valid" } */ + for (b = 1; b < 10; b++) + ; + +#pragma acc parallel loop if_present /* { dg-error "'if_present' is not valid" } */ + for (b = 1; b < 10; b++) + ; +} + diff --git a/gcc/testsuite/g++.dg/goacc/template.C b/gcc/testsuite/g++.dg/goacc/template.C index 852f42f..dae92b0 100644 --- a/gcc/testsuite/g++.dg/goacc/template.C +++ b/gcc/testsuite/g++.dg/goacc/template.C @@ -1,4 +1,4 @@ -#pragma acc routine +#pragma acc routine seq template <typename T> T accDouble(int val) { @@ -31,7 +31,7 @@ oacc_parallel_copy (T a) #pragma acc parallel num_gangs (a) if (1) { -#pragma acc loop independent collapse (2) gang +#pragma acc loop independent collapse (2) for (int i = 0; i < a; i++) for (int j = 0; j < 5; j++) b = a; @@ -86,6 +86,8 @@ oacc_parallel_copy (T a) #pragma acc update self (b) #pragma acc update device (b) #pragma acc exit data delete (b) +#pragma acc exit data finalize copyout (b) +#pragma acc exit data delete (b) finalize return b; } @@ -133,6 +135,13 @@ oacc_kernels_copy (T a) b = a; } +#pragma acc update host (b) +#pragma acc update self (b) +#pragma acc update device (b) +#pragma acc exit data delete (b) +#pragma acc exit data finalize copyout (b) +#pragma acc exit data delete (b) finalize + return b; } diff --git a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 index 42a447a..9563492 100644 --- a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 @@ -146,5 +146,5 @@ end subroutine test ! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3" 2 "gimple" } } ! { dg-final { scan-tree-dump-times "acc loop private.i. independent" 2 "gimple" } } ! { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } } -! { dg-final { scan-tree-dump-times "omp target oacc_\[^ \]+ map.force_tofrom:y" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_\[^ \]+ map.tofrom:y" 2 "gimple" } } ! { dg-final { scan-tree-dump-times "acc loop private.i. reduction..:y." 2 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 index 44efc8a..f16d62c 100644 --- a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 @@ -15,10 +15,10 @@ end program test ! { dg-final { scan-tree-dump-times "pragma acc data" 1 "original" } } ! { dg-final { scan-tree-dump-times "if" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 index aa1704f..7aa3dab 100644 --- a/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 @@ -11,11 +11,11 @@ subroutine asubr (b) !$acc declare copyout (b) ! { dg-error "Invalid clause in module" } !$acc declare present (b) ! { dg-error "Invalid clause in module" } !$acc declare present_or_copy (b) ! { dg-error "Invalid clause in module" } - !$acc declare present_or_copyin (b) ! { dg-error "Invalid clause in module" } + !$acc declare present_or_copyin (b) ! { dg-error "present on multiple" } !$acc declare present_or_copyout (b) ! { dg-error "Invalid clause in module" } - !$acc declare present_or_create (b) ! { dg-error "Invalid clause in module" } + !$acc declare present_or_create (b) ! { dg-error "present on multiple" } !$acc declare deviceptr (b) ! { dg-error "Invalid clause in module" } - !$acc declare create (b) copyin (b) ! { dg-error "present on multiple clauses" } + !$acc declare create (b) copyin (b) ! { dg-error "present on multiple" } end subroutine diff --git a/gcc/testsuite/gfortran.dg/goacc/default-4.f b/gcc/testsuite/gfortran.dg/goacc/default-4.f index 77291f4..30f411f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/default-4.f +++ b/gcc/testsuite/gfortran.dg/goacc/default-4.f @@ -8,7 +8,7 @@ REAL, DIMENSION (2) :: F1_B !$ACC DATA COPYIN (F1_A) COPYOUT (F1_B) -! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f1_a \[^\\)\]+\\) map\\(force_from:f1_b" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f1_a \[^\\)\]+\\) map\\(from:f1_b" 1 "gimple" } } !$ACC KERNELS ! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } } F1_B(1) = F1_A; @@ -26,7 +26,7 @@ REAL, DIMENSION (2) :: F2_B !$ACC DATA COPYIN (F2_A) COPYOUT (F2_B) -! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f2_a \[^\\)\]+\\) map\\(force_from:f2_b" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f2_a \[^\\)\]+\\) map\\(from:f2_b" 1 "gimple" } } !$ACC KERNELS DEFAULT (NONE) ! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } } F2_B(1) = F2_A; @@ -44,7 +44,7 @@ REAL, DIMENSION (2) :: F3_B !$ACC DATA COPYIN (F3_A) COPYOUT (F3_B) -! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f3_a \[^\\)\]+\\) map\\(force_from:f3_b" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f3_a \[^\\)\]+\\) map\\(from:f3_b" 1 "gimple" } } !$ACC KERNELS DEFAULT (PRESENT) ! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } F3_B(1) = F3_A; diff --git a/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 index 8f1715e..805459c 100644 --- a/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 @@ -84,5 +84,8 @@ contains !$acc exit data delete (tip) ! { dg-error "POINTER" } !$acc exit data delete (tia) ! { dg-error "ALLOCATABLE" } !$acc exit data copyout (i) delete (i) ! { dg-error "multiple clauses" } + !$acc exit data finalize + !$acc exit data finalize copyout (i) + !$acc exit data finalize delete (i) end subroutine foo end module test diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f new file mode 100644 index 0000000..5c7a921 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f @@ -0,0 +1,27 @@ +! Test valid usage and processing of the finalize clause. + +! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } + + SUBROUTINE f + IMPLICIT NONE + INTEGER :: del_r + REAL, DIMENSION (3) :: del_f + DOUBLE PRECISION, DIMENSION (8) :: cpo_r + LOGICAL :: cpo_f + +!$ACC EXIT DATA DELETE (del_r) +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } + +!$ACC EXIT DATA FINALIZE DELETE (del_f) +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } + +!$ACC EXIT DATA COPYOUT (cpo_r) +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } + +!$ACC EXIT DATA COPYOUT (cpo_f) FINALIZE +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_f\\) finalize;$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } + END SUBROUTINE f diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-alias-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-alias-2.f95 index 7e348dd..6a9f241 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-alias-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-alias-2.f95 @@ -15,9 +15,11 @@ program main end program main +! The xfails occur in light of the new OpenACC data semantics. + ! { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } +! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } } +! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } } +! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } } +! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } } ! { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-alias.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-alias.f95 index 8d6ccb3..62f9a71 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-alias.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-alias.f95 @@ -15,9 +15,11 @@ program main end program main +! The xfails occur in light of the new OpenACC data semantics. + ! { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } +! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } } +! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } } +! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } } +! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } } ! { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 index 7daca59..a70f1e7 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 @@ -21,10 +21,10 @@ end program test ! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } } ! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 index 2fcaa40..005193f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 @@ -25,6 +25,8 @@ contains local_a (:) = 5 local_arg = 5 + !$acc update device(local_a) if_present + !$acc kernels loop & !$acc gang(num:local_arg) worker(local_arg) vector(local_arg) & !$acc wait async(local_arg) @@ -54,12 +56,16 @@ contains enddo enddo !$acc end kernels loop + + !$acc exit data copyout(local_a) delete(local_i) finalize end subroutine local subroutine nonlocal () nonlocal_a (:) = 5 nonlocal_arg = 5 + !$acc update device(nonlocal_a) if_present + !$acc kernels loop & !$acc gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) & !$acc wait async(nonlocal_arg) @@ -89,5 +95,7 @@ contains enddo enddo !$acc end kernels loop + + !$acc exit data copyout(nonlocal_a) delete(nonlocal_i) finalize end subroutine nonlocal end program main diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 index 5b2e01d..2697bb7 100644 --- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 @@ -1,5 +1,4 @@ -! { dg-do compile } -! { dg-additional-options "-fdump-tree-original" } +! { dg-additional-options "-fdump-tree-original" } ! test for tree-dump-original and spaces-commas @@ -15,6 +14,7 @@ program test !$acc end parallel end program test + ! { dg-final { scan-tree-dump-times "pragma acc parallel" 1 "original" } } ! { dg-final { scan-tree-dump-times "if" 1 "original" } } @@ -24,10 +24,10 @@ end program test ! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } ! { dg-final { scan-tree-dump-times "reduction\\(max:q\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 b/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 index 6ff913a..1d247ca 100644 --- a/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 @@ -38,9 +38,7 @@ program test !$acc end parallel end program test -! { dg-final { scan-tree-dump-times "map.tofrom:v1" 8 "gimple" } } -! { dg-final { scan-tree-dump-times "map.tofrom:v2" 8 "gimple" } } -! { dg-final { scan-tree-dump-times "map.force_tofrom:v1" 1 "gimple" } } -! { dg-final { scan-tree-dump-times "map.force_tofrom:v2" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map.tofrom:v1" 9 "gimple" } } +! { dg-final { scan-tree-dump-times "map.tofrom:v2" 9 "gimple" } } ! { dg-final { scan-tree-dump-times "map.force_present:v1" 1 "gimple" } } ! { dg-final { scan-tree-dump-times "map.force_present:v2" 1 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90 b/gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90 new file mode 100644 index 0000000..a183aae --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90 @@ -0,0 +1,27 @@ +! Test valid usages of the if_present clause. + +! { dg-additional-options "-fdump-tree-omplower" } + +subroutine t + implicit none + integer a, b, c(10) + real, allocatable :: x, y, z(:) + + a = 5 + b = 10 + c(:) = -1 + + allocate (x, y, z(100)) + + !$acc update self(a) if_present + !$acc update device(b) if_present async + !$acc update host(c(1:3)) wait(4) if_present + !$acc update self(c) device(a) host(b) if_present async(10) if(a == 10) + + !$acc update self(x) if_present + !$acc update device(y) if_present async + !$acc update host(z(1:3)) wait(3) if_present + !$acc update self(z) device(y) host(x) if_present async(4) if(a == 1) +end subroutine t + +! { dg-final { scan-tree-dump-times " if_present" 8 "omplower" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90 b/gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90 new file mode 100644 index 0000000..e73c2dc --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90 @@ -0,0 +1,52 @@ +! Test invalid usages of the if_present clause. + +subroutine t1 + implicit none + !$acc routine gang if_present ! { dg-error "Unclassifiable OpenACC directive" } + integer a, b, c(10) + real, allocatable :: x, y, z(:) + + a = 5 + b = 10 + c(:) = -1 + + allocate (x, y, z(100)) + + !$acc enter data copyin(a) if_present ! { dg-error "Unclassifiable OpenACC directive" } + !$acc exit data copyout(a) if_present ! { dg-error "Unclassifiable OpenACC directive" } + + !$acc data copy(a) if_present ! { dg-error "Unclassifiable OpenACC directive" } + !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" } + + !$acc declare link(a) if_present ! { dg-error "Unexpected junk after" } + + !$acc init if_present ! { dg-error "Unclassifiable OpenACC directive" } + !$acc shutdown if_present ! { dg-error "Unclassifiable OpenACC directive" } + + !$acc update self(a) device_type(nvidia) device(b) if_present ! { dg-error "Unclassifiable OpenACC directive" } +end subroutine t1 + +subroutine t2 + implicit none + integer a, b, c(10) + + a = 5 + b = 10 + c(:) = -1 + + !$acc parallel + !$acc loop if_present ! { dg-error "Unclassifiable OpenACC directive" } + do b = 1, 10 + end do + !$acc end parallel + + !$acc kernels loop if_present ! { dg-error "Unclassifiable OpenACC directive" } + do b = 1, 10 + end do + !$acc end kernels loop ! { dg-error "Unexpected ..ACC END KERNELS LOOP statement" } + + !$acc parallel loop if_present ! { dg-error "Unclassifiable OpenACC directive" } + do b = 1, 10 + end do + !$acc end parallel loop ! { dg-error "Unexpected ..ACC END PARALLEL LOOP statement" } +end subroutine t2 diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 2bebb22..4a04e9e 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -454,7 +454,13 @@ enum omp_clause_code { /* OpenMP internal-only clause to specify grid dimensions of a gridified kernel. */ - OMP_CLAUSE__GRIDDIM_ + OMP_CLAUSE__GRIDDIM_, + + /* OpenACC clause: if_present. */ + OMP_CLAUSE_IF_PRESENT, + + /* OpenACC clause: finalize. */ + OMP_CLAUSE_FINALIZE }; #undef DEFTREESTRUCT diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c index b335d6b..257ceae 100644 --- a/gcc/tree-nested.c +++ b/gcc/tree-nested.c @@ -1333,6 +1333,8 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_SEQ: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: break; /* The following clause belongs to the OpenACC cache directive, which @@ -2022,6 +2024,8 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_SEQ: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: break; /* The following clause belongs to the OpenACC cache directive, which diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 63ec823..e65c40a 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -1045,6 +1045,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) false); pp_right_paren (pp); break; + case OMP_CLAUSE_IF_PRESENT: + pp_string (pp, "if_present"); + break; + case OMP_CLAUSE_FINALIZE: + pp_string (pp, "finalize"); + break; default: /* Should never happen. */ @@ -343,6 +343,8 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_VECTOR_LENGTH */ 3, /* OMP_CLAUSE_TILE */ 2, /* OMP_CLAUSE__GRIDDIM_ */ + 0, /* OMP_CLAUSE_IF_PRESENT */ + 0, /* OMP_CLAUSE_FINALIZE */ }; const char * const omp_clause_code_name[] = @@ -413,7 +415,9 @@ const char * const omp_clause_code_name[] = "num_workers", "vector_length", "tile", - "_griddim_" + "_griddim_", + "if_present", + "finalize", }; @@ -11594,6 +11598,8 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE_SEQ: case OMP_CLAUSE_TILE: case OMP_CLAUSE__SIMT_: + case OMP_CLAUSE_IF_PRESENT: + case OMP_CLAUSE_FINALIZE: WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp)); case OMP_CLAUSE_LASTPRIVATE: diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index f575180..c4ba406 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,82 @@ +2018-06-20 Chung-Lin Tang <cltang@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + Cesar Philippidis <cesar@codesourcery.com> + + * libgomp.h (struct splay_tree_key_s): Add dynamic_refcount member. + (gomp_acc_remove_pointer): Update declaration. + (gomp_acc_declare_allocate): Declare. + (gomp_remove_var): Declare. + * libgomp.map (OACC_2.5): Define. + * oacc-mem.c (acc_map_data): Update refcount. + (acc_unmap_data): Likewise. + (present_create_copy): Likewise. + (acc_create): Add FLAG_PRESENT when calling present_create_copy. + (acc_copyin): Likewise. + (FLAG_FINALIZE): Define. + (delete_copyout): Update dynamic refcounts, add support for FINALIZE. + (acc_delete_finalize): New function. + (acc_delete_finalize_async): New function. + (acc_copyout_finalize): New function. + (acc_copyout_finalize_async): New function. + (gomp_acc_insert_pointer): Update refcounts. + (gomp_acc_remove_pointer): Return if data is not present on the + accelerator. + * oacc-parallel.c (find_pset): Rename to find_pointer. + (find_pointer): Add support for GOMP_MAP_POINTER. + (handle_ftn_pointers): New function. + (GOACC_parallel_keyed): Update refcounts of variables. + (GOACC_enter_exit_data): Add support for finalized data mappings. + Add support for GOMP_MAP_{TO,ALLOC,RELESE,FROM}. Update handling + of fortran arrays. + (GOACC_update): Add support for GOMP_MAP_{ALWAYS_POINTER,TO,FROM}. + (GOACC_declare): Add support for GOMP_MAP_RELEASE, remove support + for GOMP_MAP_FORCE_FROM. + * openacc.f90 (module openacc_internal): Add + acc_copyout_finalize_{32_h,64_h,array_h,_l}, and + acc_delete_finalize_{32_h,64_h,array_h,_l}. Add interfaces for + acc_copyout_finalize and acc_delete_finalize. + (acc_copyout_finalize_32_h): New subroutine. + (acc_copyout_finalize_64_h): New subroutine. + (acc_copyout_finalize_array_h): New subroutine. + (acc_delete_finalize_32_h): New subroutine. + (acc_delete_finalize_64_h): New subroutine. + (acc_delete_finalize_array_h): New subroutine. + * openacc.h (acc_copyout_finalize): Declare. + (acc_copyout_finalize_async): Declare. + (acc_delete_finalize): Declare. + (acc_delete_finalize_async): Declare. + * openacc_lib.h (acc_copyout_finalize): New interface. + (acc_delete_finalize): New interface. + * target.c (gomp_map_vars): Update dynamic_refcount. + (gomp_remove_var): New function. + (gomp_unmap_vars): Use it. + (gomp_unload_image_from_device): Likewise. + * testsuite/libgomp.oacc-c-c++-common/data-already-1.c: Update test + case to utilize OpenACC 2.5 data clause semantics. + * testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-16.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-32.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-83.c: Likewise. + * testsuite/libgomp.oacc-fortran/data-5.f90: New test. + * testsuite/libgomp.oacc-fortran/data-already-1.f: Update test case to + utilize OpenACC 2.5 data clause semantics. + * testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise. + * testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise. + * testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise. + * testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise. + * testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise. + * testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise. + * testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise. + * testsuite/libgomp.oacc-fortran/lib-32-1.f: Likewise. + * testsuite/libgomp.oacc-fortran/lib-32-2.f: Likewise. + 2018-05-21 Janus Weil <janus@gcc.gnu.org> PR fortran/85841 diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 10ea894..3a8cc2b 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -853,6 +853,8 @@ struct splay_tree_key_s { uintptr_t tgt_offset; /* Reference count. */ uintptr_t refcount; + /* Dynamic reference count. */ + uintptr_t dynamic_refcount; /* Pointer to the original mapping of "omp declare target link" object. */ splay_tree_key link_key; }; @@ -991,7 +993,9 @@ enum gomp_map_vars_kind }; extern void gomp_acc_insert_pointer (size_t, void **, size_t *, void *); -extern void gomp_acc_remove_pointer (void *, bool, int, int); +extern void gomp_acc_remove_pointer (void *, size_t, bool, int, int, int); +extern void gomp_acc_declare_allocate (bool, size_t, void **, size_t *, + unsigned short *); extern struct target_mem_desc *gomp_map_vars (struct gomp_device_descr *, size_t, void **, void **, @@ -1001,6 +1005,7 @@ extern void gomp_unmap_vars (struct target_mem_desc *, bool); extern void gomp_init_device (struct gomp_device_descr *); extern void gomp_free_memmap (struct splay_tree_s *); extern void gomp_unload_device (struct gomp_device_descr *); +extern bool gomp_remove_var (struct gomp_device_descr *, splay_tree_key); /* work.c */ diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index 8752348..2cd3bf5 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -386,6 +386,18 @@ OACC_2.0.1 { acc_pcreate; } OACC_2.0; +OACC_2.5 { + global: + acc_copyout_finalize; + acc_copyout_finalize_32_h_; + acc_copyout_finalize_64_h_; + acc_copyout_finalize_array_h_; + acc_delete_finalize; + acc_delete_finalize_32_h_; + acc_delete_finalize_64_h_; + acc_delete_finalize_array_h_; +} OACC_2.0.1; + GOACC_2.0 { global: GOACC_data_end; diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 158f086..3787ce4 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -347,6 +347,7 @@ acc_map_data (void *h, void *d, size_t s) tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes, &kinds, true, GOMP_MAP_VARS_OPENACC); + tgt->list[0].key->refcount = REFCOUNT_INFINITY; } gomp_mutex_lock (&acc_dev->lock); @@ -389,6 +390,9 @@ acc_unmap_data (void *h) (void *) n->host_start, (int) host_size, (void *) h); } + /* Mark for removal. */ + n->refcount = 1; + t = n->tgt; if (t->refcount == 2) @@ -460,6 +464,11 @@ present_create_copy (unsigned f, void *h, size_t s) gomp_fatal ("[%p,+%d] not mapped", (void *)h, (int)s); } + if (n->refcount != REFCOUNT_INFINITY) + { + n->refcount++; + n->dynamic_refcount++; + } gomp_mutex_unlock (&acc_dev->lock); } else if (!(f & FLAG_CREATE)) @@ -483,6 +492,8 @@ present_create_copy (unsigned f, void *h, size_t s) tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, NULL, &s, &kinds, true, GOMP_MAP_VARS_OPENACC); + /* Initialize dynamic refcount. */ + tgt->list[0].key->dynamic_refcount = 1; gomp_mutex_lock (&acc_dev->lock); @@ -499,13 +510,13 @@ present_create_copy (unsigned f, void *h, size_t s) void * acc_create (void *h, size_t s) { - return present_create_copy (FLAG_CREATE, h, s); + return present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s); } void * acc_copyin (void *h, size_t s) { - return present_create_copy (FLAG_CREATE | FLAG_COPY, h, s); + return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s); } void * @@ -542,7 +553,8 @@ acc_pcopyin (void *h, size_t s) } #endif -#define FLAG_COPYOUT (1 << 0) +#define FLAG_COPYOUT (1 << 0) +#define FLAG_FINALIZE (1 << 1) static void delete_copyout (unsigned f, void *h, size_t s, const char *libfnname) @@ -581,15 +593,52 @@ delete_copyout (unsigned f, void *h, size_t s, const char *libfnname) (void *) n->host_start, (int) host_size, (void *) h, (int) s); } - gomp_mutex_unlock (&acc_dev->lock); + if (n->refcount == REFCOUNT_INFINITY) + { + n->refcount = 0; + n->dynamic_refcount = 0; + } + if (n->refcount < n->dynamic_refcount) + { + gomp_mutex_unlock (&acc_dev->lock); + gomp_fatal ("Dynamic reference counting assert fail\n"); + } - if (f & FLAG_COPYOUT) - acc_dev->dev2host_func (acc_dev->target_id, h, d, s); + if (f & FLAG_FINALIZE) + { + n->refcount -= n->dynamic_refcount; + n->dynamic_refcount = 0; + } + else if (n->dynamic_refcount) + { + n->dynamic_refcount--; + n->refcount--; + } - acc_unmap_data (h); + if (n->refcount == 0) + { + if (n->tgt->refcount == 2) + { + struct target_mem_desc *tp, *t; + for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL; + tp = t, t = t->prev) + if (n->tgt == t) + { + if (tp) + tp->prev = t->prev; + else + acc_dev->openacc.data_environ = t->prev; + break; + } + } - if (!acc_dev->free_func (acc_dev->target_id, d)) - gomp_fatal ("error in freeing device memory in %s", libfnname); + if (f & FLAG_COPYOUT) + acc_dev->dev2host_func (acc_dev->target_id, h, d, s); + + gomp_remove_var (acc_dev, n); + } + + gomp_mutex_unlock (&acc_dev->lock); } void @@ -599,11 +648,35 @@ acc_delete (void *h , size_t s) } void +acc_delete_finalize (void *h , size_t s) +{ + delete_copyout (FLAG_FINALIZE, h, s, __FUNCTION__); +} + +void +acc_delete_finalize_async (void *h , size_t s, int async) +{ + delete_copyout (FLAG_FINALIZE, h, s, __FUNCTION__); +} + +void acc_copyout (void *h, size_t s) { delete_copyout (FLAG_COPYOUT, h, s, __FUNCTION__); } +void +acc_copyout_finalize (void *h, size_t s) +{ + delete_copyout (FLAG_COPYOUT | FLAG_FINALIZE, h, s, __FUNCTION__); +} + +void +acc_copyout_finalize_async (void *h, size_t s, int async) +{ + delete_copyout (FLAG_COPYOUT | FLAG_FINALIZE, h, s, __FUNCTION__); +} + static void update_dev_host (int is_dev, void *h, size_t s) { @@ -659,11 +732,37 @@ gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes, struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (acc_is_present (*hostaddrs, *sizes)) + { + splay_tree_key n; + gomp_mutex_lock (&acc_dev->lock); + n = lookup_host (acc_dev, *hostaddrs, *sizes); + gomp_mutex_unlock (&acc_dev->lock); + + tgt = n->tgt; + for (size_t i = 0; i < tgt->list_count; i++) + if (tgt->list[i].key == n) + { + for (size_t j = 0; j < mapnum; j++) + if (i + j < tgt->list_count && tgt->list[i + j].key) + { + tgt->list[i + j].key->refcount++; + tgt->list[i + j].key->dynamic_refcount++; + } + return; + } + /* Should not reach here. */ + gomp_fatal ("Dynamic refcount incrementing failed for pointer/pset"); + } + gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__); tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true, GOMP_MAP_VARS_OPENACC); gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__); + /* Initialize dynamic refcount. */ + tgt->list[0].key->dynamic_refcount = 1; + gomp_mutex_lock (&acc_dev->lock); tgt->prev = acc_dev->openacc.data_environ; acc_dev->openacc.data_environ = tgt; @@ -671,7 +770,8 @@ gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes, } void -gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async, int mapnum) +gomp_acc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async, + int finalize, int mapnum) { struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; @@ -679,6 +779,9 @@ gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async, int mapnum) struct target_mem_desc *t; int minrefs = (mapnum == 1) ? 2 : 3; + if (!acc_is_present (h, s)) + return; + gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, h, 1); @@ -693,40 +796,65 @@ gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async, int mapnum) t = n->tgt; - struct target_mem_desc *tp; + if (n->refcount < n->dynamic_refcount) + { + gomp_mutex_unlock (&acc_dev->lock); + gomp_fatal ("Dynamic reference counting assert fail\n"); + } - if (t->refcount == minrefs) + if (finalize) { - /* This is the last reference, so pull the descriptor off the - chain. This avoids gomp_unmap_vars via gomp_unmap_tgt from - freeing the device memory. */ - t->tgt_end = 0; - t->to_free = 0; + n->refcount -= n->dynamic_refcount; + n->dynamic_refcount = 0; + } + else if (n->dynamic_refcount) + { + n->dynamic_refcount--; + n->refcount--; + } - for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL; - tp = t, t = t->prev) + gomp_mutex_unlock (&acc_dev->lock); + + if (n->refcount == 0) + { + if (t->refcount == minrefs) { - if (n->tgt == t) + /* This is the last reference, so pull the descriptor off the + chain. This prevents gomp_unmap_vars via gomp_unmap_tgt from + freeing the device memory. */ + struct target_mem_desc *tp; + for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL; + tp = t, t = t->prev) { - if (tp) - tp->prev = t->prev; - else - acc_dev->openacc.data_environ = t->prev; - break; + if (n->tgt == t) + { + if (tp) + tp->prev = t->prev; + else + acc_dev->openacc.data_environ = t->prev; + break; + } } } - } - if (force_copyfrom) - t->list[0].copy_from = 1; + /* Set refcount to 1 to allow gomp_unmap_vars to unmap it. */ + n->refcount = 1; + t->refcount = minrefs; + for (size_t i = 0; i < t->list_count; i++) + if (t->list[i].key == n) + { + t->list[i].copy_from = force_copyfrom ? 1 : 0; + break; + } - gomp_mutex_unlock (&acc_dev->lock); + /* If running synchronously, unmap immediately. */ + if (async < acc_async_noval) + gomp_unmap_vars (t, true); + else + t->device_descr->openacc.register_async_cleanup_func (t, async); + } - /* If running synchronously, unmap immediately. */ - if (async_synchronous_p (async)) - gomp_unmap_vars (t, true); - else - t->device_descr->openacc.register_async_cleanup_func (t, async); + gomp_mutex_unlock (&acc_dev->lock); gomp_debug (0, " %s: mappings restored\n", __FUNCTION__); } diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index 9eae431..b80ace5 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -38,15 +38,68 @@ #include <stdarg.h> #include <assert.h> +/* Returns the number of mappings associated with the pointer or pset. PSET + have three mappings, whereas pointer have two. */ + static int -find_pset (int pos, size_t mapnum, unsigned short *kinds) +find_pointer (int pos, size_t mapnum, unsigned short *kinds) { if (pos + 1 >= mapnum) return 0; unsigned char kind = kinds[pos+1] & 0xff; - return kind == GOMP_MAP_TO_PSET; + if (kind == GOMP_MAP_TO_PSET) + return 3; + else if (kind == GOMP_MAP_POINTER) + return 2; + + return 0; +} + +/* Handle the mapping pair that are presented when a + deviceptr clause is used with Fortran. */ + +static void +handle_ftn_pointers (size_t mapnum, void **hostaddrs, size_t *sizes, + unsigned short *kinds) +{ + int i; + + for (i = 0; i < mapnum; i++) + { + unsigned short kind1 = kinds[i] & 0xff; + + /* Handle Fortran deviceptr clause. */ + if (kind1 == GOMP_MAP_FORCE_DEVICEPTR) + { + unsigned short kind2; + + if (i < (signed)mapnum - 1) + kind2 = kinds[i + 1] & 0xff; + else + kind2 = 0xffff; + + if (sizes[i] == sizeof (void *)) + continue; + + /* At this point, we're dealing with a Fortran deviceptr. + If the next element is not what we're expecting, then + this is an instance of where the deviceptr variable was + not used within the region and the pointer was removed + by the gimplifier. */ + if (kind2 == GOMP_MAP_POINTER + && sizes[i + 1] == 0 + && hostaddrs[i] == *(void **)hostaddrs[i + 1]) + { + kinds[i+1] = kinds[i]; + sizes[i+1] = sizeof (void *); + } + + /* Invalidate the entry. */ + hostaddrs[i] = NULL; + } + } } static void goacc_wait (int async, int num_waits, va_list *ap); @@ -88,6 +141,8 @@ GOACC_parallel_keyed (int device, void (*fn) (void *), thr = goacc_thread (); acc_dev = thr->dev; + handle_ftn_pointers (mapnum, hostaddrs, sizes, kinds); + /* Host fallback if "if" clause is false or if the current device is set to the host. */ if (host_fallback) @@ -183,10 +238,29 @@ GOACC_parallel_keyed (int device, void (*fn) (void *), async, dims, tgt); /* If running synchronously, unmap immediately. */ + bool copyfrom = true; if (async_synchronous_p (async)) gomp_unmap_vars (tgt, true); else - tgt->device_descr->openacc.register_async_cleanup_func (tgt, async); + { + bool async_unmap = false; + for (size_t i = 0; i < tgt->list_count; i++) + { + splay_tree_key k = tgt->list[i].key; + if (k && k->refcount == 1) + { + async_unmap = true; + break; + } + } + if (async_unmap) + tgt->device_descr->openacc.register_async_cleanup_func (tgt, async); + else + { + copyfrom = false; + gomp_unmap_vars (tgt, copyfrom); + } + } acc_dev->openacc.async_set_async_func (acc_async_sync); } @@ -286,6 +360,17 @@ GOACC_enter_exit_data (int device, size_t mapnum, va_end (ap); } + /* Determine whether "finalize" semantics apply to all mappings of this + OpenACC directive. */ + bool finalize = false; + if (mapnum > 0) + { + unsigned char kind = kinds[0] & 0xff; + if (kind == GOMP_MAP_DELETE + || kind == GOMP_MAP_FORCE_FROM) + finalize = true; + } + acc_dev->openacc.async_set_async_func (async); /* Determine if this is an "acc enter data". */ @@ -298,13 +383,17 @@ GOACC_enter_exit_data (int device, size_t mapnum, if (kind == GOMP_MAP_FORCE_ALLOC || kind == GOMP_MAP_FORCE_PRESENT - || kind == GOMP_MAP_FORCE_TO) + || kind == GOMP_MAP_FORCE_TO + || kind == GOMP_MAP_TO + || kind == GOMP_MAP_ALLOC) { data_enter = true; break; } - if (kind == GOMP_MAP_DELETE + if (kind == GOMP_MAP_RELEASE + || kind == GOMP_MAP_DELETE + || kind == GOMP_MAP_FROM || kind == GOMP_MAP_FORCE_FROM) break; @@ -312,31 +401,39 @@ GOACC_enter_exit_data (int device, size_t mapnum, kind); } + /* In c, non-pointers and arrays are represented by a single data clause. + Dynamically allocated arrays and subarrays are represented by a data + clause followed by an internal GOMP_MAP_POINTER. + + In fortran, scalars and not allocated arrays are represented by a + single data clause. Allocated arrays and subarrays have three mappings: + 1) the original data clause, 2) a PSET 3) a pointer to the array data. + */ + if (data_enter) { for (i = 0; i < mapnum; i++) { unsigned char kind = kinds[i] & 0xff; - /* Scan for PSETs. */ - int psets = find_pset (i, mapnum, kinds); + /* Scan for pointers and PSETs. */ + int pointer = find_pointer (i, mapnum, kinds); - if (!psets) + if (!pointer) { switch (kind) { - case GOMP_MAP_POINTER: - gomp_acc_insert_pointer (1, &hostaddrs[i], &sizes[i], - &kinds[i]); + case GOMP_MAP_ALLOC: + acc_present_or_create (hostaddrs[i], sizes[i]); break; case GOMP_MAP_FORCE_ALLOC: acc_create (hostaddrs[i], sizes[i]); break; - case GOMP_MAP_FORCE_PRESENT: + case GOMP_MAP_TO: acc_present_or_copyin (hostaddrs[i], sizes[i]); break; case GOMP_MAP_FORCE_TO: - acc_present_or_copyin (hostaddrs[i], sizes[i]); + acc_copyin (hostaddrs[i], sizes[i]); break; default: gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", @@ -346,12 +443,13 @@ GOACC_enter_exit_data (int device, size_t mapnum, } else { - gomp_acc_insert_pointer (3, &hostaddrs[i], &sizes[i], &kinds[i]); + gomp_acc_insert_pointer (pointer, &hostaddrs[i], + &sizes[i], &kinds[i]); /* Increment 'i' by two because OpenACC requires fortran arrays to be contiguous, so each PSET is associated with one of MAP_FORCE_ALLOC/MAP_FORCE_PRESET/MAP_FORCE_TO, and one MAP_POINTER. */ - i += 2; + i += pointer - 1; } } } @@ -360,22 +458,28 @@ GOACC_enter_exit_data (int device, size_t mapnum, { unsigned char kind = kinds[i] & 0xff; - int psets = find_pset (i, mapnum, kinds); + int pointer = find_pointer (i, mapnum, kinds); - if (!psets) + if (!pointer) { switch (kind) { - case GOMP_MAP_POINTER: - gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff) - == GOMP_MAP_FORCE_FROM, - async, 1); - break; + case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: - acc_delete (hostaddrs[i], sizes[i]); + if (acc_is_present (hostaddrs[i], sizes[i])) + { + if (finalize) + acc_delete_finalize (hostaddrs[i], sizes[i]); + else + acc_delete (hostaddrs[i], sizes[i]); + } break; + case GOMP_MAP_FROM: case GOMP_MAP_FORCE_FROM: - acc_copyout (hostaddrs[i], sizes[i]); + if (finalize) + acc_copyout_finalize (hostaddrs[i], sizes[i]); + else + acc_copyout (hostaddrs[i], sizes[i]); break; default: gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", @@ -385,10 +489,12 @@ GOACC_enter_exit_data (int device, size_t mapnum, } else { - gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff) - == GOMP_MAP_FORCE_FROM, async, 3); + bool copyfrom = (kind == GOMP_MAP_FORCE_FROM + || kind == GOMP_MAP_FROM); + gomp_acc_remove_pointer (hostaddrs[i], sizes[i], copyfrom, async, + finalize, pointer); /* See the above comment. */ - i += 2; + i += pointer - 1; } } @@ -447,6 +553,7 @@ GOACC_update (int device, size_t mapnum, acc_dev->openacc.async_set_async_func (async); + bool update_device = false; for (i = 0; i < mapnum; ++i) { unsigned char kind = kinds[i] & 0xff; @@ -457,11 +564,46 @@ GOACC_update (int device, size_t mapnum, case GOMP_MAP_TO_PSET: break; + case GOMP_MAP_ALWAYS_POINTER: + if (update_device) + { + /* Save the contents of the host pointer. */ + void *dptr = acc_deviceptr (hostaddrs[i-1]); + uintptr_t t = *(uintptr_t *) hostaddrs[i]; + + /* Update the contents of the host pointer to reflect + the value of the allocated device memory in the + previous pointer. */ + *(uintptr_t *) hostaddrs[i] = (uintptr_t)dptr; + acc_update_device (hostaddrs[i], sizeof (uintptr_t)); + + /* Restore the host pointer. */ + *(uintptr_t *) hostaddrs[i] = t; + update_device = false; + } + break; + + case GOMP_MAP_TO: + if (!acc_is_present (hostaddrs[i], sizes[i])) + { + update_device = false; + break; + } + /* Fallthru */ case GOMP_MAP_FORCE_TO: + update_device = true; acc_update_device (hostaddrs[i], sizes[i]); break; + case GOMP_MAP_FROM: + if (!acc_is_present (hostaddrs[i], sizes[i])) + { + update_device = false; + break; + } + /* Fallthru */ case GOMP_MAP_FORCE_FROM: + update_device = false; acc_update_self (hostaddrs[i], sizes[i]); break; @@ -522,6 +664,7 @@ GOACC_declare (int device, size_t mapnum, case GOMP_MAP_FORCE_FROM: case GOMP_MAP_FORCE_TO: case GOMP_MAP_POINTER: + case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i], &kinds[i], GOMP_ASYNC_SYNC, 0); @@ -543,7 +686,6 @@ GOACC_declare (int device, size_t mapnum, break; case GOMP_MAP_FROM: - kinds[i] = GOMP_MAP_FORCE_FROM; GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i], &kinds[i], GOMP_ASYNC_SYNC, 0); break; diff --git a/libgomp/openacc.f90 b/libgomp/openacc.f90 index d201d1d..84a8700 100644 --- a/libgomp/openacc.f90 +++ b/libgomp/openacc.f90 @@ -222,6 +222,24 @@ module openacc_internal type (*), dimension (..), contiguous :: a end subroutine + subroutine acc_copyout_finalize_32_h (a, len) + use iso_c_binding, only: c_int32_t + !GCC$ ATTRIBUTES NO_ARG_CHECK :: a + type (*), dimension (*) :: a + integer (c_int32_t) len + end subroutine + + subroutine acc_copyout_finalize_64_h (a, len) + use iso_c_binding, only: c_int64_t + !GCC$ ATTRIBUTES NO_ARG_CHECK :: a + type (*), dimension (*) :: a + integer (c_int64_t) len + end subroutine + + subroutine acc_copyout_finalize_array_h (a) + type (*), dimension (..), contiguous :: a + end subroutine + subroutine acc_delete_32_h (a, len) use iso_c_binding, only: c_int32_t !GCC$ ATTRIBUTES NO_ARG_CHECK :: a @@ -240,6 +258,24 @@ module openacc_internal type (*), dimension (..), contiguous :: a end subroutine + subroutine acc_delete_finalize_32_h (a, len) + use iso_c_binding, only: c_int32_t + !GCC$ ATTRIBUTES NO_ARG_CHECK :: a + type (*), dimension (*) :: a + integer (c_int32_t) len + end subroutine + + subroutine acc_delete_finalize_64_h (a, len) + use iso_c_binding, only: c_int64_t + !GCC$ ATTRIBUTES NO_ARG_CHECK :: a + type (*), dimension (*) :: a + integer (c_int64_t) len + end subroutine + + subroutine acc_delete_finalize_array_h (a) + type (*), dimension (..), contiguous :: a + end subroutine + subroutine acc_update_device_32_h (a, len) use iso_c_binding, only: c_int32_t !GCC$ ATTRIBUTES NO_ARG_CHECK :: a @@ -426,6 +462,14 @@ module openacc_internal integer (c_size_t), value :: len end subroutine + subroutine acc_copyout_finalize_l (a, len) & + bind (C, name = "acc_copyout_finalize") + use iso_c_binding, only: c_size_t + !GCC$ ATTRIBUTES NO_ARG_CHECK :: a + type (*), dimension (*) :: a + integer (c_size_t), value :: len + end subroutine + subroutine acc_delete_l (a, len) & bind (C, name = "acc_delete") use iso_c_binding, only: c_size_t @@ -434,6 +478,14 @@ module openacc_internal integer (c_size_t), value :: len end subroutine + subroutine acc_delete_finalize_l (a, len) & + bind (C, name = "acc_delete_finalize") + use iso_c_binding, only: c_size_t + !GCC$ ATTRIBUTES NO_ARG_CHECK :: a + type (*), dimension (*) :: a + integer (c_size_t), value :: len + end subroutine + subroutine acc_update_device_l (a, len) & bind (C, name = "acc_update_device") use iso_c_binding, only: c_size_t @@ -598,12 +650,24 @@ module openacc procedure :: acc_copyout_array_h end interface + interface acc_copyout_finalize + procedure :: acc_copyout_finalize_32_h + procedure :: acc_copyout_finalize_64_h + procedure :: acc_copyout_finalize_array_h + end interface + interface acc_delete procedure :: acc_delete_32_h procedure :: acc_delete_64_h procedure :: acc_delete_array_h end interface + interface acc_delete_finalize + procedure :: acc_delete_finalize_32_h + procedure :: acc_delete_finalize_64_h + procedure :: acc_delete_finalize_array_h + end interface + interface acc_update_device procedure :: acc_update_device_32_h procedure :: acc_update_device_64_h @@ -860,6 +924,30 @@ subroutine acc_copyout_array_h (a) call acc_copyout_l (a, sizeof (a)) end subroutine +subroutine acc_copyout_finalize_32_h (a, len) + use iso_c_binding, only: c_int32_t, c_size_t + use openacc_internal, only: acc_copyout_finalize_l + !GCC$ ATTRIBUTES NO_ARG_CHECK :: a + type (*), dimension (*) :: a + integer (c_int32_t) len + call acc_copyout_finalize_l (a, int (len, kind = c_size_t)) +end subroutine + +subroutine acc_copyout_finalize_64_h (a, len) + use iso_c_binding, only: c_int64_t, c_size_t + use openacc_internal, only: acc_copyout_finalize_l + !GCC$ ATTRIBUTES NO_ARG_CHECK :: a + type (*), dimension (*) :: a + integer (c_int64_t) len + call acc_copyout_finalize_l (a, int (len, kind = c_size_t)) +end subroutine + +subroutine acc_copyout_finalize_array_h (a) + use openacc_internal, only: acc_copyout_finalize_l + type (*), dimension (..), contiguous :: a + call acc_copyout_finalize_l (a, sizeof (a)) +end subroutine + subroutine acc_delete_32_h (a, len) use iso_c_binding, only: c_int32_t, c_size_t use openacc_internal, only: acc_delete_l @@ -884,6 +972,30 @@ subroutine acc_delete_array_h (a) call acc_delete_l (a, sizeof (a)) end subroutine +subroutine acc_delete_finalize_32_h (a, len) + use iso_c_binding, only: c_int32_t, c_size_t + use openacc_internal, only: acc_delete_finalize_l + !GCC$ ATTRIBUTES NO_ARG_CHECK :: a + type (*), dimension (*) :: a + integer (c_int32_t) len + call acc_delete_finalize_l (a, int (len, kind = c_size_t)) +end subroutine + +subroutine acc_delete_finalize_64_h (a, len) + use iso_c_binding, only: c_int64_t, c_size_t + use openacc_internal, only: acc_delete_finalize_l + !GCC$ ATTRIBUTES NO_ARG_CHECK :: a + type (*), dimension (*) :: a + integer (c_int64_t) len + call acc_delete_finalize_l (a, int (len, kind = c_size_t)) +end subroutine + +subroutine acc_delete_finalize_array_h (a) + use openacc_internal, only: acc_delete_finalize_l + type (*), dimension (..), contiguous :: a + call acc_delete_finalize_l (a, sizeof (a)) +end subroutine + subroutine acc_update_device_32_h (a, len) use iso_c_binding, only: c_int32_t, c_size_t use openacc_internal, only: acc_update_device_l diff --git a/libgomp/openacc.h b/libgomp/openacc.h index b857257..02a85a0 100644 --- a/libgomp/openacc.h +++ b/libgomp/openacc.h @@ -109,6 +109,12 @@ int acc_is_present (void *, size_t) __GOACC_NOTHROW; void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW; void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW; +/* Finalize versions of copyout/delete functions, specified in OpenACC 2.5. */ +void acc_copyout_finalize (void *, size_t) __GOACC_NOTHROW; +void acc_copyout_finalize_async (void *, size_t, int) __GOACC_NOTHROW; +void acc_delete_finalize (void *, size_t) __GOACC_NOTHROW; +void acc_delete_finalize_async (void *, size_t, int) __GOACC_NOTHROW; + /* CUDA-specific routines. */ void *acc_get_current_cuda_device (void) __GOACC_NOTHROW; void *acc_get_current_cuda_context (void) __GOACC_NOTHROW; diff --git a/libgomp/openacc_lib.h b/libgomp/openacc_lib.h index 5cf743c..737c582 100644 --- a/libgomp/openacc_lib.h +++ b/libgomp/openacc_lib.h @@ -273,6 +273,26 @@ end subroutine end interface + interface acc_copyout_finalize + subroutine acc_copyout_finalize_32_h (a, len) + use iso_c_binding, only: c_int32_t + !GCC$ ATTRIBUTES NO_ARG_CHECK :: a + type (*), dimension (*) :: a + integer (c_int32_t) len + end subroutine + + subroutine acc_copyout_finalize_64_h (a, len) + use iso_c_binding, only: c_int64_t + !GCC$ ATTRIBUTES NO_ARG_CHECK :: a + type (*), dimension (*) :: a + integer (c_int64_t) len + end subroutine + + subroutine acc_copyout_finalize_array_h (a) + type (*), dimension (..), contiguous :: a + end subroutine + end interface + interface acc_delete subroutine acc_delete_32_h (a, len) use iso_c_binding, only: c_int32_t @@ -293,6 +313,26 @@ end subroutine end interface + interface acc_delete_finalize + subroutine acc_delete_finalize_32_h (a, len) + use iso_c_binding, only: c_int32_t + !GCC$ ATTRIBUTES NO_ARG_CHECK :: a + type (*), dimension (*) :: a + integer (c_int32_t) len + end subroutine + + subroutine acc_delete_finalize_64_h (a, len) + use iso_c_binding, only: c_int64_t + !GCC$ ATTRIBUTES NO_ARG_CHECK :: a + type (*), dimension (*) :: a + integer (c_int64_t) len + end subroutine + + subroutine acc_delete_finalize_array_h (a) + type (*), dimension (..), contiguous :: a + end subroutine + end interface + interface acc_update_device subroutine acc_update_device_32_h (a, len) use iso_c_binding, only: c_int32_t diff --git a/libgomp/target.c b/libgomp/target.c index 509776d..dda041c 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -859,6 +859,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, tgt->list[i].offset = 0; tgt->list[i].length = k->host_end - k->host_start; k->refcount = 1; + k->dynamic_refcount = 0; tgt->refcount++; array->left = NULL; array->right = NULL; @@ -1011,6 +1012,23 @@ gomp_unmap_tgt (struct target_mem_desc *tgt) free (tgt); } +attribute_hidden bool +gomp_remove_var (struct gomp_device_descr *devicep, splay_tree_key k) +{ + bool is_tgt_unmapped = false; + splay_tree_remove (&devicep->mem_map, k); + if (k->link_key) + splay_tree_insert (&devicep->mem_map, (splay_tree_node) k->link_key); + if (k->tgt->refcount > 1) + k->tgt->refcount--; + else + { + is_tgt_unmapped = true; + gomp_unmap_tgt (k->tgt); + } + return is_tgt_unmapped; +} + /* Unmap variables described by TGT. If DO_COPYFROM is true, copy relevant variables back from device to host: if it is false, it is assumed that this has been done already. */ @@ -1059,16 +1077,7 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom) + tgt->list[i].offset), tgt->list[i].length); if (do_unmap) - { - splay_tree_remove (&devicep->mem_map, k); - if (k->link_key) - splay_tree_insert (&devicep->mem_map, - (splay_tree_node) k->link_key); - if (k->tgt->refcount > 1) - k->tgt->refcount--; - else - gomp_unmap_tgt (k->tgt); - } + gomp_remove_var (devicep, k); } if (tgt->refcount > 1) @@ -1298,17 +1307,7 @@ gomp_unload_image_from_device (struct gomp_device_descr *devicep, else { splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &k); - splay_tree_remove (&devicep->mem_map, n); - if (n->link_key) - { - if (n->tgt->refcount > 1) - n->tgt->refcount--; - else - { - is_tgt_unmapped = true; - gomp_unmap_tgt (n->tgt); - } - } + is_tgt_unmapped = gomp_remove_var (devicep, n); } } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c index 0ed53a4..fd3b77d 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c @@ -18,5 +18,3 @@ main (int argc, char *argv[]) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c index 00adf2a..0118b25 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c @@ -18,5 +18,3 @@ main (int argc, char *argv[]) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c index 04073e3..b346c69 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c @@ -18,5 +18,3 @@ main (int argc, char *argv[]) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "already mapped to" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c index cfa5cb2..e99ad33 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c @@ -16,5 +16,3 @@ main (int argc, char *argv[]) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "already mapped to" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c index e15c3fb..f8370c0 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c @@ -16,5 +16,3 @@ main (int argc, char *argv[]) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "already mapped to" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c index 4570c74..d7f4deb 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c @@ -16,5 +16,3 @@ main (int argc, char *argv[]) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "already mapped to" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c index 467cf39..54be595 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c @@ -16,5 +16,3 @@ main (int argc, char *argv[]) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "already mapped to" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c index f41431c..e5c0f9c 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c @@ -18,5 +18,3 @@ main (int argc, char *argv[]) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c index c81a78d..9a1c9d3 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c @@ -1,8 +1,5 @@ -/* Test if duplicate data mappings with acc_copy_in. */ +/* Test if acc_copyin has present_or_ and reference counting behavior. */ -/* { dg-do run { target openacc_nvidia_accel_selected } } */ - -#include <stdio.h> #include <stdlib.h> #include <openacc.h> @@ -21,15 +18,21 @@ main (int argc, char **argv) } (void) acc_copyin (h, N); - - fprintf (stderr, "CheCKpOInT\n"); (void) acc_copyin (h, N); + acc_copyout (h, N); + + if (!acc_is_present (h, N)) + abort (); + + acc_copyout (h, N); + +#if !ACC_MEM_SHARED + if (acc_is_present (h, N)) + abort (); +#endif + free (h); return 0; } - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+256\\\] already mapped to \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c index 5f00ccb..9b42dee 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c @@ -1,8 +1,5 @@ -/* Exercise acc_create and acc_delete on nvidia targets. */ +/* Exercise acc_create and acc_delete. */ -/* { dg-do run { target openacc_nvidia_accel_selected } } */ - -#include <stdio.h> #include <stdlib.h> #include <openacc.h> @@ -19,18 +16,23 @@ main (int argc, char **argv) if (!d) abort (); - fprintf (stderr, "CheCKpOInT\n"); d = acc_create (h, N); if (!d) abort (); acc_delete (h, N); + if (!acc_is_present (h, N)) + abort (); + + acc_delete (h, N); + +#if !ACC_MEM_SHARED + if (acc_is_present (h, N)) + abort (); +#endif + free (h); return 0; } - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+256\\\] already mapped to \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c index 1696fb6..9ec3453 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c @@ -127,7 +127,7 @@ main (int argc, char **argv) h[i] = i + 10; } - acc_copyout (h, S); + acc_copyout_finalize (h, S); d = NULL; if (!shared_mem) if (acc_is_present (h, S)) @@ -236,7 +236,7 @@ main (int argc, char **argv) abort (); } - acc_delete (h, S); + acc_delete_finalize (h, S); d = NULL; if (!shared_mem) if (acc_is_present (h, S)) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-83.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-83.c index 1c2e52b..51b7ee7 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-83.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-83.c @@ -5,21 +5,19 @@ #include <stdlib.h> #include <unistd.h> #include <openacc.h> -#include "timer.h" +#include <cuda.h> +#include <sys/time.h> int main (int argc, char **argv) { - float atime; CUstream stream; CUresult r; + struct timeval tv1, tv2; + time_t t1; acc_init (acc_device_nvidia); - (void) acc_get_device_num (acc_device_nvidia); - - init_timers (1); - stream = (CUstream) acc_get_cuda_stream (0); if (stream != NULL) abort (); @@ -34,22 +32,22 @@ main (int argc, char **argv) if (!acc_set_cuda_stream (0, stream)) abort (); - start_timer (0); + gettimeofday (&tv1, NULL); acc_wait_all_async (0); acc_wait (0); - atime = stop_timer (0); + gettimeofday (&tv2, NULL); - if (0.010 < atime) + t1 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec); + + if (t1 > 1000) { - fprintf (stderr, "actual time too long\n"); + fprintf (stderr, "too long\n"); abort (); } - fini_timers (); - acc_shutdown (acc_device_nvidia); exit (0); diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/data-5.f90 new file mode 100644 index 0000000..a8843de --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/data-5.f90 @@ -0,0 +1,56 @@ +! { dg-do run } +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } + +program refcount_test + use openacc + integer, allocatable :: h(:) + integer i, N + + N = 256 + allocate (h(N)) + + do i = 1, N + h(i) = i + end do + + !$acc enter data create (h(1:N)) + !$acc enter data copyin (h(1:N)) + !$acc enter data copyin (h(1:N)) + !$acc enter data copyin (h(1:N)) + + call acc_update_self (h) + do i = 1, N + if (h(i) .eq. i) c = c + 1 + end do + ! h[] should be filled with uninitialized device values, + ! abort if it's not. + if (c .eq. N) call abort + + h(:) = 0 + + !$acc parallel present (h(1:N)) + do i = 1, N + h(i) = 111 + end do + !$acc end parallel + + ! No actual copyout should happen. + call acc_copyout (h) + do i = 1, N + if (h(i) .ne. 0) call abort + end do + + !$acc exit data delete (h(1:N)) + + ! This should not actually be deleted yet. + if (acc_is_present (h) .eqv. .FALSE.) call abort + + !$acc exit data copyout (h(1:N)) finalize + + do i = 1, N + if (h(i) .ne. 111) call abort + end do + + if (acc_is_present (h) .eqv. .TRUE.) call abort + +end program refcount_test diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f index 9e99cc6..fab0ffc 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f @@ -14,5 +14,3 @@ END ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } -! { dg-shouldfail "" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f index b908a0c..bd03062 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f @@ -14,5 +14,3 @@ END ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } -! { dg-shouldfail "" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f index d93e1c5..60ea386 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f @@ -13,5 +13,3 @@ END ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "already mapped to" } -! { dg-shouldfail "" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f index ea76e05..2abdbf0 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f @@ -12,5 +12,3 @@ END ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "already mapped to" } -! { dg-shouldfail "" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f index 19df1f8..f361d8c 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f @@ -12,5 +12,3 @@ END ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "already mapped to" } -! { dg-shouldfail "" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f index 2bd1079..a864737 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f @@ -12,5 +12,3 @@ END ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "already mapped to" } -! { dg-shouldfail "" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f index 1342360..0d89328 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f @@ -12,5 +12,3 @@ END ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "already mapped to" } -! { dg-shouldfail "" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f index b206547..7a41f67 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f @@ -14,5 +14,3 @@ END ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } -! { dg-shouldfail "" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f b/libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f index 99e8f35..3f979eb 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f @@ -90,7 +90,7 @@ H(I) = I + 10 END DO - CALL ACC_COPYOUT (H, INT (SIZEOF (H), 4)) + CALL ACC_COPYOUT_FINALIZE (H, INT (SIZEOF (H), 4)) IF (.NOT. SHARED_MEM) THEN IF (ACC_IS_PRESENT (H, INT (SIZEOF (H), 8))) STOP 11 ENDIF @@ -163,7 +163,7 @@ IF (H(I) .NE. I + MERGE (18, 17, SHARED_MEM)) STOP 23 END DO - CALL ACC_DELETE (H) + CALL ACC_DELETE_FINALIZE (H) IF (.NOT. SHARED_MEM) THEN IF (ACC_IS_PRESENT (H, INT (SIZEOF (H), 4))) STOP 24 ENDIF diff --git a/libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f b/libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f index 514c04e..642552c 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f @@ -90,7 +90,7 @@ H(I) = I + 10 END DO - CALL ACC_COPYOUT (H, INT (SIZEOF (H), 4)) + CALL ACC_COPYOUT_FINALIZE (H, INT (SIZEOF (H), 4)) IF (.NOT. SHARED_MEM) THEN IF (ACC_IS_PRESENT (H, INT (SIZEOF (H), 8))) STOP 11 ENDIF @@ -163,7 +163,7 @@ IF (H(I) .NE. I + MERGE (18, 17, SHARED_MEM)) STOP 23 END DO - CALL ACC_DELETE (H) + CALL ACC_DELETE_FINALIZE (H) IF (.NOT. SHARED_MEM) THEN IF (ACC_IS_PRESENT (H, INT (SIZEOF (H), 4))) STOP 24 ENDIF |