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