(write-ptx-element) ( elt -- )

Vocabulary

cuda.ptx

Inputs

elt

an object

Outputs

None

Definition

IN: cuda.ptx

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

Methods

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 ;