cuda.ptx: some unit tests

diff --git a/extra/cuda/ptx/ptx-tests.factor b/extra/cuda/ptx/ptx-tests.factor
new file mode 100644 (file)
index 0000000..877bc82
--- /dev/null
+++ b/extra/cuda/ptx/ptx-tests.factor
@@ -0,0 +1,114 @@
+USING: cuda.ptx tools.test ;
+IN: cuda.ptx.tests
+
+[ """  .version 2.0
+       .target sm_20
+""" ] [
+    T{ ptx
+        { version "2.0" }
+        { target T{ ptx-target { arch sm_20 } } }
+    } ptx>string
+] unit-test
+
+[ """  .version 2.0
+       .target sm_20, .texmode_independent
+""" ] [
+    T{ ptx
+        { version "2.0" }
+        { target T{ ptx-target { arch sm_20 } { texmode .texmode_independent } } }
+    } ptx>string
+] unit-test
+
+[ """  .version 2.0
+       .target sm_11, map_f64_to_f32
+""" ] [
+    T{ ptx
+        { version "2.0" }
+        { target T{ ptx-target
+            { arch sm_11 }
+            { map_f64_to_f32? t }
+        } }
+    } ptx>string
+] unit-test
+
+[ """  .version 2.0
+       .target sm_11, map_f64_to_f32, .texmode_independent
+""" ] [
+    T{ ptx
+        { version "2.0" }
+        { target T{ ptx-target
+            { arch sm_11 }
+            { map_f64_to_f32? t }
+            { texmode .texmode_independent }
+        } }
+    } ptx>string
+] unit-test
+
+[ """  .version 2.0
+       .target sm_20
+       .global .f32 foo[9000];
+       .extern .align 16 .shared .v4.f32 bar[];
+       .func (.reg .f32 sum) zap (.reg .f32 a, .reg .f32 b)
+       {
+       add.rn.f32 sum, a, b;
+       ret;
+       }
+       .func frob (.align 8 .param .u64 in, .align 8 .param .u64 out, .align 8 .param .u64 len)
+       {
+       ret;
+       }
+       .func twib
+       {
+       ret;
+       }
+""" ] [
+    T{ ptx
+        { version "2.0" }
+        { target T{ ptx-target { arch sm_20 } } }
+        { body {
+            T{ ptx-variable
+                { storage-space .global }
+                { type .f32 }
+                { name "foo" }
+                { dim 9000 }
+            }
+            T{ ptx-variable
+                { extern? t }
+                { align 16 }
+                { storage-space .shared }
+                { type T{ .v4 f .f32 } }
+                { name "bar" }
+                { dim 0 }
+            }
+            T{ ptx-func
+                { return T{ ptx-variable { storage-space .reg } { type .f32 } { name "sum" } } }
+                { name "zap" }
+                { params {
+                    T{ ptx-variable { storage-space .reg } { type .f32 } { name "a" } }
+                    T{ ptx-variable { storage-space .reg } { type .f32 } { name "b" } }
+                } }
+                { body {
+                    T{ add { round .rn } { type .f32 } { dest "sum" } { a "a" } { b "b" } }
+                    T{ ret }
+                } }
+            }
+            T{ ptx-func
+                { name "frob" }
+                { params {
+                    T{ ptx-variable { align 8 } { storage-space .param } { type .u64 } { name "in" } }
+                    T{ ptx-variable { align 8 } { storage-space .param } { type .u64 } { name "out" } }
+                    T{ ptx-variable { align 8 } { storage-space .param } { type .u64 } { name "len" } }
+                } }
+                { body {
+                    T{ ret }
+                } }
+            }
+            T{ ptx-func
+                { name "twib" }
+                { body {
+                    T{ ret }
+                } }
+            }
+        } }
+    } ptx>string
+] unit-test

diff --git a/extra/cuda/ptx/ptx.factor b/extra/cuda/ptx/ptx.factor
index 8d4925d55fe29612d8166a6212012e3bcbfb62c5..8a30659640727c203dc7882942d82479261e53af 100644 (file)
--- a/extra/cuda/ptx/ptx.factor
+++ b/extra/cuda/ptx/ptx.factor
@@ -1,6 +1,6 @@
 ! (c)2010 Joe Groff bsd license
-USING: accessors arrays combinators io kernel math math.parser
-roles sequences strings variants words ;
+USING: accessors arrays combinators io io.streams.string kernel
+math math.parser roles sequences strings variants words ;
 FROM: roles => TUPLE: ;
 IN: cuda.ptx
 
@@ -62,6 +62,7 @@ TUPLE: ptx-variable
     { parameter ?integer }
     { dim dim }
     { initializer ?string } ;
+UNION: ?ptx-variable POSTPONE: f ptx-variable ;
 
 TUPLE: ptx-predicate
     { negated? boolean }
@@ -79,7 +80,7 @@ TUPLE: ptx-entry
     body ;
 
 TUPLE: ptx-func < ptx-entry
-    { return ptx-variable } ;
+    { return ?ptx-variable } ;
 
 TUPLE: ptx-directive ;
 
@@ -331,15 +332,23 @@ TUPLE: xor       < ptx-3op-instruction ;
 GENERIC: ptx-element-label ( elt -- label )
 M: object ptx-element-label  drop f ;
 
+GENERIC: ptx-semicolon? ( elt -- ? )
+M: object ptx-semicolon? drop t ;
+M: ptx-target ptx-semicolon? drop f ;
+M: ptx-entry ptx-semicolon? drop f ;
+M: ptx-func ptx-semicolon? drop f ;
+M: .file ptx-semicolon? drop f ;
+M: .loc ptx-semicolon? drop f ;
+
 GENERIC: (write-ptx-element) ( elt -- )
 
 : write-ptx-element ( elt -- )
     dup ptx-element-label [ write ":" write ] when*
-    "\t" write (write-ptx-element) 
-    ";" print ;
+    "\t" write dup (write-ptx-element) 
+    ptx-semicolon? [ ";" print ] [ nl ] if ;
 
 : write-ptx ( ptx -- )
-    "\t.version " write dup version>> write ";" print
+    "\t.version " write dup version>> print
     dup target>> write-ptx-element
     body>> [ write-ptx-element ] each ;
 
@@ -399,9 +408,9 @@ M: ptx-variable (write-ptx-element)
     "\t}" write ;
 
 : write-entry ( entry -- )
-    dup name>> write " " write
-    dup params>> [ write-params ] when* nl
-    dup directives>> [ (write-ptx-element) ] each nl
+    dup name>> write
+    dup params>> [  " " write write-params ] when* nl
+    dup directives>> [ (write-ptx-element) nl ] each
     dup body>> write-body
     drop ;
 
@@ -754,5 +763,8 @@ M: vote (write-ptx-element)
     dup mode>> (write-ptx-element)
     write-2op ;
 M: xor (write-ptx-element)
-    "or" write-insn
+    "xor" write-insn
     write-3op ;
+
+: ptx>string ( ptx -- string )
+    [ write-ptx ] with-string-writer ;