r268178 - in /trunk/libgomp: ChangeLog plugin/p...

[vries@gcc.gnu.org]

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  <tdevries@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