https://gcc.gnu.org/bugzilla/show_bug.cgi?id=85463
Bug ID: 85463
Summary: [nvptx] "exit" in offloaded region doesn't terminate
process
Product: gcc
Version: 8.0
Status: UNCONFIRMED
Keywords: openacc, openmp
Severity: normal
Priority: P3
Component: libgomp
Assignee: unassigned at gcc dot gnu.org
Reporter: tschwinge at gcc dot gnu.org
CC: jakub at gcc dot gnu.org
Target Milestone: ---
Target: nvptx
Consider:
#include <stdlib.h>
int main(int argc, char *argv[])
{
#pragma acc parallel
{
if (argc != 1)
exit(35);
}
return 0;
}
No matter the argc value, this program always terminates normally (exit code
zero).
Is the right solution to have GOACC_parallel etc. detect that exit etc. have
been called inside offloaded regions, and then have them terminate the process
in the appropriate way?
Can we make use of PTX trap in some way for that, to avoid adding overhead of
retval checking for (the majority of) kernel launches that terminate normally?
That is, when C exit is called, set some retval (like currently only done in
stand-alone nvptx target testing), and then call PTX trap instead of PTX exit,
which will then cause cuStreamSynchronize will indicate launch error, then we
can check the retval and do the appropritate thing?
Or, similarly, a scheme where the low-overhead PTX exit is only ever used for
normal (exit code zero) C exit(0), and everything else uses the PTX trap
variant?
This relates to PR85166 "[nvptx, libgfortran] Libgomp fortran tests using stop
in offloaded fns fail to compile", where (still now) a Fortran STOP usage will
not actually terminate the process with the appropriate error code. (The
"abort" testcases in the libgomp testsuite might appear to do the right thing
nevertheless; see the (unresolved) discussion in
<http://mid.mail-archive.com/87inlx3o1a.fsf@hertz.schwinge.homeip.net>.)
