[PATCH,NVPTX] Fix PR83920

[Tom de Vries]

On 01/18/2018 02:27 PM, Tom de Vries wrote:
> On 01/18/2018 12:40 AM, Cesar Philippidis wrote:
>> In PR83920, I encountered a nvptx bug where live predicate variables
>> were clobbered before their value was broadcasted. 
> 
> Hi,
> 
> I've managed to reproduce the problem based on the description in the PR.

> I think the way to address it is using a tmp .pred reg like so:
> ...
> {
>  Â  .reg .u32 %x;
>  Â  mov.u32 %x,%tid.x;
>  Â  setp.ne.u32 %rnotvzero,%x,0;
> }
> 
> {
>  Â  .reg .pred %rcond2;
>  Â  setp.eq.u32 %rcond2, 1, 0; // workaround
> 
>  Â  @%rnotvzero bra Lskip;
>  Â  ...
>  Â  setp.<op>.<type> %rcond,op1,op2; // could be here, could be earlier
>  Â  mov.b1 %rcond2, %rcond; // used pseudo opcode mov.b1 for convenience
>  Â Lskip:
>  Â  selp.u32 %rcondu32,1,0,%rcond2;
>  Â  shfl.idx.b32 %rcondu32,%rcondu32,0,31;
>  Â  setp.ne.u32 %rcond,%rcondu32,0;
> }
> ...
> 

Hi,

this is the fix that I plan to commit (similar to the scheme listed 
above, but modified to keep the selp.u32 using rcond, which is easier in 
code generation).

Build and reg-tested on x86_64 with nvptx accelerator.

Richard, this is an 8 regression for the nvptx target. OK for stage 4 or 
defer to stage1?

Thanks,
- Tom
-------------- next part --------------
A non-text attachment was scrubbed...
Name: 0001-nvptx-Fix-bug-in-jit-bug-workaround.patch
Type: text/x-patch
Size: 3620 bytes
Desc: not available
URL: <http://gcc.gnu.org/pipermail/gcc-patches/attachments/20180119/a5aecec6/attachment.bin>