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
