[Openmp-commits] [PATCH] D107656: [OpenMP] Use events and taskyield in target nowait task to unblock host threads

[Joachim Protze via Phabricator via Openmp-commits]

protze.joachim added a comment.

Regarding the weird nature of taskwait I refer to https://link.springer.com/chapter/10.1007%2F978-3-319-98521-3_1
Not everything in the paper is applicable for your situation. The most dangerous point I see here is, that taskyield if not used with care will effectively build a recursive call stack, so that a task that called taskyield can only proceed, if all recursively called task have finished.

  #pragma omp target nowait depend(inout:a)
  {}

As I understand the current implementation, this code translates to something like:

  #pragma omp (hidden)task depend(inout:a)
  {
    a = kernel_launch_async();
    wait_async(a);
  }

As I understand your proposal, you want to replace it by something like:

  #pragma omp (hidden)task depend(inout:a)
  {
    a = kernel_launch_async();
    while (!test_async(a))
    {
     #pragma omp taskyield
    }
  }

Think of 3 ready target nowait regions: the target task for the first target region calls taskyield and schedules the second target task. The second task also calls taskyield and schedules the third task. The first task will only continue/complete after the second and third task completed.
Depending on the number of available target tasks, you might even exceed the stack limit.

My proposed code pattern would be like:

  #pragma omp (hidden)task depend(inout:a) detach(event)
  {
    a = kernel_launch_async();
    a.register_signal(omp_fulfill_event, event); // this registers omp_fulfill_event as a callback to be called, when the asynchronous execution is finished
  } //<-- the hidden helper task is done executing. the event handling in omp_fulfill_event will take care of releasing the dependent tasks 


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D107656/new/

https://reviews.llvm.org/D107656