extra/cuda/ptx/ptx.factor

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