From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 57428 invoked by alias); 17 Jul 2015 16:31: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 57415 invoked by uid 89); 17 Jul 2015 16:31:53 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL,BAYES_00,FREEMAIL_FROM,RCVD_IN_DNSWL_LOW,SPF_PASS autolearn=ham version=3.3.2 X-HELO: mail-wi0-f174.google.com Received: from mail-wi0-f174.google.com (HELO mail-wi0-f174.google.com) (209.85.212.174) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Fri, 17 Jul 2015 16:31:48 +0000 Received: by wibud3 with SMTP id ud3so47205221wib.0 for ; Fri, 17 Jul 2015 09:31:45 -0700 (PDT) X-Received: by 10.194.202.197 with SMTP id kk5mr30540975wjc.64.1437150705707; Fri, 17 Jul 2015 09:31:45 -0700 (PDT) Received: from msticlxl57.ims.intel.com ([192.55.54.40]) by smtp.gmail.com with ESMTPSA id h9sm19064294wjq.10.2015.07.17.09.31.41 (version=TLSv1 cipher=ECDHE-RSA-AES128-SHA bits=128/128); Fri, 17 Jul 2015 09:31:44 -0700 (PDT) Date: Fri, 17 Jul 2015 16:47:00 -0000 From: Ilya Verbin To: Jakub Jelinek , Thomas Schwinge Cc: gcc-patches@gcc.gnu.org, Kirill Yukhin Subject: Re: [gomp4.1] Initial support for some OpenMP 4.1 construct parsing Message-ID: <20150717163136.GB15252@msticlxl57.ims.intel.com> 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> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <20150625201058.GK10247@tucnak.redhat.com> User-Agent: Mutt/1.5.21 (2010-09-15) X-IsSubscribed: yes X-SW-Source: 2015-07/txt/msg01555.txt.bz2 On Thu, Jun 25, 2015 at 22:10:58 +0200, Jakub Jelinek wrote: > 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. 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) Currently I'm trying to figure out what liboffloadmic can do. BTW, do you plan to remove GOMP_MAP_POINTER mappings from array sections? The enter/exit patch for libgomp depends on this change. -- Ilya