[gomp4.5] depend nowait support for target

[Ilya Verbin]

On Thu, Nov 12, 2015 at 18:45:09 +0100, Jakub Jelinek wrote:
> But the testcase I wrote (target-33.c) hangs, the problem is in the
>   #pragma omp target nowait map (tofrom: a, b) depend(out: d[3])
>   {
>     #pragma omp atomic update
>     a = a + 9;
>     b -= 8;
>   }
>   #pragma omp target nowait map (tofrom: a, c) depend(out: d[4])
>   {
>     #pragma omp atomic update
>     a = a + 4;
>     c >>= 1;
>   }
>   #pragma omp task if (0) depend (in: d[3], d[4])
>   if (a != 50 || b != 4 || c != 20)
>     abort ();
> part, where (I should change that for the case of no dependencies
> eventually) the task with map_vars+async_run is queued in both cases,
> then we reach GOMP_task, which calls gomp_task_maybe_wait_for_dependencies
> which spawns the first half task (map_vars+async_run), and then
> the second half task (map_vars+async_run), but that one gets stuck somewhere
> in liboffloadmic, then some other thread (from liboffloadmic) calls
> GOMP_PLUGIN_target_task_completion and enqueues the second half of the first
> target task (unmap_vars), but as the only normal thread in the main program
> is stuck in liboffloadmic (during gomp_map_vars, trying to allocate
> target memory in the plugin), there is no thread to schedule the second half
> of first target task.  So, if liboffloadmic is stuck waiting for unmap_vars,
> it is a deadlock.  Can you please try to debug this?

I'm unable to reproduce the hang (have tried various values of OMP_NUM_THREADS).
The testcase just aborts at (a != 50 || b != 4 || c != 20), because
a == 37, b == 12, c == 40.

BTW, don't know is this a bug or not:
Conditional jump or move depends on uninitialised value(s)
   at 0x4C2083D: priority_queue_insert (priority_queue.h:347)
   by 0x4C24DF9: GOMP_PLUGIN_target_task_completion (task.c:678)

  -- Ilya