From e4b67b268c76dbd110943557732166c70ef10412 Mon Sep 17 00:00:00 2001 From: Doug Coleman Date: Tue, 20 Apr 2010 13:40:16 -0500 Subject: [PATCH 1/6] Add eval-js and eval-js-file --- extra/javascriptcore/ffi/ffi.factor | 7 ++-- extra/javascriptcore/javascriptcore.factor | 42 +++++++++++++++++++++- 2 files changed, 45 insertions(+), 4 deletions(-) diff --git a/extra/javascriptcore/ffi/ffi.factor b/extra/javascriptcore/ffi/ffi.factor index 844e169eed..6489699d4d 100644 --- a/extra/javascriptcore/ffi/ffi.factor +++ b/extra/javascriptcore/ffi/ffi.factor @@ -1,7 +1,8 @@ ! 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 ; +classes.struct combinators io.encodings.utf16n +io.encodings.utf8 system ; IN: javascriptcore.ffi << @@ -9,6 +10,7 @@ IN: javascriptcore.ffi { [ os macosx? ] [ "/System/Library/Frameworks/JavaScriptCore.framework/Versions/Current/JavaScriptCore" ] } ! { [ os winnt? ] [ "javascriptcore.dll" ] } ! { [ os unix? ] [ "libsqlite3.so" ] } + [ ] } cond cdecl add-library >> @@ -36,7 +38,6 @@ TYPEDEF: void* JSObjectHasInstanceCallback TYPEDEF: void* JSObjectConvertToTypeCallback TYPEDEF: uint unsigned TYPEDEF: ushort JSChar -! char[utf16n] for strings C-ENUM: JSPropertyAttributes { kJSPropertyAttributeNone 0 } @@ -202,7 +203,7 @@ FUNCTION: void JSPropertyNameAccumulatorAddName ( JSPropertyNameAccumulatorRef a FUNCTION: JSStringRef JSStringCreateWithCharacters ( JSChar* chars, size_t numChars ) ; -FUNCTION: JSStringRef JSStringCreateWithUTF8CString ( c-string[utf8] string ) ; +FUNCTION: JSStringRef JSStringCreateWithUTF8CString ( c-string string ) ; FUNCTION: JSStringRef JSStringRetain ( JSStringRef string ) ; diff --git a/extra/javascriptcore/javascriptcore.factor b/extra/javascriptcore/javascriptcore.factor index 773a559d2d..bfd222f9e8 100644 --- a/extra/javascriptcore/javascriptcore.factor +++ b/extra/javascriptcore/javascriptcore.factor @@ -1,8 +1,48 @@ ! Copyright (C) 2010 Doug Coleman. ! See http://factorcode.org/license.txt for BSD license. -USING: javascriptcore.ffi.hack kernel ; +USING: alien.c-types alien.data byte-arrays continuations fry +io.encodings.string io.encodings.utf8 io.files +javascriptcore.ffi javascriptcore.ffi.hack kernel namespaces +sequences ; IN: javascriptcore : with-javascriptcore ( quot -- ) set-callstack-bounds call ; inline + +SYMBOL: js-context + +: with-global-context ( quot -- ) + [ + [ f JSGlobalContextCreate ] dip + [ '[ _ @ ] ] + [ drop '[ _ JSGlobalContextRelease ] ] 2bi + [ ] cleanup + ] with-scope ; inline + +: JSString>string ( JSString -- string ) + dup JSStringGetMaximumUTF8CStringSize [ ] keep + [ JSStringGetUTF8CString drop ] [ drop ] 2bi + utf8 decode [ 0 = ] trim-tail ; + +: JSValueRef>string ( ctx JSValueRef/f -- string/f ) + [ + f JSValueToStringCopy + [ JSString>string ] [ JSStringRelease ] bi + ] [ + drop f + ] if* ; + +: eval-js ( string -- ret/f exception/f ) + [ + [ + [ + swap JSStringCreateWithUTF8CString f f 0 JSValueRef + [ JSEvaluateScript ] keep *void* + ] + [ '[ [ _ ] dip JSValueRef>string ] bi@ ] bi + ] with-global-context + ] with-javascriptcore ; + +: eval-js-path ( path -- ret/f exception/f ) utf8 file-contents eval-js ; + From 73eb31a35c36b1323e34df9f0467deb43f12a7a8 Mon Sep 17 00:00:00 2001 From: Doug Coleman Date: Tue, 20 Apr 2010 14:21:05 -0500 Subject: [PATCH 2/6] Add unit test to javascriptcore, make eval-js throw errors and return a string --- extra/javascriptcore/javascriptcore-tests.factor | 10 ++++++++++ extra/javascriptcore/javascriptcore.factor | 16 ++++++++-------- 2 files changed, 18 insertions(+), 8 deletions(-) create mode 100644 extra/javascriptcore/javascriptcore-tests.factor diff --git a/extra/javascriptcore/javascriptcore-tests.factor b/extra/javascriptcore/javascriptcore-tests.factor new file mode 100644 index 0000000000..f04ada89f2 --- /dev/null +++ b/extra/javascriptcore/javascriptcore-tests.factor @@ -0,0 +1,10 @@ +! Copyright (C) 2010 Doug Coleman. +! See http://factorcode.org/license.txt for BSD license. +USING: accessors javascriptcore kernel tools.test ; +IN: javascriptcore.tests + +[ "2" ] [ "1+1" eval-js ] unit-test + +[ "1+shoes" eval-js ] +[ error>> "ReferenceError: Can't find variable: shoes" = ] must-fail-with + diff --git a/extra/javascriptcore/javascriptcore.factor b/extra/javascriptcore/javascriptcore.factor index bfd222f9e8..eabb97df61 100644 --- a/extra/javascriptcore/javascriptcore.factor +++ b/extra/javascriptcore/javascriptcore.factor @@ -6,6 +6,8 @@ javascriptcore.ffi javascriptcore.ffi.hack kernel namespaces sequences ; IN: javascriptcore +ERROR: javascriptcore-error error ; + : with-javascriptcore ( quot -- ) set-callstack-bounds call ; inline @@ -33,16 +35,14 @@ SYMBOL: js-context drop f ] if* ; -: eval-js ( string -- ret/f exception/f ) - [ +: eval-js ( string -- result-string ) + '[ [ - [ - swap JSStringCreateWithUTF8CString f f 0 JSValueRef - [ JSEvaluateScript ] keep *void* - ] - [ '[ [ _ ] dip JSValueRef>string ] bi@ ] bi + dup _ JSStringCreateWithUTF8CString f f 0 JSValueRef + [ JSEvaluateScript ] keep *void* + dup [ nip JSValueRef>string javascriptcore-error ] [ drop JSValueRef>string ] if ] with-global-context ] with-javascriptcore ; -: eval-js-path ( path -- ret/f exception/f ) utf8 file-contents eval-js ; +: eval-js-path ( path -- result-string ) utf8 file-contents eval-js ; From b71933f4d7331d04851ac6a55aae21461c3f7018 Mon Sep 17 00:00:00 2001 From: Joe Groff Date: Tue, 20 Apr 2010 13:51:10 -0700 Subject: [PATCH 3/6] cuda.ptx: better representation of operands --- extra/cuda/ptx/ptx-tests.factor | 79 +++++++++------- extra/cuda/ptx/ptx.factor | 159 ++++++++++++++++++++------------ 2 files changed, 144 insertions(+), 94 deletions(-) diff --git a/extra/cuda/ptx/ptx-tests.factor b/extra/cuda/ptx/ptx-tests.factor index 28391a5f58..1ba7ecfcc8 100644 --- a/extra/cuda/ptx/ptx-tests.factor +++ b/extra/cuda/ptx/ptx-tests.factor @@ -1,4 +1,4 @@ -USING: cuda.ptx tools.test ; +USING: cuda.ptx io.streams.string tools.test ; IN: cuda.ptx.tests [ """ .version 2.0 @@ -113,6 +113,17 @@ IN: cuda.ptx.tests } ptx>string ] unit-test +[ "a" ] [ [ "a" write-ptx-operand ] with-string-writer ] unit-test +[ "2" ] [ [ 2 write-ptx-operand ] with-string-writer ] unit-test +[ "0d4000000000000000" ] [ [ 2.0 write-ptx-operand ] with-string-writer ] unit-test +[ "!a" ] [ [ T{ ptx-negation f "a" } write-ptx-operand ] with-string-writer ] unit-test +[ "{a, b, c, d}" ] [ [ T{ ptx-vector f { "a" "b" "c" "d" } } write-ptx-operand ] with-string-writer ] unit-test +[ "[a]" ] [ [ T{ ptx-indirect f "a" 0 } write-ptx-operand ] with-string-writer ] unit-test +[ "[a+1]" ] [ [ T{ ptx-indirect f "a" 1 } write-ptx-operand ] with-string-writer ] unit-test +[ "[a-1]" ] [ [ T{ ptx-indirect f "a" -1 } write-ptx-operand ] with-string-writer ] unit-test +[ "a[1]" ] [ [ T{ ptx-element f "a" 1 } write-ptx-operand ] with-string-writer ] unit-test +[ "{a, b[2], 3, 0d4000000000000000}" ] [ [ T{ ptx-vector f { "a" T{ ptx-element f "b" 2 } 3 2.0 } } write-ptx-operand ] with-string-writer ] unit-test + [ """ .version 2.0 .target sm_20 abs.s32 a, b; @@ -127,11 +138,11 @@ foo: abs.s32 a, b; { body { T{ abs { type .s32 } { dest "a" } { a "b" } } T{ abs - { predicate T{ ptx-predicate { variable "p" } } } + { predicate "p" } { type .s32 } { dest "a" } { a "b" } } T{ abs - { predicate T{ ptx-predicate { negated? t } { variable "p" } } } + { predicate T{ ptx-negation f "p" } } { type .s32 } { dest "a" } { a "b" } } T{ abs @@ -206,9 +217,9 @@ foo: abs.s32 a, b; { 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" } } + T{ atom { op .and } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } { b "c" } } + T{ atom { storage-space .global } { op .or } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } { b "c" } } + T{ atom { storage-space .shared } { op .cas } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } { b "c" } { c "d" } } } } } ptx>string @@ -229,8 +240,8 @@ foo: abs.s32 a, b; { 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.red { op .popc } { type .u32 } { dest "a" } { a "b" } { c T{ ptx-negation f "d" } } } + T{ bar.red { op .popc } { type .u32 } { dest "a" } { a "b" } { b "c" } { c T{ ptx-negation f "d" } } } T{ bar.sync { a "a" } } T{ bar.sync { a "a" } { b "b" } } } } @@ -327,6 +338,7 @@ foo: abs.s32 a, b; call (a), foo, (b); call (a), foo, (b, c); call (a), foo, (b, c, d); + call (a[2]), foo, (b, c, d[3]); call foo, (b, c, d); """ ] [ T{ ptx @@ -339,6 +351,7 @@ foo: abs.s32 a, b; 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 { return T{ ptx-element f "a" 2 } } { target "foo" } { params { "b" "c" T{ ptx-element f "d" 3 } } } } T{ call { target "foo" } { params { "b" "c" "d" } } } } } } ptx>string @@ -549,13 +562,13 @@ foo: abs.s32 a, b; { 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 { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } } + T{ ld { type T{ .v2 { of .u32 } } } { dest "a" } { a T{ ptx-indirect f "b" } } } + T{ ld { type T{ .v4 { of .u32 } } } { dest "a" } { a T{ ptx-indirect f "b" } } } + T{ ld { type T{ .v4 { of .u32 } } } { dest T{ ptx-vector f { "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]" } } + T{ ld { storage-space T{ .const } } { cache-op .lu } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } } + T{ ld { volatile? t } { storage-space T{ .const { bank 5 } } } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } } } } } ptx>string ] unit-test @@ -574,13 +587,13 @@ foo: abs.s32 a, b; { 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 { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } } + T{ ldu { type T{ .v2 { of .u32 } } } { dest "a" } { a T{ ptx-indirect f "b" } } } + T{ ldu { type T{ .v4 { of .u32 } } } { dest "a" } { a T{ ptx-indirect f "b" } } } + T{ ldu { type T{ .v4 { of .u32 } } } { dest T{ ptx-vector f { "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]" } } + T{ ldu { storage-space T{ .const } } { cache-op .lu } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } } + T{ ldu { volatile? t } { storage-space T{ .const { bank 5 } } } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } } } } } ptx>string ] unit-test @@ -723,9 +736,9 @@ foo: abs.s32 a, b; { 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]" } } + T{ prefetch { level .L1 } { a T{ ptx-indirect f "a" } } } + T{ prefetch { storage-space .local } { level .L2 } { a T{ ptx-indirect f "a" } } } + T{ prefetchu { level .L1 } { a T{ ptx-indirect f "a" } } } } } } ptx>string ] unit-test @@ -781,8 +794,8 @@ foo: abs.s32 a, b; { 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" } } + T{ red { op .and } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } } + T{ red { storage-space .global } { op .and } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } } } } } ptx>string ] unit-test @@ -861,7 +874,7 @@ foo: abs.s32 a, b; 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" } } + T{ set { cmp-op .gt } { bool-op .and } { ftz? t } { dest-type .u32 } { type .f32 } { dest "a" } { a "b" } { b "c" } { c T{ ptx-negation f "d" } } } } } } ptx>string ] unit-test @@ -982,13 +995,13 @@ foo: abs.s32 a, b; { 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" } } + T{ st { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } } + T{ st { type T{ .v2 { of .u32 } } } { dest T{ ptx-indirect f "a" } } { a "b" } } + T{ st { type T{ .v4 { of .u32 } } } { dest T{ ptx-indirect f "a" } } { a "b" } } + T{ st { type T{ .v4 { of .u32 } } } { dest T{ ptx-indirect f "a" } } { a T{ ptx-vector f { "b" "c" "d" "e" } } } } + T{ st { cache-op .lu } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } } + T{ st { storage-space .local } { cache-op .lu } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } } + T{ st { volatile? t } { storage-space .local } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } } } } } ptx>string ] unit-test diff --git a/extra/cuda/ptx/ptx.factor b/extra/cuda/ptx/ptx.factor index 4618f8b5b6..49a53d7fbf 100644 --- a/extra/cuda/ptx/ptx.factor +++ b/extra/cuda/ptx/ptx.factor @@ -64,14 +64,30 @@ TUPLE: ptx-variable { initializer ?string } ; UNION: ?ptx-variable POSTPONE: f ptx-variable ; -TUPLE: ptx-predicate - { negated? boolean } - { variable string } ; -UNION: ?ptx-predicate POSTPONE: f ptx-predicate ; +TUPLE: ptx-negation + { var string } ; + +TUPLE: ptx-vector + elements ; + +TUPLE: ptx-element + { var string } + { index integer } ; + +UNION: ptx-var + string ptx-element ; + +TUPLE: ptx-indirect + { base ptx-var } + { offset integer } ; + +UNION: ptx-operand + integer float ptx-var ptx-negation ptx-vector ptx-indirect ; +UNION: ?ptx-operand POSTPONE: f ptx-operand ; TUPLE: ptx-instruction { label ?string } - { predicate ?ptx-predicate } ; + { predicate ?ptx-operand } ; TUPLE: ptx-entry { name string } @@ -112,25 +128,25 @@ UNION: ?ptx-rounding-mode POSTPONE: f ptx-rounding-mode ; TUPLE: ptx-typed-instruction < ptx-instruction { type ptx-type } - { dest string } ; + { dest ptx-operand } ; TUPLE: ptx-2op-instruction < ptx-typed-instruction - { a string } ; + { a ptx-operand } ; TUPLE: ptx-3op-instruction < ptx-typed-instruction - { a string } - { b string } ; + { a ptx-operand } + { b ptx-operand } ; TUPLE: ptx-4op-instruction < ptx-typed-instruction - { a string } - { b string } - { c string } ; + { a ptx-operand } + { b ptx-operand } + { c ptx-operand } ; TUPLE: ptx-5op-instruction < ptx-typed-instruction - { a string } - { b string } - { c string } - { d string } ; + { a ptx-operand } + { b ptx-operand } + { c ptx-operand } + { d ptx-operand } ; TUPLE: ptx-addsub-instruction < ptx-3op-instruction { sat? boolean } @@ -181,7 +197,7 @@ INSTANCE: .hi ptx-cmp-op TUPLE: ptx-set-instruction < ptx-3op-instruction { cmp-op ptx-cmp-op } { bool-op ?ptx-op } - { c ?string } + { c ?ptx-operand } { ftz? boolean } ; VARIANT: ptx-cache-op @@ -216,17 +232,17 @@ TUPLE: and < ptx-3op-instruction ; TUPLE: atom < ptx-3op-instruction { storage-space ?ptx-storage-space } { op ptx-op } - { c ?string } ; + { c ?ptx-operand } ; TUPLE: bar.arrive < ptx-instruction - { a string } - { b string } ; + { a ptx-operand } + { b ptx-operand } ; TUPLE: bar.red < ptx-2op-instruction { op ptx-op } - { b ?string } - { c string } ; + { b ?ptx-operand } + { c ptx-operand } ; TUPLE: bar.sync < ptx-instruction - { a string } - { b ?string } ; + { a ptx-operand } + { b ?ptx-operand } ; TUPLE: bfe < ptx-4op-instruction ; TUPLE: bfi < ptx-5op-instruction ; TUPLE: bfind < ptx-2op-instruction @@ -235,7 +251,7 @@ TUPLE: bra < ptx-branch-instruction ; TUPLE: brev < ptx-2op-instruction ; TUPLE: brkpt < ptx-instruction ; TUPLE: call < ptx-branch-instruction - { return ?string } + { return ?ptx-operand } params ; TUPLE: clz < ptx-2op-instruction ; TUPLE: cnot < ptx-2op-instruction ; @@ -255,8 +271,8 @@ TUPLE: exit < ptx-instruction ; TUPLE: fma <{ ptx-mad-instruction ptx-float-env } ; TUPLE: isspacep < ptx-instruction { storage-space ptx-storage-space } - { dest string } - { a string } ; + { dest ptx-operand } + { a ptx-operand } ; TUPLE: ld < ptx-ldst-instruction ; TUPLE: ldu < ptx-ldst-instruction ; TUPLE: lg2 <{ ptx-2op-instruction ptx-float-env } ; @@ -273,14 +289,14 @@ TUPLE: neg <{ ptx-2op-instruction ptx-float-ftz } ; TUPLE: not < ptx-2op-instruction ; TUPLE: or < ptx-3op-instruction ; TUPLE: pmevent < ptx-instruction - { a string } ; + { a ptx-operand } ; TUPLE: popc < ptx-2op-instruction ; TUPLE: prefetch < ptx-instruction - { a string } + { a ptx-operand } { storage-space ?ptx-storage-space } { level ptx-cache-level } ; TUPLE: prefetchu < ptx-instruction - { a string } + { a ptx-operand } { level ptx-cache-level } ; TUPLE: prmt < ptx-4op-instruction { mode ?ptx-prmt-mode } ; @@ -296,7 +312,7 @@ TUPLE: selp < ptx-4op-instruction ; TUPLE: set < ptx-set-instruction { dest-type ptx-type } ; TUPLE: setp < ptx-set-instruction - { |dest ?string } ; + { |dest ?ptx-operand } ; TUPLE: shl < ptx-3op-instruction ; TUPLE: shr < ptx-3op-instruction ; TUPLE: sin <{ ptx-2op-instruction ptx-float-env } ; @@ -340,6 +356,27 @@ M: ptx-func ptx-semicolon? drop f ; M: .file ptx-semicolon? drop f ; M: .loc ptx-semicolon? drop f ; +GENERIC: write-ptx-operand ( operand -- ) + +M: string write-ptx-operand write ; +M: integer write-ptx-operand number>string write ; +M: float write-ptx-operand "0d" write double>bits >hex 16 CHAR: 0 pad-head write ; +M: ptx-negation write-ptx-operand "!" write var>> write ; +M: ptx-vector write-ptx-operand + "{" write + elements>> [ ", " write ] [ write-ptx-operand ] interleave + "}" write ; +M: ptx-element write-ptx-operand dup var>> write "[" write index>> number>string write "]" write ; +M: ptx-indirect write-ptx-operand + "[" write + dup base>> write-ptx-operand + offset>> { + { [ dup zero? ] [ drop ] } + { [ dup 0 < ] [ number>string write ] } + [ "+" write number>string write ] + } cond + "]" write ; + GENERIC: (write-ptx-element) ( elt -- ) : write-ptx-element ( elt -- ) @@ -376,7 +413,7 @@ M: ptx-target (write-ptx-element) [ arch>> [ name>> ] [ f ] if* ] [ map_f64_to_f32?>> [ "map_f64_to_f32" ] [ f ] if ] [ texmode>> [ name>> ] [ f ] if* ] tri - 3array sift ", " join write ; + 3array sift [ ", " write ] [ write ] interleave ; : write-ptx-dim ( dim -- ) { @@ -435,7 +472,7 @@ M: .maxnreg (write-ptx-element) ".maxnreg " write n>> number>string write ; M: .maxntid (write-ptx-element) ".maxntid " write - dup sequence? [ [ number>string ] map ", " join write ] [ number>string write ] if ; + dup sequence? [ [ ", " write ] [ number>string write ] interleave ] [ number>string write ] if ; M: .pragma (write-ptx-element) ".pragma \"" write pragma>> write "\"" write ; @@ -444,28 +481,28 @@ M: ptx-instruction ptx-element-label : write-insn ( insn name -- insn ) over predicate>> - [ "@" write dup negated?>> [ "!" write ] when variable>> write " " write ] when* + [ "@" write write-ptx-operand " " write ] when* write ; : write-2op ( insn -- ) dup type>> (write-ptx-element) " " write - dup dest>> write ", " write - dup a>> write + dup dest>> write-ptx-operand ", " write + dup a>> write-ptx-operand drop ; : write-3op ( insn -- ) dup write-2op ", " write - dup b>> write + dup b>> write-ptx-operand drop ; : write-4op ( insn -- ) dup write-3op ", " write - dup c>> write + dup c>> write-ptx-operand drop ; : write-5op ( insn -- ) dup write-4op ", " write - dup d>> write + dup d>> write-ptx-operand drop ; : write-ftz ( insn -- ) @@ -534,22 +571,22 @@ M: atom (write-ptx-element) dup storage-space>> (write-ptx-element) dup op>> (write-ptx-element) dup write-3op - c>> [ ", " write write ] when* ; + c>> [ ", " write write-ptx-operand ] when* ; M: bar.arrive (write-ptx-element) "bar.arrive " write-insn - dup a>> write ", " write - dup b>> write + dup a>> write-ptx-operand ", " write + dup b>> write-ptx-operand drop ; M: bar.red (write-ptx-element) "bar.red" write-insn dup op>> (write-ptx-element) dup write-2op - dup b>> [ ", " write write ] when* - ", " write c>> write ; + dup b>> [ ", " write write-ptx-operand ] when* + ", " write c>> write-ptx-operand ; M: bar.sync (write-ptx-element) "bar.sync " write-insn - dup a>> write - dup b>> [ ", " write write ] when* + dup a>> write-ptx-operand + dup b>> [ ", " write write-ptx-operand ] when* drop ; M: bfe (write-ptx-element) "bfe" write-insn @@ -573,9 +610,9 @@ M: brkpt (write-ptx-element) M: call (write-ptx-element) "call" write-insn dup write-uni " " write - dup return>> [ "(" write write "), " write ] when* + dup return>> [ "(" write write-ptx-operand "), " write ] when* dup target>> write - dup params>> [ ", (" write ", " join write ")" write ] unless-empty + dup params>> [ ", (" write [ ", " write ] [ write-ptx-operand ] interleave ")" write ] unless-empty drop ; M: clz (write-ptx-element) "clz" write-insn @@ -619,7 +656,7 @@ M: isspacep (write-ptx-element) "isspacep" write-insn dup storage-space>> (write-ptx-element) " " write - dup dest>> write ", " write a>> write ; + dup dest>> write-ptx-operand ", " write a>> write-ptx-operand ; M: ld (write-ptx-element) "ld" write-insn write-ldst ; @@ -679,19 +716,19 @@ M: prefetch (write-ptx-element) "prefetch" write-insn dup storage-space>> (write-ptx-element) dup level>> (write-ptx-element) - " " write a>> write ; + " " write a>> write-ptx-operand ; M: prefetchu (write-ptx-element) "prefetchu" write-insn dup level>> (write-ptx-element) - " " write a>> write ; + " " write a>> write-ptx-operand ; M: prmt (write-ptx-element) "prmt" write-insn 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 + dup dest>> write-ptx-operand ", " write + dup a>> write-ptx-operand ", " write + dup b>> write-ptx-operand ", " write + dup c>> write-ptx-operand drop ; M: rcp (write-ptx-element) "rcp" write-insn @@ -722,16 +759,16 @@ M: set (write-ptx-element) dup write-set dup dest-type>> (write-ptx-element) dup write-3op - c>> [ ", " write write ] when* ; + c>> [ ", " write write-ptx-operand ] when* ; M: setp (write-ptx-element) "setp" write-insn dup write-set dup type>> (write-ptx-element) " " write - dup dest>> write - dup |dest>> [ "|" write write ] when* ", " write - dup a>> write ", " write - dup b>> write - c>> [ ", " write write ] when* ; + dup dest>> write-ptx-operand + dup |dest>> [ "|" write write-ptx-operand ] when* ", " write + dup a>> write-ptx-operand ", " write + dup b>> write-ptx-operand + c>> [ ", " write write-ptx-operand ] when* ; M: shl (write-ptx-element) "shl" write-insn write-3op ; From 6ecf43b91f777670fae1125738cebba22dc88eae Mon Sep 17 00:00:00 2001 From: Doug Coleman Date: Tue, 20 Apr 2010 16:04:35 -0500 Subject: [PATCH 4/6] fix add-library in javascriptcore --- extra/javascriptcore/ffi/ffi.factor | 14 ++++++++------ 1 file changed, 8 insertions(+), 6 deletions(-) diff --git a/extra/javascriptcore/ffi/ffi.factor b/extra/javascriptcore/ffi/ffi.factor index 6489699d4d..7a038b0883 100644 --- a/extra/javascriptcore/ffi/ffi.factor +++ b/extra/javascriptcore/ffi/ffi.factor @@ -2,16 +2,18 @@ ! See http://factorcode.org/license.txt for BSD license. USING: alien alien.c-types alien.libraries alien.syntax classes.struct combinators io.encodings.utf16n -io.encodings.utf8 system ; +io.encodings.utf8 kernel 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 + { [ os macosx? ] [ + "/System/Library/Frameworks/JavaScriptCore.framework/Versions/Current/JavaScriptCore" cdecl add-library + ] } + ! { [ os winnt? ] [ "javascriptcore.dll" ] } + ! { [ os unix? ] [ "libsqlite3.so" ] } + [ drop ] +} cond >> LIBRARY: javascriptcore From b32782ac9b5729da2a73134d28828b43c5593581 Mon Sep 17 00:00:00 2001 From: Doug Coleman Date: Tue, 20 Apr 2010 17:25:28 -0500 Subject: [PATCH 5/6] Move the guts of eval-js to its own word, add eval-js-standalone for unit tests --- .../javascriptcore/javascriptcore-tests.factor | 4 ++-- extra/javascriptcore/javascriptcore.factor | 17 ++++++++--------- 2 files changed, 10 insertions(+), 11 deletions(-) diff --git a/extra/javascriptcore/javascriptcore-tests.factor b/extra/javascriptcore/javascriptcore-tests.factor index f04ada89f2..53ae12d2b0 100644 --- a/extra/javascriptcore/javascriptcore-tests.factor +++ b/extra/javascriptcore/javascriptcore-tests.factor @@ -3,8 +3,8 @@ USING: accessors javascriptcore kernel tools.test ; IN: javascriptcore.tests -[ "2" ] [ "1+1" eval-js ] unit-test +[ "2" ] [ "1+1" eval-js-standalone ] unit-test -[ "1+shoes" eval-js ] +[ "1+shoes" eval-js-standalone ] [ error>> "ReferenceError: Can't find variable: shoes" = ] must-fail-with diff --git a/extra/javascriptcore/javascriptcore.factor b/extra/javascriptcore/javascriptcore.factor index eabb97df61..65b6fe5fff 100644 --- a/extra/javascriptcore/javascriptcore.factor +++ b/extra/javascriptcore/javascriptcore.factor @@ -35,14 +35,13 @@ SYMBOL: js-context drop f ] if* ; -: eval-js ( string -- result-string ) - '[ - [ - dup _ JSStringCreateWithUTF8CString f f 0 JSValueRef - [ JSEvaluateScript ] keep *void* - dup [ nip JSValueRef>string javascriptcore-error ] [ drop JSValueRef>string ] if - ] with-global-context - ] with-javascriptcore ; +: eval-js ( context string -- result-string ) + dupd JSStringCreateWithUTF8CString f f 0 JSValueRef + [ JSEvaluateScript ] keep *void* + dup [ nip JSValueRef>string javascriptcore-error ] [ drop JSValueRef>string ] if ; -: eval-js-path ( path -- result-string ) utf8 file-contents eval-js ; +: eval-js-standalone ( string -- result-string ) + '[ [ _ eval-js ] with-global-context ] with-javascriptcore ; + +: eval-js-path-standalone ( path -- result-string ) utf8 file-contents eval-js-standalone ; From 78e26edb9a7829821caddbb0205650dd584c4887 Mon Sep 17 00:00:00 2001 From: Joe Groff Date: Tue, 20 Apr 2010 15:28:09 -0700 Subject: [PATCH 6/6] GNUmakefile: build factor shared library by default again --- GNUmakefile | 42 +++++++++++++++++++++++------------------- 1 file changed, 23 insertions(+), 19 deletions(-) diff --git a/GNUmakefile b/GNUmakefile index 30f44e9eba..300a62f71c 100755 --- a/GNUmakefile +++ b/GNUmakefile @@ -106,61 +106,63 @@ help: @echo "NO_UI=1 don't link with X11 libraries (ignored on Mac OS X)" @echo "X11=1 force link with X11 libraries instead of Cocoa (only on Mac OS X)" +ALL = factor factor-ffi-test factor-lib + openbsd-x86-32: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.openbsd.x86.32 + $(MAKE) $(ALL) CONFIG=vm/Config.openbsd.x86.32 openbsd-x86-64: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.openbsd.x86.64 + $(MAKE) $(ALL) CONFIG=vm/Config.openbsd.x86.64 freebsd-x86-32: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.freebsd.x86.32 + $(MAKE) $(ALL) CONFIG=vm/Config.freebsd.x86.32 freebsd-x86-64: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.freebsd.x86.64 + $(MAKE) $(ALL) CONFIG=vm/Config.freebsd.x86.64 netbsd-x86-32: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.netbsd.x86.32 + $(MAKE) $(ALL) CONFIG=vm/Config.netbsd.x86.32 netbsd-x86-64: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.netbsd.x86.64 + $(MAKE) $(ALL) CONFIG=vm/Config.netbsd.x86.64 macosx-ppc: - $(MAKE) factor factor-ffi-test macosx.app CONFIG=vm/Config.macosx.ppc + $(MAKE) $(ALL) macosx.app CONFIG=vm/Config.macosx.ppc macosx-x86-32: - $(MAKE) factor factor-ffi-test macosx.app CONFIG=vm/Config.macosx.x86.32 + $(MAKE) $(ALL) macosx.app CONFIG=vm/Config.macosx.x86.32 macosx-x86-64: - $(MAKE) factor factor-ffi-test macosx.app CONFIG=vm/Config.macosx.x86.64 + $(MAKE) $(ALL) macosx.app CONFIG=vm/Config.macosx.x86.64 linux-x86-32: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.linux.x86.32 + $(MAKE) $(ALL) CONFIG=vm/Config.linux.x86.32 linux-x86-64: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.linux.x86.64 + $(MAKE) $(ALL) CONFIG=vm/Config.linux.x86.64 linux-ppc: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.linux.ppc + $(MAKE) $(ALL) CONFIG=vm/Config.linux.ppc linux-arm: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.linux.arm + $(MAKE) $(ALL) CONFIG=vm/Config.linux.arm solaris-x86-32: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.solaris.x86.32 + $(MAKE) $(ALL) CONFIG=vm/Config.solaris.x86.32 solaris-x86-64: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.solaris.x86.64 + $(MAKE) $(ALL) CONFIG=vm/Config.solaris.x86.64 winnt-x86-32: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.windows.nt.x86.32 + $(MAKE) $(ALL) CONFIG=vm/Config.windows.nt.x86.32 $(MAKE) factor-console CONFIG=vm/Config.windows.nt.x86.32 winnt-x86-64: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.windows.nt.x86.64 + $(MAKE) $(ALL) CONFIG=vm/Config.windows.nt.x86.64 $(MAKE) factor-console CONFIG=vm/Config.windows.nt.x86.64 wince-arm: - $(MAKE) factor factor-ffi-test CONFIG=vm/Config.windows.ce.arm + $(MAKE) $(ALL) CONFIG=vm/Config.windows.ce.arm ifdef CONFIG @@ -173,6 +175,8 @@ macosx.app: factor $(ENGINE): $(DLL_OBJS) $(TOOLCHAIN_PREFIX)$(LINKER) $(ENGINE) $(DLL_OBJS) +factor-lib: $(ENGINE) + factor: $(EXE_OBJS) $(DLL_OBJS) $(TOOLCHAIN_PREFIX)$(CPP) $(LIBS) $(LIBPATH) -L. $(DLL_OBJS) \ $(CFLAGS) -o $(EXECUTABLE) $(EXE_OBJS) @@ -217,4 +221,4 @@ clean: tags: etags vm/*.{cpp,hpp,mm,S,c} -.PHONY: factor factor-console factor-ffi-test tags clean macosx.app +.PHONY: factor factor-lib factor-console factor-ffi-test tags clean macosx.app