From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 44021 invoked by alias); 13 Nov 2015 15:31:28 -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 44007 invoked by uid 89); 13 Nov 2015 15:31:28 -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,RP_MATCHES_RCVD,SPF_HELO_PASS 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; Fri, 13 Nov 2015 15:31:27 +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 B4F74C0C9A61; Fri, 13 Nov 2015 15:31:25 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-34.ams2.redhat.com [10.36.116.34]) by int-mx11.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id tADFVNsV026915 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Fri, 13 Nov 2015 10:31:25 -0500 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.15.2/8.15.2) with ESMTP id tADFVLqU013373; Fri, 13 Nov 2015 16:31:21 +0100 Received: (from jakub@localhost) by tucnak.zalov.cz (8.15.2/8.15.2/Submit) id tADFVI2J013372; Fri, 13 Nov 2015 16:31:18 +0100 Date: Fri, 13 Nov 2015 15:31:00 -0000 From: Jakub Jelinek To: Julian Brown Cc: James Norris , GCC Patches , "Joseph S. Myers" , Nathan Sidwell Subject: Re: [Bulk] [OpenACC 0/7] host_data construct Message-ID: <20151113153118.GR5675@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <56293476.5020801@codesourcery.com> <562A578E.4080907@codesourcery.com> <20151026183422.GW478@tucnak.redhat.com> <20151102183339.365c3d33@octopus> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <20151102183339.365c3d33@octopus> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2015-11/txt/msg01721.txt.bz2 On Mon, Nov 02, 2015 at 06:33:39PM +0000, Julian Brown wrote: > Firstly, on trunk at least, use_device_ptr variables are restricted to > pointer or array types: that restriction doesn't exist in OpenACC, nor > actually could I find it in the OpenMP 4.1 document (my guess is the > standards are supposed to match in this regard). I think that a program > such as this should work: So, after talking about this on omp-lang, it seems there is agreement that only arrays and pointer types (or reference to arrays or pointers) should be allowed in use_device_ptr clause and that for pointers/reference to pointers it should probably act the way I've coded it up, i.e. that for them it translates the pointer to point to corresponding object to the one to which it points on the host. It is too late to change the standard now, but will be changed soon, and hopefully clarified in examples. > void target_fn (int *targ_data); > > int > main (int argc, char *argv[]) > { > char out; > int myvar; > #pragma omp target enter data map(to: myvar) > > #pragma omp target data use_device_ptr(myvar) map(from:out) > { > target_fn (&myvar); > out = 5; > } > > return 0; > } That would make the above non-conforming for OpenMP. > Secondly, attempts to use use_device_ptr on (e.g. > dynamically-allocated) arrays accessed through a pointer cause an ICE > with the existing trunk OpenMP code: > > #include > > void target_fn (char *targ_data); > > int > main (int argc, char *argv[]) > { > char *myarr, out; > > myarr = malloc (1024); > > #pragma omp target data map(to: myarr[0:1024]) > { > #pragma omp target data use_device_ptr(myarr) map(from:out) > { > target_fn (myarr); > out = 5; > } > } > > return 0; > } Can't reproduce this ICE (at least not on gomp-4_5-branch, but there aren't significant changes from the trunk there). > Furthermore, this looks strange to me (006t.omplower): > > .omp_data_arr.5.out = &out; > myarr.8 = myarr; > .omp_data_arr.5.myarr = myarr.8; > #pragma omp target data map(from:out [len: 1]) use_device_ptr(myarr) > { > D.2436 = .omp_data_arr.5.myarr; > myarr = D.2436; > > That's clobbering the original myarr variable, right? Just use -fdump-tree-omplower-uid to see that it is a different variable. Basically, for OpenMP use_device_ptr creates a private copy of the pointer for the body of the target data construct, and that pointer is assigned the target device's address. For arrays the implementation creates an artificial pointer variable (holding the start of the array initially) and replaces all references to the array in the target data body with dereference of the pointer. Jakub