https://gcc.gnu.org/bugzilla/show_bug.cgi?id=80035
Alexander Monakov <amonakov at gcc dot gnu.org> changed:
What |Removed |Added
----------------------------------------------------------------------------
CC| |amonakov at gcc dot gnu.org
--- Comment #3 from Alexander Monakov <amonakov at gcc dot gnu.org> ---
Noreturn is a bit of a distraction. On the one hand, yes, on this and related
testcases it may be possible to workaround things more successfully on the gcc
side, when gcc knows that the callee _never_ returns. On the other hand, the
real issue is that ptxas gets confused around loops lacking exit edges (and
perhaps other "unusual" CFG structures). They can easily appear even in
original source and do not require presence of noreturn functions:
extern __device__ void g();
extern __device__ void h();
__device__ void f(int v)
{
if (v) goto L;
for (;;) {
g();
L:
h();
}
}
This minimized testcase leads to ptxas segfault with nvcc, so they are not
immune either. And here a compiler can't do anything to make code more
digestible to ptxas, afaict, so it really should be fixed in ptxas.
FWIW, it seems to work with -arch sm_50, so it might be fixed for newer
architectures (I suspect CUDA driver and toolkit carry separate code generators
for sm_3x and sm_5x).
Can you share if you (Mentor Embeddded) reported such issues to NVIDIA in the
past?
