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