* [gomp4] ptx assembler formatting
@ 2015-09-01 16:14 Nathan Sidwell
0 siblings, 0 replies; only message in thread
From: Nathan Sidwell @ 2015-09-01 16:14 UTC (permalink / raw)
To: GCC Patches
[-- Attachment #1: Type: text/plain, Size: 410 bytes --]
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
[-- Attachment #2: gomp4-ptx.patch --]
[-- Type: text/x-patch, Size: 5743 bytes --]
2015-09-01 Nathan Sidwell <nathan@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;")
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2015-09-01 16:14 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-09-01 16:14 [gomp4] ptx assembler formatting Nathan Sidwell
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).