Use macros to include the inline assembly instead of inline functions

The new nvcc version 4 no longer allows registers to be declared at
module scope, so we need to do all the work that uses them inside
one function.  Have cios.asm now provide macros instead of functions
to do all the asm work.

dpstream.cu

@@ -61,30 +61,26 @@ void DPstreamEnd(void)
     DPbuffer_device = NULL;
 }
 
-__device__ inline void write_dp(unsigned int*, unsigned int, unsigned int,
-    unsigned int, unsigned int, unsigned int, unsigned int);
-
 // Call this from the device to write the following into the buffer:
 //   block and thread indices (1 word total)
 //   x (WORDS words)
 //   a (2 words)
 //   b (2 words)
-__device__ inline void DPstreamWrite(unsigned int a_0, unsigned int a_1,
+__device__ inline unsigned int *DPstreamWrite(unsigned int a_0, unsigned int a_1,
     unsigned int a_2, unsigned int b_0, unsigned int b_1, unsigned int b_2)
 {
-    if (!DPbuffer) return;
+    if (!DPbuffer) return NULL;
 
     // Atomically increment the DPbuffertail
-    if (DPbuffertail >= DPbuffersize) return;
+    if (DPbuffertail >= DPbuffersize) return NULL;
     unsigned int ouroffset = atomicAdd((unsigned int *)&DPbuffertail,
 		(unsigned int) (DPrecordsize));
     if (ouroffset + DPrecordsize > DPbuffersize) {
 	atomicSub((unsigned int *)&DPbuffertail, (unsigned int)DPrecordsize);
-	return;
+	return NULL;
     }
     unsigned int *ourbuffer = DPbuffer + ouroffset;
-
-    write_dp(ourbuffer, a_0, a_1, a_2, b_0, b_1, b_2);
+    return ourbuffer;
 }
 
 static void dpcallback(void *data, unsigned short threadId,

gencios_reg_20

@@ -32,296 +32,288 @@ my $i;
 ## scope
 
 print <<EOA;
-// An evil hack that declares a bunch of registers in the module scope
-__device__ __noinline__ void evilhack() {
-    asm("ret;}\\n"
-        ".reg.u32 \$xr<$WORDSp1>;\\n"
-        ".reg.u32 \$cg<$WORDS>;\\n"
-        ".reg.u32 \$cy<$WORDS>;\\n"
-	".reg.u32 \$yi,\$crp,\$m;\\n"
-	".reg.u64 \$cs,\$ct;\\n"
-	".reg.u32 \$zr<$WORDSp2>;\\n"
-	".reg.u32 \$crho<$WORDS>;\\n"
-	".func hackend {\\n"
-EOA
-print "\t\"ret.uni;\\n\"\n";
+#define CIOS_REG_DECLARE \\
+    asm volatile ( \\
+        ".reg.u32 \$xr<$WORDSp1>;\\n" \\
+        ".reg.u32 \$cg<$WORDS>;\\n" \\
+        ".reg.u32 \$cy<$WORDS>;\\n" \\
+	".reg.u32 \$yi,\$crp,\$m;\\n" \\
+	".reg.u64 \$cs,\$ct;\\n" \\
+	".reg.u32 \$zr<$WORDSp2>;\\n" \\
+	".reg.u32 \$crho<$WORDS>;\\n" \\
+    )
 
-print "    );\n";
-print "}\n";
+EOA
 
 print <<'EOA';
-__device__ inline void loadx(unsigned int *x)
-{
-    asm(
+#define CIOS_LOADX(x) \
+    asm volatile ( \
 EOA
 
 for ($i=0;$i<$WORDS;++$i) {
     my $xoff = $i*4;
-    print "\t\"ld.global.u32 \$xr$i, [%0+$xoff];\\n\"\n";
+    print "\t\"ld.global.u32 \$xr$i, [%0+$xoff];\\n\" \\\n";
 }
-print "\t\"// Load c_rho into registers\\n\"\n";
+print "\t\"// Load c_rho into registers\\n\" \\\n";
 my ($xoff, $xoffm1);
 my $croff;
 for($i=0;$i<$WORDS;++$i) {
     $croff = $i*4;
-    print "\t\"ld.const.u32 \$crho$i, [c_rho+$croff];\\n\"\n";
+    print "\t\"ld.const.u32 \$crho$i, [c_rho+$croff];\\n\" \\\n";
 }
-print "\t\"mov.u32 \$crp, %1;\\n\"\n";
-print "\t\"// Load c_g into registers\\n\"\n";
+print "\t\"mov.u32 \$crp, %1;\\n\" \\\n";
+print "\t\"// Load c_g into registers\\n\" \\\n";
 for($i=0;$i<$WORDS;++$i) {
     $croff = $i*4;
-    print "\t\"ld.const.u32 \$cg$i, [c_g+$croff];\\n\"\n";
+    print "\t\"ld.const.u32 \$cg$i, [c_g+$croff];\\n\" \\\n";
 }
-print "\t\"// Load c_y into registers\\n\"\n";
+print "\t\"// Load c_y into registers\\n\" \\\n";
 for($i=0;$i<$WORDS;++$i) {
     $croff = $i*4;
-    print "\t\"ld.const.u32 \$cy$i, [c_y+$croff];\\n\"\n";
+    print "\t\"ld.const.u32 \$cy$i, [c_y+$croff];\\n\" \\\n";
 }
 
 
 print <<'EOA';
-    : : "l" (x), "r" (c_rho_prime));
-}
+    : : "l" (x), "r" (c_rho_prime))
 
-__device__ inline void savex(unsigned int *x)
-{
-    asm(
+EOA
+
+print <<'EOA';
+#define CIOS_SAVEX(x) \
+    asm volatile ( \
 EOA
 
 for ($i=0;$i<$WORDS;++$i) {
     my $xoff = $i*4;
-    print "\t\"st.global.u32 [%0+$xoff], \$xr$i;\\n\"\n";
+    print "\t\"st.global.u32 [%0+$xoff], \$xr$i;\\n\" \\\n";
 }
 
 print <<'EOA';
-    : : "l" (x));
-}
+    : : "l" (x))
+
 EOA
 
 print <<'EOA';
-__device__ void modmul(int multtype)
 // x <- (x * y) * r_inv   (mod rho);
 // where y is (multtype == 0 ? c_g : multtype == 1 ? c_y : x)
-{
-    asm(
-    ".reg.pred $mtp0,$mtp1;\n"
+#define CIOS_MODMUL(multtype) \
+    asm volatile ( \
+    ".reg.pred $mtp0,$mtp1;\n" \
 EOA
-print "    \".reg.pred \$retaddr<", $WORDS-2, ">;\\n\"\n";
+print "    \".reg.pred \$retaddr<", $WORDS-2, ">;\\n\" \\\n";
 print <<'EOA';
-    "mov.u32 $m, 0;\n"
-    "setp.eq.u32 $mtp0, %0, $m;\n"
-    "mov.u32 $m, 1;\n"
-    "setp.eq.u32 $mtp1, %0, $m;\n"
-    : : "r" (multtype));
+    "mov.u32 $m, 0;\n" \
+    "setp.eq.u32 $mtp0, %0, $m;\n" \
+    "mov.u32 $m, 1;\n" \
+    "setp.eq.u32 $mtp1, %0, $m;\n" \
+    : : "r" (multtype)); \
 EOA
 
-print "    asm(\n";
-print "\t\"// First iteration: i=0 and z is unset\\n\"\n";
+print "    asm volatile ( \\\n";
+print "\t\"// First iteration: i=0 and z is unset\\n\" \\\n";
 
 # Pick the correct value of $yi
-print "\t\"selp.u32 \$yi, \$cg0, \$xr0, \$mtp0;\\n\"\n";
-print "\t\"selp.u32 \$yi, \$cy0, \$yi, \$mtp1;\\n\"\n";
+print "\t\"selp.u32 \$yi, \$cg0, \$xr0, \$mtp0;\\n\" \\\n";
+print "\t\"selp.u32 \$yi, \$cy0, \$yi, \$mtp1;\\n\" \\\n";
 
-print "\t\"mul.wide.u32 \$cs, \$xr0, \$yi;\\n\"\n";
-print "\t\"cvt.u32.u64 \$zr0, \$cs;\\n\"\n";
-print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
+print "\t\"mul.wide.u32 \$cs, \$xr0, \$yi;\\n\" \\\n";
+print "\t\"cvt.u32.u64 \$zr0, \$cs;\\n\" \\\n";
+print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
 
 for($i=1; $i<$WORDS; ++$i) {
     $xoff = $i*32*4;
-    print "\t\"mad.wide.u32 \$cs, \$xr$i, \$yi, \$cs;\\n\"\n";
-    print "\t\"cvt.u32.u64 \$zr$i, \$cs;\\n\"\n";
-    print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
+    print "\t\"mad.wide.u32 \$cs, \$xr$i, \$yi, \$cs;\\n\" \\\n";
+    print "\t\"cvt.u32.u64 \$zr$i, \$cs;\\n\" \\\n";
+    print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
 }
-print "\t\"cvt.u32.u64 \$zr$WORDS, \$cs;\\n\"\n";
+print "\t\"cvt.u32.u64 \$zr$WORDS, \$cs;\\n\" \\\n";
 
-print "\t\"mul.lo.u32 \$m, \$zr0, \$crp;\\n\"\n";
-print "\t\"cvt.u64.u32 \$ct, \$zr0;\\n\"\n";
-print "\t\"mad.wide.u32 \$cs, \$m, \$crho0, \$ct;\\n\"\n";
+print "\t\"mul.lo.u32 \$m, \$zr0, \$crp;\\n\" \\\n";
+print "\t\"cvt.u64.u32 \$ct, \$zr0;\\n\" \\\n";
+print "\t\"mad.wide.u32 \$cs, \$m, \$crho0, \$ct;\\n\" \\\n";
 
 for($i=1;$i<$WORDS;++$i) {
     my $im1 = ($i-1);
-    print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
-    print "\t\"cvt.u64.u32 \$ct, \$zr$i;\\n\"\n";
-    print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\"\n";
-    print "\t\"mad.wide.u32 \$cs, \$m, \$crho$i, \$cs;\\n\"\n";
-    print "\t\"cvt.u32.u64 \$zr$im1, \$cs;\\n\"\n";
+    print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
+    print "\t\"cvt.u64.u32 \$ct, \$zr$i;\\n\" \\\n";
+    print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\" \\\n";
+    print "\t\"mad.wide.u32 \$cs, \$m, \$crho$i, \$cs;\\n\" \\\n";
+    print "\t\"cvt.u32.u64 \$zr$im1, \$cs;\\n\" \\\n";
 }
-print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
-print "\t\"cvt.u64.u32 \$ct, \$zr$WORDS;\\n\"\n";
-print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\"\n";
-print "\t\"cvt.u32.u64 \$zr$WORDSm1, \$cs;\\n\"\n";
-print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
-print "\t\"cvt.u32.u64 \$zr$WORDS, \$cs;\\n\"\n";
+print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
+print "\t\"cvt.u64.u32 \$ct, \$zr$WORDS;\\n\" \\\n";
+print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\" \\\n";
+print "\t\"cvt.u32.u64 \$zr$WORDSm1, \$cs;\\n\" \\\n";
+print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
+print "\t\"cvt.u32.u64 \$zr$WORDS, \$cs;\\n\" \\\n";
 
 my $j;
 for ($j=1; $j<$WORDS-1; ++$j) {
 
 print <<'EOA';
-	"// Middle iterations\n"
+	"// Middle iterations\n" \
 EOA
 
 # Pick the correct value of $yi
 my $yioff = $j*4;
-print "\t\"selp.u32 \$yi, \$cg$j, \$xr$j, \$mtp0;\\n\"\n";
-print "\t\"selp.u32 \$yi, \$cy$j, \$yi, \$mtp1;\\n\"\n";
+print "\t\"selp.u32 \$yi, \$cg$j, \$xr$j, \$mtp0;\\n\" \\\n";
+print "\t\"selp.u32 \$yi, \$cy$j, \$yi, \$mtp1;\\n\" \\\n";
 my $jm1 = $j-1;
 my $jm2 = $j-2;
 if ($j==1) {
-    print "\t\"setp.eq.u32 \$retaddr0,\$xr0,\$xr0;\\n\"\n";
+    print "\t\"setp.eq.u32 \$retaddr0,\$xr0,\$xr0;\\n\" \\\n";
 } else {
-    print "\t\"setp.eq.u32 \$retaddr$jm1|\$retaddr$jm2,\$xr0,\$xr0;\\n\"\n";
+    print "\t\"setp.eq.u32 \$retaddr$jm1|\$retaddr$jm2,\$xr0,\$xr0;\\n\" \\\n";
 }
-print "\t\"bra.uni innerloop;\\n\"\n";
-print "\t\"innerret$jm1:\\n\"\n";
+print "\t\"bra.uni innerloop;\\n\" \\\n";
+print "\t\"innerret$jm1:\\n\" \\\n";
 }
 
 print <<'EOA';
-	"// Last iteration: store the output in x instead of z\n"
+	"// Last iteration: store the output in x instead of z\n" \
 EOA
 
 # Pick the correct value of $yi
 my $yioff = ($WORDSm1)*4;
-print "\t\"selp.u32 \$yi, \$cg$WORDSm1, \$xr$WORDSm1, \$mtp0;\\n\"\n";
-print "\t\"selp.u32 \$yi, \$cy$WORDSm1, \$yi, \$mtp1;\\n\"\n";
+print "\t\"selp.u32 \$yi, \$cg$WORDSm1, \$xr$WORDSm1, \$mtp0;\\n\" \\\n";
+print "\t\"selp.u32 \$yi, \$cy$WORDSm1, \$yi, \$mtp1;\\n\" \\\n";
 
-print "\t\"cvt.u64.u32 \$ct, \$zr0;\\n\"\n";
-print "\t\"mad.wide.u32 \$cs, \$xr0, \$yi, \$ct;\\n\"\n";
-print "\t\"cvt.u32.u64 \$zr0, \$cs;\\n\"\n";
-print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
+print "\t\"cvt.u64.u32 \$ct, \$zr0;\\n\" \\\n";
+print "\t\"mad.wide.u32 \$cs, \$xr0, \$yi, \$ct;\\n\" \\\n";
+print "\t\"cvt.u32.u64 \$zr0, \$cs;\\n\" \\\n";
+print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
 
 for($i=1; $i<$WORDS; ++$i) {
     $xoff = $i*32*4;
-    print "\t\"cvt.u64.u32 \$ct, \$zr$i;\\n\"\n";
-    print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\"\n";
-    print "\t\"mad.wide.u32 \$cs, \$xr$i, \$yi, \$cs;\\n\"\n";
-    print "\t\"cvt.u32.u64 \$zr$i, \$cs;\\n\"\n";
-    print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
+    print "\t\"cvt.u64.u32 \$ct, \$zr$i;\\n\" \\\n";
+    print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\" \\\n";
+    print "\t\"mad.wide.u32 \$cs, \$xr$i, \$yi, \$cs;\\n\" \\\n";
+    print "\t\"cvt.u32.u64 \$zr$i, \$cs;\\n\" \\\n";
+    print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
 }
-print "\t\"cvt.u64.u32 \$ct, \$zr$WORDS;\\n\"\n";
-print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\"\n";
-print "\t\"cvt.u32.u64 \$zr$WORDS, \$cs;\\n\"\n";
-print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
-print "\t\"cvt.u32.u64 \$zr$WORDSp1, \$cs;\\n\"\n";
+print "\t\"cvt.u64.u32 \$ct, \$zr$WORDS;\\n\" \\\n";
+print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\" \\\n";
+print "\t\"cvt.u32.u64 \$zr$WORDS, \$cs;\\n\" \\\n";
+print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
+print "\t\"cvt.u32.u64 \$zr$WORDSp1, \$cs;\\n\" \\\n";
 
-print "\t\"mul.lo.u32 \$m, \$zr0, \$crp;\\n\"\n";
-print "\t\"cvt.u64.u32 \$ct, \$zr0;\\n\"\n";
-print "\t\"mad.wide.u32 \$cs, \$m, \$crho0, \$ct;\\n\"\n";
+print "\t\"mul.lo.u32 \$m, \$zr0, \$crp;\\n\" \\\n";
+print "\t\"cvt.u64.u32 \$ct, \$zr0;\\n\" \\\n";
+print "\t\"mad.wide.u32 \$cs, \$m, \$crho0, \$ct;\\n\" \\\n";
 
 for($i=1;$i<$WORDS;++$i) {
     my $im1 = $i-1;
-    print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
-    print "\t\"cvt.u64.u32 \$ct, \$zr$i;\\n\"\n";
-    print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\"\n";
-    print "\t\"mad.wide.u32 \$cs, \$m, \$crho$i, \$cs;\\n\"\n";
-    print "\t\"cvt.u32.u64 \$xr$im1, \$cs;\\n\"\n";
+    print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
+    print "\t\"cvt.u64.u32 \$ct, \$zr$i;\\n\" \\\n";
+    print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\" \\\n";
+    print "\t\"mad.wide.u32 \$cs, \$m, \$crho$i, \$cs;\\n\" \\\n";
+    print "\t\"cvt.u32.u64 \$xr$im1, \$cs;\\n\" \\\n";
 }
-print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
-print "\t\"cvt.u64.u32 \$ct, \$zr$WORDS;\\n\"\n";
-print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\"\n";
-print "\t\"cvt.u32.u64 \$xr$WORDSm1, \$cs;\\n\"\n";
-print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
-print "\t\"cvt.u64.u32 \$ct, \$zr$WORDSp1;\\n\"\n";
-print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\"\n";
-print "\t\"cvt.u32.u64 \$xr$WORDS, \$cs;\\n\"\n";
+print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
+print "\t\"cvt.u64.u32 \$ct, \$zr$WORDS;\\n\" \\\n";
+print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\" \\\n";
+print "\t\"cvt.u32.u64 \$xr$WORDSm1, \$cs;\\n\" \\\n";
+print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
+print "\t\"cvt.u64.u32 \$ct, \$zr$WORDSp1;\\n\" \\\n";
+print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\" \\\n";
+print "\t\"cvt.u32.u64 \$xr$WORDS, \$cs;\\n\" \\\n";
 
 print <<'EOA';
-        );
+        ); \
 EOA
 
 # If x >= c_rho, x <- x - c_rho
 
-print "    asm(\n";
-print "\t\".reg.pred \$gtP,\$ltP;\\n\"\n";
-print "\t\"mov.u32 \$m, 0;\\n\"\n";
-print "\t\"setp.gt.u32 \$gtP, \$xr$WORDS, \$m;\\n\"\n";
-print "\t\"\@\$gtP bra \$comp_done;\\n\"\n";
+print "    asm volatile ( \\\n";
+print "\t\".reg.pred \$gtP,\$ltP;\\n\" \\\n";
+print "\t\"mov.u32 \$m, 0;\\n\" \\\n";
+print "\t\"setp.gt.u32 \$gtP, \$xr$WORDS, \$m;\\n\" \\\n";
+print "\t\"\@\$gtP bra \$comp_done;\\n\" \\\n";
 for ($i = $WORDS-1; $i >= 1; --$i) {
-    print "\t\"setp.gt.u32 \$gtP, \$xr$i, \$crho$i;\\n\"\n";
-    print "\t\"\@\$gtP bra \$comp_done;\\n\"\n";
-    print "\t\"setp.lt.u32 \$ltP, \$xr$i, \$crho$i;\\n\"\n";
-    print "\t\"\@\$ltP bra \$comp_done;\\n\"\n";
+    print "\t\"setp.gt.u32 \$gtP, \$xr$i, \$crho$i;\\n\" \\\n";
+    print "\t\"\@\$gtP bra \$comp_done;\\n\" \\\n";
+    print "\t\"setp.lt.u32 \$ltP, \$xr$i, \$crho$i;\\n\" \\\n";
+    print "\t\"\@\$ltP bra \$comp_done;\\n\" \\\n";
 }
-print "\t\"setp.ge.u32 \$gtP, \$xr0, \$crho0;\\n\"\n";
-print "\t\"\$comp_done:\\n\"\n";
+print "\t\"setp.ge.u32 \$gtP, \$xr0, \$crho0;\\n\" \\\n";
+print "\t\"\$comp_done:\\n\" \\\n";
 
 # Do the subtraction
-print "\t\"@\$gtP sub.cc.u32 \$xr0, \$xr0, \$crho0;\\n\"\n";
+print "\t\"@\$gtP sub.cc.u32 \$xr0, \$xr0, \$crho0;\\n\" \\\n";
 for ($i=1; $i<$WORDS; ++$i) {
-    print "\t\"@\$gtP subc.cc.u32 \$xr$i, \$xr$i, \$crho$i;\\n\"\n";
+    print "\t\"@\$gtP subc.cc.u32 \$xr$i, \$xr$i, \$crho$i;\\n\" \\\n";
 }
-print "\t\"@\$gtP subc.u32 \$xr$WORDS, \$xr$WORDS, \$m;\\n\"\n";
+print "\t\"@\$gtP subc.u32 \$xr$WORDS, \$xr$WORDS, \$m;\\n\" \\\n";
 
-print "\t\"bra.uni mulmoddone;\\n\"\n";
+print "\t\"bra.uni mulmoddone;\\n\" \\\n";
 
-print "\t\"innerloop:\\n\"\n";
-print "\t\"cvt.u64.u32 \$ct, \$zr0;\\n\"\n";
-print "\t\"mad.wide.u32 \$cs, \$xr0, \$yi, \$ct;\\n\"\n";
-print "\t\"cvt.u32.u64 \$zr0, \$cs;\\n\"\n";
-print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
+print "\t\"innerloop:\\n\" \\\n";
+print "\t\"cvt.u64.u32 \$ct, \$zr0;\\n\" \\\n";
+print "\t\"mad.wide.u32 \$cs, \$xr0, \$yi, \$ct;\\n\" \\\n";
+print "\t\"cvt.u32.u64 \$zr0, \$cs;\\n\" \\\n";
+print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
 
 my $i;
 for($i=1; $i<$WORDS; ++$i) {
-    print "\t\"cvt.u64.u32 \$ct, \$zr$i;\\n\"\n";
-    print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\"\n";
-    print "\t\"mad.wide.u32 \$cs, \$xr$i, \$yi, \$cs;\\n\"\n";
-    print "\t\"cvt.u32.u64 \$zr$i, \$cs;\\n\"\n";
-    print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
+    print "\t\"cvt.u64.u32 \$ct, \$zr$i;\\n\" \\\n";
+    print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\" \\\n";
+    print "\t\"mad.wide.u32 \$cs, \$xr$i, \$yi, \$cs;\\n\" \\\n";
+    print "\t\"cvt.u32.u64 \$zr$i, \$cs;\\n\" \\\n";
+    print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
 }
-print "\t\"cvt.u64.u32 \$ct, \$zr$WORDS;\\n\"\n";
-print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\"\n";
-print "\t\"cvt.u32.u64 \$zr$WORDS, \$cs;\\n\"\n";
-print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
-print "\t\"cvt.u32.u64 \$zr$WORDSp1, \$cs;\\n\"\n";
+print "\t\"cvt.u64.u32 \$ct, \$zr$WORDS;\\n\" \\\n";
+print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\" \\\n";
+print "\t\"cvt.u32.u64 \$zr$WORDS, \$cs;\\n\" \\\n";
+print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
+print "\t\"cvt.u32.u64 \$zr$WORDSp1, \$cs;\\n\" \\\n";
 
-print "\t\"mul.lo.u32 \$m, \$zr0, \$crp;\\n\"\n";
-print "\t\"cvt.u64.u32 \$ct, \$zr0;\\n\"\n";
-print "\t\"mad.wide.u32 \$cs, \$m, \$crho0, \$ct;\\n\"\n";
+print "\t\"mul.lo.u32 \$m, \$zr0, \$crp;\\n\" \\\n";
+print "\t\"cvt.u64.u32 \$ct, \$zr0;\\n\" \\\n";
+print "\t\"mad.wide.u32 \$cs, \$m, \$crho0, \$ct;\\n\" \\\n";
 
 for($i=1;$i<$WORDS;++$i) {
     my $im1 = ($i-1);
-    print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
-    print "\t\"cvt.u64.u32 \$ct, \$zr$i;\\n\"\n";
-    print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\"\n";
-    print "\t\"mad.wide.u32 \$cs, \$m, \$crho$i, \$cs;\\n\"\n";
-    print "\t\"cvt.u32.u64 \$zr$im1, \$cs;\\n\"\n";
+    print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
+    print "\t\"cvt.u64.u32 \$ct, \$zr$i;\\n\" \\\n";
+    print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\" \\\n";
+    print "\t\"mad.wide.u32 \$cs, \$m, \$crho$i, \$cs;\\n\" \\\n";
+    print "\t\"cvt.u32.u64 \$zr$im1, \$cs;\\n\" \\\n";
 }
-print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
-print "\t\"cvt.u64.u32 \$ct, \$zr$WORDS;\\n\"\n";
-print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\"\n";
-print "\t\"cvt.u32.u64 \$zr$WORDSm1, \$cs;\\n\"\n";
-print "\t\"shr.u64 \$cs, \$cs, 32;\\n\"\n";
-print "\t\"cvt.u32.u64 \$zr$WORDS, \$cs;\\n\"\n";
+print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
+print "\t\"cvt.u64.u32 \$ct, \$zr$WORDS;\\n\" \\\n";
+print "\t\"add.u64 \$cs, \$cs, \$ct;\\n\" \\\n";
+print "\t\"cvt.u32.u64 \$zr$WORDSm1, \$cs;\\n\" \\\n";
+print "\t\"shr.u64 \$cs, \$cs, 32;\\n\" \\\n";
+print "\t\"cvt.u32.u64 \$zr$WORDS, \$cs;\\n\" \\\n";
 
 for ($i=0;$i<$WORDS-2;++$i) {
-    print "\t\"\@\$retaddr$i bra.uni innerret$i;\\n\"\n";
+    print "\t\"\@\$retaddr$i bra.uni innerret$i;\\n\" \\\n";
 }
-print "\t\"mulmoddone:\\n\"\n";
+print "\t\"mulmoddone:\\n\" \\\n";
 
-print "    );\n";
-print "}\n";
+print "    )\n";
 
 print <<'EOA';
 
-__device__ inline void write_dp(unsigned int *buf, unsigned int a_0,
-    unsigned int a_1, unsigned int a_2, unsigned int b_0, unsigned int b_1,
-    unsigned int b_2)
-{
-    unsigned short blockid = blockIdx.x + gridDim.x*blockIdx.y;
-    unsigned short threadid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
-    buf[0] = (((unsigned int)blockid) << 16) + threadid;
-    asm(
+#define CIOS_WRITE_DP(buf, a_0, a_1, a_2, b_0, b_1, b_2) \
+do { \
+    unsigned short blockid = blockIdx.x + gridDim.x*blockIdx.y; \
+    unsigned short threadid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z; \
+    buf[0] = (((unsigned int)blockid) << 16) + threadid; \
+    asm volatile ( \
 EOA
 for($i=0;$i<$WORDS;++$i) {
     my $off = 4*($i+1);
-    print "\t\"st.global.u32 [%0+$off], \$xr$i;\\n\"\n";
+    print "\t\"st.u32 [%0+$off], \$xr$i;\\n\" \\\n";
 }
 print <<'EOA';
-    : : "l" (buf));
-    buf[WORDS+1] = a_0;
-    buf[WORDS+2] = a_1;
-    buf[WORDS+3] = a_2;
-    buf[WORDS+4] = b_0;
-    buf[WORDS+5] = b_1;
-    buf[WORDS+6] = b_2;
-}
+    : : "l" (buf)); \
+    buf[WORDS+1] = a_0; \
+    buf[WORDS+2] = a_1; \
+    buf[WORDS+3] = a_2; \
+    buf[WORDS+4] = b_0; \
+    buf[WORDS+5] = b_1; \
+    buf[WORDS+6] = b_2; \
+} while(0)
 EOA

parrhoasm.cu

@@ -42,6 +42,11 @@
 #define SIZE_Y WORDS
 #define SIZE_X (WORDS + 1)
 
+__device__ __constant__ unsigned int c_rho[WORDS + 1], c_r_inv[WORDS + 1], c_rho_prime;
+__device__ __constant__ unsigned int c_y[SIZE_Y], c_g[SIZE_G];
+
+#include "cios.asm"
+
 #include "dpstream.cu"
 
 // #define X(idx) x[idx<<5]
@@ -85,9 +90,6 @@ typedef struct
 } GlobalThreadState;
 
 
-__device__ __constant__ unsigned int c_rho[WORDS + 1], c_r_inv[WORDS + 1], c_rho_prime;
-__device__ __constant__ unsigned int c_y[SIZE_Y], c_g[SIZE_G];
-
 /*
 __device__ unsigned int getgy(int i)
 {
@@ -95,8 +97,6 @@ __device__ unsigned int getgy(int i)
 }
 */
 
-#include "cios.asm"
-
 #if 0
 __device__ void _sub(unsigned int *x)
 // x <- x - y
@@ -191,6 +191,8 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
 	unsigned int order_0, unsigned int order_1, unsigned int order_2)
 // x <- x * y mod rho
 {
+    CIOS_REG_DECLARE;
+
     // extern __shared__ unsigned int sharedmem[];
     register int i;
     register int tid = threadIdx.x + threadDimx * threadIdx.y;
@@ -205,7 +207,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]);
-    loadx(global_x_base);
+    CIOS_LOADX(global_x_base);
     a_2 = global_a_base[2];
     a_1 = global_a_base[1];
     a_0 = global_a_base[0];
@@ -267,7 +269,7 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
 	    asm("subc.u32 %0, %1, %2;" : "=r"(b_2) : "r"(b_2), "r"(order_2));
 	}
 		
-        modmul(multtype);
+        CIOS_MODMUL(multtype);
         //memcpy(ts[tid].x, ts[tid].z, (WORDS + 1) * sizeof(unsigned int));
 	/*
         if (_gt(shared_x_base)) {
@@ -281,10 +283,13 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
 	// Check for a distinguished point
 	asm("mov.u32 %0, $xr0;" : "=r" (xlow));
 	if ((xlow & 0x000003ff) == 0 || !(order_2|order_1) && order_0 < (1<<20)) {
-	    DPstreamWrite(a_0,a_1,a_2,b_0,b_1,b_2);
+	    unsigned int *ourbuffer = DPstreamWrite(a_0,a_1,a_2,b_0,b_1,b_2);
+	    if (ourbuffer) {
+		CIOS_WRITE_DP(ourbuffer, a_0, a_1, a_2, b_0, b_1, b_2);
+	    }
 	}
     }
-    savex(global_x_base);
+    CIOS_SAVEX(global_x_base);
     global_a_base[0] = a_0;
     global_a_base[1] = a_1;
     global_a_base[2] = a_2;
@@ -296,7 +301,6 @@ __global__ void cudaMulmod(GlobalThreadState *global_ts,
     // cuPrintf("d_B = %08X\n", (unsigned int)b);
     // cuPrintf("d_a = %08X%08X\n", global_a_base[1], global_a_base[0]);
     // cuPrintf("d_b = %08X%08X\n", global_b_base[1], global_b_base[0]);
-    if (!(order_0|order_1|order_2)) evilhack();
 }
 
 int nthreads = 25600;