IN: cuda.ptx

GENERIC: (write-ptx-element) ( elt -- )

USING: accessors cuda.ptx io kernel math.parser ;

M: .const (write-ptx-element)
".const" write bank>>
[ "[" write number>string write "]" write ] when* ;

USING: accessors cuda.ptx io ;

M: .file (write-ptx-element) ".file " write info>> write ;

USING: accessors cuda.ptx io ;

M: .loc (write-ptx-element) ".loc " write info>> write ;

USING: accessors cuda.ptx io math.parser ;

M: .maxnctapersm (write-ptx-element)
".maxnctapersm " write ncta>> number>string write ;

USING: accessors cuda.ptx io math.parser ;

M: .maxnreg (write-ptx-element)
".maxnreg " write n>> number>string write ;

USING: cuda.ptx io kernel math.parser sequences ;

M: .maxntid (write-ptx-element)
".maxntid " write dup sequence?
[ [ ", " write ] [ number>string write ] interleave ]
[ number>string write ] if ;

USING: accessors cuda.ptx io math.parser ;

M: .minnctapersm (write-ptx-element)
".minnctapersm " write ncta>> number>string write ;

USING: accessors cuda.ptx io ;

M: .pragma (write-ptx-element)
".pragma \"" write pragma>> write "\"" write ;

USING: accessors cuda.ptx io ;

M: .struct (write-ptx-element) ".struct " write name>> write ;

USING: accessors cuda.ptx io ;

M: .v2 (write-ptx-element)
".v2" write of>> (write-ptx-element) ;

USING: accessors cuda.ptx io ;

M: .v4 (write-ptx-element)
".v4" write of>> (write-ptx-element) ;

USING: cuda.ptx kernel ;

M: abs (write-ptx-element)
"abs" write-insn dup write-ftz write-2op ;

USING: cuda.ptx ;

M: add (write-ptx-element) "add" write-insn write-addsub ;

USING: cuda.ptx ;

M: addc (write-ptx-element)
"addc" write-insn write-int-addsub ;

USING: cuda.ptx ;

M: and (write-ptx-element) "and" write-insn write-3op ;

USING: accessors cuda.ptx io kernel ;

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-ptx-operand ] when* ;

USING: accessors cuda.ptx io kernel ;

M: bar.arrive (write-ptx-element)
"bar.arrive " write-insn dup a>> write-ptx-operand
", " write dup b>> write-ptx-operand drop ;

USING: accessors cuda.ptx io kernel ;

M: bar.red (write-ptx-element)
"bar.red" write-insn dup op>> (write-ptx-element)
dup write-2op dup b>>
[ ", " write write-ptx-operand ] when* ", " write
c>> write-ptx-operand ;

USING: accessors cuda.ptx io kernel ;

M: bar.sync (write-ptx-element)
"bar.sync " write-insn dup a>> write-ptx-operand dup b>>
[ ", " write write-ptx-operand ] when* drop ;

USING: cuda.ptx ;

M: bfe (write-ptx-element) "bfe" write-insn write-4op ;

USING: cuda.ptx ;

M: bfi (write-ptx-element) "bfi" write-insn write-5op ;

USING: accessors cuda.ptx io kernel ;

M: bfind (write-ptx-element)
"bfind" write-insn dup shiftamt?>>
[ ".shiftamt" write ] when write-2op ;

USING: accessors cuda.ptx io kernel ;

M: bra (write-ptx-element)
"bra" write-insn dup write-uni bl target>> write ;

USING: cuda.ptx ;

M: brev (write-ptx-element) "brev" write-insn write-2op ;

USING: cuda.ptx kernel ;

M: brkpt (write-ptx-element) "brkpt" write-insn drop ;

USING: accessors cuda.ptx io kernel sequences ;

M: call (write-ptx-element)
"call" write-insn dup write-uni bl dup return>>
[ "(" write write-ptx-operand "), " write ] when*
dup target>> write dup params>> [
", (" write
[ ", " write ] [ write-ptx-operand ] interleave
")" write
] unless-empty drop ;

USING: cuda.ptx ;

M: clz (write-ptx-element) "clz" write-insn write-2op ;

USING: cuda.ptx ;

M: cnot (write-ptx-element) "cnot" write-insn write-2op ;

USING: cuda.ptx ;

M: copysign (write-ptx-element)
"copysign" write-insn write-3op ;

USING: cuda.ptx kernel ;

M: cos (write-ptx-element)
"cos" write-insn dup write-float-env write-2op ;

USING: accessors cuda.ptx kernel ;

M: cvt (write-ptx-element)
"cvt" write-insn dup round>> (write-ptx-element)
dup write-ftz dup write-sat
dup dest-type>> (write-ptx-element) write-2op ;

USING: accessors cuda.ptx io kernel ;

M: cvta (write-ptx-element)
"cvta" write-insn dup to?>> [ ".to" write ] when
dup storage-space>> (write-ptx-element) write-2op ;

USING: cuda.ptx kernel ;

M: div (write-ptx-element)
"div" write-insn dup write-float-env write-3op ;

USING: cuda.ptx kernel ;

M: ex2 (write-ptx-element)
"ex2" write-insn dup write-float-env write-2op ;

USING: cuda.ptx kernel ;

M: exit (write-ptx-element) "exit" write-insn drop ;

USING: cuda.ptx kernel ;

M: f (write-ptx-element) drop ;

USING: cuda.ptx ;

M: fma (write-ptx-element) "fma" write-insn write-mad ;

USING: accessors cuda.ptx io kernel ;

M: isspacep (write-ptx-element)
"isspacep" write-insn
dup storage-space>> (write-ptx-element)
bl dup dest>> write-ptx-operand ", " write
a>> write-ptx-operand ;

USING: cuda.ptx ;

M: ld (write-ptx-element) "ld" write-insn write-ldst ;

USING: cuda.ptx ;

M: ldu (write-ptx-element) "ldu" write-insn write-ldst ;

USING: cuda.ptx kernel ;

M: lg2 (write-ptx-element)
"lg2" write-insn dup write-float-env write-2op ;

USING: cuda.ptx kernel ;

M: mad24 (write-ptx-element)
"mad24" write-insn dup (write-mul) dup write-sat write-4op
;

USING: cuda.ptx ;

M: mad (write-ptx-element) "mad" write-insn write-mad ;

USING: cuda.ptx kernel ;

M: max (write-ptx-element)
"max" write-insn dup write-ftz write-3op ;

USING: accessors cuda.ptx kernel ;

M: membar (write-ptx-element)
"membar" write-insn dup level>> (write-ptx-element) drop ;

USING: cuda.ptx kernel ;

M: min (write-ptx-element)
"min" write-insn dup write-ftz write-3op ;

USING: cuda.ptx ;

M: mov (write-ptx-element) "mov" write-insn write-2op ;

USING: cuda.ptx kernel ;

M: mul24 (write-ptx-element)
"mul24" write-insn dup (write-mul) write-3op ;

USING: cuda.ptx ;

M: mul (write-ptx-element) "mul" write-insn write-mul ;

USING: cuda.ptx kernel ;

M: neg (write-ptx-element)
"neg" write-insn dup write-ftz write-2op ;

USING: cuda.ptx ;

M: not (write-ptx-element) "not" write-insn write-2op ;

USING: cuda.ptx ;

M: or (write-ptx-element) "or" write-insn write-3op ;

USING: accessors cuda.ptx io ;

M: pmevent (write-ptx-element)
"pmevent" write-insn bl a>> write ;

USING: cuda.ptx ;

M: popc (write-ptx-element) "popc" write-insn write-2op ;

USING: accessors cuda.ptx io kernel ;

M: prefetch (write-ptx-element)
"prefetch" write-insn
dup storage-space>> (write-ptx-element)
dup level>> (write-ptx-element) bl a>> write-ptx-operand ;

USING: accessors cuda.ptx io kernel ;

M: prefetchu (write-ptx-element)
"prefetchu" write-insn dup level>> (write-ptx-element)
bl a>> write-ptx-operand ;

USING: accessors cuda.ptx io kernel ;

M: prmt (write-ptx-element)
"prmt" write-insn dup type>> (write-ptx-element)
dup mode>> (write-ptx-element)
bl dup dest>> write-ptx-operand ", " write
dup a>> write-ptx-operand ", " write
dup b>> write-ptx-operand ", " write
dup c>> write-ptx-operand drop ;

USING: cuda.ptx io ;

M: ptx-entry (write-ptx-element) ".entry " write write-entry ;

USING: accessors cuda.ptx io kernel ;

M: ptx-func (write-ptx-element)
".func " write dup return>>
[ "(" write (write-ptx-element) ") " write ] when*
write-entry ;

USING: accessors arrays cuda.ptx io kernel sequences ;

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
[ ", " write ] [ write ] interleave ;

USING: accessors cuda.ptx io kernel math.parser ;

M: ptx-variable (write-ptx-element)
dup extern?>> [ ".extern " write ] when dup visible?>>
[ ".visible " write ] when dup align>>
[ ".align " write number>string write bl ] when*
dup storage-space>> (write-ptx-element)
bl dup type>> (write-ptx-element) bl dup name>> write
dup parameter>>
[ "<" write number>string write ">" write ] when* dup dim>>
[ write-ptx-dim ] when* dup initializer>>
[ " = " write write ] when* drop ;

USING: cuda.ptx kernel ;

M: rcp (write-ptx-element)
"rcp" write-insn dup write-float-env write-2op ;

USING: accessors cuda.ptx kernel ;

M: red (write-ptx-element)
"red" write-insn dup storage-space>> (write-ptx-element)
dup op>> (write-ptx-element) write-2op ;

USING: cuda.ptx ;

M: rem (write-ptx-element) "rem" write-insn write-3op ;

USING: cuda.ptx kernel ;

M: ret (write-ptx-element) "ret" write-insn drop ;

USING: cuda.ptx kernel ;

M: rsqrt (write-ptx-element)
"rsqrt" write-insn dup write-float-env write-2op ;

USING: cuda.ptx ;

M: sad (write-ptx-element) "sad" write-insn write-4op ;

USING: cuda.ptx ;

M: selp (write-ptx-element) "selp" write-insn write-4op ;

USING: accessors cuda.ptx io kernel ;

M: set (write-ptx-element)
"set" write-insn dup write-set
dup dest-type>> (write-ptx-element) dup write-3op c>>
[ ", " write write-ptx-operand ] when* ;

USING: accessors cuda.ptx io kernel ;

M: setp (write-ptx-element)
"setp" write-insn dup write-set
dup type>> (write-ptx-element)
bl dup dest>> write-ptx-operand dup |dest>>
[ "|" write write-ptx-operand ] when* ", " write
dup a>> write-ptx-operand ", " write
dup b>> write-ptx-operand c>>
[ ", " write write-ptx-operand ] when* ;

USING: cuda.ptx ;

M: shl (write-ptx-element) "shl" write-insn write-3op ;

USING: cuda.ptx ;

M: shr (write-ptx-element) "shr" write-insn write-3op ;

USING: cuda.ptx kernel ;

M: sin (write-ptx-element)
"sin" write-insn dup write-float-env write-2op ;

USING: accessors cuda.ptx kernel ;

M: slct (write-ptx-element)
"slct" write-insn dup write-ftz
dup dest-type>> (write-ptx-element) write-4op ;

USING: cuda.ptx kernel ;

M: sqrt (write-ptx-element)
"sqrt" write-insn dup write-float-env write-2op ;

USING: cuda.ptx ;

M: st (write-ptx-element) "st" write-insn write-ldst ;

USING: cuda.ptx ;

M: sub (write-ptx-element) "sub" write-insn write-addsub ;

USING: cuda.ptx ;

M: subc (write-ptx-element)
"subc" write-insn write-int-addsub ;

USING: accessors cuda.ptx kernel ;

M: testp (write-ptx-element)
"testp" write-insn dup op>> (write-ptx-element) write-2op ;

USING: cuda.ptx kernel ;

M: trap (write-ptx-element) "trap" write-insn drop ;

USING: accessors cuda.ptx kernel ;

M: vote (write-ptx-element)
"vote" write-insn dup mode>> (write-ptx-element) write-2op
;

USING: accessors cuda.ptx io words ;

M: word (write-ptx-element) name>> write ;

USING: cuda.ptx ;

M: xor (write-ptx-element) "xor" write-insn write-3op ;