From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 16641 invoked by alias); 23 Aug 2013 11:19:42 -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 16632 invoked by uid 89); 23 Aug 2013 11:19:41 -0000 X-Spam-SWARE-Status: No, score=-7.6 required=5.0 tests=AWL,BAYES_00,RCVD_IN_HOSTKARMA_W,RCVD_IN_HOSTKARMA_WL,RP_MATCHES_RCVD,SPF_HELO_PASS,SPF_PASS autolearn=ham version=3.3.2 Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.84/v0.84-167-ge50287c) with ESMTP; Fri, 23 Aug 2013 11:19:41 +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 r7NBHZxC021315 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-SHA bits=256 verify=OK); Fri, 23 Aug 2013 07:17:35 -0400 Received: from zalov.cz (vpn1-4-247.ams2.redhat.com [10.36.4.247]) by int-mx10.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id r7NBHXsC023879 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-SHA bits=256 verify=NO); Fri, 23 Aug 2013 07:17:35 -0400 Received: from zalov.cz (localhost [127.0.0.1]) by zalov.cz (8.14.5/8.14.5) with ESMTP id r7NBHX23005194; Fri, 23 Aug 2013 13:17:33 +0200 Received: (from jakub@localhost) by zalov.cz (8.14.5/8.14.5/Submit) id r7NBHVJI005193; Fri, 23 Aug 2013 13:17:31 +0200 Date: Fri, 23 Aug 2013 14:38:00 -0000 From: Jakub Jelinek To: Ilya Verbin Cc: Richard Biener , Uday Khedker , hubicka@ucw.cz, rth@redhat.com, kirill.yukhin@gmail.com, "Michael V. Zolotukhin" , gcc@gcc.gnu.org Subject: Re: Questions about LTO infrastructure and pragma omp target Message-ID: <20130823111731.GK1814@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <20130815133639.GA42355@msticlxl57.ims.intel.com> <20130815134437.GB1814@tucnak.redhat.com> <4df844f6-f385-4e63-9413-8ea341992b77@email.android.com> <20130823105527.GA6976@msticlxl57.ims.intel.com> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <20130823105527.GA6976@msticlxl57.ims.intel.com> User-Agent: Mutt/1.5.21 (2010-09-15) X-SW-Source: 2013-08/txt/msg00273.txt.bz2 On Fri, Aug 23, 2013 at 02:55:27PM +0400, Ilya Verbin wrote: > I'm trying to implement the approach with modified lto-wrapper. > Suppose we have a bytecode of the routine foo, streamed during ompexp pass into some section, say .gnu.omptarget_foo. > In function lto.c:do_whole_program_analysis() an extra partition should be created, that will contain bytecode from .gnu.omptarget_foo, right? > As far as I understood, in addition to the bytecode of foo, we should also stream extra symtab_nodes, and read them somewhere in lto-cgraph.c:input_symtab(). > This means we should maintain 2 symtabs inside WPA stage - original for host and new for target? I don't think we should stream into more than one target section. There should be just .gnu.target_lto section (or whatever other suitable name) and should stream into it: 1) all functions and variables with "omp declare target" attribute 2) the outlined bodies of #pragma omp target turned into *.ompfn functions 3) all the types, symtab etc. needed for that If compiling with -flto, you'll also get everything from the CU streamed into the normal LTO section, otherwise you'll get assembly for the host variables/functions/etc. Then the question is what the plugin should perform with these sections, whether it will compile each input .gnu.target_lto section hunk separately (as in non-LTO mode), or with -flto also LTO them together. Jakub