Added comments in the CUDA kernel code.

parrhoasm.cu

@@ -214,6 +214,7 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
     unsigned int xlow;
 
     // cuPrintf("d_z = %08X%08X%08X\n", global_x_base[2], global_x_base[1], global_x_base[0]);
+
     CIOS_LOADX(global_x_base);
     a_2 = global_a_base[2];
     a_1 = global_a_base[1];
@@ -221,11 +222,13 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
     b_2 = global_b_base[2];
     b_1 = global_b_base[1];
     b_0 = global_b_base[0];
+
     // cuPrintf("s_A = %08X\n", (unsigned int)a);
     // cuPrintf("s_B = %08X\n", (unsigned int)b);
     // cuPrintf("s_a = %08X%08X\n", global_a_base[1], global_a_base[0]);
     // cuPrintf("s_b = %08X%08X\n", global_b_base[1], global_b_base[0]);
-    #pragma unroll 1 // to allow 'nmult' to be 2 or 3
+
+    #pragma unroll 1 // to prevent namespace conflicts caused by loop unrolling when 'nmult' is small
     for (i = 0; i < nmult; ++i)
     {
         //memset(ts[tid].z, 0, (WORDS + 2) * sizeof(unsigned int));
@@ -233,8 +236,12 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
         //cuPrintf("x    = %p %08X%08X%08X\n", ts[tid].x, ts[tid].x[2], ts[tid].x[1], ts[tid].x[0]);
         //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 volatile("mov.u32 %0, $xr0;" : "=r" (xlow));
+	// volatile qualifier needed to prevent the compiler from moving it outside of the loop
+	// change can be seen by reading 'cudadl.o' after compiling with the nvcc '-ptx' option
+
         if (xlow < TWO_32_DIV_3)
         {
             multtype = 0;
@@ -278,6 +285,7 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
 	}
 		
         CIOS_MODMUL(multtype);
+
         //memcpy(ts[tid].x, ts[tid].z, (WORDS + 1) * sizeof(unsigned int));
 	/*
         if (_gt(shared_x_base)) {
@@ -297,6 +305,7 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
 	    }
 	}
     }
+
     CIOS_SAVEX(global_x_base);
     global_a_base[0] = a_0;
     global_a_base[1] = a_1;
@@ -304,6 +313,7 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
     global_b_base[0] = b_0;
     global_b_base[1] = b_1;
     global_b_base[2] = b_2;
+
     // cuPrintf("d_z = %08X%08X%08X\n", global_x_base[2], global_x_base[1], global_x_base[0]);
     // cuPrintf("d_A = %08X\n", (unsigned int)a);
     // cuPrintf("d_B = %08X\n", (unsigned int)b);