From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 27321 invoked by alias); 3 Sep 2013 18:54:39 -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 27309 invoked by uid 89); 3 Sep 2013 18:54:39 -0000 Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 03 Sep 2013 18:54:39 +0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-4.3 required=5.0 tests=AWL,BAYES_00,RP_MATCHES_RCVD autolearn=ham version=3.3.2 X-HELO: mx1.redhat.com Received: from int-mx02.intmail.prod.int.phx2.redhat.com (int-mx02.intmail.prod.int.phx2.redhat.com [10.5.11.12]) by mx1.redhat.com (8.14.4/8.14.4) with ESMTP id r83Is7NG006801 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-SHA bits=256 verify=OK); Tue, 3 Sep 2013 14:54:08 -0400 Received: from tucnak.zalov.cz (vpn1-6-94.ams2.redhat.com [10.36.6.94]) by int-mx02.intmail.prod.int.phx2.redhat.com (8.13.8/8.13.8) with ESMTP id r83Is4rp022880 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-SHA bits=256 verify=NO); Tue, 3 Sep 2013 14:54:05 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.14.7/8.14.7) with ESMTP id r83Is3YI017804; Tue, 3 Sep 2013 20:54:03 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.14.7/8.14.7/Submit) id r83Is2vU017803; Tue, 3 Sep 2013 20:54:02 +0200 Date: Tue, 03 Sep 2013 18:54:00 -0000 From: Jakub Jelinek To: "Michael V. Zolotukhin" Cc: Thomas Schwinge , Ilya Verbin , Richard Biener , Uday Khedker , hubicka@ucw.cz, rth@redhat.com, kirill.yukhin@gmail.com, gcc@gcc.gnu.org Subject: Re: Questions about LTO infrastructure and pragma omp target Message-ID: <20130903185402.GK21876@tucnak.zalov.cz> Reply-To: Jakub Jelinek References: <20130823105527.GA6976@msticlxl57.ims.intel.com> <85e37f42-69fe-4bbf-bf1d-f73194e7c444@email.android.com> <20130823123638.GL1814@tucnak.redhat.com> <20130823171514.GB6976@msticlxl57.ims.intel.com> <20130823172347.GP1814@tucnak.redhat.com> <20130903135935.GC43295@msticlxl57.ims.intel.com> <20130903141837.GF21876@tucnak.zalov.cz> <20130903151801.GE43295@msticlxl57.ims.intel.com> <87ioyhhkk9.fsf@kepler.schwinge.homeip.net> <20130903182956.GH43295@msticlxl57.ims.intel.com> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <20130903182956.GH43295@msticlxl57.ims.intel.com> User-Agent: Mutt/1.5.21 (2010-09-15) X-IsSubscribed: yes X-SW-Source: 2013-09/txt/msg00026.txt.bz2 On Tue, Sep 03, 2013 at 10:29:56PM +0400, Michael V. Zolotukhin wrote: > > The idea, as we discussed it at the GNU Tools Cauldron's Acceleration > > BoF, is that the host program (for at least some acceleration devices) > > will be responsible for loading the acceleration device's code to the > > device, using some support library that is specific to each acceleration > > device > Unfortunately, I missed the Cauldron, though I'm familiar with the > general idea and now I'm trying to clarify details. > > > and for that it is useful to have the the code readily accessible > > in the host program, and thus link it in as "data". > Oh, if we just link the target binary as a data section into the host > binary, then I see no problems in that, it seems absolutely feasible > with the existing infrastructure. I just thought (seemingly it was > incorrect) that we're speaking about linking of target code with the > host code. No. The rough idea is that you emit the accelerator related subset of CUs into the (special named) LTO sections, and when linking a binary you collect all those sections from all the input object files, compile those (without -flto ideally separately), link together and finally embed into the executable into a data section. Similarly when linking a shared library, you do a target shared library and embed it in a data section of the host shared library. It is kind of fat binaries/shared libraries. Each of the accelerators would use different name of the data sections, so they could coexist. For the MIC, you'd then use COI to create the binary or shared libraries from the (ro)data section memory image. For others whatever they support. Perhaps it should be for MIC a shared library even in the binary and have some binary in a data section of the libgomp plugin, because it really should work also if the host binary doesn't have any #pragma omp target in it at all, but shared libraries do. Jakub