From: Tom de Vries <tdevries@suse.de>
Date: Wed, 19 Dec 2018 10:17:01 +0000 (+0000)
Subject: [nvptx] Generalize bar.sync instruction
X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=1dcf26882b1a8c4f388ed6a618d6e4b1e1330057;p=gcc.git

[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
---

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  <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
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"