extra/cuda/ptx/ptx.factor

   1 ! Copyright (C) 2010 Joe Groff.
   2 ! See http://factorcode.org/license.txt for BSD license.
   3 USING: accessors arrays combinators io io.streams.string kernel
   4 math math.parser roles sequences strings variants words ;
   5 IN: cuda.ptx
   6
   7 UNION: dim integer sequence ;
   8
   9 VARIANT: ptx-type
  10     .s8 .s16 .s32 .s64
  11     .u8 .u16 .u32 .u64
  12         .f16 .f32 .f64
  13     .b8 .b16 .b32 .b64
  14     .pred
  15     .texref .samplerref .surfref
  16     .v2: { { of ptx-type } }
  17     .v4: { { of ptx-type } }
  18     .struct: { { name string } } ;
  19
  20 VARIANT: ptx-arch
  21     sm_10 sm_11 sm_12 sm_13 sm_20 ;
  22
  23 VARIANT: ptx-texmode
  24     .texmode_unified .texmode_independent ;
  25
  26 VARIANT: ptx-storage-space
  27     .reg
  28     .sreg
  29     .const: { { bank maybe{ integer } } }
  30     .global
  31     .local
  32     .param
  33     .shared
  34     .tex ;
  35
  36 ROLE-TUPLE: ptx-target
  37     { arch maybe{ ptx-arch } }
  38     { map_f64_to_f32? boolean }
  39     { texmode maybe{ ptx-texmode } } ;
  40
  41 ROLE-TUPLE: ptx
  42     { version string }
  43     { target ptx-target }
  44     body ;
  45
  46 ROLE-TUPLE: ptx-struct-definition
  47     { name string }
  48     members ;
  49
  50 ROLE-TUPLE: ptx-variable
  51     { extern? boolean }
  52     { visible? boolean }
  53     { align maybe{ integer } }
  54     { storage-space ptx-storage-space }
  55     { type ptx-type }
  56     { name string }
  57     { parameter maybe{ integer } }
  58     { dim dim }
  59     { initializer maybe{ string } } ;
  60
  61 ROLE-TUPLE: ptx-negation
  62     { var string } ;
  63
  64 ROLE-TUPLE: ptx-vector
  65     elements ;
  66
  67 ROLE-TUPLE: ptx-element
  68     { var string }
  69     { index integer } ;
  70
  71 UNION: ptx-var
  72     string ptx-element ;
  73
  74 ROLE-TUPLE: ptx-indirect
  75     { base ptx-var }
  76     { offset integer } ;
  77
  78 UNION: ptx-operand
  79     integer float ptx-var ptx-negation ptx-vector ptx-indirect ;
  80
  81 ROLE-TUPLE: ptx-instruction
  82     { label maybe{ string } }
  83     { predicate maybe{ ptx-operand } } ;
  84
  85 ROLE-TUPLE: ptx-entry
  86     { name string }
  87     params
  88     directives
  89     body ;
  90
  91 ROLE-TUPLE: ptx-func < ptx-entry
  92     { return maybe{ ptx-variable } } ;
  93
  94 ROLE-TUPLE: ptx-directive ;
  95
  96 ROLE-TUPLE: .file         < ptx-directive
  97     { info string } ;
  98 ROLE-TUPLE: .loc          < ptx-directive
  99     { info string } ;
 100 ROLE-TUPLE: .maxnctapersm < ptx-directive
 101     { ncta integer } ;
 102 ROLE-TUPLE: .minnctapersm < ptx-directive
 103     { ncta integer } ;
 104 ROLE-TUPLE: .maxnreg      < ptx-directive
 105     { n integer } ;
 106 ROLE-TUPLE: .maxntid      < ptx-directive
 107     { dim dim } ;
 108 ROLE-TUPLE: .pragma       < ptx-directive
 109     { pragma string } ;
 110
 111 VARIANT: ptx-float-rounding-mode
 112     .rn .rz .rm .rp .approx .full ;
 113 VARIANT: ptx-int-rounding-mode
 114     .rni .rzi .rmi .rpi ;
 115
 116 UNION: ptx-rounding-mode
 117     ptx-float-rounding-mode ptx-int-rounding-mode ;
 118
 119 ROLE-TUPLE: ptx-typed-instruction < ptx-instruction
 120     { type ptx-type }
 121     { dest ptx-operand } ;
 122
 123 ROLE-TUPLE: ptx-2op-instruction < ptx-typed-instruction
 124     { a ptx-operand } ;
 125
 126 ROLE-TUPLE: ptx-3op-instruction < ptx-typed-instruction
 127     { a ptx-operand }
 128     { b ptx-operand } ;
 129
 130 ROLE-TUPLE: ptx-4op-instruction < ptx-typed-instruction
 131     { a ptx-operand }
 132     { b ptx-operand }
 133     { c ptx-operand } ;
 134
 135 ROLE-TUPLE: ptx-5op-instruction < ptx-typed-instruction
 136     { a ptx-operand }
 137     { b ptx-operand }
 138     { c ptx-operand }
 139     { d ptx-operand } ;
 140
 141 ROLE-TUPLE: ptx-addsub-instruction < ptx-3op-instruction
 142     { sat? boolean }
 143     { cc? boolean } ;
 144
 145 VARIANT: ptx-mul-mode
 146     .wide ;
 147
 148 ROLE-TUPLE: ptx-mul-instruction < ptx-3op-instruction
 149     { mode maybe{ ptx-mul-mode } } ;
 150
 151 ROLE-TUPLE: ptx-mad-instruction < ptx-4op-instruction
 152     { mode maybe{ ptx-mul-mode } }
 153     { sat? boolean } ;
 154
 155 VARIANT: ptx-prmt-mode
 156     .f4e .b4e .rc8 .ecl .ecr .rc16 ;
 157
 158 ROLE: ptx-float-ftz
 159     { ftz? boolean } ;
 160 ROLE: ptx-float-env < ptx-float-ftz
 161     { round maybe{ ptx-float-rounding-mode } } ;
 162
 163 VARIANT: ptx-testp-op
 164     .finite .infinite .number .notanumber .normal .subnormal ;
 165
 166 VARIANT: ptx-cmp-op
 167     .eq .ne
 168     .lt .le .gt .ge
 169     .ls .hs
 170     .equ .neu
 171     .ltu .leu .gtu .geu
 172     .num .nan ;
 173
 174 VARIANT: ptx-op
 175     .and .or .xor .cas .exch .add .inc .dec .min .max
 176     .popc ;
 177
 178 SINGLETONS: .lo .hi ;
 179 INSTANCE: .lo ptx-mul-mode
 180 INSTANCE: .lo ptx-cmp-op
 181 INSTANCE: .hi ptx-mul-mode
 182 INSTANCE: .hi ptx-cmp-op
 183
 184 ROLE-TUPLE: ptx-set-instruction < ptx-3op-instruction
 185     { cmp-op ptx-cmp-op }
 186     { bool-op maybe{ ptx-op } }
 187     { c maybe{ ptx-operand } }
 188     { ftz? boolean } ;
 189
 190 VARIANT: ptx-cache-op
 191     .ca .cg .cs .lu .cv
 192     .wb .wt ;
 193
 194 ROLE-TUPLE: ptx-ldst-instruction < ptx-2op-instruction
 195     { volatile? boolean }
 196     { storage-space maybe{ ptx-storage-space } }
 197     { cache-op maybe{ ptx-cache-op } } ;
 198
 199 VARIANT: ptx-cache-level
 200     .L1 .L2 ;
 201
 202 ROLE-TUPLE: ptx-branch-instruction < ptx-instruction
 203     { target string }
 204     { uni? boolean } ;
 205
 206 VARIANT: ptx-membar-level
 207     .cta .gl .sys ;
 208
 209 VARIANT: ptx-vote-mode
 210     .all .any .uni .ballot ;
 211
 212 ROLE-TUPLE: ptx-instruction-not-supported-yet < ptx-instruction ;
 213
 214 ROLE-TUPLE: abs       <{ ptx-2op-instruction ptx-float-ftz } ;
 215 ROLE-TUPLE: add       <{ ptx-addsub-instruction ptx-float-env } ;
 216 ROLE-TUPLE: addc      < ptx-addsub-instruction ;
 217 ROLE-TUPLE: and       < ptx-3op-instruction ;
 218 ROLE-TUPLE: atom      < ptx-3op-instruction
 219     { storage-space maybe{ ptx-storage-space } }
 220     { op ptx-op }
 221     { c maybe{ ptx-operand } } ;
 222 ROLE-TUPLE: bar.arrive < ptx-instruction
 223     { a ptx-operand }
 224     { b ptx-operand } ;
 225 ROLE-TUPLE: bar.red   < ptx-2op-instruction
 226     { op ptx-op }
 227     { b maybe{ ptx-operand } }
 228     { c ptx-operand } ;
 229 ROLE-TUPLE: bar.sync  < ptx-instruction
 230     { a ptx-operand }
 231     { b maybe{ ptx-operand } } ;
 232 ROLE-TUPLE: bfe       < ptx-4op-instruction ;
 233 ROLE-TUPLE: bfi       < ptx-5op-instruction ;
 234 ROLE-TUPLE: bfind     < ptx-2op-instruction
 235     { shiftamt? boolean } ;
 236 ROLE-TUPLE: bra       < ptx-branch-instruction ;
 237 ROLE-TUPLE: brev      < ptx-2op-instruction ;
 238 ROLE-TUPLE: brkpt     < ptx-instruction ;
 239 ROLE-TUPLE: call      < ptx-branch-instruction
 240     { return maybe{ ptx-operand } }
 241     params ;
 242 ROLE-TUPLE: clz       < ptx-2op-instruction ;
 243 ROLE-TUPLE: cnot      < ptx-2op-instruction ;
 244 ROLE-TUPLE: copysign  < ptx-3op-instruction ;
 245 ROLE-TUPLE: cos       <{ ptx-2op-instruction ptx-float-env } ;
 246 ROLE-TUPLE: cvt       < ptx-2op-instruction
 247     { round maybe{ ptx-rounding-mode } }
 248     { ftz? boolean }
 249     { sat? boolean }
 250     { dest-type ptx-type } ;
 251 ROLE-TUPLE: cvta      < ptx-2op-instruction
 252     { to? boolean }
 253     { storage-space maybe{ ptx-storage-space } } ;
 254 ROLE-TUPLE: div       <{ ptx-3op-instruction ptx-float-env } ;
 255 ROLE-TUPLE: ex2       <{ ptx-2op-instruction ptx-float-env } ;
 256 ROLE-TUPLE: exit      < ptx-instruction ;
 257 ROLE-TUPLE: fma       <{ ptx-mad-instruction ptx-float-env } ;
 258 ROLE-TUPLE: isspacep  < ptx-instruction
 259     { storage-space ptx-storage-space }
 260     { dest ptx-operand }
 261     { a ptx-operand } ;
 262 ROLE-TUPLE: ld        < ptx-ldst-instruction ;
 263 ROLE-TUPLE: ldu       < ptx-ldst-instruction ;
 264 ROLE-TUPLE: lg2       <{ ptx-2op-instruction ptx-float-env } ;
 265 ROLE-TUPLE: mad       <{ ptx-mad-instruction ptx-float-env } ;
 266 ROLE-TUPLE: mad24     < ptx-mad-instruction ;
 267 ROLE-TUPLE: max       <{ ptx-3op-instruction ptx-float-ftz } ;
 268 ROLE-TUPLE: membar    < ptx-instruction
 269     { level ptx-membar-level } ;
 270 ROLE-TUPLE: min       <{ ptx-3op-instruction ptx-float-ftz } ;
 271 ROLE-TUPLE: mov       < ptx-2op-instruction ;
 272 ROLE-TUPLE: mul       <{ ptx-mul-instruction ptx-float-env } ;
 273 ROLE-TUPLE: mul24     < ptx-mul-instruction ;
 274 ROLE-TUPLE: neg       <{ ptx-2op-instruction ptx-float-ftz } ;
 275 ROLE-TUPLE: not       < ptx-2op-instruction ;
 276 ROLE-TUPLE: or        < ptx-3op-instruction ;
 277 ROLE-TUPLE: pmevent   < ptx-instruction
 278     { a ptx-operand } ;
 279 ROLE-TUPLE: popc      < ptx-2op-instruction ;
 280 ROLE-TUPLE: prefetch  < ptx-instruction
 281     { a ptx-operand }
 282     { storage-space maybe{ ptx-storage-space } }
 283     { level ptx-cache-level } ;
 284 ROLE-TUPLE: prefetchu < ptx-instruction
 285     { a ptx-operand }
 286     { level ptx-cache-level } ;
 287 ROLE-TUPLE: prmt      < ptx-4op-instruction
 288     { mode maybe{ ptx-prmt-mode } } ;
 289 ROLE-TUPLE: rcp       <{ ptx-2op-instruction ptx-float-env } ;
 290 ROLE-TUPLE: red       < ptx-2op-instruction
 291     { storage-space maybe{ ptx-storage-space } }
 292     { op ptx-op } ;
 293 ROLE-TUPLE: rem       < ptx-3op-instruction ;
 294 ROLE-TUPLE: ret       < ptx-instruction ;
 295 ROLE-TUPLE: rsqrt     <{ ptx-2op-instruction ptx-float-env } ;
 296 ROLE-TUPLE: sad       < ptx-4op-instruction ;
 297 ROLE-TUPLE: selp      < ptx-4op-instruction ;
 298 ROLE-TUPLE: set       < ptx-set-instruction
 299     { dest-type ptx-type } ;
 300 ROLE-TUPLE: setp      < ptx-set-instruction
 301     { |dest maybe{ ptx-operand } } ;
 302 ROLE-TUPLE: shl       < ptx-3op-instruction ;
 303 ROLE-TUPLE: shr       < ptx-3op-instruction ;
 304 ROLE-TUPLE: sin       <{ ptx-2op-instruction ptx-float-env } ;
 305 ROLE-TUPLE: slct      < ptx-4op-instruction
 306     { dest-type ptx-type }
 307     { ftz? boolean } ;
 308 ROLE-TUPLE: sqrt      <{ ptx-2op-instruction ptx-float-env } ;
 309 ROLE-TUPLE: st        < ptx-ldst-instruction ;
 310 ROLE-TUPLE: sub       <{ ptx-addsub-instruction ptx-float-env } ;
 311 ROLE-TUPLE: subc      < ptx-addsub-instruction  ;
 312 ROLE-TUPLE: suld      < ptx-instruction-not-supported-yet ;
 313 ROLE-TUPLE: sured     < ptx-instruction-not-supported-yet ;
 314 ROLE-TUPLE: sust      < ptx-instruction-not-supported-yet ;
 315 ROLE-TUPLE: suq       < ptx-instruction-not-supported-yet ;
 316 ROLE-TUPLE: testp     < ptx-2op-instruction
 317     { op ptx-testp-op } ;
 318 ROLE-TUPLE: tex       < ptx-instruction-not-supported-yet ;
 319 ROLE-TUPLE: txq       < ptx-instruction-not-supported-yet ;
 320 ROLE-TUPLE: trap      < ptx-instruction ;
 321 ROLE-TUPLE: vabsdiff  < ptx-instruction-not-supported-yet ;
 322 ROLE-TUPLE: vadd      < ptx-instruction-not-supported-yet ;
 323 ROLE-TUPLE: vmad      < ptx-instruction-not-supported-yet ;
 324 ROLE-TUPLE: vmax      < ptx-instruction-not-supported-yet ;
 325 ROLE-TUPLE: vmin      < ptx-instruction-not-supported-yet ;
 326 ROLE-TUPLE: vset      < ptx-instruction-not-supported-yet ;
 327 ROLE-TUPLE: vshl      < ptx-instruction-not-supported-yet ;
 328 ROLE-TUPLE: vshr      < ptx-instruction-not-supported-yet ;
 329 ROLE-TUPLE: vsub      < ptx-instruction-not-supported-yet ;
 330 ROLE-TUPLE: vote      < ptx-2op-instruction
 331     { mode ptx-vote-mode } ;
 332 ROLE-TUPLE: xor       < ptx-3op-instruction ;
 333
 334 GENERIC: ptx-element-label ( elt -- label )
 335 M: object ptx-element-label  drop f ;
 336
 337 GENERIC: ptx-semicolon? ( elt -- ? )
 338 M: object ptx-semicolon? drop t ;
 339 M: ptx-target ptx-semicolon? drop f ;
 340 M: ptx-entry ptx-semicolon? drop f ;
 341 M: ptx-func ptx-semicolon? drop f ;
 342 M: .file ptx-semicolon? drop f ;
 343 M: .loc ptx-semicolon? drop f ;
 344
 345 GENERIC: write-ptx-operand ( operand -- )
 346
 347 M: string write-ptx-operand write ;
 348 M: integer write-ptx-operand number>string write ;
 349 M: float write-ptx-operand "0d" write double>bits >hex 16 CHAR: 0 pad-head write ;
 350 M: ptx-negation write-ptx-operand "!" write var>> write ;
 351 M: ptx-vector write-ptx-operand
 352     "{" write
 353     elements>> [ ", " write ] [ write-ptx-operand ] interleave
 354     "}" write ;
 355 M: ptx-element write-ptx-operand dup var>> write "[" write index>> number>string write "]" write ;
 356 M: ptx-indirect write-ptx-operand
 357     "[" write
 358     dup base>> write-ptx-operand
 359     offset>> {
 360         { [ dup zero? ] [ drop ] }
 361         { [ dup 0 < ] [ number>string write ] }
 362         [ "+" write number>string write ]
 363     } cond
 364     "]" write ;
 365
 366 GENERIC: (write-ptx-element) ( elt -- )
 367
 368 : write-ptx-element ( elt -- )
 369     dup ptx-element-label [ write ":" write ] when*
 370     "\t" write dup (write-ptx-element)
 371     ptx-semicolon? [ ";" print ] [ nl ] if ;
 372
 373 : write-ptx ( ptx -- )
 374     "\t.version " write dup version>> print
 375     dup target>> write-ptx-element
 376     body>> [ write-ptx-element ] each ;
 377
 378 : write-ptx-symbol ( symbol/f -- )
 379     [ name>> write ] when* ;
 380
 381 M: f (write-ptx-element)
 382     drop ;
 383
 384 M: word (write-ptx-element)
 385     name>> write ;
 386
 387 M: .const (write-ptx-element)
 388     ".const" write
 389     bank>> [ "[" write number>string write "]" write ] when* ;
 390 M: .v2 (write-ptx-element)
 391     ".v2" write of>> (write-ptx-element) ;
 392 M: .v4 (write-ptx-element)
 393     ".v4" write of>> (write-ptx-element) ;
 394 M: .struct (write-ptx-element)
 395     ".struct " write name>> write ;
 396
 397 M: ptx-target (write-ptx-element)
 398     ".target " write
 399     [ arch>> [ name>> ] [ f ] if* ]
 400     [ map_f64_to_f32?>> [ "map_f64_to_f32" ] [ f ] if ]
 401     [ texmode>> [ name>> ] [ f ] if* ] tri
 402     3array sift [ ", " write ] [ write ] interleave ;
 403
 404 : write-ptx-dim ( dim -- )
 405     {
 406         { [ dup zero? ] [ drop "[]" write ] }
 407         { [ dup sequence? ] [ [ "[" write number>string write "]" write ] each ] }
 408         [ "[" write number>string write "]" write ]
 409     } cond ;
 410
 411 M: ptx-variable (write-ptx-element)
 412     dup extern?>> [ ".extern " write ] when
 413     dup visible?>> [ ".visible " write ] when
 414     dup align>> [ ".align " write number>string write bl ] when*
 415     dup storage-space>> (write-ptx-element) bl
 416     dup type>> (write-ptx-element) bl
 417     dup name>> write
 418     dup parameter>> [ "<" write number>string write ">" write ] when*
 419     dup dim>> [ write-ptx-dim ] when*
 420     dup initializer>> [ " = " write write ] when*
 421     drop ;
 422
 423 : write-params ( params -- )
 424     "(" write unclip (write-ptx-element)
 425     [ ", " write (write-ptx-element) ] each
 426     ")" write ;
 427
 428 : write-body ( params -- )
 429     "\t{" print
 430     [ write-ptx-element ] each
 431     "\t}" write ;
 432
 433 : write-entry ( entry -- )
 434     dup name>> write
 435     dup params>> [  bl write-params ] when* nl
 436     dup directives>> [ (write-ptx-element) nl ] each
 437     dup body>> write-body
 438     drop ;
 439
 440 M: ptx-entry (write-ptx-element)
 441     ".entry " write
 442     write-entry ;
 443
 444 M: ptx-func (write-ptx-element)
 445     ".func " write
 446     dup return>> [ "(" write (write-ptx-element) ") " write ] when*
 447     write-entry ;
 448
 449 M: .file (write-ptx-element)
 450     ".file " write info>> write ;
 451 M: .loc (write-ptx-element)
 452     ".loc " write info>> write ;
 453 M: .maxnctapersm (write-ptx-element)
 454     ".maxnctapersm " write ncta>> number>string write ;
 455 M: .minnctapersm (write-ptx-element)
 456     ".minnctapersm " write ncta>> number>string write ;
 457 M: .maxnreg (write-ptx-element)
 458     ".maxnreg " write n>> number>string write ;
 459 M: .maxntid (write-ptx-element)
 460     ".maxntid " write
 461     dup sequence? [ [ ", " write ] [ number>string write ] interleave ] [ number>string write ] if ;
 462 M: .pragma (write-ptx-element)
 463     ".pragma \"" write pragma>> write "\"" write ;
 464
 465 M: ptx-instruction ptx-element-label
 466     label>> ;
 467
 468 : write-insn ( insn name -- insn )
 469     over predicate>>
 470     [ "@" write write-ptx-operand bl ] when*
 471     write ;
 472
 473 : write-2op ( insn -- )
 474     dup type>> (write-ptx-element) bl
 475     dup dest>> write-ptx-operand ", " write
 476     dup a>> write-ptx-operand
 477     drop ;
 478
 479 : write-3op ( insn -- )
 480     dup write-2op ", " write
 481     dup b>> write-ptx-operand
 482     drop ;
 483
 484 : write-4op ( insn -- )
 485     dup write-3op ", " write
 486     dup c>> write-ptx-operand
 487     drop ;
 488
 489 : write-5op ( insn -- )
 490     dup write-4op ", " write
 491     dup d>> write-ptx-operand
 492     drop ;
 493
 494 : write-ftz ( insn -- )
 495     ftz?>> [ ".ftz" write ] when ;
 496
 497 : write-sat ( insn -- )
 498     sat?>> [ ".sat" write ] when ;
 499
 500 : write-float-env ( insn -- )
 501     dup round>> (write-ptx-element)
 502     write-ftz ;
 503
 504 : write-int-addsub ( insn -- )
 505     dup write-sat
 506     dup cc?>>  [ ".cc"  write ] when
 507     write-3op ;
 508
 509 : write-addsub ( insn -- )
 510     dup write-float-env
 511     write-int-addsub ;
 512
 513 : write-ldst ( insn -- )
 514     dup volatile?>> [ ".volatile" write ] when
 515     dup storage-space>> (write-ptx-element)
 516     dup cache-op>> (write-ptx-element)
 517     write-2op ;
 518
 519 : (write-mul) ( insn -- )
 520     dup mode>> (write-ptx-element)
 521     drop ;
 522
 523 : write-mul ( insn -- )
 524     dup write-float-env
 525     dup (write-mul)
 526     write-3op ;
 527
 528 : write-mad ( insn -- )
 529     dup write-float-env
 530     dup (write-mul)
 531     dup write-sat
 532     write-4op ;
 533
 534 : write-uni ( insn -- )
 535     uni?>> [ ".uni" write ] when ;
 536
 537 : write-set ( insn -- )
 538     dup cmp-op>> (write-ptx-element)
 539     dup bool-op>> (write-ptx-element)
 540     write-ftz ;
 541
 542 M: abs (write-ptx-element)
 543     "abs" write-insn
 544     dup write-ftz
 545     write-2op ;
 546 M: add (write-ptx-element)
 547     "add" write-insn
 548     write-addsub ;
 549 M: addc (write-ptx-element)
 550     "addc" write-insn
 551     write-int-addsub ;
 552 M: and (write-ptx-element)
 553     "and" write-insn
 554     write-3op ;
 555 M: atom (write-ptx-element)
 556     "atom" write-insn
 557     dup storage-space>> (write-ptx-element)
 558     dup op>> (write-ptx-element)
 559     dup write-3op
 560     c>> [ ", " write write-ptx-operand ] when* ;
 561 M: bar.arrive (write-ptx-element)
 562     "bar.arrive " write-insn
 563     dup a>> write-ptx-operand ", " write
 564     dup b>> write-ptx-operand
 565     drop ;
 566 M: bar.red (write-ptx-element)
 567     "bar.red" write-insn
 568     dup op>> (write-ptx-element)
 569     dup write-2op
 570     dup b>> [ ", " write write-ptx-operand ] when*
 571     ", " write c>> write-ptx-operand ;
 572 M: bar.sync (write-ptx-element)
 573     "bar.sync " write-insn
 574     dup a>> write-ptx-operand
 575     dup b>> [ ", " write write-ptx-operand ] when*
 576     drop ;
 577 M: bfe (write-ptx-element)
 578     "bfe" write-insn
 579     write-4op ;
 580 M: bfi (write-ptx-element)
 581     "bfi" write-insn
 582     write-5op ;
 583 M: bfind (write-ptx-element)
 584     "bfind" write-insn
 585     dup shiftamt?>> [ ".shiftamt" write ] when
 586     write-2op ;
 587 M: bra (write-ptx-element)
 588     "bra" write-insn
 589     dup write-uni bl
 590     target>> write ;
 591 M: brev (write-ptx-element)
 592     "brev" write-insn
 593     write-2op ;
 594 M: brkpt (write-ptx-element)
 595     "brkpt" write-insn drop ;
 596 M: call (write-ptx-element)
 597     "call" write-insn
 598     dup write-uni bl
 599     dup return>> [ "(" write write-ptx-operand "), " write ] when*
 600     dup target>> write
 601     dup params>> [ ", (" write [ ", " write ] [ write-ptx-operand ] interleave ")" write ] unless-empty
 602     drop ;
 603 M: clz (write-ptx-element)
 604     "clz" write-insn
 605     write-2op ;
 606 M: cnot (write-ptx-element)
 607     "cnot" write-insn
 608     write-2op ;
 609 M: copysign (write-ptx-element)
 610     "copysign" write-insn
 611     write-3op ;
 612 M: cos (write-ptx-element)
 613     "cos" write-insn
 614     dup write-float-env
 615     write-2op ;
 616 M: cvt (write-ptx-element)
 617     "cvt" write-insn
 618     dup round>> (write-ptx-element)
 619     dup write-ftz
 620     dup write-sat
 621     dup dest-type>> (write-ptx-element)
 622     write-2op ;
 623 M: cvta (write-ptx-element)
 624     "cvta" write-insn
 625     dup to?>> [ ".to" write ] when
 626     dup storage-space>> (write-ptx-element)
 627     write-2op ;
 628 M: div (write-ptx-element)
 629     "div" write-insn
 630     dup write-float-env
 631     write-3op ;
 632 M: ex2 (write-ptx-element)
 633     "ex2" write-insn
 634     dup write-float-env
 635     write-2op ;
 636 M: exit (write-ptx-element)
 637     "exit" write-insn drop ;
 638 M: fma (write-ptx-element)
 639     "fma" write-insn
 640     write-mad ;
 641 M: isspacep (write-ptx-element)
 642     "isspacep" write-insn
 643     dup storage-space>> (write-ptx-element)
 644     bl
 645     dup dest>> write-ptx-operand ", " write a>> write-ptx-operand ;
 646 M: ld (write-ptx-element)
 647     "ld" write-insn
 648     write-ldst ;
 649 M: ldu (write-ptx-element)
 650     "ldu" write-insn
 651     write-ldst ;
 652 M: lg2 (write-ptx-element)
 653     "lg2" write-insn
 654     dup write-float-env
 655     write-2op ;
 656 M: mad (write-ptx-element)
 657     "mad" write-insn
 658     write-mad ;
 659 M: mad24 (write-ptx-element)
 660     "mad24" write-insn
 661     dup (write-mul)
 662     dup write-sat
 663     write-4op ;
 664 M: max (write-ptx-element)
 665     "max" write-insn
 666     dup write-ftz
 667     write-3op ;
 668 M: membar (write-ptx-element)
 669     "membar" write-insn
 670     dup level>> (write-ptx-element)
 671     drop ;
 672 M: min (write-ptx-element)
 673     "min" write-insn
 674     dup write-ftz
 675     write-3op ;
 676 M: mov (write-ptx-element)
 677     "mov" write-insn
 678     write-2op ;
 679 M: mul (write-ptx-element)
 680     "mul" write-insn
 681     write-mul ;
 682 M: mul24 (write-ptx-element)
 683     "mul24" write-insn
 684     dup (write-mul)
 685     write-3op ;
 686 M: neg (write-ptx-element)
 687     "neg" write-insn
 688     dup write-ftz
 689     write-2op ;
 690 M: not (write-ptx-element)
 691     "not" write-insn
 692     write-2op ;
 693 M: or (write-ptx-element)
 694     "or" write-insn
 695     write-3op ;
 696 M: pmevent (write-ptx-element)
 697     "pmevent" write-insn bl a>> write ;
 698 M: popc (write-ptx-element)
 699     "popc" write-insn
 700     write-2op ;
 701 M: prefetch (write-ptx-element)
 702     "prefetch" write-insn
 703     dup storage-space>> (write-ptx-element)
 704     dup level>> (write-ptx-element)
 705     bl a>> write-ptx-operand ;
 706 M: prefetchu (write-ptx-element)
 707     "prefetchu" write-insn
 708     dup level>> (write-ptx-element)
 709     bl a>> write-ptx-operand ;
 710 M: prmt (write-ptx-element)
 711     "prmt" write-insn
 712     dup type>> (write-ptx-element)
 713     dup mode>> (write-ptx-element) bl
 714     dup dest>> write-ptx-operand ", " write
 715     dup a>> write-ptx-operand ", " write
 716     dup b>> write-ptx-operand ", " write
 717     dup c>> write-ptx-operand
 718     drop ;
 719 M: rcp (write-ptx-element)
 720     "rcp" write-insn
 721     dup write-float-env
 722     write-2op ;
 723 M: red (write-ptx-element)
 724     "red" write-insn
 725     dup storage-space>> (write-ptx-element)
 726     dup op>> (write-ptx-element)
 727     write-2op ;
 728 M: rem (write-ptx-element)
 729     "rem" write-insn
 730     write-3op ;
 731 M: ret (write-ptx-element)
 732     "ret" write-insn drop ;
 733 M: rsqrt (write-ptx-element)
 734     "rsqrt" write-insn
 735     dup write-float-env
 736     write-2op ;
 737 M: sad (write-ptx-element)
 738     "sad" write-insn
 739     write-4op ;
 740 M: selp (write-ptx-element)
 741     "selp" write-insn
 742     write-4op ;
 743 M: set (write-ptx-element)
 744     "set" write-insn
 745     dup write-set
 746     dup dest-type>> (write-ptx-element)
 747     dup write-3op
 748     c>> [ ", " write write-ptx-operand ] when* ;
 749 M: setp (write-ptx-element)
 750     "setp" write-insn
 751     dup write-set
 752     dup type>> (write-ptx-element) bl
 753     dup dest>> write-ptx-operand
 754     dup |dest>> [ "|" write write-ptx-operand ] when* ", " write
 755     dup a>> write-ptx-operand ", " write
 756     dup b>> write-ptx-operand
 757     c>> [ ", " write write-ptx-operand ] when* ;
 758 M: shl (write-ptx-element)
 759     "shl" write-insn
 760     write-3op ;
 761 M: shr (write-ptx-element)
 762     "shr" write-insn
 763     write-3op ;
 764 M: sin (write-ptx-element)
 765     "sin" write-insn
 766     dup write-float-env
 767     write-2op ;
 768 M: slct (write-ptx-element)
 769     "slct" write-insn
 770     dup write-ftz
 771     dup dest-type>> (write-ptx-element)
 772     write-4op ;
 773 M: sqrt (write-ptx-element)
 774     "sqrt" write-insn
 775     dup write-float-env
 776     write-2op ;
 777 M: st (write-ptx-element)
 778     "st" write-insn
 779     write-ldst ;
 780 M: sub (write-ptx-element)
 781     "sub" write-insn
 782     write-addsub ;
 783 M: subc (write-ptx-element)
 784     "subc" write-insn
 785     write-int-addsub ;
 786 M: testp (write-ptx-element)
 787     "testp" write-insn
 788     dup op>> (write-ptx-element)
 789     write-2op ;
 790 M: trap (write-ptx-element)
 791     "trap" write-insn drop ;
 792 M: vote (write-ptx-element)
 793     "vote" write-insn
 794     dup mode>> (write-ptx-element)
 795     write-2op ;
 796 M: xor (write-ptx-element)
 797     "xor" write-insn
 798     write-3op ;
 799
 800 : ptx>string ( ptx -- string )
 801     [ write-ptx ] with-string-writer ;