[nvptx, PR85204] Fix neutering of bb with only cond jump
Tom de Vries
Tom_deVries@mentor.com
Thu Apr 5 08:41:00 GMT 2018
Hi,
When compiling the test-case in the patch, the following ptx code is
generated:
...
$L4:
@ %r91 bra.uni $L24;
selp.u32 %r95,1,0,%r80;
st.shared.u32 [__worker_bcast],%r95;
$L25:
$L24:
@ %r92 bra $L25;
...
There's an eternal loop starting at the last insn, and unsurprisingly
the test-case hangs.
The last insn is a vector neutering branch, which should have been
inserted after the worker neutering branch (the first insn).
In other words, we want:
...
$L4:
@ %r91 bra.uni $L24;
+ @ %r92 bra $L25;
selp.u32 %r95,1,0,%r80;
st.shared.u32 [__worker_bcast],%r95;
$L25:
$L24:
- @ %r92 bra $L25;
...
This minimal stage4 patch fixes this problem. [ I filed a PR85223
"[nvptx] nvptx_single needs rewrite" for a stage1 rewrite of nvptx_single. ]
Build x86_64 with nvptx accelerator, and tested libgomp.
Committed to stage4 trunk.
Thanks,
- Tom
-------------- next part --------------
A non-text attachment was scrubbed...
Name: 0001-nvptx-Fix-neutering-of-bb-with-only-cond-jump.patch
Type: text/x-patch
Size: 2438 bytes
Desc: not available
URL: <http://gcc.gnu.org/pipermail/gcc-patches/attachments/20180405/51383d12/attachment.bin>
More information about the Gcc-patches
mailing list