From 1ae8cdc5876cefd6e6bfe330c07ceb9228cbc50e Mon Sep 17 00:00:00 2001 From: Slava Pestov Date: Sat, 17 Apr 2010 14:51:29 -0500 Subject: [PATCH 1/5] mason: use web service instead of shell script for status notifications, to scale better in the cloud --- extra/mason/config/config.factor | 10 +-- extra/mason/notify/notify.factor | 51 +++++------- extra/mason/server/notify/authors.txt | 1 - extra/mason/server/notify/notify.factor | 80 ------------------- extra/mason/server/server.factor | 5 +- extra/webapps/mason/download-package.xml | 2 +- extra/webapps/mason/mason.factor | 6 +- extra/webapps/mason/status-update/authors.txt | 1 + .../mason/status-update/status-update.factor | 74 +++++++++++++++++ 9 files changed, 108 insertions(+), 122 deletions(-) delete mode 100644 extra/mason/server/notify/authors.txt delete mode 100644 extra/mason/server/notify/notify.factor create mode 100644 extra/webapps/mason/status-update/authors.txt create mode 100644 extra/webapps/mason/status-update/status-update.factor diff --git a/extra/mason/config/config.factor b/extra/mason/config/config.factor index 5ec44df0a9..48f4d307c8 100644 --- a/extra/mason/config/config.factor +++ b/extra/mason/config/config.factor @@ -1,4 +1,4 @@ -! Copyright (C) 2008 Eduardo Cavazos, Slava Pestov. +! Copyright (C) 2008, 2010 Eduardo Cavazos, Slava Pestov. ! See http://factorcode.org/license.txt for BSD license. USING: system io.files io.pathnames namespaces kernel accessors assocs ; @@ -39,11 +39,11 @@ target-os get-global [ ! Keep test-log around? SYMBOL: builder-debug -! Host to send status notifications to. -SYMBOL: status-host +! URL for status notifications. +SYMBOL: status-url -! Username to log in. -SYMBOL: status-username +! Password for status notifications. +SYMBOL: status-secret SYMBOL: upload-help? diff --git a/extra/mason/notify/notify.factor b/extra/mason/notify/notify.factor index d7319c0f20..144f0de122 100644 --- a/extra/mason/notify/notify.factor +++ b/extra/mason/notify/notify.factor @@ -1,57 +1,50 @@ ! Copyright (C) 2009, 2010 Slava Pestov. ! See http://factorcode.org/license.txt for BSD license. -USING: arrays accessors io io.sockets io.encodings.utf8 io.files -io.launcher kernel make mason.config mason.common mason.email -mason.twitter namespaces sequences prettyprint fry ; +USING: accessors fry http.client io io.encodings.utf8 io.files +kernel mason.common mason.config mason.email mason.twitter +namespaces prettyprint sequences ; IN: mason.notify -: status-notify ( input-file args -- ) - status-host get [ - [ - "ssh" , status-host get , "-l" , status-username get , - "./mason-notify" , - short-host-name , - target-cpu get , - target-os get , - ] { } make prepend - [ 5 ] 2dip '[ - - _ >>stdin - _ >>command - short-running-process - ] retry - ] [ 2drop ] if ; +: status-notify ( report arg message -- ) + [ + short-host-name "host-name" set + target-cpu get "target-cpu" set + target-os get "target-os" set + status-secret get "secret" set + "message" set + "arg" set + "report" set + ] H{ } make-assoc + [ 5 ] dip '[ _ status-url get http-post 2drop ] retry ; : notify-heartbeat ( -- ) - f { "heartbeat" } status-notify ; + f f "heartbeat" status-notify ; : notify-begin-build ( git-id -- ) [ "Starting build of GIT ID " write print flush ] - [ f swap "git-id" swap 2array status-notify ] + [ f swap "git-id" status-notify ] bi ; : notify-make-vm ( -- ) "Compiling VM" print flush - f { "make-vm" } status-notify ; + f f "make-vm" status-notify ; : notify-boot ( -- ) "Bootstrapping" print flush - f { "boot" } status-notify ; + f f "boot" status-notify ; : notify-test ( -- ) "Running tests" print flush - f { "test" } status-notify ; + f f "test" status-notify ; : notify-report ( status -- ) [ "Build finished with status: " write . flush ] [ - [ "report" ] dip - [ [ utf8 file-contents ] dip email-report ] - [ "report" swap name>> 2array status-notify ] - 2bi + [ "report" utf8 file-contents ] dip + [ name>> "report" status-notify ] [ email-report ] 2bi ] bi ; : notify-release ( archive-name -- ) [ "Uploaded " prepend [ print flush ] [ mason-tweet ] bi ] - [ f swap "release" swap 2array status-notify ] + [ f swap "release" status-notify ] bi ; diff --git a/extra/mason/server/notify/authors.txt b/extra/mason/server/notify/authors.txt deleted file mode 100644 index d4f5d6b3ae..0000000000 --- a/extra/mason/server/notify/authors.txt +++ /dev/null @@ -1 +0,0 @@ -Slava Pestov \ No newline at end of file diff --git a/extra/mason/server/notify/notify.factor b/extra/mason/server/notify/notify.factor deleted file mode 100644 index bfa1027d92..0000000000 --- a/extra/mason/server/notify/notify.factor +++ /dev/null @@ -1,80 +0,0 @@ -! Copyright (C) 2009, 2010 Slava Pestov. -! See http://factorcode.org/license.txt for BSD license. -USING: accessors calendar combinators combinators.smart -command-line db.tuples io io.encodings.utf8 io.files kernel -mason.server namespaces present sequences ; -IN: mason.server.notify - -SYMBOLS: host-name target-os target-cpu message message-arg ; - -: parse-args ( command-line -- ) - dup last message-arg set - [ - { - [ host-name set ] - [ target-cpu set ] - [ target-os set ] - [ message set ] - } spread - ] input>host-name - target-os get >>os - target-cpu get >>cpu - dup select-tuple [ ] [ dup insert-tuple ] ?if ; - -: heartbeat ( builder -- ) now >>heartbeat-timestamp drop ; - -: git-id ( builder id -- ) >>current-git-id +starting+ >>status drop ; - -: make-vm ( builder -- ) +make-vm+ >>status drop ; - -: boot ( builder -- ) +boot+ >>status drop ; - -: test ( builder -- ) +test+ >>status drop ; - -: report ( builder status content -- ) - [ >>status ] [ >>last-report ] bi* - dup status>> +clean+ = [ - dup current-git-id>> >>clean-git-id - dup current-timestamp>> >>clean-timestamp - ] when - dup current-git-id>> >>last-git-id - dup current-timestamp>> >>last-timestamp - drop ; - -: release ( builder name -- ) - >>last-release - dup clean-git-id>> >>release-git-id - drop ; - -: update-builder ( builder -- ) - message get { - { "heartbeat" [ heartbeat ] } - { "git-id" [ message-arg get git-id ] } - { "make-vm" [ make-vm ] } - { "boot" [ boot ] } - { "test" [ test ] } - { "report" [ message-arg get contents report ] } - { "release" [ message-arg get release ] } - } case ; - -: handle-update ( command-line timestamp -- ) - [ - [ parse-args find-builder ] dip >>current-timestamp - [ update-builder ] [ update-tuple ] bi - ] with-mason-db ; - -CONSTANT: log-file "resource:mason.log" - -: log-update ( command-line timestamp -- ) - log-file utf8 [ - present write ": " write " " join print - ] with-file-appender ; - -: main ( -- ) - command-line get now [ log-update ] [ handle-update ] 2bi ; - -MAIN: main diff --git a/extra/mason/server/server.factor b/extra/mason/server/server.factor index 26be4df57c..d0fe29b917 100644 --- a/extra/mason/server/server.factor +++ b/extra/mason/server/server.factor @@ -17,8 +17,7 @@ clean-git-id clean-timestamp last-release release-git-id last-git-id last-timestamp last-report current-git-id current-timestamp -status -heartbeat-timestamp ; +status ; builder "BUILDERS" { { "host-name" "HOST_NAME" TEXT +user-assigned-id+ } @@ -39,8 +38,6 @@ builder "BUILDERS" { ! Can't name it CURRENT_TIMESTAMP because of bug in db library { "current-timestamp" "CURR_TIMESTAMP" TIMESTAMP } { "status" "STATUS" TEXT } - - { "heartbeat-timestamp" "HEARTBEAT_TIMESTAMP" TIMESTAMP } } define-persistent : mason-db ( -- db ) "resource:mason.db" ; diff --git a/extra/webapps/mason/download-package.xml b/extra/webapps/mason/download-package.xml index 43212cfc61..27102056f8 100644 --- a/extra/webapps/mason/download-package.xml +++ b/extra/webapps/mason/download-package.xml @@ -28,7 +28,7 @@ - + diff --git a/extra/webapps/mason/mason.factor b/extra/webapps/mason/mason.factor index ecb1348532..81eb36a17d 100644 --- a/extra/webapps/mason/mason.factor +++ b/extra/webapps/mason/mason.factor @@ -4,7 +4,7 @@ USING: accessors furnace.auth furnace.db http.server.dispatchers mason.server webapps.mason.grids webapps.mason.make-release webapps.mason.package webapps.mason.release webapps.mason.report -webapps.mason.downloads ; +webapps.mason.downloads webapps.mason.status-update ; IN: webapps.mason TUPLE: mason-app < dispatcher ; @@ -35,5 +35,7 @@ can-make-releases? define-capability "make releases" >>description { can-make-releases? } >>capabilities + "make-release" add-responder - "make-release" add-responder ; + + "status-update" add-responder ; diff --git a/extra/webapps/mason/status-update/authors.txt b/extra/webapps/mason/status-update/authors.txt new file mode 100644 index 0000000000..1901f27a24 --- /dev/null +++ b/extra/webapps/mason/status-update/authors.txt @@ -0,0 +1 @@ +Slava Pestov diff --git a/extra/webapps/mason/status-update/status-update.factor b/extra/webapps/mason/status-update/status-update.factor new file mode 100644 index 0000000000..5156b1ef70 --- /dev/null +++ b/extra/webapps/mason/status-update/status-update.factor @@ -0,0 +1,74 @@ +! Copyright (C) 2010 Slava Pestov. +! See http://factorcode.org/license.txt for BSD license. +USING: accessors calendar combinators db.tuples furnace.actions +furnace.redirection html.forms http.server.responses io kernel +mason.config mason.server namespaces validators ; +IN: webapps.mason.status-update + +: find-builder ( -- builder ) + builder new + "host-name" value >>host-name + "target-os" value >>os + "target-cpu" value >>cpu + dup select-tuple [ ] [ dup insert-tuple ] ?if ; + +: git-id ( builder id -- ) >>current-git-id +starting+ >>status drop ; + +: make-vm ( builder -- ) +make-vm+ >>status drop ; + +: boot ( builder -- ) +boot+ >>status drop ; + +: test ( builder -- ) +test+ >>status drop ; + +: report ( builder status content -- ) + [ >>status ] [ >>last-report ] bi* + dup status>> +clean+ = [ + dup current-git-id>> >>clean-git-id + dup current-timestamp>> >>clean-timestamp + ] when + dup current-git-id>> >>last-git-id + dup current-timestamp>> >>last-timestamp + drop ; + +: release ( builder name -- ) + >>last-release + dup clean-git-id>> >>release-git-id + drop ; + +: update-builder ( builder -- ) + "message" value { + { "heartbeat" [ drop ] } + { "git-id" [ "arg" value git-id ] } + { "make-vm" [ make-vm ] } + { "boot" [ boot ] } + { "test" [ test ] } + { "report" [ "arg" value "report" value report ] } + { "release" [ "arg" value release ] } + } case ; + +: ( -- action ) + + [ + { + { "host-name" [ v-one-line ] } + { "target-cpu" [ v-one-line ] } + { "target-os" [ v-one-line ] } + { "message" [ v-one-line ] } + { "arg" [ [ v-one-line ] v-optional ] } + { "report" [ ] } + { "secret" [ v-one-line ] } + } validate-params + + "secret" value status-secret get = [ validation-failed ] unless + ] >>validate + + [ + [ + [ + find-builder + now >>current-timestamp + [ update-builder ] [ update-tuple ] bi + ] with-mason-db + "OK" "text/html" + ] if-secure + ] >>submit ; From dedc448f553f659dc9c1ca658e498faac6cb0b66 Mon Sep 17 00:00:00 2001 From: Slava Pestov Date: Sat, 17 Apr 2010 15:52:40 -0500 Subject: [PATCH 2/5] webapps.mason.package: fix --- extra/webapps/mason/package/package.factor | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/extra/webapps/mason/package/package.factor b/extra/webapps/mason/package/package.factor index 5c36a7f23a..504ba7093f 100644 --- a/extra/webapps/mason/package/package.factor +++ b/extra/webapps/mason/package/package.factor @@ -66,7 +66,7 @@ IN: webapps.mason.package [ current-status "status" set-value ] [ last-build-status "last-build" set-value ] [ clean-build-status "last-clean-build" set-value ] - [ heartbeat-timestamp>> "heartbeat-timestamp" set-value ] + [ current-timestamp>> "current-timestamp" set-value ] [ packages-link "binaries" set-value ] [ clean-image-link "clean-images" set-value ] [ report-link "last-report" set-value ] From f2de2222c7e235465b2b54e09b930e952195ea91 Mon Sep 17 00:00:00 2001 From: Joe Groff Date: Sat, 17 Apr 2010 15:58:51 -0700 Subject: [PATCH 3/5] cuda: AST representation for PTX code --- extra/cuda/ptx/ptx.factor | 753 ++++++++++++++++++++++++++++++++++++++ 1 file changed, 753 insertions(+) create mode 100644 extra/cuda/ptx/ptx.factor diff --git a/extra/cuda/ptx/ptx.factor b/extra/cuda/ptx/ptx.factor new file mode 100644 index 0000000000..12b132e117 --- /dev/null +++ b/extra/cuda/ptx/ptx.factor @@ -0,0 +1,753 @@ +! (c)2010 Joe Groff bsd license +USING: accessors arrays combinators io kernel math math.parser +roles sequences strings variants words ; +FROM: roles => TUPLE: ; +IN: cuda.ptx + +UNION: dim integer sequence ; +UNION: ?integer integer POSTPONE: f ; +UNION: ?string string POSTPONE: f ; + +VARIANT: ptx-type + .s8 .s16 .s32 .s64 + .u8 .u16 .u32 .u64 + .f16 .f32 .f64 + .b8 .b16 .b32 .b64 + .pred + .texref .samplerref .surfref + .v2: { { of ptx-type } } + .v4: { { of ptx-type } } + .struct: { { name string } } ; + +VARIANT: ptx-arch + sm_10 sm_11 sm_12 sm_13 sm_20 ; +UNION: ?ptx-arch ptx-arch POSTPONE: f ; + +VARIANT: ptx-texmode + .texmode_unified .texmode_independent ; +UNION: ?ptx-texmode ptx-texmode POSTPONE: f ; + +VARIANT: ptx-storage-space + .reg + .sreg + .const: { { bank ?integer } } + .global + .local + .param + .shared + .tex ; +UNION: ?ptx-storage-space ptx-storage-space POSTPONE: f ; + +TUPLE: ptx-target + { arch ?ptx-arch } + { map_f64_to_f32? boolean } + { texmode ?ptx-texmode } ; + +TUPLE: ptx + { version string } + { target ptx-target } + body ; + +TUPLE: ptx-struct-definition + { name string } + members ; + +TUPLE: ptx-variable + { extern? boolean } + { visible? boolean } + { align ?integer } + { storage-space ptx-storage-space } + { type ptx-type } + { name string } + { parameter ?integer } + { dim dim } + { initializer ?string } ; + +TUPLE: ptx-predicate + { negated? boolean } + { variable string } ; +UNION: ?ptx-predicate ptx-predicate POSTPONE: f ; + +TUPLE: ptx-instruction + { label ?string } + { predicate ?ptx-predicate } ; + +TUPLE: ptx-entry + { name string } + params + directives + body ; + +TUPLE: ptx-func < ptx-entry + { return ptx-variable } ; + +TUPLE: ptx-directive ; + +TUPLE: .file < ptx-directive + { info string } ; +TUPLE: .loc < ptx-directive + { info string } ; +TUPLE: .maxnctapersm < ptx-directive + { ncta integer } ; +TUPLE: .minnctapersm < ptx-directive + { ncta integer } ; +TUPLE: .maxnreg < ptx-directive + { n integer } ; +TUPLE: .maxntid < ptx-directive + { dim dim } ; +TUPLE: .pragma < ptx-directive + { pragma string } ; + +VARIANT: ptx-float-rounding-mode + .rn .rz .rm .rp .approx .full ; +VARIANT: ptx-int-rounding-mode + .rni .rzi .rmi .rpi ; +UNION: ?ptx-float-rounding-mode ptx-float-rounding-mode POSTPONE: f ; +UNION: ?ptx-int-rounding-mode ptx-int-rounding-mode POSTPONE: f ; + +UNION: ptx-rounding-mode + ptx-float-rounding-mode ptx-int-rounding-mode ; +UNION: ?ptx-rounding-mode ptx-rounding-mode POSTPONE: f ; + +TUPLE: ptx-typed-instruction < ptx-instruction + { type ptx-type } + { dest string } ; + +TUPLE: ptx-2op-instruction < ptx-typed-instruction + { a string } ; + +TUPLE: ptx-3op-instruction < ptx-typed-instruction + { a string } + { b string } ; + +TUPLE: ptx-4op-instruction < ptx-typed-instruction + { a string } + { b string } + { c string } ; + +TUPLE: ptx-5op-instruction < ptx-typed-instruction + { a string } + { b string } + { c string } + { d string } ; + +TUPLE: ptx-addsub-instruction < ptx-3op-instruction + { sat? boolean } + { cc? boolean } ; + +VARIANT: ptx-mul-mode + .wide ; +UNION: ?ptx-mul-mode ptx-mul-mode POSTPONE: f ; + +TUPLE: ptx-mul-instruction < ptx-3op-instruction + { mode ?ptx-mul-mode } ; + +TUPLE: ptx-mad-instruction < ptx-4op-instruction + { mode ?ptx-mul-mode } + { sat? boolean } ; + +VARIANT: ptx-prmt-mode + .f4e .b4e .rc8 .ecl .ecr .rc16 ; +UNION: ?ptx-prmt-mode ptx-prmt-mode POSTPONE: f ; + +ROLE: ptx-float-ftz + { ftz? boolean } ; +ROLE: ptx-float-env < ptx-float-ftz + { round ?ptx-float-rounding-mode } ; + +VARIANT: ptx-testp-op + .finite .infinite .number .notanumber .normal .subnormal ; + +VARIANT: ptx-cmp-op + .eq .ne + .lt .le .gt .ge + .ls .hs + .equ .neu + .ltu .leu .gtu .geu + .num .nan ; + +VARIANT: ptx-op + .and .or .xor .cas .exch .add .inc .dec .min .max + .popc ; + +SINGLETONS: .lo .hi ; +INSTANCE: .lo ptx-mul-mode +INSTANCE: .lo ptx-cmp-op +INSTANCE: .hi ptx-mul-mode +INSTANCE: .hi ptx-cmp-op + +TUPLE: ptx-set-instruction < ptx-3op-instruction + { cmp-op ptx-cmp-op } + { bool-op ptx-op } + { c ?string } + { ftz? boolean } ; + +VARIANT: ptx-cache-op + .ca .cg .cs .lu .cv + .wb .wt ; +UNION: ?ptx-cache-op ptx-cache-op POSTPONE: f ; + +TUPLE: ptx-ldst-instruction < ptx-2op-instruction + { volatile? boolean } + { storage-space ?ptx-storage-space } + { cache-op ?ptx-cache-op } ; + +VARIANT: ptx-cache-level + .L1 .L2 ; + +TUPLE: ptx-branch-instruction < ptx-instruction + { target string } + { uni? boolean } ; + +VARIANT: ptx-membar-level + .cta .gl .sys ; + +VARIANT: ptx-vote-mode + .all .any .uni .ballot ; + +TUPLE: ptx-instruction-not-supported-yet < ptx-instruction ; + +TUPLE: abs <{ ptx-2op-instruction ptx-float-ftz } ; +TUPLE: add <{ ptx-addsub-instruction ptx-float-env } ; +TUPLE: addc < ptx-addsub-instruction ; +TUPLE: and < ptx-3op-instruction ; +TUPLE: atom < ptx-3op-instruction + { storage-space ?ptx-storage-space } + { op ptx-op } + { c ?string } ; +TUPLE: bar.arrive < ptx-instruction + { a string } + { b string } ; +TUPLE: bar.red < ptx-2op-instruction + { op ptx-op } + { b ?string } + { c string } ; +TUPLE: bar.sync < ptx-instruction + { a string } + { b ?string } ; +TUPLE: bfe < ptx-4op-instruction ; +TUPLE: bfi < ptx-5op-instruction ; +TUPLE: bfind < ptx-2op-instruction + { shiftamt? boolean } ; +TUPLE: bra < ptx-branch-instruction ; +TUPLE: brev < ptx-2op-instruction ; +TUPLE: brkpt < ptx-instruction ; +TUPLE: call < ptx-branch-instruction + { return ?string } + params ; +TUPLE: clz < ptx-2op-instruction ; +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 } + { ftz? boolean } + { sat? boolean } + { dest-type ptx-type } ; +TUPLE: cvta < ptx-2op-instruction + { to? boolean } + { storage-space ?ptx-storage-space } ; +TUPLE: div <{ ptx-3op-instruction ptx-float-env } ; +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 } + { dest string } + { a string } ; +TUPLE: ld < ptx-ldst-instruction ; +TUPLE: ldu < ptx-ldst-instruction ; +TUPLE: lg2 <{ ptx-2op-instruction ptx-float-env } ; +TUPLE: mad <{ ptx-mad-instruction ptx-float-env } ; +TUPLE: mad24 < ptx-mad-instruction ; +TUPLE: max <{ ptx-3op-instruction ptx-float-ftz } ; +TUPLE: membar < ptx-instruction + { level ptx-membar-level } ; +TUPLE: min <{ ptx-3op-instruction ptx-float-ftz } ; +TUPLE: mov < ptx-2op-instruction ; +TUPLE: mul <{ ptx-mul-instruction ptx-float-env } ; +TUPLE: mul24 < ptx-mul-instruction ; +TUPLE: neg <{ ptx-2op-instruction ptx-float-ftz } ; +TUPLE: not < ptx-2op-instruction ; +TUPLE: or < ptx-3op-instruction ; +TUPLE: pmevent < ptx-instruction + { a string } ; +TUPLE: popc < ptx-2op-instruction ; +TUPLE: prefetch < ptx-instruction + { a string } + { storage-space ?ptx-storage-space } + { level ptx-cache-level } ; +TUPLE: prefetchu < ptx-instruction + { a string } + { level ptx-cache-level } ; +TUPLE: prmt < ptx-4op-instruction + { mode ?ptx-prmt-mode } ; +TUPLE: rcp <{ ptx-2op-instruction ptx-float-env } ; +TUPLE: red < ptx-2op-instruction + { storage-space ?ptx-storage-space } + { op ptx-op } ; +TUPLE: rem < ptx-3op-instruction ; +TUPLE: ret < ptx-instruction ; +TUPLE: rsqrt <{ ptx-2op-instruction ptx-float-env } ; +TUPLE: sad < ptx-4op-instruction ; +TUPLE: selp < ptx-4op-instruction ; +TUPLE: set < ptx-set-instruction + { dest-type ptx-type } ; +TUPLE: setp < ptx-set-instruction + { |dest ?string } ; +TUPLE: shl < ptx-3op-instruction ; +TUPLE: shr < ptx-3op-instruction ; +TUPLE: sin <{ ptx-2op-instruction ptx-float-env } ; +TUPLE: slct < ptx-4op-instruction + { dest-type ptx-type } + { ftz? boolean } ; +TUPLE: sqrt <{ ptx-2op-instruction ptx-float-env } ; +TUPLE: st < ptx-ldst-instruction ; +TUPLE: sub <{ ptx-addsub-instruction ptx-float-env } ; +TUPLE: subc < ptx-addsub-instruction ; +TUPLE: suld < ptx-instruction-not-supported-yet ; +TUPLE: sured < ptx-instruction-not-supported-yet ; +TUPLE: sust < ptx-instruction-not-supported-yet ; +TUPLE: suq < ptx-instruction-not-supported-yet ; +TUPLE: testp < ptx-2op-instruction + { op ptx-testp-op } ; +TUPLE: tex < ptx-instruction-not-supported-yet ; +TUPLE: txq < ptx-instruction-not-supported-yet ; +TUPLE: trap < ptx-instruction ; +TUPLE: vabsdiff < ptx-instruction-not-supported-yet ; +TUPLE: vadd < ptx-instruction-not-supported-yet ; +TUPLE: vmad < ptx-instruction-not-supported-yet ; +TUPLE: vmax < ptx-instruction-not-supported-yet ; +TUPLE: vmin < ptx-instruction-not-supported-yet ; +TUPLE: vset < ptx-instruction-not-supported-yet ; +TUPLE: vshl < ptx-instruction-not-supported-yet ; +TUPLE: vshr < ptx-instruction-not-supported-yet ; +TUPLE: vsub < ptx-instruction-not-supported-yet ; +TUPLE: vote < ptx-2op-instruction + { mode ptx-vote-mode } ; +TUPLE: xor < ptx-3op-instruction ; + +GENERIC: ptx-element-label ( elt -- label ) +M: object ptx-element-label drop f ; + +GENERIC: (write-ptx-element) ( elt -- ) + +: write-ptx-element ( elt -- ) + dup ptx-element-label [ write ":" write ] when* + "\t" write (write-ptx-element) + ";" print ; + +: write-ptx ( ptx -- ) + "\t.version " write dup version>> write ";" print + dup target>> write-ptx-element + body>> [ write-ptx-element ] each ; + +: write-ptx-symbol ( symbol/f -- ) + [ name>> write ] when* ; + +M: f (write-ptx-element) + drop ; + +M: word (write-ptx-element) + name>> write ; + +M: .const (write-ptx-element) + ".const" write + bank>> [ "[" write number>string write "]" write ] when* ; +M: .v2 (write-ptx-element) + ".v2" write of>> (write-ptx-element) ; +M: .v4 (write-ptx-element) + ".v4" write of>> (write-ptx-element) ; +M: .struct (write-ptx-element) + ".struct " write name>> write ; + +M: ptx-target (write-ptx-element) + ".target " write + [ arch>> [ name>> ] [ f ] if* ] + [ map_f64_to_f32?>> [ "map_f64_to_f32" ] [ f ] if ] + [ texmode>> [ name>> ] [ f ] if* ] tri + 3array sift ", " join write ; + +: write-ptx-dim ( dim -- ) + { + { [ dup zero? ] [ drop "[]" write ] } + { [ dup sequence? ] [ [ "[" write number>string write "]" write ] each ] } + [ "[" write number>string write "]" write ] + } cond ; + +M: ptx-variable (write-ptx-element) + dup extern?>> [ ".extern " write ] when + dup visible?>> [ ".visible " write ] when + dup align>> [ ".align " write number>string write " " write ] when* + dup storage-space>> (write-ptx-element) " " write + dup type>> (write-ptx-element) " " write + dup name>> write + dup parameter>> [ "<" write number>string write ">" write ] when* + dup dim>> [ write-ptx-dim ] when* + dup initializer>> [ " = " write write ] when* + drop ; + +: write-params ( params -- ) + "(" write unclip (write-ptx-element) + [ ", " write (write-ptx-element) ] each + ")" write ; + +: write-body ( params -- ) + "\t{" print + [ write-ptx-element ] each + "\t}" write ; + +: write-entry ( entry -- ) + dup name>> write " " write + dup params>> [ write-params ] when* nl + dup directives>> [ (write-ptx-element) ] each nl + dup body>> write-body + drop ; + +M: ptx-entry (write-ptx-element) + ".entry " write + write-entry ; + +M: ptx-func (write-ptx-element) + ".func " write + dup return>> [ "(" write (write-ptx-element) ") " write ] when* + write-entry ; + +M: .file (write-ptx-element) + ".file " write info>> write ; +M: .loc (write-ptx-element) + ".loc " write info>> write ; +M: .maxnctapersm (write-ptx-element) + ".maxnctapersm " write ncta>> number>string write ; +M: .minnctapersm (write-ptx-element) + ".minnctapersm " write ncta>> number>string write ; +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 ; +M: .pragma (write-ptx-element) + ".pragma \"" write pragma>> write "\"" write ; + +M: ptx-instruction ptx-element-label + label>> ; + +: write-insn ( insn name -- insn ) + over predicate>> + [ "@" write dup negated?>> [ "!" write ] when variable>> write ] when* + write ; + +: write-2op ( insn -- ) + dup type>> (write-ptx-element) " " write + dup dest>> write ", " write + dup a>> write + drop ; + +: write-3op ( insn -- ) + dup write-2op ", " write + dup b>> write + drop ; + +: write-4op ( insn -- ) + dup write-3op ", " write + dup c>> write + drop ; + +: write-5op ( insn -- ) + dup write-4op ", " write + dup d>> write + drop ; + +: write-ftz ( insn -- ) + ftz?>> [ ".ftz" write ] when ; + +: write-sat ( insn -- ) + sat?>> [ ".sat" write ] when ; + +: write-float-env ( insn -- ) + dup round>> (write-ptx-element) + write-ftz ; + +: write-int-addsub ( insn -- ) + dup write-sat + dup cc?>> [ ".cc" write ] when + write-3op ; + +: write-addsub ( insn -- ) + dup write-float-env + write-int-addsub ; + +: write-ldst ( insn -- ) + dup volatile?>> [ ".volatile" write ] when + dup storage-space>> (write-ptx-element) + dup cache-op>> (write-ptx-element) + write-2op ; + +: (write-mul) ( insn -- ) + dup mode>> (write-ptx-element) + drop ; + +: write-mul ( insn -- ) + dup write-float-env + dup (write-mul) + write-3op ; + +: write-mad ( insn -- ) + dup write-float-env + dup (write-mul) + dup write-sat + write-4op ; + +: write-uni ( insn -- ) + uni?>> [ ".uni" write ] when ; + +: write-set ( insn -- ) + dup cmp-op>> (write-ptx-element) + dup bool-op>> (write-ptx-element) + write-ftz ; + +M: abs (write-ptx-element) + "abs" write-insn + dup write-ftz + write-2op ; +M: add (write-ptx-element) + "add" write-insn + write-addsub ; +M: addc (write-ptx-element) + "addc" write-insn + write-int-addsub ; +M: and (write-ptx-element) + "and" write-insn + write-3op ; +M: atom (write-ptx-element) + "atom" write-insn + dup storage-space>> (write-ptx-element) + dup op>> (write-ptx-element) + dup write-3op + c>> [ ", " write write ] when* ; +M: bar.arrive (write-ptx-element) + "bar.arrive " write-insn + dup a>> write ", " write + dup b>> write + 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 ; +M: bar.sync (write-ptx-element) + "bar.arrive " write-insn + dup a>> write + dup b>> [ ", " write write ] when* + drop ; +M: bfe (write-ptx-element) + "bfe" write-insn + write-4op ; +M: bfi (write-ptx-element) + "bfi" write-insn + write-5op ; +M: bfind (write-ptx-element) + "bfind" write-insn + dup shiftamt?>> [ ".shiftamt" write ] when + write-2op ; +M: bra (write-ptx-element) + "bra" write-insn + 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 + dup return>> [ "(" write write "), " write ] when* + dup target>> write + dup params>> [ ", (" write ", " join write ")" write ] unless-empty + drop ; +M: clz (write-ptx-element) + "clz" write-insn + write-2op ; +M: cnot (write-ptx-element) + "cnot" write-insn + write-2op ; +M: copysign (write-ptx-element) + "copysign" write-insn + write-3op ; +M: cos (write-ptx-element) + "cos" write-insn + dup write-float-env + write-2op ; +M: cvt (write-ptx-element) + "cvt" write-insn + dup rounding-mode>> (write-ptx-element) + dup write-ftz + dup write-sat + dup dest-type>> (write-ptx-element) + write-2op ; +M: cvta (write-ptx-element) + "cvta" write-insn + dup to?>> [ ".to" write ] when + dup storage-space>> (write-ptx-element) + write-2op ; +M: div (write-ptx-element) + "div" write-insn + dup write-float-env + write-3op ; +M: ex2 (write-ptx-element) + "ex2" write-insn + dup write-float-env + write-2op ; +M: exit (write-ptx-element) + "exit" write-insn drop ; +M: fma (write-ptx-element) + "fma" write-insn + write-mad ; +M: isspacep (write-ptx-element) + "isspacep" write-insn + dup storage-space>> (write-ptx-element) + " " write + dup dest>> write ", " write a>> write ; +M: ld (write-ptx-element) + "ld" write-insn + write-ldst ; +M: ldu (write-ptx-element) + "ldu" write-insn + write-ldst ; +M: lg2 (write-ptx-element) + "lg2" write-insn + dup write-float-env + write-2op ; +M: mad (write-ptx-element) + "mad" write-insn + write-mad ; +M: mad24 (write-ptx-element) + "mad24" write-insn + dup (write-mul) + dup write-sat + write-4op ; +M: max (write-ptx-element) + "max" write-insn + dup write-ftz + write-3op ; +M: membar (write-ptx-element) + "membar" write-insn + dup level>> (write-ptx-element) + drop ; +M: min (write-ptx-element) + "min" write-insn + dup write-ftz + write-3op ; +M: mov (write-ptx-element) + "mov" write-insn + write-2op ; +M: mul (write-ptx-element) + "mul" write-insn + write-mul ; +M: mul24 (write-ptx-element) + "mul24" write-insn + dup (write-mul) + write-3op ; +M: neg (write-ptx-element) + "neg" write-insn + dup write-ftz + write-2op ; +M: not (write-ptx-element) + "not" write-insn + write-2op ; +M: or (write-ptx-element) + "or" write-insn + write-3op ; +M: pmevent (write-ptx-element) + "pmevent" write-insn " " write a>> write ; +M: popc (write-ptx-element) + "popc" write-insn + write-2op ; +M: prefetch (write-ptx-element) + "prefetch" write-insn + dup storage-space>> (write-ptx-element) + dup level>> (write-ptx-element) + " " write a>> write ; +M: prefetchu (write-ptx-element) + "prefetchu" write-insn + dup level>> (write-ptx-element) + " " write a>> write ; +M: prmt (write-ptx-element) + "prmt" write-insn + dup mode>> (write-ptx-element) + write-4op ; +M: rcp (write-ptx-element) + "rcp" write-insn + dup write-float-env + write-3op ; +M: red (write-ptx-element) + "red" write-insn + dup storage-space>> (write-ptx-element) + dup op>> (write-ptx-element) + write-2op ; +M: rem (write-ptx-element) + "rem" write-insn + write-3op ; +M: ret (write-ptx-element) + "ret" write-insn drop ; +M: rsqrt (write-ptx-element) + "rsqrt" write-insn + dup write-float-env + write-2op ; +M: sad (write-ptx-element) + "sad" write-insn + write-4op ; +M: selp (write-ptx-element) + "selp" write-insn + write-4op ; +M: set (write-ptx-element) + "set" write-insn + dup write-set + dup dest-type>> (write-ptx-element) + dup write-3op + c>> [ ", " write write ] when* ; +M: setp (write-ptx-element) + "setp" write-insn + dup write-set + dup write-3op + c>> [ ", " write write ] when* ; +M: shl (write-ptx-element) + "shl" write-insn + write-3op ; +M: shr (write-ptx-element) + "shr" write-insn + write-3op ; +M: sin (write-ptx-element) + "sin" write-insn + dup write-float-env + write-2op ; +M: slct (write-ptx-element) + "slct" write-insn + dup write-ftz + dup dest-type>> (write-ptx-element) + write-4op ; +M: sqrt (write-ptx-element) + "sqrt" write-insn + dup write-float-env + write-2op ; +M: st (write-ptx-element) + "st" write-insn + write-ldst ; +M: sub (write-ptx-element) + "sub" write-insn + write-addsub ; +M: subc (write-ptx-element) + "subc" write-insn + write-int-addsub ; +M: testp (write-ptx-element) + "testp" write-insn + dup op>> (write-ptx-element) + write-2op ; +M: vote (write-ptx-element) + "vote" write-insn + dup mode>> (write-ptx-element) + write-2op ; +M: xor (write-ptx-element) + "or" write-insn + write-3op ; From 0f2f54a195c53d6852410ff10e33e68d1e32fa3d Mon Sep 17 00:00:00 2001 From: Joe Groff Date: Sat, 17 Apr 2010 16:10:53 -0700 Subject: [PATCH 4/5] cuda.ptx: rearrange maybe types so that tuple slots default to f properly, and adjust output of some insns --- extra/cuda/ptx/ptx.factor | 35 ++++++++++++++++++++--------------- 1 file changed, 20 insertions(+), 15 deletions(-) diff --git a/extra/cuda/ptx/ptx.factor b/extra/cuda/ptx/ptx.factor index 12b132e117..8d4925d55f 100644 --- a/extra/cuda/ptx/ptx.factor +++ b/extra/cuda/ptx/ptx.factor @@ -5,8 +5,8 @@ FROM: roles => TUPLE: ; IN: cuda.ptx UNION: dim integer sequence ; -UNION: ?integer integer POSTPONE: f ; -UNION: ?string string POSTPONE: f ; +UNION: ?integer POSTPONE: f integer ; +UNION: ?string POSTPONE: f string ; VARIANT: ptx-type .s8 .s16 .s32 .s64 @@ -21,11 +21,11 @@ VARIANT: ptx-type VARIANT: ptx-arch sm_10 sm_11 sm_12 sm_13 sm_20 ; -UNION: ?ptx-arch ptx-arch POSTPONE: f ; +UNION: ?ptx-arch POSTPONE: f ptx-arch ; VARIANT: ptx-texmode .texmode_unified .texmode_independent ; -UNION: ?ptx-texmode ptx-texmode POSTPONE: f ; +UNION: ?ptx-texmode POSTPONE: f ptx-texmode ; VARIANT: ptx-storage-space .reg @@ -36,7 +36,7 @@ VARIANT: ptx-storage-space .param .shared .tex ; -UNION: ?ptx-storage-space ptx-storage-space POSTPONE: f ; +UNION: ?ptx-storage-space POSTPONE: f ptx-storage-space ; TUPLE: ptx-target { arch ?ptx-arch } @@ -66,7 +66,7 @@ TUPLE: ptx-variable TUPLE: ptx-predicate { negated? boolean } { variable string } ; -UNION: ?ptx-predicate ptx-predicate POSTPONE: f ; +UNION: ?ptx-predicate POSTPONE: f ptx-predicate ; TUPLE: ptx-instruction { label ?string } @@ -102,12 +102,12 @@ VARIANT: ptx-float-rounding-mode .rn .rz .rm .rp .approx .full ; VARIANT: ptx-int-rounding-mode .rni .rzi .rmi .rpi ; -UNION: ?ptx-float-rounding-mode ptx-float-rounding-mode POSTPONE: f ; -UNION: ?ptx-int-rounding-mode ptx-int-rounding-mode POSTPONE: f ; +UNION: ?ptx-float-rounding-mode POSTPONE: f ptx-float-rounding-mode ; +UNION: ?ptx-int-rounding-mode POSTPONE: f ptx-int-rounding-mode ; UNION: ptx-rounding-mode ptx-float-rounding-mode ptx-int-rounding-mode ; -UNION: ?ptx-rounding-mode ptx-rounding-mode POSTPONE: f ; +UNION: ?ptx-rounding-mode POSTPONE: f ptx-rounding-mode ; TUPLE: ptx-typed-instruction < ptx-instruction { type ptx-type } @@ -137,7 +137,7 @@ TUPLE: ptx-addsub-instruction < ptx-3op-instruction VARIANT: ptx-mul-mode .wide ; -UNION: ?ptx-mul-mode ptx-mul-mode POSTPONE: f ; +UNION: ?ptx-mul-mode POSTPONE: f ptx-mul-mode ; TUPLE: ptx-mul-instruction < ptx-3op-instruction { mode ?ptx-mul-mode } ; @@ -148,7 +148,7 @@ TUPLE: ptx-mad-instruction < ptx-4op-instruction VARIANT: ptx-prmt-mode .f4e .b4e .rc8 .ecl .ecr .rc16 ; -UNION: ?ptx-prmt-mode ptx-prmt-mode POSTPONE: f ; +UNION: ?ptx-prmt-mode POSTPONE: f ptx-prmt-mode ; ROLE: ptx-float-ftz { ftz? boolean } ; @@ -169,6 +169,7 @@ VARIANT: ptx-cmp-op VARIANT: ptx-op .and .or .xor .cas .exch .add .inc .dec .min .max .popc ; +UNION: ?ptx-op POSTPONE: f ptx-op ; SINGLETONS: .lo .hi ; INSTANCE: .lo ptx-mul-mode @@ -178,14 +179,14 @@ INSTANCE: .hi ptx-cmp-op TUPLE: ptx-set-instruction < ptx-3op-instruction { cmp-op ptx-cmp-op } - { bool-op ptx-op } + { bool-op ?ptx-op } { c ?string } { ftz? boolean } ; VARIANT: ptx-cache-op .ca .cg .cs .lu .cv .wb .wt ; -UNION: ?ptx-cache-op ptx-cache-op POSTPONE: f ; +UNION: ?ptx-cache-op POSTPONE: f ptx-cache-op ; TUPLE: ptx-ldst-instruction < ptx-2op-instruction { volatile? boolean } @@ -434,7 +435,7 @@ M: ptx-instruction ptx-element-label : write-insn ( insn name -- insn ) over predicate>> - [ "@" write dup negated?>> [ "!" write ] when variable>> write ] when* + [ "@" write dup negated?>> [ "!" write ] when variable>> write " " write ] when* write ; : write-2op ( insn -- ) @@ -710,7 +711,11 @@ M: set (write-ptx-element) M: setp (write-ptx-element) "setp" write-insn dup write-set - dup write-3op + 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* ; M: shl (write-ptx-element) "shl" write-insn From 5c17e6ee98424db6e9f6a55b655a7e6589707472 Mon Sep 17 00:00:00 2001 From: Doug Coleman Date: Sun, 18 Apr 2010 18:33:18 -0500 Subject: [PATCH 5/5] CUDA-FUNCTION: works, splitting up CUDA into more vocabs --- extra/cuda/cuda.factor | 131 +++++++++++++----- extra/cuda/demos/hello-world/authors.txt | 1 + .../cuda/demos/hello-world/hello-world.factor | 30 ++++ extra/cuda/{ => demos/hello-world}/hello.cu | 0 extra/cuda/{ => demos/hello-world}/hello.ptx | 0 extra/cuda/demos/prefix-sum/authors.txt | 2 + .../cuda/{ => demos/prefix-sum}/prefix-sum.cu | 0 extra/cuda/demos/prefix-sum/prefix-sum.factor | 21 +++ .../{ => demos/prefix-sum}/prefix-sum.ptx | 0 extra/cuda/syntax/authors.txt | 1 + extra/cuda/syntax/syntax.factor | 15 ++ 11 files changed, 164 insertions(+), 37 deletions(-) create mode 100644 extra/cuda/demos/hello-world/authors.txt create mode 100644 extra/cuda/demos/hello-world/hello-world.factor rename extra/cuda/{ => demos/hello-world}/hello.cu (100%) rename extra/cuda/{ => demos/hello-world}/hello.ptx (100%) create mode 100644 extra/cuda/demos/prefix-sum/authors.txt rename extra/cuda/{ => demos/prefix-sum}/prefix-sum.cu (100%) create mode 100644 extra/cuda/demos/prefix-sum/prefix-sum.factor rename extra/cuda/{ => demos/prefix-sum}/prefix-sum.ptx (100%) create mode 100644 extra/cuda/syntax/authors.txt create mode 100644 extra/cuda/syntax/syntax.factor diff --git a/extra/cuda/cuda.factor b/extra/cuda/cuda.factor index 6b343fb1cc..d8b6f2e2ce 100644 --- a/extra/cuda/cuda.factor +++ b/extra/cuda/cuda.factor @@ -1,11 +1,13 @@ ! 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 +USING: accessors alien alien.data alien.parser alien.strings +alien.syntax 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 ; +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 @@ -15,13 +17,32 @@ SYMBOL: cuda-function SYMBOL: cuda-launcher SYMBOL: cuda-memory-hashtable +SYMBOL: cuda-libraries +cuda-libraries [ H{ } clone ] initialize + +SYMBOL: cuda-functions + +TUPLE: cuda-library name path ; + +: ( name path -- obj ) + \ cuda-library new + swap >>path + swap >>name ; + +: add-cuda-library ( name path -- ) + normalize-path + dup name>> cuda-libraries get set-at ; + +: cuda-library ( name -- cuda-library ) + cuda-libraries get at ; + 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 ; + a:int [ cuDriverGetVersion cuda-error ] keep a:*int ; : init-cuda ( -- ) 0 cuInit cuda-error ; @@ -29,12 +50,19 @@ ERROR: throw-cuda-error n ; TUPLE: launcher { device integer initial: 0 } { device-flags initial: 0 } -path block-shape shared-size grid ; +path ; + +TUPLE: function-launcher +dim-block +dim-grid +shared-size +stream ; : with-cuda-context ( flags device quot -- ) + H{ } clone cuda-functions set [ [ CUcontext ] 2dip - [ cuCtxCreate cuda-error ] 3keep 2drop *void* + [ cuCtxCreate cuda-error ] 3keep 2drop a:*void* ] dip [ '[ _ @ ] ] [ drop '[ _ cuCtxDestroy cuda-error ] ] 2bi @@ -44,7 +72,7 @@ path block-shape shared-size grid ; [ normalize-path [ CUmodule ] dip - [ cuModuleLoad cuda-error ] 2keep drop *void* + [ cuModuleLoad cuda-error ] 2keep drop a:*void* ] dip [ '[ _ @ ] ] [ drop '[ _ cuModuleUnload cuda-error ] ] 2bi @@ -74,10 +102,10 @@ path block-shape shared-size grid ; [ cuDeviceGetCount cuda-error ] keep *int ; + a:int [ cuDeviceGetCount cuda-error ] keep a:*int ; : n>cuda-device ( n -- device ) - [ CUdevice ] dip [ cuDeviceGet cuda-error ] 2keep drop *int ; + [ CUdevice ] dip [ cuDeviceGet cuda-error ] 2keep drop a:*int ; : enumerate-cuda-devices ( -- devices ) #cuda-devices iota [ n>cuda-device ] map ; @@ -98,27 +126,30 @@ PRIVATE> [ 2drop utf8 alien>string ] 3bi ; : cuda-device-capability ( n -- pair ) - [ int int ] dip + [ a:int a:int ] dip [ cuDeviceComputeCapability cuda-error ] - [ drop [ *int ] bi@ ] 3bi 2array ; + [ drop [ a:*int ] bi@ ] 3bi 2array ; : cuda-device-memory ( n -- bytes ) - [ uint ] dip + [ a:uint ] dip [ cuDeviceTotalMem cuda-error ] - [ drop *uint ] 2bi ; + [ drop a:*uint ] 2bi ; -: get-cuda-function* ( module string -- function ) +: get-function-ptr* ( module string -- function ) [ CUfunction ] 2dip - [ cuModuleGetFunction cuda-error ] 3keep 2drop *void* ; + [ cuModuleGetFunction cuda-error ] 3keep 2drop a:*void* ; -: get-cuda-function ( string -- function ) - [ cuda-module get ] dip get-cuda-function* ; +: get-function-ptr ( string -- function ) + [ cuda-module get ] dip get-function-ptr* ; : with-cuda-function ( string quot -- ) [ - get-cuda-function cuda-function set + get-function-ptr* cuda-function set ] dip call ; inline +: cached-cuda-function ( string -- alien ) + cuda-functions get [ get-function-ptr ] cache ; + : launch-function* ( function -- ) cuLaunch cuda-error ; : launch-function ( -- ) cuda-function get cuLaunch cuda-error ; @@ -157,7 +188,7 @@ M: cuda-memory byte-length length>> ; : cuda-malloc ( n -- ptr ) [ CUdeviceptr ] dip [ cuMemAlloc cuda-error ] 2keep - [ *int ] dip add-cuda-memory ; + [ a:*int ] dip add-cuda-memory ; : cuda-free* ( ptr -- ) cuMemFree cuda-error ; @@ -237,9 +268,9 @@ ERROR: bad-cuda-parameter parameter ; offset param-size ; : cuda-device-attribute ( attribute dev -- n ) - [ int ] 2dip + [ a:int ] 2dip [ cuDeviceGetAttribute cuda-error ] - [ 2drop *int ] 3bi ; + [ 2drop a:*int ] 3bi ; : function-block-shape* ( function x y z -- ) cuFuncSetBlockShape cuda-error ; @@ -289,20 +320,46 @@ ERROR: bad-cuda-parameter parameter ; "CUDA Version: " write cuda-version number>string print nl #cuda-devices iota [ nl ] [ cuda-device. ] interleave ; +: c-type>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 ; -: 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 +: run-function-launcher ( function-launcher function -- ) + swap + { + [ dim-block>> first3 function-block-shape* ] + [ shared-size>> function-shared-size* ] + [ + dim-grid>> [ + launch-function* + ] [ + first2 launch-function-grid* + ] if-empty + ] + } 2cleave ; - [ 1array set-parameters ] - [ drop launch ] - [ device>host utf8 alien>string . ] tri - ] with-cuda-function - ] with-cuda ; +: 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 string arguments -- ) + [ + '[ + _ get-function-ptr + [ nip _ cuda-arguments ] + [ run-function-launcher ] 2bi + ] + ] + [ nip \ function-launcher suffix a:void function-effect ] + 2bi 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..6a598dda44 --- /dev/null +++ b/extra/cuda/demos/hello-world/hello-world.factor @@ -0,0 +1,30 @@ +! Copyright (C) 2010 Doug Coleman. +! See http://factorcode.org/license.txt for BSD license. +USING: alien.c-types alien.strings cuda cuda.syntax destructors +io.encodings.utf8 kernel locals math prettyprint sequences ; +IN: cuda.hello-world + +CUDA-LIBRARY: hello vocab:cuda/hello.ptx + +CUDA-FUNCTION: helloWorld ( char* string-ptr ) ; + +:: cuda-hello-world ( -- ) + T{ launcher + { device 0 } + { path "vocab:cuda/hello.ptx" } + } [ + "Hello World!" [ - ] map-index malloc-device-string &dispose dup :> str + + T{ function-launcher + { dim-block { 6 1 1 } } + { dim-grid { 2 1 } } + { shared-size 0 } + } + helloWorld + + ! <<< { 6 1 1 } { 2 1 } 1 >>> helloWorld + + str device>host utf8 alien>string . + ] 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..2cd8eba166 --- /dev/null +++ b/extra/cuda/demos/prefix-sum/prefix-sum.factor @@ -0,0 +1,21 @@ +! 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 } + { path "vocab:cuda/demos/prefix-sum/prefix-sum.ptx" } + } [ + + + ! { 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/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..b8df30f61c --- /dev/null +++ b/extra/cuda/syntax/syntax.factor @@ -0,0 +1,15 @@ +! Copyright (C) 2010 Doug Coleman. +! See http://factorcode.org/license.txt for BSD license. +USING: alien.parser cuda kernel lexer parser ; +IN: cuda.syntax + +SYNTAX: CUDA-LIBRARY: scan scan add-cuda-library ; + +SYNTAX: CUDA-FUNCTION: + scan [ create-in ] [ ] 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 ;
Host name:
Last heartbeat:
Last heartbeat:
Current status:
Last build:
Last clean build: