diff options
| author | Tom de Vries <tdevries@suse.de> | 2018-12-19 10:17:01 +0000 |
|---|---|---|
| committer | Tom de Vries <vries@gcc.gnu.org> | 2018-12-19 10:17:01 +0000 |
| commit | 1dcf26882b1a8c4f388ed6a618d6e4b1e1330057 (patch) | |
| tree | 819d7d0613f2f27c82e00fcf8178b5235c86cbef | |
| parent | 22aa0613532c867b252ebfa681a0b74231a82efa (diff) | |
| download | gcc-1dcf26882b1a8c4f388ed6a618d6e4b1e1330057.zip gcc-1dcf26882b1a8c4f388ed6a618d6e4b1e1330057.tar.gz gcc-1dcf26882b1a8c4f388ed6a618d6e4b1e1330057.tar.bz2 | |
[nvptx] Generalize bar.sync instruction
Allow the logical barrier operand of nvptx_barsync to be a register, and add a
thread count operand.
Build and reg-tested on x86_64 with nvptx accelerator.
2018-12-19 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.md (nvptx_barsync): Add and handle operand.
* config/nvptx/nvptx.c (nvptx_wsync): Update call to gen_nvptx_barsync.
From-SVN: r267258
| -rw-r--r-- | gcc/ChangeLog | 5 | ||||
| -rw-r--r-- | gcc/config/nvptx/nvptx.c | 2 | ||||
| -rw-r--r-- | gcc/config/nvptx/nvptx.md | 10 |
3 files changed, 14 insertions, 3 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 9596b26..a9f41a5 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,5 +1,10 @@ 2018-12-19 Tom de Vries <tdevries@suse.de> + * config/nvptx/nvptx.md (nvptx_barsync): Add and handle operand. + * config/nvptx/nvptx.c (nvptx_wsync): Update call to gen_nvptx_barsync. + +2018-12-19 Tom de Vries <tdevries@suse.de> + * config/nvptx/nvptx.c (nvptx_single): Always pass false to nvptx_wsync. (nvptx_process_pars): Likewise. diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index a354811..1ad3ba9 100644 --- a/gcc/config/nvptx/nvptx.c +++ b/gcc/config/nvptx/nvptx.c @@ -3974,7 +3974,7 @@ nvptx_wpropagate (bool pre_p, bool is_call, basic_block block, rtx_insn *insn) static rtx nvptx_wsync (bool after) { - return gen_nvptx_barsync (GEN_INT (after)); + return gen_nvptx_barsync (GEN_INT (after), GEN_INT (0)); } #if WORKAROUND_PTXJIT_BUG diff --git a/gcc/config/nvptx/nvptx.md b/gcc/config/nvptx/nvptx.md index ca00b1d..f1f6fe0 100644 --- a/gcc/config/nvptx/nvptx.md +++ b/gcc/config/nvptx/nvptx.md @@ -1454,10 +1454,16 @@ [(set_attr "atomic" "true")]) (define_insn "nvptx_barsync" - [(unspec_volatile [(match_operand:SI 0 "const_int_operand" "")] + [(unspec_volatile [(match_operand:SI 0 "nvptx_nonmemory_operand" "Ri") + (match_operand:SI 1 "const_int_operand")] UNSPECV_BARSYNC)] "" - "\\tbar.sync\\t%0;" + { + if (INTVAL (operands[1]) == 0) + return "\\tbar.sync\\t%0;"; + else + return "\\tbar.sync\\t%0, %1;"; + } [(set_attr "predicable" "false")]) (define_expand "memory_barrier" |