cuda.ptx: better representation of operands

diff --git a/extra/cuda/ptx/ptx-tests.factor b/extra/cuda/ptx/ptx-tests.factor
index 28391a5f58dbf12c030fea8626e5c4d387e720d8..1ba7ecfcc8c1ae96a400e9546a16229fb7b3de8b 100644 (file)
--- a/extra/cuda/ptx/ptx-tests.factor
+++ b/extra/cuda/ptx/ptx-tests.factor
@@ -1,4 +1,4 @@
-USING: cuda.ptx tools.test ;
+USING: cuda.ptx io.streams.string tools.test ;
 IN: cuda.ptx.tests
 
 [ """  .version 2.0
@@ -113,6 +113,17 @@ IN: cuda.ptx.tests
     } ptx>string
 ] unit-test
 
+[ "a" ] [ [ "a" write-ptx-operand ] with-string-writer ] unit-test
+[ "2" ] [ [ 2 write-ptx-operand ] with-string-writer ] unit-test
+[ "0d4000000000000000" ] [ [ 2.0 write-ptx-operand ] with-string-writer ] unit-test
+[ "!a" ] [ [ T{ ptx-negation f "a" } write-ptx-operand ] with-string-writer ] unit-test
+[ "{a, b, c, d}" ] [ [ T{ ptx-vector f { "a" "b" "c" "d" } } write-ptx-operand ] with-string-writer ] unit-test
+[ "[a]" ] [ [ T{ ptx-indirect f "a" 0 } write-ptx-operand ] with-string-writer ] unit-test
+[ "[a+1]" ] [ [ T{ ptx-indirect f "a" 1 } write-ptx-operand ] with-string-writer ] unit-test
+[ "[a-1]" ] [ [ T{ ptx-indirect f "a" -1 } write-ptx-operand ] with-string-writer ] unit-test
+[ "a[1]" ] [ [ T{ ptx-element f "a" 1 } write-ptx-operand ] with-string-writer ] unit-test
+[ "{a, b[2], 3, 0d4000000000000000}" ] [ [ T{ ptx-vector f { "a" T{ ptx-element f "b" 2 } 3 2.0 } } write-ptx-operand ] with-string-writer ] unit-test
+
 [ """  .version 2.0
        .target sm_20
        abs.s32 a, b;
@@ -127,11 +138,11 @@ foo:      abs.s32 a, b;
         { body {
             T{ abs { type .s32 } { dest "a" } { a "b" } }
             T{ abs
-                { predicate T{ ptx-predicate { variable "p" } } }
+                { predicate "p" }
                 { type .s32 } { dest "a" } { a "b" }
             }
             T{ abs
-                { predicate T{ ptx-predicate { negated? t } { variable "p" } } }
+                { predicate T{ ptx-negation f "p" } }
                 { type .s32 } { dest "a" } { a "b" }
             }
             T{ abs
@@ -206,9 +217,9 @@ foo:        abs.s32 a, b;
         { version "2.0" }
         { target T{ ptx-target { arch sm_20 } } }
         { body {
-            T{ atom { op .and } { type .u32 } { dest "a" } { a "[b]" } { b "c" } }
-            T{ atom { storage-space .global } { op .or } { type .u32 } { dest "a" } { a "[b]" } { b "c" } }
-            T{ atom { storage-space .shared } { op .cas } { type .u32 } { dest "a" } { a "[b]" } { b "c" } { c "d" } }
+            T{ atom { op .and } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } { b "c" } }
+            T{ atom { storage-space .global } { op .or } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } { b "c" } }
+            T{ atom { storage-space .shared } { op .cas } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } { b "c" } { c "d" } }
 
         } }
     } ptx>string
@@ -229,8 +240,8 @@ foo:        abs.s32 a, b;
         { body {
             T{ bar.arrive { a "a" } { b "b" } }
             T{ bar.red { op .popc } { type .u32 } { dest "a" } { a "b" } { c "d" } }
-            T{ bar.red { op .popc } { type .u32 } { dest "a" } { a "b" } { c "!d" } }
-            T{ bar.red { op .popc } { type .u32 } { dest "a" } { a "b" } { b "c" } { c "!d" } }
+            T{ bar.red { op .popc } { type .u32 } { dest "a" } { a "b" } { c T{ ptx-negation f "d" } } }
+            T{ bar.red { op .popc } { type .u32 } { dest "a" } { a "b" } { b "c" } { c T{ ptx-negation f "d" } } }
             T{ bar.sync { a "a" } }
             T{ bar.sync { a "a" } { b "b" } }
         } }
@@ -327,6 +338,7 @@ foo:        abs.s32 a, b;
        call (a), foo, (b);
        call (a), foo, (b, c);
        call (a), foo, (b, c, d);
+       call (a[2]), foo, (b, c, d[3]);
        call foo, (b, c, d);
 """ ] [
     T{ ptx
@@ -339,6 +351,7 @@ foo:        abs.s32 a, b;
             T{ call { return "a" } { target "foo" } { params { "b" } } }
             T{ call { return "a" } { target "foo" } { params { "b" "c" } } }
             T{ call { return "a" } { target "foo" } { params { "b" "c" "d" } } }
+            T{ call { return T{ ptx-element f "a" 2 } } { target "foo" } { params { "b" "c" T{ ptx-element f "d" 3 } } } }
             T{ call { target "foo" } { params { "b" "c" "d" } } }
         } }
     } ptx>string
@@ -549,13 +562,13 @@ foo:      abs.s32 a, b;
         { version "2.0" }
         { target T{ ptx-target { arch sm_20 } } }
         { body {
-            T{ ld { type .u32 } { dest "a" } { a "[b]" } }
-            T{ ld { type T{ .v2 { of .u32 } } } { dest "a" } { a "[b]" } }
-            T{ ld { type T{ .v4 { of .u32 } } } { dest "a" } { a "[b]" } }
-            T{ ld { type T{ .v4 { of .u32 } } } { dest "{a, b, c, d}" } { a "[e]" } }
+            T{ ld { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } }
+            T{ ld { type T{ .v2 { of .u32 } } } { dest "a" } { a T{ ptx-indirect f "b" } } }
+            T{ ld { type T{ .v4 { of .u32 } } } { dest "a" } { a T{ ptx-indirect f "b" } } }
+            T{ ld { type T{ .v4 { of .u32 } } } { dest T{ ptx-vector f { "a" "b" "c" "d" } } } { a "[e]" } }
             T{ ld { cache-op .lu } { type .u32 } { dest "a" } { a "[b]" } }
-            T{ ld { storage-space T{ .const } } { cache-op .lu } { type .u32 } { dest "a" } { a "[b]" } }
-            T{ ld { volatile? t } { storage-space T{ .const { bank 5 } } } { type .u32 } { dest "a" } { a "[b]" } }
+            T{ ld { storage-space T{ .const } } { cache-op .lu } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } }
+            T{ ld { volatile? t } { storage-space T{ .const { bank 5 } } } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } }
         } }
     } ptx>string
 ] unit-test
@@ -574,13 +587,13 @@ foo:      abs.s32 a, b;
         { version "2.0" }
         { target T{ ptx-target { arch sm_20 } } }
         { body {
-            T{ ldu { type .u32 } { dest "a" } { a "[b]" } }
-            T{ ldu { type T{ .v2 { of .u32 } } } { dest "a" } { a "[b]" } }
-            T{ ldu { type T{ .v4 { of .u32 } } } { dest "a" } { a "[b]" } }
-            T{ ldu { type T{ .v4 { of .u32 } } } { dest "{a, b, c, d}" } { a "[e]" } }
+            T{ ldu { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } }
+            T{ ldu { type T{ .v2 { of .u32 } } } { dest "a" } { a T{ ptx-indirect f "b" } } }
+            T{ ldu { type T{ .v4 { of .u32 } } } { dest "a" } { a T{ ptx-indirect f "b" } } }
+            T{ ldu { type T{ .v4 { of .u32 } } } { dest T{ ptx-vector f { "a" "b" "c" "d" } } } { a "[e]" } }
             T{ ldu { cache-op .lu } { type .u32 } { dest "a" } { a "[b]" } }
-            T{ ldu { storage-space T{ .const } } { cache-op .lu } { type .u32 } { dest "a" } { a "[b]" } }
-            T{ ldu { volatile? t } { storage-space T{ .const { bank 5 } } } { type .u32 } { dest "a" } { a "[b]" } }
+            T{ ldu { storage-space T{ .const } } { cache-op .lu } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } }
+            T{ ldu { volatile? t } { storage-space T{ .const { bank 5 } } } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } }
         } }
     } ptx>string
 ] unit-test
@@ -723,9 +736,9 @@ foo:        abs.s32 a, b;
         { version "2.0" }
         { target T{ ptx-target { arch sm_20 } } }
         { body {
-            T{ prefetch { level .L1 } { a "[a]" } }
-            T{ prefetch { storage-space .local } { level .L2 } { a "[a]" } }
-            T{ prefetchu { level .L1 } { a "[a]" } }
+            T{ prefetch { level .L1 } { a T{ ptx-indirect f "a" } } }
+            T{ prefetch { storage-space .local } { level .L2 } { a T{ ptx-indirect f "a" } } }
+            T{ prefetchu { level .L1 } { a T{ ptx-indirect f "a" } } }
         } }
     } ptx>string
 ] unit-test
@@ -781,8 +794,8 @@ foo:        abs.s32 a, b;
         { version "2.0" }
         { target T{ ptx-target { arch sm_20 } } }
         { body {
-            T{ red { op .and } { type .u32 } { dest "[a]" } { a "b" } }
-            T{ red { storage-space .global } { op .and } { type .u32 } { dest "[a]" } { a "b" } }
+            T{ red { op .and } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } }
+            T{ red { storage-space .global } { op .and } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } }
         } }
     } ptx>string
 ] unit-test
@@ -861,7 +874,7 @@ foo:        abs.s32 a, b;
             T{ set { cmp-op .gt } { dest-type .u32 } { type .s32 } { dest "a" } { a "b" } { b "c" } }
             T{ set { cmp-op .gt } { ftz? t } { dest-type .u32 } { type .f32 } { dest "a" } { a "b" } { b "c" } }
             T{ set { cmp-op .gt } { bool-op .and } { ftz? t } { dest-type .u32 } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
-            T{ set { cmp-op .gt } { bool-op .and } { ftz? t } { dest-type .u32 } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "!d" } }
+            T{ set { cmp-op .gt } { bool-op .and } { ftz? t } { dest-type .u32 } { type .f32 } { dest "a" } { a "b" } { b "c" } { c T{ ptx-negation f "d" } } }
         } }
     } ptx>string
 ] unit-test
@@ -982,13 +995,13 @@ foo:      abs.s32 a, b;
         { version "2.0" }
         { target T{ ptx-target { arch sm_20 } } }
         { body {
-            T{ st { type .u32 } { dest "[a]" } { a "b" } }
-            T{ st { type T{ .v2 { of .u32 } } } { dest "[a]" } { a "b" } }
-            T{ st { type T{ .v4 { of .u32 } } } { dest "[a]" } { a "b" } }
-            T{ st { type T{ .v4 { of .u32 } } } { dest "[a]" } { a "{b, c, d, e}" } }
-            T{ st { cache-op .lu } { type .u32 } { dest "[a]" } { a "b" } }
-            T{ st { storage-space .local } { cache-op .lu } { type .u32 } { dest "[a]" } { a "b" } }
-            T{ st { volatile? t } { storage-space .local } { type .u32 } { dest "[a]" } { a "b" } }
+            T{ st { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } }
+            T{ st { type T{ .v2 { of .u32 } } } { dest T{ ptx-indirect f "a" } } { a "b" } }
+            T{ st { type T{ .v4 { of .u32 } } } { dest T{ ptx-indirect f "a" } } { a "b" } }
+            T{ st { type T{ .v4 { of .u32 } } } { dest T{ ptx-indirect f "a" } } { a T{ ptx-vector f { "b" "c" "d" "e" } } } }
+            T{ st { cache-op .lu } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } }
+            T{ st { storage-space .local } { cache-op .lu } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } }
+            T{ st { volatile? t } { storage-space .local } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } }
         } }
     } ptx>string
 ] unit-test

diff --git a/extra/cuda/ptx/ptx.factor b/extra/cuda/ptx/ptx.factor
index 4618f8b5b6197a0eadf07da0e4c06e851bd2d464..49a53d7fbf9bec4529893105b00cd9ea841c44b8 100644 (file)
--- a/extra/cuda/ptx/ptx.factor
+++ b/extra/cuda/ptx/ptx.factor
@@ -64,14 +64,30 @@ TUPLE: ptx-variable
     { initializer ?string } ;
 UNION: ?ptx-variable POSTPONE: f ptx-variable ;
 
-TUPLE: ptx-predicate
-    { negated? boolean }
-    { variable string } ; 
-UNION: ?ptx-predicate POSTPONE: f ptx-predicate ;
+TUPLE: ptx-negation
+    { var string } ; 
+
+TUPLE: ptx-vector
+    elements ;
+
+TUPLE: ptx-element
+    { var string }
+    { index integer } ;
+
+UNION: ptx-var
+    string ptx-element ;
+
+TUPLE: ptx-indirect
+    { base ptx-var }
+    { offset integer } ;
+
+UNION: ptx-operand
+    integer float ptx-var ptx-negation ptx-vector ptx-indirect ;
+UNION: ?ptx-operand POSTPONE: f ptx-operand ;
 
 TUPLE: ptx-instruction
     { label ?string }
-    { predicate ?ptx-predicate } ;
+    { predicate ?ptx-operand } ;
 
 TUPLE: ptx-entry
     { name string }
@@ -112,25 +128,25 @@ UNION: ?ptx-rounding-mode POSTPONE: f ptx-rounding-mode ;
 
 TUPLE: ptx-typed-instruction < ptx-instruction
     { type ptx-type }
-    { dest string } ;
+    { dest ptx-operand } ;
 
 TUPLE: ptx-2op-instruction < ptx-typed-instruction
-    { a string } ;
+    { a ptx-operand } ;
 
 TUPLE: ptx-3op-instruction < ptx-typed-instruction
-    { a string }
-    { b string } ;
+    { a ptx-operand }
+    { b ptx-operand } ;
 
 TUPLE: ptx-4op-instruction < ptx-typed-instruction
-    { a string }
-    { b string }
-    { c string } ;
+    { a ptx-operand }
+    { b ptx-operand }
+    { c ptx-operand } ;
 
 TUPLE: ptx-5op-instruction < ptx-typed-instruction
-    { a string }
-    { b string }
-    { c string }
-    { d string } ;
+    { a ptx-operand }
+    { b ptx-operand }
+    { c ptx-operand }
+    { d ptx-operand } ;
 
 TUPLE: ptx-addsub-instruction < ptx-3op-instruction
     { sat? boolean }
@@ -181,7 +197,7 @@ INSTANCE: .hi ptx-cmp-op
 TUPLE: ptx-set-instruction < ptx-3op-instruction
     { cmp-op ptx-cmp-op }
     { bool-op ?ptx-op }
-    { c ?string }
+    { c ?ptx-operand }
     { ftz? boolean } ;
 
 VARIANT: ptx-cache-op
@@ -216,17 +232,17 @@ TUPLE: and       < ptx-3op-instruction ;
 TUPLE: atom      < ptx-3op-instruction
     { storage-space ?ptx-storage-space }
     { op ptx-op }
-    { c ?string } ;
+    { c ?ptx-operand } ;
 TUPLE: bar.arrive < ptx-instruction
-    { a string }
-    { b string } ;
+    { a ptx-operand }
+    { b ptx-operand } ;
 TUPLE: bar.red   < ptx-2op-instruction
     { op ptx-op }
-    { b ?string }
-    { c string } ;
+    { b ?ptx-operand }
+    { c ptx-operand } ;
 TUPLE: bar.sync  < ptx-instruction
-    { a string }
-    { b ?string } ;
+    { a ptx-operand }
+    { b ?ptx-operand } ;
 TUPLE: bfe       < ptx-4op-instruction ;
 TUPLE: bfi       < ptx-5op-instruction ;
 TUPLE: bfind     < ptx-2op-instruction
@@ -235,7 +251,7 @@ TUPLE: bra       < ptx-branch-instruction ;
 TUPLE: brev      < ptx-2op-instruction ;
 TUPLE: brkpt     < ptx-instruction ;
 TUPLE: call      < ptx-branch-instruction
-    { return ?string }
+    { return ?ptx-operand }
     params ;
 TUPLE: clz       < ptx-2op-instruction ;
 TUPLE: cnot      < ptx-2op-instruction ;
@@ -255,8 +271,8 @@ TUPLE: exit      < ptx-instruction ;
 TUPLE: fma       <{ ptx-mad-instruction ptx-float-env } ;
 TUPLE: isspacep  < ptx-instruction
     { storage-space ptx-storage-space }
-    { dest string }
-    { a string } ;
+    { dest ptx-operand }
+    { a ptx-operand } ;
 TUPLE: ld        < ptx-ldst-instruction ;
 TUPLE: ldu       < ptx-ldst-instruction ;
 TUPLE: lg2       <{ ptx-2op-instruction ptx-float-env } ;
@@ -273,14 +289,14 @@ TUPLE: neg       <{ ptx-2op-instruction ptx-float-ftz } ;
 TUPLE: not       < ptx-2op-instruction ;
 TUPLE: or        < ptx-3op-instruction ;
 TUPLE: pmevent   < ptx-instruction
-    { a string } ;
+    { a ptx-operand } ;
 TUPLE: popc      < ptx-2op-instruction ;
 TUPLE: prefetch  < ptx-instruction
-    { a string }
+    { a ptx-operand }
     { storage-space ?ptx-storage-space }
     { level ptx-cache-level } ;
 TUPLE: prefetchu < ptx-instruction
-    { a string }
+    { a ptx-operand }
     { level ptx-cache-level } ;
 TUPLE: prmt      < ptx-4op-instruction
     { mode ?ptx-prmt-mode } ;
@@ -296,7 +312,7 @@ TUPLE: selp      < ptx-4op-instruction ;
 TUPLE: set       < ptx-set-instruction
     { dest-type ptx-type } ;
 TUPLE: setp      < ptx-set-instruction
-    { |dest ?string } ;
+    { |dest ?ptx-operand } ;
 TUPLE: shl       < ptx-3op-instruction ;
 TUPLE: shr       < ptx-3op-instruction ;
 TUPLE: sin       <{ ptx-2op-instruction ptx-float-env } ;
@@ -340,6 +356,27 @@ M: ptx-func ptx-semicolon? drop f ;
 M: .file ptx-semicolon? drop f ;
 M: .loc ptx-semicolon? drop f ;
 
+GENERIC: write-ptx-operand ( operand -- )
+
+M: string write-ptx-operand write ;
+M: integer write-ptx-operand number>string write ;
+M: float write-ptx-operand "0d" write double>bits >hex 16 CHAR: 0 pad-head write ;
+M: ptx-negation write-ptx-operand "!" write var>> write ;
+M: ptx-vector write-ptx-operand
+    "{" write
+    elements>> [ ", " write ] [ write-ptx-operand ] interleave
+    "}" write ;
+M: ptx-element write-ptx-operand dup var>> write "[" write index>> number>string write "]" write ;
+M: ptx-indirect write-ptx-operand
+    "[" write
+    dup base>> write-ptx-operand
+    offset>> {
+        { [ dup zero? ] [ drop ] }
+        { [ dup 0 < ] [ number>string write ] }
+        [ "+" write number>string write ]
+    } cond
+    "]" write ;
+
 GENERIC: (write-ptx-element) ( elt -- )
 
 : write-ptx-element ( elt -- )
@@ -376,7 +413,7 @@ M: ptx-target (write-ptx-element)
     [ arch>> [ name>> ] [ f ] if* ]
     [ map_f64_to_f32?>> [ "map_f64_to_f32" ] [ f ] if ]
     [ texmode>> [ name>> ] [ f ] if* ] tri
-    3array sift ", " join write ;
+    3array sift [ ", " write ] [ write ] interleave ;
 
 : write-ptx-dim ( dim -- )
     {
@@ -435,7 +472,7 @@ M: .maxnreg (write-ptx-element)
     ".maxnreg " write n>> number>string write ;
 M: .maxntid (write-ptx-element)
     ".maxntid " write
-    dup sequence? [ [ number>string ] map ", " join write ] [ number>string write ] if ;
+    dup sequence? [ [ ", " write ] [ number>string write ] interleave ] [ number>string write ] if ;
 M: .pragma (write-ptx-element)
     ".pragma \"" write pragma>> write "\"" write ;
 
@@ -444,28 +481,28 @@ M: ptx-instruction ptx-element-label
 
 : write-insn ( insn name -- insn )
     over predicate>>
-    [ "@" write dup negated?>> [ "!" write ] when variable>> write " " write ] when*
+    [ "@" write write-ptx-operand " " write ] when*
     write ;
 
 : write-2op ( insn -- )
     dup type>> (write-ptx-element) " " write
-    dup dest>> write ", " write
-    dup a>> write
+    dup dest>> write-ptx-operand ", " write
+    dup a>> write-ptx-operand
     drop ;
 
 : write-3op ( insn -- )
     dup write-2op ", " write
-    dup b>> write
+    dup b>> write-ptx-operand
     drop ;
 
 : write-4op ( insn -- )
     dup write-3op ", " write
-    dup c>> write
+    dup c>> write-ptx-operand
     drop ;
 
 : write-5op ( insn -- )
     dup write-4op ", " write
-    dup d>> write
+    dup d>> write-ptx-operand
     drop ;
 
 : write-ftz ( insn -- )
@@ -534,22 +571,22 @@ M: atom (write-ptx-element)
     dup storage-space>> (write-ptx-element)
     dup op>> (write-ptx-element)
     dup write-3op
-    c>> [ ", " write write ] when* ;
+    c>> [ ", " write write-ptx-operand ] when* ;
 M: bar.arrive (write-ptx-element)
     "bar.arrive " write-insn
-    dup a>> write ", " write
-    dup b>> write
+    dup a>> write-ptx-operand ", " write
+    dup b>> write-ptx-operand
     drop ;
 M: bar.red (write-ptx-element)
     "bar.red" write-insn
     dup op>> (write-ptx-element)
     dup write-2op
-    dup b>> [ ", " write write ] when*
-    ", " write c>> write ;
+    dup b>> [ ", " write write-ptx-operand ] when*
+    ", " write c>> write-ptx-operand ;
 M: bar.sync (write-ptx-element)
     "bar.sync " write-insn
-    dup a>> write
-    dup b>> [ ", " write write ] when*
+    dup a>> write-ptx-operand
+    dup b>> [ ", " write write-ptx-operand ] when*
     drop ;
 M: bfe (write-ptx-element)
     "bfe" write-insn
@@ -573,9 +610,9 @@ M: brkpt (write-ptx-element)
 M: call (write-ptx-element)
     "call" write-insn
     dup write-uni " " write
-    dup return>> [ "(" write write "), " write ] when*
+    dup return>> [ "(" write write-ptx-operand "), " write ] when*
     dup target>> write
-    dup params>> [ ", (" write ", " join write ")" write ] unless-empty
+    dup params>> [ ", (" write [ ", " write ] [ write-ptx-operand ] interleave ")" write ] unless-empty
     drop ;
 M: clz (write-ptx-element)
     "clz" write-insn
@@ -619,7 +656,7 @@ M: isspacep (write-ptx-element)
     "isspacep" write-insn
     dup storage-space>> (write-ptx-element)
     " " write
-    dup dest>> write ", " write a>> write ;
+    dup dest>> write-ptx-operand ", " write a>> write-ptx-operand ;
 M: ld (write-ptx-element)
     "ld" write-insn
     write-ldst ;
@@ -679,19 +716,19 @@ M: prefetch (write-ptx-element)
     "prefetch" write-insn
     dup storage-space>> (write-ptx-element)
     dup level>> (write-ptx-element)
-    " " write a>> write ;
+    " " write a>> write-ptx-operand ;
 M: prefetchu (write-ptx-element)
     "prefetchu" write-insn
     dup level>> (write-ptx-element)
-    " " write a>> write ;
+    " " write a>> write-ptx-operand ;
 M: prmt (write-ptx-element)
     "prmt" write-insn
     dup type>> (write-ptx-element)
     dup mode>> (write-ptx-element) " " write
-    dup dest>> write ", " write
-    dup a>> write ", " write
-    dup b>> write ", " write
-    dup c>> write
+    dup dest>> write-ptx-operand ", " write
+    dup a>> write-ptx-operand ", " write
+    dup b>> write-ptx-operand ", " write
+    dup c>> write-ptx-operand
     drop ;
 M: rcp (write-ptx-element)
     "rcp" write-insn
@@ -722,16 +759,16 @@ M: set (write-ptx-element)
     dup write-set
     dup dest-type>> (write-ptx-element)
     dup write-3op
-    c>> [ ", " write write ] when* ;
+    c>> [ ", " write write-ptx-operand ] when* ;
 M: setp (write-ptx-element)
     "setp" write-insn
     dup write-set
     dup type>> (write-ptx-element) " " write
-    dup dest>> write
-    dup |dest>> [ "|" write write ] when* ", " write
-    dup a>> write ", " write
-    dup b>> write
-    c>> [ ", " write write ] when* ;
+    dup dest>> write-ptx-operand
+    dup |dest>> [ "|" write write-ptx-operand ] when* ", " write
+    dup a>> write-ptx-operand ", " write
+    dup b>> write-ptx-operand
+    c>> [ ", " write write-ptx-operand ] when* ;
 M: shl (write-ptx-element)
     "shl" write-insn
     write-3op ;