From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 85048 invoked by alias); 24 Aug 2015 12:27: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 85027 invoked by uid 89); 24 Aug 2015 12:27:03 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.1 required=5.0 tests=AWL,BAYES_00,KAM_LAZY_DOMAIN_SECURITY,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; Mon, 24 Aug 2015 12:27:02 +0000 Received: from int-mx14.intmail.prod.int.phx2.redhat.com (int-mx14.intmail.prod.int.phx2.redhat.com [10.5.11.27]) by mx1.redhat.com (Postfix) with ESMTPS id 7DE898E4E6; Mon, 24 Aug 2015 12:27:01 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-26.ams2.redhat.com [10.36.116.26]) by int-mx14.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t7OCQxL0021392 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Mon, 24 Aug 2015 08:27:01 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.14.9/8.14.9) with ESMTP id t7OCQwGu014750; Mon, 24 Aug 2015 14:26:58 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.14.9/8.14.9/Submit) id t7OCQv1I014749; Mon, 24 Aug 2015 14:26:57 +0200 Date: Mon, 24 Aug 2015 12:38:00 -0000 From: Jakub Jelinek To: Ilya Verbin Cc: gcc-patches@gcc.gnu.org, Kirill Yukhin Subject: Re: [gomp4.1] Initial support for some OpenMP 4.1 construct parsing Message-ID: <20150824122657.GJ9425@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <20150429111406.GE1751@tucnak.redhat.com> <874mnzrw1z.fsf@schwinge.name> <20150429120644.GG1751@tucnak.redhat.com> <20150609183608.GA47936@msticlxl57.ims.intel.com> <20150609202426.GG10247@tucnak.redhat.com> <20150625194529.GB33078@msticlxl57.ims.intel.com> <20150625201058.GK10247@tucnak.redhat.com> <20150717163136.GB15252@msticlxl57.ims.intel.com> <20150717164306.GT1780@tucnak.redhat.com> <20150720185342.GA350@msticlxl57.ims.intel.com> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <20150720185342.GA350@msticlxl57.ims.intel.com> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2015-08/txt/msg01406.txt.bz2 On Mon, Jul 20, 2015 at 09:53:42PM +0300, Ilya Verbin wrote: > On Fri, Jul 17, 2015 at 18:43:06 +0200, Jakub Jelinek wrote: > > On Fri, Jul 17, 2015 at 07:31:36PM +0300, Ilya Verbin wrote: > > > One big question is who will maintain the list of scheduled job, its > > > dependencies, etc. - libgomp or each target plugin? > > > > > > > > > OpenACC has async queues: > > > #pragma acc parallel async(2) wait(1) > > > > > > But it's not possible to have 2 waits like: > > > #pragma acc parallel async(3) wait(1) wait(2) > > > > > > (GOMP_OFFLOAD_openacc_async_wait_async has only one argument with the number of > > > queue to wait) > > > > > > Thomas, please correct me if I'm wrong. > > > > > > In this regard, OpenMP is more complicated, since it allows e.g.: > > > #pragma omp target nowait depend(in: a, b) depend(out: c, d) > > > > If it is each plugin, then supposedly it should use (if possible) some > > common libgomp routine to maintain the queues, duplicating the dependency > > graph handling code in each plugins might be too ugly. > > > > > Currently I'm trying to figure out what liboffloadmic can do. > > Latest liboffloadmic (I'm preparing an update for trunk) can take some pointer > *ptr* as argument of __offload_offload, which is used for execution and data > transfer. When given job is finished, it will call some callback in libgomp on > host, passing *ptr* back to it, thus libgomp can distinguish which job has > been finished. BTW, which word to use here to avoid confusion? (task? job?) > > I'm going to prototype something in libgomp using this interface. Based on the recent largish thread on omp-lang, it seems that the intent is that depend clause on target* constructs is for an implicit "target task" on the host side, so sorry for understanding it as dependencies on the implicit task on the device instead before. As the "target task" is supposed to be merged (i.e. sharing ICVs and variables), for constructs without nowait and depend clauses we can do what we do right now, perhaps later with some optimizations trying to invoke other tasks if waiting for the hardware. For target* constructs with nowait and/or depend clauses, we'll need to pass in the depend info (which is a pointer to array of addresses, with first "address" being total count and second "address" being number of out/inout dependencies at the beginning, and whether nowait is present or not (a bit somewhere), then create some struct gomp_task for it supposedly with a new GOMP_TASK_TARGET kind, and handle the dependency waiting etc. like for other tasks, except that it will not have a real body, but saved arguments from the GOMP_target* call (at least for not immediately satisfied dependency my understanding from the mail thread is that the data transfer operations are supposed to start only when the dependency is met). Jakub