From 5677b75c32c836da0fd19eca9d8b692574f73d3e Mon Sep 17 00:00:00 2001 From: "Maciej W. Rozycki" Date: Thu, 20 Dec 2018 14:10:19 +0000 Subject: [PATCH] Add OpenACC 2.6 `no_create' clause support 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. 2018-12-19 Julian Brown Maciej W. Rozycki 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, OACC_SERIAL_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_SERIAL_CLAUSES, OACC_DATA_CLAUSES): Add OMP_CLAUSE_NO_CREATE. * trans-openmp.c (gfc_trans_omp_clauses_1): Support OMP_MAP_NO_ALLOC. 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/nocreate-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/nocreate-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/nocreate-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/nocreate-4.c: New test. * testsuite/libgomp.oacc-fortran/nocreate-1.f90: New test. * testsuite/libgomp.oacc-fortran/nocreate-2.f90: New test. (cherry picked from openacc-gcc-9-branch commit 8e74c2ec2b90819c995444370e742864a685209f) --- gcc/ChangeLog.omp | 6 ++ gcc/c-family/ChangeLog.omp | 6 ++ gcc/c-family/c-pragma.h | 1 + gcc/c/ChangeLog.omp | 12 ++++ gcc/c/c-parser.c | 20 +++++- gcc/c/c-typeck.c | 1 + gcc/cp/ChangeLog.omp | 11 ++++ gcc/cp/parser.c | 20 +++++- gcc/cp/semantics.c | 1 + gcc/fortran/ChangeLog.omp | 12 ++++ gcc/fortran/gfortran.h | 1 + gcc/fortran/openmp.c | 21 ++++--- gcc/fortran/trans-openmp.c | 3 + gcc/omp-low.c | 2 + gcc/tree-pretty-print.c | 3 + include/ChangeLog.omp | 5 ++ include/gomp-constants.h | 2 + libgomp/ChangeLog.omp | 12 ++++ libgomp/target.c | 53 ++++++++++++++++ .../libgomp.oacc-c-c++-common/nocreate-1.c | 40 ++++++++++++ .../libgomp.oacc-c-c++-common/nocreate-2.c | 28 +++++++++ .../libgomp.oacc-c-c++-common/nocreate-3.c | 38 ++++++++++++ .../libgomp.oacc-c-c++-common/nocreate-4.c | 42 +++++++++++++ .../libgomp.oacc-fortran/nocreate-1.f90 | 29 +++++++++ .../libgomp.oacc-fortran/nocreate-2.f90 | 61 +++++++++++++++++++ 25 files changed, 419 insertions(+), 11 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-3.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-4.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90 diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 43370cf90f6..8477373c1c6 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,9 @@ +2018-12-19 Julian Brown + Maciej W. Rozycki + + * omp-low.c (lower_omp_target): Support GOMP_MAP_NO_ALLOC. + * tree-pretty-print.c (dump_omp_clause): Likewise. + 2018-12-20 Maciej W. Rozycki * gimple.h (gf_mask): Add GF_OMP_TARGET_KIND_OACC_SERIAL diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp index f0ea8c5ce3f..40b8e3b9dff 100644 --- a/gcc/c-family/ChangeLog.omp +++ b/gcc/c-family/ChangeLog.omp @@ -1,3 +1,9 @@ +2018-12-19 Julian Brown + Maciej W. Rozycki + + * c-pragma.h (pragma_omp_clause): Add + PRAGMA_OACC_CLAUSE_NO_CREATE. + 2018-12-20 Maciej W. Rozycki * c-pragma.h (pragma_kind): Add PRAGMA_OACC_SERIAL enumeration diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index d9a75fc80ca..a6bae05995f 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -150,6 +150,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_NOHOST, PRAGMA_OACC_CLAUSE_NUM_GANGS, PRAGMA_OACC_CLAUSE_NUM_WORKERS, diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp index 68a1d77c010..c9341355d1a 100644 --- a/gcc/c/ChangeLog.omp +++ b/gcc/c/ChangeLog.omp @@ -1,3 +1,15 @@ +2018-12-19 Julian Brown + Maciej W. Rozycki + + * 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. + 2018-12-20 Maciej W. Rozycki * c-parser.c (OACC_SERIAL_CLAUSE_MASK): New macro. diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 2baddd0e464..54fc6b32aa8 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -11741,7 +11741,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; @@ -12210,7 +12212,10 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, create ( variable-list ) delete ( variable-list ) detach ( 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, @@ -12252,6 +12257,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_NO_ALLOC; + break; case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; @@ -14930,6 +14938,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_NOHOST: clauses = c_parser_oacc_simple_clause (here, OMP_CLAUSE_NOHOST, clauses); @@ -15373,6 +15385,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 @@ -15710,6 +15723,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) \ @@ -15726,6 +15740,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) \ @@ -15745,6 +15760,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 a4a78cf8f36..2acd12d849f 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13429,6 +13429,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_NO_ALLOC: case GOMP_MAP_TO: case GOMP_MAP_FROM: case GOMP_MAP_TOFROM: diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp index 10484632cca..9d16175e1f2 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,3 +1,14 @@ +2018-12-19 Julian Brown + Maciej W. Rozycki + + * 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, OACC_SERIAL_CLAUSE_MASK): Add + PRAGMA_OACC_CLAUSE_NO_CREATE. + * semantics.c (handle_omp_array_sections): Support no_create. + 2018-12-20 Maciej W. Rozycki * constexpr.c (potential_constant_expression_1): Handle diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index bd278ee0cbd..143ad7144fa 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -32283,7 +32283,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 ("nohost", p)) result = PRAGMA_OACC_CLAUSE_NOHOST; @@ -32661,7 +32663,10 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list, create ( variable-list ) delete ( variable-list ) detach ( 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, @@ -32703,6 +32708,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_NO_ALLOC; + break; case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; @@ -35130,6 +35138,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_NOHOST: clauses = cp_parser_oacc_simple_clause (here, OMP_CLAUSE_NOHOST, clauses); @@ -38743,6 +38755,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) ) static tree @@ -39069,6 +39082,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) \ @@ -39086,6 +39100,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (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) \ @@ -39104,6 +39119,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 12ab7ceed9d..657f23c6a58 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5238,6 +5238,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_NO_ALLOC: case GOMP_MAP_TO: case GOMP_MAP_FROM: case GOMP_MAP_TOFROM: diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp index 9e83b622879..bb948709d8a 100644 --- a/gcc/fortran/ChangeLog.omp +++ b/gcc/fortran/ChangeLog.omp @@ -1,3 +1,15 @@ +2018-12-19 Julian Brown + Maciej W. Rozycki + + * 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_SERIAL_CLAUSES, OACC_DATA_CLAUSES): Add + OMP_CLAUSE_NO_CREATE. + * trans-openmp.c (gfc_trans_omp_clauses_1): Support + OMP_MAP_NO_ALLOC. + 2018-12-20 Maciej W. Rozycki * gfortran.h (gfc_statement): Add ST_OACC_SERIAL_LOOP, diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index b47b98133d5..0c3eb1e2929 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1189,6 +1189,7 @@ enum gfc_omp_depend_op enum gfc_omp_map_op { OMP_MAP_ALLOC, + OMP_MAP_NO_ALLOC, OMP_MAP_ATTACH, OMP_MAP_TO, OMP_MAP_FROM, diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 82ac0fa8523..679f99714b0 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -797,6 +797,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, @@ -1465,6 +1466,12 @@ 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_NO_ALLOC, true, + allow_derived)) + continue; if ((mask & OMP_CLAUSE_NOGROUP) && !c->nogroup && gfc_match ("nogroup") == MATCH_YES) @@ -1979,28 +1986,28 @@ 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_ATTACH) + | 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 | OMP_CLAUSE_ATTACH) #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_ATTACH) + | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \ + | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH) #define OACC_SERIAL_CLAUSES \ (omp_mask (OMP_CLAUSE_ASYNC) | OMP_CLAUSE_WAIT \ | OMP_CLAUSE_IF \ | OMP_CLAUSE_REDUCTION \ | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ - | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT \ + | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \ | OMP_CLAUSE_DEVICEPTR \ | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_ATTACH) #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_ATTACH) + | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_ATTACH) #define OACC_LOOP_CLAUSES \ (omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER \ | OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT \ diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index ed132ac7fa6..f2d8997124b 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2491,6 +2491,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_NO_ALLOC: + OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_NO_ALLOC); + break; case OMP_MAP_ATTACH: OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ATTACH); break; diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 92d118a3c51..7cdbfbb2ed4 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -9944,6 +9944,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_NO_ALLOC: case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_TO: case GOMP_MAP_FORCE_FROM: @@ -10431,6 +10432,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) switch (tkind) { case GOMP_MAP_ALLOC: + case GOMP_MAP_NO_ALLOC: case GOMP_MAP_TO: case GOMP_MAP_FROM: case GOMP_MAP_TOFROM: diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 5e1483f77cb..d6171f09f0e 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -765,6 +765,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_NO_ALLOC: + pp_string (pp, "no_alloc"); + break; case GOMP_MAP_TO: case GOMP_MAP_TO_PSET: pp_string (pp, "to"); diff --git a/include/ChangeLog.omp b/include/ChangeLog.omp index e632fdb4ebf..bd8dba0544c 100644 --- a/include/ChangeLog.omp +++ b/include/ChangeLog.omp @@ -1,3 +1,8 @@ +2018-12-19 Julian Brown + Maciej W. Rozycki + + * gomp-constants.h (gomp_map_kind): Support GOMP_MAP_NO_ALLOC. + 2018-12-21 Gergö Barany * gomp-constants.h (GOACC_FLAG_HOST_DATA_IF_PRESENT): New constant. diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 90b18091a13..dae4eea66de 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -80,6 +80,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_NO_ALLOC = (GOMP_MAP_FLAG_SPECIAL_1 | 3), /* Allocate. */ GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0), /* Similarly, but store the value in the pointer rather than diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index c063ddab27e..6c1defe2011 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,15 @@ +2018-12-19 Julian Brown + Maciej W. Rozycki + + * target.c (gomp_map_vars_async): Support GOMP_MAP_NO_ALLOC. + + * testsuite/libgomp.oacc-c-c++-common/nocreate-1.c: New test. + * testsuite/libgomp.oacc-c-c++-common/nocreate-2.c: New test. + * testsuite/libgomp.oacc-c-c++-common/nocreate-3.c: New test. + * testsuite/libgomp.oacc-c-c++-common/nocreate-4.c: New test. + * testsuite/libgomp.oacc-fortran/nocreate-1.f90: New test. + * testsuite/libgomp.oacc-fortran/nocreate-2.f90: New test. + 2018-12-20 Maciej W. Rozycki * testsuite/libgomp.oacc-c-c++-common/serial-dims.c: New test. diff --git a/libgomp/target.c b/libgomp/target.c index 91d0f25cde5..62b8ee4759e 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1180,6 +1180,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, has_firstprivate = true; continue; } + else if ((kind & typemask) == GOMP_MAP_NO_ALLOC) + { + tgt->list[i].key = NULL; + tgt->list[i].offset = 0; + continue; + } cur_node.host_start = (uintptr_t) hostaddrs[i]; if (!GOMP_MAP_POINTER_P (kind & typemask) && (kind & typemask) != GOMP_MAP_ATTACH) @@ -1468,6 +1474,53 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, cbufp); continue; } + case GOMP_MAP_NO_ALLOC: + { + cur_node.host_start = (uintptr_t) hostaddrs[i]; + cur_node.host_end = cur_node.host_start + sizes[i]; + splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); + if (n != NULL) + { + tgt->list[i].key = n; + tgt->list[i].offset = cur_node.host_start - n->host_start; + tgt->list[i].length = n->host_end - n->host_start; + tgt->list[i].copy_from = false; + tgt->list[i].always_copy_from = false; + tgt->list[i].do_detach = false; + n->refcount++; + } + else + { + tgt->list[i].key = NULL; + tgt->list[i].offset = OFFSET_INLINED; + tgt->list[i].length = sizes[i]; + tgt->list[i].copy_from = false; + tgt->list[i].always_copy_from = false; + tgt->list[i].do_detach = false; + if (i + 1 < mapnum) + { + int kind2 = get_kind (short_mapkind, kinds, i + 1); + switch (kind2 & typemask) + { + case GOMP_MAP_ATTACH: + case GOMP_MAP_POINTER: + /* The data is not present but we have an attach + or pointer clause next. Skip over it. */ + i++; + tgt->list[i].key = NULL; + tgt->list[i].offset = OFFSET_INLINED; + tgt->list[i].length = sizes[i]; + tgt->list[i].copy_from = false; + tgt->list[i].always_copy_from = false; + tgt->list[i].do_detach = false; + break; + default: + break; + } + } + } + continue; + } default: break; } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-1.c new file mode 100644 index 00000000000..c7a1bd9c015 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-1.c @@ -0,0 +1,40 @@ +/* Test no_create clause when data is present on the device. */ + +#include +#include +#include + +#define N 128 + +int +main (int argc, char *argv[]) +{ + int *arr = (int *) malloc (N * sizeof (*arr)); + int *devptr; + + acc_copyin (arr, N * sizeof (*arr)); + + #pragma acc parallel no_create(arr[0:N]) copyout(devptr) + { + devptr = &arr[2]; + } + +#if !ACC_MEM_SHARED + if (acc_hostptr (devptr) != (void *) &arr[2]) + __builtin_abort (); +#endif + + acc_delete (arr, N * sizeof (*arr)); + +#if ACC_MEM_SHARED + if (&arr[2] != devptr) + __builtin_abort (); +#else + if (&arr[2] == devptr) + __builtin_abort (); +#endif + + free (arr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-2.c new file mode 100644 index 00000000000..2964a40b217 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-2.c @@ -0,0 +1,28 @@ +/* Test no_create clause when data is not present on the device. */ + +#include +#include + +#define N 128 + +int +main (int argc, char *argv[]) +{ + int *arr = (int *) malloc (N * sizeof (*arr)); + int *devptr; + + #pragma acc data no_create(arr[0:N]) + { + #pragma acc parallel copyout(devptr) + { + devptr = &arr[2]; + } + } + + if (devptr != &arr[2]) + __builtin_abort (); + + free (arr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-3.c new file mode 100644 index 00000000000..618af6e927c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-3.c @@ -0,0 +1,38 @@ +/* Test no_create clause with attach/detach when data is not present on the + device. */ + +#include +#include +#include + +#define N 128 + +typedef struct { + int x; + int *y; +} mystruct; + +int +main (int argc, char *argv[]) +{ + int *devptr; + mystruct s; + + s.x = 5; + s.y = (int *) malloc (N * sizeof (int)); + + #pragma acc data copyin(s) + { + #pragma acc parallel no_create(s.y[0:N]) copyout(devptr) + { + devptr = &s.y[2]; + } + } + + if (devptr != &s.y[2]) + __builtin_abort (); + + free (s.y); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-4.c new file mode 100644 index 00000000000..75ab616f6ce --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-4.c @@ -0,0 +1,42 @@ +/* Test no_create clause with attach/detach when data is present on the + device. */ + +#include +#include +#include + +#define N 128 + +typedef struct { + int x; + int *y; +} mystruct; + +int +main (int argc, char *argv[]) +{ + int *devptr; + mystruct s; + + s.x = 5; + s.y = (int *) malloc (N * sizeof (int)); + + #pragma acc data copyin(s) + { + #pragma acc enter data copyin(s.y[0:N]) + + #pragma acc parallel no_create(s.y[0:N]) copyout(devptr) + { + devptr = &s.y[2]; + } + } + + if (devptr != acc_deviceptr (&s.y[2])) + __builtin_abort (); + + #pragma acc exit data delete(s.y[0:N]) + + free (s.y); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90 new file mode 100644 index 00000000000..f048355d7df --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90 @@ -0,0 +1,29 @@ +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } + +! Test no_create clause with data construct when data is present/not present. + +program nocreate + use openacc + implicit none + integer, parameter :: n = 512 + integer :: myarr(n) + integer i + + do i = 1, n + myarr(i) = 0 + end do + + !$acc data no_create (myarr) + if (acc_is_present (myarr)) stop 1 + !$acc end data + + !$acc enter data copyin (myarr) + !$acc data no_create (myarr) + if (acc_is_present (myarr) .eqv. .false.) stop 2 + !$acc end data + !$acc exit data copyout (myarr) + + do i = 1, n + if (myarr(i) .ne. 0) stop 3 + end do +end program nocreate diff --git a/libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90 new file mode 100644 index 00000000000..34444ecf5b0 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90 @@ -0,0 +1,61 @@ +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } + +! Test no_create clause with data/parallel constructs. + +program nocreate + use openacc + implicit none + integer, parameter :: n = 512 + integer :: myarr(n) + integer i + + do i = 1, n + myarr(i) = 0 + end do + + call do_on_target(myarr, n) + + do i = 1, n + if (myarr(i) .ne. i) stop 1 + end do + + do i = 1, n + myarr(i) = 0 + end do + + !$acc enter data copyin(myarr) + call do_on_target(myarr, n) + !$acc exit data copyout(myarr) + + do i = 1, n + if (myarr(i) .ne. i * 2) stop 2 + end do +end program nocreate + +subroutine do_on_target (arr, n) + use openacc + implicit none + integer :: n, arr(n) + integer :: i + +!$acc data no_create (arr) + +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 -- 2.47.3