Merge branch 'master' of git://factorcode.org/git/factor into s3
commit
e7a09eb80f
|
@ -1,11 +1,13 @@
|
||||||
! Copyright (C) 2010 Doug Coleman.
|
! Copyright (C) 2010 Doug Coleman.
|
||||||
! See http://factorcode.org/license.txt for BSD license.
|
! See http://factorcode.org/license.txt for BSD license.
|
||||||
USING: accessors alien alien.c-types alien.data alien.parser
|
USING: accessors alien alien.data alien.parser alien.strings
|
||||||
alien.strings arrays assocs byte-arrays classes.struct
|
alien.syntax arrays assocs byte-arrays classes.struct
|
||||||
combinators continuations cuda.ffi destructors fry io
|
combinators continuations cuda.ffi destructors fry io
|
||||||
io.backend io.encodings.string io.encodings.utf8 kernel lexer
|
io.backend io.encodings.string io.encodings.utf8 kernel lexer
|
||||||
locals math math.parser namespaces opengl.gl.extensions
|
locals macros math math.parser namespaces nested-comments
|
||||||
prettyprint quotations sequences ;
|
opengl.gl.extensions parser prettyprint quotations sequences
|
||||||
|
words ;
|
||||||
|
QUALIFIED-WITH: alien.c-types a
|
||||||
IN: cuda
|
IN: cuda
|
||||||
|
|
||||||
SYMBOL: cuda-device
|
SYMBOL: cuda-device
|
||||||
|
@ -15,13 +17,32 @@ SYMBOL: cuda-function
|
||||||
SYMBOL: cuda-launcher
|
SYMBOL: cuda-launcher
|
||||||
SYMBOL: cuda-memory-hashtable
|
SYMBOL: cuda-memory-hashtable
|
||||||
|
|
||||||
|
SYMBOL: cuda-libraries
|
||||||
|
cuda-libraries [ H{ } clone ] initialize
|
||||||
|
|
||||||
|
SYMBOL: cuda-functions
|
||||||
|
|
||||||
|
TUPLE: cuda-library name path ;
|
||||||
|
|
||||||
|
: <cuda-library> ( name path -- obj )
|
||||||
|
\ cuda-library new
|
||||||
|
swap >>path
|
||||||
|
swap >>name ;
|
||||||
|
|
||||||
|
: add-cuda-library ( name path -- )
|
||||||
|
normalize-path <cuda-library>
|
||||||
|
dup name>> cuda-libraries get set-at ;
|
||||||
|
|
||||||
|
: cuda-library ( name -- cuda-library )
|
||||||
|
cuda-libraries get at ;
|
||||||
|
|
||||||
ERROR: throw-cuda-error n ;
|
ERROR: throw-cuda-error n ;
|
||||||
|
|
||||||
: cuda-error ( n -- )
|
: cuda-error ( n -- )
|
||||||
dup CUDA_SUCCESS = [ drop ] [ throw-cuda-error ] if ;
|
dup CUDA_SUCCESS = [ drop ] [ throw-cuda-error ] if ;
|
||||||
|
|
||||||
: cuda-version ( -- n )
|
: cuda-version ( -- n )
|
||||||
int <c-object> [ cuDriverGetVersion cuda-error ] keep *int ;
|
a:int <c-object> [ cuDriverGetVersion cuda-error ] keep a:*int ;
|
||||||
|
|
||||||
: init-cuda ( -- )
|
: init-cuda ( -- )
|
||||||
0 cuInit cuda-error ;
|
0 cuInit cuda-error ;
|
||||||
|
@ -29,12 +50,19 @@ ERROR: throw-cuda-error n ;
|
||||||
TUPLE: launcher
|
TUPLE: launcher
|
||||||
{ device integer initial: 0 }
|
{ device integer initial: 0 }
|
||||||
{ device-flags 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 -- )
|
: with-cuda-context ( flags device quot -- )
|
||||||
|
H{ } clone cuda-functions set
|
||||||
[
|
[
|
||||||
[ CUcontext <c-object> ] 2dip
|
[ CUcontext <c-object> ] 2dip
|
||||||
[ cuCtxCreate cuda-error ] 3keep 2drop *void*
|
[ cuCtxCreate cuda-error ] 3keep 2drop a:*void*
|
||||||
] dip
|
] dip
|
||||||
[ '[ _ @ ] ]
|
[ '[ _ @ ] ]
|
||||||
[ drop '[ _ cuCtxDestroy cuda-error ] ] 2bi
|
[ drop '[ _ cuCtxDestroy cuda-error ] ] 2bi
|
||||||
|
@ -44,7 +72,7 @@ path block-shape shared-size grid ;
|
||||||
[
|
[
|
||||||
normalize-path
|
normalize-path
|
||||||
[ CUmodule <c-object> ] dip
|
[ CUmodule <c-object> ] dip
|
||||||
[ cuModuleLoad cuda-error ] 2keep drop *void*
|
[ cuModuleLoad cuda-error ] 2keep drop a:*void*
|
||||||
] dip
|
] dip
|
||||||
[ '[ _ @ ] ]
|
[ '[ _ @ ] ]
|
||||||
[ drop '[ _ cuModuleUnload cuda-error ] ] 2bi
|
[ drop '[ _ cuModuleUnload cuda-error ] ] 2bi
|
||||||
|
@ -74,10 +102,10 @@ path block-shape shared-size grid ;
|
||||||
<PRIVATE
|
<PRIVATE
|
||||||
|
|
||||||
: #cuda-devices ( -- n )
|
: #cuda-devices ( -- n )
|
||||||
int <c-object> [ cuDeviceGetCount cuda-error ] keep *int ;
|
a:int <c-object> [ cuDeviceGetCount cuda-error ] keep a:*int ;
|
||||||
|
|
||||||
: n>cuda-device ( n -- device )
|
: n>cuda-device ( n -- device )
|
||||||
[ CUdevice <c-object> ] dip [ cuDeviceGet cuda-error ] 2keep drop *int ;
|
[ CUdevice <c-object> ] dip [ cuDeviceGet cuda-error ] 2keep drop a:*int ;
|
||||||
|
|
||||||
: enumerate-cuda-devices ( -- devices )
|
: enumerate-cuda-devices ( -- devices )
|
||||||
#cuda-devices iota [ n>cuda-device ] map ;
|
#cuda-devices iota [ n>cuda-device ] map ;
|
||||||
|
@ -98,27 +126,30 @@ PRIVATE>
|
||||||
[ 2drop utf8 alien>string ] 3bi ;
|
[ 2drop utf8 alien>string ] 3bi ;
|
||||||
|
|
||||||
: cuda-device-capability ( n -- pair )
|
: cuda-device-capability ( n -- pair )
|
||||||
[ int <c-object> int <c-object> ] dip
|
[ a:int <c-object> a:int <c-object> ] dip
|
||||||
[ cuDeviceComputeCapability cuda-error ]
|
[ cuDeviceComputeCapability cuda-error ]
|
||||||
[ drop [ *int ] bi@ ] 3bi 2array ;
|
[ drop [ a:*int ] bi@ ] 3bi 2array ;
|
||||||
|
|
||||||
: cuda-device-memory ( n -- bytes )
|
: cuda-device-memory ( n -- bytes )
|
||||||
[ uint <c-object> ] dip
|
[ a:uint <c-object> ] dip
|
||||||
[ cuDeviceTotalMem cuda-error ]
|
[ cuDeviceTotalMem cuda-error ]
|
||||||
[ drop *uint ] 2bi ;
|
[ drop a:*uint ] 2bi ;
|
||||||
|
|
||||||
: get-cuda-function* ( module string -- function )
|
: get-function-ptr* ( module string -- function )
|
||||||
[ CUfunction <c-object> ] 2dip
|
[ CUfunction <c-object> ] 2dip
|
||||||
[ cuModuleGetFunction cuda-error ] 3keep 2drop *void* ;
|
[ cuModuleGetFunction cuda-error ] 3keep 2drop a:*void* ;
|
||||||
|
|
||||||
: get-cuda-function ( string -- function )
|
: get-function-ptr ( string -- function )
|
||||||
[ cuda-module get ] dip get-cuda-function* ;
|
[ cuda-module get ] dip get-function-ptr* ;
|
||||||
|
|
||||||
: with-cuda-function ( string quot -- )
|
: with-cuda-function ( string quot -- )
|
||||||
[
|
[
|
||||||
get-cuda-function cuda-function set
|
get-function-ptr* cuda-function set
|
||||||
] dip call ; inline
|
] dip call ; inline
|
||||||
|
|
||||||
|
: cached-cuda-function ( string -- alien )
|
||||||
|
cuda-functions get [ get-function-ptr ] cache ;
|
||||||
|
|
||||||
: launch-function* ( function -- ) cuLaunch cuda-error ;
|
: launch-function* ( function -- ) cuLaunch cuda-error ;
|
||||||
|
|
||||||
: launch-function ( -- ) cuda-function get 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 )
|
: cuda-malloc ( n -- ptr )
|
||||||
[ CUdeviceptr <c-object> ] dip
|
[ CUdeviceptr <c-object> ] dip
|
||||||
[ cuMemAlloc cuda-error ] 2keep
|
[ cuMemAlloc cuda-error ] 2keep
|
||||||
[ *int ] dip <cuda-memory> add-cuda-memory ;
|
[ a:*int ] dip <cuda-memory> add-cuda-memory ;
|
||||||
|
|
||||||
: cuda-free* ( ptr -- )
|
: cuda-free* ( ptr -- )
|
||||||
cuMemFree cuda-error ;
|
cuMemFree cuda-error ;
|
||||||
|
@ -237,9 +268,9 @@ ERROR: bad-cuda-parameter parameter ;
|
||||||
offset param-size ;
|
offset param-size ;
|
||||||
|
|
||||||
: cuda-device-attribute ( attribute dev -- n )
|
: cuda-device-attribute ( attribute dev -- n )
|
||||||
[ int <c-object> ] 2dip
|
[ a:int <c-object> ] 2dip
|
||||||
[ cuDeviceGetAttribute cuda-error ]
|
[ cuDeviceGetAttribute cuda-error ]
|
||||||
[ 2drop *int ] 3bi ;
|
[ 2drop a:*int ] 3bi ;
|
||||||
|
|
||||||
: function-block-shape* ( function x y z -- )
|
: function-block-shape* ( function x y z -- )
|
||||||
cuFuncSetBlockShape cuda-error ;
|
cuFuncSetBlockShape cuda-error ;
|
||||||
|
@ -289,20 +320,46 @@ ERROR: bad-cuda-parameter parameter ;
|
||||||
"CUDA Version: " write cuda-version number>string print nl
|
"CUDA Version: " write cuda-version number>string print nl
|
||||||
#cuda-devices iota [ nl ] [ cuda-device. ] interleave ;
|
#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 ( -- )
|
: run-function-launcher ( function-launcher function -- )
|
||||||
T{ launcher
|
swap
|
||||||
{ path "vocab:cuda/hello.ptx" }
|
{
|
||||||
{ block-shape { 6 6 6 } }
|
[ dim-block>> first3 function-block-shape* ]
|
||||||
{ shared-size 2 }
|
[ shared-size>> function-shared-size* ]
|
||||||
{ grid { 2 6 } }
|
[
|
||||||
} [
|
dim-grid>> [
|
||||||
"helloWorld" [
|
launch-function*
|
||||||
"Hello World!" [ - ] map-index
|
] [
|
||||||
malloc-device-string &dispose
|
first2 launch-function-grid*
|
||||||
|
] if-empty
|
||||||
|
]
|
||||||
|
} 2cleave ;
|
||||||
|
|
||||||
[ 1array set-parameters ]
|
: cuda-argument-setter ( offset c-type -- offset' quot )
|
||||||
[ drop launch ]
|
c-type>cuda-setter
|
||||||
[ device>host utf8 alien>string . ] tri
|
[ over [ + ] dip ] dip
|
||||||
] with-cuda-function
|
'[ swap _ swap _ call ] ;
|
||||||
] with-cuda ;
|
|
||||||
|
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 ;
|
||||||
|
|
|
@ -0,0 +1 @@
|
||||||
|
Doug Coleman
|
|
@ -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
|
|
@ -0,0 +1,2 @@
|
||||||
|
Doug Coleman
|
||||||
|
Joe Groff
|
|
@ -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
|
|
@ -0,0 +1,758 @@
|
||||||
|
! (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 POSTPONE: f integer ;
|
||||||
|
UNION: ?string POSTPONE: f string ;
|
||||||
|
|
||||||
|
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 POSTPONE: f ptx-arch ;
|
||||||
|
|
||||||
|
VARIANT: ptx-texmode
|
||||||
|
.texmode_unified .texmode_independent ;
|
||||||
|
UNION: ?ptx-texmode POSTPONE: f ptx-texmode ;
|
||||||
|
|
||||||
|
VARIANT: ptx-storage-space
|
||||||
|
.reg
|
||||||
|
.sreg
|
||||||
|
.const: { { bank ?integer } }
|
||||||
|
.global
|
||||||
|
.local
|
||||||
|
.param
|
||||||
|
.shared
|
||||||
|
.tex ;
|
||||||
|
UNION: ?ptx-storage-space POSTPONE: f ptx-storage-space ;
|
||||||
|
|
||||||
|
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 POSTPONE: f ptx-predicate ;
|
||||||
|
|
||||||
|
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 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 POSTPONE: f ptx-rounding-mode ;
|
||||||
|
|
||||||
|
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 POSTPONE: f ptx-mul-mode ;
|
||||||
|
|
||||||
|
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 POSTPONE: f ptx-prmt-mode ;
|
||||||
|
|
||||||
|
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 ;
|
||||||
|
UNION: ?ptx-op POSTPONE: f ptx-op ;
|
||||||
|
|
||||||
|
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 POSTPONE: f ptx-cache-op ;
|
||||||
|
|
||||||
|
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 " " 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 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
|
||||||
|
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 ;
|
|
@ -0,0 +1 @@
|
||||||
|
Doug Coleman
|
|
@ -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 ;
|
|
@ -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.
|
! See http://factorcode.org/license.txt for BSD license.
|
||||||
USING: system io.files io.pathnames namespaces kernel accessors
|
USING: system io.files io.pathnames namespaces kernel accessors
|
||||||
assocs ;
|
assocs ;
|
||||||
|
@ -39,11 +39,11 @@ target-os get-global [
|
||||||
! Keep test-log around?
|
! Keep test-log around?
|
||||||
SYMBOL: builder-debug
|
SYMBOL: builder-debug
|
||||||
|
|
||||||
! Host to send status notifications to.
|
! URL for status notifications.
|
||||||
SYMBOL: status-host
|
SYMBOL: status-url
|
||||||
|
|
||||||
! Username to log in.
|
! Password for status notifications.
|
||||||
SYMBOL: status-username
|
SYMBOL: status-secret
|
||||||
|
|
||||||
SYMBOL: upload-help?
|
SYMBOL: upload-help?
|
||||||
|
|
||||||
|
|
|
@ -1,57 +1,50 @@
|
||||||
! Copyright (C) 2009, 2010 Slava Pestov.
|
! Copyright (C) 2009, 2010 Slava Pestov.
|
||||||
! See http://factorcode.org/license.txt for BSD license.
|
! See http://factorcode.org/license.txt for BSD license.
|
||||||
USING: arrays accessors io io.sockets io.encodings.utf8 io.files
|
USING: accessors fry http.client io io.encodings.utf8 io.files
|
||||||
io.launcher kernel make mason.config mason.common mason.email
|
kernel mason.common mason.config mason.email mason.twitter
|
||||||
mason.twitter namespaces sequences prettyprint fry ;
|
namespaces prettyprint sequences ;
|
||||||
IN: mason.notify
|
IN: mason.notify
|
||||||
|
|
||||||
: status-notify ( input-file args -- )
|
: status-notify ( report arg message -- )
|
||||||
status-host get [
|
[
|
||||||
[
|
short-host-name "host-name" set
|
||||||
"ssh" , status-host get , "-l" , status-username get ,
|
target-cpu get "target-cpu" set
|
||||||
"./mason-notify" ,
|
target-os get "target-os" set
|
||||||
short-host-name ,
|
status-secret get "secret" set
|
||||||
target-cpu get ,
|
"message" set
|
||||||
target-os get ,
|
"arg" set
|
||||||
] { } make prepend
|
"report" set
|
||||||
[ 5 ] 2dip '[
|
] H{ } make-assoc
|
||||||
<process>
|
[ 5 ] dip '[ _ status-url get http-post 2drop ] retry ;
|
||||||
_ >>stdin
|
|
||||||
_ >>command
|
|
||||||
short-running-process
|
|
||||||
] retry
|
|
||||||
] [ 2drop ] if ;
|
|
||||||
|
|
||||||
: notify-heartbeat ( -- )
|
: notify-heartbeat ( -- )
|
||||||
f { "heartbeat" } status-notify ;
|
f f "heartbeat" status-notify ;
|
||||||
|
|
||||||
: notify-begin-build ( git-id -- )
|
: notify-begin-build ( git-id -- )
|
||||||
[ "Starting build of GIT ID " write print flush ]
|
[ "Starting build of GIT ID " write print flush ]
|
||||||
[ f swap "git-id" swap 2array status-notify ]
|
[ f swap "git-id" status-notify ]
|
||||||
bi ;
|
bi ;
|
||||||
|
|
||||||
: notify-make-vm ( -- )
|
: notify-make-vm ( -- )
|
||||||
"Compiling VM" print flush
|
"Compiling VM" print flush
|
||||||
f { "make-vm" } status-notify ;
|
f f "make-vm" status-notify ;
|
||||||
|
|
||||||
: notify-boot ( -- )
|
: notify-boot ( -- )
|
||||||
"Bootstrapping" print flush
|
"Bootstrapping" print flush
|
||||||
f { "boot" } status-notify ;
|
f f "boot" status-notify ;
|
||||||
|
|
||||||
: notify-test ( -- )
|
: notify-test ( -- )
|
||||||
"Running tests" print flush
|
"Running tests" print flush
|
||||||
f { "test" } status-notify ;
|
f f "test" status-notify ;
|
||||||
|
|
||||||
: notify-report ( status -- )
|
: notify-report ( status -- )
|
||||||
[ "Build finished with status: " write . flush ]
|
[ "Build finished with status: " write . flush ]
|
||||||
[
|
[
|
||||||
[ "report" ] dip
|
[ "report" utf8 file-contents ] dip
|
||||||
[ [ utf8 file-contents ] dip email-report ]
|
[ name>> "report" status-notify ] [ email-report ] 2bi
|
||||||
[ "report" swap name>> 2array status-notify ]
|
|
||||||
2bi
|
|
||||||
] bi ;
|
] bi ;
|
||||||
|
|
||||||
: notify-release ( archive-name -- )
|
: notify-release ( archive-name -- )
|
||||||
[ "Uploaded " prepend [ print flush ] [ mason-tweet ] bi ]
|
[ "Uploaded " prepend [ print flush ] [ mason-tweet ] bi ]
|
||||||
[ f swap "release" swap 2array status-notify ]
|
[ f swap "release" status-notify ]
|
||||||
bi ;
|
bi ;
|
||||||
|
|
|
@ -1 +0,0 @@
|
||||||
Slava Pestov
|
|
|
@ -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<sequence ;
|
|
||||||
|
|
||||||
: find-builder ( -- builder )
|
|
||||||
builder new
|
|
||||||
host-name get >>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
|
|
|
@ -17,8 +17,7 @@ clean-git-id clean-timestamp
|
||||||
last-release release-git-id
|
last-release release-git-id
|
||||||
last-git-id last-timestamp last-report
|
last-git-id last-timestamp last-report
|
||||||
current-git-id current-timestamp
|
current-git-id current-timestamp
|
||||||
status
|
status ;
|
||||||
heartbeat-timestamp ;
|
|
||||||
|
|
||||||
builder "BUILDERS" {
|
builder "BUILDERS" {
|
||||||
{ "host-name" "HOST_NAME" TEXT +user-assigned-id+ }
|
{ "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
|
! Can't name it CURRENT_TIMESTAMP because of bug in db library
|
||||||
{ "current-timestamp" "CURR_TIMESTAMP" TIMESTAMP }
|
{ "current-timestamp" "CURR_TIMESTAMP" TIMESTAMP }
|
||||||
{ "status" "STATUS" TEXT }
|
{ "status" "STATUS" TEXT }
|
||||||
|
|
||||||
{ "heartbeat-timestamp" "HEARTBEAT_TIMESTAMP" TIMESTAMP }
|
|
||||||
} define-persistent
|
} define-persistent
|
||||||
|
|
||||||
: mason-db ( -- db ) "resource:mason.db" <sqlite-db> ;
|
: mason-db ( -- db ) "resource:mason.db" <sqlite-db> ;
|
||||||
|
|
|
@ -28,7 +28,7 @@
|
||||||
|
|
||||||
<table border="1">
|
<table border="1">
|
||||||
<tr><td>Host name:</td><td><t:xml t:name="host-name" /></td></tr>
|
<tr><td>Host name:</td><td><t:xml t:name="host-name" /></td></tr>
|
||||||
<tr><td>Last heartbeat:</td><td><t:label t:name="heartbeat-timestamp" /></td></tr>
|
<tr><td>Last heartbeat:</td><td><t:label t:name="current-timestamp" /></td></tr>
|
||||||
<tr><td>Current status:</td><td><t:xml t:name="status" /></td></tr>
|
<tr><td>Current status:</td><td><t:xml t:name="status" /></td></tr>
|
||||||
<tr><td>Last build:</td><td><t:xml t:name="last-build" /></td></tr>
|
<tr><td>Last build:</td><td><t:xml t:name="last-build" /></td></tr>
|
||||||
<tr><td>Last clean build:</td><td><t:xml t:name="last-clean-build" /></td></tr>
|
<tr><td>Last clean build:</td><td><t:xml t:name="last-clean-build" /></td></tr>
|
||||||
|
|
|
@ -4,7 +4,7 @@ USING: accessors furnace.auth furnace.db
|
||||||
http.server.dispatchers mason.server webapps.mason.grids
|
http.server.dispatchers mason.server webapps.mason.grids
|
||||||
webapps.mason.make-release webapps.mason.package
|
webapps.mason.make-release webapps.mason.package
|
||||||
webapps.mason.release webapps.mason.report
|
webapps.mason.release webapps.mason.report
|
||||||
webapps.mason.downloads ;
|
webapps.mason.downloads webapps.mason.status-update ;
|
||||||
IN: webapps.mason
|
IN: webapps.mason
|
||||||
|
|
||||||
TUPLE: mason-app < dispatcher ;
|
TUPLE: mason-app < dispatcher ;
|
||||||
|
@ -35,5 +35,7 @@ can-make-releases? define-capability
|
||||||
<protected>
|
<protected>
|
||||||
"make releases" >>description
|
"make releases" >>description
|
||||||
{ can-make-releases? } >>capabilities
|
{ can-make-releases? } >>capabilities
|
||||||
|
"make-release" add-responder
|
||||||
|
|
||||||
"make-release" add-responder ;
|
<status-update-action>
|
||||||
|
"status-update" add-responder ;
|
||||||
|
|
|
@ -66,7 +66,7 @@ IN: webapps.mason.package
|
||||||
[ current-status "status" set-value ]
|
[ current-status "status" set-value ]
|
||||||
[ last-build-status "last-build" set-value ]
|
[ last-build-status "last-build" set-value ]
|
||||||
[ clean-build-status "last-clean-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 ]
|
[ packages-link "binaries" set-value ]
|
||||||
[ clean-image-link "clean-images" set-value ]
|
[ clean-image-link "clean-images" set-value ]
|
||||||
[ report-link "last-report" set-value ]
|
[ report-link "last-report" set-value ]
|
||||||
|
|
|
@ -0,0 +1 @@
|
||||||
|
Slava Pestov
|
|
@ -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 ;
|
||||||
|
|
||||||
|
: <status-update-action> ( -- action )
|
||||||
|
<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" <content>
|
||||||
|
] if-secure
|
||||||
|
] >>submit ;
|
Loading…
Reference in New Issue