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?