From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 104354 invoked by alias); 5 Nov 2015 17:13:08 -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 104323 invoked by uid 89); 5 Nov 2015 17:13:07 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.4 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; Thu, 05 Nov 2015 17:13:07 +0000 Received: from svr-orw-fem-05.mgc.mentorg.com ([147.34.97.43]) by relay1.mentorg.com with esmtp id 1ZuO5w-0003UX-1U from Nathan_Sidwell@mentor.com ; Thu, 05 Nov 2015 09:13:04 -0800 Received: from [127.0.0.1] (147.34.91.1) by svr-orw-fem-05.mgc.mentorg.com (147.34.97.43) with Microsoft SMTP Server id 14.3.224.2; Thu, 5 Nov 2015 09:13:03 -0800 Subject: Re: [openacc] tile, independent, default, private and firstprivate support in c/++ To: Thomas Schwinge , Cesar Philippidis References: <5639325B.90006@codesourcery.com> <87k2pwzlqa.fsf@kepler.schwinge.homeip.net> <563B6C1E.8070408@codesourcery.com> <87twp02xdr.fsf@schwinge.name> CC: "gcc-patches@gcc.gnu.org" , Jakub Jelinek From: Nathan Sidwell Message-ID: <563B8E1E.1020400@codesourcery.com> Date: Thu, 05 Nov 2015 17:13:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 In-Reply-To: <87twp02xdr.fsf@schwinge.name> Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: 7bit X-SW-Source: 2015-11/txt/msg00498.txt.bz2 On 11/05/15 12:01, Thomas Schwinge wrote: > On Thu, 5 Nov 2015 06:47:58 -0800, Cesar Philippidis wrote: >> On 11/05/2015 04:14 AM, Thomas Schwinge wrote: >> Sorry, I must have mis-phrased it. The spec is unclear here. There are >> three possible ways to interpret 'acc parallel loop reduction': >> >> 1. acc parallel reduction >> acc loop > > This is what you propose in your patch, but I don't think that makes > sense, or does it? I'm happy to learn otherwise, but in my current > understanding, a reduction clause needs to be attached (at least) to the > innermost construct where reductions are to be processed. (Let's also Correct, the above interpretation must be wrong. > consider multi-level gang/worker/vector loops/reductions.) So, either: > >> 2. acc parallel >> acc loop reduction > > ... this, or even this: > >> 3. acc parallel reduction >> acc loop reduction > > ..., which I'm not sure what the execution model implementation requires. > (Nathan?) interpretation #2 is sufficient, I think. However, both are lacking a 'copy (reduction_var)', clause as otherwise there's nothing changing the default data attribute of 'firstprivate' (working on that patch). Perhaps 'reduction' on 'parallel' is meant to imply that (because that's what makes sense), but the std doesn't say it. In summary it's probably safe to implement interpretation #3. That way we can implement the hypothesis that reductions at the outer construct imply copy. > And while we're at it: the very same question also applies to the private > clause, which -- contrary to all other (as far as I remember) clauses -- > also is applicable to both the parallel and loop constructs: > > #pragma acc parallel loop private([...]) > > ... is to be decomposed into which of the following: > > #pragma acc parallel private([...]) > #pragma acc loop > > #pragma acc parallel > #pragma acc loop private([...]) > > #pragma acc parallel private([...]) > #pragma acc loop private([...]) > > (There is no private clause allowed to be specified with the kernels > construct for what it's worth, but that doesn't mean we couldn't use it > internally, of course, if so required.) I think interpretation #2 or #3 make sense, and I suspect result in the same emitted code. nathan -- Nathan Sidwell