Merge branch 'master' into symbolic-enums
Joe Groff
commit
5b98e96081
42
GNUmakefile
42
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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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 ;
|
||||
|
|
|
|||
|
|
@ -1,15 +1,19 @@
|
|||
! 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 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
|
||||
|
|
@ -36,7 +40,6 @@ TYPEDEF: void* JSObjectHasInstanceCallback
|
|||
TYPEDEF: void* JSObjectConvertToTypeCallback
|
||||
TYPEDEF: uint unsigned
|
||||
TYPEDEF: ushort JSChar
|
||||
! char[utf16n] for strings
|
||||
|
||||
C-ENUM: JSPropertyAttributes
|
||||
{ kJSPropertyAttributeNone 0 }
|
||||
|
|
@ -202,7 +205,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 ) ;
|
||||
|
||||
|
|
|
|||
|
|
@ -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-standalone ] unit-test
|
||||
|
||||
[ "1+shoes" eval-js-standalone ]
|
||||
[ error>> "ReferenceError: Can't find variable: shoes" = ] must-fail-with
|
||||
|
||||
|
|
@ -1,8 +1,47 @@
|
|||
! 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
|
||||
|
||||
ERROR: javascriptcore-error error ;
|
||||
|
||||
: 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 [ <byte-array> ] 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 ( context string -- result-string )
|
||||
dupd JSStringCreateWithUTF8CString f f 0 JSValueRef <c-object>
|
||||
[ JSEvaluateScript ] keep *void*
|
||||
dup [ nip JSValueRef>string javascriptcore-error ] [ drop JSValueRef>string ] if ;
|
||||
|
||||
: 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 ;
|
||||
|
||||
|
|
|
|||
Loading…
Reference in New Issue