https://gcc.gnu.org/bugzilla/show_bug.cgi?id=88946
Bug ID: 88946
Summary: [nvptx, openacc, libgomp] cuMemAlloc error for two
empty asynchronous parallels
Product: gcc
Version: 9.0
Status: UNCONFIRMED
Severity: normal
Priority: P3
Component: target
Assignee: unassigned at gcc dot gnu.org
Reporter: vries at gcc dot gnu.org
Target Milestone: ---
Consider test-case:
...
/* { dg-do run } */
int
main (void)
{
#pragma acc parallel async
;
#pragma acc parallel async
;
#pragma acc wait
return 0;
}
...
This fails in execution:
...
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/test.c \
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O0 execution test
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/test.c \
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O2 execution test
...
with:
...
libgomp: cuMemAlloc error: invalid argument
Segmentation fault (core dumped)
...
because we're trying to allocate zero bytes:
...
Thread 1 "test.exe" hit Breakpoint 1, GOMP_PLUGIN_error (
msg=msg@entry=0x7ffff6ea9987 "cuMemAlloc error: %s")
at libgomp/libgomp-plugin.c:64
64 {
(gdb) up
#1 0x00007ffff6ea647c in cuda_map_create (size=0)
at /home/vries/oacc/trunk/source-gcc/libgomp/plugin/plugin-nvptx.c:231
231 CUDA_CALL_ERET (NULL, cuMemAlloc, &map->d, size);
...
This causes cuda_map_create to return NULL, after which map in map_push map is
NULL, so we segfault at:
...
s->map = map;
s->map->active = true;
...
If we remove the async clauses from the test-case, it passes. Still we call
map_push with size == 0, but that doesn't result in a cuda_map_create with size
== 0.
