From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 52861 invoked by alias); 19 Mar 2015 14:57:18 -0000 Mailing-List: contact gcc-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-owner@gcc.gnu.org Received: (qmail 52851 invoked by uid 89); 19 Mar 2015 14:57:17 -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,SPF_HELO_PASS,SPF_PASS,T_RP_MATCHES_RCVD autolearn=ham version=3.3.2 X-HELO: mx1.redhat.com Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Thu, 19 Mar 2015 14:57:16 +0000 Received: from int-mx11.intmail.prod.int.phx2.redhat.com (int-mx11.intmail.prod.int.phx2.redhat.com [10.5.11.24]) by mx1.redhat.com (Postfix) with ESMTPS id E074632B8E6; Thu, 19 Mar 2015 14:57:14 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-63.ams2.redhat.com [10.36.116.63]) by int-mx11.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t2JEvC09006381 (version=TLSv1/SSLv3 cipher=AES128-GCM-SHA256 bits=128 verify=NO); Thu, 19 Mar 2015 10:57:14 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.14.9/8.14.9) with ESMTP id t2JEvBqU016405; Thu, 19 Mar 2015 15:57:11 +0100 Received: (from jakub@localhost) by tucnak.zalov.cz (8.14.9/8.14.9/Submit) id t2JEvA1K016377; Thu, 19 Mar 2015 15:57:10 +0100 Date: Thu, 19 Mar 2015 14:57:00 -0000 From: Jakub Jelinek To: Ilya Verbin Cc: gcc@gcc.gnu.org, Kirill Yukhin Subject: Re: [gomp4] Questions about "declare target" and "target update" pragmas Message-ID: <20150319145710.GY1746@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <20140122155151.GA50489@msticlxl57.ims.intel.com> <20150310165252.GC37666@msticlxl57.ims.intel.com> <20150316184153.GA42550@msticlxl57.ims.intel.com> <20150319134744.GW1746@tucnak.redhat.com> <20150319144947.GA20881@msticlxl57.ims.intel.com> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <20150319144947.GA20881@msticlxl57.ims.intel.com> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2015-03/txt/msg00235.txt.bz2 On Thu, Mar 19, 2015 at 05:49:47PM +0300, Ilya Verbin wrote: > > > void foo (int a1[]) > > > { > > > #pragma omp target > > > { > > > a1[10]++; > > > a2[10]++; > > > } > > > } > > > > That is a buggy test. int a1[] function argument is changed > > into int *a1, so it is actually > > #pragma omp target map(tofrom:a1, a2) > > Actually, it copies only a1 pointer, since a2 points to the global array. Sure, the implicit map(tofrom:a2) doesn't really do anything, since a2 is clearly already mapped and will stay to be mapped, and is a global var. > > > { > > a1[10]++; > > a2[10]++; > > } > > which copies the a1 pointer to the device by value (no pointer > > transformation). > > Perhaps the testcase writer meant to use #pragma omp target map(a1[10]) > > instead (or map(a1[0:50])? > > If I understand correctly, it's not allowed to map global target arrays this > way, since it's already present in the initial device data environment: It of course is allowed. It just means that it doesn't allocate new memory (sizeof(int) large in the map(a1[10]) case, sizeof(int)*50 large in the map(a1[0:50]) case), nor copy the bytes around, all it will do is allocate memory for the target copy of the a1 pointer, and do pointer transformation such that the a1 pointer on the target will point to the global target a1 array. Without the map(a1[10]) or map(a1[0:50]) clauses (i.e. implicit map(tofrom:a1)) it does similar thing, except it copies the pointer value to the target (and back at the end of the region) instead, which is not what you want... > So, to fix this testcase I can just remove the "int a1[]" function argument, and > add some "#pragma omp target update" where needed. Well, supposedly the test was meant to test it with a pointer parameter, otherwise why would there be both a1 and a2 when a2 would be enough? Jakub