From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 63744 invoked by alias); 6 Nov 2015 20:16:20 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 63735 invoked by uid 89); 6 Nov 2015 20:16:19 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL,BAYES_00,RCVD_IN_DNSWL_LOW,SPF_PASS autolearn=ham version=3.3.2 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; Fri, 06 Nov 2015 20:16:18 +0000 Received: from svr-orw-fem-02x.mgc.mentorg.com ([147.34.96.206] helo=SVR-ORW-FEM-02.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1ZunQk-0000mp-Hp from James_Norris@mentor.com ; Fri, 06 Nov 2015 12:16:14 -0800 Received: from [172.30.80.115] (147.34.91.1) by svr-orw-fem-02.mgc.mentorg.com (147.34.96.168) with Microsoft SMTP Server id 14.3.224.2; Fri, 6 Nov 2015 12:16:13 -0800 Message-ID: <563D0A8C.3070003@codesourcery.com> Date: Fri, 06 Nov 2015 20:16:00 -0000 From: James Norris User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:31.0) Gecko/20100101 Thunderbird/31.7.0 MIME-Version: 1.0 To: Jakub Jelinek CC: Thomas Schwinge , GCC Patches Subject: Re: [OpenACC] declare directive References: <562FDBF8.1040105@mentor.com> <5638E164.5010207@codesourcery.com> <87611h1zi7.fsf@kepler.schwinge.homeip.net> <563CD07A.3000404@codesourcery.com> <20151106190352.GG5675@tucnak.redhat.com> In-Reply-To: <20151106190352.GG5675@tucnak.redhat.com> X-TagToolbar-Keys: D20151106141611993 Content-Type: text/plain; charset="windows-1252"; format=flowed Content-Transfer-Encoding: 7bit X-SW-Source: 2015-11/txt/msg00725.txt.bz2 Jakub, On 11/06/2015 01:03 PM, Jakub Jelinek wrote: > On Fri, Nov 06, 2015 at 10:08:26AM -0600, James Norris wrote: >> 2015-10-27 James Norris >> Joseph Myers >> >> gcc/ >> * c-family/c-pragma.c (oacc_pragmas): Add entry for declare directive. >> * c-family/c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DECLARE. > c-family/, c/, cp/ have all their own ChangeLog files, and the entries > that go in there should not have the path prefixes. Will fix. > >> + for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t)) >> + { >> + location_t loc = OMP_CLAUSE_LOCATION (t); >> + tree decl = OMP_CLAUSE_DECL (t); >> + tree devres = NULL_TREE; >> + if (!DECL_P (decl)) >> + { >> + error_at (loc, "subarray in %<#pragma acc declare%>"); > Is that the OpenACC term that you use everywhere? In OpenMP > those are array sections. I can change the term so that both OpenACC and OpenMP use the same terminology. > >> + case GOMP_MAP_LINK: >> + if (!global_bindings_p () && !DECL_EXTERNAL (decl)) >> + { >> + error_at (loc, >> + "%qD must be a global variable in" >> + "%<#pragma acc declare link%>", >> + decl); >> + error = true; >> + continue; > What about TREE_STATIC? Will fix. > >> diff --git a/gcc/gimplify.c b/gcc/gimplify.c >> index fa34858..fbd4a69 100644 >> --- a/gcc/gimplify.c >> +++ b/gcc/gimplify.c >> @@ -157,6 +157,7 @@ struct gimplify_omp_ctx >> bool target_map_scalars_firstprivate; >> bool target_map_pointers_as_0len_arrays; >> bool target_firstprivatize_array_bases; >> + gomp_target *declare_returns; >> }; > This declare_returns stuff looks broken. > What exactly do you want to achieve? Why do you record it > in gimplify_omp_ctx, but then only look at it in gimplify_body? > The way you abuse gimplify_omp_ctx for that is just too ugly. > All the gimplification code expects that when you enter some OpenMP/OpenACC > construct, you create new context if it needs one, then process the body > of the construct, then pop it up. The declare_returns stuff > violates all of this. Do you want to enqueue all the statements > at the end of the body? Then just stick it into some gimple_seq > outside of the gimplify_omp_ctx, and process in there. Or if you > want to process it when leaving some construct, arrange for that. For example: void func (int *a) { int b[10]; #pragma acc declare copyout (b) . . . } What I was trying to do was for a copyout to be done at the time the function return'ed. It also may be the case where a mapping operation must be done at entry and exit to a function. I think the approach you outline below (target data/ acc data) is the approach to use. >> @@ -5841,6 +5863,8 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl, >> flags |= GOVD_FIRSTPRIVATE; >> break; >> case OMP_CLAUSE_DEFAULT_UNSPECIFIED: >> + if (is_global_var (decl) && device_resident_p (decl)) >> + flags |= GOVD_MAP_TO_ONLY | GOVD_MAP; > I don't think you want to do this except for (selected or all?) > OpenACC contexts. Say, I don't see why one couldn't e.g. try to mix > OpenMP host parallelization or tasking with OpenACC offloading, > and that affecting in weird way OpenMP semantics. Will fix. > >> + /* Any clauses that affect the state of a variable prior >> + to return are saved and dealt with after the body has >> + been gimplified. */ >> + >> + gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, >> + OACC_DECLARE); >> + >> + c = gimplify_omp_ctxp; >> + gimplify_omp_ctxp = c->outer_context; >> + delete_omp_context (c); > Why don't you call gimplify_adjust_omp_clauses instead? What I was attempting to do was to create to sets of clauses. One used for mapping operations at entry to a function and the other at exit. Obviously, I've made a mess of things with this approach. > >> + stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE, >> + clauses); >> + gimplify_omp_ctxp->declare_returns = stmt; > But the above is the main thing I have trouble with. > What happens if you have multiple #pragma acc declare directives? > Is the stmt from the other ones lost? > I'd expect something like gimple_seq in there instead and pushing > stmts into the sequence. > I don't know what is the semantics of the construct, if it is something > close to say target data in OpenMP or acc data in OpenACC, the addition > of code to unmap the variables is performed using a cleanup, otherwise > how do you handle exceptions, or goto and all other kinds of abnormal > returns from the function. I think the approach you've outlined is the way to go. > >> @@ -160,6 +151,25 @@ varpool_node::get_create (tree decl) >> node->force_output = 1; >> #endif >> } >> +} >> + >> +/* Return varpool node assigned to DECL. Create new one when needed. */ >> +varpool_node * >> +varpool_node::get_create (tree decl) >> +{ >> + varpool_node *node = varpool_node::get (decl); >> + gcc_checking_assert (TREE_CODE (decl) == VAR_DECL); >> + if (node) >> + { >> + if (!node->offloadable) >> + make_offloadable (node, decl); > I don't like this change at all, that is significant extra work > every time somebody calls this function. > Just do what we do on the OpenMP side, if you add some "omp declare target" > or similar attribute and the varpool node for it has been already created, > set offloadable in the FE. Yes setting node->offloadable = 1 is much easier. > > Jakub Thank you for taking time for the review. Will resubmit patch for review after corrections. Thanks! Jim