From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id B1770396ECDC for ; Fri, 25 Sep 2020 15:22:53 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org B1770396ECDC Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Andrew_Stubbs@mentor.com IronPort-SDR: XNTCDMWOTsBQk8n7wyrjvXVlBwGcw1Rhr0YOciZGIDqjpEvkLoRzVrYAiJC47xzfmESxS+fQML 7Jn98ourXY0i+d71SkK09bdX/aSDZnZb9EhmTYml1ZIvYyaDlChC8BAU5yhgCSdctk8/h7s7Xc A2Q1qPnkteLs4LDQbuGBzUuI6BQcFB+xqZn/ZUYVpU3CP0Q8wUYqfCrziUceh6s/2qOK15euh4 l6baOA74uotUmIQa5iKrK3WuhrRYSqaXDAQTeMXgx7k66DwFyxfBMGtjvv2Nio+RiMlah2WygR 4B4= X-IronPort-AV: E=Sophos;i="5.77,302,1596528000"; d="scan'208";a="55517732" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 25 Sep 2020 07:22:52 -0800 IronPort-SDR: 6nQxjnlJg0fao/9QoBGWj2hbeFc5++1A+rpBaUGr0vzbXO7Z5npEvUhAMIkgN0xDWaKxUXRBsx i2X8EJsH2e0noAw6kqYAbpLspIg4K+DhDPxBu0V7lPbolqtoz1qGtqHldf19YUMysZK7b/ngA1 FbP/dsUOGMlwszIu9bXwvi26N4W705UpX4iICe8aQPLzvn6c6tWoNPXxMq2ygFiesBAtNbP4zp fH77Oght/h2Af232iRcUzWPrb9ddij477uNj3Csgt5wVLPtry/dpVBR6nNaiYFriqCTF9UQsiJ vSc= Subject: Re: [PATCH] OpenACC: Separate enter/exit data APIs From: Andrew Stubbs To: "gcc-patches@gcc.gnu.org" References: <048c22ac-051c-415a-4a86-0fc0b0354d0a@codesourcery.com> Message-ID: <8787206d-f7f6-0201-d1fb-4defca282697@codesourcery.com> Date: Fri, 25 Sep 2020 16:22:47 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:68.0) Gecko/20100101 Thunderbird/68.10.0 MIME-Version: 1.0 In-Reply-To: <048c22ac-051c-415a-4a86-0fc0b0354d0a@codesourcery.com> Content-Type: multipart/mixed; boundary="------------F2E91FACC6F887459CD1C26C" Content-Language: en-GB X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-08.mgc.mentorg.com (139.181.222.8) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=3.4 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, NICE_REPLY_A, SPF_HELO_PASS, SPF_PASS, TXREP, UNWANTED_LANGUAGE_BODY autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Level: *** X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , X-List-Received-Date: Fri, 25 Sep 2020 15:22:57 -0000 --------------F2E91FACC6F887459CD1C26C Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: 7bit On 30/07/2020 12:10, Andrew Stubbs wrote: > On 29/07/2020 15:05, Andrew Stubbs wrote: >> This patch does not implement anything new, but simply separates >> OpenACC 'enter data' and 'exit data' into two libgomp API functions. >> The original API name is kept for backward compatibility, but no >> longer referenced by the compiler. >> >> The previous implementation assumed that it would always be possible >> to infer which kind of pragma it was dealing with from the context, >> but there are a few exceptions, and I want to add one more: >> zero-length arrays. >> >> By cleaning this up I will be free to add the new feature without the >> reference counting getting broken. This update fixes a new conflict and updates the patterns in a number of testcases that were affected. OK to commit? Andrew --------------F2E91FACC6F887459CD1C26C Content-Type: text/x-patch; charset="UTF-8"; name="200925-separate-enter-exit.patch" Content-Transfer-Encoding: 7bit Content-Disposition: inline; filename="200925-separate-enter-exit.patch" OpenACC: Separate enter/exit data APIs Move the OpenACC enter and exit data directives from using a single builtin to having one each. For most purposes it was easy to tell which was which, from the directives given, but there are some exceptions. In particular, zero-length array copies are indistiguishable, but we still want reference counting to work. gcc/ChangeLog: * gimple-pretty-print.c (dump_gimple_omp_target): Replace GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA with GF_OMP_TARGET_KIND_OACC_ENTER_DATA and GF_OMP_TARGET_KIND_OACC_EXIT_DATA. * gimple.h (enum gf_mask): Likewise. (is_gimple_omp_oacc): Likewise. * gimplify.c (gimplify_omp_target_update): Likewise. * omp-builtins.def (BUILT_IN_GOACC_ENTER_EXIT_DATA): Delete. (BUILT_IN_GOACC_ENTER_DATA): Add new. (BUILT_IN_GOACC_EXIT_DATA): Add new. * omp-expand.c (expand_omp_target): Replace GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA with GF_OMP_TARGET_KIND_OACC_ENTER_DATA and GF_OMP_TARGET_KIND_OACC_EXIT_DATA. (build_omp_regions_1): Likewise. (omp_make_gimple_edges): Likewise. * omp-low.c (check_omp_nesting_restrictions): Likewise. (lower_omp_target): Likewise. gcc/testsuite/ChangeLog: * c-c++-common/goacc-gomp/nesting-fail-1.c: Adjust patterns. * c-c++-common/goacc/finalize-1.c: Adjust patterns. * c-c++-common/goacc/mdc-1.c: Adjust patterns. * c-c++-common/goacc/nesting-fail-1.c: Adjust patterns. * c-c++-common/goacc/struct-enter-exit-data-1.c: Adjust patterns. libgomp/ChangeLog: * libgomp.map: Add GOACC_enter_data and GOACC_exit_data. * libgomp_g.h (GOACC_enter_exit_data): Delete. (GOACC_enter_data): New prototype. (GOACC_exit_data) New prototype.: * oacc-mem.c (GOACC_enter_exit_data): Move most of the content ... (GOACC_enter_exit_data_internal): ... here. (GOACC_enter_data): New function. (GOACC_exit_data) New function.: * oacc-parallel.c (GOACC_declare): Replace GOACC_enter_exit_data with GOACC_enter_data and GOACC_exit_data. * testsuite/libgomp.oacc-c-c++-common/lib-26.c: Delete file. * testsuite/libgomp.oacc-c-c++-common/lib-36.c: Delete file. * testsuite/libgomp.oacc-c-c++-common/lib-40.c: Delete file. diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c index a01bf901657..26978ec1ab5 100644 --- a/gcc/gimple-pretty-print.c +++ b/gcc/gimple-pretty-print.c @@ -1691,8 +1691,11 @@ dump_gimple_omp_target (pretty_printer *buffer, const gomp_target *gs, case GF_OMP_TARGET_KIND_OACC_UPDATE: kind = " oacc_update"; break; - case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: - kind = " oacc_enter_exit_data"; + case GF_OMP_TARGET_KIND_OACC_ENTER_DATA: + kind = " oacc_enter_data"; + break; + case GF_OMP_TARGET_KIND_OACC_EXIT_DATA: + kind = " oacc_exit_data"; break; case GF_OMP_TARGET_KIND_OACC_DECLARE: kind = " oacc_declare"; diff --git a/gcc/gimple.h b/gcc/gimple.h index 6cc7e66059d..3f17b1c0739 100644 --- a/gcc/gimple.h +++ b/gcc/gimple.h @@ -171,9 +171,10 @@ enum gf_mask { GF_OMP_TARGET_KIND_OACC_SERIAL = 7, GF_OMP_TARGET_KIND_OACC_DATA = 8, GF_OMP_TARGET_KIND_OACC_UPDATE = 9, - GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 10, + GF_OMP_TARGET_KIND_OACC_ENTER_DATA = 10, GF_OMP_TARGET_KIND_OACC_DECLARE = 11, GF_OMP_TARGET_KIND_OACC_HOST_DATA = 12, + GF_OMP_TARGET_KIND_OACC_EXIT_DATA = 13, GF_OMP_TEAMS_HOST = 1 << 0, /* True on an GIMPLE_OMP_RETURN statement if the return does not require @@ -6482,7 +6483,8 @@ is_gimple_omp_oacc (const gimple *stmt) case GF_OMP_TARGET_KIND_OACC_SERIAL: case GF_OMP_TARGET_KIND_OACC_DATA: case GF_OMP_TARGET_KIND_OACC_UPDATE: - case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_ENTER_DATA: + case GF_OMP_TARGET_KIND_OACC_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: case GF_OMP_TARGET_KIND_OACC_HOST_DATA: return true; diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 2dea03cce3d..8fcba8b5b18 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -12976,8 +12976,11 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p) switch (TREE_CODE (expr)) { case OACC_ENTER_DATA: + kind = GF_OMP_TARGET_KIND_OACC_ENTER_DATA; + ort = ORT_ACC; + break; case OACC_EXIT_DATA: - kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA; + kind = GF_OMP_TARGET_KIND_OACC_EXIT_DATA; ort = ORT_ACC; break; case OACC_UPDATE: diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index f461d60e52b..ab45eecb752 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -35,7 +35,10 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start", BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end", BT_FN_VOID, ATTR_NOTHROW_LIST) -DEF_GOACC_BUILTIN (BUILT_IN_GOACC_ENTER_EXIT_DATA, "GOACC_enter_exit_data", +DEF_GOACC_BUILTIN (BUILT_IN_GOACC_ENTER_DATA, "GOACC_enter_data", + BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR, + ATTR_NOTHROW_LIST) +DEF_GOACC_BUILTIN (BUILT_IN_GOACC_EXIT_DATA, "GOACC_exit_data", BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR, ATTR_NOTHROW_LIST) DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel_keyed", diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index 8f1286e3176..70880b78ffc 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -8917,7 +8917,8 @@ expand_omp_target (struct omp_region *region) case GF_OMP_TARGET_KIND_OACC_KERNELS: case GF_OMP_TARGET_KIND_OACC_SERIAL: case GF_OMP_TARGET_KIND_OACC_UPDATE: - case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_ENTER_DATA: + case GF_OMP_TARGET_KIND_OACC_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: data_region = false; break; @@ -9182,8 +9183,11 @@ expand_omp_target (struct omp_region *region) case GF_OMP_TARGET_KIND_OACC_UPDATE: start_ix = BUILT_IN_GOACC_UPDATE; break; - case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: - start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA; + case GF_OMP_TARGET_KIND_OACC_ENTER_DATA: + start_ix = BUILT_IN_GOACC_ENTER_DATA; + break; + case GF_OMP_TARGET_KIND_OACC_EXIT_DATA: + start_ix = BUILT_IN_GOACC_EXIT_DATA; break; case GF_OMP_TARGET_KIND_OACC_DECLARE: start_ix = BUILT_IN_GOACC_DECLARE; @@ -9381,7 +9385,8 @@ expand_omp_target (struct omp_region *region) oacc_set_fn_attrib (child_fn, clauses, &args); tagging = true; /* FALLTHRU */ - case BUILT_IN_GOACC_ENTER_EXIT_DATA: + case BUILT_IN_GOACC_ENTER_DATA: + case BUILT_IN_GOACC_EXIT_DATA: case BUILT_IN_GOACC_UPDATE: { tree t_async = NULL_TREE; @@ -9654,7 +9659,8 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent, case GF_OMP_TARGET_KIND_ENTER_DATA: case GF_OMP_TARGET_KIND_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_UPDATE: - case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_ENTER_DATA: + case GF_OMP_TARGET_KIND_OACC_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: /* ..., other than for those stand-alone directives... */ region = NULL; @@ -9908,7 +9914,8 @@ omp_make_gimple_edges (basic_block bb, struct omp_region **region, case GF_OMP_TARGET_KIND_ENTER_DATA: case GF_OMP_TARGET_KIND_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_UPDATE: - case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_ENTER_DATA: + case GF_OMP_TARGET_KIND_OACC_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: cur_region = cur_region->outer; break; diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 3d2a9d77c1c..12346689f36 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -3387,8 +3387,10 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) case GF_OMP_TARGET_KIND_OACC_SERIAL: stmt_name = "serial"; break; case GF_OMP_TARGET_KIND_OACC_DATA: stmt_name = "data"; break; case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break; - case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: - stmt_name = "enter/exit data"; break; + case GF_OMP_TARGET_KIND_OACC_ENTER_DATA: + stmt_name = "enter data"; break; + case GF_OMP_TARGET_KIND_OACC_EXIT_DATA: + stmt_name = "exit data"; break; case GF_OMP_TARGET_KIND_OACC_DECLARE: stmt_name = "declare"; break; case GF_OMP_TARGET_KIND_OACC_HOST_DATA: stmt_name = "host_data"; break; @@ -11344,7 +11346,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GF_OMP_TARGET_KIND_OACC_KERNELS: case GF_OMP_TARGET_KIND_OACC_SERIAL: case GF_OMP_TARGET_KIND_OACC_UPDATE: - case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: + case GF_OMP_TARGET_KIND_OACC_ENTER_DATA: + case GF_OMP_TARGET_KIND_OACC_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: data_region = false; break; diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c index 1a3324200e2..ddbd247342f 100644 --- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c +++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c @@ -162,8 +162,8 @@ f_omp (void) #pragma acc data /* { dg-error "OpenACC .data. construct inside of OpenMP .target. region" } */ ; #pragma acc update host(i) /* { dg-error "OpenACC .update. construct inside of OpenMP .target. region" } */ -#pragma acc enter data copyin(i) /* { dg-error "OpenACC .enter/exit data. construct inside of OpenMP .target. region" } */ -#pragma acc exit data delete(i) /* { dg-error "OpenACC .enter/exit data. construct inside of OpenMP .target. region" } */ +#pragma acc enter data copyin(i) /* { dg-error "OpenACC .enter data. construct inside of OpenMP .target. region" } */ +#pragma acc exit data delete(i) /* { dg-error "OpenACC .exit data. construct inside of OpenMP .target. region" } */ #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 0; i < 2; ++i) ; diff --git a/gcc/testsuite/c-c++-common/goacc/finalize-1.c b/gcc/testsuite/c-c++-common/goacc/finalize-1.c index 3d64b2e7cb3..54bf1b76a1b 100644 --- a/gcc/testsuite/c-c++-common/goacc/finalize-1.c +++ b/gcc/testsuite/c-c++-common/goacc/finalize-1.c @@ -13,25 +13,25 @@ void f () { #pragma acc exit data delete (del_r) /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } } - { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ #pragma acc exit data finalize delete (del_f) /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } } - { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } */ + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } */ #pragma acc exit data finalize delete (del_f_p[2:5]) /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(del_f_p \\+ 2\\) \\\[len: 5\\\]\\) map\\(firstprivate:del_f_p \\\[pointer assign, bias: 2\\\]\\) finalize;$" 1 "original" } } - { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\[^ \]+ \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */ + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:\[^ \]+ \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */ #pragma acc exit data copyout (cpo_r) /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } } - { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ #pragma acc exit data copyout (cpo_f) finalize /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:cpo_f\\);$" 1 "original" } } - { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ #pragma acc exit data copyout (cpo_f_p[4:10]) finalize /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:\\*\\(cpo_f_p \\+ 4\\) \\\[len: 10\\\]\\) map\\(firstprivate:cpo_f_p \\\[pointer assign, bias: 4\\\]\\);$" 1 "original" } } - { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:\[^ \]+ \\\[len: 10\\\]\\)$" 1 "gimple" } } */ + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data finalize map\\(force_from:\[^ \]+ \\\[len: 10\\\]\\)$" 1 "gimple" } } */ } diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c index 337c1f7cc77..d961c77e3bb 100644 --- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c @@ -43,14 +43,14 @@ t1 () } } -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.to:s .len: 32.." 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 0.." 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 0.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.attach:a .bias: 0.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.to:a .len: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.detach:s.e .bias: 0.." 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 0.." 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .bias: 0.." 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .bias: 0.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.release:a .len: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data finalize map.force_detach:a .bias: 0.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data finalize map.force_detach:s.a .bias: 0.." 1 "omplower" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c index 93a911120d4..5cfb327f4b6 100644 --- a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c +++ b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c @@ -14,8 +14,8 @@ f_acc_parallel (void) #pragma acc data /* { dg-error ".data. construct inside of .parallel. region" } */ ; #pragma acc update host(i) /* { dg-error ".update. construct inside of .parallel. region" } */ -#pragma acc enter data copyin(i) /* { dg-error ".enter/exit data. construct inside of .parallel. region" } */ -#pragma acc exit data delete(i) /* { dg-error ".enter/exit data. construct inside of .parallel. region" } */ +#pragma acc enter data copyin(i) /* { dg-error ".enter data. construct inside of .parallel. region" } */ +#pragma acc exit data delete(i) /* { dg-error ".exit data. construct inside of .parallel. region" } */ } } @@ -33,8 +33,8 @@ f_acc_kernels (void) #pragma acc data /* { dg-error ".data. construct inside of .kernels. region" } */ ; #pragma acc update host(i) /* { dg-error ".update. construct inside of .kernels. region" } */ -#pragma acc enter data copyin(i) /* { dg-error ".enter/exit data. construct inside of .kernels. region" } */ -#pragma acc exit data delete(i) /* { dg-error ".enter/exit data. construct inside of .kernels. region" } */ +#pragma acc enter data copyin(i) /* { dg-error ".enter data. construct inside of .kernels. region" } */ +#pragma acc exit data delete(i) /* { dg-error ".exit data. construct inside of .kernels. region" } */ } } diff --git a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c index df405e448b2..9e5d3f2c9d2 100644 --- a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c +++ b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c @@ -20,8 +20,8 @@ test (int *b, int *c, int *e) struct str s = { .a = 0, .b = b, .c = c, .d = 0, .e = e, .f = 0 }; #pragma acc enter data copyin(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f) - /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */ + /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */ #pragma acc exit data copyout(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f) - /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */ + /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */ } diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index c808e810702..3965f036c43 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -517,6 +517,8 @@ GOACC_2.0 { global: GOACC_data_end; GOACC_data_start; + GOACC_enter_data; + GOACC_exit_data; GOACC_enter_exit_data; GOACC_parallel; GOACC_update; diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h index 59e3697bfd8..51df36bc8db 100644 --- a/libgomp/libgomp_g.h +++ b/libgomp/libgomp_g.h @@ -363,8 +363,10 @@ extern void GOACC_wait (int, int, ...); /* oacc-mem.c */ -extern void GOACC_enter_exit_data (int, size_t, void **, size_t *, - unsigned short *, int, int, ...); +extern void GOACC_enter_data (int, size_t, void **, size_t *, + unsigned short *, int, int, ...); +extern void GOACC_exit_data (int, size_t, void **, size_t *, + unsigned short *, int, int, ...); /* oacc-parallel.c */ diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 65757ab2ffc..a46c6a02626 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1317,56 +1317,22 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, gomp_mutex_unlock (&acc_dev->lock); } -void -GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, - size_t *sizes, unsigned short *kinds, int async, - int num_waits, ...) +static void +GOACC_enter_exit_data_internal (int flags_m, size_t mapnum, void **hostaddrs, + size_t *sizes, unsigned short *kinds, + int async, bool data_enter, int num_waits, + va_list *ap) { int flags = GOACC_FLAGS_UNMARSHAL (flags_m); struct goacc_thread *thr; struct gomp_device_descr *acc_dev; - bool data_enter = false; - size_t i; goacc_lazy_initialize (); thr = goacc_thread (); acc_dev = thr->dev; - /* Determine if this is an "acc enter data". */ - for (i = 0; i < mapnum; ++i) - { - unsigned char kind = kinds[i] & 0xff; - - if (kind == GOMP_MAP_POINTER - || kind == GOMP_MAP_TO_PSET - || kind == GOMP_MAP_STRUCT) - continue; - - if (kind == GOMP_MAP_FORCE_ALLOC - || kind == GOMP_MAP_FORCE_PRESENT - || kind == GOMP_MAP_ATTACH - || kind == GOMP_MAP_FORCE_TO - || kind == GOMP_MAP_TO - || kind == GOMP_MAP_ALLOC) - { - data_enter = true; - break; - } - - if (kind == GOMP_MAP_RELEASE - || kind == GOMP_MAP_DELETE - || kind == GOMP_MAP_DETACH - || kind == GOMP_MAP_FORCE_DETACH - || kind == GOMP_MAP_FROM - || kind == GOMP_MAP_FORCE_FROM) - break; - - gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", - kind); - } - bool profiling_p = GOACC_PROFILING_DISPATCH_P (true); acc_prof_info prof_info; @@ -1430,13 +1396,7 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, } if (num_waits) - { - va_list ap; - - va_start (ap, num_waits); - goacc_wait (async, num_waits, &ap); - va_end (ap); - } + goacc_wait (async, num_waits, ap); goacc_aq aq = get_goacc_asyncqueue (async); @@ -1458,3 +1418,77 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, thr->api_info = NULL; } } + +void +GOACC_enter_data (int flags_m, size_t mapnum, void **hostaddrs, + size_t *sizes, unsigned short *kinds, int async, + int num_waits, ...) +{ + va_list ap; + va_start (ap, num_waits); + GOACC_enter_exit_data_internal (flags_m, mapnum, hostaddrs, sizes, kinds, + async, true, num_waits, &ap); + va_end (ap); +} + +void +GOACC_exit_data (int flags_m, size_t mapnum, void **hostaddrs, + size_t *sizes, unsigned short *kinds, int async, + int num_waits, ...) +{ + va_list ap; + va_start (ap, num_waits); + GOACC_enter_exit_data_internal (flags_m, mapnum, hostaddrs, sizes, kinds, + async, false, num_waits, &ap); + va_end (ap); +} + +/* This function is not used. It is provided for backwards compatibility + with older user-binaries only. */ + +void +GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, + size_t *sizes, unsigned short *kinds, int async, + int num_waits, ...) +{ + bool data_enter = false; + + /* Determine if this is an "acc enter data". */ + for (int i = 0; i < mapnum; ++i) + { + unsigned char kind = kinds[i] & 0xff; + + if (kind == GOMP_MAP_POINTER + || kind == GOMP_MAP_TO_PSET + || kind == GOMP_MAP_STRUCT) + continue; + + if (kind == GOMP_MAP_FORCE_ALLOC + || kind == GOMP_MAP_FORCE_PRESENT + || kind == GOMP_MAP_ATTACH + || kind == GOMP_MAP_FORCE_TO + || kind == GOMP_MAP_TO + || kind == GOMP_MAP_ALLOC) + { + data_enter = true; + break; + } + + if (kind == GOMP_MAP_RELEASE + || kind == GOMP_MAP_DELETE + || kind == GOMP_MAP_DETACH + || kind == GOMP_MAP_FORCE_DETACH + || kind == GOMP_MAP_FROM + || kind == GOMP_MAP_FORCE_FROM) + break; + + gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", + kind); + } + + va_list ap; + va_start (ap, num_waits); + GOACC_enter_exit_data_internal (flags_m, mapnum, hostaddrs, sizes, kinds, + async, data_enter, num_waits, &ap); + va_end (ap); +} diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index c7e46e35bd6..bca31b51427 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -745,12 +745,15 @@ GOACC_declare (int flags_m, size_t mapnum, switch (kind) { case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_FROM: case GOMP_MAP_FORCE_TO: - case GOMP_MAP_POINTER: + GOACC_enter_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); + break; + + case GOMP_MAP_FORCE_FROM: case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], + GOACC_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], &kinds[i], GOMP_ASYNC_SYNC, 0); break; @@ -759,19 +762,19 @@ GOACC_declare (int flags_m, size_t mapnum, case GOMP_MAP_ALLOC: if (!acc_is_present (hostaddrs[i], sizes[i])) - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); + GOACC_enter_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); break; case GOMP_MAP_TO: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); + GOACC_enter_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); break; case GOMP_MAP_FROM: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); + GOACC_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); break; case GOMP_MAP_FORCE_PRESENT: diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-26.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-26.c deleted file mode 100644 index 8e1a911abd2..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-26.c +++ /dev/null @@ -1,30 +0,0 @@ -/* { dg-do run } */ - -#include -#include -#include - -int -main (int argc, char **argv) -{ - const int N = 256; - unsigned char *h; - void *d; - - h = (unsigned char *) malloc (N); - - fprintf (stderr, "CheCKpOInT\n"); - d = acc_create (h, 0); - if (!d) - abort (); - - acc_delete (h, N); - - free (h); - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+0\\\] is a bad range" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-36.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-36.c deleted file mode 100644 index 8ff61c76491..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-36.c +++ /dev/null @@ -1,30 +0,0 @@ -/* { dg-do run } */ - -#include -#include -#include - -int -main (int argc, char **argv) -{ - const int N = 256; - unsigned char *h; - void *d; - - h = (unsigned char *) malloc (N); - - fprintf (stderr, "CheCKpOInT\n"); - d = acc_present_or_create (h, 0); - if (!d) - abort (); - - acc_delete (h, N); - - free (h); - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+0\\\] is a bad range" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-40.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-40.c deleted file mode 100644 index f4ab6d8fc3e..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-40.c +++ /dev/null @@ -1,46 +0,0 @@ -/* { dg-do run } */ - -#include -#include -#include -#include -#include - -int -main (int argc, char **argv) -{ - const int N = 256; - int i; - unsigned char *h; - void *d; - - h = (unsigned char *) malloc (N); - - for (i = 0; i < N; i++) - { - h[i] = i; - } - - fprintf (stderr, "CheCKpOInT\n"); - d = acc_present_or_copyin (h, 0); - if (!d) - abort (); - - memset (&h[0], 0, N); - - acc_copyout (h, N); - - for (i = 0; i < N; i++) - { - if (h[i] != i) - abort (); - } - - free (h); - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+0\\\] is a bad range" } */ -/* { dg-shouldfail "" } */ --------------F2E91FACC6F887459CD1C26C--