Fixed an inline PTX bug.

The compiler was moving the 'asm("mov.u32 %0, $xr0;" : "=r" (xlow));' line out of the parent for loop, so specifying 'volatile' keeps it in place. This change can be seen by reading 'cudadl.o' after compiling with the nvcc '-ptx' option. The 'volatile' keyword was added to all of the inline PTX calls in 'parrhoasm.cu', but only the first one (line 237) is needed to fix the bug.

parrhoasm.cu

@@ -233,47 +233,47 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
         //cuPrintf("y    = %p %08X%08X\n", c_y, c_y[1], c_y[0]);
         //cuPrintf("z    = %p %08X%08X%08X%08X\n", ts[tid].z, ts[tid].z[3], ts[tid].z[2],  ts[tid].z[1], ts[tid].z[0]);
         int multtype;
-	asm("mov.u32 %0, $xr0;" : "=r" (xlow));
+	asm volatile("mov.u32 %0, $xr0;" : "=r" (xlow));
         if (xlow < TWO_32_DIV_3)
         {
             multtype = 0;
-	    asm("add.cc.u32 %0, %1, %2;" : "=r"(a_0) : "r"(a_0), "r"((unsigned int)1U));
-	    asm("addc.cc.u32 %0, %1, %2;" : "=r"(a_1) : "r"(a_1), "r"((unsigned int)0U));
-	    asm("addc.u32 %0, %1, %2;" : "=r"(a_2) : "r"(a_2), "r"((unsigned int)0U));
+	    asm volatile("add.cc.u32 %0, %1, %2;" : "=r"(a_0) : "r"(a_0), "r"((unsigned int)1U));
+	    asm volatile("addc.cc.u32 %0, %1, %2;" : "=r"(a_1) : "r"(a_1), "r"((unsigned int)0U));
+	    asm volatile("addc.u32 %0, %1, %2;" : "=r"(a_2) : "r"(a_2), "r"((unsigned int)0U));
 //	    a += 1;
 	    // cuPrintf("inc a\n");
         }
         else if (xlow < TWO_32_DIV_3_X2)
         {
             multtype = 1;
-	    asm("add.cc.u32 %0, %1, %2;" : "=r"(b_0) : "r"(b_0), "r"((unsigned int)1U));
-	    asm("addc.cc.u32 %0, %1, %2;" : "=r"(b_1) : "r"(b_1), "r"((unsigned int)0U));
-	    asm("addc.u32 %0, %1, %2;" : "=r"(b_2) : "r"(b_2), "r"((unsigned int)0U));
+	    asm volatile("add.cc.u32 %0, %1, %2;" : "=r"(b_0) : "r"(b_0), "r"((unsigned int)1U));
+	    asm volatile("addc.cc.u32 %0, %1, %2;" : "=r"(b_1) : "r"(b_1), "r"((unsigned int)0U));
+	    asm volatile("addc.u32 %0, %1, %2;" : "=r"(b_2) : "r"(b_2), "r"((unsigned int)0U));
 //	    b += 1;
 	    // cuPrintf("inc b\n");
         }
         else
         {
             multtype = 2;
-	    asm("add.cc.u32 %0, %1, %2;" : "=r"(a_0) : "r"(a_0), "r"(a_0));
-	    asm("addc.cc.u32 %0, %1, %2;" : "=r"(a_1) : "r"(a_1), "r"(a_1));
-	    asm("addc.u32 %0, %1, %2;" : "=r"(a_2) : "r"(a_2), "r"(a_2));
-	    asm("add.cc.u32 %0, %1, %2;" : "=r"(b_0) : "r"(b_0), "r"(b_0));
-	    asm("addc.cc.u32 %0, %1, %2;" : "=r"(b_1) : "r"(b_1), "r"(b_1));
-	    asm("addc.u32 %0, %1, %2;" : "=r"(b_2) : "r"(b_2), "r"(b_2));
+	    asm volatile("add.cc.u32 %0, %1, %2;" : "=r"(a_0) : "r"(a_0), "r"(a_0));
+	    asm volatile("addc.cc.u32 %0, %1, %2;" : "=r"(a_1) : "r"(a_1), "r"(a_1));
+	    asm volatile("addc.u32 %0, %1, %2;" : "=r"(a_2) : "r"(a_2), "r"(a_2));
+	    asm volatile("add.cc.u32 %0, %1, %2;" : "=r"(b_0) : "r"(b_0), "r"(b_0));
+	    asm volatile("addc.cc.u32 %0, %1, %2;" : "=r"(b_1) : "r"(b_1), "r"(b_1));
+	    asm volatile("addc.u32 %0, %1, %2;" : "=r"(b_2) : "r"(b_2), "r"(b_2));
 //	    a += a;
 //	    b += b;
 	    // cuPrintf("double\n");
         }
 	if (a_2 > order_2 || ((a_2 == order_2) && (a_1 > order_1)) || (((a_2 == order_2) && (a_1 == order_1) && (a_0 > order_0)))) {
-	    asm("sub.cc.u32 %0, %1, %2;" : "=r"(a_0) : "r"(a_0), "r"(order_0));
-	    asm("subc.cc.u32 %0, %1, %2;" : "=r"(a_1) : "r"(a_1), "r"(order_1));
-	    asm("subc.u32 %0, %1, %2;" : "=r"(a_2) : "r"(a_2), "r"(order_2));
+	    asm volatile("sub.cc.u32 %0, %1, %2;" : "=r"(a_0) : "r"(a_0), "r"(order_0));
+	    asm volatile("subc.cc.u32 %0, %1, %2;" : "=r"(a_1) : "r"(a_1), "r"(order_1));
+	    asm volatile("subc.u32 %0, %1, %2;" : "=r"(a_2) : "r"(a_2), "r"(order_2));
 	}
 	if (b_2 > order_2 || ((b_2 == order_2) && (b_1 > order_1)) || (((b_2 == order_2) && (b_1 == order_1) && (b_0 > order_0)))) {
-	    asm("sub.cc.u32 %0, %1, %2;" : "=r"(b_0) : "r"(b_0), "r"(order_0));
-	    asm("subc.cc.u32 %0, %1, %2;" : "=r"(b_1) : "r"(b_1), "r"(order_1));
-	    asm("subc.u32 %0, %1, %2;" : "=r"(b_2) : "r"(b_2), "r"(order_2));
+	    asm volatile("sub.cc.u32 %0, %1, %2;" : "=r"(b_0) : "r"(b_0), "r"(order_0));
+	    asm volatile("subc.cc.u32 %0, %1, %2;" : "=r"(b_1) : "r"(b_1), "r"(order_1));
+	    asm volatile("subc.u32 %0, %1, %2;" : "=r"(b_2) : "r"(b_2), "r"(order_2));
 	}
 		
         CIOS_MODMUL(multtype);
@@ -288,7 +288,7 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
         //cuPrintf("y    = %p %08X%08X\n", c_y, c_y[1], c_y[0]);
 
 	// Check for a distinguished point
-	asm("mov.u32 %0, $xr0;" : "=r" (xlow));
+	asm volatile("mov.u32 %0, $xr0;" : "=r" (xlow));
 	if (xlow <= dpfreq) {
 	    unsigned int *ourbuffer = DPstreamAlloc();
 	    if (ourbuffer) {