diff --git a/basis/compiler/tree/propagation/branches/branches.factor b/basis/compiler/tree/propagation/branches/branches.factor
index 28f34cb425..ef9e4e8f0b 100644
--- a/basis/compiler/tree/propagation/branches/branches.factor
+++ b/basis/compiler/tree/propagation/branches/branches.factor
@@ -1,8 +1,8 @@
-! Copyright (C) 2008, 2009 Slava Pestov.
+! Copyright (C) 2008, 2010 Slava Pestov.
 ! See http://factorcode.org/license.txt for BSD license.
-USING: fry kernel sequences assocs accessors namespaces
+USING: fry kernel sequences assocs accessors
 math.intervals arrays classes.algebra combinators columns
-stack-checker.branches locals math
+stack-checker.branches locals math namespaces
 compiler.utilities
 compiler.tree
 compiler.tree.combinators
@@ -10,6 +10,8 @@ compiler.tree.propagation.info
 compiler.tree.propagation.nodes
 compiler.tree.propagation.simple
 compiler.tree.propagation.constraints ;
+FROM: sets => union ;
+FROM: assocs => change-at ;
 IN: compiler.tree.propagation.branches
 
 ! For conditionals, an assoc of child node # --> constraint
@@ -90,7 +92,7 @@ M: #phi propagate-before ( #phi -- )
     bi ;
 
 :: update-constraints ( new old -- )
-    new [| key value | key old [ value append ] change-at ] assoc-each ;
+    new [| key value | key old [ value union ] change-at ] assoc-each ;
 
 : include-child-constraints ( i -- )
     infer-children-data get nth constraints swap at last
diff --git a/basis/compiler/tree/propagation/constraints/constraints.factor b/basis/compiler/tree/propagation/constraints/constraints.factor
index 617352d699..f9988ba220 100644
--- a/basis/compiler/tree/propagation/constraints/constraints.factor
+++ b/basis/compiler/tree/propagation/constraints/constraints.factor
@@ -1,4 +1,4 @@
-! Copyright (C) 2008 Slava Pestov.
+! Copyright (C) 2008, 2010 Slava Pestov.
 ! See http://factorcode.org/license.txt for BSD license.
 USING: arrays assocs math math.intervals kernel accessors
 sequences namespaces classes classes.algebra
@@ -87,8 +87,11 @@ TUPLE: implication p q ;
 
 C: --> implication
 
+: maybe-add ( elt seq -- seq' )
+    2dup member? [ nip ] [ swap suffix ] if ;
+
 : assume-implication ( q p -- )
-    [ constraints get [ assoc-stack swap suffix ] 2keep last set-at ]
+    [ constraints get [ assoc-stack maybe-add ] 2keep last set-at ]
     [ satisfied? [ assume ] [ drop ] if ] 2bi ;
 
 M: implication assume*
diff --git a/basis/formatting/formatting-tests.factor b/basis/formatting/formatting-tests.factor
index 5710ceb985..35b1dfff4a 100644
--- a/basis/formatting/formatting-tests.factor
+++ b/basis/formatting/formatting-tests.factor
@@ -16,6 +16,7 @@ IN: formatting.tests
 [ t ] [ "+0010" 10 "%+05d" sprintf = ] unit-test
 [ t ] [ "123.456000" 123.456 "%f" sprintf = ] unit-test
 [ t ] [ "2.44" 2.436 "%.2f" sprintf = ] unit-test
+[ t ] [ "8.950" 8.950179003580072 "%.3f" sprintf = ] unit-test
 [ t ] [ "123.10" 123.1 "%01.2f" sprintf = ] unit-test
 [ t ] [ "1.2346" 1.23456789 "%.4f" sprintf = ] unit-test
 [ t ] [ "  1.23" 1.23456789 "%6.2f" sprintf = ] unit-test
diff --git a/basis/formatting/formatting.factor b/basis/formatting/formatting.factor
index ec3c9f1d8e..5abcb12916 100644
--- a/basis/formatting/formatting.factor
+++ b/basis/formatting/formatting.factor
@@ -3,7 +3,9 @@
 USING: accessors arrays assocs calendar combinators fry kernel
 generalizations io io.streams.string macros math math.functions
 math.parser peg.ebnf quotations sequences splitting strings
-unicode.categories unicode.case vectors combinators.smart ;
+unicode.categories unicode.case vectors combinators.smart
+present ;
+FROM: math.parser.private => format-float ;
 IN: formatting
 
 <PRIVATE
@@ -26,31 +28,15 @@ IN: formatting
 : >digits ( string -- digits )
     [ 0 ] [ string>number ] if-empty ;
 
-: pad-digits ( string digits -- string' )
-    [ "." split1 ] dip [ CHAR: 0 pad-tail ] [ head-slice ] bi "." glue ;
+: format-simple ( x digits string -- string )
+    [ [ >float ] [ number>string ] bi* "%." ] dip
+    surround format-float ;
 
-: max-digits ( n digits -- n' )
-    10^ [ * round ] keep / ; inline
+: format-scientific ( x digits -- string ) "e" format-simple ;
 
-: >exp ( x -- exp base )
-    [
-        abs 0 swap
-        [ dup [ 10.0 >= ] [ 1.0 < ] bi or ]
-        [ dup 10.0 >=
-          [ 10.0 / [ 1 + ] dip ]
-          [ 10.0 * [ 1 - ] dip ] if
-        ] while
-     ] keep 0 < [ neg ] when ;
+: format-decimal ( x digits -- string ) "f" format-simple ;
 
-: exp>string ( exp base digits -- string )
-    [ max-digits ] keep -rot
-    [
-        [ 0 < "-" "+" ? ]
-        [ abs number>string 2 CHAR: 0 pad-head ] bi
-        "e" -rot 3append
-    ]
-    [ number>string ] bi*
-    rot pad-digits prepend ;
+ERROR: unknown-printf-directive ;
 
 EBNF: parse-printf
 
@@ -73,15 +59,15 @@ digits    = (digits_)?           => [[ 6 or ]]
 fmt-%     = "%"                  => [[ [ "%" ] ]]
 fmt-c     = "c"                  => [[ [ 1string ] ]]
 fmt-C     = "C"                  => [[ [ 1string >upper ] ]]
-fmt-s     = "s"                  => [[ [ dup number? [ number>string ] when ] ]]
-fmt-S     = "S"                  => [[ [ dup number? [ number>string ] when >upper ] ]]
-fmt-d     = "d"                  => [[ [ >fixnum number>string ] ]]
-fmt-e     = digits "e"           => [[ first '[ >exp _ exp>string ] ]]
-fmt-E     = digits "E"           => [[ first '[ >exp _ exp>string >upper ] ]]
-fmt-f     = digits "f"           => [[ first dup '[ >float _ max-digits number>string _ pad-digits ] ]]
+fmt-s     = "s"                  => [[ [ present ] ]]
+fmt-S     = "S"                  => [[ [ present >upper ] ]]
+fmt-d     = "d"                  => [[ [ >integer number>string ] ]]
+fmt-e     = digits "e"           => [[ first '[ _ format-scientific ] ]]
+fmt-E     = digits "E"           => [[ first '[ _ format-scientific >upper ] ]]
+fmt-f     = digits "f"           => [[ first '[ _ format-decimal ] ]]
 fmt-x     = "x"                  => [[ [ >hex ] ]]
 fmt-X     = "X"                  => [[ [ >hex >upper ] ]]
-unknown   = (.)*                 => [[ "Unknown directive" throw ]]
+unknown   = (.)*                 => [[ unknown-printf-directive ]]
 
 strings_  = fmt-c|fmt-C|fmt-s|fmt-S
 strings   = pad width strings_   => [[ reverse compose-all ]]
diff --git a/basis/stack-checker/known-words/known-words.factor b/basis/stack-checker/known-words/known-words.factor
index 15895184df..1fa9a94677 100644
--- a/basis/stack-checker/known-words/known-words.factor
+++ b/basis/stack-checker/known-words/known-words.factor
@@ -289,7 +289,7 @@ M: bad-executable summary
 \ (dlsym) { byte-array object } { c-ptr } define-primitive
 \ (exists?) { string } { object } define-primitive
 \ (exit) { integer } { } define-primitive
-\ (float>string) { float } { byte-array } define-primitive \ (float>string) make-foldable
+\ (format-float) { float byte-array } { byte-array } define-primitive \ (format-float) make-foldable
 \ (fopen) { byte-array byte-array } { alien } define-primitive
 \ (identity-hashcode) { object } { fixnum } define-primitive
 \ (save-image) { byte-array byte-array } { } define-primitive
diff --git a/core/bootstrap/primitives.factor b/core/bootstrap/primitives.factor
index 87963848bf..c466b0c1f8 100644
--- a/core/bootstrap/primitives.factor
+++ b/core/bootstrap/primitives.factor
@@ -470,7 +470,7 @@ tuple
     { "byte-array>bignum" "math" "primitive_byte_array_to_bignum" (( x -- y )) }
     { "double>bits" "math" "primitive_double_bits" (( x -- n )) }
     { "float>bits" "math" "primitive_float_bits" (( x -- n )) }
-    { "(float>string)" "math.parser.private" "primitive_float_to_str" (( n -- str )) }
+    { "(format-float)" "math.parser.private" "primitive_format_float" (( n format -- byte-array )) }
     { "bignum*" "math.private" "primitive_bignum_multiply" (( x y -- z )) }
     { "bignum+" "math.private" "primitive_bignum_add" (( x y -- z )) }
     { "bignum-" "math.private" "primitive_bignum_subtract" (( x y -- z )) }
diff --git a/core/math/parser/parser.factor b/core/math/parser/parser.factor
index 5bb024db9d..14fd6a2983 100644
--- a/core/math/parser/parser.factor
+++ b/core/math/parser/parser.factor
@@ -1,6 +1,7 @@
 ! (c)2009 Joe Groff bsd license
-USING: accessors combinators kernel kernel.private math
-namespaces sequences sequences.private splitting strings make ;
+USING: accessors byte-arrays combinators kernel kernel.private
+math namespaces sequences sequences.private splitting strings
+make ;
 IN: math.parser
 
 : digit> ( ch -- n )
@@ -356,15 +357,15 @@ M: ratio >base
         mantissa-expt [ float>hex-value ] [ float>hex-expt ] bi*
     ] bi 3append ;
 
-: float>decimal ( n -- str )
-    (float>string)
-    [ 0 = ] trim-tail >string
+: format-float ( n format -- string )
+    0 suffix >byte-array (format-float)
+    dup [ 0 = ] find drop head >string
     fix-float ;
 
 : float>base ( n base -- str )
     {
         { 16 [ float>hex ] }
-        [ drop float>decimal ]
+        [ drop "%.16g" format-float ]
     } case ; inline
 
 PRIVATE>
diff --git a/extra/benchmark/knucleotide/knucleotide.factor b/extra/benchmark/knucleotide/knucleotide.factor
index 70fa1bb061..f03e26675e 100644
--- a/extra/benchmark/knucleotide/knucleotide.factor
+++ b/extra/benchmark/knucleotide/knucleotide.factor
@@ -1,16 +1,9 @@
-USING: kernel locals io io.files splitting strings io.encodings.ascii
-       hashtables sequences assocs math namespaces prettyprint
-       math.parser combinators arrays sorting unicode.case ;
-
+USING: ascii kernel io io.files splitting strings
+io.encodings.ascii hashtables sequences assocs math
+math.statistics namespaces prettyprint math.parser combinators
+arrays sorting formatting grouping fry ;
 IN: benchmark.knucleotide
 
-: float>string ( float places -- string )
-    swap >float number>string
-    "." split1 rot
-    over length over <
-    [ CHAR: 0 pad-tail ] 
-    [ head ] if "." glue ;
-
 : discard-lines ( -- )
     readln
     [ ">THREE" head? [ discard-lines ] unless ] when* ;
@@ -20,37 +13,25 @@ IN: benchmark.knucleotide
     ">" read-until drop
     CHAR: \n swap remove >upper ;
 
-: tally ( x exemplar -- b )
-    clone [ [ inc-at ] curry each ] keep ;
-
-: small-groups ( x n -- b )
-    swap
-    [ length swap - 1 + iota ] 2keep
-    [ [ over + ] dip subseq ] 2curry map ;
-
 : handle-table ( inputs n -- )
-    small-groups
-    [ length ] keep
-    H{ } tally >alist
-    sort-values reverse
-    [
-      dup first write bl
-      second 100 * over / 3 float>string print
-    ] each
-    drop ;
+    clump
+    [ histogram >alist sort-values reverse ] [ length ] bi
+    '[
+        [ first write bl ]
+        [ second 100 * _ /f "%.3f" printf nl ] bi
+    ] each ;
 
-:: handle-n ( inputs x -- )
-    inputs x length small-groups :> groups
-    groups H{ } tally :> b
-    x b at [ 0 ] unless*
-    number>string 8 CHAR: \s pad-tail write ;
+: handle-n ( input x -- )
+    [ nip ] [ length clump histogram ] 2bi at 0 or "%d\t" printf ;
 
 : process-input ( input -- )
-    dup 1 handle-table nl
-    dup 2 handle-table nl
-    { "GGT" "GGTA" "GGTATT" "GGTATTTTAATT" "GGTATTTTAATTTATAGT" }
-    [ [ dupd handle-n ] keep print ] each
-    drop ;
+    [ 1 handle-table nl ]
+    [ 2 handle-table nl ]
+    [
+        { "GGT" "GGTA" "GGTATT" "GGTATTTTAATT" "GGTATTTTAATTTATAGT" }
+        [ [ handle-n ] keep print ] with each
+    ]
+    tri ;
 
 : knucleotide ( -- )
     "resource:extra/benchmark/knucleotide/knucleotide-input.txt"
diff --git a/extra/cuda/cuda.factor b/extra/cuda/cuda.factor
index 887740d542..6b343fb1cc 100644
--- a/extra/cuda/cuda.factor
+++ b/extra/cuda/cuda.factor
@@ -1,17 +1,24 @@
 ! Copyright (C) 2010 Doug Coleman.
 ! See http://factorcode.org/license.txt for BSD license.
-USING: alien.c-types alien.data assocs classes.struct
-combinators continuations cuda.ffi fry io.backend kernel
-sequences ;
+USING: accessors alien alien.c-types alien.data alien.parser
+alien.strings arrays assocs byte-arrays classes.struct
+combinators continuations cuda.ffi destructors fry io
+io.backend io.encodings.string io.encodings.utf8 kernel lexer
+locals math math.parser namespaces opengl.gl.extensions
+prettyprint quotations sequences ;
 IN: cuda
 
+SYMBOL: cuda-device
+SYMBOL: cuda-context
+SYMBOL: cuda-module
+SYMBOL: cuda-function
+SYMBOL: cuda-launcher
+SYMBOL: cuda-memory-hashtable
+
 ERROR: throw-cuda-error n ;
 
 : cuda-error ( n -- )
-    {
-        { CUDA_SUCCESS [ ] }
-        [ throw-cuda-error ]
-    } case ;
+    dup CUDA_SUCCESS = [ drop ] [ throw-cuda-error ] if ;
 
 : cuda-version ( -- n )
     int <c-object> [ cuDriverGetVersion cuda-error ] keep *int ;
@@ -19,32 +26,10 @@ ERROR: throw-cuda-error n ;
 : init-cuda ( -- )
     0 cuInit cuda-error ;
 
-: with-cuda ( quot -- )
-    init-cuda [ ] [ ] cleanup ; inline
-
-<PRIVATE
-
-: #cuda-devices ( -- n )
-    int <c-object> [ cuDeviceGetCount cuda-error ] keep *int ;
-
-: n>cuda-device ( n -- device )
-    [ CUdevice <c-object> ] dip [ cuDeviceGet cuda-error ] 2keep drop *int ;
-
-: enumerate-cuda-devices ( -- devices )
-    #cuda-devices iota [ n>cuda-device ] map ;
-
-: cuda-device>properties ( device -- properties )
-    [ CUdevprop <c-object> ] dip
-    [ cuDeviceGetProperties cuda-error ] 2keep drop
-    CUdevprop memory>struct ;
-
-: cuda-device-properties ( -- properties )
-    enumerate-cuda-devices [ cuda-device>properties ] map ;
-
-PRIVATE>
-
-: cuda-devices ( -- assoc )
-    enumerate-cuda-devices [ dup cuda-device>properties ] { } map>assoc ;
+TUPLE: launcher
+{ device integer initial: 0 }
+{ device-flags initial: 0 }
+path block-shape shared-size grid ;
 
 : with-cuda-context ( flags device quot -- )
     [
@@ -65,13 +50,259 @@ PRIVATE>
     [ drop '[ _ cuModuleUnload cuda-error ] ] 2bi
     [ ] cleanup ; inline
 
-: get-cuda-function ( module string -- function )
+: with-cuda-program ( flags device path quot -- )
+    [ dup cuda-device set ] 2dip
+    '[
+        cuda-context set
+        _ [
+            cuda-module set
+            _ call
+        ] with-cuda-module
+    ] with-cuda-context ; inline
+
+: with-cuda ( launcher quot -- )
+    [
+        init-cuda
+        H{ } clone cuda-memory-hashtable
+    ] 2dip '[
+        _ 
+        [ cuda-launcher set ]
+        [ [ device>> ] [ device-flags>> ] [ path>> ] tri ] bi
+        _ with-cuda-program
+    ] with-variable ; inline
+
+<PRIVATE
+
+: #cuda-devices ( -- n )
+    int <c-object> [ cuDeviceGetCount cuda-error ] keep *int ;
+
+: n>cuda-device ( n -- device )
+    [ CUdevice <c-object> ] dip [ cuDeviceGet cuda-error ] 2keep drop *int ;
+
+: enumerate-cuda-devices ( -- devices )
+    #cuda-devices iota [ n>cuda-device ] map ;
+
+: cuda-device-properties ( device -- properties )
+    [ CUdevprop <c-object> ] dip
+    [ cuDeviceGetProperties cuda-error ] 2keep drop
+    CUdevprop memory>struct ;
+
+PRIVATE>
+
+: cuda-devices ( -- assoc )
+    enumerate-cuda-devices [ dup cuda-device-properties ] { } map>assoc ;
+
+: cuda-device-name ( n -- string )
+    [ 256 [ <byte-array> ] keep ] dip
+    [ cuDeviceGetName cuda-error ]
+    [ 2drop utf8 alien>string ] 3bi ;
+
+: cuda-device-capability ( n -- pair )
+    [ int <c-object> int <c-object> ] dip
+    [ cuDeviceComputeCapability cuda-error ]
+    [ drop [ *int ] bi@ ] 3bi 2array ;
+
+: cuda-device-memory ( n -- bytes )
+    [ uint <c-object> ] dip
+    [ cuDeviceTotalMem cuda-error ]
+    [ drop *uint ] 2bi ;
+
+: get-cuda-function* ( module string -- function )
     [ CUfunction <c-object> ] 2dip
     [ cuModuleGetFunction cuda-error ] 3keep 2drop *void* ;
 
+: get-cuda-function ( string -- function )
+    [ cuda-module get ] dip get-cuda-function* ;
+
+: with-cuda-function ( string quot -- )
+    [
+        get-cuda-function cuda-function set
+    ] dip call ; inline
+
+: launch-function* ( function -- ) cuLaunch cuda-error ;
+
+: launch-function ( -- ) cuda-function get cuLaunch cuda-error ;
+
+: launch-function-grid* ( function width height -- )
+    cuLaunchGrid cuda-error ;
+
+: launch-function-grid ( width height -- )
+    [ cuda-function get ] 2dip
+    cuLaunchGrid cuda-error ;
+
+TUPLE: cuda-memory < disposable ptr length ;
+
+: <cuda-memory> ( ptr length -- obj )
+    cuda-memory new-disposable
+        swap >>length
+        swap >>ptr ;
+
+: add-cuda-memory ( obj -- obj )
+    dup dup ptr>> cuda-memory-hashtable get set-at ;
+
+: delete-cuda-memory ( obj -- )
+    cuda-memory-hashtable delete-at ;
+
+ERROR: invalid-cuda-memory ptr ;
+
+: cuda-memory-length ( cuda-memory -- n )
+    ptr>> cuda-memory-hashtable get ?at [
+        length>>
+    ] [
+        invalid-cuda-memory
+    ] if ;
+
+M: cuda-memory byte-length length>> ;
+
 : cuda-malloc ( n -- ptr )
     [ CUdeviceptr <c-object> ] dip
-    [ cuMemAlloc cuda-error ] 2keep drop *int ;
+    [ cuMemAlloc cuda-error ] 2keep
+    [ *int ] dip <cuda-memory> add-cuda-memory ;
 
-: cuda-free ( ptr -- )
+: cuda-free* ( ptr -- )
     cuMemFree cuda-error ;
+
+M: cuda-memory dispose ( ptr -- )
+    ptr>> cuda-free* ;
+
+: host>device ( dest-ptr src-ptr -- )
+    [ ptr>> ] dip dup length cuMemcpyHtoD cuda-error ;
+
+:: device>host ( ptr -- seq )
+    ptr byte-length <byte-array>
+    [ ptr [ ptr>> ] [ byte-length ] bi cuMemcpyDtoH cuda-error ] keep ;
+
+: memcpy-device>device ( dest-ptr src-ptr count -- )
+    cuMemcpyDtoD cuda-error ;
+
+: memcpy-device>array ( dest-array dest-index src-ptr count -- )
+    cuMemcpyDtoA cuda-error ;
+
+: memcpy-array>device ( dest-ptr src-array src-index count -- )
+    cuMemcpyAtoD cuda-error ;
+
+: memcpy-array>host ( dest-ptr src-array src-index count -- )
+    cuMemcpyAtoH cuda-error ;
+
+: memcpy-host>array ( dest-array dest-index src-ptr count -- )
+    cuMemcpyHtoA cuda-error ;
+
+: memcpy-array>array ( dest-array dest-index src-array src-ptr count -- )
+    cuMemcpyAtoA cuda-error ;
+
+: cuda-int* ( function offset value -- )
+    cuParamSeti cuda-error ;
+
+: cuda-int ( offset value -- )
+    [ cuda-function get ] 2dip cuda-int* ;
+
+: cuda-float* ( function offset value -- )
+    cuParamSetf cuda-error ;
+
+: cuda-float ( offset value -- )
+    [ cuda-function get ] 2dip cuda-float* ;
+
+: cuda-vector* ( function offset ptr n -- )
+    cuParamSetv cuda-error ;
+
+: cuda-vector ( offset ptr n -- )
+    [ cuda-function get ] 3dip cuda-vector* ;
+
+: param-size* ( function n -- )
+    cuParamSetSize cuda-error ;
+
+: param-size ( n -- )
+    [ cuda-function get ] dip param-size* ;
+
+: malloc-device-string ( string -- n )
+    utf8 encode
+    [ length cuda-malloc ] keep
+    [ host>device ] [ drop ] 2bi ;
+
+ERROR: bad-cuda-parameter parameter ;
+
+:: set-parameters ( seq -- )
+    cuda-function get :> function
+    0 :> offset!
+    seq [
+        [ offset ] dip
+        {
+            { [ dup cuda-memory? ] [ ptr>> cuda-int ] }
+            { [ dup float? ] [ cuda-float ] }
+            { [ dup integer? ] [ cuda-int ] }
+            [ bad-cuda-parameter ]
+        } cond
+        offset 4 + offset!
+    ] each
+    offset param-size ;
+
+: cuda-device-attribute ( attribute dev -- n )
+    [ int <c-object> ] 2dip
+    [ cuDeviceGetAttribute cuda-error ]
+    [ 2drop *int ] 3bi ;
+
+: function-block-shape* ( function x y z -- )
+    cuFuncSetBlockShape cuda-error ;
+
+: function-block-shape ( x y z -- )
+    [ cuda-function get ] 3dip
+    cuFuncSetBlockShape cuda-error ;
+
+: function-shared-size* ( function n -- )
+    cuFuncSetSharedSize cuda-error ;
+
+: function-shared-size ( n -- )
+    [ cuda-function get ] dip
+    cuFuncSetSharedSize cuda-error ;
+
+: launch ( -- )
+    cuda-launcher get {
+        [ block-shape>> first3 function-block-shape ]
+        [ shared-size>> function-shared-size ]
+        [
+            grid>> [
+                launch-function
+            ] [
+                first2 launch-function-grid
+            ] if-empty
+        ]
+    } cleave ;
+
+: cuda-device. ( n -- )
+    {
+        [ "Device: " write number>string print ]
+        [ "Name: " write cuda-device-name print ]
+        [ "Memory: " write cuda-device-memory number>string print ]
+        [
+            "Capability: " write
+            cuda-device-capability [ number>string ] map " " join print
+        ]
+        [ "Properties: " write cuda-device-properties . ]
+        [
+            "CU_DEVICE_ATTRIBUTE_GPU_OVERLAP: " write
+            CU_DEVICE_ATTRIBUTE_GPU_OVERLAP swap
+            cuda-device-attribute number>string print
+        ]
+    } cleave ;
+
+: cuda. ( -- )
+    "CUDA Version: " write cuda-version number>string print nl
+    #cuda-devices iota [ nl ] [ cuda-device. ] interleave ;
+
+
+: test-cuda0 ( -- )
+    T{ launcher
+        { path "vocab:cuda/hello.ptx" }
+        { block-shape { 6 6 6 } }
+        { shared-size 2 }
+        { grid { 2 6 } }
+    } [
+        "helloWorld" [
+            "Hello World!" [ - ] map-index
+            malloc-device-string &dispose
+
+            [ 1array set-parameters ]
+            [ drop launch ]
+            [ device>host utf8 alien>string . ] tri
+        ] with-cuda-function
+    ] with-cuda ;
diff --git a/extra/cuda/ffi/tags.txt b/extra/cuda/ffi/tags.txt
index 700f0dc9a5..f74dbeec64 100644
--- a/extra/cuda/ffi/tags.txt
+++ b/extra/cuda/ffi/tags.txt
@@ -1 +1,2 @@
 not tested
+bindings
diff --git a/extra/cuda/hello.cu b/extra/cuda/hello.cu
new file mode 100644
index 0000000000..1f3cd677f9
--- /dev/null
+++ b/extra/cuda/hello.cu
@@ -0,0 +1,65 @@
+/*
+ World using CUDA
+** 
+** The string "Hello World!" is mangled then restored using a common CUDA idiom
+**
+** Byron Galbraith
+** 2009-02-18
+*/
+#include <cuda.h>
+#include <stdio.h>
+
+// Prototypes
+extern "C" __global__ void helloWorld(char*);
+
+// Host function
+int
+main(int argc, char** argv)
+{
+  int i;
+
+  // desired output
+  char str[] = "Hello World!";
+
+  // mangle contents of output
+  // the null character is left intact for simplicity
+  for(i = 0; i < 12; i++)
+    str[i] -= i;
+
+  // allocate memory on the device 
+  char *d_str;
+  size_t size = sizeof(str);
+  cudaMalloc((void**)&d_str, size);
+
+  // copy the string to the device
+  cudaMemcpy(d_str, str, size, cudaMemcpyHostToDevice);
+
+  // set the grid and block sizes
+  dim3 dimGrid(2);   // one block per word  
+  dim3 dimBlock(6); // one thread per character
+  
+  // invoke the kernel
+  helloWorld<<< dimGrid, dimBlock >>>(d_str);
+
+  // retrieve the results from the device
+  cudaMemcpy(str, d_str, size, cudaMemcpyDeviceToHost);
+
+  // free up the allocated memory on the device
+  cudaFree(d_str);
+  
+  // everyone's favorite part
+  printf("%s\n", str);
+
+  return 0;
+}
+
+// Device kernel
+__global__ void
+helloWorld(char* str)
+{
+  // determine where in the thread grid we are
+  int idx = blockIdx.x * blockDim.x + threadIdx.x;
+
+  // unmangle output
+  str[idx] += idx;
+}
diff --git a/extra/cuda/hello.ptx b/extra/cuda/hello.ptx
new file mode 100644
index 0000000000..049bb5e9a5
--- /dev/null
+++ b/extra/cuda/hello.ptx
@@ -0,0 +1,71 @@
+	.version 1.4
+	.target sm_10, map_f64_to_f32
+	// compiled with /usr/local/cuda/bin/../open64/lib//be
+	// nvopencc 3.0 built on 2010-03-11
+
+	//-----------------------------------------------------------
+	// Compiling /tmp/tmpxft_00000eab_00000000-7_hello.cpp3.i (/var/folders/KD/KDnx4D80Eh0fsORqNrFWBE+++TI/-Tmp-/ccBI#.AYqbdQ)
+	//-----------------------------------------------------------
+
+	//-----------------------------------------------------------
+	// Options:
+	//-----------------------------------------------------------
+	//  Target:ptx, ISA:sm_10, Endian:little, Pointer Size:32
+	//  -O3	(Optimization level)
+	//  -g0	(Debug level)
+	//  -m2	(Report advisories)
+	//-----------------------------------------------------------
+
+	.file	1	"<command-line>"
+	.file	2	"/tmp/tmpxft_00000eab_00000000-6_hello.cudafe2.gpu"
+	.file	3	"/usr/lib/gcc/i686-apple-darwin10/4.2.1/include/stddef.h"
+	.file	4	"/usr/local/cuda/bin/../include/crt/device_runtime.h"
+	.file	5	"/usr/local/cuda/bin/../include/host_defines.h"
+	.file	6	"/usr/local/cuda/bin/../include/builtin_types.h"
+	.file	7	"/usr/local/cuda/bin/../include/device_types.h"
+	.file	8	"/usr/local/cuda/bin/../include/driver_types.h"
+	.file	9	"/usr/local/cuda/bin/../include/texture_types.h"
+	.file	10	"/usr/local/cuda/bin/../include/vector_types.h"
+	.file	11	"/usr/local/cuda/bin/../include/device_launch_parameters.h"
+	.file	12	"/usr/local/cuda/bin/../include/crt/storage_class.h"
+	.file	13	"/usr/include/i386/_types.h"
+	.file	14	"/usr/include/time.h"
+	.file	15	"/usr/local/cuda/bin/../include/texture_fetch_functions.h"
+	.file	16	"/usr/local/cuda/bin/../include/common_functions.h"
+	.file	17	"/usr/local/cuda/bin/../include/crt/func_macro.h"
+	.file	18	"/usr/local/cuda/bin/../include/math_functions.h"
+	.file	19	"/usr/local/cuda/bin/../include/device_functions.h"
+	.file	20	"/usr/local/cuda/bin/../include/math_constants.h"
+	.file	21	"/usr/local/cuda/bin/../include/sm_11_atomic_functions.h"
+	.file	22	"/usr/local/cuda/bin/../include/sm_12_atomic_functions.h"
+	.file	23	"/usr/local/cuda/bin/../include/sm_13_double_functions.h"
+	.file	24	"/usr/local/cuda/bin/../include/common_types.h"
+	.file	25	"/usr/local/cuda/bin/../include/sm_20_atomic_functions.h"
+	.file	26	"/usr/local/cuda/bin/../include/sm_20_intrinsics.h"
+	.file	27	"/usr/local/cuda/bin/../include/math_functions_dbl_ptx1.h"
+	.file	28	"hello.cu"
+
+
+	.entry helloWorld (
+		.param .u32 __cudaparm_helloWorld_str)
+	{
+	.reg .u16 %rh<4>;
+	.reg .u32 %r<9>;
+	.loc	28	58	0
+$LBB1_helloWorld:
+	.loc	28	64	0
+	mov.u16 	%rh1, %ctaid.x;
+	mov.u16 	%rh2, %ntid.x;
+	mul.wide.u16 	%r1, %rh1, %rh2;
+	cvt.u32.u16 	%r2, %tid.x;
+	add.u32 	%r3, %r2, %r1;
+	ld.param.u32 	%r4, [__cudaparm_helloWorld_str];
+	add.u32 	%r5, %r4, %r3;
+	ld.global.s8 	%r6, [%r5+0];
+	add.s32 	%r7, %r6, %r3;
+	st.global.s8 	[%r5+0], %r7;
+	.loc	28	65	0
+	exit;
+$LDWend_helloWorld:
+	} // helloWorld
+
diff --git a/vm/callstack.cpp b/vm/callstack.cpp
index ad7528ab84..eae976219f 100755
--- a/vm/callstack.cpp
+++ b/vm/callstack.cpp
@@ -122,7 +122,7 @@ cell factor_vm::frame_scan(stack_frame *frame)
 			if(obj.type_p(QUOTATION_TYPE))
 			{
 				char *return_addr = (char *)FRAME_RETURN_ADDRESS(frame,this);
-				char *quot_entry_point = (char *)(frame_code(frame) + 1);
+				char *quot_entry_point = (char *)frame_code(frame)->entry_point();
 
 				return tag_fixnum(quot_code_offset_to_scan(
 					obj.value(),(cell)(return_addr - quot_entry_point)));
diff --git a/vm/io.cpp b/vm/io.cpp
index 8ce7ff5256..94e6e64d1d 100755
--- a/vm/io.cpp
+++ b/vm/io.cpp
@@ -216,7 +216,7 @@ void factor_vm::primitive_fread()
 		if(feof(file))
 		{
 			byte_array *new_buf = allot_byte_array(c);
-			memcpy(new_buf + 1, buf.untagged() + 1,c);
+			memcpy(new_buf->data<char>(), buf->data<char>(),c);
 			buf = new_buf;
 		}
 
diff --git a/vm/math.cpp b/vm/math.cpp
index a462232344..e64db2690e 100755
--- a/vm/math.cpp
+++ b/vm/math.cpp
@@ -260,10 +260,12 @@ void factor_vm::primitive_bignum_to_float()
 	ctx->replace(allot_float(bignum_to_float(ctx->peek())));
 }
 
-void factor_vm::primitive_float_to_str()
+void factor_vm::primitive_format_float()
 {
-	byte_array *array = allot_byte_array(33);
-	SNPRINTF((char *)(array + 1),32,"%.16g",untag_float_check(ctx->pop()));
+	byte_array *array = allot_byte_array(100);
+	char *format = alien_offset(ctx->pop());
+	double value = untag_float_check(ctx->pop());
+	SNPRINTF(array->data<char>(),99,format,value);
 	ctx->push(tag<byte_array>(array));
 }
 
diff --git a/vm/primitives.hpp b/vm/primitives.hpp
index ff0947912c..e98cf508b6 100644
--- a/vm/primitives.hpp
+++ b/vm/primitives.hpp
@@ -82,8 +82,8 @@ namespace factor
 	_(float_subtract) \
 	_(float_to_bignum) \
 	_(float_to_fixnum) \
-	_(float_to_str) \
 	_(fopen) \
+	_(format_float) \
 	_(fputc) \
 	_(fread) \
 	_(fseek) \
diff --git a/vm/vm.hpp b/vm/vm.hpp
index 36ec3260d6..dd1d48cf03 100755
--- a/vm/vm.hpp
+++ b/vm/vm.hpp
@@ -464,7 +464,7 @@ struct factor_vm
 	cell unbox_array_size_slow();
 	void primitive_fixnum_to_float();
 	void primitive_bignum_to_float();
-	void primitive_float_to_str();
+	void primitive_format_float();
 	void primitive_float_eq();
 	void primitive_float_add();
 	void primitive_float_subtract();