[Bug target/83920] [nvptx] bad predicate reset

vries at gcc dot gnu.org gcc-bugzilla@gcc.gnu.org
Thu Jan 18 10:41:00 GMT 2018


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

Tom de Vries <vries at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
             Status|UNCONFIRMED                 |NEW
   Last reconfirmed|                            |2018-01-18
     Ever confirmed|0                           |1

--- Comment #2 from Tom de Vries <vries at gcc dot gnu.org> ---
Confirmed. The workaround is intended to work on code like this:
...
                {                                                               
                    .reg .u32 %x;                                               
                    mov.u32 %x,%tid.x;                                          
                    setp.ne.u32 %rnotvzero,%x,0;                                
                 }                                                              

                 @%rnotvzero bra Lskip;                                         
                 setp.<op>.<type> %rcond,op1,op2;                               
                 Lskip:                                                         
                 selp.u32 %rcondu32,1,0,%rcond;                                 
                 shfl.idx.b32 %rcondu32,%rcondu32,0,31;                         
                 setp.ne.u32 %rcond,%rcondu32,0;                                
...
and adds 'setp.eq.u32 %rcond, 1, 0;' before "bra Lskip". However, if the branch
condition is not calculated in the basic block containing the conditional jump,
then the workaround overwrites the branch condition.

In the case of comment 1, it's the most extreme case: we're neutering an empty
block:
...
        @%r341  bra     $L33;
$L33:
...
so the condition must be defined elsewhere.


More information about the Gcc-bugs mailing list