[Bug target/81069] nvptx offloading: "-O1" execution test of "libgomp.oacc-fortran/nested-function-1.f90" timeout/FAIL

["vries at gcc dot gnu.org" <gcc-bugzilla at gcc dot gnu dot org>]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=81069

--- Comment #3 from Tom de Vries <vries at gcc dot gnu.org> ---
The immediate reason for the failure is a bug in the ptx JIT compiler.

This piece of ptx code (containing two functionally equivalent but differently
layed out variants, DETOUR=0 and DETOUR=1) is a minimal example, and is
translated correctly for DETOUR=0 and incorrectly for DETOUR=1:
...
.version 3.1
.target sm_30
.address_size 64

.entry test_function (.param .u64 %in_ar0);

.entry test_function (.param .u64 %in_ar0)
{
        // Init %ra to 0 in all lanes.
        .reg .u32 %ra;
        mov.u32 %ra,0;

        // Branch around if not lane 0.
        {
                .reg .u32 %l;
                mov.u32 %l,%laneid;
                .reg .pred %lane0;
                setp.eq.u32 %lane0,%l,0;
                @ ! %lane0 bra $L20;
        }
        bra $L3;

#if DETOUR == 0
$L3:
        // Set %ra to 1 in lane zero
        mov.u32 %ra,1;
$L20:   
        bra $L8;
#endif

$L8:
        // Broadcast %ra from lane 0 to all lanes
        shfl.idx.b32 %ra,%ra,0,31;

        // Branch to trap if %ra not 1 in all lanes
        {
                .reg .pred %ra_is_one;
                setp.eq.u32 %ra_is_one,%ra,1;
                @ ! %ra_is_one bra $L6;
        }

        ret;

$L6:
        trap;

#if DETOUR == 1
$L3:
        // Set %ra to 1 in lane zero
        mov.u32 %ra,1;
$L20:   
        bra $L8;
#endif
}
...

The problem for the DETOUR=1 case is that while the divergent branch '@ !
%lane0 bra $L20' is post-dominated by $L20, the synchronization point is
inserted much later, after the shfl. This means the shfl is executed in
divergent state, which invokes undefined behaviour, causing an undefined result
%ra in one of the threads in the warp, which causes the example to call the
trap.