From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 25885 invoked by alias); 24 Oct 2019 12:48:48 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 25736 invoked by uid 89); 24 Oct 2019 12:48:30 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-20.5 required=5.0 tests=AWL,BAYES_00,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,KAM_SHORT,SPF_PASS autolearn=ham version=3.3.1 spammy=tgt, __builtin_abort, UD:c-typeck.c, tkind X-HELO: esa2.mentor.iphmx.com Received: from esa2.mentor.iphmx.com (HELO esa2.mentor.iphmx.com) (68.232.141.98) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Thu, 24 Oct 2019 12:48:25 +0000 IronPort-SDR: fJypxLw39PvINL7a8LsEUhzurhfyzXlXM2VWeA+nV8SrqpDsOmdpZ8K3YTrdKwwp1fAOqDMWcM QlxCngLzQ/EaZCSp8RpNaVABLhImTRY83Y2LygRRkDUgyRt0WIJRyKpauELzlXJY5B/YS0vG8i ZRdrtw713xwqKO1j1mR38tgQ9dVlHzlQE14UCDx3ix0lYadsLy8hve0H6PsnbzvxjwPb5NBl2c FONOhTua/EBrYRBj+eKxUrKiZ0hXRxCFIO9LnZ5tM3nYkxOPzljSheMM5j5+6tnvGV0RtdzPXY 3Yo= Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 24 Oct 2019 04:48:11 -0800 IronPort-SDR: b8r5MzqaWD+fiAJP0VF33tP0vvVLxRKYQxQuPpmDnxrYEtLWZzItxtut3xfm1dMv1mtPg1RH0F vwMVIvJv1VuoDyyF58zw/JlN82nH1jU0Y6Txv45F3Kzzkhlx5NmlRYszjajXSmthki9Dfts8aC mEl7slnv+FX+hIE++HW5qmCQB7SVhtf/jOHvs1wHeNo9S/rEX+Wrggz8scMj7BSA4bBY55+bPN Jt4Ib1U2Ss9yo2SRryGJ5UZyFVQU32vTuUNnKvnQ2I6AAOMv+GZZq30ND7pMJOTIHmFZaEaxKT XL0= To: gcc-patches , fortran From: Tobias Burnus Subject: [Patch] Add OpenACC 2.6's no_create Message-ID: <7589420a-ae8a-16e5-7c81-9f4ce38a287e@mentor.com> Date: Thu, 24 Oct 2019 13:26:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:68.0) Gecko/20100101 Thunderbird/68.1.2 MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="------------DAF53AE5BDBF888EA1198CC0" Return-Path: Tobias_Burnus@mentor.com X-SW-Source: 2019-10/txt/msg01737.txt.bz2 --------------DAF53AE5BDBF888EA1198CC0 Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: 8bit Content-length: 1694 The clause (new in OpenACC 2.6) 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. – Or in words of OpenACC 2.7 (in Sect. 2.7.9 no_create clause): "The no_create clause may appear on structured data and compute constructs." / "For each var in varlist, if var is in shared memory, no action is taken; if var is not in shared memory, the no_create clause behaves as follows:" [digest: if present, update present count, if pointer attach/detach; if not not present, device-local memory used.] "The restrictions regarding subarrays in the present clause apply to this clause." Note: The "no_create" maps to the (new) GOMP_MAP_NO_ALLOC in the middle end – and all action in libgomp/target.c but only applies to GOMP_MAP_NO_ALLOC; hence, the code should only affect OpenACC. OK for the trunk? Cheers, Tobias PS: This patch is a re-diffed version of the OG9/OG8 version; as some other features are not yet on trunk, it misses a test case for "no_create(s.y…)" (i.e. the struct component-ref; libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-{3,4}.c); trunk also lacks 'acc serial' and, hence, the attach patch lacks the OACC_SERIAL_CLAUSE_MASK updates – and gfc_match_omp_map_clause needs later to be updated for the allow_derived and allow_common arguments. Furthermore, some 'do_detach = false' are missing in libgomp/target.c as they do not yet exist on trunk, either. The openacc-gcc-9 /…-8 branch patch is commit 8e74c2ec2b90819c995444370e742864a685209f of Dec 20, 2018. It has been posted as https://gcc.gnu.org/ml/gcc-patches/2018-12/msg01418.html --------------DAF53AE5BDBF888EA1198CC0 Content-Type: text/x-patch; charset="UTF-8"; name="openacc_no_create2.diff" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="openacc_no_create2.diff" Content-length: 20946 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. 2019-10-24 Julian Brown Maciej W. Rozycki Tobias Burnus 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. 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-fortran/nocreate-1.f90: New test. * testsuite/libgomp.oacc-fortran/nocreate-2.f90: New test. diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index e0aa774555a..da6cfdb8b98 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -153,6 +153,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/c-parser.c b/gcc/c/c-parser.c index 7618a46c8bc..1004a2e5579 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -11833,7 +11833,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; @@ -12296,7 +12298,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, @@ -12332,6 +12337,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; @@ -15130,6 +15138,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, @@ -15598,6 +15610,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 @@ -15925,6 +15938,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) \ @@ -15940,6 +15954,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) \ diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index c7339509bd1..3be6b654598 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13409,6 +13409,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/parser.c b/gcc/cp/parser.c index 3857fe47d67..8d7de8bc33b 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -33019,7 +33019,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; @@ -33385,7 +33387,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, @@ -33421,6 +33426,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; @@ -35983,6 +35991,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"; @@ -39788,6 +39800,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 @@ -40105,6 +40118,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) \ @@ -40119,8 +40133,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) \ diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 59def3170ab..db7cac82312 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5291,6 +5291,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/gfortran.h b/gcc/fortran/gfortran.h index 920acdafc6b..5c930097a66 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1191,6 +1191,7 @@ enum gfc_omp_depend_op enum gfc_omp_map_op { OMP_MAP_ALLOC, + OMP_MAP_NO_ALLOC, OMP_MAP_TO, OMP_MAP_FROM, OMP_MAP_TOFROM, diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 5c91fcdfd31..00575fd37aa 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, @@ -1444,6 +1445,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_NO_ALLOC)) + continue; if ((mask & OMP_CLAUSE_NOGROUP) && !c->nogroup && gfc_match ("nogroup") == MATCH_YES) @@ -1954,19 +1960,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_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_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 \ diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index dad11a24430..979f83c234e 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2351,6 +2351,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_TO: OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO); break; diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 279b6ef893a..8cf43852b2a 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -11315,6 +11315,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: @@ -11724,6 +11725,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 53b3f55a3e6..51cfa837c91 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_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/gomp-constants.h b/include/gomp-constants.h index 82e9094c934..e9f5441d9da 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_NO_ALLOC = (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/target.c b/libgomp/target.c index 84d6daa76ca..cc79bb4dd06 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -667,6 +667,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)) cur_node.host_end = cur_node.host_start + sizes[i]; @@ -892,6 +898,49 @@ 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_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; + 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; + if (i + 1 < mapnum) + { + int kind2 = get_kind (short_mapkind, kinds, i + 1); + switch (kind2 & typemask) + { + 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; + 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-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 --------------DAF53AE5BDBF888EA1198CC0--