Remove symbolic operand references from inline asm

gencios_reg_20

@@ -57,7 +57,7 @@ EOA
 
 for ($i=0;$i<$WORDS;++$i) {
     my $xoff = $i*4;
-    print "\t\"ld.global.u32 \$xr$i, [%[x]+$xoff];\\n\"\n";
+    print "\t\"ld.global.u32 \$xr$i, [%0+$xoff];\\n\"\n";
 }
 print "\t\"// Load c_rho into registers\\n\"\n";
 my ($xoff, $xoffm1);
@@ -66,7 +66,7 @@ for($i=0;$i<$WORDS;++$i) {
     $croff = $i*4;
     print "\t\"ld.const.u32 \$crho$i, [c_rho+$croff];\\n\"\n";
 }
-print "\t\"mov.u32 \$crp, %[crp];\\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;
@@ -80,7 +80,7 @@ for($i=0;$i<$WORDS;++$i) {
 
 
 print <<'EOA';
-    : : [x] "l" (x), [crp] "r" (c_rho_prime));
+    : : "l" (x), "r" (c_rho_prime));
 }
 
 __device__ inline void savex(unsigned int *x)
@@ -90,11 +90,11 @@ EOA
 
 for ($i=0;$i<$WORDS;++$i) {
     my $xoff = $i*4;
-    print "\t\"st.global.u32 [%[x]+$xoff], \$xr$i;\\n\"\n";
+    print "\t\"st.global.u32 [%0+$xoff], \$xr$i;\\n\"\n";
 }
 
 print <<'EOA';
-    : : [x] "l" (x));
+    : : "l" (x));
 }
 EOA
 
@@ -109,10 +109,10 @@ EOA
 print "    \".reg.pred \$retaddr<", $WORDS-2, ">;\\n\"\n";
 print <<'EOA';
     "mov.u32 $m, 0;\n"
-    "setp.eq.u32 $mtp0, %[tp], $m;\n"
+    "setp.eq.u32 $mtp0, %0, $m;\n"
     "mov.u32 $m, 1;\n"
-    "setp.eq.u32 $mtp1, %[tp], $m;\n"
-    : : [tp] "r" (multtype));
+    "setp.eq.u32 $mtp1, %0, $m;\n"
+    : : "r" (multtype));
 EOA
 
 print "    asm(\n";