From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 27916 invoked by alias); 22 Oct 2014 08:33:04 -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 27905 invoked by uid 89); 22 Oct 2014 08:33:03 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.8 required=5.0 tests=AWL,BAYES_00,RP_MATCHES_RCVD,SPF_HELO_PASS,SPF_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; Wed, 22 Oct 2014 08:33:02 +0000 Received: from int-mx10.intmail.prod.int.phx2.redhat.com (int-mx10.intmail.prod.int.phx2.redhat.com [10.5.11.23]) by mx1.redhat.com (8.14.4/8.14.4) with ESMTP id s9M8WxpT008617 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=FAIL); Wed, 22 Oct 2014 04:32:59 -0400 Received: from tucnak.zalov.cz (ovpn-116-116.ams2.redhat.com [10.36.116.116]) by int-mx10.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id s9M8Wvrx023695 (version=TLSv1/SSLv3 cipher=AES128-GCM-SHA256 bits=128 verify=NO); Wed, 22 Oct 2014 04:32:59 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.14.9/8.14.9) with ESMTP id s9M8Wtnp014573; Wed, 22 Oct 2014 10:32:56 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.14.9/8.14.9/Submit) id s9M8Ws4I014572; Wed, 22 Oct 2014 10:32:54 +0200 Date: Wed, 22 Oct 2014 08:34:00 -0000 From: Jakub Jelinek To: Richard Biener Cc: Bernd Schmidt , Jeff Law , GCC Patches Subject: Re: The nvptx port [1/11+] indirect jumps Message-ID: <20141022083253.GI10376@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <54451994.9070209@codesourcery.com> <544519D8.70606@codesourcery.com> <5446A55A.6060807@redhat.com> <5446C973.3040102@codesourcery.com> <20141021213025.GE10376@tucnak.redhat.com> <5446D0D4.9000706@codesourcery.com> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2014-10/txt/msg02214.txt.bz2 On Wed, Oct 22, 2014 at 10:18:49AM +0200, Richard Biener wrote: > On Tue, Oct 21, 2014 at 11:32 PM, Bernd Schmidt wrote: > > On 10/21/2014 11:30 PM, Jakub Jelinek wrote: > >> > >> At least for OpenMP, the best would be if the #pragma omp target regions > >> and/or #pragma omp declare target functions contain anything a particular > >> offloading accelerator can't handle, instead of failing the whole > >> compilation perhaps just emit some at least by default non-fatal warning > >> and not emit anything for the particular offloading target, which would > >> mean > >> either host fallback, or, if some other offloading target succeeded, just > >> that target. > > > > > > I guess a test could be added to mkoffload if gcc were to return a different > > value for a sorry vs. any other compilation failure. The tool could then > > choose not to produce offloading support for that target. > > But that would be for the whole file instead of for the specific region? > > So maybe we should produce one LTO offload object for each offload > function and make the symbols they are supposed to provide weak > so a fail doesn't end up failing to link the main program? > > Looks like this gets somewhat awkward with the LTO setup. I don't think we want to do a fine-grained granularity here, it will only lead to significant nightmares. E.g. a target region can call other target functions, if a target function it calls (perhaps directly through a series of other target functions, perhaps indirectly through function pointers etc.) can't be supported by the host, you'd need to give up on offloading all target regions that do or could invoke that. That can be in another TU within the same shared library etc. And, if some regions are emitted and others are not, #pragma omp target data will behave less predictably and more confusingly, right now it can test, does this library have usable offloading for everything it provides (i.e. libgomp would ask the plugin to initialize offloading from the current shared library if not already done, and if successful, say it supports offloading for the particular device and map variables to that device as requested, otherwise it would just assume only host fallback is possible and not really map anything). When a target region is hit, from either within the target data region or elsewhere, it is already figured out if it has to fallback to host or not. Now, if you have fine-grained offloading, 33.2% of target regions being offloadable, the rest not, what would you actually do in target data region? It doesn't generically know what target regions will be encountered. So act as if offloading perhaps was possible? But then at each target region find out if it is really possible? IMHO people that care about performance will use target regions with care, with the offloading targets that they care about in mind, for those that don't care about that, either they will be lucky and things will work out all, or they will just end up with host fallback. Jakub