From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 54958 invoked by alias); 9 Nov 2015 14:46:17 -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 54942 invoked by uid 89); 9 Nov 2015 14:46:16 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.6 required=5.0 tests=BAYES_00,FREEMAIL_FROM,RCVD_IN_DNSWL_LOW,SPF_PASS autolearn=ham version=3.3.2 X-HELO: mail-qg0-f49.google.com Received: from mail-qg0-f49.google.com (HELO mail-qg0-f49.google.com) (209.85.192.49) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Mon, 09 Nov 2015 14:46:15 +0000 Received: by qgea14 with SMTP id a14so62879395qge.0 for ; Mon, 09 Nov 2015 06:46:13 -0800 (PST) X-Received: by 10.140.230.6 with SMTP id a6mr33007395qhc.31.1447080373183; Mon, 09 Nov 2015 06:46:13 -0800 (PST) Received: from ?IPv6:2601:181:c000:c497:a2a8:cdff:fe3e:b48? ([2601:181:c000:c497:a2a8:cdff:fe3e:b48]) by smtp.googlemail.com with ESMTPSA id l6sm4995036qgd.16.2015.11.09.06.46.12 (version=TLSv1.2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Mon, 09 Nov 2015 06:46:12 -0800 (PST) Subject: Re: OpenACC Firstprivate To: Jakub Jelinek References: <563E01A4.20607@acm.org> <20151109134619.GQ5675@tucnak.redhat.com> <5640A6B3.3030409@acm.org> <20151109141034.GS5675@tucnak.redhat.com> Cc: GCC Patches , Cesar Philippidis From: Nathan Sidwell Message-ID: <5640B1B4.2070701@acm.org> Date: Mon, 09 Nov 2015 14:46: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: <20151109141034.GS5675@tucnak.redhat.com> Content-Type: text/plain; charset=windows-1252; format=flowed Content-Transfer-Encoding: 7bit X-SW-Source: 2015-11/txt/msg00974.txt.bz2 On 11/09/15 09:10, Jakub Jelinek wrote: > On Mon, Nov 09, 2015 at 08:59:15AM -0500, Nathan Sidwell wrote: >>> This I'm afraid performs often two copies rather than just one (one to copy >>> the host value to the present_copyin mapped value, another one in the >>> region), >> >> I don't think that can be avoided. The host doesn't have control over when >> the CTAs (a gang) start -- they may even be serialized onto the same >> physical HW. So each gang has to initialize its own instance. Or did you >> mean something else? > > So, what is the scope of the private and firstprivate vars in OpenACC? > In OpenMP if a variable is private or firstprivate on the target construct, > unless further privatized in inner constructs it is really shared among all > the threads in all the teams (ro one var per all CTAs/workers in PTX terms). > Is that the case for OpenACC too, or are the vars e.g. private to each CTA > already or to each thread in each CTA, something different? > If they are shared by all CTAs, then you should hopefully be able to use the > GOMP_MAP_FIRSTPRIVATE{,_INT}, if not, then I'd say you should at least use > those to provide you the initializer data to initialize your private vars > from as a cheaper alternative to mapping. I'm going to try and get clarification, but I think the intent is to initialize with the value seen on the device. Consider: int foo = 0; #pragma acc data copyin(foo) { #pragma acc parallel present(foo) { foo = 2; } if (expr){ #pragma update host (foo) } #pragma acc parallel firstprivate (foo) { // which initialization value? } } Here we copy data to the device, then set it a distinct value there. We conditionally update the host's instance from the device. My thinking is that the intent of the firstprivate is to initialize with the value known on the device (and behave as-if copyin, if it's not there). Not the value most recently seen on the host -- the update clause could change that, and may well be being used as a debugging aide, so it seems bizarre that it can change program semantics in such a way.