diff --git a/basis/alien/data/data.factor b/basis/alien/data/data.factor index a0450d5122..af1ed24663 100644 --- a/basis/alien/data/data.factor +++ b/basis/alien/data/data.factor @@ -1,8 +1,7 @@ ! (c)2009, 2010 Slava Pestov, Joe Groff bsd license -USING: accessors alien alien.c-types alien.arrays alien.strings arrays -byte-arrays cpu.architecture fry io io.encodings.binary -io.files io.streams.memory kernel libc math sequences words -byte-vectors ; +USING: accessors alien alien.c-types alien.arrays alien.strings +arrays byte-arrays cpu.architecture fry io io.encodings.binary +io.files io.streams.memory kernel libc math sequences words ; IN: alien.data GENERIC: require-c-array ( c-type -- ) @@ -63,13 +62,6 @@ M: memory-stream stream-read swap memory>byte-array ] [ [ + ] change-index drop ] 2bi ; -M: byte-vector stream-write - [ dup byte-length tail-slice ] - [ [ [ byte-length ] bi@ + ] keep lengthen ] - [ drop byte-length ] - 2tri - [ >c-ptr swap >c-ptr ] dip memcpy ; - M: value-type c-type-rep drop int-rep ; M: value-type c-type-getter @@ -83,4 +75,3 @@ M: array c-type-boxer-quot unclip [ array-length ] dip [ ] 2curry ; M: array c-type-unboxer-quot drop [ >c-ptr ] ; - diff --git a/basis/biassocs/biassocs.factor b/basis/biassocs/biassocs.factor index 7daa478f54..ab3157d400 100644 --- a/basis/biassocs/biassocs.factor +++ b/basis/biassocs/biassocs.factor @@ -13,9 +13,9 @@ TUPLE: biassoc from to ; M: biassoc assoc-size from>> assoc-size ; -M: biassoc at* from>> at* ; +M: biassoc at* from>> at* ; inline -M: biassoc value-at* to>> at* ; +M: biassoc value-at* to>> at* ; inline : once-at ( value key assoc -- ) 2dup key? [ 3drop ] [ set-at ] if ; diff --git a/basis/binary-search/binary-search.factor b/basis/binary-search/binary-search.factor index 36e983a1c8..db40408d5e 100644 --- a/basis/binary-search/binary-search.factor +++ b/basis/binary-search/binary-search.factor @@ -1,14 +1,14 @@ ! Copyright (C) 2008, 2010 Slava Pestov. ! See http://factorcode.org/license.txt for BSD license. USING: accessors arrays combinators hints kernel locals math -math.order sequences ; +math.order sequences sequences.private ; IN: binary-search ) -- i elt ) from to + 2/ :> midpoint@ - midpoint@ seq nth :> midpoint + midpoint@ seq nth-unsafe :> midpoint to from - 1 <= [ midpoint@ midpoint diff --git a/basis/bit-sets/bit-sets-tests.factor b/basis/bit-sets/bit-sets-tests.factor index 4e97e703d0..0d4543f8f2 100644 --- a/basis/bit-sets/bit-sets-tests.factor +++ b/basis/bit-sets/bit-sets-tests.factor @@ -11,6 +11,9 @@ IN: bit-sets.tests T{ bit-set f ?{ f f t f t f } } intersect ] unit-test +[ f ] [ T{ bit-set f ?{ t f f f t f } } null? ] unit-test +[ t ] [ T{ bit-set f ?{ f f f f f f } } null? ] unit-test + [ T{ bit-set f ?{ t f t f f f } } ] [ T{ bit-set f ?{ t t t f f f } } T{ bit-set f ?{ f t f f t t } } diff diff --git a/basis/bootstrap/compiler/compiler.factor b/basis/bootstrap/compiler/compiler.factor index 0237ed99ee..56109e2de6 100644 --- a/basis/bootstrap/compiler/compiler.factor +++ b/basis/bootstrap/compiler/compiler.factor @@ -20,8 +20,8 @@ IN: bootstrap.compiler "alien.remote-control" require ] unless -"prettyprint" "alien.prettyprint" require-when -"debugger" "alien.debugger" require-when +{ "boostrap.compiler" "prettyprint" } "alien.prettyprint" require-when +{ "boostrap.compiler" "debugger" } "alien.debugger" require-when "cpu." cpu name>> append require @@ -35,7 +35,7 @@ gc [ optimized? not ] filter compile ; "debug-compiler" get [ - + nl "Compiling..." write flush @@ -57,7 +57,7 @@ gc curry compose uncurry - array-nth set-array-nth length>> + array-nth set-array-nth wrap probe @@ -117,4 +117,6 @@ gc " done" print flush + "io.streams.byte-array.fast" require + ] unless diff --git a/basis/bootstrap/handbook/handbook.factor b/basis/bootstrap/handbook/handbook.factor index 11f7349b79..ef7a456b7b 100644 --- a/basis/bootstrap/handbook/handbook.factor +++ b/basis/bootstrap/handbook/handbook.factor @@ -1,4 +1,4 @@ USING: vocabs.loader vocabs kernel ; IN: bootstrap.handbook -"bootstrap.help" "help.handbook" require-when +{ "boostrap.handbook" "bootstrap.help" } "help.handbook" require-when diff --git a/basis/bootstrap/threads/threads.factor b/basis/bootstrap/threads/threads.factor index 3a8fe98cf4..2bc8d612b6 100644 --- a/basis/bootstrap/threads/threads.factor +++ b/basis/bootstrap/threads/threads.factor @@ -4,6 +4,6 @@ USING: vocabs.loader kernel io.thread threads compiler.utilities namespaces ; IN: bootstrap.threads -"debugger" "debugger.threads" require-when +{ "bootstrap.threads" "debugger" } "debugger.threads" require-when [ yield ] yield-hook set-global diff --git a/basis/bootstrap/ui/tools/tools.factor b/basis/bootstrap/ui/tools/tools.factor index 7db69ce9c1..3efd156983 100644 --- a/basis/bootstrap/ui/tools/tools.factor +++ b/basis/bootstrap/ui/tools/tools.factor @@ -4,7 +4,7 @@ USING: kernel vocabs vocabs.loader sequences system ; [ "bootstrap." prepend vocab ] all? [ "ui.tools" require - "ui.backend.cocoa" "ui.backend.cocoa.tools" require-when + { "ui.backend.cocoa" } "ui.backend.cocoa.tools" require-when "ui.tools.walker" require ] when diff --git a/basis/classes/struct/struct.factor b/basis/classes/struct/struct.factor index ffde233748..605ee573f5 100644 --- a/basis/classes/struct/struct.factor +++ b/basis/classes/struct/struct.factor @@ -404,4 +404,4 @@ FUNCTOR-SYNTAX: STRUCT: USING: vocabs vocabs.loader ; -"prettyprint" "classes.struct.prettyprint" require-when +{ "classes.struct" "prettyprint" } "classes.struct.prettyprint" require-when diff --git a/basis/compiler/cfg/alias-analysis/alias-analysis.factor b/basis/compiler/cfg/alias-analysis/alias-analysis.factor index 44326c179f..2e0684c5d0 100644 --- a/basis/compiler/cfg/alias-analysis/alias-analysis.factor +++ b/basis/compiler/cfg/alias-analysis/alias-analysis.factor @@ -287,7 +287,7 @@ M: ##copy analyze-aliases* M: ##compare analyze-aliases* call-next-method dup useless-compare? [ - dst>> \ f type-number \ ##load-immediate new-insn + dst>> f \ ##load-constant new-insn analyze-aliases* ] when ; diff --git a/basis/compiler/cfg/builder/builder.factor b/basis/compiler/cfg/builder/builder.factor index 529c3b5ae6..370f3d053f 100644 --- a/basis/compiler/cfg/builder/builder.factor +++ b/basis/compiler/cfg/builder/builder.factor @@ -123,7 +123,7 @@ M: #recursive emit-node and ; : emit-trivial-if ( -- ) - ds-pop \ f type-number cc/= ^^compare-imm ds-push ; + ds-pop f cc/= ^^compare-imm ds-push ; : trivial-not-if? ( #if -- ? ) children>> first2 @@ -132,12 +132,12 @@ M: #recursive emit-node and ; : emit-trivial-not-if ( -- ) - ds-pop \ f type-number cc= ^^compare-imm ds-push ; + ds-pop f cc= ^^compare-imm ds-push ; : emit-actual-if ( #if -- ) ! Inputs to the final instruction need to be copied because of ! loc>vreg sync - ds-pop any-rep ^^copy \ f type-number cc/= ##compare-imm-branch emit-if ; + ds-pop any-rep ^^copy f cc/= ##compare-imm-branch emit-if ; M: #if emit-node { diff --git a/basis/compiler/cfg/hats/hats.factor b/basis/compiler/cfg/hats/hats.factor index 9d1945c525..fb89b36efa 100644 --- a/basis/compiler/cfg/hats/hats.factor +++ b/basis/compiler/cfg/hats/hats.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: accessors arrays byte-arrays kernel layouts math -namespaces sequences combinators splitting parser effects -words cpu.architecture compiler.cfg.registers +USING: accessors arrays byte-arrays combinators.short-circuit +kernel layouts math namespaces sequences combinators splitting +parser effects words cpu.architecture compiler.cfg.registers compiler.cfg.instructions compiler.cfg.instructions.syntax ; IN: compiler.cfg.hats @@ -41,11 +41,13 @@ insn-classes get [ >> +: immutable? ( obj -- ? ) + { [ float? ] [ word? ] [ not ] } 1|| ; inline + : ^^load-literal ( obj -- dst ) [ next-vreg dup ] dip { - { [ dup not ] [ drop \ f type-number ##load-immediate ] } { [ dup fixnum? ] [ tag-fixnum ##load-immediate ] } - { [ dup float? ] [ ##load-constant ] } + { [ dup immutable? ] [ ##load-constant ] } [ ##load-reference ] } cond ; diff --git a/basis/compiler/cfg/instructions/instructions.factor b/basis/compiler/cfg/instructions/instructions.factor index c015cb640b..5ddf7b4db5 100644 --- a/basis/compiler/cfg/instructions/instructions.factor +++ b/basis/compiler/cfg/instructions/instructions.factor @@ -33,6 +33,10 @@ INSN: ##load-constant def: dst/int-rep constant: obj ; +INSN: ##load-double +def: dst/double-rep +constant: val ; + INSN: ##peek def: dst/int-rep literal: loc ; diff --git a/basis/compiler/cfg/intrinsics/fixnum/fixnum.factor b/basis/compiler/cfg/intrinsics/fixnum/fixnum.factor index e4d1735eae..ad7e02df8a 100644 --- a/basis/compiler/cfg/intrinsics/fixnum/fixnum.factor +++ b/basis/compiler/cfg/intrinsics/fixnum/fixnum.factor @@ -20,9 +20,6 @@ IN: compiler.cfg.intrinsics.fixnum 0 cc= ^^compare-imm ds-push ; -: tag-literal ( n -- tagged ) - literal>> [ tag-fixnum ] [ \ f type-number ] if* ; - : emit-fixnum-op ( insn -- ) [ 2inputs ] dip call ds-push ; inline @@ -44,7 +41,7 @@ IN: compiler.cfg.intrinsics.fixnum { [ dup 0 [-inf,a] interval-subset? ] [ drop emit-fixnum-right-shift ] } [ drop emit-fixnum-shift-general ] } cond ; - + : emit-fixnum-bitnot ( -- ) ds-pop ^^not tag-mask get ^^xor-imm ds-push ; diff --git a/basis/compiler/cfg/representations/preferred/preferred.factor b/basis/compiler/cfg/representations/preferred/preferred.factor index ffb8f9a390..e4114c9249 100644 --- a/basis/compiler/cfg/representations/preferred/preferred.factor +++ b/basis/compiler/cfg/representations/preferred/preferred.factor @@ -68,23 +68,23 @@ PRIVATE> tri ] with-compilation-unit -: each-def-rep ( ... insn vreg-quot: ( ... vreg rep -- ... ) -- ... ) +: each-def-rep ( insn vreg-quot: ( vreg rep -- ) -- ) [ [ defs-vreg ] [ defs-vreg-rep ] bi ] dip with when* ; inline -: each-use-rep ( ... insn vreg-quot: ( ... vreg rep -- ... ) -- ... ) +: each-use-rep ( insn vreg-quot: ( vreg rep -- ) -- ) [ [ uses-vregs ] [ uses-vreg-reps ] bi ] dip 2each ; inline -: each-temp-rep ( ... insn vreg-quot: ( ... vreg rep -- ... ) -- ... ) +: each-temp-rep ( insn vreg-quot: ( vreg rep -- ) -- ) [ [ temp-vregs ] [ temp-vreg-reps ] bi ] dip 2each ; inline +: each-rep ( insn vreg-quot: ( vreg rep -- ) -- ) + [ each-def-rep ] [ each-use-rep ] [ each-temp-rep ] 2tri ; inline + : with-vreg-reps ( ..a cfg vreg-quot: ( ..a vreg rep -- ..b ) -- ..b ) '[ [ basic-block set ] [ [ - _ - [ each-def-rep ] - [ each-use-rep ] - [ each-temp-rep ] 2tri + _ each-rep ] each-non-phi ] bi ] each-basic-block ; inline diff --git a/basis/compiler/cfg/representations/representations-tests.factor b/basis/compiler/cfg/representations/representations-tests.factor index c50cfc4c86..a00f65e075 100644 --- a/basis/compiler/cfg/representations/representations-tests.factor +++ b/basis/compiler/cfg/representations/representations-tests.factor @@ -1,6 +1,7 @@ -USING: tools.test cpu.architecture -compiler.cfg.registers compiler.cfg.instructions -compiler.cfg.representations.preferred ; +USING: accessors compiler.cfg compiler.cfg.debugger +compiler.cfg.instructions compiler.cfg.registers +compiler.cfg.representations.preferred cpu.architecture kernel +namespaces tools.test sequences arrays system ; IN: compiler.cfg.representations [ { double-rep double-rep } ] [ @@ -16,4 +17,111 @@ IN: compiler.cfg.representations { dst 5 } { src 3 } } defs-vreg-rep -] unit-test \ No newline at end of file +] unit-test + +: test-representations ( -- ) + cfg new 0 get >>entry dup cfg set select-representations drop ; + +! Make sure cost calculation isn't completely wrong +V{ + T{ ##prologue } + T{ ##branch } +} 0 test-bb + +V{ + T{ ##peek f 1 D 0 } + T{ ##peek f 2 D 1 } + T{ ##add-float f 3 1 2 } + T{ ##replace f 3 D 0 } + T{ ##replace f 3 D 1 } + T{ ##replace f 3 D 2 } + T{ ##branch } +} 1 test-bb + +V{ + T{ ##epilogue } + T{ ##return } +} 2 test-bb + +0 1 edge +1 2 edge + +[ ] [ test-representations ] unit-test + +[ 1 ] [ 1 get instructions>> [ ##allot? ] count ] unit-test + +cpu x86.32? [ + + ! Make sure load-constant is converted into load-double + V{ + T{ ##prologue } + T{ ##branch } + } 0 test-bb + + V{ + T{ ##peek f 1 D 0 } + T{ ##load-constant f 2 0.5 } + T{ ##add-float f 3 1 2 } + T{ ##replace f 3 D 0 } + T{ ##branch } + } 1 test-bb + + V{ + T{ ##epilogue } + T{ ##return } + } 2 test-bb + + 0 1 edge + 1 2 edge + + [ ] [ test-representations ] unit-test + + [ t ] [ 1 get instructions>> second ##load-double? ] unit-test + + ! Make sure phi nodes are handled in a sane way + V{ + T{ ##prologue } + T{ ##branch } + } 0 test-bb + + V{ + T{ ##peek f 1 D 0 } + T{ ##compare-imm-branch f 1 2 } + } 1 test-bb + + V{ + T{ ##load-constant f 2 1.5 } + T{ ##branch } + } 2 test-bb + + V{ + T{ ##load-constant f 3 2.5 } + T{ ##branch } + } 3 test-bb + + V{ + T{ ##phi f 4 } + T{ ##peek f 5 D 0 } + T{ ##add-float f 6 4 5 } + T{ ##replace f 6 D 0 } + } 4 test-bb + + V{ + T{ ##epilogue } + T{ ##return } + } 5 test-bb + + test-diamond + 4 5 edge + + 2 get 2 2array + 3 get 3 2array 2array 4 get instructions>> first (>>inputs) + + [ ] [ test-representations ] unit-test + + [ t ] [ 2 get instructions>> first ##load-double? ] unit-test + + [ t ] [ 3 get instructions>> first ##load-double? ] unit-test + + [ t ] [ 4 get instructions>> first ##phi? ] unit-test +] when \ No newline at end of file diff --git a/basis/compiler/cfg/representations/representations.factor b/basis/compiler/cfg/representations/representations.factor index 05e365e5e4..f202dc4c6a 100644 --- a/basis/compiler/cfg/representations/representations.factor +++ b/basis/compiler/cfg/representations/representations.factor @@ -1,4 +1,4 @@ -! Copyright (C) 2009 Slava Pestov +! Copyright (C) 2009, 2010 Slava Pestov ! See http://factorcode.org/license.txt for BSD license. USING: kernel fry accessors sequences assocs sets namespaces arrays combinators combinators.short-circuit math make locals @@ -91,8 +91,8 @@ SYMBOL: possibilities : possible ( vreg -- reps ) possibilities get at ; : compute-possibilities ( cfg -- ) - H{ } clone [ '[ swap _ conjoin-at ] with-vreg-reps ] keep - [ keys ] assoc-map possibilities set ; + H{ } clone [ '[ swap _ adjoin-at ] with-vreg-reps ] keep + [ members ] assoc-map possibilities set ; ! Compute vregs which must remain tagged for their lifetime. SYMBOL: always-boxed @@ -119,15 +119,18 @@ SYMBOL: always-boxed SYMBOL: costs : init-costs ( -- ) - possibilities get [ [ 0 ] H{ } map>assoc ] assoc-map costs set ; + possibilities get [ drop H{ } clone ] assoc-map costs set ; + +: record-possibility ( rep vreg -- ) + costs get at [ 0 or ] change-at ; : increase-cost ( rep vreg -- ) ! Increase cost of keeping vreg in rep, making a choice of rep less ! likely. - [ basic-block get loop-nesting-at ] 2dip costs get at at+ ; + costs get at [ 0 or basic-block get loop-nesting-at 1 + + ] change-at ; : maybe-increase-cost ( possible vreg preferred -- ) - pick eq? [ 2drop ] [ increase-cost ] if ; + pick eq? [ record-possibility ] [ increase-cost ] if ; : representation-cost ( vreg preferred -- ) ! 'preferred' is a representation that the instruction can accept with no cost. @@ -137,11 +140,29 @@ SYMBOL: costs [ '[ _ _ maybe-increase-cost ] ] 2bi each ; +GENERIC: compute-insn-costs ( insn -- ) + +M: ##load-constant compute-insn-costs + ! There's no cost to unboxing the result of a ##load-constant + drop ; + +M: insn compute-insn-costs [ representation-cost ] each-rep ; + : compute-costs ( cfg -- costs ) - init-costs [ representation-cost ] with-vreg-reps costs get ; + init-costs + [ + [ basic-block set ] + [ + [ + compute-insn-costs + ] each-non-phi + ] bi + ] each-basic-block + costs get ; ! For every vreg, compute preferred representation, that minimizes costs. : minimize-costs ( costs -- representations ) + [ nip assoc-empty? not ] assoc-filter [ >alist alist-min first ] assoc-map ; : compute-representations ( cfg -- ) @@ -150,6 +171,54 @@ SYMBOL: costs bi assoc-union representations set ; +! PHI nodes require special treatment +! If the output of a phi instruction is only used as the input to another +! phi instruction, then we want to use the same representation for both +! if possible. +SYMBOL: phis + +: collect-phis ( cfg -- ) + H{ } clone phis set + [ + phis get + '[ [ inputs>> values ] [ dst>> ] bi _ set-at ] each-phi + ] each-basic-block ; + +SYMBOL: work-list + +: add-to-work-list ( vregs -- ) + work-list get push-all-front ; + +: rep-assigned ( vregs -- vregs' ) + representations get '[ _ key? ] filter ; + +: rep-not-assigned ( vregs -- vregs' ) + representations get '[ _ key? not ] filter ; + +: add-ready-phis ( -- ) + phis get keys rep-assigned add-to-work-list ; + +: process-phi ( dst -- ) + ! If dst = phi(src1,src2,...) and dst's representation has been + ! determined, assign that representation to each one of src1,... + ! that does not have a representation yet, and process those, too. + dup phis get at* [ + [ rep-of ] [ rep-not-assigned ] bi* + [ [ set-rep-of ] with each ] [ add-to-work-list ] bi + ] [ 2drop ] if ; + +: remaining-phis ( -- ) + phis get keys rep-not-assigned { } assert-sequence= ; + +: process-phis ( -- ) + work-list set + add-ready-phis + work-list get [ process-phi ] slurp-deque + remaining-phis ; + +: compute-phi-representations ( cfg -- ) + collect-phis process-phis ; + ! Insert conversions. This introduces new temporaries, so we need ! to rename opearands too. @@ -188,7 +257,7 @@ SYMBOLS: renaming-set needs-renaming? ; : record-renaming ( from to -- ) 2array renaming-set get push needs-renaming? on ; -:: (compute-renaming-set) ( ..a vreg required quot: ( ..a vreg preferred required -- ..b ) -- ..b ) +:: (compute-renaming-set) ( vreg required quot: ( vreg preferred required -- new-vreg ) -- ) vreg rep-of :> preferred preferred required eq? [ vreg no-renaming ] @@ -217,15 +286,16 @@ RENAMING: convert [ converted-value ] [ converted-value ] [ ] GENERIC: conversions-for-insn ( insn -- ) -SYMBOL: phi-mappings +M: ##phi conversions-for-insn , ; -! compiler.cfg.cssa inserts conversions which convert phi inputs into -! the representation of the output. However, we still have to do some -! processing here, because if the only node that uses the output of -! the phi instruction is another phi instruction then this phi node's -! output won't have a representation assigned. -M: ##phi conversions-for-insn - [ , ] [ [ inputs>> values ] [ dst>> ] bi phi-mappings get set-at ] bi ; +! When a float is unboxed, we replace the ##load-constant with a ##load-double +! if the architecture supports it +: convert-to-load-double? ( insn -- ? ) + { + [ drop load-double? ] + [ dst>> rep-of double-rep? ] + [ obj>> float? ] + } 1&& ; ! When a literal zeroes/ones vector is unboxed, we replace the ##load-reference ! with a ##zero-vector or ##fill-vector instruction since this is more efficient. @@ -234,17 +304,25 @@ M: ##phi conversions-for-insn [ dst>> rep-of vector-rep? ] [ obj>> B{ 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 } = ] } 1&& ; + : convert-to-fill-vector? ( insn -- ? ) { [ dst>> rep-of vector-rep? ] [ obj>> B{ 255 255 255 255 255 255 255 255 255 255 255 255 255 255 255 255 } = ] } 1&& ; +: (convert-to-load-double) ( insn -- dst val ) + [ dst>> ] [ obj>> ] bi ; inline + : (convert-to-zero/fill-vector) ( insn -- dst rep ) dst>> dup rep-of ; inline : conversions-for-load-insn ( insn -- ?insn ) { + { + [ dup convert-to-load-double? ] + [ (convert-to-load-double) ##load-double f ] + } { [ dup convert-to-zero-vector? ] [ (convert-to-zero/fill-vector) ##zero-vector f ] @@ -277,46 +355,8 @@ M: insn conversions-for-insn , ; ] change-instructions drop ] if ; -! If the output of a phi instruction is only used as the input to another -! phi instruction, then we want to use the same representation for both -! if possible. -SYMBOL: work-list - -: add-to-work-list ( vregs -- ) - work-list get push-all-front ; - -: rep-assigned ( vregs -- vregs' ) - representations get '[ _ key? ] filter ; - -: rep-not-assigned ( vregs -- vregs' ) - representations get '[ _ key? not ] filter ; - -: add-ready-phis ( -- ) - phi-mappings get keys rep-assigned add-to-work-list ; - -: process-phi-mapping ( dst -- ) - ! If dst = phi(src1,src2,...) and dst's representation has been - ! determined, assign that representation to each one of src1,... - ! that does not have a representation yet, and process those, too. - dup phi-mappings get at* [ - [ rep-of ] [ rep-not-assigned ] bi* - [ [ set-rep-of ] with each ] [ add-to-work-list ] bi - ] [ 2drop ] if ; - -: remaining-phi-mappings ( -- ) - phi-mappings get keys rep-not-assigned - [ [ int-rep ] dip set-rep-of ] each ; - -: process-phi-mappings ( -- ) - work-list set - add-ready-phis - work-list get [ process-phi-mapping ] slurp-deque - remaining-phi-mappings ; - : insert-conversions ( cfg -- ) - H{ } clone phi-mappings set - [ conversions-for-block ] each-basic-block - process-phi-mappings ; + [ conversions-for-block ] each-basic-block ; PRIVATE> @@ -326,6 +366,7 @@ PRIVATE> { [ compute-possibilities ] [ compute-representations ] + [ compute-phi-representations ] [ insert-conversions ] [ ] } cleave diff --git a/basis/compiler/cfg/value-numbering/rewrite/rewrite.factor b/basis/compiler/cfg/value-numbering/rewrite/rewrite.factor index 0fa0314c3e..81f39d7da2 100644 --- a/basis/compiler/cfg/value-numbering/rewrite/rewrite.factor +++ b/basis/compiler/cfg/value-numbering/rewrite/rewrite.factor @@ -27,6 +27,12 @@ IN: compiler.cfg.value-numbering.rewrite [ value>> immediate-bitwise? ] } 1&& ; +: vreg-immediate-comparand? ( vreg -- ? ) + vreg>expr { + [ constant-expr? ] + [ value>> immediate-comparand? ] + } 1&& ; + ! Outputs f to mean no change GENERIC: rewrite ( insn -- insn/f ) @@ -35,10 +41,7 @@ M: insn rewrite drop f ; : ##branch-t? ( insn -- ? ) dup ##compare-imm-branch? [ - { - [ cc>> cc/= eq? ] - [ src2>> \ f type-number eq? ] - } 1&& + { [ cc>> cc/= eq? ] [ src2>> not ] } 1&& ] [ drop f ] if ; inline : general-compare-expr? ( insn -- ? ) @@ -118,8 +121,8 @@ M: ##compare-imm rewrite-tagged-comparison : rewrite-redundant-comparison? ( insn -- ? ) { [ src1>> vreg>expr general-compare-expr? ] - [ src2>> \ f type-number = ] - [ cc>> { cc= cc/= } member-eq? ] + [ src2>> not ] + [ cc>> { cc= cc/= } member? ] } 1&& ; inline : rewrite-redundant-comparison ( insn -- insn' ) @@ -131,17 +134,12 @@ M: ##compare-imm rewrite-tagged-comparison } cond swap cc= eq? [ [ negate-cc ] change-cc ] when ; -ERROR: bad-comparison ; - : (fold-compare-imm) ( insn -- ? ) - [ [ src1>> vreg>constant ] [ src2>> ] bi ] [ cc>> ] bi - pick integer? - [ [ <=> ] dip evaluate-cc ] - [ - 2nip { - { cc= [ f ] } - { cc/= [ t ] } - [ bad-comparison ] + [ src1>> vreg>constant ] [ src2>> ] [ cc>> ] tri + 2over [ integer? ] both? [ [ <=> ] dip evaluate-cc ] [ + { + { cc= [ eq? ] } + { cc/= [ eq? not ] } } case ] if ; @@ -189,8 +187,8 @@ M: ##compare-imm-branch rewrite M: ##compare-branch rewrite { - { [ dup src1>> vreg-immediate-arithmetic? ] [ t >compare-imm-branch ] } - { [ dup src2>> vreg-immediate-arithmetic? ] [ f >compare-imm-branch ] } + { [ dup src1>> vreg-immediate-comparand? ] [ t >compare-imm-branch ] } + { [ dup src2>> vreg-immediate-comparand? ] [ f >compare-imm-branch ] } { [ dup self-compare? ] [ rewrite-self-compare-branch ] } [ drop f ] } cond ; @@ -209,19 +207,15 @@ M: ##compare-branch rewrite next-vreg \ ##compare-imm new-insn ; inline : >boolean-insn ( insn ? -- insn' ) - [ dst>> ] dip - { - { t [ t \ ##load-constant new-insn ] } - { f [ \ f type-number \ ##load-immediate new-insn ] } - } case ; + [ dst>> ] dip \ ##load-constant new-insn ; : rewrite-self-compare ( insn -- insn' ) dup (rewrite-self-compare) >boolean-insn ; M: ##compare rewrite { - { [ dup src1>> vreg-immediate-arithmetic? ] [ t >compare-imm ] } - { [ dup src2>> vreg-immediate-arithmetic? ] [ f >compare-imm ] } + { [ dup src1>> vreg-immediate-comparand? ] [ t >compare-imm ] } + { [ dup src2>> vreg-immediate-comparand? ] [ f >compare-imm ] } { [ dup self-compare? ] [ rewrite-self-compare ] } [ drop f ] } cond ; @@ -254,7 +248,12 @@ M: ##shl-imm constant-fold* drop shift ; : constant-fold ( insn -- insn' ) [ dst>> ] - [ [ src1>> vreg>constant ] [ src2>> ] [ ] tri constant-fold* ] bi + [ + [ src1>> vreg>constant \ f type-number or ] + [ src2>> ] + [ ] + tri constant-fold* + ] bi \ ##load-immediate new-insn ; inline : unary-constant-fold? ( insn -- ? ) @@ -380,7 +379,7 @@ M: ##sar-imm rewrite [ drop f ] } cond ; -: insn>imm-insn ( insn op swap? -- ) +: insn>imm-insn ( insn op swap? -- new-insn ) swap [ [ [ dst>> ] [ src1>> ] [ src2>> ] tri ] dip [ swap ] when vreg>constant @@ -390,13 +389,13 @@ M: ##sar-imm rewrite arithmetic-op? [ vreg-immediate-arithmetic? ] [ vreg-immediate-bitwise? ] if ; -: rewrite-arithmetic ( insn op -- ? ) +: rewrite-arithmetic ( insn op -- insn/f ) { { [ over src2>> over vreg-immediate? ] [ f insn>imm-insn ] } [ 2drop f ] } cond ; inline -: rewrite-arithmetic-commutative ( insn op -- ? ) +: rewrite-arithmetic-commutative ( insn op -- insn/f ) { { [ over src2>> over vreg-immediate? ] [ f insn>imm-insn ] } { [ over src1>> over vreg-immediate? ] [ t insn>imm-insn ] } diff --git a/basis/compiler/cfg/value-numbering/value-numbering-tests.factor b/basis/compiler/cfg/value-numbering/value-numbering-tests.factor index ac992ff98d..f835200702 100644 --- a/basis/compiler/cfg/value-numbering/value-numbering-tests.factor +++ b/basis/compiler/cfg/value-numbering/value-numbering-tests.factor @@ -4,7 +4,8 @@ cpu.architecture tools.test kernel math combinators.short-circuit accessors sequences compiler.cfg.predecessors locals compiler.cfg.dce compiler.cfg.ssa.destruction compiler.cfg.loop-detection compiler.cfg.representations compiler.cfg assocs vectors arrays -layouts literals namespaces alien compiler.cfg.value-numbering.simd ; +layouts literals namespaces alien compiler.cfg.value-numbering.simd +system ; IN: compiler.cfg.value-numbering.tests : trim-temps ( insns -- insns ) @@ -82,7 +83,7 @@ IN: compiler.cfg.value-numbering.tests T{ ##load-reference f 1 + } T{ ##peek f 2 D 0 } T{ ##compare f 4 2 1 cc> } - T{ ##compare-imm f 6 4 $[ \ f type-number ] cc/= } + T{ ##compare-imm f 6 4 f cc/= } T{ ##replace f 6 D 0 } } value-numbering-step trim-temps ] unit-test @@ -100,7 +101,7 @@ IN: compiler.cfg.value-numbering.tests T{ ##load-reference f 1 + } T{ ##peek f 2 D 0 } T{ ##compare f 4 2 1 cc<= } - T{ ##compare-imm f 6 4 $[ \ f type-number ] cc= } + T{ ##compare-imm f 6 4 f cc= } T{ ##replace f 6 D 0 } } value-numbering-step trim-temps ] unit-test @@ -118,7 +119,7 @@ IN: compiler.cfg.value-numbering.tests T{ ##peek f 8 D 0 } T{ ##peek f 9 D -1 } T{ ##compare-float-unordered f 12 8 9 cc< } - T{ ##compare-imm f 14 12 $[ \ f type-number ] cc= } + T{ ##compare-imm f 14 12 f cc= } T{ ##replace f 14 D 0 } } value-numbering-step trim-temps ] unit-test @@ -135,7 +136,7 @@ IN: compiler.cfg.value-numbering.tests T{ ##peek f 29 D -1 } T{ ##peek f 30 D -2 } T{ ##compare f 33 29 30 cc<= } - T{ ##compare-imm-branch f 33 $[ \ f type-number ] cc/= } + T{ ##compare-imm-branch f 33 f cc/= } } value-numbering-step trim-temps ] unit-test @@ -149,7 +150,7 @@ IN: compiler.cfg.value-numbering.tests { T{ ##peek f 1 D -1 } T{ ##test-vector f 2 1 f float-4-rep vcc-any } - T{ ##compare-imm-branch f 2 $[ \ f type-number ] cc/= } + T{ ##compare-imm-branch f 2 f cc/= } } value-numbering-step trim-temps ] unit-test @@ -418,6 +419,36 @@ IN: compiler.cfg.value-numbering.tests } value-numbering-step trim-temps ] unit-test +cpu x86.32? [ + [ + { + T{ ##peek f 0 D 0 } + T{ ##load-constant f 1 + } + T{ ##compare-imm f 2 0 + cc= } + } + ] [ + { + T{ ##peek f 0 D 0 } + T{ ##load-constant f 1 + } + T{ ##compare f 2 0 1 cc= } + } value-numbering-step trim-temps + ] unit-test + + [ + { + T{ ##peek f 0 D 0 } + T{ ##load-constant f 1 + } + T{ ##compare-imm-branch f 0 + cc= } + } + ] [ + { + T{ ##peek f 0 D 0 } + T{ ##load-constant f 1 + } + T{ ##compare-branch f 0 1 cc= } + } value-numbering-step trim-temps + ] unit-test +] when + [ { T{ ##peek f 0 D 0 } @@ -432,6 +463,20 @@ IN: compiler.cfg.value-numbering.tests } value-numbering-step trim-temps ] unit-test +[ + { + T{ ##peek f 0 D 0 } + T{ ##load-constant f 1 3.5 } + T{ ##compare-branch f 0 1 cc= } + } +] [ + { + T{ ##peek f 0 D 0 } + T{ ##load-constant f 1 3.5 } + T{ ##compare-branch f 0 1 cc= } + } value-numbering-step trim-temps +] unit-test + [ { T{ ##peek f 0 D 0 } @@ -460,20 +505,6 @@ IN: compiler.cfg.value-numbering.tests } value-numbering-step ] unit-test -[ - { - T{ ##peek f 0 D 0 } - T{ ##load-constant f 1 3.5 } - T{ ##compare-branch f 0 1 cc= } - } -] [ - { - T{ ##peek f 0 D 0 } - T{ ##load-constant f 1 3.5 } - T{ ##compare-branch f 0 1 cc= } - } value-numbering-step trim-temps -] unit-test - [ { T{ ##peek f 0 D 0 } @@ -488,6 +519,59 @@ IN: compiler.cfg.value-numbering.tests } value-numbering-step trim-temps ] unit-test +! Branch folding +[ + { + T{ ##load-immediate f 1 100 } + T{ ##load-immediate f 2 200 } + T{ ##load-constant f 3 t } + } +] [ + { + T{ ##load-immediate f 1 100 } + T{ ##load-immediate f 2 200 } + T{ ##compare f 3 1 2 cc<= } + } value-numbering-step trim-temps +] unit-test + +[ + { + T{ ##load-immediate f 1 100 } + T{ ##load-immediate f 2 200 } + T{ ##load-constant f 3 f } + } +] [ + { + T{ ##load-immediate f 1 100 } + T{ ##load-immediate f 2 200 } + T{ ##compare f 3 1 2 cc= } + } value-numbering-step trim-temps +] unit-test + +[ + { + T{ ##load-immediate f 1 100 } + T{ ##load-constant f 2 f } + } +] [ + { + T{ ##load-immediate f 1 100 } + T{ ##compare-imm f 2 1 f cc= } + } value-numbering-step trim-temps +] unit-test + +[ + { + T{ ##load-constant f 1 f } + T{ ##load-constant f 2 t } + } +] [ + { + T{ ##load-constant f 1 f } + T{ ##compare-imm f 2 1 f cc= } + } value-numbering-step trim-temps +] unit-test + ! Reassociation [ { @@ -1011,6 +1095,19 @@ cell 8 = [ } value-numbering-step ] unit-test +! Stupid constant folding corner case +[ + { + T{ ##load-constant f 1 f } + T{ ##load-immediate f 2 $[ \ f type-number ] } + } +] [ + { + T{ ##load-constant f 1 f } + T{ ##and-imm f 2 1 15 } + } value-numbering-step +] unit-test + ! Displaced alien optimizations 3 vreg-counter set-global @@ -1073,7 +1170,7 @@ cell 8 = [ { T{ ##load-immediate f 1 10 } T{ ##load-immediate f 2 20 } - T{ ##load-immediate f 3 $[ \ f type-number ] } + T{ ##load-constant f 3 f } } ] [ { @@ -1115,7 +1212,7 @@ cell 8 = [ { T{ ##load-immediate f 1 10 } T{ ##load-immediate f 2 20 } - T{ ##load-immediate f 3 $[ \ f type-number ] } + T{ ##load-constant f 3 f } } ] [ { @@ -1128,7 +1225,7 @@ cell 8 = [ [ { T{ ##peek f 0 D 0 } - T{ ##load-immediate f 1 $[ \ f type-number ] } + T{ ##load-constant f 1 f } } ] [ { @@ -1152,7 +1249,7 @@ cell 8 = [ [ { T{ ##peek f 0 D 0 } - T{ ##load-immediate f 1 $[ \ f type-number ] } + T{ ##load-constant f 1 f } } ] [ { @@ -1176,7 +1273,7 @@ cell 8 = [ [ { T{ ##peek f 0 D 0 } - T{ ##load-immediate f 1 $[ \ f type-number ] } + T{ ##load-constant f 1 f } } ] [ { @@ -1557,7 +1654,7 @@ cell 8 = [ { T{ ##peek f 0 D 0 } T{ ##compare f 1 0 0 cc<= } - T{ ##compare-imm-branch f 1 $[ \ f type-number ] cc/= } + T{ ##compare-imm-branch f 1 f cc/= } } test-branch-folding ] unit-test @@ -1659,7 +1756,7 @@ V{ T{ ##copy { dst 21 } { src 20 } { rep any-rep } } T{ ##compare-imm-branch { src1 21 } - { src2 $[ \ f type-number ] } + { src2 f } { cc cc/= } } } 1 test-bb diff --git a/basis/compiler/codegen/codegen.factor b/basis/compiler/codegen/codegen.factor index b16f471d11..99564b7e0e 100755 --- a/basis/compiler/codegen/codegen.factor +++ b/basis/compiler/codegen/codegen.factor @@ -81,6 +81,7 @@ SYNTAX: CODEGEN: CODEGEN: ##load-immediate %load-immediate CODEGEN: ##load-reference %load-reference CODEGEN: ##load-constant %load-reference +CODEGEN: ##load-double %load-double CODEGEN: ##peek %peek CODEGEN: ##replace %replace CODEGEN: ##inc-d %inc-d diff --git a/basis/compiler/codegen/fixup/fixup.factor b/basis/compiler/codegen/fixup/fixup.factor index eef517a2bb..fa8dfc2149 100644 --- a/basis/compiler/codegen/fixup/fixup.factor +++ b/basis/compiler/codegen/fixup/fixup.factor @@ -70,9 +70,12 @@ MEMO: cached-string>symbol ( symbol -- obj ) string>symbol ; : rel-word-pic-tail ( word class -- ) [ add-literal ] dip rt-entry-point-pic-tail rel-fixup ; -: rel-immediate ( literal class -- ) +: rel-literal ( literal class -- ) [ add-literal ] dip rt-literal rel-fixup ; +: rel-float ( literal class -- ) + [ add-literal ] dip rt-float rel-fixup ; + : rel-this ( class -- ) rt-this rel-fixup ; diff --git a/basis/compiler/constants/constants.factor b/basis/compiler/constants/constants.factor index 2fec5ca190..0e2fc3041b 100644 --- a/basis/compiler/constants/constants.factor +++ b/basis/compiler/constants/constants.factor @@ -68,7 +68,8 @@ C-ENUM: f rt-vm rt-cards-offset rt-decks-offset - rt-exception-handler ; + rt-exception-handler + rt-float ; : rc-absolute? ( n -- ? ) ${ diff --git a/basis/compiler/tests/low-level-ir.factor b/basis/compiler/tests/low-level-ir.factor index bc7f3fa2f2..5f00d251cf 100644 --- a/basis/compiler/tests/low-level-ir.factor +++ b/basis/compiler/tests/low-level-ir.factor @@ -33,10 +33,10 @@ IN: compiler.tests.low-level-ir compile-test-cfg execute( -- result ) ; -! loading immediates +! loading constants [ f ] [ V{ - T{ ##load-immediate f 0 $[ \ f type-number ] } + T{ ##load-constant f 0 f } } compile-test-bb ] unit-test diff --git a/basis/compiler/tree/propagation/recursive/recursive-tests.factor b/basis/compiler/tree/propagation/recursive/recursive-tests.factor index 42325d97ca..af2bdbda60 100644 --- a/basis/compiler/tree/propagation/recursive/recursive-tests.factor +++ b/basis/compiler/tree/propagation/recursive/recursive-tests.factor @@ -8,7 +8,7 @@ IN: compiler.tree.propagation.recursive.tests integer generalize-counter-interval ] unit-test -[ T{ interval f { 0 t } { $[ most-positive-fixnum ] t } } ] [ +[ T{ interval f { 0 t } { $[ max-array-capacity ] t } } ] [ T{ interval f { 1 t } { 1 t } } T{ interval f { 0 t } { 0 t } } fixnum generalize-counter-interval diff --git a/basis/compiler/tree/propagation/recursive/recursive.factor b/basis/compiler/tree/propagation/recursive/recursive.factor index d4ab697e21..854e730662 100644 --- a/basis/compiler/tree/propagation/recursive/recursive.factor +++ b/basis/compiler/tree/propagation/recursive/recursive.factor @@ -1,7 +1,7 @@ -! Copyright (C) 2008, 2009 Slava Pestov. +! Copyright (C) 2008, 2010 Slava Pestov. ! See http://factorcode.org/license.txt for BSD license. -USING: kernel sequences accessors arrays fry math math.intervals -layouts combinators namespaces locals +USING: kernel classes.algebra sequences accessors arrays fry +math math.intervals layouts combinators namespaces locals stack-checker.inlining compiler.tree compiler.tree.combinators @@ -11,6 +11,7 @@ compiler.tree.propagation.nodes compiler.tree.propagation.simple compiler.tree.propagation.branches compiler.tree.propagation.constraints ; +FROM: sequences.private => array-capacity ; IN: compiler.tree.propagation.recursive : check-fixed-point ( node infos1 infos2 -- ) @@ -24,7 +25,14 @@ IN: compiler.tree.propagation.recursive [ label>> calls>> [ node>> node-input-infos ] map flip ] [ latest-input-infos ] bi ; +: counter-class ( interval class -- class' ) + dup fixnum class<= [ + swap array-capacity-interval interval-subset? + [ drop array-capacity ] when + ] [ nip ] if ; + :: generalize-counter-interval ( interval initial-interval class -- interval' ) + interval class counter-class :> class { { [ interval initial-interval interval-subset? ] [ initial-interval ] } { [ interval empty-interval eq? ] [ initial-interval ] } diff --git a/basis/compiler/tree/propagation/transforms/transforms.factor b/basis/compiler/tree/propagation/transforms/transforms.factor index 4f0eea9cbb..f8d43e37c4 100644 --- a/basis/compiler/tree/propagation/transforms/transforms.factor +++ b/basis/compiler/tree/propagation/transforms/transforms.factor @@ -1,12 +1,13 @@ ! Copyright (C) 2008, 2010 Slava Pestov, Daniel Ehrenberg. ! See http://factorcode.org/license.txt for BSD license. -USING: alien.c-types kernel sequences words fry generic accessors -classes.tuple classes classes.algebra definitions -stack-checker.dependencies quotations classes.tuple.private math -math.partial-dispatch math.private math.intervals sets.private -math.floats.private math.integers.private layouts math.order -vectors hashtables combinators effects generalizations assocs -sets combinators.short-circuit sequences.private locals growable +USING: alien.c-types kernel sequences words fry generic +generic.single accessors classes.tuple classes classes.algebra +definitions stack-checker.dependencies quotations +classes.tuple.private math math.partial-dispatch math.private +math.intervals sets.private math.floats.private +math.integers.private layouts math.order vectors hashtables +combinators effects generalizations assocs sets +combinators.short-circuit sequences.private locals growable stack-checker namespaces compiler.tree.propagation.info ; FROM: math => float ; FROM: sets => set ; @@ -299,6 +300,12 @@ M\ set intersect [ intersect-quot ] 1 define-partial-eval [ \ push def>> ] [ f ] if ] "custom-inlining" set-word-prop +! Speeds up fasta benchmark +\ >fixnum [ + in-d>> first value-info class>> fixnum \ f class-or class<= + [ [ dup [ \ >fixnum no-method ] unless ] ] [ f ] if +] "custom-inlining" set-word-prop + ! We want to constant-fold calls to heap-size, and recompile those ! calls when a C type is redefined \ heap-size [ diff --git a/basis/cpu/architecture/architecture.factor b/basis/cpu/architecture/architecture.factor index 1aaf1bf2ea..a98b5cbafb 100644 --- a/basis/cpu/architecture/architecture.factor +++ b/basis/cpu/architecture/architecture.factor @@ -202,8 +202,9 @@ M: ulonglong-2-rep scalar-rep-of drop ulonglong-scalar-rep ; ! Mapping from register class to machine registers HOOK: machine-registers cpu ( -- assoc ) -HOOK: %load-immediate cpu ( reg obj -- ) +HOOK: %load-immediate cpu ( reg val -- ) HOOK: %load-reference cpu ( reg obj -- ) +HOOK: %load-double cpu ( reg val -- ) HOOK: %peek cpu ( vreg loc -- ) HOOK: %replace cpu ( vreg loc -- ) @@ -496,15 +497,32 @@ M: reg-class param-reg param-regs nth ; M: stack-params param-reg 2drop ; -! Is this integer small enough to be an immediate operand for -! %add-imm, %sub-imm, and %mul-imm? +! Does this architecture support %load-double? +HOOK: load-double? cpu ( -- ? ) + +M: object load-double? f ; + +! Can this value be an immediate operand for %add-imm, %sub-imm, +! or %mul-imm? HOOK: immediate-arithmetic? cpu ( n -- ? ) -! Is this integer small enough to be an immediate operand for -! %and-imm, %or-imm, and %xor-imm? +! Can this value be an immediate operand for %and-imm, %or-imm, +! or %xor-imm? HOOK: immediate-bitwise? cpu ( n -- ? ) -! What c-type describes the implicit struct return pointer for large structs? +! Can this value be an immediate operand for %compare-imm or +! %compare-imm-branch? +HOOK: immediate-comparand? cpu ( n -- ? ) + +M: object immediate-comparand? ( n -- ? ) + { + { [ dup integer? ] [ immediate-arithmetic? ] } + { [ dup not ] [ drop t ] } + [ drop f ] + } cond ; + +! What c-type describes the implicit struct return pointer for +! large structs? HOOK: struct-return-pointer-type cpu ( -- c-type ) ! Is this structure small enough to be returned in registers? diff --git a/basis/cpu/ppc/ppc.factor b/basis/cpu/ppc/ppc.factor index 551693d5c7..8adae2ae99 100644 --- a/basis/cpu/ppc/ppc.factor +++ b/basis/cpu/ppc/ppc.factor @@ -47,7 +47,7 @@ CONSTANT: fp-scratch-reg 30 M: ppc %load-immediate ( reg n -- ) swap LOAD ; M: ppc %load-reference ( reg obj -- ) - [ 0 swap LOAD32 ] [ rc-absolute-ppc-2/2 rel-immediate ] bi* ; + [ 0 swap LOAD32 ] [ rc-absolute-ppc-2/2 rel-literal ] bi* ; M: ppc %alien-global ( register symbol dll -- ) [ 0 swap LOAD32 ] 2dip rc-absolute-ppc-2/2 rel-dlsym ; @@ -492,7 +492,7 @@ M: ppc %epilogue ( n -- ) } case ; : (%compare) ( src1 src2 -- ) [ 0 ] dip CMP ; inline -: (%compare-imm) ( src1 src2 -- ) [ 0 ] 2dip CMPI ; inline +: (%compare-imm) ( src1 src2 -- ) [ 0 ] [ ] [ \ f type-number or ] tri* CMPI ; inline : (%compare-float-unordered) ( src1 src2 -- ) [ 0 ] dip FCMPU ; inline : (%compare-float-ordered) ( src1 src2 -- ) [ 0 ] dip FCMPO ; inline diff --git a/basis/cpu/x86/32/32.factor b/basis/cpu/x86/32/32.factor index 05c627fb99..c567c1e1f0 100755 --- a/basis/cpu/x86/32/32.factor +++ b/basis/cpu/x86/32/32.factor @@ -2,13 +2,13 @@ ! See http://factorcode.org/license.txt for BSD license. USING: locals alien alien.c-types alien.libraries alien.syntax arrays kernel fry math namespaces sequences system layouts io -vocabs.loader accessors init classes.struct combinators command-line -make compiler compiler.units compiler.constants compiler.alien -compiler.codegen compiler.codegen.fixup -compiler.cfg.instructions compiler.cfg.builder -compiler.cfg.intrinsics compiler.cfg.stack-frame -cpu.x86.assembler cpu.x86.assembler.operands cpu.x86 -cpu.architecture vm ; +vocabs.loader accessors init classes.struct combinators +command-line make words compiler compiler.units +compiler.constants compiler.alien compiler.codegen +compiler.codegen.fixup compiler.cfg.instructions +compiler.cfg.builder compiler.cfg.intrinsics +compiler.cfg.stack-frame cpu.x86.assembler +cpu.x86.assembler.operands cpu.x86 cpu.architecture vm ; FROM: layouts => cell ; IN: cpu.x86.32 @@ -24,6 +24,14 @@ M: x86.32 stack-reg ESP ; M: x86.32 frame-reg EBP ; M: x86.32 temp-reg ECX ; +M: x86.32 immediate-comparand? ( n -- ? ) + [ call-next-method ] [ word? ] bi or ; + +M: x86.32 load-double? ( -- ? ) t ; + +M: x86.32 %load-double ( dst val -- ) + [ 0 [] MOVSD ] dip rc-absolute rel-float ; + M: x86.32 %mov-vm-ptr ( reg -- ) 0 MOV 0 rc-absolute-cell rel-vm ; diff --git a/basis/cpu/x86/x86.factor b/basis/cpu/x86/x86.factor index 028cca48e3..7bb33dec9a 100644 --- a/basis/cpu/x86/x86.factor +++ b/basis/cpu/x86/x86.factor @@ -66,7 +66,7 @@ HOOK: pic-tail-reg cpu ( -- reg ) M: x86 %load-immediate dup 0 = [ drop dup XOR ] [ MOV ] if ; -M: x86 %load-reference swap 0 MOV rc-absolute-cell rel-immediate ; +M: x86 %load-reference swap 0 MOV rc-absolute-cell rel-literal ; HOOK: ds-reg cpu ( -- reg ) HOOK: rs-reg cpu ( -- reg ) @@ -491,43 +491,60 @@ M: x86 %push-context-stack ( -- ) M: x86 %epilogue ( n -- ) cell - incr-stack-reg ; -:: %boolean ( dst temp word -- ) +:: (%boolean) ( dst temp insn -- ) dst \ f type-number MOV - temp 0 MOV \ t rc-absolute-cell rel-immediate - dst temp word execute ; inline + temp 0 MOV \ t rc-absolute-cell rel-literal + dst temp insn execute ; inline -: (%compare) ( src1 src2 cc -- ) - 2over [ { cc= cc/= } member? ] [ register? ] [ 0 = ] tri* and and - [ drop dup TEST ] - [ CMP ] if ; +: %boolean ( dst cc temp -- ) + swap order-cc { + { cc< [ \ CMOVL (%boolean) ] } + { cc<= [ \ CMOVLE (%boolean) ] } + { cc> [ \ CMOVG (%boolean) ] } + { cc>= [ \ CMOVGE (%boolean) ] } + { cc= [ \ CMOVE (%boolean) ] } + { cc/= [ \ CMOVNE (%boolean) ] } + } case ; M:: x86 %compare ( dst src1 src2 cc temp -- ) - src1 src2 cc (%compare) - cc order-cc { - { cc< [ dst temp \ CMOVL %boolean ] } - { cc<= [ dst temp \ CMOVLE %boolean ] } - { cc> [ dst temp \ CMOVG %boolean ] } - { cc>= [ dst temp \ CMOVGE %boolean ] } - { cc= [ dst temp \ CMOVE %boolean ] } - { cc/= [ dst temp \ CMOVNE %boolean ] } - } case ; + src1 src2 CMP + dst cc temp %boolean ; -M: x86 %compare-imm ( dst src1 src2 cc temp -- ) - %compare ; +: use-test? ( src1 src2 cc -- ? ) + [ register? ] [ 0 = ] [ { cc= cc/= } member? ] tri* and and ; + +: (%compare-tagged) ( src1 src2 -- ) + [ HEX: ffffffff CMP ] dip rc-absolute rel-literal ; + +: (%compare-imm) ( src1 src2 cc -- ) + { + { [ 3dup use-test? ] [ 2drop dup TEST ] } + { [ over integer? ] [ drop CMP ] } + { [ over word? ] [ drop (%compare-tagged) ] } + { [ over not ] [ 2drop \ f type-number CMP ] } + } cond ; + +M:: x86 %compare-imm ( dst src1 src2 cc temp -- ) + src1 src2 cc (%compare-imm) + dst cc temp %boolean ; + +: %branch ( label cc -- ) + order-cc { + { cc< [ JL ] } + { cc<= [ JLE ] } + { cc> [ JG ] } + { cc>= [ JGE ] } + { cc= [ JE ] } + { cc/= [ JNE ] } + } case ; M:: x86 %compare-branch ( label src1 src2 cc -- ) - src1 src2 cc (%compare) - cc order-cc { - { cc< [ label JL ] } - { cc<= [ label JLE ] } - { cc> [ label JG ] } - { cc>= [ label JGE ] } - { cc= [ label JE ] } - { cc/= [ label JNE ] } - } case ; + src1 src2 CMP + label cc %branch ; -M: x86 %compare-imm-branch ( label src1 src2 cc -- ) - %compare-branch ; +M:: x86 %compare-imm-branch ( label src1 src2 cc -- ) + src1 src2 cc (%compare-imm) + label cc %branch ; M: x86 %add-float double-rep two-operand ADDSD ; M: x86 %sub-float double-rep two-operand SUBSD ; @@ -569,20 +586,20 @@ M: x86 %float>integer CVTTSD2SI ; :: (%compare-float) ( dst src1 src2 cc temp compare -- ) cc { - { cc< [ src2 src1 \ compare execute( a b -- ) dst temp \ CMOVA %boolean ] } - { cc<= [ src2 src1 \ compare execute( a b -- ) dst temp \ CMOVAE %boolean ] } - { cc> [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVA %boolean ] } - { cc>= [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVAE %boolean ] } - { cc= [ src1 src2 \ compare execute( a b -- ) dst temp \ %cmov-float= %boolean ] } - { cc<> [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVNE %boolean ] } - { cc<>= [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVNP %boolean ] } - { cc/< [ src2 src1 \ compare execute( a b -- ) dst temp \ CMOVBE %boolean ] } - { cc/<= [ src2 src1 \ compare execute( a b -- ) dst temp \ CMOVB %boolean ] } - { cc/> [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVBE %boolean ] } - { cc/>= [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVB %boolean ] } - { cc/= [ src1 src2 \ compare execute( a b -- ) dst temp \ %cmov-float/= %boolean ] } - { cc/<> [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVE %boolean ] } - { cc/<>= [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVP %boolean ] } + { cc< [ src2 src1 \ compare execute( a b -- ) dst temp \ CMOVA (%boolean) ] } + { cc<= [ src2 src1 \ compare execute( a b -- ) dst temp \ CMOVAE (%boolean) ] } + { cc> [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVA (%boolean) ] } + { cc>= [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVAE (%boolean) ] } + { cc= [ src1 src2 \ compare execute( a b -- ) dst temp \ %cmov-float= (%boolean) ] } + { cc<> [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVNE (%boolean) ] } + { cc<>= [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVNP (%boolean) ] } + { cc/< [ src2 src1 \ compare execute( a b -- ) dst temp \ CMOVBE (%boolean) ] } + { cc/<= [ src2 src1 \ compare execute( a b -- ) dst temp \ CMOVB (%boolean) ] } + { cc/> [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVBE (%boolean) ] } + { cc/>= [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVB (%boolean) ] } + { cc/= [ src1 src2 \ compare execute( a b -- ) dst temp \ %cmov-float/= (%boolean) ] } + { cc/<> [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVE (%boolean) ] } + { cc/<>= [ src1 src2 \ compare execute( a b -- ) dst temp \ CMOVP (%boolean) ] } } case ; inline M: x86 %compare-float-ordered ( dst src1 src2 cc temp -- ) @@ -954,10 +971,10 @@ M: x86 %compare-vector-ccs :: %test-vector-mask ( dst temp mask vcc -- ) vcc { - { vcc-any [ dst dst TEST dst temp \ CMOVNE %boolean ] } - { vcc-none [ dst dst TEST dst temp \ CMOVE %boolean ] } - { vcc-all [ dst mask CMP dst temp \ CMOVE %boolean ] } - { vcc-notall [ dst mask CMP dst temp \ CMOVNE %boolean ] } + { vcc-any [ dst dst TEST dst temp \ CMOVNE (%boolean) ] } + { vcc-none [ dst dst TEST dst temp \ CMOVE (%boolean) ] } + { vcc-all [ dst mask CMP dst temp \ CMOVE (%boolean) ] } + { vcc-notall [ dst mask CMP dst temp \ CMOVNE (%boolean) ] } } case ; : %move-vector-mask ( dst src rep -- mask ) diff --git a/basis/hints/hints.factor b/basis/hints/hints.factor index 558f7dd8a4..dc16cf8b24 100644 --- a/basis/hints/hints.factor +++ b/basis/hints/hints.factor @@ -2,10 +2,10 @@ ! See http://factorcode.org/license.txt for BSD license. USING: accessors arrays assocs byte-arrays byte-vectors classes combinators definitions effects fry generic generic.single -generic.standard hashtables io.binary io.streams.string kernel -kernel.private math math.integers.private math.parser -namespaces parser sbufs sequences splitting splitting.private strings -vectors words ; +generic.standard hashtables io.binary io.encodings +io.streams.string kernel kernel.private math +math.integers.private math.parser namespaces parser sbufs +sequences splitting splitting.private strings vectors words ; IN: hints GENERIC: specializer-predicate ( spec -- quot ) @@ -131,3 +131,5 @@ M\ hashtable at* { { fixnum object } { word object } } "specializer" set-word-pr M\ hashtable set-at { { object fixnum object } { object word object } } "specializer" set-word-prop \ bignum/f { { bignum bignum } { bignum fixnum } { fixnum bignum } { fixnum fixnum } } "specializer" set-word-prop + +\ encode-string { string object object } "specializer" set-word-prop diff --git a/basis/http/client/client.factor b/basis/http/client/client.factor index 1221ee39f3..aa2fc8962b 100644 --- a/basis/http/client/client.factor +++ b/basis/http/client/client.factor @@ -194,6 +194,6 @@ ERROR: download-failed response ; : http-delete ( url -- response data ) http-request ; -USING: vocabs vocabs.loader ; +USE: vocabs.loader -"debugger" "http.client.debugger" require-when +{ "http.client" "debugger" } "http.client.debugger" require-when diff --git a/basis/io/encodings/8-bit/8-bit.factor b/basis/io/encodings/8-bit/8-bit.factor index 7f92028c31..db269c319d 100644 --- a/basis/io/encodings/8-bit/8-bit.factor +++ b/basis/io/encodings/8-bit/8-bit.factor @@ -1,10 +1,10 @@ ! Copyright (C) 2008 Daniel Ehrenberg, Doug Coleman. ! See http://factorcode.org/license.txt for BSD license. -USING: math.parser arrays io.encodings sequences kernel assocs -hashtables io.encodings.ascii generic parser classes.tuple words -words.symbol io io.files splitting namespaces math -compiler.units accessors classes.singleton classes.mixin -io.encodings.iana fry simple-flat-file lexer ; +USING: arrays assocs biassocs kernel io.encodings math.parser +sequences hashtables io.encodings.ascii generic parser +classes.tuple words words.symbol io io.files splitting +namespaces math compiler.units accessors classes.singleton +classes.mixin io.encodings.iana fry simple-flat-file lexer ; IN: io.encodings.8-bit > value-at [ encode-error ] unless* ; inline -M: 8-bit encode-char biassoc>> encode-8-bit ; +M: 8-bit encode-char + swap [ 8-bit-encode ] dip stream-write1 ; -: decode-8-bit ( stream assoc -- char/f ) - swap stream-read1 - [ swap at [ replacement-char ] unless* ] - [ drop f ] if* ; inline +M: 8-bit encode-string + swap [ '[ _ 8-bit-encode ] B{ } map-as ] dip stream-write ; -M: 8-bit decode-char biassoc>> decode-8-bit ; +M: 8-bit decode-char + swap stream-read1 dup + [ swap biassoc>> at [ replacement-char ] unless* ] + [ 2drop f ] + if ; MIXIN: 8-bit-encoding diff --git a/basis/io/encodings/ascii/ascii.factor b/basis/io/encodings/ascii/ascii.factor index 00d3bc7509..2b5640489f 100644 --- a/basis/io/encodings/ascii/ascii.factor +++ b/basis/io/encodings/ascii/ascii.factor @@ -1,22 +1,27 @@ ! Copyright (C) 2008 Daniel Ehrenberg. ! See http://factorcode.org/license.txt for BSD license. -USING: io io.encodings kernel math io.encodings.private ; +USING: accessors byte-arrays io io.encodings +io.encodings.private kernel math sequences ; IN: io.encodings.ascii - ] 2bi [ >fixnum ] [ drop replacement-char ] if ] - [ 2drop f ] if ; inline -PRIVATE> - SINGLETON: ascii M: ascii encode-char - 128 encode-if< ; inline + drop + over 127 <= [ stream-write1 ] [ encode-error ] if ; inline + +M: ascii encode-string + drop + [ + dup aux>> + [ [ dup 127 <= [ encode-error ] unless ] B{ } map-as ] + [ >byte-array ] + if + ] dip + stream-write ; M: ascii decode-char - 128 decode-if< ; inline + drop + stream-read1 dup [ + dup 127 <= [ >fixnum ] [ drop replacement-char ] if + ] when ; inline diff --git a/basis/io/ports/ports.factor b/basis/io/ports/ports.factor index 0927e7e480..cd0843a70b 100644 --- a/basis/io/ports/ports.factor +++ b/basis/io/ports/ports.factor @@ -114,7 +114,7 @@ M: output-port stream-write1 : write-in-groups ( byte-array port -- ) [ binary-object ] dip - [ buffer>> size>> ] [ '[ _ stream-write ] ] bi + [ buffer>> size>> ] [ '[ _ stream-write ] ] bi each ; M: output-port stream-write @@ -198,5 +198,3 @@ io.encodings.private ; HINTS: decoder-read-until { string input-port utf8 } { string input-port ascii } ; HINTS: decoder-readln { input-port utf8 } { input-port ascii } ; - -HINTS: encoder-write { object output-port utf8 } { object output-port ascii } ; diff --git a/basis/io/streams/byte-array/fast/authors.txt b/basis/io/streams/byte-array/fast/authors.txt new file mode 100644 index 0000000000..1901f27a24 --- /dev/null +++ b/basis/io/streams/byte-array/fast/authors.txt @@ -0,0 +1 @@ +Slava Pestov diff --git a/basis/io/streams/byte-array/fast/fast.factor b/basis/io/streams/byte-array/fast/fast.factor new file mode 100644 index 0000000000..e231335bfd --- /dev/null +++ b/basis/io/streams/byte-array/fast/fast.factor @@ -0,0 +1,15 @@ +! Copyright (C) 2010 Slava Pestov. +! See http://factorcode.org/license.txt for BSD license. +USING: alien byte-vectors io kernel libc math sequences ; +IN: io.streams.byte-array.fast + +! This is split off from io.streams.byte-array because it uses +! memcpy, which is a non-core word that only works after the +! optimizing compiler has been loaded. + +M: byte-vector stream-write + [ dup byte-length tail-slice ] + [ [ [ byte-length ] bi@ + ] keep lengthen ] + [ drop byte-length ] + 2tri + [ >c-ptr swap >c-ptr ] dip memcpy ; diff --git a/basis/locals/locals.factor b/basis/locals/locals.factor index 7d67881c47..5fd12e2fb3 100644 --- a/basis/locals/locals.factor +++ b/basis/locals/locals.factor @@ -26,5 +26,5 @@ SYNTAX: MEMO:: (::) define-memoized ; "locals.fry" } [ require ] each -"prettyprint" "locals.definitions" require-when -"prettyprint" "locals.prettyprint" require-when +{ "locals" "prettyprint" } "locals.definitions" require-when +{ "locals" "prettyprint" } "locals.prettyprint" require-when diff --git a/basis/math/rectangles/rectangles.factor b/basis/math/rectangles/rectangles.factor index 78ac5457bc..15f4d5376d 100644 --- a/basis/math/rectangles/rectangles.factor +++ b/basis/math/rectangles/rectangles.factor @@ -64,4 +64,4 @@ M: rect contains-point? USE: vocabs.loader -"prettyprint" "math.rectangles.prettyprint" require-when +{ "math.rectangles" "prettyprint" } "math.rectangles.prettyprint" require-when diff --git a/basis/math/vectors/simd/simd.factor b/basis/math/vectors/simd/simd.factor index 65d6e113bf..c845a4df63 100644 --- a/basis/math/vectors/simd/simd.factor +++ b/basis/math/vectors/simd/simd.factor @@ -339,4 +339,4 @@ M: short-8 v*hs+ M: int-4 v*hs+ int-4-rep [ (simd-v*hs+) ] [ call-next-method ] vv->v-op longlong-2-cast ; inline -"mirrors" "math.vectors.simd.mirrors" require-when +{ "math.vectors.simd" "mirrors" } "math.vectors.simd.mirrors" require-when diff --git a/basis/peg/peg.factor b/basis/peg/peg.factor index ca7d28bb97..e50c1d8d95 100644 --- a/basis/peg/peg.factor +++ b/basis/peg/peg.factor @@ -628,6 +628,6 @@ SYNTAX: PEG: ] append! ] ; -USING: vocabs vocabs.loader ; +USE: vocabs.loader -"debugger" "peg.debugger" require-when +{ "debugger" "peg" } "peg.debugger" require-when diff --git a/basis/regexp/regexp.factor b/basis/regexp/regexp.factor index eea0a26ea5..bbfe440967 100644 --- a/basis/regexp/regexp.factor +++ b/basis/regexp/regexp.factor @@ -216,6 +216,6 @@ SYNTAX: R` CHAR: ` parsing-regexp ; SYNTAX: R{ CHAR: } parsing-regexp ; SYNTAX: R| CHAR: | parsing-regexp ; -USING: vocabs vocabs.loader ; +USE: vocabs.loader -"prettyprint" "regexp.prettyprint" require-when +{ "prettyprint" "regexp" } "regexp.prettyprint" require-when diff --git a/basis/specialized-arrays/specialized-arrays.factor b/basis/specialized-arrays/specialized-arrays.factor index c82ebd78c8..38f97303ba 100644 --- a/basis/specialized-arrays/specialized-arrays.factor +++ b/basis/specialized-arrays/specialized-arrays.factor @@ -173,6 +173,6 @@ SYNTAX: SPECIALIZED-ARRAYS: SYNTAX: SPECIALIZED-ARRAY: scan-c-type define-array-vocab use-vocab ; -"prettyprint" "specialized-arrays.prettyprint" require-when +{ "specialized-arrays" "prettyprint" } "specialized-arrays.prettyprint" require-when -"mirrors" "specialized-arrays.mirrors" require-when +{ "specialized-arrays" "mirrors" } "specialized-arrays.mirrors" require-when diff --git a/basis/stack-checker/errors/errors.factor b/basis/stack-checker/errors/errors.factor index 5eca37ffbe..f3aeb7bb64 100644 --- a/basis/stack-checker/errors/errors.factor +++ b/basis/stack-checker/errors/errors.factor @@ -35,4 +35,4 @@ ERROR: bad-declaration-error < inference-error declaration ; ERROR: unbalanced-branches-error < inference-error word quots declareds actuals ; -"debugger" "stack-checker.errors.prettyprint" require-when +{ "stack-checker.errors" "debugger" } "stack-checker.errors.prettyprint" require-when diff --git a/basis/stack-checker/known-words/known-words.factor b/basis/stack-checker/known-words/known-words.factor index 1fa9a94677..c0d4b6c543 100644 --- a/basis/stack-checker/known-words/known-words.factor +++ b/basis/stack-checker/known-words/known-words.factor @@ -349,6 +349,7 @@ M: bad-executable summary \ both-fixnums? { object object } { object } define-primitive \ byte-array>bignum { byte-array } { bignum } define-primitive \ byte-array>bignum make-foldable \ callstack { } { callstack } define-primitive \ callstack make-flushable +\ callstack-bounds { } { alien alien } define-primitive \ callstack-bounds make-flushable \ callstack-for { c-ptr } { callstack } define-primitive \ callstack make-flushable \ callstack>array { callstack } { array } define-primitive \ callstack>array make-flushable \ check-datastack { array integer integer } { object } define-primitive \ check-datastack make-flushable diff --git a/basis/typed/typed.factor b/basis/typed/typed.factor index df46303b79..65b21fcc38 100644 --- a/basis/typed/typed.factor +++ b/basis/typed/typed.factor @@ -164,6 +164,6 @@ SYNTAX: TYPED: SYNTAX: TYPED:: (::) define-typed ; -USING: vocabs vocabs.loader ; +USE: vocabs.loader -"prettyprint" "typed.prettyprint" require-when +{ "typed" "prettyprint" } "typed.prettyprint" require-when diff --git a/basis/ui/gadgets/gadgets.factor b/basis/ui/gadgets/gadgets.factor index dca340cd3b..3c1ece1f5e 100644 --- a/basis/ui/gadgets/gadgets.factor +++ b/basis/ui/gadgets/gadgets.factor @@ -395,4 +395,4 @@ M: f request-focus-on 2drop ; USE: vocabs.loader -"prettyprint" "ui.gadgets.prettyprint" require-when +{ "ui.gadgets" "prettyprint" } "ui.gadgets.prettyprint" require-when diff --git a/basis/unix/unix.factor b/basis/unix/unix.factor index dbbfbcce6e..d860bf490e 100644 --- a/basis/unix/unix.factor +++ b/basis/unix/unix.factor @@ -72,6 +72,6 @@ M: unix open-file [ open ] unix-system-call ; << -"debugger" "unix.debugger" require-when +{ "unix" "debugger" } "unix.debugger" require-when >> diff --git a/basis/urls/urls.factor b/basis/urls/urls.factor index cd470a451a..0f89ba0d9f 100644 --- a/basis/urls/urls.factor +++ b/basis/urls/urls.factor @@ -185,4 +185,4 @@ SYNTAX: URL" lexer get skip-blank parse-string >url suffix! ; USE: vocabs.loader -"prettyprint" "urls.prettyprint" require-when +{ "urls" "prettyprint" } "urls.prettyprint" require-when diff --git a/basis/windows/com/syntax/syntax.factor b/basis/windows/com/syntax/syntax.factor index 9d74ac49f8..dc6a0604fb 100644 --- a/basis/windows/com/syntax/syntax.factor +++ b/basis/windows/com/syntax/syntax.factor @@ -96,4 +96,4 @@ SYNTAX: GUID: scan string>guid suffix! ; USE: vocabs.loader -"prettyprint" "windows.com.prettyprint" require-when +{ "windows.com" "prettyprint" } "windows.com.prettyprint" require-when diff --git a/basis/x11/x11.factor b/basis/x11/x11.factor index e91c6a6909..67c94c88ea 100644 --- a/basis/x11/x11.factor +++ b/basis/x11/x11.factor @@ -33,4 +33,4 @@ SYMBOL: root : with-x ( display-string quot -- ) [ init-x ] dip [ close-x ] [ ] cleanup ; inline -"io.backend.unix" "x11.io.unix" require-when +{ "x11" "io.backend.unix" } "x11.io.unix" require-when diff --git a/basis/xml/syntax/syntax.factor b/basis/xml/syntax/syntax.factor index a58526faa3..e7e8714b29 100644 --- a/basis/xml/syntax/syntax.factor +++ b/basis/xml/syntax/syntax.factor @@ -177,4 +177,4 @@ SYNTAX: [XML USE: vocabs.loader -"inverse" "xml.syntax.inverse" require-when +{ "xml.syntax" "inverse" } "xml.syntax.inverse" require-when diff --git a/core/bootstrap/primitives.factor b/core/bootstrap/primitives.factor index c466b0c1f8..27699725f1 100644 --- a/core/bootstrap/primitives.factor +++ b/core/bootstrap/primitives.factor @@ -451,6 +451,7 @@ tuple { "retainstack" "kernel" "primitive_retainstack" (( -- array )) } { "(identity-hashcode)" "kernel.private" "primitive_identity_hashcode" (( obj -- code )) } { "become" "kernel.private" "primitive_become" (( old new -- )) } + { "callstack-bounds" "kernel.private" "primitive_callstack_bounds" (( -- start end )) } { "check-datastack" "kernel.private" "primitive_check_datastack" (( array in# out# -- ? )) } { "compute-identity-hashcode" "kernel.private" "primitive_compute_identity_hashcode" (( obj -- )) } { "context-object" "kernel.private" "primitive_context_object" (( n -- obj )) } diff --git a/core/hash-sets/hash-sets-tests.factor b/core/hash-sets/hash-sets-tests.factor index 5b7ffafc8b..ca995a38e6 100644 --- a/core/hash-sets/hash-sets-tests.factor +++ b/core/hash-sets/hash-sets-tests.factor @@ -31,3 +31,6 @@ IN: hash-sets.tests [ f ] [ HS{ 1 2 3 } HS{ 2 3 } set= ] unit-test [ HS{ 1 2 } HS{ 1 2 3 } ] [ HS{ 1 2 } clone dup clone [ 3 swap adjoin ] keep ] unit-test + +[ t ] [ HS{ } null? ] unit-test +[ f ] [ HS{ 1 } null? ] unit-test diff --git a/core/hash-sets/hash-sets.factor b/core/hash-sets/hash-sets.factor index 3ca2cce93c..ac198a2ca2 100644 --- a/core/hash-sets/hash-sets.factor +++ b/core/hash-sets/hash-sets.factor @@ -18,6 +18,7 @@ M: hash-set delete table>> delete-at ; inline M: hash-set members table>> keys ; inline M: hash-set set-like drop dup hash-set? [ members ] unless ; M: hash-set clone table>> clone hash-set boa ; +M: hash-set null? table>> assoc-empty? ; M: sequence fast-set ; M: f fast-set drop H{ } clone hash-set boa ; diff --git a/core/io/encodings/encodings.factor b/core/io/encodings/encodings.factor index 03e8723d20..1880859db1 100644 --- a/core/io/encodings/encodings.factor +++ b/core/io/encodings/encodings.factor @@ -1,4 +1,4 @@ -! Copyright (C) 2008 Daniel Ehrenberg. +! Copyright (C) 2008, 2010 Daniel Ehrenberg, Slava Pestov. ! See http://factorcode.org/license.txt for BSD license. USING: math kernel sequences sbufs vectors namespaces growable strings io classes continuations destructors combinators @@ -12,6 +12,10 @@ GENERIC: decode-char ( stream encoding -- char/f ) GENERIC: encode-char ( char stream encoding -- ) +GENERIC: encode-string ( string stream encoding -- ) + +M: object encode-string [ encode-char ] 2curry each ; inline + GENERIC: ( stream encoding -- newstream ) CONSTANT: replacement-char HEX: fffd @@ -134,13 +138,8 @@ M: encoder stream-element-type M: encoder stream-write1 >encoder< encode-char ; -GENERIC# encoder-write 2 ( string stream encoding -- ) - -M: string encoder-write - [ encode-char ] 2curry each ; - M: encoder stream-write - >encoder< encoder-write ; + >encoder< encode-string ; M: encoder dispose stream>> dispose ; diff --git a/core/io/encodings/utf8/utf8.factor b/core/io/encodings/utf8/utf8.factor index 2911385c09..c78a86c072 100644 --- a/core/io/encodings/utf8/utf8.factor +++ b/core/io/encodings/utf8/utf8.factor @@ -1,7 +1,8 @@ ! Copyright (C) 2006, 2008 Daniel Ehrenberg. ! See http://factorcode.org/license.txt for BSD license. -USING: math math.order kernel sequences sbufs vectors growable io -continuations namespaces io.encodings combinators strings ; +USING: accessors byte-arrays math math.order kernel sequences +sbufs vectors growable io continuations namespaces io.encodings +combinators strings ; IN: io.encodings.utf8 ! Decoding UTF-8 @@ -45,10 +46,10 @@ M: utf8 decode-char ! Encoding UTF-8 : encoded ( stream char -- ) - BIN: 111111 bitand BIN: 10000000 bitor swap stream-write1 ; + BIN: 111111 bitand BIN: 10000000 bitor swap stream-write1 ; inline -: char>utf8 ( stream char -- ) - { +: char>utf8 ( char stream -- ) + swap { { [ dup -7 shift zero? ] [ swap stream-write1 ] } { [ dup -11 shift zero? ] [ 2dup -6 shift BIN: 11000000 bitor swap stream-write1 @@ -65,10 +66,16 @@ M: utf8 decode-char 2dup -6 shift encoded encoded ] - } cond ; + } cond ; inline M: utf8 encode-char - drop swap char>utf8 ; + drop char>utf8 ; + +M: utf8 encode-string + drop + over aux>> + [ [ char>utf8 ] curry each ] + [ [ >byte-array ] dip stream-write ] if ; PRIVATE> diff --git a/core/sets/sets-docs.factor b/core/sets/sets-docs.factor index 5bde8a1feb..5ae9641734 100644 --- a/core/sets/sets-docs.factor +++ b/core/sets/sets-docs.factor @@ -23,6 +23,8 @@ ARTICLE: "set-operations" "Operations on sets" adjoin delete } +"To test if a set is the empty set:" +{ $subsections null? } "Basic mathematical operations, which any type of set may override for efficiency:" { $subsections diff @@ -178,3 +180,7 @@ HELP: within HELP: without { $values { "seq" sequence } { "set" set } { "subseq" sequence } } { $description "Returns the subsequence of the given sequence consisting of things that are not members of the set. This may contain duplicates, if the sequence has duplicates." } ; + +HELP: null? +{ $values { "set" set } { "?" "a boolean" } } +{ $description "Tests whether the given set is empty. This outputs " { $snippet "t" } " when given a null set of any type." } ; diff --git a/core/sets/sets-tests.factor b/core/sets/sets-tests.factor index e4bc762512..9a48acc4cf 100644 --- a/core/sets/sets-tests.factor +++ b/core/sets/sets-tests.factor @@ -61,3 +61,6 @@ IN: sets.tests [ f ] [ HS{ 1 2 3 1 2 1 } duplicates ] unit-test [ H{ { 3 HS{ 1 2 } } } ] [ H{ } clone 1 3 pick adjoin-at 2 3 pick adjoin-at ] unit-test + +[ t ] [ f null? ] unit-test +[ f ] [ { 4 } null? ] unit-test diff --git a/core/sets/sets.factor b/core/sets/sets.factor index d279f036d4..9c1870aa2e 100644 --- a/core/sets/sets.factor +++ b/core/sets/sets.factor @@ -21,10 +21,13 @@ GENERIC: subset? ( set1 set2 -- ? ) GENERIC: set= ( set1 set2 -- ? ) GENERIC: duplicates ( set -- seq ) GENERIC: all-unique? ( set -- ? ) +GENERIC: null? ( set -- ? ) ! Defaults for some methods. ! Override them for efficiency +M: set null? members null? ; inline + M: set set-like drop ; inline M: set union @@ -91,6 +94,9 @@ M: sequence set-like M: sequence members [ pruned ] keep like ; + +M: sequence null? + empty? ; inline : combine ( sets -- set ) [ f ] diff --git a/core/strings/strings.factor b/core/strings/strings.factor index 18af08b3f6..50d79a4d8a 100644 --- a/core/strings/strings.factor +++ b/core/strings/strings.factor @@ -1,8 +1,7 @@ ! Copyright (C) 2003, 2008 Slava Pestov. ! See http://factorcode.org/license.txt for BSD license. USING: accessors kernel math.private sequences kernel.private -math sequences.private slots.private byte-arrays -alien.accessors ; +math sequences.private slots.private alien.accessors ; IN: strings >source-loaded? ] dip [ % ] [ call( -- ) ] if-bootstrapping +done+ >>source-loaded? - vocab-name load-conditional-requires + load-conditional-requires ] [ ] [ f >>source-loaded? ] cleanup ; : load-docs ( vocab -- ) @@ -97,10 +106,12 @@ PRIVATE> load-vocab drop ; : require-when ( if then -- ) - over vocab - [ nip require ] - [ swap conditional-requires get [ swap suffix ] change-at ] - if ; + over [ vocab ] all? [ + require drop + ] [ + [ drop [ require-when-vocabs get adjoin ] each ] + [ 2array require-when-table get push ] 2bi + ] if ; : reload ( name -- ) dup vocab diff --git a/core/vocabs/loader/test/m/m.factor b/core/vocabs/loader/test/m/m.factor index d6d3bd8a7a..cd35d83e4f 100644 --- a/core/vocabs/loader/test/m/m.factor +++ b/core/vocabs/loader/test/m/m.factor @@ -1,4 +1,5 @@ USE: vocabs.loader IN: vocabs.loader.test.m -"vocabs.loader.test.o" "vocabs.loader.test.n" require-when +{ "vocabs.loader.test.o" "vocabs.loader.test.m" } +"vocabs.loader.test.n" require-when diff --git a/core/vocabs/vocabs.factor b/core/vocabs/vocabs.factor index e48d6c3031..38881673e9 100644 --- a/core/vocabs/vocabs.factor +++ b/core/vocabs/vocabs.factor @@ -1,7 +1,7 @@ ! Copyright (C) 2007, 2009 Eduardo Cavazos, Slava Pestov. ! See http://factorcode.org/license.txt for BSD license. USING: accessors assocs strings kernel sorting namespaces -sequences definitions sets ; +sequences definitions sets combinators ; IN: vocabs SYMBOL: dictionary @@ -83,9 +83,6 @@ ERROR: bad-vocab-name name ; : check-vocab-name ( name -- name ) dup string? [ bad-vocab-name ] unless ; -SYMBOL: conditional-requires -conditional-requires [ H{ } clone ] initialize - : create-vocab ( name -- vocab ) check-vocab-name dictionary get [ ] cache diff --git a/extra/benchmark/fasta/fasta.factor b/extra/benchmark/fasta/fasta.factor index 8c06716ddb..f1ebc2aa9f 100644 --- a/extra/benchmark/fasta/fasta.factor +++ b/extra/benchmark/fasta/fasta.factor @@ -91,10 +91,13 @@ TYPED:: make-repeat-fasta ( k: fixnum len: fixnum alu: string -- k': fixnum ) n 2 * ALU "Homo sapiens alu" "ONE" write-repeat-fasta initial-seed + n 3 * homo-sapiens-chars homo-sapiens-floats "IUB ambiguity codes" "TWO" write-random-fasta + n 5 * IUB-chars IUB-floats "Homo sapiens frequency" "THREE" write-random-fasta + drop ] with-file-writer ] ; diff --git a/extra/cuda/cuda.factor b/extra/cuda/cuda.factor index 6b343fb1cc..94e10a96dd 100644 --- a/extra/cuda/cuda.factor +++ b/extra/cuda/cuda.factor @@ -1,308 +1,83 @@ ! Copyright (C) 2010 Doug Coleman. ! See http://factorcode.org/license.txt for BSD license. -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 ; +USING: accessors alien alien.data alien.parser alien.strings +alien.syntax arrays assocs byte-arrays classes.struct +combinators continuations cuda.ffi cuda.memory cuda.utils +destructors fry io io.backend io.encodings.string +io.encodings.utf8 kernel lexer locals macros math math.parser +namespaces nested-comments opengl.gl.extensions parser +prettyprint quotations sequences words ; +QUALIFIED-WITH: alien.c-types a 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 -- ) - dup CUDA_SUCCESS = [ drop ] [ throw-cuda-error ] if ; - -: cuda-version ( -- n ) - int [ cuDriverGetVersion cuda-error ] keep *int ; - -: init-cuda ( -- ) - 0 cuInit cuda-error ; - TUPLE: launcher { device integer initial: 0 } -{ device-flags initial: 0 } -path block-shape shared-size grid ; +{ device-flags initial: 0 } ; + +TUPLE: function-launcher +dim-block dim-grid shared-size stream ; : with-cuda-context ( flags device quot -- ) - [ - [ CUcontext ] 2dip - [ cuCtxCreate cuda-error ] 3keep 2drop *void* - ] dip + H{ } clone cuda-modules set-global + H{ } clone cuda-functions set + [ create-context ] dip [ '[ _ @ ] ] - [ drop '[ _ cuCtxDestroy cuda-error ] ] 2bi + [ drop '[ _ destroy-context ] ] 2bi [ ] cleanup ; inline -: with-cuda-module ( path quot -- ) - [ - normalize-path - [ CUmodule ] dip - [ cuModuleLoad cuda-error ] 2keep drop *void* - ] dip - [ '[ _ @ ] ] - [ drop '[ _ cuModuleUnload cuda-error ] ] 2bi - [ ] cleanup ; inline - -: with-cuda-program ( flags device path quot -- ) +: with-cuda-program ( flags device quot -- ) [ dup cuda-device set ] 2dip - '[ - cuda-context set - _ [ - cuda-module set - _ call - ] with-cuda-module - ] with-cuda-context ; inline + '[ cuda-context set _ call ] with-cuda-context ; inline : with-cuda ( launcher quot -- ) - [ - init-cuda - H{ } clone cuda-memory-hashtable - ] 2dip '[ + init-cuda + [ H{ } clone cuda-memory-hashtable ] 2dip '[ _ [ cuda-launcher set ] - [ [ device>> ] [ device-flags>> ] [ path>> ] tri ] bi + [ [ device>> ] [ device-flags>> ] bi ] bi _ with-cuda-program ] with-variable ; inline -cuda-setter ( c-type -- n cuda-type ) + { + { [ dup a:int = ] [ drop 4 [ cuda-int* ] ] } + { [ dup a:uint = ] [ drop 4 [ cuda-int* ] ] } + { [ dup a:float = ] [ drop 4 [ cuda-float* ] ] } + { [ dup a:pointer? ] [ drop 4 [ ptr>> cuda-int* ] ] } + { [ dup a:void* = ] [ drop 4 [ ptr>> cuda-int* ] ] } + } cond ; -: #cuda-devices ( -- n ) - int [ cuDeviceGetCount cuda-error ] keep *int ; - -: n>cuda-device ( n -- device ) - [ CUdevice ] dip [ cuDeviceGet cuda-error ] 2keep drop *int ; - -: enumerate-cuda-devices ( -- devices ) - #cuda-devices iota [ n>cuda-device ] map ; - -: cuda-device-properties ( device -- properties ) - [ CUdevprop ] 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 [ ] keep ] dip - [ cuDeviceGetName cuda-error ] - [ 2drop utf8 alien>string ] 3bi ; - -: cuda-device-capability ( n -- pair ) - [ int int ] dip - [ cuDeviceComputeCapability cuda-error ] - [ drop [ *int ] bi@ ] 3bi 2array ; - -: cuda-device-memory ( n -- bytes ) - [ uint ] dip - [ cuDeviceTotalMem cuda-error ] - [ drop *uint ] 2bi ; - -: get-cuda-function* ( module string -- function ) - [ CUfunction ] 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 ; - -: ( 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 ] dip - [ cuMemAlloc cuda-error ] 2keep - [ *int ] dip add-cuda-memory ; - -: 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 - [ 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 ] 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 ] +: run-function-launcher ( function-launcher function -- ) + swap + { + [ dim-block>> first3 function-block-shape* ] + [ shared-size>> function-shared-size* ] [ - grid>> [ - launch-function + dim-grid>> [ + launch-function* ] [ - first2 launch-function-grid + first2 launch-function-grid* ] if-empty ] - } cleave ; + } 2cleave ; -: 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 +: cuda-argument-setter ( offset c-type -- offset' quot ) + c-type>cuda-setter + [ over [ + ] dip ] dip + '[ swap _ swap _ call ] ; + +MACRO: cuda-arguments ( c-types -- quot: ( args... function -- ) ) + [ 0 ] dip [ cuda-argument-setter ] map reverse + swap '[ _ param-size* ] suffix + '[ _ cleave ] ; + +: define-cuda-word ( word module-name function-name arguments -- ) + [ + '[ + _ _ cached-function + [ nip _ cuda-arguments ] + [ run-function-launcher ] 2bi ] - [ "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 ; + ] + [ 2nip \ function-launcher suffix a:void function-effect ] + 3bi define-declared ; diff --git a/extra/cuda/demos/hello-world/authors.txt b/extra/cuda/demos/hello-world/authors.txt new file mode 100644 index 0000000000..7c1b2f2279 --- /dev/null +++ b/extra/cuda/demos/hello-world/authors.txt @@ -0,0 +1 @@ +Doug Coleman diff --git a/extra/cuda/demos/hello-world/hello-world.factor b/extra/cuda/demos/hello-world/hello-world.factor new file mode 100644 index 0000000000..8855ce6fea --- /dev/null +++ b/extra/cuda/demos/hello-world/hello-world.factor @@ -0,0 +1,21 @@ +! Copyright (C) 2010 Doug Coleman. +! See http://factorcode.org/license.txt for BSD license. +USING: alien.c-types alien.strings cuda cuda.memory cuda.syntax +destructors io io.encodings.utf8 kernel locals math sequences ; +IN: cuda.demos.hello-world + +CUDA-LIBRARY: hello vocab:cuda/demos/hello-world/hello.ptx + +CUDA-FUNCTION: helloWorld ( char* string-ptr ) ; + +:: cuda-hello-world ( -- ) + T{ launcher { device 0 } } [ + "Hello World!" [ - ] map-index malloc-device-string + &dispose dup :> str + + { 6 1 1 } { 2 1 } 1 3<<< helloWorld + + str device>host utf8 alien>string print + ] with-cuda ; + +MAIN: cuda-hello-world diff --git a/extra/cuda/hello.cu b/extra/cuda/demos/hello-world/hello.cu similarity index 100% rename from extra/cuda/hello.cu rename to extra/cuda/demos/hello-world/hello.cu diff --git a/extra/cuda/hello.ptx b/extra/cuda/demos/hello-world/hello.ptx similarity index 100% rename from extra/cuda/hello.ptx rename to extra/cuda/demos/hello-world/hello.ptx diff --git a/extra/cuda/demos/prefix-sum/authors.txt b/extra/cuda/demos/prefix-sum/authors.txt new file mode 100644 index 0000000000..2d6d4567d3 --- /dev/null +++ b/extra/cuda/demos/prefix-sum/authors.txt @@ -0,0 +1,2 @@ +Doug Coleman +Joe Groff diff --git a/extra/cuda/prefix-sum.cu b/extra/cuda/demos/prefix-sum/prefix-sum.cu similarity index 100% rename from extra/cuda/prefix-sum.cu rename to extra/cuda/demos/prefix-sum/prefix-sum.cu diff --git a/extra/cuda/demos/prefix-sum/prefix-sum.factor b/extra/cuda/demos/prefix-sum/prefix-sum.factor new file mode 100644 index 0000000000..c7e59b515a --- /dev/null +++ b/extra/cuda/demos/prefix-sum/prefix-sum.factor @@ -0,0 +1,16 @@ +! Copyright (C) 2010 Doug Coleman. +! See http://factorcode.org/license.txt for BSD license. +USING: alien.c-types cuda cuda.syntax locals ; +IN: cuda.demos.prefix-sum + +CUDA-LIBRARY: prefix-sum vocab:cuda/demos/prefix-sum/prefix-sum.ptx + +CUDA-FUNCTION: prefix_sum_block ( uint* in, uint* out, uint n ) ; + +:: cuda-prefix-sum ( -- ) + T{ launcher { device 0 } } + [ + ! { 1 1 1 } { 2 1 } 0 3<<< prefix_sum_block + ] with-cuda ; + +MAIN: cuda-prefix-sum diff --git a/extra/cuda/prefix-sum.ptx b/extra/cuda/demos/prefix-sum/prefix-sum.ptx similarity index 100% rename from extra/cuda/prefix-sum.ptx rename to extra/cuda/demos/prefix-sum/prefix-sum.ptx diff --git a/extra/cuda/devices/authors.txt b/extra/cuda/devices/authors.txt new file mode 100644 index 0000000000..7c1b2f2279 --- /dev/null +++ b/extra/cuda/devices/authors.txt @@ -0,0 +1 @@ +Doug Coleman diff --git a/extra/cuda/devices/devices.factor b/extra/cuda/devices/devices.factor new file mode 100644 index 0000000000..37e199e74e --- /dev/null +++ b/extra/cuda/devices/devices.factor @@ -0,0 +1,65 @@ +! Copyright (C) 2010 Doug Coleman. +! See http://factorcode.org/license.txt for BSD license. +USING: alien.c-types alien.data alien.strings arrays assocs +byte-arrays classes.struct combinators cuda.ffi cuda.utils io +io.encodings.utf8 kernel math.parser prettyprint sequences ; +IN: cuda.devices + +: #cuda-devices ( -- n ) + int [ cuDeviceGetCount cuda-error ] keep *int ; + +: n>cuda-device ( n -- device ) + [ CUdevice ] dip [ cuDeviceGet cuda-error ] 2keep drop *int ; + +: enumerate-cuda-devices ( -- devices ) + #cuda-devices iota [ n>cuda-device ] map ; + +: cuda-device-properties ( device -- properties ) + [ CUdevprop ] dip + [ cuDeviceGetProperties cuda-error ] 2keep drop + CUdevprop memory>struct ; + +: cuda-devices ( -- assoc ) + enumerate-cuda-devices [ dup cuda-device-properties ] { } map>assoc ; + +: cuda-device-name ( n -- string ) + [ 256 [ ] keep ] dip + [ cuDeviceGetName cuda-error ] + [ 2drop utf8 alien>string ] 3bi ; + +: cuda-device-capability ( n -- pair ) + [ int int ] dip + [ cuDeviceComputeCapability cuda-error ] + [ drop [ *int ] bi@ ] 3bi 2array ; + +: cuda-device-memory ( n -- bytes ) + [ uint ] dip + [ cuDeviceTotalMem cuda-error ] + [ drop *uint ] 2bi ; + +: cuda-device-attribute ( attribute dev -- n ) + [ int ] 2dip + [ cuDeviceGetAttribute cuda-error ] + [ 2drop *int ] 3bi ; + +: 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 ; + diff --git a/extra/cuda/memory/authors.txt b/extra/cuda/memory/authors.txt new file mode 100644 index 0000000000..7c1b2f2279 --- /dev/null +++ b/extra/cuda/memory/authors.txt @@ -0,0 +1 @@ +Doug Coleman diff --git a/extra/cuda/memory/memory.factor b/extra/cuda/memory/memory.factor new file mode 100644 index 0000000000..c3dfe56a53 --- /dev/null +++ b/extra/cuda/memory/memory.factor @@ -0,0 +1,74 @@ +! Copyright (C) 2010 Doug Coleman. +! See http://factorcode.org/license.txt for BSD license. +USING: accessors alien alien.data assocs byte-arrays cuda.ffi +cuda.utils destructors io.encodings.string io.encodings.utf8 +kernel locals namespaces sequences ; +QUALIFIED-WITH: alien.c-types a +IN: cuda.memory + +SYMBOL: cuda-memory-hashtable + +TUPLE: cuda-memory < disposable ptr length ; + +: ( 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 ] dip + [ cuMemAlloc cuda-error ] 2keep + [ a:*int ] dip add-cuda-memory ; + +: cuda-free* ( ptr -- ) + cuMemFree cuda-error ; + +M: cuda-memory dispose ( ptr -- ) + ptr>> cuda-free* ; + +: 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 ; + +: host>device ( dest-ptr src-ptr -- ) + [ ptr>> ] dip dup length cuMemcpyHtoD cuda-error ; + +:: device>host ( ptr -- seq ) + ptr byte-length + [ ptr [ ptr>> ] [ byte-length ] bi cuMemcpyDtoH cuda-error ] keep ; + +: malloc-device-string ( string -- n ) + utf8 encode + [ length cuda-malloc ] keep + [ host>device ] [ drop ] 2bi ; diff --git a/extra/cuda/ptx/ptx-tests.factor b/extra/cuda/ptx/ptx-tests.factor new file mode 100644 index 0000000000..28391a5f58 --- /dev/null +++ b/extra/cuda/ptx/ptx-tests.factor @@ -0,0 +1,1091 @@ +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 + +[ """ .version 2.0 + .target sm_20 + abs.s32 a, b; + @p abs.s32 a, b; + @!p abs.s32 a, b; +foo: abs.s32 a, b; + abs.ftz.f32 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ abs { type .s32 } { dest "a" } { a "b" } } + T{ abs + { predicate T{ ptx-predicate { variable "p" } } } + { type .s32 } { dest "a" } { a "b" } + } + T{ abs + { predicate T{ ptx-predicate { negated? t } { variable "p" } } } + { type .s32 } { dest "a" } { a "b" } + } + T{ abs + { label "foo" } + { type .s32 } { dest "a" } { a "b" } + } + T{ abs { type .f32 } { dest "a" } { a "b" } { ftz? t } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + add.s32 a, b, c; + add.cc.s32 a, b, c; + add.sat.s32 a, b, c; + add.ftz.f32 a, b, c; + add.ftz.sat.f32 a, b, c; + add.rz.sat.f32 a, b, c; + add.rz.ftz.sat.f32 a, b, c; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ add { type .s32 } { dest "a" } { a "b" } { b "c" } } + T{ add { cc? t } { type .s32 } { dest "a" } { a "b" } { b "c" } } + T{ add { sat? t } { type .s32 } { dest "a" } { a "b" } { b "c" } } + T{ add { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ add { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ add { round .rz } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ add { round .rz } { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + addc.s32 a, b, c; + addc.cc.s32 a, b, c; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ addc { type .s32 } { dest "a" } { a "b" } { b "c" } } + T{ addc { cc? t } { type .s32 } { dest "a" } { a "b" } { b "c" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + and.b32 a, b, c; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ and { type .b32 } { dest "a" } { a "b" } { b "c" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + atom.and.u32 a, [b], c; + atom.global.or.u32 a, [b], c; + atom.shared.cas.u32 a, [b], c, d; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ atom { op .and } { type .u32 } { dest "a" } { a "[b]" } { b "c" } } + T{ atom { storage-space .global } { op .or } { type .u32 } { dest "a" } { a "[b]" } { b "c" } } + T{ atom { storage-space .shared } { op .cas } { type .u32 } { dest "a" } { a "[b]" } { b "c" } { c "d" } } + + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + bar.arrive a, b; + bar.red.popc.u32 a, b, d; + bar.red.popc.u32 a, b, !d; + bar.red.popc.u32 a, b, c, !d; + bar.sync a; + bar.sync a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ bar.arrive { a "a" } { b "b" } } + T{ bar.red { op .popc } { type .u32 } { dest "a" } { a "b" } { c "d" } } + T{ bar.red { op .popc } { type .u32 } { dest "a" } { a "b" } { c "!d" } } + T{ bar.red { op .popc } { type .u32 } { dest "a" } { a "b" } { b "c" } { c "!d" } } + T{ bar.sync { a "a" } } + T{ bar.sync { a "a" } { b "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + bfe.u32 a, b, c, d; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ bfe { type .u32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + bfi.u32 a, b, c, d, e; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ bfi { type .u32 } { dest "a" } { a "b" } { b "c" } { c "d" } { d "e" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + bfind.u32 a, b; + bfind.shiftamt.u32 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ bfind { type .u32 } { dest "a" } { a "b" } } + T{ bfind { type .u32 } { shiftamt? t } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + bra foo; + bra.uni bar; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ bra { target "foo" } } + T{ bra { uni? t } { target "bar" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + brev.b32 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ brev { type .b32 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + brkpt; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ brkpt } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + call foo; + call.uni foo; + call (a), foo; + call (a), foo, (b); + call (a), foo, (b, c); + call (a), foo, (b, c, d); + call foo, (b, c, d); +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ call { target "foo" } } + T{ call { uni? t } { target "foo" } } + T{ call { return "a" } { target "foo" } } + T{ call { return "a" } { target "foo" } { params { "b" } } } + T{ call { return "a" } { target "foo" } { params { "b" "c" } } } + T{ call { return "a" } { target "foo" } { params { "b" "c" "d" } } } + T{ call { target "foo" } { params { "b" "c" "d" } } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + clz.b32 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ clz { type .b32 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + cnot.b32 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ cnot { type .b32 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + copysign.f64 a, b, c; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ copysign { type .f64 } { dest "a" } { a "b" } { b "c" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + cos.approx.f32 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ cos { round .approx } { type .f32 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + cvt.f32.s32 a, b; + cvt.s32.f32 a, b; + cvt.rp.f32.f64 a, b; + cvt.rpi.s32.f32 a, b; + cvt.ftz.f32.f64 a, b; + cvt.sat.f32.f64 a, b; + cvt.ftz.sat.f32.f64 a, b; + cvt.rp.ftz.sat.f32.f64 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ cvt { dest-type .f32 } { type .s32 } { dest "a" } { a "b" } } + T{ cvt { dest-type .s32 } { type .f32 } { dest "a" } { a "b" } } + T{ cvt { round .rp } { dest-type .f32 } { type .f64 } { dest "a" } { a "b" } } + T{ cvt { round .rpi } { dest-type .s32 } { type .f32 } { dest "a" } { a "b" } } + T{ cvt { ftz? t } { dest-type .f32 } { type .f64 } { dest "a" } { a "b" } } + T{ cvt { sat? t } { dest-type .f32 } { type .f64 } { dest "a" } { a "b" } } + T{ cvt { ftz? t } { sat? t } { dest-type .f32 } { type .f64 } { dest "a" } { a "b" } } + T{ cvt { round .rp } { ftz? t } { sat? t } { dest-type .f32 } { type .f64 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + cvta.global.u64 a, b; + cvta.shared.u64 a, b; + cvta.to.shared.u64 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ cvta { storage-space .global } { type .u64 } { dest "a" } { a "b" } } + T{ cvta { storage-space .shared } { type .u64 } { dest "a" } { a "b" } } + T{ cvta { to? t } { storage-space .shared } { type .u64 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + div.u32 a, b, c; + div.approx.f32 a, b, c; + div.approx.ftz.f32 a, b, c; + div.full.f32 a, b, c; + div.full.ftz.f32 a, b, c; + div.f32 a, b, c; + div.rz.f32 a, b, c; + div.ftz.f32 a, b, c; + div.rz.ftz.f32 a, b, c; + div.f64 a, b, c; + div.rz.f64 a, b, c; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ div { type .u32 } { dest "a" } { a "b" } { b "c" } } + T{ div { round .approx } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ div { round .approx } { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ div { round .full } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ div { round .full } { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ div { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ div { round .rz } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ div { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ div { round .rz } { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ div { type .f64 } { dest "a" } { a "b" } { b "c" } } + T{ div { round .rz } { type .f64 } { dest "a" } { a "b" } { b "c" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + ex2.approx.f32 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ ex2 { round .approx } { type .f32 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + exit; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ exit } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + fma.f32 a, b, c, d; + fma.sat.f32 a, b, c, d; + fma.ftz.f32 a, b, c, d; + fma.ftz.sat.f32 a, b, c, d; + fma.rz.sat.f32 a, b, c, d; + fma.rz.ftz.sat.f32 a, b, c, d; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ fma { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ fma { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ fma { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ fma { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ fma { round .rz } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ fma { round .rz } { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + isspacep.shared a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ isspacep { storage-space .shared } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + ld.u32 a, [b]; + ld.v2.u32 a, [b]; + ld.v4.u32 a, [b]; + ld.v4.u32 {a, b, c, d}, [e]; + ld.lu.u32 a, [b]; + ld.const.lu.u32 a, [b]; + ld.volatile.const[5].u32 a, [b]; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ ld { type .u32 } { dest "a" } { a "[b]" } } + T{ ld { type T{ .v2 { of .u32 } } } { dest "a" } { a "[b]" } } + T{ ld { type T{ .v4 { of .u32 } } } { dest "a" } { a "[b]" } } + T{ ld { type T{ .v4 { of .u32 } } } { dest "{a, b, c, d}" } { a "[e]" } } + T{ ld { cache-op .lu } { type .u32 } { dest "a" } { a "[b]" } } + T{ ld { storage-space T{ .const } } { cache-op .lu } { type .u32 } { dest "a" } { a "[b]" } } + T{ ld { volatile? t } { storage-space T{ .const { bank 5 } } } { type .u32 } { dest "a" } { a "[b]" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + ldu.u32 a, [b]; + ldu.v2.u32 a, [b]; + ldu.v4.u32 a, [b]; + ldu.v4.u32 {a, b, c, d}, [e]; + ldu.lu.u32 a, [b]; + ldu.const.lu.u32 a, [b]; + ldu.volatile.const[5].u32 a, [b]; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ ldu { type .u32 } { dest "a" } { a "[b]" } } + T{ ldu { type T{ .v2 { of .u32 } } } { dest "a" } { a "[b]" } } + T{ ldu { type T{ .v4 { of .u32 } } } { dest "a" } { a "[b]" } } + T{ ldu { type T{ .v4 { of .u32 } } } { dest "{a, b, c, d}" } { a "[e]" } } + T{ ldu { cache-op .lu } { type .u32 } { dest "a" } { a "[b]" } } + T{ ldu { storage-space T{ .const } } { cache-op .lu } { type .u32 } { dest "a" } { a "[b]" } } + T{ ldu { volatile? t } { storage-space T{ .const { bank 5 } } } { type .u32 } { dest "a" } { a "[b]" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + lg2.approx.f32 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ lg2 { round .approx } { type .f32 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + mad.s32 a, b, c, d; + mad.lo.s32 a, b, c, d; + mad.sat.s32 a, b, c, d; + mad.hi.sat.s32 a, b, c, d; + mad.ftz.f32 a, b, c, d; + mad.ftz.sat.f32 a, b, c, d; + mad.rz.sat.f32 a, b, c, d; + mad.rz.ftz.sat.f32 a, b, c, d; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ mad { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ mad { mode .lo } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ mad { sat? t } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ mad { mode .hi } { sat? t } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ mad { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ mad { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ mad { round .rz } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ mad { round .rz } { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + mad24.s32 a, b, c, d; + mad24.lo.s32 a, b, c, d; + mad24.sat.s32 a, b, c, d; + mad24.hi.sat.s32 a, b, c, d; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ mad24 { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ mad24 { mode .lo } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ mad24 { sat? t } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ mad24 { mode .hi } { sat? t } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + neg.s32 a, b; + neg.f32 a, b; + neg.ftz.f32 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ neg { type .s32 } { dest "a" } { a "b" } } + T{ neg { type .f32 } { dest "a" } { a "b" } } + T{ neg { ftz? t } { type .f32 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + not.b32 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ not { type .b32 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + or.b32 a, b, c; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ or { type .b32 } { dest "a" } { a "b" } { b "c" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + pmevent a; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ pmevent { a "a" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + popc.b64 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ popc { type .b64 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + prefetch.L1 [a]; + prefetch.local.L2 [a]; + prefetchu.L1 [a]; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ prefetch { level .L1 } { a "[a]" } } + T{ prefetch { storage-space .local } { level .L2 } { a "[a]" } } + T{ prefetchu { level .L1 } { a "[a]" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + prmt.b32 a, b, c, d; + prmt.b32.f4e a, b, c, d; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ prmt { type .b32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ prmt { type .b32 } { mode .f4e } { dest "a" } { a "b" } { b "c" } { c "d" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + rcp.approx.f32 a, b; + rcp.approx.ftz.f32 a, b; + rcp.f32 a, b; + rcp.rz.f32 a, b; + rcp.ftz.f32 a, b; + rcp.rz.ftz.f32 a, b; + rcp.f64 a, b; + rcp.rz.f64 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ rcp { round .approx } { type .f32 } { dest "a" } { a "b" } } + T{ rcp { round .approx } { ftz? t } { type .f32 } { dest "a" } { a "b" } } + T{ rcp { type .f32 } { dest "a" } { a "b" } } + T{ rcp { round .rz } { type .f32 } { dest "a" } { a "b" } } + T{ rcp { ftz? t } { type .f32 } { dest "a" } { a "b" } } + T{ rcp { round .rz } { ftz? t } { type .f32 } { dest "a" } { a "b" } } + T{ rcp { type .f64 } { dest "a" } { a "b" } } + T{ rcp { round .rz } { type .f64 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + red.and.u32 [a], b; + red.global.and.u32 [a], b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ red { op .and } { type .u32 } { dest "[a]" } { a "b" } } + T{ red { storage-space .global } { op .and } { type .u32 } { dest "[a]" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + rsqrt.approx.f32 a, b; + rsqrt.approx.ftz.f32 a, b; + rsqrt.approx.f64 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ rsqrt { round .approx } { type .f32 } { dest "a" } { a "b" } } + T{ rsqrt { round .approx } { ftz? t } { type .f32 } { dest "a" } { a "b" } } + T{ rsqrt { round .approx } { type .f64 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + rsqrt.approx.f32 a, b; + rsqrt.approx.ftz.f32 a, b; + rsqrt.approx.f64 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ rsqrt { round .approx } { type .f32 } { dest "a" } { a "b" } } + T{ rsqrt { round .approx } { ftz? t } { type .f32 } { dest "a" } { a "b" } } + T{ rsqrt { round .approx } { type .f64 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + sad.u32 a, b, c, d; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ sad { type .u32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + selp.u32 a, b, c, d; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ selp { type .u32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + set.gt.u32.s32 a, b, c; + set.gt.ftz.u32.f32 a, b, c; + set.gt.and.ftz.u32.f32 a, b, c, d; + set.gt.and.ftz.u32.f32 a, b, c, !d; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ set { cmp-op .gt } { dest-type .u32 } { type .s32 } { dest "a" } { a "b" } { b "c" } } + T{ set { cmp-op .gt } { ftz? t } { dest-type .u32 } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ set { cmp-op .gt } { bool-op .and } { ftz? t } { dest-type .u32 } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ set { cmp-op .gt } { bool-op .and } { ftz? t } { dest-type .u32 } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "!d" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + setp.gt.s32 a, b, c; + setp.gt.s32 a|z, b, c; + setp.gt.ftz.f32 a, b, c; + setp.gt.and.ftz.f32 a, b, c, d; + setp.gt.and.ftz.f32 a, b, c, !d; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ setp { cmp-op .gt } { type .s32 } { dest "a" } { a "b" } { b "c" } } + T{ setp { cmp-op .gt } { type .s32 } { dest "a" } { |dest "z" } { a "b" } { b "c" } } + T{ setp { cmp-op .gt } { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ setp { cmp-op .gt } { bool-op .and } { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ setp { cmp-op .gt } { bool-op .and } { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "!d" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + shl.b32 a, b, c; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ shl { type .b32 } { dest "a" } { a "b" } { b "c" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + shr.b32 a, b, c; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ shr { type .b32 } { dest "a" } { a "b" } { b "c" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + sin.approx.f32 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ sin { round .approx } { type .f32 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + slct.f32.s32 a, b, c, d; + slct.ftz.f32.s32 a, b, c, d; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ slct { dest-type .f32 } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + T{ slct { ftz? t } { dest-type .f32 } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + sqrt.approx.f32 a, b; + sqrt.approx.ftz.f32 a, b; + sqrt.f32 a, b; + sqrt.rz.f32 a, b; + sqrt.ftz.f32 a, b; + sqrt.rz.ftz.f32 a, b; + sqrt.f64 a, b; + sqrt.rz.f64 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ sqrt { round .approx } { type .f32 } { dest "a" } { a "b" } } + T{ sqrt { round .approx } { ftz? t } { type .f32 } { dest "a" } { a "b" } } + T{ sqrt { type .f32 } { dest "a" } { a "b" } } + T{ sqrt { round .rz } { type .f32 } { dest "a" } { a "b" } } + T{ sqrt { ftz? t } { type .f32 } { dest "a" } { a "b" } } + T{ sqrt { round .rz } { ftz? t } { type .f32 } { dest "a" } { a "b" } } + T{ sqrt { type .f64 } { dest "a" } { a "b" } } + T{ sqrt { round .rz } { type .f64 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + st.u32 [a], b; + st.v2.u32 [a], b; + st.v4.u32 [a], b; + st.v4.u32 [a], {b, c, d, e}; + st.lu.u32 [a], b; + st.local.lu.u32 [a], b; + st.volatile.local.u32 [a], b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ st { type .u32 } { dest "[a]" } { a "b" } } + T{ st { type T{ .v2 { of .u32 } } } { dest "[a]" } { a "b" } } + T{ st { type T{ .v4 { of .u32 } } } { dest "[a]" } { a "b" } } + T{ st { type T{ .v4 { of .u32 } } } { dest "[a]" } { a "{b, c, d, e}" } } + T{ st { cache-op .lu } { type .u32 } { dest "[a]" } { a "b" } } + T{ st { storage-space .local } { cache-op .lu } { type .u32 } { dest "[a]" } { a "b" } } + T{ st { volatile? t } { storage-space .local } { type .u32 } { dest "[a]" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + sub.s32 a, b, c; + sub.cc.s32 a, b, c; + sub.sat.s32 a, b, c; + sub.ftz.f32 a, b, c; + sub.ftz.sat.f32 a, b, c; + sub.rz.sat.f32 a, b, c; + sub.rz.ftz.sat.f32 a, b, c; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ sub { type .s32 } { dest "a" } { a "b" } { b "c" } } + T{ sub { cc? t } { type .s32 } { dest "a" } { a "b" } { b "c" } } + T{ sub { sat? t } { type .s32 } { dest "a" } { a "b" } { b "c" } } + T{ sub { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ sub { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ sub { round .rz } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } } + T{ sub { round .rz } { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + subc.s32 a, b, c; + subc.cc.s32 a, b, c; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ subc { type .s32 } { dest "a" } { a "b" } { b "c" } } + T{ subc { cc? t } { type .s32 } { dest "a" } { a "b" } { b "c" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + testp.finite.f32 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ testp { op .finite } { type .f32 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + trap; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ trap } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + vote.all.pred a, b; + vote.all.pred a, !b; + vote.ballot.b32 a, b; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ vote { mode .all } { type .pred } { dest "a" } { a "b" } } + T{ vote { mode .all } { type .pred } { dest "a" } { a "!b" } } + T{ vote { mode .ballot } { type .b32 } { dest "a" } { a "b" } } + } } + } ptx>string +] unit-test + +[ """ .version 2.0 + .target sm_20 + xor.b32 a, b, c; +""" ] [ + T{ ptx + { version "2.0" } + { target T{ ptx-target { arch sm_20 } } } + { body { + T{ xor { type .b32 } { dest "a" } { a "b" } { b "c" } } + } } + } ptx>string +] unit-test + diff --git a/extra/cuda/ptx/ptx.factor b/extra/cuda/ptx/ptx.factor index 8d4925d55f..4618f8b5b6 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 ; @@ -241,7 +242,7 @@ TUPLE: cnot < ptx-2op-instruction ; TUPLE: copysign < ptx-3op-instruction ; TUPLE: cos <{ ptx-2op-instruction ptx-float-env } ; TUPLE: cvt < ptx-2op-instruction - { rounding-mode ?ptx-rounding-mode } + { round ?ptx-rounding-mode } { ftz? boolean } { sat? boolean } { dest-type ptx-type } ; @@ -253,7 +254,7 @@ TUPLE: ex2 <{ ptx-2op-instruction ptx-float-env } ; TUPLE: exit < ptx-instruction ; TUPLE: fma <{ ptx-mad-instruction ptx-float-env } ; TUPLE: isspacep < ptx-instruction - { storage-space ?ptx-storage-space } + { storage-space ptx-storage-space } { dest string } { a string } ; TUPLE: ld < ptx-ldst-instruction ; @@ -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 ; @@ -538,7 +547,7 @@ M: bar.red (write-ptx-element) dup b>> [ ", " write write ] when* ", " write c>> write ; M: bar.sync (write-ptx-element) - "bar.arrive " write-insn + "bar.sync " write-insn dup a>> write dup b>> [ ", " write write ] when* drop ; @@ -554,15 +563,16 @@ M: bfind (write-ptx-element) write-2op ; M: bra (write-ptx-element) "bra" write-insn - dup write-uni - " " write target>> write ; + dup write-uni " " write + target>> write ; M: brev (write-ptx-element) "brev" write-insn write-2op ; M: brkpt (write-ptx-element) "brkpt" write-insn drop ; M: call (write-ptx-element) - "call" write-insn " " write + "call" write-insn + dup write-uni " " write dup return>> [ "(" write write "), " write ] when* dup target>> write dup params>> [ ", (" write ", " join write ")" write ] unless-empty @@ -582,7 +592,7 @@ M: cos (write-ptx-element) write-2op ; M: cvt (write-ptx-element) "cvt" write-insn - dup rounding-mode>> (write-ptx-element) + dup round>> (write-ptx-element) dup write-ftz dup write-sat dup dest-type>> (write-ptx-element) @@ -676,12 +686,17 @@ M: prefetchu (write-ptx-element) " " write a>> write ; M: prmt (write-ptx-element) "prmt" write-insn - dup mode>> (write-ptx-element) - write-4op ; + dup type>> (write-ptx-element) + dup mode>> (write-ptx-element) " " write + dup dest>> write ", " write + dup a>> write ", " write + dup b>> write ", " write + dup c>> write + drop ; M: rcp (write-ptx-element) "rcp" write-insn dup write-float-env - write-3op ; + write-2op ; M: red (write-ptx-element) "red" write-insn dup storage-space>> (write-ptx-element) @@ -749,10 +764,15 @@ M: testp (write-ptx-element) "testp" write-insn dup op>> (write-ptx-element) write-2op ; +M: trap (write-ptx-element) + "trap" write-insn drop ; M: vote (write-ptx-element) "vote" write-insn 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 ; diff --git a/extra/cuda/syntax/authors.txt b/extra/cuda/syntax/authors.txt new file mode 100644 index 0000000000..7c1b2f2279 --- /dev/null +++ b/extra/cuda/syntax/authors.txt @@ -0,0 +1 @@ +Doug Coleman diff --git a/extra/cuda/syntax/syntax.factor b/extra/cuda/syntax/syntax.factor new file mode 100644 index 0000000000..1cd5edb9d4 --- /dev/null +++ b/extra/cuda/syntax/syntax.factor @@ -0,0 +1,20 @@ +! Copyright (C) 2010 Doug Coleman. +! See http://factorcode.org/license.txt for BSD license. +USING: alien.parser cuda cuda.utils io.backend kernel lexer +namespaces parser ; +IN: cuda.syntax + +SYNTAX: CUDA-LIBRARY: + scan scan normalize-path + [ add-cuda-library ] + [ drop current-cuda-library set-global ] 2bi ; + +SYNTAX: CUDA-FUNCTION: + scan [ create-in current-cuda-library get ] [ ] bi + ";" scan-c-args drop define-cuda-word ; + +: 3<<< ( dim-block dim-grid shared-size -- function-launcher ) + f function-launcher boa ; + +: 4<<< ( dim-block dim-grid shared-size stream -- function-launcher ) + function-launcher boa ; diff --git a/extra/cuda/utils/authors.txt b/extra/cuda/utils/authors.txt new file mode 100644 index 0000000000..7c1b2f2279 --- /dev/null +++ b/extra/cuda/utils/authors.txt @@ -0,0 +1 @@ +Doug Coleman diff --git a/extra/cuda/utils/utils.factor b/extra/cuda/utils/utils.factor new file mode 100644 index 0000000000..912b9e2e92 --- /dev/null +++ b/extra/cuda/utils/utils.factor @@ -0,0 +1,143 @@ +! Copyright (C) 2010 Doug Coleman. +! See http://factorcode.org/license.txt for BSD license. +USING: accessors alien.c-types alien.data alien.strings arrays +assocs byte-arrays classes.struct combinators cuda.ffi io +io.backend io.encodings.utf8 kernel math.parser namespaces +prettyprint sequences ; +IN: cuda.utils + +SYMBOL: cuda-device +SYMBOL: cuda-context +SYMBOL: cuda-module +SYMBOL: cuda-function +SYMBOL: cuda-launcher + +SYMBOL: cuda-modules +SYMBOL: cuda-functions + +ERROR: throw-cuda-error n ; + +: cuda-error ( n -- ) + dup CUDA_SUCCESS = [ drop ] [ throw-cuda-error ] if ; + +: init-cuda ( -- ) + 0 cuInit cuda-error ; + +: cuda-version ( -- n ) + int [ cuDriverGetVersion cuda-error ] keep *int ; + +: get-function-ptr* ( module string -- function ) + [ CUfunction ] 2dip + [ cuModuleGetFunction cuda-error ] 3keep 2drop *void* ; + +: get-function-ptr ( string -- function ) + [ cuda-module get ] dip get-function-ptr* ; + +: with-cuda-function ( string quot -- ) + [ + get-function-ptr* cuda-function set + ] dip call ; inline + +: create-context ( flags device -- context ) + [ CUcontext ] 2dip + [ cuCtxCreate cuda-error ] 3keep 2drop *void* ; + +: destroy-context ( context -- ) cuCtxDestroy cuda-error ; + +SYMBOL: cuda-libraries +cuda-libraries [ H{ } clone ] initialize + +SYMBOL: current-cuda-library + +TUPLE: cuda-library name path handle ; + +: ( name path -- obj ) + \ cuda-library new + swap >>path + swap >>name ; + +: add-cuda-library ( name path -- ) + normalize-path + dup name>> cuda-libraries get-global set-at ; + +: ?delete-at ( key assoc -- old/key ? ) + 2dup delete-at* [ 2nip t ] [ 2drop f ] if ; inline + +ERROR: no-cuda-library name ; + +: load-module ( path -- module ) + [ CUmodule ] dip + [ cuModuleLoad cuda-error ] 2keep drop *void* ; + +: unload-module ( module -- ) + cuModuleUnload cuda-error ; + +: load-cuda-library ( library -- handle ) + path>> load-module ; + +: lookup-cuda-library ( name -- cuda-library ) + cuda-libraries get ?at [ no-cuda-library ] unless ; + +: remove-cuda-library ( name -- library ) + cuda-libraries get ?delete-at [ no-cuda-library ] unless ; + +: unload-cuda-library ( name -- ) + remove-cuda-library handle>> unload-module ; + + +: cached-module ( module-name -- alien ) + lookup-cuda-library + cuda-modules get-global [ load-cuda-library ] cache ; + +: cached-function ( module-name function-name -- alien ) + [ cached-module ] dip + 2array cuda-functions get [ first2 get-function-ptr* ] cache ; + +: launch-function* ( function -- ) cuLaunch cuda-error ; + +: launch-function ( -- ) cuda-function get cuLaunch 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* ; + +: launch-function-grid* ( function width height -- ) + cuLaunchGrid cuda-error ; + +: launch-function-grid ( width height -- ) + [ cuda-function get ] 2dip + cuLaunchGrid cuda-error ; + +: 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 ; diff --git a/extra/game/loop/loop.factor b/extra/game/loop/loop.factor index ffe5acd879..312d7dbd1c 100644 --- a/extra/game/loop/loop.factor +++ b/extra/game/loop/loop.factor @@ -112,6 +112,6 @@ PRIVATE> M: game-loop dispose stop-loop ; -USING: vocabs vocabs.loader ; +USE: vocabs.loader -"prettyprint" "game.loop.prettyprint" require-when +{ "game.loop" "prettyprint" } "game.loop.prettyprint" require-when diff --git a/extra/gpu/shaders/shaders.factor b/extra/gpu/shaders/shaders.factor index 974f2f8070..8a2931e431 100755 --- a/extra/gpu/shaders/shaders.factor +++ b/extra/gpu/shaders/shaders.factor @@ -632,4 +632,4 @@ M: program-instance dispose [ world>> ] [ program>> instances>> ] [ ] tri ?delete-at reset-memos ; -"prettyprint" "gpu.shaders.prettyprint" require-when +{ "gpu.shaders" "prettyprint" } "gpu.shaders.prettyprint" require-when diff --git a/extra/javascriptcore/authors.txt b/extra/javascriptcore/authors.txt new file mode 100644 index 0000000000..7c1b2f2279 --- /dev/null +++ b/extra/javascriptcore/authors.txt @@ -0,0 +1 @@ +Doug Coleman diff --git a/extra/javascriptcore/core-foundation/authors.txt b/extra/javascriptcore/core-foundation/authors.txt new file mode 100644 index 0000000000..7c1b2f2279 --- /dev/null +++ b/extra/javascriptcore/core-foundation/authors.txt @@ -0,0 +1 @@ +Doug Coleman diff --git a/extra/javascriptcore/core-foundation/core-foundation.factor b/extra/javascriptcore/core-foundation/core-foundation.factor new file mode 100644 index 0000000000..9dfc93b101 --- /dev/null +++ b/extra/javascriptcore/core-foundation/core-foundation.factor @@ -0,0 +1,11 @@ +! Copyright (C) 2010 Doug Coleman. +! See http://factorcode.org/license.txt for BSD license. +USING: alien.syntax core-foundation core-foundation.strings +javascriptcore.ffi ; +IN: javascriptcore.core-foundation + +FUNCTION: JSStringRef JSStringCreateWithCFString ( CFStringRef string ) ; + +FUNCTION: CFStringRef JSStringCopyCFString ( CFAllocatorRef alloc, JSStringRef string ) ; + + diff --git a/extra/javascriptcore/core-foundation/platforms.txt b/extra/javascriptcore/core-foundation/platforms.txt new file mode 100644 index 0000000000..6e806f449e --- /dev/null +++ b/extra/javascriptcore/core-foundation/platforms.txt @@ -0,0 +1 @@ +macosx diff --git a/extra/javascriptcore/ffi/authors.txt b/extra/javascriptcore/ffi/authors.txt new file mode 100644 index 0000000000..7c1b2f2279 --- /dev/null +++ b/extra/javascriptcore/ffi/authors.txt @@ -0,0 +1 @@ +Doug Coleman diff --git a/extra/javascriptcore/ffi/ffi.factor b/extra/javascriptcore/ffi/ffi.factor new file mode 100644 index 0000000000..844e169eed --- /dev/null +++ b/extra/javascriptcore/ffi/ffi.factor @@ -0,0 +1,266 @@ +! Copyright (C) 2010 Doug Coleman. +! See http://factorcode.org/license.txt for BSD license. +USING: alien alien.c-types alien.libraries alien.syntax +classes.struct combinators io.encodings.utf8 system ; +IN: javascriptcore.ffi + +<< +"javascriptcore" { + { [ os macosx? ] [ "/System/Library/Frameworks/JavaScriptCore.framework/Versions/Current/JavaScriptCore" ] } + ! { [ os winnt? ] [ "javascriptcore.dll" ] } + ! { [ os unix? ] [ "libsqlite3.so" ] } + } cond cdecl add-library +>> + +LIBRARY: javascriptcore + +TYPEDEF: void* JSContextGroupRef +TYPEDEF: void* JSContextRef +TYPEDEF: void* JSGlobalContextRef +TYPEDEF: void* JSStringRef +TYPEDEF: void* JSClassRef +TYPEDEF: void* JSPropertyNameArrayRef +TYPEDEF: void* JSPropertyNameAccumulatorRef +TYPEDEF: void* JSValueRef +TYPEDEF: void* JSObjectRef +TYPEDEF: void* JSObjectInitializeCallback +TYPEDEF: void* JSObjectFinalizeCallback +TYPEDEF: void* JSObjectHasPropertyCallback +TYPEDEF: void* JSObjectGetPropertyCallback +TYPEDEF: void* JSObjectSetPropertyCallback +TYPEDEF: void* JSObjectDeletePropertyCallback +TYPEDEF: void* JSObjectGetPropertyNamesCallback +TYPEDEF: void* JSObjectCallAsFunctionCallback +TYPEDEF: void* JSObjectCallAsConstructorCallback +TYPEDEF: void* JSObjectHasInstanceCallback +TYPEDEF: void* JSObjectConvertToTypeCallback +TYPEDEF: uint unsigned +TYPEDEF: ushort JSChar +! char[utf16n] for strings + +C-ENUM: JSPropertyAttributes + { kJSPropertyAttributeNone 0 } + { kJSPropertyAttributeReadOnly 2 } + { kJSPropertyAttributeDontEnum 4 } + { kJSPropertyAttributeDontDelete 8 } ; + +C-ENUM: JSClassAttributes + { kJSClassAttributeNone 0 } + { kJSClassAttributeNoAutomaticPrototype 2 } ; + +C-ENUM: JSType + kJSTypeUndefined, + kJSTypeNull, + kJSTypeBoolean, + kJSTypeNumber, + kJSTypeString, + kJSTypeObject ; + +STRUCT: JSStaticValue + { name c-string } + { getProperty JSObjectGetPropertyCallback } + { setProperty JSObjectSetPropertyCallback } + { attributes JSPropertyAttributes } ; + +STRUCT: JSStaticFunction + { name c-string } + { callAsFunction JSObjectCallAsFunctionCallback } ; + +STRUCT: JSClassDefinition + { version int } + { attributes JSClassAttributes } + { className c-string } + { parentClass JSClassRef } + { staticValues JSStaticValue* } + { staticFunctions JSStaticFunction* } + { initialize JSObjectInitializeCallback } + { finalize JSObjectFinalizeCallback } + { hasProperty JSObjectHasPropertyCallback } + { getProperty JSObjectGetPropertyCallback } + { setProperty JSObjectSetPropertyCallback } + { deleteProperty JSObjectDeletePropertyCallback } + { getPropertyNames JSObjectGetPropertyNamesCallback } + { callAsFunction JSObjectCallAsFunctionCallback } + { callAsConstructor JSObjectCallAsConstructorCallback } + { hasInstance JSObjectHasInstanceCallback } + { convertToType JSObjectConvertToTypeCallback } ; + +ALIAS: kJSClassDefinitionEmpty JSClassDefinition + +FUNCTION: JSValueRef JSEvaluateScript ( + JSContextRef ctx, + JSStringRef script, + JSObjectRef thisObject, + JSStringRef sourceURL, + int startingLineNumber, + JSValueRef* exception ) ; + +FUNCTION: bool JSCheckScriptSyntax ( + JSContextRef ctx, + JSStringRef script, + JSStringRef sourceURL, + int startingLineNumber, + JSValueRef* exception ) ; + +FUNCTION: void JSGarbageCollect + ( JSContextRef ctx ) ; + +FUNCTION: JSContextGroupRef JSContextGroupCreate + ( ) ; + +FUNCTION: JSContextGroupRef JSContextGroupRetain + ( JSContextGroupRef group ) ; + +FUNCTION: void JSContextGroupRelease + ( JSContextGroupRef group ) ; + +FUNCTION: JSGlobalContextRef JSGlobalContextCreate + ( JSClassRef globalObjectClass ) ; + +FUNCTION: JSGlobalContextRef JSGlobalContextCreateInGroup ( + JSContextGroupRef group, + JSClassRef globalObjectClass ) ; + +FUNCTION: JSGlobalContextRef JSGlobalContextRetain + ( JSGlobalContextRef ctx ) ; + +FUNCTION: void JSGlobalContextRelease + ( JSGlobalContextRef ctx ) ; + +FUNCTION: JSObjectRef JSContextGetGlobalObject + ( JSContextRef ctx ) ; + +FUNCTION: JSContextGroupRef JSContextGetGroup + ( JSContextRef ctx ) ; + +FUNCTION: JSClassRef JSClassCreate + ( JSClassDefinition* definition ) ; + +FUNCTION: JSClassRef JSClassRetain + ( JSClassRef jsClass ) ; + +FUNCTION: void JSClassRelease + ( JSClassRef jsClass ) ; + +FUNCTION: JSObjectRef JSObjectMake + ( JSContextRef ctx, + JSClassRef jsClass, void* data ) ; + +FUNCTION: JSObjectRef JSObjectMakeFunctionWithCallback ( JSContextRef ctx, JSStringRef name, JSObjectCallAsFunctionCallback callAsFunction ) ; + +FUNCTION: JSObjectRef JSObjectMakeConstructor ( JSContextRef ctx, JSClassRef jsClass, JSObjectCallAsConstructorCallback callAsConstructor ) ; + +FUNCTION: JSObjectRef JSObjectMakeArray ( JSContextRef ctx, size_t argumentCount, JSValueRef arguments[], JSValueRef* exception ) ; + +FUNCTION: JSObjectRef JSObjectMakeDate ( JSContextRef ctx, size_t argumentCount, JSValueRef arguments[], JSValueRef* exception ) ; + +FUNCTION: JSObjectRef JSObjectMakeError ( JSContextRef ctx, size_t argumentCount, JSValueRef arguments[], JSValueRef* exception ) ; + +FUNCTION: JSObjectRef JSObjectMakeRegExp ( JSContextRef ctx, size_t argumentCount, JSValueRef arguments[], JSValueRef* exception ) ; + +FUNCTION: JSObjectRef JSObjectMakeFunction ( JSContextRef ctx, JSStringRef name, unsigned parameterCount, JSStringRef parameterNames[], JSStringRef body, JSStringRef sourceURL, int startingLineNumber, JSValueRef* exception ) ; + +FUNCTION: JSValueRef JSObjectGetPrototype ( JSContextRef ctx, JSObjectRef object ) ; + +FUNCTION: void JSObjectSetPrototype ( JSContextRef ctx, JSObjectRef object, JSValueRef value ) ; + +FUNCTION: bool JSObjectHasProperty ( JSContextRef ctx, JSObjectRef object, JSStringRef propertyName ) ; + +FUNCTION: JSValueRef JSObjectGetProperty ( JSContextRef ctx, JSObjectRef object, JSStringRef propertyName, JSValueRef* exception ) ; + +FUNCTION: void JSObjectSetProperty ( JSContextRef ctx, JSObjectRef object, JSStringRef propertyName, JSValueRef value, JSPropertyAttributes attributes, JSValueRef* exception ) ; + +FUNCTION: bool JSObjectDeleteProperty ( JSContextRef ctx, JSObjectRef object, JSStringRef propertyName, JSValueRef* exception ) ; + +FUNCTION: JSValueRef JSObjectGetPropertyAtIndex ( JSContextRef ctx, JSObjectRef object, unsigned propertyIndex, JSValueRef* exception ) ; + +FUNCTION: void JSObjectSetPropertyAtIndex ( JSContextRef ctx, JSObjectRef object, unsigned propertyIndex, JSValueRef value, JSValueRef* exception ) ; + +FUNCTION: void* JSObjectGetPrivate ( JSObjectRef object ) ; + +FUNCTION: bool JSObjectSetPrivate ( JSObjectRef object, void* data ) ; + +FUNCTION: bool JSObjectIsFunction ( JSContextRef ctx, JSObjectRef object ) ; + +FUNCTION: JSValueRef JSObjectCallAsFunction ( JSContextRef ctx, JSObjectRef object, JSObjectRef thisObject, size_t argumentCount, JSValueRef arguments[], JSValueRef* exception ) ; + +FUNCTION: bool JSObjectIsConstructor ( JSContextRef ctx, JSObjectRef object ) ; + +FUNCTION: JSObjectRef JSObjectCallAsConstructor ( JSContextRef ctx, JSObjectRef object, size_t argumentCount, JSValueRef arguments[], JSValueRef* exception ) ; + +FUNCTION: JSPropertyNameArrayRef JSObjectCopyPropertyNames ( JSContextRef ctx, JSObjectRef object ) ; + +FUNCTION: JSPropertyNameArrayRef JSPropertyNameArrayRetain ( JSPropertyNameArrayRef array ) ; + +FUNCTION: void JSPropertyNameArrayRelease ( JSPropertyNameArrayRef array ) ; + +FUNCTION: size_t JSPropertyNameArrayGetCount ( JSPropertyNameArrayRef array ) ; + +FUNCTION: JSStringRef JSPropertyNameArrayGetNameAtIndex ( JSPropertyNameArrayRef array, size_t index ) ; + +FUNCTION: void JSPropertyNameAccumulatorAddName ( JSPropertyNameAccumulatorRef accumulator, JSStringRef propertyName ) ; + +FUNCTION: JSStringRef JSStringCreateWithCharacters ( JSChar* chars, size_t numChars ) ; + +FUNCTION: JSStringRef JSStringCreateWithUTF8CString ( c-string[utf8] string ) ; + +FUNCTION: JSStringRef JSStringRetain ( JSStringRef string ) ; + +FUNCTION: void JSStringRelease ( JSStringRef string ) ; + +FUNCTION: size_t JSStringGetLength ( JSStringRef string ) ; + +FUNCTION: JSChar* JSStringGetCharactersPtr ( JSStringRef string ) ; + +FUNCTION: size_t JSStringGetMaximumUTF8CStringSize ( JSStringRef string ) ; + +FUNCTION: size_t JSStringGetUTF8CString ( JSStringRef string, char* buffer, size_t bufferSize ) ; + +FUNCTION: bool JSStringIsEqual ( JSStringRef a, JSStringRef b ) ; + +FUNCTION: bool JSStringIsEqualToUTF8CString ( JSStringRef a, char* b ) ; + +FUNCTION: JSType JSValueGetType ( JSContextRef ctx, JSValueRef value ) ; + +FUNCTION: bool JSValueIsUndefined ( JSContextRef ctx, JSValueRef value ) ; + +FUNCTION: bool JSValueIsNull ( JSContextRef ctx, JSValueRef value ) ; + +FUNCTION: bool JSValueIsBoolean ( JSContextRef ctx, JSValueRef value ) ; + +FUNCTION: bool JSValueIsNumber ( JSContextRef ctx, JSValueRef value ) ; + +FUNCTION: bool JSValueIsString ( JSContextRef ctx, JSValueRef value ) ; + +FUNCTION: bool JSValueIsObject ( JSContextRef ctx, JSValueRef value ) ; + +FUNCTION: bool JSValueIsObjectOfClass ( JSContextRef ctx, JSValueRef value, JSClassRef jsClass ) ; + +FUNCTION: bool JSValueIsEqual ( JSContextRef ctx, JSValueRef a, JSValueRef b, JSValueRef* exception ) ; + +FUNCTION: bool JSValueIsStrictEqual ( JSContextRef ctx, JSValueRef a, JSValueRef b ) ; + +FUNCTION: bool JSValueIsInstanceOfConstructor ( JSContextRef ctx, JSValueRef value, JSObjectRef constructor, JSValueRef* exception ) ; + +FUNCTION: JSValueRef JSValueMakeUndefined ( JSContextRef ctx ) ; + +FUNCTION: JSValueRef JSValueMakeNull ( JSContextRef ctx ) ; + +FUNCTION: JSValueRef JSValueMakeBoolean ( JSContextRef ctx, bool boolean ) ; + +FUNCTION: JSValueRef JSValueMakeNumber ( JSContextRef ctx, double number ) ; + +FUNCTION: JSValueRef JSValueMakeString ( JSContextRef ctx, JSStringRef string ) ; + +FUNCTION: bool JSValueToBoolean ( JSContextRef ctx, JSValueRef value ) ; + +FUNCTION: double JSValueToNumber ( JSContextRef ctx, JSValueRef value, JSValueRef* exception ) ; + +FUNCTION: JSStringRef JSValueToStringCopy ( JSContextRef ctx, JSValueRef value, JSValueRef* exception ) ; + +FUNCTION: JSObjectRef JSValueToObject ( JSContextRef ctx, JSValueRef value, JSValueRef* exception ) ; + +FUNCTION: void JSValueProtect ( JSContextRef ctx, JSValueRef value ) ; + +FUNCTION: void JSValueUnprotect ( JSContextRef ctx, JSValueRef value ) ; + diff --git a/extra/javascriptcore/ffi/hack/authors.txt b/extra/javascriptcore/ffi/hack/authors.txt new file mode 100644 index 0000000000..7c1b2f2279 --- /dev/null +++ b/extra/javascriptcore/ffi/hack/authors.txt @@ -0,0 +1 @@ +Doug Coleman diff --git a/extra/javascriptcore/ffi/hack/hack.factor b/extra/javascriptcore/ffi/hack/hack.factor new file mode 100644 index 0000000000..1866a24e22 --- /dev/null +++ b/extra/javascriptcore/ffi/hack/hack.factor @@ -0,0 +1,29 @@ +! Copyright (C) 2010 Doug Coleman. +! See http://factorcode.org/license.txt for BSD license. +USING: alien alien.accessors alien.syntax kernel kernel.private +math system ; +IN: javascriptcore.ffi.hack + +HOOK: set-callstack-bounds os ( -- ) + +HOOK: macosx-callstack-start-offset cpu ( -- address ) +HOOK: macosx-callstack-size-offset cpu ( -- address ) + +M: ppc macosx-callstack-start-offset HEX: 188 ; +M: ppc macosx-callstack-size-offset HEX: 18c ; + +M: x86.32 macosx-callstack-start-offset HEX: c48 ; +M: x86.32 macosx-callstack-size-offset HEX: c4c ; + +M: x86.64 macosx-callstack-start-offset HEX: 1860 ; +M: x86.64 macosx-callstack-size-offset HEX: 1868 ; + +M: object set-callstack-bounds ; + +FUNCTION: void* pthread_self ( ) ; + +M: macosx set-callstack-bounds + callstack-bounds over [ alien-address ] bi@ - + pthread_self + [ macosx-callstack-size-offset set-alien-unsigned-cell ] + [ macosx-callstack-start-offset set-alien-cell ] bi ; diff --git a/extra/javascriptcore/javascriptcore.factor b/extra/javascriptcore/javascriptcore.factor new file mode 100644 index 0000000000..773a559d2d --- /dev/null +++ b/extra/javascriptcore/javascriptcore.factor @@ -0,0 +1,8 @@ +! Copyright (C) 2010 Doug Coleman. +! See http://factorcode.org/license.txt for BSD license. +USING: javascriptcore.ffi.hack kernel ; +IN: javascriptcore + +: with-javascriptcore ( quot -- ) + set-callstack-bounds + call ; inline diff --git a/vm/callstack.cpp b/vm/callstack.cpp index eae976219f..bb716cbc6d 100755 --- a/vm/callstack.cpp +++ b/vm/callstack.cpp @@ -214,4 +214,10 @@ void factor_vm::primitive_set_innermost_stack_frame_quot() FRAME_RETURN_ADDRESS(inner,this) = (char *)quot->entry_point + offset; } +void factor_vm::primitive_callstack_bounds() +{ + ctx->push(allot_alien((void*)ctx->callstack_seg->start)); + ctx->push(allot_alien((void*)ctx->callstack_seg->end)); +} + } diff --git a/vm/code_blocks.cpp b/vm/code_blocks.cpp index de103cda12..2e7b8d4f09 100755 --- a/vm/code_blocks.cpp +++ b/vm/code_blocks.cpp @@ -265,6 +265,9 @@ struct initial_code_block_visitor { case RT_LITERAL: op.store_value(next_literal()); break; + case RT_FLOAT: + op.store_float(next_literal()); + break; case RT_ENTRY_POINT: op.store_value(parent->compute_entry_point_address(next_literal())); break; diff --git a/vm/compaction.cpp b/vm/compaction.cpp index 5e52c70b0c..34398e3d88 100644 --- a/vm/compaction.cpp +++ b/vm/compaction.cpp @@ -111,6 +111,9 @@ struct code_block_compaction_relocation_visitor { case RT_LITERAL: op.store_value(slot_forwarder.visit_pointer(op.load_value(old_offset))); break; + case RT_FLOAT: + op.store_float(slot_forwarder.visit_pointer(op.load_float(old_offset))); + break; case RT_ENTRY_POINT: case RT_ENTRY_POINT_PIC: case RT_ENTRY_POINT_PIC_TAIL: diff --git a/vm/image.cpp b/vm/image.cpp index ccce96a952..4dfdc4242e 100755 --- a/vm/image.cpp +++ b/vm/image.cpp @@ -185,6 +185,9 @@ struct code_block_fixup_relocation_visitor { case RT_LITERAL: op.store_value(data_visitor.visit_pointer(op.load_value(old_offset))); break; + case RT_FLOAT: + op.store_float(data_visitor.visit_pointer(op.load_float(old_offset))); + break; case RT_ENTRY_POINT: case RT_ENTRY_POINT_PIC: case RT_ENTRY_POINT_PIC_TAIL: diff --git a/vm/instruction_operands.cpp b/vm/instruction_operands.cpp index b11db279a5..af7d363aef 100644 --- a/vm/instruction_operands.cpp +++ b/vm/instruction_operands.cpp @@ -62,6 +62,16 @@ fixnum instruction_operand::load_value() return load_value(pointer); } +cell instruction_operand::load_float() +{ + return (cell)load_value() - boxed_float_offset; +} + +cell instruction_operand::load_float(cell pointer) +{ + return (cell)load_value(pointer) - boxed_float_offset; +} + code_block *instruction_operand::load_code_block(cell relative_to) { return ((code_block *)load_value(relative_to) - 1); @@ -135,6 +145,11 @@ void instruction_operand::store_value(fixnum absolute_value) } } +void instruction_operand::store_float(cell value) +{ + store_value((fixnum)value + boxed_float_offset); +} + void instruction_operand::store_code_block(code_block *compiled) { store_value((cell)compiled->entry_point()); diff --git a/vm/instruction_operands.hpp b/vm/instruction_operands.hpp index 5dda411c8b..5c120c2ec7 100644 --- a/vm/instruction_operands.hpp +++ b/vm/instruction_operands.hpp @@ -30,6 +30,9 @@ enum relocation_type { type since its used in a situation where relocation arguments cannot be passed in, and so RT_DLSYM is inappropriate (Windows only) */ RT_EXCEPTION_HANDLER, + /* pointer to a float's payload */ + RT_FLOAT, + }; enum relocation_class { @@ -112,6 +115,7 @@ struct relocation_entry { case RT_CARDS_OFFSET: case RT_DECKS_OFFSET: case RT_EXCEPTION_HANDLER: + case RT_FLOAT: return 0; default: critical_error("Bad rel type",rel_type()); @@ -152,12 +156,15 @@ struct instruction_operand { fixnum load_value_masked(cell mask, cell bits, cell shift); fixnum load_value(cell relative_to); fixnum load_value(); + cell load_float(cell relative_to); + cell load_float(); code_block *load_code_block(cell relative_to); code_block *load_code_block(); void store_value_2_2(fixnum value); void store_value_masked(fixnum value, cell mask, cell shift); void store_value(fixnum value); + void store_float(cell value); void store_code_block(code_block *compiled); }; diff --git a/vm/layouts.hpp b/vm/layouts.hpp index 9b574e554d..3e51d1fa4d 100644 --- a/vm/layouts.hpp +++ b/vm/layouts.hpp @@ -246,6 +246,8 @@ struct wrapper : public object { cell object; }; +const fixnum boxed_float_offset = 8 - FLOAT_TYPE; + /* Assembly code makes assumptions about the layout of this struct */ struct boxed_float : object { static const cell type_number = FLOAT_TYPE; diff --git a/vm/primitives.hpp b/vm/primitives.hpp index e98cf508b6..a2bf912749 100644 --- a/vm/primitives.hpp +++ b/vm/primitives.hpp @@ -35,6 +35,7 @@ namespace factor _(byte_array_to_bignum) \ _(callback) \ _(callstack) \ + _(callstack_bounds) \ _(callstack_for) \ _(callstack_to_array) \ _(check_datastack) \ diff --git a/vm/slot_visitor.hpp b/vm/slot_visitor.hpp index d4dd44bed1..cb2db1c705 100644 --- a/vm/slot_visitor.hpp +++ b/vm/slot_visitor.hpp @@ -192,8 +192,17 @@ struct literal_references_visitor { void operator()(instruction_operand op) { - if(op.rel_type() == RT_LITERAL) + switch(op.rel_type()) + { + case RT_LITERAL: op.store_value(visitor->visit_pointer(op.load_value())); + break; + case RT_FLOAT: + op.store_float(visitor->visit_pointer(op.load_float())); + break; + default: + break; + } } }; diff --git a/vm/vm.hpp b/vm/vm.hpp index dd1d48cf03..d9bd17fa51 100755 --- a/vm/vm.hpp +++ b/vm/vm.hpp @@ -606,6 +606,7 @@ struct factor_vm void primitive_innermost_stack_frame_executing(); void primitive_innermost_stack_frame_scan(); void primitive_set_innermost_stack_frame_quot(); + void primitive_callstack_bounds(); template void iterate_callstack(context *ctx, Iterator &iterator); // alien