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
