From mboxrd@z Thu Jan  1 00:00:00 1970
Return-Path: <gcc-patches-return-413246-listarch-gcc-patches=gcc.gnu.org@gcc.gnu.org>
Received: (qmail 52195 invoked by alias); 9 Nov 2015 14:06:14 -0000
Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm
Precedence: bulk
List-Id: <gcc-patches.gcc.gnu.org>
List-Archive: <http://gcc.gnu.org/ml/gcc-patches/>
List-Post: <mailto:gcc-patches@gcc.gnu.org>
List-Help: <mailto:gcc-patches-help@gcc.gnu.org>
Sender: gcc-patches-owner@gcc.gnu.org
Received: (qmail 52182 invoked by uid 89); 9 Nov 2015 14:06:13 -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-f52.google.com
Received: from mail-qg0-f52.google.com (HELO mail-qg0-f52.google.com) (209.85.192.52) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Mon, 09 Nov 2015 14:06:08 +0000
Received: by qgea14 with SMTP id a14so61823941qge.0        for <gcc-patches@gcc.gnu.org>; Mon, 09 Nov 2015 06:06:06 -0800 (PST)
X-Received: by 10.140.145.150 with SMTP id 144mr31624977qhr.82.1447077965995;        Mon, 09 Nov 2015 06:06:05 -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 a84sm4933693qkb.43.2015.11.09.06.06.05        (version=TLSv1.2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128);        Mon, 09 Nov 2015 06:06:05 -0800 (PST)
Subject: Re: OpenACC Firstprivate
To: Jakub Jelinek <jakub@redhat.com>
References: <563E01A4.20607@acm.org> <20151109134619.GQ5675@tucnak.redhat.com> <5640A6B3.3030409@acm.org>
Cc: GCC Patches <gcc-patches@gcc.gnu.org>, Cesar Philippidis <cesar_philippidis@mentor.com>
From: Nathan Sidwell <nathan@acm.org>
Message-ID: <5640A84D.30902@acm.org>
Date: Mon, 09 Nov 2015 14:06: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: <5640A6B3.3030409@acm.org>
Content-Type: text/plain; charset=windows-1252; format=flowed
Content-Transfer-Encoding: 7bit
X-SW-Source: 2015-11/txt/msg00962.txt.bz2

On 11/09/15 08:59, Nathan Sidwell wrote:
> On 11/09/15 08:46, Jakub Jelinek wrote:
>> On Sat, Nov 07, 2015 at 08:50:28AM -0500, Nathan Sidwell wrote:

>
>> Say
>>    int arr[64];
>> // initialize arr
>> #pragma acc data copyin (arr)
>> {
>>    // modify arr on the host
>>    # pragma acc parallel firstprivate (arr)
>>    {
>>      ...
>>    }
>> }
>
> Hm, I suspect that is either ill formed or the std does not contemplate.

just realized, there are two ways to consider the above.

1) it's  ill formed.   Once you've transferred data to the device, modifying it 
on the host is unspecified.  I'm having trouble finding words in the std that 
actually say that though :(

2) on a system with shared physical global memory, the host modification would 
be visiable on the device (possibly at an arbitrary point due to lack of 
synchronization primitive?)

I don't think this changes 'why not use OpenMP's ...' question, because IIUC you 
think that can be made to DTRT anyway?

nathan