https://gcc.gnu.org/bugzilla/show_bug.cgi?id=88946
Tom de Vries <vries at gcc dot gnu.org> changed:
What |Removed |Added
----------------------------------------------------------------------------
Assignee|unassigned at gcc dot gnu.org |vries at gcc dot gnu.org
Target Milestone|--- |9.0
--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
https://gcc.gnu.org/ml/gcc-cvs/2019-01/msg00697.html:
Author: vries
Date: Wed Jan 23 08:16:56 2019
New Revision: 268178
URL: https://gcc.gnu.org/viewcvs?rev=268178&root=gcc&view=rev
Log:
[nvptx, libgomp] Fix cuMemAlloc with size zero
Consider test-case:
...
int
main (void)
{
#pragma acc parallel async
;
#pragma acc parallel async
;
#pragma acc wait
return 0;
}
...
This fails with:
...
libgomp: cuMemAlloc error: invalid argument
Segmentation fault (core dumped)
...
The cuMemAlloc error is due to the fact that we're try to allocate 0 bytes.
Fix this by preventing calling map_push with size zero argument in nvptx_exec.
This also has the consequence that for the abort-1.c test-case, we end up
calling cuMemFree during map_fini for the struct cuda_map allocated in
map_init, which fails because an abort happened. Fix this by calling
cuMemFree with CUDA_CALL_NOCHECK in cuda_map_destroy.
2019-01-23 Tom de Vries <tdevr...@suse.de>
PR target/PR88946
* plugin/plugin-nvptx.c (cuda_map_destroy): Use CUDA_CALL_NOCHECK for
cuMemFree.
(nvptx_exec): Don't call map_push if mapnum == 0.
* testsuite/libgomp.oacc-c-c++-common/pr88946.c: New test.
Added:
trunk/libgomp/testsuite/libgomp.oacc-c-c++-common/pr88946.c
Modified:
trunk/libgomp/ChangeLog
trunk/libgomp/plugin/plugin-nvptx.c
