From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 44311 invoked by alias); 25 Jun 2015 20:11:06 -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 44293 invoked by uid 89); 25 Jun 2015 20:11:05 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.6 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; Thu, 25 Jun 2015 20:11:04 +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 A9A253A8107; Thu, 25 Jun 2015 20:11:03 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-82.ams2.redhat.com [10.36.116.82]) by int-mx13.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t5PKB2th030589 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Thu, 25 Jun 2015 16:11:03 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.14.9/8.14.9) with ESMTP id t5PKB0B1001776; Thu, 25 Jun 2015 22:11:00 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.14.9/8.14.9/Submit) id t5PKAwrR001775; Thu, 25 Jun 2015 22:10:58 +0200 Date: Thu, 25 Jun 2015 20:31: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: <20150625201058.GK10247@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> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <20150625194529.GB33078@msticlxl57.ims.intel.com> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2015-06/txt/msg01869.txt.bz2 On Thu, Jun 25, 2015 at 10:45:29PM +0300, Ilya Verbin wrote: > So, as I understood, three tasks will be generated almost simultaneously in > foo1: one on host and two on target. > Target task 1 will be executed immediately. > Host task will wait for task 1 to be completed on target. > (Or it is not possible to mix "omp target" and "omp task" dependencies?) > And task 2 will wait on target for task 1. My understanding is that you don't create any extra tasks, but rather you pointer translate the host address from the start of the variable (or array section; thus the depend clause argument) into target address, and check if it can be offloaded right away (no need to wait for dependencies). If yes, you just offload it, with nowait without waiting in the caller till it finishes. If not, you arrange that when some other offloaded job finishes that provides the dependency, your scheduled job is executed. So, the task on the target is the implicit one, what executes the body of the target region. In tasking (task.c) dependencies are only honored for sibling tasks, whether the different target implicit tasks are sibling is questionable and supposedly should be clarified, but I can't imagine they aren't meant to. So, you don't really need to care about the task.c dependencies, target.c could have its own ones if it is easier to write it that way. Supposedly for nowait you want to spawn or queue the job and return right away, and for queued job stick it into some data structure (supposedly inside of libgomp on the host) that when the library is (asynchronously) notified that some offloaded job finished you check the data structures and spawn something different. Or have the data structures on the offloaded device instead? In any case, I'd look what the Mentor folks are doing for OpenACC async offloading, what libmicoffload allows you to do and figure out something from that. It is not entirely clear to me if the host tasks and threads that execute the async #pragma omp target {,enter data,exit data,update} depend(...) {,nowait} constructs are relevant to it at all, but I bet one task/thread could supposedly invoke something and a different one invoke something that depends on it, or waits for its completion (exit data?). And, wonder how is this really meant to work with host fallback. Bet that really needs to be clarified on omp-lang. Jakub