This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [gomp4.5] depend nowait support for target
- From: Ilya Verbin <iverbin at gmail dot com>
- To: Jakub Jelinek <jakub at redhat dot com>
- Cc: Aldy Hernandez <aldyh at redhat dot com>, gcc-patches at gcc dot gnu dot org, Kirill Yukhin <kirill dot yukhin at gmail dot com>, Thomas Schwinge <thomas at codesourcery dot com>, Alexander Monakov <amonakov at ispras dot ru>, Martin Jambor <mjambor at suse dot cz>
- Date: Thu, 12 Nov 2015 23:51:33 +0300
- Subject: Re: [gomp4.5] depend nowait support for target
- Authentication-results: sourceware.org; auth=none
- References: <20150908092014 dot GA1847 at tucnak dot redhat dot com> <20151002192801 dot GA24765 at msticlxl57 dot ims dot intel dot com> <20151015140156 dot GE478 at tucnak dot redhat dot com> <20151019194754 dot GB1855 at msticlxl57 dot ims dot intel dot com> <20151111165222 dot GL5675 at tucnak dot redhat dot com> <20151112174509 dot GG5675 at tucnak dot redhat dot com>
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