From patchwork Thu Apr 13 11:11:16 2017 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 750413 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org 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 3w3dQt0RrTz9sN9 for ; Thu, 13 Apr 2017 21:11:49 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="DkX62WdJ"; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :subject:to:cc:message-id:date:mime-version:content-type; q=dns; s=default; b=cT4IfOv6EK6t35ViCoIPtN67mRdMUSLfp7wfNheA/Rrm9Lg1u/ S6etln5QRg3EQBb3Pbp9EDoZxUEwdxOhPfpb+2IkshH2HLVwjG26JMnmGPfi8t20 ieiutFg78IBKCNZXjAxBmAy3IyfKMwlJgmYv5DaKEnlSJhFbw4TDEbK58= 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:from :subject:to:cc:message-id:date:mime-version:content-type; s= default; bh=ba9pbKjXKrf3HjII4SJETinlBT8=; b=DkX62WdJMr6qHUeNTk9v 8CyXEeJ++KE9T6ka1NFiIeDL09eDCf1x0qWMUgI5KMKD1SMJ7crCna7Ev4NW9103 sCKcex30c0FQ97fJQwyjiQgY+WwtTo5WJswLtvZ4wG3dZcQzEJFqs2LFEge/aWuB K/Ll2qxhrv33YR4ei4pK5Rs= Received: (qmail 74159 invoked by alias); 13 Apr 2017 11:11:33 -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 74026 invoked by uid 89); 13 Apr 2017 11:11:31 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-10.1 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_2, GIT_PATCH_3, KAM_ASCII_DIVIDERS, RCVD_IN_DNSWL_NONE, SPF_PASS, URIBL_RED autolearn=ham version=3.3.2 spammy=325, 1665, H*r:ip*0.0.0.0, H*r:0.0.0 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; Thu, 13 Apr 2017 11:11:23 +0000 Received: from svr-orw-mbx-06.mgc.mentorg.com ([147.34.90.206]) by relay1.mentorg.com with esmtp id 1cycen-0004kR-UK from ChungLin_Tang@mentor.com for gcc-patches@gcc.gnu.org; Thu, 13 Apr 2017 04:11:21 -0700 Received: from svr-orw-mbx-08.mgc.mentorg.com (147.34.90.208) by SVR-ORW-MBX-06.mgc.mentorg.com (147.34.90.206) with Microsoft SMTP Server (TLS) id 15.0.1210.3; Thu, 13 Apr 2017 04:11:18 -0700 Received: from [0.0.0.0] (147.34.91.1) by svr-orw-mbx-08.mgc.mentorg.com (147.34.90.208) with Microsoft SMTP Server (TLS) id 15.0.1210.3 via Frontend Transport; Thu, 13 Apr 2017 04:11:17 -0700 From: Chung-Lin Tang Subject: [patch, gomp4, committed] Adjust copy/copyin/copyout/create for OpenACC 2.5 To: gcc-patches CC: Thomas Schwinge , Cesar Philippidis Message-ID: Date: Thu, 13 Apr 2017 19:11:16 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.11; rv:45.0) Gecko/20100101 Thunderbird/45.8.0 MIME-Version: 1.0 The behavior of the copy/copyin/copyout/create clauses has been changed in OpenACC 2.5 to be like the present_or_* variants, and the original present_or_* syntax relegated to legacy status. This patch removes the presence of any PRESENT_OR_* symbols, and changes the mapping of the copy/copyin/copyout/create clauses to map kinds without the FORCE flag. Library routines acc_[present_or_]copy, etc. has also been updated. This patch has been applied to the gomp-4_0-branch. Chung-Lin 2017-04-13 Chung-Lin Tang gcc/ * gimplify.c (gimplify_oacc_declare_1): Remove GOMP_MAP_FORCE_* cases. gcc/c-family/ * c-pragma.h (enum pragma_omp_clauses): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY/COPYIN/COPYOUT/CREATE. gcc/c/ * c-parser.c (c_parser_omp_clause_name): Remove occurences of PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY/COPYIN/COPYOUT/CREATE, adjust them to non-PRESENT_OR values. (c_parser_oacc_data_clause): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_* cases, remove FORCE from PRAGMA_OACC_CLAUSE_COPY/COPYIN/COPYOUT/CREATE kinds, update comment description. (c_parser_oacc_all_clauses): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_* cases. (OACC_DATA_CLAUSE_MASK): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_*. (OACC_DECLARE_CLAUSE_MASK): Likewise. (c_parser_oacc_declare): Remove GOMP_MAP_FORCE_ALLOC/TO, change to COMP_MAP_ALLOC/TO. (OACC_ENTER_DATA_CLAUSE_MASK): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_*. (OACC_KERNELS_CLAUSE_MASK): Likewise. (OACC_PARALLEL_CLAUSE_MASK): Likewise. gcc/cp/ * parser.c (cp_parser_omp_clause_name): Remove occurences of PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY/COPYIN/COPYOUT/CREATE, adjust them to non-PRESENT_OR values. (cp_parser_oacc_data_clause): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_* cases, remove FORCE from PRAGMA_OACC_CLAUSE_COPY/COPYIN/COPYOUT/CREATE kinds, update comment description. (cp_parser_oacc_all_clauses): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_* cases. (OACC_DATA_CLAUSE_MASK): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_*. (OACC_DECLARE_CLAUSE_MASK): Likewise. (cp_parser_oacc_declare): Remove GOMP_MAP_FORCE_ALLOC/TO, change to COMP_MAP_ALLOC/TO. (OACC_ENTER_DATA_CLAUSE_MASK): Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_*. (OACC_KERNELS_CLAUSE_MASK): Likewise. (OACC_PARALLEL_CLAUSE_MASK): Likewise. gcc/fortran/ * openmp.c (enum omp_mask2): Remove OMP_CLAUSE_PRESENT_OR_COPY/COPYIN/COPYOUT/CREATE. (gfc_match_omp_clauses): Remove FORCE flag from OpenACC OMP_MAP_* cases, adjust and remove PRESENT_OR_ values. (OACC_PARALLEL_CLAUSES): Remove OMP_CLAUSE_PRESENT_OR_*. (OACC_KERNELS_CLAUSES): Likewise. (OACC_DATA_CLAUSES): Likewise. (OACC_DECLARE_CLAUSES): Likewise. (OACC_ENTER_DATA_CLAUSES): Likewise. gcc/testsuite/ * gfortran.dg/goacc/data-tree.f95: Remove "force_" from dump scan. * gfortran.dg/goacc/kernels-tree.f95: Likewise. * gfortran.dg/goacc/data-tree.f95: Likewise. * gfortran.dg/goacc/reduction-promotions.f90: Likewise. * gfortran.dg/goacc/combined-directives.f90: Likewise. * gfortran.dg/goacc/default-4.f: Likewise. * gfortran.dg/declare-2.f95: Adjust error scan. * c-c++-common/goacc/acc-data-chain.c: Remove "force_" from dump scan. * c-c++-common/goacc/default-4.c: Likewise. * c-c++-common/goacc/declare-2.c: Move present_or_copyin/create tests from here to... * c-c++-common/goacc/declare-1.c: ...here. libgomp/ * oacc-mem.c (acc_create): Add FLAG_PRESENT to call of present_create_copy. (acc_create_async): Likewise. (acc_copyin): Likewise. (acc_copyin_async): Likewise. (acc_present_or_create): Remove definition and change to alias of acc_create. (acc_present_or_copyin): Remove definition and change to alias of acc_copyin. * oacc-parallel.c (GOACC_enter_exit_data): Add GOMP_MAP_FROM as handled map kind. * testsuite/libgomp.oacc-c-c++-common/data-already-1.c: Delete. * testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Delete. * testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Delete. * testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Delete. * testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Delete. * testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Delete. * testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Delete. * testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Delete. * testsuite/libgomp.oacc-fortran/data-already-1.f: Delete. * testsuite/libgomp.oacc-fortran/data-already-2.f: Delete. * testsuite/libgomp.oacc-fortran/data-already-3.f: Delete. * testsuite/libgomp.oacc-fortran/data-already-4.f: Delete. * testsuite/libgomp.oacc-fortran/data-already-5.f: Delete. * testsuite/libgomp.oacc-fortran/data-already-6.f: Delete. * testsuite/libgomp.oacc-fortran/data-already-7.f: Delete. * testsuite/libgomp.oacc-fortran/data-already-8.f: Delete. * testsuite/libgomp.oacc-c-c++-common/lib-16.c: Remove dg-shouldfail. * testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise. Index: gcc/gimplify.c =================================================================== --- gcc/gimplify.c (revision 246808) +++ gcc/gimplify.c (working copy) @@ -8584,24 +8584,10 @@ gimplify_oacc_declare_1 (tree clause) switch (kind) { case GOMP_MAP_ALLOC: - case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_TO: new_op = GOMP_MAP_DELETE; ret = true; break; - case GOMP_MAP_FORCE_FROM: - OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC); - new_op = GOMP_MAP_FORCE_FROM; - ret = true; - break; - - case GOMP_MAP_FORCE_TOFROM: - OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_TO); - new_op = GOMP_MAP_FORCE_FROM; - ret = true; - break; - case GOMP_MAP_FROM: OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC); new_op = GOMP_MAP_FROM; Index: gcc/c/c-parser.c =================================================================== --- gcc/c/c-parser.c (revision 246808) +++ gcc/c/c-parser.c (working copy) @@ -10447,18 +10447,20 @@ c_parser_omp_clause_name (c_parser *parser, bool c result = PRAGMA_OMP_CLAUSE_PARALLEL; else if (!strcmp ("present", p)) result = PRAGMA_OACC_CLAUSE_PRESENT; + /* As of OpenACC 2.5, these are now aliases of the non-present_or + clauses. */ else if (!strcmp ("present_or_copy", p) || !strcmp ("pcopy", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY; + result = PRAGMA_OACC_CLAUSE_COPY; else if (!strcmp ("present_or_copyin", p) || !strcmp ("pcopyin", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN; + result = PRAGMA_OACC_CLAUSE_COPYIN; else if (!strcmp ("present_or_copyout", p) || !strcmp ("pcopyout", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT; + result = PRAGMA_OACC_CLAUSE_COPYOUT; else if (!strcmp ("present_or_create", p) || !strcmp ("pcreate", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE; + result = PRAGMA_OACC_CLAUSE_CREATE; else if (!strcmp ("priority", p)) result = PRAGMA_OMP_CLAUSE_PRIORITY; else if (!strcmp ("private", p)) @@ -10752,7 +10754,7 @@ c_parser_omp_var_list_parens (c_parser *parser, en return list; } -/* OpenACC 2.0: +/* OpenACC 2.5: copy ( variable-list ) copyin ( variable-list ) copyout ( variable-list ) @@ -10760,15 +10762,7 @@ c_parser_omp_var_list_parens (c_parser *parser, en delete ( variable-list ) device_resident ( variable-list ) link ( variable-list ) - present ( variable-list ) - present_or_copy ( variable-list ) - pcopy ( variable-list ) - present_or_copyin ( variable-list ) - pcopyin ( variable-list ) - present_or_copyout ( variable-list ) - pcopyout ( variable-list ) - present_or_create ( variable-list ) - pcreate ( variable-list ) */ + present ( variable-list ) */ static tree c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, @@ -10778,16 +10772,16 @@ c_parser_oacc_data_clause (c_parser *parser, pragm switch (c_kind) { case PRAGMA_OACC_CLAUSE_COPY: - kind = GOMP_MAP_FORCE_TOFROM; + kind = GOMP_MAP_TOFROM; break; case PRAGMA_OACC_CLAUSE_COPYIN: - kind = GOMP_MAP_FORCE_TO; + kind = GOMP_MAP_TO; break; case PRAGMA_OACC_CLAUSE_COPYOUT: - kind = GOMP_MAP_FORCE_FROM; + kind = GOMP_MAP_FROM; break; case PRAGMA_OACC_CLAUSE_CREATE: - kind = GOMP_MAP_FORCE_ALLOC; + kind = GOMP_MAP_ALLOC; break; case PRAGMA_OACC_CLAUSE_DELETE: kind = GOMP_MAP_DELETE; @@ -10807,18 +10801,6 @@ c_parser_oacc_data_clause (c_parser *parser, pragm case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY: - kind = GOMP_MAP_TOFROM; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN: - kind = GOMP_MAP_TO; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT: - kind = GOMP_MAP_FROM; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE: - kind = GOMP_MAP_ALLOC; - break; default: gcc_unreachable (); } @@ -13285,22 +13267,6 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_c clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "present"; break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY: - clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_copy"; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN: - clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_copyin"; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT: - clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_copyout"; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE: - clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_create"; - break; case PRAGMA_OACC_CLAUSE_PRIVATE: clauses = c_parser_omp_clause_private (parser, clauses); c_name = "private"; @@ -13706,11 +13672,7 @@ c_parser_oacc_cache (location_t loc, c_parser *par | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) ) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)) static tree c_parser_oacc_data (location_t loc, c_parser *parser, bool *if_p) @@ -13741,11 +13703,7 @@ c_parser_oacc_data (location_t loc, c_parser *pars | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) ) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)) static void c_parser_oacc_declare (c_parser *parser) @@ -13780,8 +13738,8 @@ c_parser_oacc_declare (c_parser *parser) switch (OMP_CLAUSE_MAP_KIND (t)) { case GOMP_MAP_FIRSTPRIVATE_POINTER: - case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_TO: + case GOMP_MAP_ALLOC: + case GOMP_MAP_TO: case GOMP_MAP_FORCE_DEVICEPTR: case GOMP_MAP_DEVICE_RESIDENT: break; @@ -13891,8 +13849,6 @@ c_parser_oacc_declare (c_parser *parser) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) #define OACC_EXIT_DATA_CLAUSE_MASK \ @@ -14051,10 +14007,6 @@ c_parser_oacc_loop (location_t loc, c_parser *pars | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) #define OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK \ @@ -14076,10 +14028,6 @@ c_parser_oacc_loop (location_t loc, c_parser *pars | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) Index: gcc/c-family/c-pragma.h =================================================================== --- gcc/c-family/c-pragma.h (revision 246808) +++ gcc/c-family/c-pragma.h (working copy) @@ -164,10 +164,6 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_NUM_GANGS, PRAGMA_OACC_CLAUSE_NUM_WORKERS, PRAGMA_OACC_CLAUSE_PRESENT, - PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY, - PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN, - PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT, - PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE, PRAGMA_OACC_CLAUSE_SEQ, PRAGMA_OACC_CLAUSE_TILE, PRAGMA_OACC_CLAUSE_USE_DEVICE, Index: gcc/fortran/openmp.c =================================================================== --- gcc/fortran/openmp.c (revision 246808) +++ gcc/fortran/openmp.c (working copy) @@ -818,10 +818,6 @@ enum omp_mask2 OMP_CLAUSE_COPYOUT, OMP_CLAUSE_CREATE, OMP_CLAUSE_PRESENT, - OMP_CLAUSE_PRESENT_OR_COPY, - OMP_CLAUSE_PRESENT_OR_COPYIN, - OMP_CLAUSE_PRESENT_OR_COPYOUT, - OMP_CLAUSE_PRESENT_OR_CREATE, OMP_CLAUSE_DEVICEPTR, OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER, @@ -1082,7 +1078,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m if ((mask & OMP_CLAUSE_COPY) && gfc_match ("copy ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_TOFROM, openacc, + OMP_MAP_TOFROM, openacc, allow_derived)) continue; if (mask & OMP_CLAUSE_COPYIN) @@ -1091,7 +1087,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m { if (gfc_match ("copyin ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_TO, true, + OMP_MAP_TO, true, allow_derived)) continue; } @@ -1103,7 +1099,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m if ((mask & OMP_CLAUSE_COPYOUT) && gfc_match ("copyout ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_FROM, true, + OMP_MAP_FROM, true, allow_derived)) continue; if ((mask & OMP_CLAUSE_COPYPRIVATE) @@ -1114,7 +1110,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m if ((mask & OMP_CLAUSE_CREATE) && gfc_match ("create ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_FORCE_ALLOC, true, + OMP_MAP_ALLOC, true, allow_derived)) continue; break; @@ -1616,22 +1612,22 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m } break; case 'p': - if ((mask & OMP_CLAUSE_PRESENT_OR_COPY) + if ((mask & OMP_CLAUSE_COPY) && gfc_match ("pcopy ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_TOFROM, true, allow_derived)) continue; - if ((mask & OMP_CLAUSE_PRESENT_OR_COPYIN) + if ((mask & OMP_CLAUSE_COPYIN) && gfc_match ("pcopyin ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_TO, true, allow_derived)) continue; - if ((mask & OMP_CLAUSE_PRESENT_OR_COPYOUT) + if ((mask & OMP_CLAUSE_COPYOUT) && gfc_match ("pcopyout ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_FROM, true, allow_derived)) continue; - if ((mask & OMP_CLAUSE_PRESENT_OR_CREATE) + if ((mask & OMP_CLAUSE_CREATE) && gfc_match ("pcreate ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_ALLOC, true, allow_derived)) @@ -1642,22 +1638,22 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m OMP_MAP_FORCE_PRESENT, false, allow_derived)) continue; - if ((mask & OMP_CLAUSE_PRESENT_OR_COPY) + if ((mask & OMP_CLAUSE_COPY) && gfc_match ("present_or_copy ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_TOFROM, true, allow_derived)) continue; - if ((mask & OMP_CLAUSE_PRESENT_OR_COPYIN) + if ((mask & OMP_CLAUSE_COPYIN) && gfc_match ("present_or_copyin ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_TO, true, allow_derived)) continue; - if ((mask & OMP_CLAUSE_PRESENT_OR_COPYOUT) + if ((mask & OMP_CLAUSE_COPYOUT) && gfc_match ("present_or_copyout ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_FROM, true, allow_derived)) continue; - if ((mask & OMP_CLAUSE_PRESENT_OR_CREATE) + if ((mask & OMP_CLAUSE_CREATE) && gfc_match ("present_or_create ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_ALLOC, true, allow_derived)) @@ -2040,24 +2036,18 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \ | OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_REDUCTION \ | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ - | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \ - | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \ - | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE \ - | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT \ - | OMP_CLAUSE_DEVICE_TYPE) + | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \ + | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \ + | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE) #define OACC_KERNELS_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_DEVICEPTR \ | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ - | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \ - | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \ - | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT \ - | OMP_CLAUSE_DEVICE_TYPE) + | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \ + | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE) #define OACC_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \ | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \ - | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \ - | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \ - | OMP_CLAUSE_PRESENT_OR_CREATE) + | OMP_CLAUSE_PRESENT) #define OACC_LOOP_CLAUSES \ (omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER \ | OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT \ @@ -2071,16 +2061,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_m #define OACC_DECLARE_CLAUSES \ (omp_mask (OMP_CLAUSE_COPY) | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ | OMP_CLAUSE_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_DEVICE_RESIDENT \ - | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \ - | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \ - | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_LINK) + | OMP_CLAUSE_PRESENT | OMP_CLAUSE_LINK) #define OACC_UPDATE_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST_SELF \ | OMP_CLAUSE_DEVICE | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE) #define OACC_ENTER_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \ - | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT_OR_COPYIN \ - | OMP_CLAUSE_PRESENT_OR_CREATE) + | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE) #define OACC_EXIT_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \ | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE) @@ -2203,8 +2190,7 @@ gfc_match_oacc_declare (void) if (s->ns->proc_name && s->ns->proc_name->attr.proc == PROC_MODULE) { - if (n->u.map_op != OMP_MAP_FORCE_ALLOC - && n->u.map_op != OMP_MAP_FORCE_TO) + if (n->u.map_op != OMP_MAP_ALLOC && n->u.map_op != OMP_MAP_TO) { gfc_error ("Invalid clause in module with $!ACC DECLARE at %L", &where); Index: gcc/cp/parser.c =================================================================== --- gcc/cp/parser.c (revision 246808) +++ gcc/cp/parser.c (working copy) @@ -29887,18 +29887,20 @@ cp_parser_omp_clause_name (cp_parser *parser, bool result = PRAGMA_OMP_CLAUSE_PARALLEL; else if (!strcmp ("present", p)) result = PRAGMA_OACC_CLAUSE_PRESENT; + /* As of OpenACC 2.5, these are now aliases of the non-present_or + clauses. */ else if (!strcmp ("present_or_copy", p) || !strcmp ("pcopy", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY; + result = PRAGMA_OACC_CLAUSE_COPY; else if (!strcmp ("present_or_copyin", p) || !strcmp ("pcopyin", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN; + result = PRAGMA_OACC_CLAUSE_COPYIN; else if (!strcmp ("present_or_copyout", p) || !strcmp ("pcopyout", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT; + result = PRAGMA_OACC_CLAUSE_COPYOUT; else if (!strcmp ("present_or_create", p) || !strcmp ("pcreate", p)) - result = PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE; + result = PRAGMA_OACC_CLAUSE_CREATE; else if (!strcmp ("priority", p)) result = PRAGMA_OMP_CLAUSE_PRIORITY; else if (!strcmp ("proc_bind", p)) @@ -30169,7 +30171,7 @@ cp_parser_omp_var_list (cp_parser *parser, enum om return list; } -/* OpenACC 2.0: +/* OpenACC 2.5: copy ( variable-list ) copyin ( variable-list ) copyout ( variable-list ) @@ -30178,15 +30180,7 @@ cp_parser_omp_var_list (cp_parser *parser, enum om device_resident ( variable-list ) firstprivate (variable-list ) link ( variable-list ) - present ( variable-list ) - present_or_copy ( variable-list ) - pcopy ( variable-list ) - present_or_copyin ( variable-list ) - pcopyin ( variable-list ) - present_or_copyout ( variable-list ) - pcopyout ( variable-list ) - present_or_create ( variable-list ) - pcreate ( variable-list ) */ + present ( variable-list ) */ static tree cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, @@ -30196,16 +30190,16 @@ cp_parser_oacc_data_clause (cp_parser *parser, pra switch (c_kind) { case PRAGMA_OACC_CLAUSE_COPY: - kind = GOMP_MAP_FORCE_TOFROM; + kind = GOMP_MAP_TOFROM; break; case PRAGMA_OACC_CLAUSE_COPYIN: - kind = GOMP_MAP_FORCE_TO; + kind = GOMP_MAP_TO; break; case PRAGMA_OACC_CLAUSE_COPYOUT: - kind = GOMP_MAP_FORCE_FROM; + kind = GOMP_MAP_FROM; break; case PRAGMA_OACC_CLAUSE_CREATE: - kind = GOMP_MAP_FORCE_ALLOC; + kind = GOMP_MAP_ALLOC; break; case PRAGMA_OACC_CLAUSE_DELETE: kind = GOMP_MAP_DELETE; @@ -30225,18 +30219,6 @@ cp_parser_oacc_data_clause (cp_parser *parser, pra case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY: - kind = GOMP_MAP_TOFROM; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN: - kind = GOMP_MAP_TO; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT: - kind = GOMP_MAP_FROM; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE: - kind = GOMP_MAP_ALLOC; - break; default: gcc_unreachable (); } @@ -32438,22 +32420,6 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "present"; break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY: - clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_copy"; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN: - clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_copyin"; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT: - clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_copyout"; - break; - case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE: - clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "present_or_create"; - break; case PRAGMA_OACC_CLAUSE_PRIVATE: clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_PRIVATE, clauses); @@ -35387,11 +35353,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) ) static tree cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) @@ -35446,11 +35408,7 @@ cp_parser_oacc_host_data (cp_parser *parser, cp_to | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) ) static tree cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) @@ -35482,8 +35440,8 @@ cp_parser_oacc_declare (cp_parser *parser, cp_toke switch (OMP_CLAUSE_MAP_KIND (t)) { case GOMP_MAP_FIRSTPRIVATE_POINTER: - case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_TO: + case GOMP_MAP_ALLOC: + case GOMP_MAP_TO: case GOMP_MAP_FORCE_DEVICEPTR: case GOMP_MAP_DEVICE_RESIDENT: break; @@ -35592,8 +35550,6 @@ cp_parser_oacc_declare (cp_parser *parser, cp_toke | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) #define OACC_EXIT_DATA_CLAUSE_MASK \ @@ -35724,11 +35680,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token * | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) #define OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ @@ -35748,10 +35700,6 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token * | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ Index: gcc/testsuite/c-c++-common/goacc/default-4.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/default-4.c (revision 246808) +++ gcc/testsuite/c-c++-common/goacc/default-4.c (working copy) @@ -8,7 +8,7 @@ void f1 () float f1_b[2]; #pragma acc data copyin (f1_a) copyout (f1_b) - /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f1_b \[^\\)\]+\\) map\\(force_to:f1_a" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f1_b \[^\\)\]+\\) map\\(to:f1_a" 1 "gimple" } } */ { #pragma acc kernels /* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } } */ @@ -29,7 +29,7 @@ void f2 () float f2_b[2]; #pragma acc data copyin (f2_a) copyout (f2_b) - /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f2_b \[^\\)\]+\\) map\\(force_to:f2_a" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f2_b \[^\\)\]+\\) map\\(to:f2_a" 1 "gimple" } } */ { #pragma acc kernels default (none) /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } } */ @@ -50,7 +50,7 @@ void f3 () float f3_b[2]; #pragma acc data copyin (f3_a) copyout (f3_b) - /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f3_b \[^\\)\]+\\) map\\(force_to:f3_a" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f3_b \[^\\)\]+\\) map\\(to:f3_a" 1 "gimple" } } */ { #pragma acc kernels default (present) /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */ Index: gcc/testsuite/c-c++-common/goacc/declare-1.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/declare-1.c (revision 246808) +++ gcc/testsuite/c-c++-common/goacc/declare-1.c (working copy) @@ -19,6 +19,12 @@ int v4; int v5, v6, v7, v8; #pragma acc declare create(v5, v6) copyin(v7, v8) +int v9; +#pragma acc declare present_or_copyin(v9) + +int v10; +#pragma acc declare present_or_create(v10) + void f (void) { @@ -49,6 +55,12 @@ f (void) extern int ve4; #pragma acc declare link(ve4) + extern int ve5; +#pragma acc declare present_or_copyin(ve5) + + extern int ve6; +#pragma acc declare present_or_create(ve6) + int va5; #pragma acc declare copy(va5) Index: gcc/testsuite/c-c++-common/goacc/acc-data-chain.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/acc-data-chain.c (revision 246808) +++ gcc/testsuite/c-c++-common/goacc/acc-data-chain.c (working copy) @@ -20,5 +20,5 @@ int main(int argc, char *argv[]) return 0; } -// { dg-final { scan-tree-dump-times "omp target oacc_data map.force_from:b.0. .len: 400.. map.force_to:a.0. .len: 400.." 1 "gimple" } } +// { dg-final { scan-tree-dump-times "omp target oacc_data map.from:b.0. .len: 400.. map.to:a.0. .len: 400.." 1 "gimple" } } // { dg-final { scan-tree-dump-times "omp target oacc_parallel map.force_present:b.0. .len: 400.. map.firstprivate:b .pointer assign, bias: 0.. map.force_present:a.0. .len: 400.. map.firstprivate:a .pointer assign, bias: 0.." 1 "gimple" } } Index: gcc/testsuite/c-c++-common/goacc/declare-2.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/declare-2.c (revision 246808) +++ gcc/testsuite/c-c++-common/goacc/declare-2.c (working copy) @@ -29,14 +29,8 @@ int v6; #pragma acc declare present_or_copy(v6) /* { dg-error "at file scope" } */ int v7; -#pragma acc declare present_or_copyin(v7) /* { dg-error "at file scope" } */ +#pragma acc declare present_or_copyout(v7) /* { dg-error "at file scope" } */ -int v8; -#pragma acc declare present_or_copyout(v8) /* { dg-error "at file scope" } */ - -int v9; -#pragma acc declare present_or_create(v9) /* { dg-error "at file scope" } */ - int va10; #pragma acc declare create (va10) #pragma acc declare link (va10) /* { dg-error "more than once" } */ @@ -67,13 +61,7 @@ f (void) #pragma acc declare present_or_copy(ve3) /* { dg-error "invalid use of" } */ extern int ve4; -#pragma acc declare present_or_copyin(ve4) /* { dg-error "invalid use of" } */ +#pragma acc declare present_or_copyout(ve4) /* { dg-error "invalid use of" } */ - extern int ve5; -#pragma acc declare present_or_copyout(ve5) /* { dg-error "invalid use of" } */ - - extern int ve6; -#pragma acc declare present_or_create(ve6) /* { dg-error "invalid use of" } */ - -#pragma acc declare present (v9) /* { dg-error "invalid use of" } */ +#pragma acc declare present (v2) /* { dg-error "invalid use of" } */ } Index: gcc/testsuite/gfortran.dg/goacc/declare-2.f95 =================================================================== --- gcc/testsuite/gfortran.dg/goacc/declare-2.f95 (revision 246808) +++ gcc/testsuite/gfortran.dg/goacc/declare-2.f95 (working copy) @@ -11,9 +11,9 @@ subroutine asubr (b) !$acc declare copyout (b) ! { dg-error "Invalid clause in module" } !$acc declare present (b) ! { dg-error "Invalid clause in module" } !$acc declare present_or_copy (b) ! { dg-error "Invalid clause in module" } - !$acc declare present_or_copyin (b) ! { dg-error "Invalid clause in module" } + !$acc declare present_or_copyin (b) ! { dg-error "present on multiple clauses" } !$acc declare present_or_copyout (b) ! { dg-error "Invalid clause in module" } - !$acc declare present_or_create (b) ! { dg-error "Invalid clause in module" } + !$acc declare present_or_create (b) ! { dg-error "present on multiple clauses" } !$acc declare deviceptr (b) ! { dg-error "Invalid clause in module" } !$acc declare create (b) copyin (b) ! { dg-error "present on multiple clauses" } Index: gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 =================================================================== --- gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 (revision 246808) +++ gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 (working copy) @@ -27,10 +27,10 @@ end program test ! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } ! { dg-final { scan-tree-dump-times "reduction\\(max:q\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } } Index: gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 =================================================================== --- gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 (revision 246808) +++ gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 (working copy) @@ -166,5 +166,5 @@ end subroutine test ! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3" 2 "gimple" { xfail *-*-* } } } ! { dg-final { scan-tree-dump-times "acc loop private.i. independent" 2 "gimple" { xfail *-*-* } } } ! { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } } -! { dg-final { scan-tree-dump-times "omp target oacc_\[^ \]+ map.force_tofrom:y" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_\[^ \]+ map.tofrom:y" 2 "gimple" } } ! { dg-final { scan-tree-dump-times "acc loop private.i. reduction..:y." 2 "gimple" { xfail *-*-* } } } Index: gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 =================================================================== --- gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 (revision 246808) +++ gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 (working copy) @@ -38,9 +38,7 @@ program test !$acc end parallel end program test -! { dg-final { scan-tree-dump-times "map.tofrom:v1" 8 "gimple" } } -! { dg-final { scan-tree-dump-times "map.tofrom:v2" 8 "gimple" } } -! { dg-final { scan-tree-dump-times "map.force_tofrom:v1" 1 "gimple" } } -! { dg-final { scan-tree-dump-times "map.force_tofrom:v2" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map.tofrom:v1" 9 "gimple" } } +! { dg-final { scan-tree-dump-times "map.tofrom:v2" 9 "gimple" } } ! { dg-final { scan-tree-dump-times "map.force_present:v1" 1 "gimple" } } ! { dg-final { scan-tree-dump-times "map.force_present:v2" 1 "gimple" } } Index: gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 =================================================================== --- gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 (revision 246808) +++ gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 (working copy) @@ -17,10 +17,10 @@ end program test ! { dg-final { scan-tree-dump-times "if" 1 "original" } } ! { dg-final { scan-tree-dump-times "async" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } } Index: gcc/testsuite/gfortran.dg/goacc/default-4.f =================================================================== --- gcc/testsuite/gfortran.dg/goacc/default-4.f (revision 246808) +++ gcc/testsuite/gfortran.dg/goacc/default-4.f (working copy) @@ -8,7 +8,7 @@ REAL, DIMENSION (2) :: F1_B !$ACC DATA COPYIN (F1_A) COPYOUT (F1_B) -! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f1_a \[^\\)\]+\\) map\\(force_from:f1_b" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f1_a \[^\\)\]+\\) map\\(from:f1_b" 1 "gimple" } } !$ACC KERNELS ! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } } F1_B(1) = F1_A; @@ -26,7 +26,7 @@ REAL, DIMENSION (2) :: F2_B !$ACC DATA COPYIN (F2_A) COPYOUT (F2_B) -! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f2_a \[^\\)\]+\\) map\\(force_from:f2_b" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f2_a \[^\\)\]+\\) map\\(from:f2_b" 1 "gimple" } } !$ACC KERNELS DEFAULT (NONE) ! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } } F2_B(1) = F2_A; @@ -44,7 +44,7 @@ REAL, DIMENSION (2) :: F3_B !$ACC DATA COPYIN (F3_A) COPYOUT (F3_B) -! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f3_a \[^\\)\]+\\) map\\(force_from:f3_b" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f3_a \[^\\)\]+\\) map\\(from:f3_b" 1 "gimple" } } !$ACC KERNELS DEFAULT (PRESENT) ! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } F3_B(1) = F3_A; Index: gcc/testsuite/gfortran.dg/goacc/data-tree.f95 =================================================================== --- gcc/testsuite/gfortran.dg/goacc/data-tree.f95 (revision 246808) +++ gcc/testsuite/gfortran.dg/goacc/data-tree.f95 (working copy) @@ -15,10 +15,10 @@ end program test ! { dg-final { scan-tree-dump-times "pragma acc data" 1 "original" } } ! { dg-final { scan-tree-dump-times "if" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } } Index: libgomp/oacc-parallel.c =================================================================== --- libgomp/oacc-parallel.c (revision 246808) +++ libgomp/oacc-parallel.c (working copy) @@ -492,6 +492,7 @@ GOACC_enter_exit_data (int device, size_t mapnum, } if (kind == GOMP_MAP_DELETE + || kind == GOMP_MAP_FROM || kind == GOMP_MAP_FORCE_FROM || kind == GOMP_MAP_DECLARE_DEALLOCATE) break; @@ -573,6 +574,7 @@ GOACC_enter_exit_data (int device, size_t mapnum, if (acc_is_present (hostaddrs[i], sizes[i])) acc_delete (hostaddrs[i], sizes[i]); break; + case GOMP_MAP_FROM: case GOMP_MAP_FORCE_FROM: acc_copyout (hostaddrs[i], sizes[i]); break; @@ -589,9 +591,9 @@ GOACC_enter_exit_data (int device, size_t mapnum, &sizes[i], &kinds[i]); else if (acc_is_present (hostaddrs[i], sizes[i])) { - gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff) - == GOMP_MAP_FORCE_FROM, async, - pointer); + bool copyfrom = (kind == GOMP_MAP_FORCE_FROM + || kind == GOMP_MAP_FROM); + gomp_acc_remove_pointer (hostaddrs[i], copyfrom, async, pointer); /* See the above comment. */ } i += pointer - 1; Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f (revision 246808) +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f (working copy) @@ -1,16 +0,0 @@ -! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } - - IMPLICIT NONE - INCLUDE "openacc_lib.h" - - INTEGER I - - CALL ACC_PRESENT_OR_COPYIN (I) - WRITE(0, *) "CheCKpOInT" - CALL ACC_COPYIN (I) - - END - -! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "already mapped to" } -! { dg-shouldfail "" } Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f (revision 246808) +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f (working copy) @@ -1,18 +0,0 @@ -! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } - - IMPLICIT NONE - - INTEGER I - -!$ACC DATA CREATE (I) - WRITE(0, *) "CheCKpOInT" -!$ACC PARALLEL COPYIN (I) - I = 0 -!$ACC END PARALLEL -!$ACC END DATA - - END - -! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } -! { dg-shouldfail "" } Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f (revision 246808) +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f (working copy) @@ -1,18 +0,0 @@ -! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } - - IMPLICIT NONE - INCLUDE "openacc_lib.h" - - INTEGER I - - CALL ACC_COPYIN (I) - WRITE(0, *) "CheCKpOInT" -!$ACC DATA COPY (I) - I = 0 -!$ACC END DATA - - END - -! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } -! { dg-shouldfail "" } Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f (revision 246808) +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f (working copy) @@ -1,16 +0,0 @@ -! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } - - IMPLICIT NONE - INCLUDE "openacc_lib.h" - - INTEGER I - -!$ACC ENTER DATA CREATE (I) - WRITE(0, *) "CheCKpOInT" - CALL ACC_COPYIN (I) - - END - -! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "already mapped to" } -! { dg-shouldfail "" } Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f (revision 246808) +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f (working copy) @@ -1,18 +0,0 @@ -! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } - - IMPLICIT NONE - - INTEGER I - -!$ACC DATA PRESENT_OR_COPY (I) - WRITE(0, *) "CheCKpOInT" -!$ACC DATA COPYOUT (I) - I = 0 -!$ACC END DATA -!$ACC END DATA - - END - -! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } -! { dg-shouldfail "" } Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f (revision 246808) +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f (working copy) @@ -1,16 +0,0 @@ -! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } - - IMPLICIT NONE - INCLUDE "openacc_lib.h" - - INTEGER I - - CALL ACC_PRESENT_OR_COPYIN (I) - WRITE(0, *) "CheCKpOInT" -!$ACC ENTER DATA CREATE (I) - - END - -! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "already mapped to" } -! { dg-shouldfail "" } Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f (revision 246808) +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f (working copy) @@ -1,17 +0,0 @@ -! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } - - IMPLICIT NONE - INCLUDE "openacc_lib.h" - - INTEGER I - -!$ACC DATA PRESENT_OR_COPY (I) - WRITE(0, *) "CheCKpOInT" - CALL ACC_COPYIN (I) -!$ACC END DATA - - END - -! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "already mapped to" } -! { dg-shouldfail "" } Index: libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f (revision 246808) +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f (working copy) @@ -1,16 +0,0 @@ -! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } - - IMPLICIT NONE - INCLUDE "openacc_lib.h" - - INTEGER I - -!$ACC ENTER DATA CREATE (I) - WRITE(0, *) "CheCKpOInT" - CALL ACC_CREATE (I) - - END - -! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } -! { dg-output "already mapped to" } -! { dg-shouldfail "" } Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c (revision 246808) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c (working copy) @@ -1,20 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ - -#include -#include - -int -main (int argc, char *argv[]) -{ - int i; - - acc_present_or_copyin (&i, sizeof i); - fprintf (stderr, "CheCKpOInT\n"); - acc_copyin (&i, sizeof i); - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "already mapped to" } */ -/* { dg-shouldfail "" } */ Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c (revision 246808) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c (working copy) @@ -1,20 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ - -#include -#include - -int -main (int argc, char *argv[]) -{ - int i; - -#pragma acc enter data create (i) - fprintf (stderr, "CheCKpOInT\n"); - acc_copyin (&i, sizeof i); - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "already mapped to" } */ -/* { dg-shouldfail "" } */ Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c (revision 246808) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c (working copy) @@ -1,20 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ - -#include -#include - -int -main (int argc, char *argv[]) -{ - int i; - - acc_present_or_copyin (&i, sizeof i); - fprintf (stderr, "CheCKpOInT\n"); -#pragma acc enter data create (i) - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "already mapped to" } */ -/* { dg-shouldfail "" } */ Index: libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c (revision 246808) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c (working copy) @@ -1,4 +1,4 @@ -/* Test if duplicate data mappings with acc_copy_in. */ +/* Test if acc_copyin has present_or_ behavior. */ /* { dg-do run { target openacc_nvidia_accel_selected } } */ @@ -31,5 +31,3 @@ main (int argc, char **argv) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+256\\\] already mapped to \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */ -/* { dg-shouldfail "" } */ Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c (revision 246808) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c (working copy) @@ -1,20 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ - -#include -#include - -int -main (int argc, char *argv[]) -{ - int i; - -#pragma acc enter data create (i) - fprintf (stderr, "CheCKpOInT\n"); - acc_create (&i, sizeof i); - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "already mapped to" } */ -/* { dg-shouldfail "" } */ Index: libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c (revision 246808) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c (working copy) @@ -32,5 +32,3 @@ main (int argc, char **argv) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+256\\\] already mapped to \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */ -/* { dg-shouldfail "" } */ Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c (revision 246808) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c (working copy) @@ -1,22 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ - -#include - -int -main (int argc, char *argv[]) -{ - int i; - -#pragma acc data create (i) - { - fprintf (stderr, "CheCKpOInT\n"); -#pragma acc parallel copyin (i) - ++i; - } - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */ -/* { dg-shouldfail "" } */ Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c (revision 246808) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c (working copy) @@ -1,22 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ - -#include -#include - -int -main (int argc, char *argv[]) -{ - int i; - - acc_copyin (&i, sizeof i); - - fprintf (stderr, "CheCKpOInT\n"); -#pragma acc data copy (i) - ++i; - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */ -/* { dg-shouldfail "" } */ Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c (revision 246808) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c (working copy) @@ -1,22 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ - -#include - -int -main (int argc, char *argv[]) -{ - int i; - -#pragma acc data present_or_copy (i) - { - fprintf (stderr, "CheCKpOInT\n"); -#pragma acc data copyout (i) - ++i; - } - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */ -/* { dg-shouldfail "" } */ Index: libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c (revision 246808) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c (working copy) @@ -1,22 +0,0 @@ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ - -#include -#include - -int -main (int argc, char *argv[]) -{ - int i; - -#pragma acc data present_or_copy (i) - { - fprintf (stderr, "CheCKpOInT\n"); - acc_copyin (&i, sizeof i); - } - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "already mapped to" } */ -/* { dg-shouldfail "" } */ Index: libgomp/oacc-mem.c =================================================================== --- libgomp/oacc-mem.c (revision 246808) +++ libgomp/oacc-mem.c (working copy) @@ -524,39 +524,49 @@ present_create_copy (unsigned f, void *h, size_t s void * acc_create (void *h, size_t s) { - return present_create_copy (FLAG_CREATE, h, s, acc_async_sync); + return present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s, acc_async_sync); } void acc_create_async (void *h, size_t s, int async) { - present_create_copy (FLAG_CREATE, h, s, async); + present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s, async); } +#ifdef HAVE_ATTRIBUTE_ALIAS +extern void *acc_present_or_create (void *, size_t) + __attribute__((alias ("acc_create"))); +#else void * +acc_present_or_create (void *h, size_t s) +{ + return acc_create (h, s); +} +#endif + +void * acc_copyin (void *h, size_t s) { - return present_create_copy (FLAG_CREATE | FLAG_COPY, h, s, acc_async_sync); + return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s, + acc_async_sync); } void acc_copyin_async (void *h, size_t s, int async) { - present_create_copy (FLAG_CREATE | FLAG_COPY, h, s, async); + present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s, async); } +#ifdef HAVE_ATTRIBUTE_ALIAS +extern void *acc_present_or_copyin (void *, size_t) + __attribute__((alias ("acc_copyin"))); +#else void * -acc_present_or_create (void *h, size_t s) -{ - return present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s, acc_async_sync); -} - -void * acc_present_or_copyin (void *h, size_t s) { - return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s, - acc_async_sync); + return acc_copyin (h, s); } +#endif #define FLAG_COPYOUT (1 << 0)