In looking at some ptx output, OCD kicked in and I couldn't tolerate the
formatting inconsistencies. Fixed with this patch.
Also, the mechanism of providing scratch regs to the spin lock and reset insns,
caused the optimizers to want to insert initializations. Fixed by making these
patterns SETs. (We can't use the usual method of naming clobbers, because the
register allocator is disabled).
nathan
2015-09-01 Nathan Sidwell <nat...@codesourcery.com>
* config/nvptx/nvptx.md: Use tabs and operand spacing consistently
throughout.
(nvptx_spin_lock): Use set.
(nvptx_spin_reset): Likewise.
Index: config/nvptx/nvptx.md
===================================================================
--- config/nvptx/nvptx.md (revision 227369)
+++ config/nvptx/nvptx.md (working copy)
@@ -799,7 +799,7 @@
[(match_operand:HSDIM 2 "nvptx_register_operand" "R")
(match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")]))]
""
- "%.\\tsetp%c1 %0,%2,%3;")
+ "%.\\tsetp%c1\\t%0, %2, %3;")
(define_insn "*cmp<mode>"
[(set (match_operand:BI 0 "nvptx_register_operand" "=R")
@@ -807,7 +807,7 @@
[(match_operand:SDFM 2 "nvptx_register_operand" "R")
(match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")]))]
""
- "%.\\tsetp%c1 %0,%2,%3;")
+ "%.\\tsetp%c1\\t%0, %2, %3;")
(define_insn "jump"
[(set (pc)
@@ -941,7 +941,7 @@
[(match_operand:HSDIM 2 "nvptx_register_operand" "R")
(match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")]))]
""
- "%.\\tset%t0%c1 %0,%2,%3;")
+ "%.\\tset%t0%c1\\t%0, %2, %3;")
(define_insn "setcc_int<mode>"
[(set (match_operand:SI 0 "nvptx_register_operand" "=R")
@@ -949,7 +949,7 @@
[(match_operand:SDFM 2 "nvptx_register_operand" "R")
(match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")]))]
""
- "%.\\tset%t0%c1 %0,%2,%3;")
+ "%.\\tset%t0%c1\\t%0, %2, %3;")
(define_insn "setcc_float<mode>"
[(set (match_operand:SF 0 "nvptx_register_operand" "=R")
@@ -957,7 +957,7 @@
[(match_operand:HSDIM 2 "nvptx_register_operand" "R")
(match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")]))]
""
- "%.\\tset%t0%c1 %0,%2,%3;")
+ "%.\\tset%t0%c1\\t%0, %2, %3;")
(define_insn "setcc_float<mode>"
[(set (match_operand:SF 0 "nvptx_register_operand" "=R")
@@ -965,7 +965,7 @@
[(match_operand:SDFM 2 "nvptx_register_operand" "R")
(match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")]))]
""
- "%.\\tset%t0%c1 %0,%2,%3;")
+ "%.\\tset%t0%c1\\t%0, %2, %3;")
(define_expand "cstorebi4"
[(set (match_operand:SI 0 "nvptx_register_operand")
@@ -1343,9 +1343,9 @@
{
static const char *const asms[] =
{ /* Must match oacc_loop_levels ordering. */
- "%.\\tmov.u32 %0, %%nctaid.x;",/* gang */
- "%.\\tmov.u32 %0, %%ntid.y;", /* worker */
- "%.\\tmov.u32 %0, %%ntid.x;", /* vector */
+ "%.\\tmov.u32\\t%0, %%nctaid.x;", /* gang */
+ "%.\\tmov.u32\\t%0, %%ntid.y;", /* worker */
+ "%.\\tmov.u32\\t%0, %%ntid.x;", /* vector */
};
return asms[INTVAL (operands[1])];
})
@@ -1358,9 +1358,9 @@
{
static const char *const asms[] =
{ /* Must match oacc_loop_levels ordering. */
- "%.\\tmov.u32 %0, %%ctaid.x;",/* gang */
- "%.\\tmov.u32 %0, %%tid.y;", /* worker */
- "%.\\tmov.u32 %0, %%tid.x;", /* vector */
+ "%.\\tmov.u32\\t%0, %%ctaid.x;", /* gang */
+ "%.\\tmov.u32\\t%0, %%tid.y;", /* worker */
+ "%.\\tmov.u32\\t%0, %%tid.x;", /* vector */
};
return asms[INTVAL (operands[1])];
})
@@ -1460,7 +1460,7 @@
(set (match_operand:SI 1 "nvptx_register_operand" "=R")
(unspec:SI [(match_dup 2) (const_int 1)] UNSPEC_BIT_CONV))]
""
- "%.\\tmov.b64 {%0,%1}, %2;")
+ "%.\\tmov.b64\\t{%0,%1}, %2;")
;; pack 2 32-bit ints into a 64 bit object
(define_insn "packsi<mode>2"
@@ -1469,21 +1469,21 @@
(match_operand:SI 2 "nvptx_register_operand" "R")]
UNSPEC_BIT_CONV))]
""
- "%.\\tmov.b64 %0, {%1,%2};")
+ "%.\\tmov.b64\\t%0, {%1,%2};")
(define_insn "worker_load<mode>"
[(set (match_operand:SDISDFM 0 "nvptx_register_operand" "=R")
(unspec:SDISDFM [(match_operand:SDISDFM 1 "memory_operand" "m")]
UNSPEC_SHARED_DATA))]
""
- "%.\\tld.shared%u0\\t%0,%1;")
+ "%.\\tld.shared%u0\\t%0, %1;")
(define_insn "worker_store<mode>"
[(set (unspec:SDISDFM [(match_operand:SDISDFM 1 "memory_operand" "=m")]
UNSPEC_SHARED_DATA)
(match_operand:SDISDFM 0 "nvptx_register_operand" "R"))]
""
- "%.\\tst.shared%u1\\t%1,%0;")
+ "%.\\tst.shared%u1\\t%1, %0;")
;; Atomic insns.
@@ -1577,30 +1577,30 @@
[(unspec_volatile [(match_operand:SI 0 "const_int_operand" "")]
UNSPECV_BARSYNC)]
""
- "bar.sync\\t%0;")
+ "\\tbar.sync\\t%0;")
(define_insn "nvptx_membar"
[(unspec_volatile [(match_operand:SI 0 "const_int_operand" "")]
UNSPECV_MEMBAR)]
""
- "membar%B0;")
+ "%.\\tmembar%B0;")
;; spin lock and reset
(define_insn "nvptx_spin_lock"
[(parallel
- [(unspec_volatile [(match_operand:SI 0 "memory_operand" "m")
- (match_operand:SI 1 "const_int_operand" "i")]
- UNSPECV_LOCK)
- (match_operand:SI 2 "register_operand" "=R")
- (match_operand:BI 3 "register_operand" "=R")
+ [(set (match_operand:SI 2 "register_operand" "=R")
+ (unspec_volatile:SI [(match_operand:SI 0 "memory_operand" "m")
+ (match_operand:SI 1 "const_int_operand" "i")]
+ UNSPECV_LOCK))
+ (set (match_operand:BI 3 "register_operand" "=R") (const_int 0))
(label_ref (match_operand 4 "" ""))])]
""
- "%4:\\tatom%R1.cas.b32 %2,%0,0,1;setp.ne.u32 %3,%2,0;@%3 bra.uni %4;")
+ "%4:\\tatom%R1.cas.b32\\t%2, %0, 0, 1;\\n\\t\\tsetp.ne.u32\\t%3, %2, 0;\\n\\t@%3\\tbra.uni\\t%4;")
(define_insn "nvptx_spin_reset"
- [(unspec_volatile [(match_operand:SI 0 "memory_operand" "m")
- (match_operand:SI 1 "const_int_operand" "i")]
- UNSPECV_LOCK)
- (match_operand:SI 2 "register_operand" "=R")]
+ [(set (match_operand:SI 2 "register_operand" "=R")
+ (unspec_volatile:SI [(match_operand:SI 0 "memory_operand" "m")
+ (match_operand:SI 1 "const_int_operand" "i")]
+ UNSPECV_LOCK))]
""
- "atom%R1.exch.b32 %2,%0,0;")
+ "%.\\tatom%R1.exch.b32\\t%2, %0, 0;")
