From: Julian Brown Date: Wed, 18 Dec 2019 23:10:08 +0000 (+0000) Subject: Add OpenACC 2.6's no_create X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=a6163563f2ce502bd4ef444bd5de33570bb8eeb1;p=gcc.git Add OpenACC 2.6's no_create The clause makes any device code use the local memory address for each of the variables specified unless the given variable is already present on the current device. 2019-12-19 Julian Brown Maciej W. Rozycki Tobias Burnus Thomas Schwinge gcc/ * omp-low.c (lower_omp_target): Support GOMP_MAP_NO_ALLOC. * tree-pretty-print.c (dump_omp_clause): Likewise. gcc/c-family/ * c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_NO_CREATE. gcc/c/ * c-parser.c (c_parser_omp_clause_name): Support no_create. (c_parser_oacc_data_clause): Likewise. (c_parser_oacc_all_clauses): Likewise. (OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK) (OACC_PARALLEL_CLAUSE_MASK, OACC_SERIAL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_NO_CREATE. * c-typeck.c (handle_omp_array_sections): Support GOMP_MAP_NO_ALLOC. gcc/cp/ * parser.c (cp_parser_omp_clause_name): Support no_create. (cp_parser_oacc_data_clause): Likewise. (cp_parser_oacc_all_clauses): Likewise. (OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK) (OACC_PARALLEL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_NO_CREATE. * semantics.c (handle_omp_array_sections): Support no_create. gcc/fortran/ * gfortran.h (gfc_omp_map_op): Add OMP_MAP_NO_ALLOC. * openmp.c (omp_mask2): Add OMP_CLAUSE_NO_CREATE. (gfc_match_omp_clauses): Support no_create. (OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES) (OACC_DATA_CLAUSES): Add OMP_CLAUSE_NO_CREATE. * trans-openmp.c (gfc_trans_omp_clauses_1): Support OMP_MAP_NO_ALLOC. gcc/testsuite/ * gfortran.dg/goacc/common-block-1.f90: Add no_create-clause tests. * gfortran.dg/goacc/common-block-1.f90: Likewise. * gfortran.dg/goacc/data-clauses.f95: Likewise. * gfortran.dg/goacc/data-tree.f95: Likewise. * gfortran.dg/goacc/kernels-tree.f95: Likewise. * gfortran.dg/goacc/parallel-tree.f95: Likewise. include/ * gomp-constants.h (gomp_map_kind): Support GOMP_MAP_NO_ALLOC. libgomp/ * target.c (gomp_map_vars_async): Support GOMP_MAP_NO_ALLOC. * testsuite/libgomp.oacc-c-c++-common/no_create-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/no_create-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/no_create-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/no_create-4.c: New test. * testsuite/libgomp.oacc-c-c++-common/no_create-5.c: New test. * testsuite/libgomp.oacc-fortran/no_create-1.f90: New test. * testsuite/libgomp.oacc-fortran/no_create-2.f90: New test. * testsuite/libgomp.oacc-fortran/no_create-3.F90: New test. Reviewed-by: Thomas Schwinge Co-Authored-By: Maciej W. Rozycki Co-Authored-By: Thomas Schwinge Co-Authored-By: Tobias Burnus From-SVN: r279551 --- diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 0f1a0b6a2ea..825b0776333 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,11 @@ +2019-12-19 Julian Brown + Maciej W. Rozycki + Tobias Burnus + Thomas Schwinge + + * omp-low.c (lower_omp_target): Support GOMP_MAP_NO_ALLOC. + * tree-pretty-print.c (dump_omp_clause): Likewise. + 2019-12-18 Eric Botcazou * ira.c (ira): Use simple LRA algorithm when not optimizing. diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog index 966607f59c9..77d928a3078 100644 --- a/gcc/c-family/ChangeLog +++ b/gcc/c-family/ChangeLog @@ -1,3 +1,11 @@ +2019-12-19 Julian Brown + Maciej W. Rozycki + Tobias Burnus + Thomas Schwinge + + * c-pragma.h (pragma_omp_clause): Add + PRAGMA_OACC_CLAUSE_NO_CREATE. + 2019-12-17 Martin Sebor PR c++/61339 diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index bfe681bb430..3754c5fda45 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -154,6 +154,7 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_HOST, PRAGMA_OACC_CLAUSE_INDEPENDENT, + PRAGMA_OACC_CLAUSE_NO_CREATE, PRAGMA_OACC_CLAUSE_NUM_GANGS, PRAGMA_OACC_CLAUSE_NUM_WORKERS, PRAGMA_OACC_CLAUSE_PRESENT, diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index 5371e9cc186..f4a088a9511 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,17 @@ +2019-12-19 Julian Brown + Maciej W. Rozycki + Tobias Burnus + Thomas Schwinge + + * c-parser.c (c_parser_omp_clause_name): Support no_create. + (c_parser_oacc_data_clause): Likewise. + (c_parser_oacc_all_clauses): Likewise. + (OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK) + (OACC_PARALLEL_CLAUSE_MASK, OACC_SERIAL_CLAUSE_MASK): Add + PRAGMA_OACC_CLAUSE_NO_CREATE. + * c-typeck.c (handle_omp_array_sections): Support + GOMP_MAP_NO_ALLOC. + 2019-12-09 David Malcolm * c-objc-common.c (range_label_for_type_mismatch::get_text): diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index bfe56998996..9b8008816d2 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -12650,7 +12650,9 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_MERGEABLE; break; case 'n': - if (!strcmp ("nogroup", p)) + if (!strcmp ("no_create", p)) + result = PRAGMA_OACC_CLAUSE_NO_CREATE; + else if (!strcmp ("nogroup", p)) result = PRAGMA_OMP_CLAUSE_NOGROUP; else if (!strcmp ("nontemporal", p)) result = PRAGMA_OMP_CLAUSE_NONTEMPORAL; @@ -13113,7 +13115,10 @@ 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 ( variable-list ) + + OpenACC 2.6: + no_create ( variable-list ) */ static tree c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, @@ -13149,6 +13154,9 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_LINK: kind = GOMP_MAP_LINK; break; + case PRAGMA_OACC_CLAUSE_NO_CREATE: + kind = GOMP_MAP_IF_PRESENT; + break; case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; @@ -15947,6 +15955,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "link"; break; + case PRAGMA_OACC_CLAUSE_NO_CREATE: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "no_create"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: clauses = c_parser_oacc_single_int_clause (parser, OMP_CLAUSE_NUM_GANGS, @@ -16415,6 +16427,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_NO_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)) static tree @@ -16747,6 +16760,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \ | (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) \ @@ -16762,6 +16776,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \ @@ -16780,6 +16795,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 36aedc063d2..ce5e6495fb1 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13422,6 +13422,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) switch (OMP_CLAUSE_MAP_KIND (c)) { case GOMP_MAP_ALLOC: + case GOMP_MAP_IF_PRESENT: case GOMP_MAP_TO: case GOMP_MAP_FROM: case GOMP_MAP_TOFROM: diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index 90a6d18a7e9..05ca5ec72a2 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,15 @@ +2019-12-19 Julian Brown + Maciej W. Rozycki + Tobias Burnus + Thomas Schwinge + + * parser.c (cp_parser_omp_clause_name): Support no_create. + (cp_parser_oacc_data_clause): Likewise. + (cp_parser_oacc_all_clauses): Likewise. + (OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK) + (OACC_PARALLEL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_NO_CREATE. + * semantics.c (handle_omp_array_sections): Support no_create. + 2019-12-18 Paolo Carlini * typeck.c (cxx_sizeof_or_alignof_type): Add location_t parameter diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 1f7526ea6e0..ce2e4b5b125 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -33622,7 +33622,9 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_MERGEABLE; break; case 'n': - if (!strcmp ("nogroup", p)) + if (!strcmp ("no_create", p)) + result = PRAGMA_OACC_CLAUSE_NO_CREATE; + else if (!strcmp ("nogroup", p)) result = PRAGMA_OMP_CLAUSE_NOGROUP; else if (!strcmp ("nontemporal", p)) result = PRAGMA_OMP_CLAUSE_NONTEMPORAL; @@ -33988,7 +33990,10 @@ 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 ( variable-list ) + + OpenACC 2.6: + no_create ( variable-list ) */ static tree cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, @@ -34024,6 +34029,9 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_LINK: kind = GOMP_MAP_LINK; break; + case PRAGMA_OACC_CLAUSE_NO_CREATE: + kind = GOMP_MAP_IF_PRESENT; + break; case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; @@ -36586,6 +36594,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "link"; break; + case PRAGMA_OACC_CLAUSE_NO_CREATE: + clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "no_create"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: code = OMP_CLAUSE_NUM_GANGS; c_name = "num_gangs"; @@ -40391,6 +40403,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_NO_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) ) static tree @@ -40712,6 +40725,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \ | (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) \ @@ -40726,8 +40740,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \ | (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) \ @@ -40745,6 +40760,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 1d664af7150..69010dc85cd 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5292,6 +5292,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) switch (OMP_CLAUSE_MAP_KIND (c)) { case GOMP_MAP_ALLOC: + case GOMP_MAP_IF_PRESENT: case GOMP_MAP_TO: case GOMP_MAP_FROM: case GOMP_MAP_TOFROM: diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog index b3a5b0d9a57..e45ec22bd9e 100644 --- a/gcc/fortran/ChangeLog +++ b/gcc/fortran/ChangeLog @@ -1,3 +1,16 @@ +2019-12-19 Julian Brown + Maciej W. Rozycki + Tobias Burnus + Thomas Schwinge + + * gfortran.h (gfc_omp_map_op): Add OMP_MAP_NO_ALLOC. + * openmp.c (omp_mask2): Add OMP_CLAUSE_NO_CREATE. + (gfc_match_omp_clauses): Support no_create. + (OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES) + (OACC_DATA_CLAUSES): Add OMP_CLAUSE_NO_CREATE. + * trans-openmp.c (gfc_trans_omp_clauses_1): Support + OMP_MAP_NO_ALLOC. + 2019-12-18 Harald Anlauf PR fortran/70853 diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index f4a2b99bdc4..3907d1407ac 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1192,6 +1192,7 @@ enum gfc_omp_depend_op enum gfc_omp_map_op { OMP_MAP_ALLOC, + OMP_MAP_IF_PRESENT, OMP_MAP_TO, OMP_MAP_FROM, OMP_MAP_TOFROM, diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index dc0521b40f0..576003d7ff8 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -807,6 +807,7 @@ enum omp_mask2 OMP_CLAUSE_COPY, OMP_CLAUSE_COPYOUT, OMP_CLAUSE_CREATE, + OMP_CLAUSE_NO_CREATE, OMP_CLAUSE_PRESENT, OMP_CLAUSE_DEVICEPTR, OMP_CLAUSE_GANG, @@ -1445,6 +1446,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, } break; case 'n': + if ((mask & OMP_CLAUSE_NO_CREATE) + && gfc_match ("no_create ( ") == MATCH_YES + && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], + OMP_MAP_IF_PRESENT, true)) + continue; if ((mask & OMP_CLAUSE_NOGROUP) && !c->nogroup && gfc_match ("nogroup") == MATCH_YES) @@ -1955,25 +1961,25 @@ 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_DEVICEPTR \ - | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \ - | OMP_CLAUSE_WAIT) + | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_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_DEFAULT \ - | OMP_CLAUSE_WAIT) + | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \ + | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT) #define OACC_SERIAL_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_REDUCTION \ | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ - | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \ - | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \ - | OMP_CLAUSE_WAIT) + | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \ + | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \ + | 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_NO_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 \ @@ -2509,7 +2515,7 @@ cleanup: #define OMP_TASKLOOP_CLAUSES \ (omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \ | OMP_CLAUSE_LASTPRIVATE | OMP_CLAUSE_SHARED | OMP_CLAUSE_IF \ - | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_UNTIED | OMP_CLAUSE_FINAL \ + | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_UNTIED | OMP_CLAUSE_FINAL \ | OMP_CLAUSE_MERGEABLE | OMP_CLAUSE_PRIORITY | OMP_CLAUSE_GRAINSIZE \ | OMP_CLAUSE_NUM_TASKS | OMP_CLAUSE_COLLAPSE | OMP_CLAUSE_NOGROUP) #define OMP_TARGET_CLAUSES \ @@ -2531,7 +2537,7 @@ cleanup: | OMP_CLAUSE_FROM | OMP_CLAUSE_DEPEND | OMP_CLAUSE_NOWAIT) #define OMP_TEAMS_CLAUSES \ (omp_mask (OMP_CLAUSE_NUM_TEAMS) | OMP_CLAUSE_THREAD_LIMIT \ - | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \ + | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \ | OMP_CLAUSE_SHARED | OMP_CLAUSE_REDUCTION) #define OMP_DISTRIBUTE_CLAUSES \ (omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \ diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index b6da7b983d5..7153491a460 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2624,6 +2624,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, case OMP_MAP_ALLOC: OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC); break; + case OMP_MAP_IF_PRESENT: + OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_IF_PRESENT); + break; case OMP_MAP_TO: OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO); break; diff --git a/gcc/omp-low.c b/gcc/omp-low.c index d422c205836..deed83b8c33 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -11431,6 +11431,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_STRUCT: case GOMP_MAP_ALWAYS_POINTER: break; + case GOMP_MAP_IF_PRESENT: case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_TO: case GOMP_MAP_FORCE_FROM: @@ -11842,6 +11843,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) switch (tkind) { case GOMP_MAP_ALLOC: + case GOMP_MAP_IF_PRESENT: case GOMP_MAP_TO: case GOMP_MAP_FROM: case GOMP_MAP_TOFROM: diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 340f9096520..55a550cb56f 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,15 @@ +2019-12-19 Julian Brown + Maciej W. Rozycki + Tobias Burnus + Thomas Schwinge + + * gfortran.dg/goacc/common-block-1.f90: Add no_create-clause tests. + * gfortran.dg/goacc/common-block-1.f90: Likewise. + * gfortran.dg/goacc/data-clauses.f95: Likewise. + * gfortran.dg/goacc/data-tree.f95: Likewise. + * gfortran.dg/goacc/kernels-tree.f95: Likewise. + * gfortran.dg/goacc/parallel-tree.f95: Likewise. + 2019-12-18 Paolo Carlini * g++.dg/diagnostic/alignof2.C: New. diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 index 228637f5883..6df5aa65e70 100644 --- a/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 @@ -51,6 +51,9 @@ program test !$acc data pcopyout(/blockA/, /blockB/, e, v) !$acc end data + !$acc data no_create(/blockA/, /blockB/, e, v) + !$acc end data + !$acc parallel private(/blockA/, /blockB/, e, v) !$acc end parallel diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 index 5d49f6195b8..30c87a91f36 100644 --- a/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 @@ -39,6 +39,9 @@ program test !$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } !$acc end data + !$acc data no_create(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } + !$acc end data + !$acc parallel private(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" } !$acc end parallel diff --git a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 index b94214e8b63..30930a0cf1c 100644 --- a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 @@ -111,6 +111,27 @@ contains !$acc end data + !$acc parallel no_create (tip) ! { dg-error "POINTER" } + !$acc end parallel + !$acc parallel no_create (tia) ! { dg-error "ALLOCATABLE" } + !$acc end parallel + !$acc parallel deviceptr (i) no_create (i) ! { dg-error "multiple clauses" } + !$acc end parallel + !$acc parallel copy (i) no_create (i) ! { dg-error "multiple clauses" } + !$acc end parallel + !$acc parallel copyin (i) no_create (i) ! { dg-error "multiple clauses" } + !$acc end parallel + !$acc parallel copyout (i) no_create (i) ! { dg-error "multiple clauses" } + !$acc end parallel + + !$acc parallel no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa) + !$acc end parallel + !$acc kernels no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa) + !$acc end kernels + !$acc data no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa) + !$acc end data + + !$acc parallel present (tip) ! { dg-error "POINTER" } !$acc end parallel !$acc parallel present (tia) ! { dg-error "ALLOCATABLE" } diff --git a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 index f16d62cce69..454417d6a05 100644 --- a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 @@ -7,6 +7,7 @@ program test logical :: l = .true. !$acc data if(l) copy(i), copyin(j), copyout(k), create(m) & + !$acc no_create(n) & !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) & !$acc deviceptr(u) !$acc end data @@ -19,7 +20,7 @@ end program test ! { 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\\(no_alloc:n\\)" 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" } } ! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 index a70f1e737bd..5583ffb4d04 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 @@ -8,6 +8,7 @@ program test !$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) & !$acc copy(i), copyin(j), copyout(k), create(m) & + !$acc no_create(n) & !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) & !$acc deviceptr(u) !$acc end kernels @@ -25,7 +26,7 @@ end program test ! { 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\\(no_alloc:n\\)" 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" } } ! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 index 2697bb79e7f..e33653bdd78 100644 --- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 @@ -9,6 +9,7 @@ program test !$acc parallel if(l) async num_gangs(i) num_workers(i) vector_length(i) & !$acc reduction(max:q), copy(i), copyin(j), copyout(k), create(m) & + !$acc no_create(n) & !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) & !$acc deviceptr(u), private(v), firstprivate(w) !$acc end parallel @@ -28,7 +29,7 @@ end program test ! { 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\\(no_alloc:n\\)" 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" } } ! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } } diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 1cf7a912133..603617358ae 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -788,6 +788,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_POINTER: pp_string (pp, "alloc"); break; + case GOMP_MAP_IF_PRESENT: + pp_string (pp, "no_alloc"); + break; case GOMP_MAP_TO: case GOMP_MAP_TO_PSET: pp_string (pp, "to"); diff --git a/include/ChangeLog b/include/ChangeLog index 6364ab4de61..faeb5c4b0ee 100644 --- a/include/ChangeLog +++ b/include/ChangeLog @@ -1,3 +1,10 @@ +2019-12-19 Julian Brown + Maciej W. Rozycki + Tobias Burnus + Thomas Schwinge + + * gomp-constants.h (gomp_map_kind): Support GOMP_MAP_NO_ALLOC. + 2019-11-16 Tim Ruehsen * demangle.h (struct demangle_component): Add member diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 9e356cdfeec..79c5de38db5 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -75,6 +75,8 @@ enum gomp_map_kind GOMP_MAP_DEVICE_RESIDENT = (GOMP_MAP_FLAG_SPECIAL_1 | 1), /* OpenACC link. */ GOMP_MAP_LINK = (GOMP_MAP_FLAG_SPECIAL_1 | 2), + /* Use device data if present, fall back to host address otherwise. */ + GOMP_MAP_IF_PRESENT = (GOMP_MAP_FLAG_SPECIAL_1 | 3), /* Do not map, copy bits for firstprivate instead. */ GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0), /* Similarly, but store the value in the pointer rather than diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 779f276ca99..e1f6bb4c622 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,18 @@ +2019-12-19 Julian Brown + Maciej W. Rozycki + Tobias Burnus + Thomas Schwinge + + * target.c (gomp_map_vars_async): Support GOMP_MAP_NO_ALLOC. + * testsuite/libgomp.oacc-c-c++-common/no_create-1.c: New test. + * testsuite/libgomp.oacc-c-c++-common/no_create-2.c: New test. + * testsuite/libgomp.oacc-c-c++-common/no_create-3.c: New test. + * testsuite/libgomp.oacc-c-c++-common/no_create-4.c: New test. + * testsuite/libgomp.oacc-c-c++-common/no_create-5.c: New test. + * testsuite/libgomp.oacc-fortran/no_create-1.f90: New test. + * testsuite/libgomp.oacc-fortran/no_create-2.f90: New test. + * testsuite/libgomp.oacc-fortran/no_create-3.F90: New test. + 2019-12-18 Thomas Schwinge * oacc-mem.c (goacc_enter_data): Refactor, so that it can be diff --git a/libgomp/target.c b/libgomp/target.c index 67cd80a3c35..d83b353d13e 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -707,6 +707,21 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, { tgt->list[i].key = NULL; + if ((kind & typemask) == GOMP_MAP_IF_PRESENT) + { + /* Not present, hence, skip entry - including its MAP_POINTER, + when existing. */ + tgt->list[i].offset = 0; + if (i + 1 < mapnum + && ((typemask & get_kind (short_mapkind, kinds, i + 1)) + == GOMP_MAP_POINTER)) + { + ++i; + tgt->list[i].key = NULL; + tgt->list[i].offset = 0; + } + continue; + } size_t align = (size_t) 1 << (kind >> rshift); not_found_cnt++; if (tgt_align < align) @@ -893,6 +908,14 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start - n->host_start; continue; + case GOMP_MAP_IF_PRESENT: + /* Not present - otherwise handled above. Skip over its + MAP_POINTER as well. */ + if (i + 1 < mapnum + && ((typemask & get_kind (short_mapkind, kinds, i + 1)) + == GOMP_MAP_POINTER)) + ++i; + continue; default: break; } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c new file mode 100644 index 00000000000..22e0c20cce9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c @@ -0,0 +1,49 @@ +/* Test 'no_create' clause on compute construct, with data present on the + device. */ + +#include +#include +#include + +#define N 128 + +int +main (int argc, char *argv[]) +{ + int var; + int *arr = (int *) malloc (N * sizeof (*arr)); + int *devptr[2]; + + acc_copyin (&var, sizeof (var)); + acc_copyin (arr, N * sizeof (*arr)); + +#pragma acc parallel no_create(var, arr[0:N]) copyout(devptr) + { + devptr[0] = &var; + devptr[1] = &arr[2]; + } + + if (acc_hostptr (devptr[0]) != (void *) &var) + __builtin_abort (); + if (acc_hostptr (devptr[1]) != (void *) &arr[2]) + __builtin_abort (); + + acc_delete (&var, sizeof (var)); + acc_delete (arr, N * sizeof (*arr)); + +#if ACC_MEM_SHARED + if (devptr[0] != &var) + __builtin_abort (); + if (devptr[1] != &arr[2]) + __builtin_abort (); +#else + if (devptr[0] == &var) + __builtin_abort (); + if (devptr[1] == &arr[2]) + __builtin_abort (); +#endif + + free (arr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c new file mode 100644 index 00000000000..fbd01a25956 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c @@ -0,0 +1,30 @@ +/* Test 'no_create' clause on compute construct, with data not present on the + device. */ + +#include +#include + +#define N 128 + +int +main (int argc, char *argv[]) +{ + int var; + int *arr = (int *) malloc (N * sizeof (*arr)); + int *devptr[2]; + +#pragma acc parallel no_create(var, arr[0:N]) copyout(devptr) + { + devptr[0] = &var; + devptr[1] = &arr[2]; + } + + if (devptr[0] != &var) + __builtin_abort (); + if (devptr[1] != &arr[2]) + __builtin_abort (); + + free (arr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c new file mode 100644 index 00000000000..18466b88b5c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c @@ -0,0 +1,25 @@ +#include /* For FLT_EPSILON. */ +#include /* For fabs. */ +#include /* For abort. */ + + +int main() +{ +#define N 100 + float b[N]; + float c[N]; + +#pragma acc enter data create(b) + +#pragma acc parallel loop no_create(b) no_create(c) + for (int i = 0; i < N; ++i) + b[i] = i; + +#pragma acc exit data copyout(b) + + for (int i = 0; i < N; ++i) + if (fabs (b[i] - i) > 10.0*FLT_EPSILON) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-4.c new file mode 100644 index 00000000000..963cb3a68f6 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-4.c @@ -0,0 +1,82 @@ +/* Test 'no_create' clause on 'data' construct and nested compute construct, + with data present on the device. */ + +#include +#include +#include + +#define N 128 + +int +main (int argc, char *argv[]) +{ + int var; + int *arr = (int *) malloc (N * sizeof (*arr)); + int *devptr[2]; + + acc_copyin (&var, sizeof (var)); + acc_copyin (arr, N * sizeof (*arr)); + +#pragma acc data no_create(var, arr[0:N]) + { + devptr[0] = (int *) acc_deviceptr (&var); + devptr[1] = (int *) acc_deviceptr (&arr[2]); + + if (devptr[0] == NULL) + __builtin_abort (); + if (devptr[1] == NULL) + __builtin_abort (); + + if (acc_hostptr (devptr[0]) != (void *) &var) + __builtin_abort (); + if (acc_hostptr (devptr[1]) != (void *) &arr[2]) + __builtin_abort (); + +#if ACC_MEM_SHARED + if (devptr[0] != &var) + __builtin_abort (); + if (devptr[1] != &arr[2]) + __builtin_abort (); +#else + if (devptr[0] == &var) + __builtin_abort (); + if (devptr[1] == &arr[2]) + __builtin_abort (); +#endif + +#pragma acc parallel copyout(devptr) + { + devptr[0] = &var; + devptr[1] = &arr[2]; + } + + if (devptr[0] == NULL) + __builtin_abort (); + if (devptr[1] == NULL) + __builtin_abort (); + + if (acc_hostptr (devptr[0]) != (void *) &var) + __builtin_abort (); + if (acc_hostptr (devptr[1]) != (void *) &arr[2]) + __builtin_abort (); + +#if ACC_MEM_SHARED + if (devptr[0] != &var) + __builtin_abort (); + if (devptr[1] != &arr[2]) + __builtin_abort (); +#else + if (devptr[0] == &var) + __builtin_abort (); + if (devptr[1] == &arr[2]) + __builtin_abort (); +#endif + } + + acc_delete (&var, sizeof (var)); + acc_delete (arr, N * sizeof (*arr)); + + free (arr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-5.c new file mode 100644 index 00000000000..6f0ace501cf --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-5.c @@ -0,0 +1,49 @@ +/* Test 'no_create' clause on 'data' construct and nested compute construct, + with data not present on the device. */ + +#include +#include +#include + +#define N 128 + +int +main (int argc, char *argv[]) +{ + int var; + int *arr = (int *) malloc (N * sizeof (*arr)); + int *devptr[2]; + +#pragma acc data no_create(var, arr[0:N]) + { + devptr[0] = (int *) acc_deviceptr (&var); + devptr[1] = (int *) acc_deviceptr (&arr[2]); + +#if ACC_MEM_SHARED + if (devptr[0] == NULL) + __builtin_abort (); + if (devptr[1] == NULL) + __builtin_abort (); +#else + if (devptr[0] != NULL) + __builtin_abort (); + if (devptr[1] != NULL) + __builtin_abort (); +#endif + +#pragma acc parallel copyout(devptr) // TODO implicit 'copy(var)' -- huh?! + { + devptr[0] = &var; + devptr[1] = &arr[2]; + } + + if (devptr[0] != &var) + __builtin_abort (); // { dg-xfail-run-if "TODO" { *-*-* } { "-DACC_MEM_SHARED=0" } } + if (devptr[1] != &arr[2]) + __builtin_abort (); + } + + free (arr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90 new file mode 100644 index 00000000000..4a1d5da98aa --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90 @@ -0,0 +1,39 @@ +! { dg-do run } + +! Test no_create clause with data construct when data is present/not present. + +program no_create + use openacc + implicit none + logical :: shared_memory + integer, parameter :: n = 512 + integer :: myvar, myarr(n) + integer i + + shared_memory = .false. + !$acc kernels copyin (shared_memory) + shared_memory = .true. + !$acc end kernels + + myvar = 77 + do i = 1, n + myarr(i) = 0 + end do + + !$acc data no_create (myvar, myarr) + if (acc_is_present (myvar) .neqv. shared_memory) stop 10 + if (acc_is_present (myarr) .neqv. shared_memory) stop 11 + !$acc end data + + !$acc enter data copyin (myvar, myarr) + !$acc data no_create (myvar, myarr) + if (acc_is_present (myvar) .eqv. .false.) stop 20 + if (acc_is_present (myarr) .eqv. .false.) stop 21 + !$acc end data + !$acc exit data copyout (myvar, myarr) + + if (myvar .ne. 77) stop 30 + do i = 1, n + if (myarr(i) .ne. 0) stop 31 + end do +end program no_create diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90 new file mode 100644 index 00000000000..0b11f454aca --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90 @@ -0,0 +1,90 @@ +! { dg-do run } + +! Test no_create clause with data/parallel constructs. + +program no_create + use openacc + implicit none + logical :: shared_memory + integer, parameter :: n = 512 + integer :: myvar, myarr(n) + integer i + + shared_memory = .false. + !$acc kernels copyin (shared_memory) + shared_memory = .true. + !$acc end kernels + + myvar = 55 + do i = 1, n + myarr(i) = 0 + end do + + call do_on_target(myvar, n, myarr) + + if (shared_memory) then + if (myvar .ne. 44) stop 10 + else + if (myvar .ne. 33) stop 11 + end if + do i = 1, n + if (shared_memory) then + if (myarr(i) .ne. i * 2) stop 20 + else + if (myarr(i) .ne. i) stop 21 + end if + end do + + myvar = 55 + do i = 1, n + myarr(i) = 0 + end do + + !$acc enter data copyin(myvar, myarr) + call do_on_target(myvar, n, myarr) + !$acc exit data copyout(myvar, myarr) + + if (myvar .ne. 44) stop 30 + do i = 1, n + if (myarr(i) .ne. i * 2) stop 31 + end do +end program no_create + +subroutine do_on_target (var, n, arr) + use openacc + implicit none + integer :: var, n, arr(n) + integer :: i + +!$acc data no_create (var, arr) + +if (acc_is_present(var)) then + ! The no_create clause is meant for partially shared-memory machines. This + ! test is written to work on non-shared-memory machines, though this is not + ! necessarily a useful way to use the no_create clause in practice. + + !$acc parallel !no_create (var) + var = 44 + !$acc end parallel +else + var = 33 +end if +if (acc_is_present(arr)) then + ! The no_create clause is meant for partially shared-memory machines. This + ! test is written to work on non-shared-memory machines, though this is not + ! necessarily a useful way to use the no_create clause in practice. + + !$acc parallel loop !no_create (arr) + do i = 1, n + arr(i) = i * 2 + end do + !$acc end parallel loop +else + do i = 1, n + arr(i) = i + end do +end if + +!$acc end data + +end subroutine do_on_target diff --git a/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.F90 b/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.F90 new file mode 100644 index 00000000000..4362688e579 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/no_create-3.F90 @@ -0,0 +1,39 @@ +! { dg-do run } + +program main + use iso_c_binding, only: c_sizeof + use openacc, only: acc_is_present + implicit none + integer i + integer, parameter :: n = 100 + real*4 b(n), c(n) + real :: d(n), e(n) + common /BLOCK/ d, e + + !$acc enter data create(b) create(d) + + if (.not. acc_is_present(b, c_sizeof(b))) stop 1 + if (.not. acc_is_present(d, c_sizeof(d))) stop 2 +#if !ACC_MEM_SHARED + if (acc_is_present(c, 1) .or. acc_is_present(c, c_sizeof(c))) stop 3 + if (acc_is_present(e, 1) .or. acc_is_present(e, c_sizeof(d))) stop 4 +#endif + + !$acc parallel loop no_create(b) no_create(c) no_create(/BLOCK/) + do i = 1, n + b(i) = i + d(i) = -i + end do + !$acc end parallel loop + + if (.not. acc_is_present(b, c_sizeof(b))) stop 5 + if (.not. acc_is_present(d, c_sizeof(d))) stop 6 +#if !ACC_MEM_SHARED + if (acc_is_present(c, 1) .or. acc_is_present(c, c_sizeof(c))) stop 7 + if (acc_is_present(e, 1) .or. acc_is_present(e, c_sizeof(e))) stop 8 +#endif + + !$acc exit data copyout(b) copyout(d) + if (any(abs(b - [(real(i), i = 1, n)]) > 10*epsilon(b))) stop 9 + if (any(abs(d - [(real(-i), i = 1, n)]) > 10*epsilon(d))) stop 10 +end program main