Hi! On Tue, 20 May 2014 13:21:06 +0200, Tom de Vries <tom_devr...@mentor.com> wrote: > On 17/05/14 17:53, Tom de Vries wrote: > > On 16/05/14 22:28, Thomas Schwinge wrote: > >> > It may make sense to have an abort function available in PTX land.
... which we've now had for a long time. However, I recently noticed that it doesn't actually behave as expected: <https://github.com/MentorEmbedded/nvptx-newlib/issues/5> 'The nvptx "abort" implementation doesn't actually abort execution #5'. ;-| That is, for the last year or so, the nvptx "abort" implementation silently "exit"ed instead of "trap"ping. However code such as: > #pragma acc parallel > { > abort (); > } ... still caused process termination, and while I have not analyzed that in detail, the reason seems to be along the lines that GCC propagates the "abort"'s "noreturn" attribute to the OMP child function ("main._omp_fn.0"), and then places a "__builtin_unreachable" after the "GOACC_parallel_keyed" -- which is the cause of the process termination. (Quite surprising to me, that is done even for "-O0"!) See abort-1.c.019t.fixup_cfg1: [...] Introduced new external node (__builtin_unreachable/8). [...] GOACC_parallel_keyed (-1, main._omp_fn.0, 0, 0B, 0B, 0B, 0); __builtin_unreachable (); [...] Now, I could do something like: #pragma acc parallel { - abort (); + int cond = 1; + asm volatile ("" : : : "memory"); + if (cond) + abort (); } ... to prevent that from happening -- but isn't there a generic way (command-line flag?) to disable the analysis that causes the "noreturn" attribute propagation/handling? Grüße Thomas
