From patchwork Wed Dec 19 21:30:19 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: "Maciej W. Rozycki" X-Patchwork-Id: 1016356 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-492844-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="dmTf645L"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 43Kp454pm2z9sB5 for ; Thu, 20 Dec 2018 08:31:33 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:cc:subject:message-id:mime-version:content-type; q=dns; s=default; b=BqUTeSaXXc79GM9auJsNQQ8lIum+6hfKJcg1QVdQNgp2iex9AM mBdkZXGrGEA6eH97LdbXLm78gV5WiyZ0CmKOIBJSchhX7I2sVLtM7xwdNFO0LzEC oEy4hRKwhTx6CRHTqj8bp0DoLKCI8sAxH9MzclUockcLdNU7RvCJ0p3XY= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:cc:subject:message-id:mime-version:content-type; s= default; bh=3f/mY6JLILrOGPxcqXi6UISLsYM=; b=dmTf645LzoZMk7Wd5lju uo5hfzUXrQUDnZ8OuvWcwxhrUWpQztyS2mhU5gSVgf/siMoO8uJB5vzhG4XfE79R 56I7WodLu2YtIJNkjYVURu0K0hzVrRBgqr698QL8xTyZJJcMjQMje+euptxlYpK/ oxx3v6Ab4rppYG/jdzFZ8lw= Received: (qmail 75697 invoked by alias); 19 Dec 2018 21:31:20 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 75403 invoked by uid 89); 19 Dec 2018 21:31:00 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-11.9 required=5.0 tests=BAYES_00, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS, TIME_LIMIT_EXCEEDED autolearn=unavailable version=3.3.2 spammy=H*Ad:U*thomas, sy, juliancodesourcerycom, julian@codesourcery.com X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 19 Dec 2018 21:30:41 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-MBX-04.mgc.mentorg.com) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1gZjQF-0002yt-88 from Maciej_Rozycki@mentor.com ; Wed, 19 Dec 2018 13:30:31 -0800 Received: from build7-trusty-cs.sje.mentorg.com (137.202.0.90) by SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Wed, 19 Dec 2018 21:30:26 +0000 Date: Wed, 19 Dec 2018 21:30:19 +0000 From: "Maciej W. Rozycki" To: CC: Julian Brown , Thomas Schwinge , Chung-Lin Tang , Jakub Jelinek , Catherine Moore Subject: [PATCH, og8] Add OpenACC 2.6 `no_create' clause support Message-ID: User-Agent: Alpine 2.21.9999 (DEB 301 2018-08-15) MIME-Version: 1.0 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. --- Hi, This has passed regression-testing with the `x86_64-linux-gnu' target and the `nvptx-none' offload target, across the `gcc', `g++', `gfortran' and `libgomp' test suites. I will appreciate feedback and if none has been given shortly, then I will commit this change to the og8 branch. Maciej --- gcc/c-family/c-pragma.h | 1 gcc/c/c-parser.c | 20 ++++ gcc/c/c-typeck.c | 1 gcc/cp/parser.c | 20 ++++ gcc/cp/semantics.c | 1 gcc/fortran/gfortran.h | 1 gcc/fortran/openmp.c | 15 ++- gcc/fortran/trans-openmp.c | 3 gcc/omp-low.c | 2 gcc/tree-pretty-print.c | 3 include/gomp-constants.h | 2 libgomp/target.c | 53 +++++++++++++ libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-1.c | 40 +++++++++ libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-2.c | 28 ++++++ libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-3.c | 38 +++++++++ libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-4.c | 42 ++++++++++ libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90 | 29 +++++++ libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90 | 61 +++++++++++++++ 18 files changed, 352 insertions(+), 8 deletions(-) gcc-openacc-no-create.diff Index: gcc-openacc-gcc-8-branch/gcc/c-family/c-pragma.h =================================================================== --- gcc-openacc-gcc-8-branch.orig/gcc/c-family/c-pragma.h +++ gcc-openacc-gcc-8-branch/gcc/c-family/c-pragma.h @@ -147,6 +147,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, Index: gcc-openacc-gcc-8-branch/gcc/c/c-parser.c =================================================================== --- gcc-openacc-gcc-8-branch.orig/gcc/c/c-parser.c +++ gcc-openacc-gcc-8-branch/gcc/c/c-parser.c @@ -11315,7 +11315,9 @@ c_parser_omp_clause_name (c_parser *pars 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 ("notinbranch", p)) result = PRAGMA_OMP_CLAUSE_NOTINBRANCH; @@ -11689,7 +11691,10 @@ c_parser_omp_var_list_parens (c_parser * 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, @@ -11731,6 +11736,9 @@ c_parser_oacc_data_clause (c_parser *par 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; @@ -14194,6 +14202,10 @@ c_parser_oacc_all_clauses (c_parser *par 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 (parser, here, OMP_CLAUSE_NOHOST, clauses); @@ -14619,6 +14631,7 @@ c_parser_oacc_cache (location_t loc, c_p | (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 @@ -14968,6 +14981,7 @@ c_parser_oacc_loop (location_t loc, c_pa | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ | (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) \ @@ -14992,6 +15006,7 @@ c_parser_oacc_loop (location_t loc, c_pa | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ | (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) \ @@ -15019,6 +15034,7 @@ c_parser_oacc_loop (location_t loc, c_pa | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ | (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) \ Index: gcc-openacc-gcc-8-branch/gcc/c/c-typeck.c =================================================================== --- gcc-openacc-gcc-8-branch.orig/gcc/c/c-typeck.c +++ gcc-openacc-gcc-8-branch/gcc/c/c-typeck.c @@ -12978,6 +12978,7 @@ handle_omp_array_sections (tree c, enum 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: Index: gcc-openacc-gcc-8-branch/gcc/cp/parser.c =================================================================== --- gcc-openacc-gcc-8-branch.orig/gcc/cp/parser.c +++ gcc-openacc-gcc-8-branch/gcc/cp/parser.c @@ -31353,7 +31353,9 @@ cp_parser_omp_clause_name (cp_parser *pa 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; @@ -31694,7 +31696,10 @@ cp_parser_omp_var_list (cp_parser *parse 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, @@ -31736,6 +31741,9 @@ cp_parser_oacc_data_clause (cp_parser *p 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; @@ -33964,6 +33972,10 @@ cp_parser_oacc_all_clauses (cp_parser *p 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 (parser, OMP_CLAUSE_NOHOST, clauses, here); @@ -36936,6 +36948,7 @@ cp_parser_oacc_cache (cp_parser *parser, | (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 @@ -37272,6 +37285,7 @@ cp_parser_oacc_loop (cp_parser *parser, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ | (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) \ @@ -37297,6 +37311,7 @@ cp_parser_oacc_loop (cp_parser *parser, | (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) \ @@ -37323,6 +37338,7 @@ cp_parser_oacc_loop (cp_parser *parser, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ | (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) \ Index: gcc-openacc-gcc-8-branch/gcc/cp/semantics.c =================================================================== --- gcc-openacc-gcc-8-branch.orig/gcc/cp/semantics.c +++ gcc-openacc-gcc-8-branch/gcc/cp/semantics.c @@ -5096,6 +5096,7 @@ handle_omp_array_sections (tree c, enum 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: Index: gcc-openacc-gcc-8-branch/gcc/fortran/gfortran.h =================================================================== --- gcc-openacc-gcc-8-branch.orig/gcc/fortran/gfortran.h +++ gcc-openacc-gcc-8-branch/gcc/fortran/gfortran.h @@ -1184,6 +1184,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, Index: gcc-openacc-gcc-8-branch/gcc/fortran/openmp.c =================================================================== --- gcc-openacc-gcc-8-branch.orig/gcc/fortran/openmp.c +++ gcc-openacc-gcc-8-branch/gcc/fortran/openmp.c @@ -818,6 +818,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, @@ -1559,6 +1560,12 @@ gfc_match_omp_clauses (gfc_omp_clauses * } 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) @@ -2070,7 +2077,7 @@ gfc_match_omp_clauses (gfc_omp_clauses * | 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) @@ -2081,7 +2088,7 @@ gfc_match_omp_clauses (gfc_omp_clauses * | OMP_CLAUSE_DEVICE_TYPE \ | OMP_CLAUSE_IF \ | 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_DEFAULT | OMP_CLAUSE_ATTACH) #define OACC_SERIAL_CLAUSES \ @@ -2090,14 +2097,14 @@ gfc_match_omp_clauses (gfc_omp_clauses * | 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_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_ATTACH) #define OACC_HOST_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_USE_DEVICE)) Index: gcc-openacc-gcc-8-branch/gcc/fortran/trans-openmp.c =================================================================== --- gcc-openacc-gcc-8-branch.orig/gcc/fortran/trans-openmp.c +++ gcc-openacc-gcc-8-branch/gcc/fortran/trans-openmp.c @@ -2348,6 +2348,9 @@ gfc_trans_omp_clauses_1 (stmtblock_t *bl 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; Index: gcc-openacc-gcc-8-branch/gcc/omp-low.c =================================================================== --- gcc-openacc-gcc-8-branch.orig/gcc/omp-low.c +++ gcc-openacc-gcc-8-branch/gcc/omp-low.c @@ -8184,6 +8184,7 @@ lower_omp_target (gimple_stmt_iterator * 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: @@ -8681,6 +8682,7 @@ lower_omp_target (gimple_stmt_iterator * switch (tkind) { case GOMP_MAP_ALLOC: + case GOMP_MAP_NO_ALLOC: case GOMP_MAP_TO: case GOMP_MAP_FROM: case GOMP_MAP_TOFROM: Index: gcc-openacc-gcc-8-branch/gcc/tree-pretty-print.c =================================================================== --- gcc-openacc-gcc-8-branch.orig/gcc/tree-pretty-print.c +++ gcc-openacc-gcc-8-branch/gcc/tree-pretty-print.c @@ -684,6 +684,9 @@ dump_omp_clause (pretty_printer *pp, tre 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"); Index: gcc-openacc-gcc-8-branch/include/gomp-constants.h =================================================================== --- gcc-openacc-gcc-8-branch.orig/include/gomp-constants.h +++ gcc-openacc-gcc-8-branch/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 Index: gcc-openacc-gcc-8-branch/libgomp/target.c =================================================================== --- gcc-openacc-gcc-8-branch.orig/libgomp/target.c +++ gcc-openacc-gcc-8-branch/libgomp/target.c @@ -1212,6 +1212,12 @@ gomp_map_vars_async (struct gomp_device_ 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) @@ -1496,6 +1502,53 @@ gomp_map_vars_async (struct gomp_device_ 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; } Index: gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-1.c =================================================================== --- /dev/null +++ gcc-openacc-gcc-8-branch/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; +} Index: gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-2.c =================================================================== --- /dev/null +++ gcc-openacc-gcc-8-branch/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; +} Index: gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-3.c =================================================================== --- /dev/null +++ gcc-openacc-gcc-8-branch/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; +} Index: gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/nocreate-4.c =================================================================== --- /dev/null +++ gcc-openacc-gcc-8-branch/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; +} Index: gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-fortran/nocreate-1.f90 =================================================================== --- /dev/null +++ gcc-openacc-gcc-8-branch/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 Index: gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-fortran/nocreate-2.f90 =================================================================== --- /dev/null +++ gcc-openacc-gcc-8-branch/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