From 1dcf26882b1a8c4f388ed6a618d6e4b1e1330057 Mon Sep 17 00:00:00 2001 From: Tom de Vries Date: Wed, 19 Dec 2018 10:17:01 +0000 Subject: [PATCH] [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 * 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 --- gcc/ChangeLog | 5 +++++ gcc/config/nvptx/nvptx.c | 2 +- gcc/config/nvptx/nvptx.md | 10 ++++++++-- 3 files changed, 14 insertions(+), 3 deletions(-) diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 9596b26860c..a9f41a52f92 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2018-12-19 Tom de Vries + + * 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 * config/nvptx/nvptx.c (nvptx_single): Always pass false to diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index a354811194c..1ad3ba92caa 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 ca00b1d8073..f1f6fe0c404 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" -- 2.30.2