--- /dev/null
+! (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 ;
projects / factor.git / commitdiff