https://gcc.gnu.org/bugzilla/show_bug.cgi?id=88941
Bug ID: 88941
Summary: [nvptx, openacc, libgomp] Assertion `!s->map->active'
failed for empty asynchronous parallel
Product: gcc
Version: 9.0
Status: UNCONFIRMED
Severity: normal
Priority: P3
Component: libgomp
Assignee: unassigned at gcc dot gnu.org
Reporter: vries at gcc dot gnu.org
CC: jakub at gcc dot gnu.org
Target Milestone: ---
Consider test-case:
...
$ cat libgomp/testsuite/libgomp.oacc-c-c++-common/empty-async.c
/* { dg-do run } */
#include <stdlib.h>
int
main (void)
{
#pragma acc parallel async
;
/* no #pragma acc wait */
return 0;
}
...
It fails in execution:
...
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/empty-async.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O0 execution test
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/empty-async.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O2 execution test
...
because:
...
empty-async.exe: libgomp/plugin/plugin-nvptx.c:271: map_fini: Assertion
`!s->map->active' failed.
...
The following scenario happens:
- map_push allocates a struct cuda_map for the kernel launch parameters
- the kernel is launched
- an event is registered to call the corresponding map_pop after the kernel
completes
- main reaches return, and atexit processing starts
- during atexit processing, fini_streams_for_device is called, which calls
map_fini
- map_fini asserts, because map_pop hasn't been called.
So either:
- the kernel hasn't finished yet, or
- the kernel has finished, but the event has not been processed
