-! 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
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
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
-! 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
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*
[ 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
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
: >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 ;
-
-: 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 ;
+: format-decimal ( x digits -- string ) "f" format-simple ;
+
+ERROR: unknown-printf-directive ;
EBNF: parse-printf
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 ]]
\ (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
{ "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 )) }
! (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 )
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>
-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* ;
">" 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"
! 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 ;
: init-cuda ( -- )
0 cuInit cuda-error ;
-: with-cuda ( quot -- )
- init-cuda [ ] [ ] cleanup ; inline
+TUPLE: launcher
+{ device integer initial: 0 }
+{ device-flags initial: 0 }
+path block-shape shared-size grid ;
+
+: with-cuda-context ( flags device quot -- )
+ [
+ [ CUcontext <c-object> ] 2dip
+ [ cuCtxCreate cuda-error ] 3keep 2drop *void*
+ ] dip
+ [ '[ _ @ ] ]
+ [ drop '[ _ cuCtxDestroy cuda-error ] ] 2bi
+ [ ] cleanup ; inline
+
+: with-cuda-module ( path quot -- )
+ [
+ normalize-path
+ [ CUmodule <c-object> ] dip
+ [ cuModuleLoad cuda-error ] 2keep drop *void*
+ ] dip
+ [ '[ _ @ ] ]
+ [ drop '[ _ cuModuleUnload cuda-error ] ] 2bi
+ [ ] cleanup ; inline
+
+: 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
: enumerate-cuda-devices ( -- devices )
#cuda-devices iota [ n>cuda-device ] map ;
-: cuda-device>properties ( device -- properties )
+: 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 ;
+ enumerate-cuda-devices [ dup cuda-device-properties ] { } map>assoc ;
-: with-cuda-context ( flags device quot -- )
- [
- [ CUcontext <c-object> ] 2dip
- [ cuCtxCreate cuda-error ] 3keep 2drop *void*
- ] dip
- [ '[ _ @ ] ]
- [ drop '[ _ cuCtxDestroy cuda-error ] ] 2bi
- [ ] cleanup ; inline
+: cuda-device-name ( n -- string )
+ [ 256 [ <byte-array> ] keep ] dip
+ [ cuDeviceGetName cuda-error ]
+ [ 2drop utf8 alien>string ] 3bi ;
-: with-cuda-module ( path quot -- )
- [
- normalize-path
- [ CUmodule <c-object> ] dip
- [ cuModuleLoad cuda-error ] 2keep drop *void*
- ] dip
- [ '[ _ @ ] ]
- [ drop '[ _ cuModuleUnload cuda-error ] ] 2bi
- [ ] cleanup ; inline
+: 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 )
+: 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 ;
--- /dev/null
+/*
+ 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;
+}
--- /dev/null
+ .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
+
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)));
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;
}
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));
}
_(float_subtract) \
_(float_to_bignum) \
_(float_to_fixnum) \
- _(float_to_str) \
_(fopen) \
+ _(format_float) \
_(fputc) \
_(fread) \
_(fseek) \
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();
projects / factor.git / commitdiff