From 23cf6413dc9a34c4f0f0d381bb6d0e05f428e098 Mon Sep 17 00:00:00 2001 From: Joe Groff Date: Sun, 18 Apr 2010 11:54:22 -0700 Subject: [PATCH] cuda.ptx: some unit tests --- extra/cuda/ptx/ptx-tests.factor | 114 ++++++++++++++++++++++++++++++++ extra/cuda/ptx/ptx.factor | 32 ++++++--- 2 files changed, 136 insertions(+), 10 deletions(-) create mode 100644 extra/cuda/ptx/ptx-tests.factor diff --git a/extra/cuda/ptx/ptx-tests.factor b/extra/cuda/ptx/ptx-tests.factor new file mode 100644 index 0000000000..877bc82811 --- /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 8d4925d55f..8a30659640 100644 --- 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 ;