[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