From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 91023 invoked by alias); 21 Oct 2015 09:06:55 -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 89203 invoked by uid 89); 21 Oct 2015 09:06:53 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.8 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; Wed, 21 Oct 2015 09:06:52 +0000 Received: from int-mx13.intmail.prod.int.phx2.redhat.com (int-mx13.intmail.prod.int.phx2.redhat.com [10.5.11.26]) by mx1.redhat.com (Postfix) with ESMTPS id 2069EC0C18AA; Wed, 21 Oct 2015 09:06:51 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-53.ams2.redhat.com [10.36.116.53]) by int-mx13.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t9L96njB029369 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Wed, 21 Oct 2015 05:06:50 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.15.2/8.15.2) with ESMTP id t9L96lW0019080; Wed, 21 Oct 2015 11:06:48 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.15.2/8.15.2/Submit) id t9L96kHp019079; Wed, 21 Oct 2015 11:06:46 +0200 Date: Wed, 21 Oct 2015 09:07:00 -0000 From: Jakub Jelinek To: Bernd Schmidt Cc: Alexander Monakov , gcc-patches@gcc.gnu.org, Dmitry Melnik Subject: Re: [gomp4 07/14] libgomp nvptx plugin: launch target functions via gomp_nvptx_main Message-ID: <20151021090646.GL478@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <1445366076-16082-1-git-send-email-amonakov@ispras.ru> <1445366076-16082-8-git-send-email-amonakov@ispras.ru> <5626ACBC.2090409@redhat.com> <5626AFCA.5090502@redhat.com> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <5626AFCA.5090502@redhat.com> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2015-10/txt/msg02035.txt.bz2 On Tue, Oct 20, 2015 at 11:19:06PM +0200, Bernd Schmidt wrote: > On 10/20/2015 11:13 PM, Alexander Monakov wrote: > >On Tue, 20 Oct 2015, Bernd Schmidt wrote: > > > >>On 10/20/2015 08:34 PM, Alexander Monakov wrote: > >>>2. Make gomp_nvptx_main a device (.func) function. To have that work, we'd > >>>need to additionally emit a "trampoline" of sorts in the NVPTX backend. For > >>>each OpenMP target entrypoint foo$_omp_fn$0, we'd have to additionally emit > >>> > >>>__global__ void foo$_omp_fn$0$entry(void *args) > >>>{ > >>> gomp_nvptx_main(foo$_omp_fn$0, args); > >>>} > >> > >>Wouldn't it be simpler to generate a .kernel for every target region function > >>(as OpenACC does)? That could be a small stub in each case which just calls > >>gomp_nvptx_main with the right function pointer. We already have the machinery > >>to look up the right kernel corresponding to a host address and invoke it, so > >>I think we should just reuse that functionality. > > > >As I see we are describing the same thing in different words. > > > >In what you describe, and in my quoted paragraph, both gomp_nvptx_main and the > >function originally outlined for a target region are device-only (.func) > >functions. The .kernel function that the plugin looks up and launches is a > >small piece of code that calls gomp_nvptx_main, passing it a pointer to the > >target region function. > > > >Unless I didn't fully catch what you say? Like I said in the email, I do like > >this approach more. > > Could be that we're talking about the same thing. I think I was confused by > a reference to .func vs .kernel and sm_30 vs sm_35 in patch 2/14. So let's > go for this approach. But you'd better then rename the .kernel stub to the name that libgomp looks up and rename the actual outlined body to some other name (add some suffix). Or maybe another possibility is to use the normal outlined target body function, and just inject some gomp_nvptx_start and gomp_nvptx_end functions to the beginning and end of "omp target entrypoint" function. Jakub