From 8c8e9a6bb656fc50b9a166ba452d252fe05a2c38 Mon Sep 17 00:00:00 2001 From: Tom de Vries Date: Fri, 19 Jan 2018 16:29:41 +0000 Subject: [nvptx] Fix bug in jit bug workaround 2018-01-19 Tom de Vries Cesar Philippidis 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 From-SVN: r256894 --- gcc/ChangeLog | 6 ++++++ gcc/config/nvptx/nvptx.c | 28 ++++++++++++++++++++++++++-- 2 files changed, 32 insertions(+), 2 deletions(-) (limited to 'gcc') 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 + Cesar Philippidis + + PR target/83920 + * config/nvptx/nvptx.c (nvptx_single): Fix jit workaround. + 2018-01-19 Cesar Philippidis 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.. %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); } -- cgit v1.1