diff options
| author | Tom de Vries <tom@codesourcery.com> | 2018-01-19 16:29:41 +0000 |
|---|---|---|
| committer | Tom de Vries <vries@gcc.gnu.org> | 2018-01-19 16:29:41 +0000 |
| commit | 8c8e9a6bb656fc50b9a166ba452d252fe05a2c38 (patch) | |
| tree | 07ec11d9f2a5ce4fdfbaeff61f7818ad75d626bd /gcc | |
| parent | 6c7c47081ac481c28f31c30e8845d0337556e8a2 (diff) | |
| download | gcc-8c8e9a6bb656fc50b9a166ba452d252fe05a2c38.zip gcc-8c8e9a6bb656fc50b9a166ba452d252fe05a2c38.tar.gz gcc-8c8e9a6bb656fc50b9a166ba452d252fe05a2c38.tar.bz2 | |
[nvptx] Fix bug in jit bug workaround
2018-01-19 Tom de Vries <tom@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
PR target/83920
* config/nvptx/nvptx.c (nvptx_single): Fix jit workaround.
* testsuite/libgomp.oacc-c-c++-common/pr83920.c: New test.
* testsuite/libgomp.oacc-fortran/pr83920.f90: New test.
Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
From-SVN: r256894
Diffstat (limited to 'gcc')
| -rw-r--r-- | gcc/ChangeLog | 6 | ||||
| -rw-r--r-- | gcc/config/nvptx/nvptx.c | 28 |
2 files changed, 32 insertions, 2 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 7aa2919..162ae1b 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,9 @@ +2018-01-19 Tom de Vries <tom@codesourcery.com> + Cesar Philippidis <cesar@codesourcery.com> + + PR target/83920 + * config/nvptx/nvptx.c (nvptx_single): Fix jit workaround. + 2018-01-19 Cesar Philippidis <cesar@codesourcery.com> PR target/83790 diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index 4cb87c8..f5bb438 100644 --- a/gcc/config/nvptx/nvptx.c +++ b/gcc/config/nvptx/nvptx.c @@ -4102,9 +4102,33 @@ nvptx_single (unsigned mask, basic_block from, basic_block to) There is nothing in the PTX spec to suggest that this is wrong, or to explain why the extra initialization is needed. So, we classify - it as a JIT bug, and the extra initialization as workaround. */ - emit_insn_before (gen_movbi (pvar, const0_rtx), + it as a JIT bug, and the extra initialization as workaround: + + { + .reg .u32 %x; + mov.u32 %x,%tid.x; + setp.ne.u32 %rnotvzero,%x,0; + } + + +.reg .pred %rcond2; + +setp.eq.u32 %rcond2, 1, 0; + + @%rnotvzero bra Lskip; + setp.<op>.<type> %rcond,op1,op2; + +mov.pred %rcond2, %rcond; + Lskip: + +mov.pred %rcond, %rcond2; + selp.u32 %rcondu32,1,0,%rcond; + shfl.idx.b32 %rcondu32,%rcondu32,0,31; + setp.ne.u32 %rcond,%rcondu32,0; + */ + rtx_insn *label = PREV_INSN (tail); + gcc_assert (label && LABEL_P (label)); + rtx tmp = gen_reg_rtx (BImode); + emit_insn_before (gen_movbi (tmp, const0_rtx), bb_first_real_insn (from)); + emit_insn_before (gen_rtx_SET (tmp, pvar), label); + emit_insn_before (gen_rtx_SET (pvar, tmp), tail); #endif emit_insn_before (nvptx_gen_vcast (pvar), tail); } |