This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[nvptx, committed] Generalize bar.sync instruction
- From: Tom de Vries <tdevries at suse dot de>
- To: "Schwinge, Thomas" <Thomas_Schwinge at mentor dot com>
- Cc: "gcc-patches at gcc dot gnu dot org" <gcc-patches at gcc dot gnu dot org>
- Date: Wed, 19 Dec 2018 11:33:22 +0100
- Subject: [nvptx, committed] Generalize bar.sync instruction
- References: <d101a77b-fc5f-0396-b8d1-c13f34fd1c38@codesourcery.com> <2ece5d7b-3675-84ab-f255-3c56a2ffd7dc@suse.de> <91b927af-d854-2865-7cbd-9a9a835ab5cc@codesourcery.com> <1394d89c-896e-f6a3-5f9a-78e98b16e85c@suse.de>
[ was: Re: [nvptx] vector length patch series ]
On 14-12-18 20:58, Tom de Vries wrote:
> 0011-nvptx-Add-thread-count-parm-to-bar.sync.patch
Factored out this patch, committed.
Thanks,
- Tom
[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-17 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.
---
gcc/config/nvptx/nvptx.c | 2 +-
gcc/config/nvptx/nvptx.md | 10 ++++++++--
2 files changed, 9 insertions(+), 3 deletions(-)
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"