Merge branch 'master' of factorcode.org:/git/factor

release
Joe Groff 2010-04-14 19:37:54 -07:00
commit 4aa62ffb24
17 changed files with 469 additions and 125 deletions

View File

@ -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. ! 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 math.intervals arrays classes.algebra combinators columns
stack-checker.branches locals math stack-checker.branches locals math namespaces
compiler.utilities compiler.utilities
compiler.tree compiler.tree
compiler.tree.combinators compiler.tree.combinators
@ -10,6 +10,8 @@ compiler.tree.propagation.info
compiler.tree.propagation.nodes compiler.tree.propagation.nodes
compiler.tree.propagation.simple compiler.tree.propagation.simple
compiler.tree.propagation.constraints ; compiler.tree.propagation.constraints ;
FROM: sets => union ;
FROM: assocs => change-at ;
IN: compiler.tree.propagation.branches IN: compiler.tree.propagation.branches
! For conditionals, an assoc of child node # --> constraint ! For conditionals, an assoc of child node # --> constraint
@ -90,7 +92,7 @@ M: #phi propagate-before ( #phi -- )
bi ; bi ;
:: update-constraints ( new old -- ) :: 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 -- ) : include-child-constraints ( i -- )
infer-children-data get nth constraints swap at last infer-children-data get nth constraints swap at last

View File

@ -1,4 +1,4 @@
! Copyright (C) 2008 Slava Pestov. ! Copyright (C) 2008, 2010 Slava Pestov.
! See http://factorcode.org/license.txt for BSD license. ! See http://factorcode.org/license.txt for BSD license.
USING: arrays assocs math math.intervals kernel accessors USING: arrays assocs math math.intervals kernel accessors
sequences namespaces classes classes.algebra sequences namespaces classes classes.algebra
@ -87,8 +87,11 @@ TUPLE: implication p q ;
C: --> implication C: --> implication
: maybe-add ( elt seq -- seq' )
2dup member? [ nip ] [ swap suffix ] if ;
: assume-implication ( q p -- ) : 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 ; [ satisfied? [ assume ] [ drop ] if ] 2bi ;
M: implication assume* M: implication assume*

View File

@ -16,6 +16,7 @@ IN: formatting.tests
[ t ] [ "+0010" 10 "%+05d" sprintf = ] unit-test [ t ] [ "+0010" 10 "%+05d" sprintf = ] unit-test
[ t ] [ "123.456000" 123.456 "%f" sprintf = ] unit-test [ t ] [ "123.456000" 123.456 "%f" sprintf = ] unit-test
[ t ] [ "2.44" 2.436 "%.2f" 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 ] [ "123.10" 123.1 "%01.2f" sprintf = ] unit-test
[ t ] [ "1.2346" 1.23456789 "%.4f" sprintf = ] unit-test [ t ] [ "1.2346" 1.23456789 "%.4f" sprintf = ] unit-test
[ t ] [ " 1.23" 1.23456789 "%6.2f" sprintf = ] unit-test [ t ] [ " 1.23" 1.23456789 "%6.2f" sprintf = ] unit-test

View File

@ -3,7 +3,9 @@
USING: accessors arrays assocs calendar combinators fry kernel USING: accessors arrays assocs calendar combinators fry kernel
generalizations io io.streams.string macros math math.functions generalizations io io.streams.string macros math math.functions
math.parser peg.ebnf quotations sequences splitting strings 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 IN: formatting
<PRIVATE <PRIVATE
@ -26,31 +28,15 @@ IN: formatting
: >digits ( string -- digits ) : >digits ( string -- digits )
[ 0 ] [ string>number ] if-empty ; [ 0 ] [ string>number ] if-empty ;
: pad-digits ( string digits -- string' ) : format-simple ( x digits string -- string )
[ "." split1 ] dip [ CHAR: 0 pad-tail ] [ head-slice ] bi "." glue ; [ [ >float ] [ number>string ] bi* "%." ] dip
surround format-float ;
: max-digits ( n digits -- n' ) : format-scientific ( x digits -- string ) "e" format-simple ;
10^ [ * round ] keep / ; inline
: >exp ( x -- exp base ) : format-decimal ( x digits -- string ) "f" format-simple ;
[
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 ;
: exp>string ( exp base digits -- string ) ERROR: unknown-printf-directive ;
[ max-digits ] keep -rot
[
[ 0 < "-" "+" ? ]
[ abs number>string 2 CHAR: 0 pad-head ] bi
"e" -rot 3append
]
[ number>string ] bi*
rot pad-digits prepend ;
EBNF: parse-printf EBNF: parse-printf
@ -73,15 +59,15 @@ digits = (digits_)? => [[ 6 or ]]
fmt-% = "%" => [[ [ "%" ] ]] fmt-% = "%" => [[ [ "%" ] ]]
fmt-c = "c" => [[ [ 1string ] ]] fmt-c = "c" => [[ [ 1string ] ]]
fmt-C = "C" => [[ [ 1string >upper ] ]] fmt-C = "C" => [[ [ 1string >upper ] ]]
fmt-s = "s" => [[ [ dup number? [ number>string ] when ] ]] fmt-s = "s" => [[ [ present ] ]]
fmt-S = "S" => [[ [ dup number? [ number>string ] when >upper ] ]] fmt-S = "S" => [[ [ present >upper ] ]]
fmt-d = "d" => [[ [ >fixnum number>string ] ]] fmt-d = "d" => [[ [ >integer number>string ] ]]
fmt-e = digits "e" => [[ first '[ >exp _ exp>string ] ]] fmt-e = digits "e" => [[ first '[ _ format-scientific ] ]]
fmt-E = digits "E" => [[ first '[ >exp _ exp>string >upper ] ]] fmt-E = digits "E" => [[ first '[ _ format-scientific >upper ] ]]
fmt-f = digits "f" => [[ first dup '[ >float _ max-digits number>string _ pad-digits ] ]] fmt-f = digits "f" => [[ first '[ _ format-decimal ] ]]
fmt-x = "x" => [[ [ >hex ] ]] fmt-x = "x" => [[ [ >hex ] ]]
fmt-X = "X" => [[ [ >hex >upper ] ]] fmt-X = "X" => [[ [ >hex >upper ] ]]
unknown = (.)* => [[ "Unknown directive" throw ]] unknown = (.)* => [[ unknown-printf-directive ]]
strings_ = fmt-c|fmt-C|fmt-s|fmt-S strings_ = fmt-c|fmt-C|fmt-s|fmt-S
strings = pad width strings_ => [[ reverse compose-all ]] strings = pad width strings_ => [[ reverse compose-all ]]

View File

@ -289,7 +289,7 @@ M: bad-executable summary
\ (dlsym) { byte-array object } { c-ptr } define-primitive \ (dlsym) { byte-array object } { c-ptr } define-primitive
\ (exists?) { string } { object } define-primitive \ (exists?) { string } { object } define-primitive
\ (exit) { integer } { } 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 \ (fopen) { byte-array byte-array } { alien } define-primitive
\ (identity-hashcode) { object } { fixnum } define-primitive \ (identity-hashcode) { object } { fixnum } define-primitive
\ (save-image) { byte-array byte-array } { } define-primitive \ (save-image) { byte-array byte-array } { } define-primitive

View File

@ -470,7 +470,7 @@ tuple
{ "byte-array>bignum" "math" "primitive_byte_array_to_bignum" (( x -- y )) } { "byte-array>bignum" "math" "primitive_byte_array_to_bignum" (( x -- y )) }
{ "double>bits" "math" "primitive_double_bits" (( x -- n )) } { "double>bits" "math" "primitive_double_bits" (( x -- n )) }
{ "float>bits" "math" "primitive_float_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_multiply" (( x y -- z )) }
{ "bignum+" "math.private" "primitive_bignum_add" (( x y -- z )) } { "bignum+" "math.private" "primitive_bignum_add" (( x y -- z )) }
{ "bignum-" "math.private" "primitive_bignum_subtract" (( x y -- z )) } { "bignum-" "math.private" "primitive_bignum_subtract" (( x y -- z )) }

View File

@ -1,6 +1,7 @@
! (c)2009 Joe Groff bsd license ! (c)2009 Joe Groff bsd license
USING: accessors combinators kernel kernel.private math USING: accessors byte-arrays combinators kernel kernel.private
namespaces sequences sequences.private splitting strings make ; math namespaces sequences sequences.private splitting strings
make ;
IN: math.parser IN: math.parser
: digit> ( ch -- n ) : digit> ( ch -- n )
@ -356,15 +357,15 @@ M: ratio >base
mantissa-expt [ float>hex-value ] [ float>hex-expt ] bi* mantissa-expt [ float>hex-value ] [ float>hex-expt ] bi*
] bi 3append ; ] bi 3append ;
: float>decimal ( n -- str ) : format-float ( n format -- string )
(float>string) 0 suffix >byte-array (format-float)
[ 0 = ] trim-tail >string dup [ 0 = ] find drop head >string
fix-float ; fix-float ;
: float>base ( n base -- str ) : float>base ( n base -- str )
{ {
{ 16 [ float>hex ] } { 16 [ float>hex ] }
[ drop float>decimal ] [ drop "%.16g" format-float ]
} case ; inline } case ; inline
PRIVATE> PRIVATE>

View File

@ -1,16 +1,9 @@
USING: kernel locals io io.files splitting strings io.encodings.ascii USING: ascii kernel io io.files splitting strings
hashtables sequences assocs math namespaces prettyprint io.encodings.ascii hashtables sequences assocs math
math.parser combinators arrays sorting unicode.case ; math.statistics namespaces prettyprint math.parser combinators
arrays sorting formatting grouping fry ;
IN: benchmark.knucleotide 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 ( -- ) : discard-lines ( -- )
readln readln
[ ">THREE" head? [ discard-lines ] unless ] when* ; [ ">THREE" head? [ discard-lines ] unless ] when* ;
@ -20,37 +13,25 @@ IN: benchmark.knucleotide
">" read-until drop ">" read-until drop
CHAR: \n swap remove >upper ; 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 -- ) : handle-table ( inputs n -- )
small-groups clump
[ length ] keep [ histogram >alist sort-values reverse ] [ length ] bi
H{ } tally >alist '[
sort-values reverse [ first write bl ]
[ [ second 100 * _ /f "%.3f" printf nl ] bi
dup first write bl ] each ;
second 100 * over / 3 float>string print
] each
drop ;
:: handle-n ( inputs x -- ) : handle-n ( input x -- )
inputs x length small-groups :> groups [ nip ] [ length clump histogram ] 2bi at 0 or "%d\t" printf ;
groups H{ } tally :> b
x b at [ 0 ] unless*
number>string 8 CHAR: \s pad-tail write ;
: process-input ( input -- ) : process-input ( input -- )
dup 1 handle-table nl [ 1 handle-table nl ]
dup 2 handle-table nl [ 2 handle-table nl ]
{ "GGT" "GGTA" "GGTATT" "GGTATTTTAATT" "GGTATTTTAATTTATAGT" } [
[ [ dupd handle-n ] keep print ] each { "GGT" "GGTA" "GGTATT" "GGTATTTTAATT" "GGTATTTTAATTTATAGT" }
drop ; [ [ handle-n ] keep print ] with each
]
tri ;
: knucleotide ( -- ) : knucleotide ( -- )
"resource:extra/benchmark/knucleotide/knucleotide-input.txt" "resource:extra/benchmark/knucleotide/knucleotide-input.txt"

View File

@ -1,17 +1,24 @@
! Copyright (C) 2010 Doug Coleman. ! Copyright (C) 2010 Doug Coleman.
! See http://factorcode.org/license.txt for BSD license. ! See http://factorcode.org/license.txt for BSD license.
USING: alien.c-types alien.data assocs classes.struct USING: accessors alien alien.c-types alien.data alien.parser
combinators continuations cuda.ffi fry io.backend kernel alien.strings arrays assocs byte-arrays classes.struct
sequences ; 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 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 ; ERROR: throw-cuda-error n ;
: cuda-error ( n -- ) : cuda-error ( n -- )
{ dup CUDA_SUCCESS = [ drop ] [ throw-cuda-error ] if ;
{ CUDA_SUCCESS [ ] }
[ throw-cuda-error ]
} case ;
: cuda-version ( -- n ) : cuda-version ( -- n )
int <c-object> [ cuDriverGetVersion cuda-error ] keep *int ; int <c-object> [ cuDriverGetVersion cuda-error ] keep *int ;
@ -19,32 +26,10 @@ ERROR: throw-cuda-error n ;
: init-cuda ( -- ) : init-cuda ( -- )
0 cuInit cuda-error ; 0 cuInit cuda-error ;
: with-cuda ( quot -- ) TUPLE: launcher
init-cuda [ ] [ ] cleanup ; inline { device integer initial: 0 }
{ device-flags initial: 0 }
<PRIVATE path block-shape shared-size grid ;
: #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 ;
: with-cuda-context ( flags device quot -- ) : with-cuda-context ( flags device quot -- )
[ [
@ -65,13 +50,259 @@ PRIVATE>
[ drop '[ _ cuModuleUnload cuda-error ] ] 2bi [ drop '[ _ cuModuleUnload cuda-error ] ] 2bi
[ ] cleanup ; inline [ ] 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 [ CUfunction <c-object> ] 2dip
[ cuModuleGetFunction cuda-error ] 3keep 2drop *void* ; [ 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 ) : cuda-malloc ( n -- ptr )
[ CUdeviceptr <c-object> ] dip [ 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 ; 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 ;

View File

@ -1 +1,2 @@
not tested not tested
bindings

65
extra/cuda/hello.cu Normal file
View File

@ -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;
}

71
extra/cuda/hello.ptx Normal file
View File

@ -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

View File

@ -122,7 +122,7 @@ cell factor_vm::frame_scan(stack_frame *frame)
if(obj.type_p(QUOTATION_TYPE)) if(obj.type_p(QUOTATION_TYPE))
{ {
char *return_addr = (char *)FRAME_RETURN_ADDRESS(frame,this); 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( return tag_fixnum(quot_code_offset_to_scan(
obj.value(),(cell)(return_addr - quot_entry_point))); obj.value(),(cell)(return_addr - quot_entry_point)));

View File

@ -216,7 +216,7 @@ void factor_vm::primitive_fread()
if(feof(file)) if(feof(file))
{ {
byte_array *new_buf = allot_byte_array(c); 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; buf = new_buf;
} }

View File

@ -260,10 +260,12 @@ void factor_vm::primitive_bignum_to_float()
ctx->replace(allot_float(bignum_to_float(ctx->peek()))); 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); byte_array *array = allot_byte_array(100);
SNPRINTF((char *)(array + 1),32,"%.16g",untag_float_check(ctx->pop())); 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)); ctx->push(tag<byte_array>(array));
} }

View File

@ -82,8 +82,8 @@ namespace factor
_(float_subtract) \ _(float_subtract) \
_(float_to_bignum) \ _(float_to_bignum) \
_(float_to_fixnum) \ _(float_to_fixnum) \
_(float_to_str) \
_(fopen) \ _(fopen) \
_(format_float) \
_(fputc) \ _(fputc) \
_(fread) \ _(fread) \
_(fseek) \ _(fseek) \

View File

@ -464,7 +464,7 @@ struct factor_vm
cell unbox_array_size_slow(); cell unbox_array_size_slow();
void primitive_fixnum_to_float(); void primitive_fixnum_to_float();
void primitive_bignum_to_float(); void primitive_bignum_to_float();
void primitive_float_to_str(); void primitive_format_float();
void primitive_float_eq(); void primitive_float_eq();
void primitive_float_add(); void primitive_float_add();
void primitive_float_subtract(); void primitive_float_subtract();