extra/cuda/ptx/ptx-tests.factor

   1 USING: cuda.ptx io.streams.string tools.test ;
   2 IN: cuda.ptx.tests
   3
   4 [ """   .version 2.0
   5         .target sm_20
   6 """ ] [
   7     T{ ptx
   8         { version "2.0" }
   9         { target T{ ptx-target { arch sm_20 } } }
  10     } ptx>string
  11 ] unit-test
  12
  13 [ """   .version 2.0
  14         .target sm_20, .texmode_independent
  15 """ ] [
  16     T{ ptx
  17         { version "2.0" }
  18         { target T{ ptx-target { arch sm_20 } { texmode .texmode_independent } } }
  19     } ptx>string
  20 ] unit-test
  21
  22 [ """   .version 2.0
  23         .target sm_11, map_f64_to_f32
  24 """ ] [
  25     T{ ptx
  26         { version "2.0" }
  27         { target T{ ptx-target
  28             { arch sm_11 }
  29             { map_f64_to_f32? t }
  30         } }
  31     } ptx>string
  32 ] unit-test
  33
  34 [ """   .version 2.0
  35         .target sm_11, map_f64_to_f32, .texmode_independent
  36 """ ] [
  37     T{ ptx
  38         { version "2.0" }
  39         { target T{ ptx-target
  40             { arch sm_11 }
  41             { map_f64_to_f32? t }
  42             { texmode .texmode_independent }
  43         } }
  44     } ptx>string
  45 ] unit-test
  46
  47 [ """   .version 2.0
  48         .target sm_20
  49         .global .f32 foo[9000];
  50         .extern .align 16 .shared .v4.f32 bar[];
  51         .func (.reg .f32 sum) zap (.reg .f32 a, .reg .f32 b)
  52         {
  53         add.rn.f32 sum, a, b;
  54         ret;
  55         }
  56         .func frob (.align 8 .param .u64 in, .align 8 .param .u64 out, .align 8 .param .u64 len)
  57         {
  58         ret;
  59         }
  60         .func twib
  61         {
  62         ret;
  63         }
  64 """ ] [
  65     T{ ptx
  66         { version "2.0" }
  67         { target T{ ptx-target { arch sm_20 } } }
  68         { body {
  69             T{ ptx-variable
  70                 { storage-space .global }
  71                 { type .f32 }
  72                 { name "foo" }
  73                 { dim 9000 }
  74             }
  75             T{ ptx-variable
  76                 { extern? t }
  77                 { align 16 }
  78                 { storage-space .shared }
  79                 { type T{ .v4 f .f32 } }
  80                 { name "bar" }
  81                 { dim 0 }
  82             }
  83             T{ ptx-func
  84                 { return T{ ptx-variable { storage-space .reg } { type .f32 } { name "sum" } } }
  85                 { name "zap" }
  86                 { params {
  87                     T{ ptx-variable { storage-space .reg } { type .f32 } { name "a" } }
  88                     T{ ptx-variable { storage-space .reg } { type .f32 } { name "b" } }
  89                 } }
  90                 { body {
  91                     T{ add { round .rn } { type .f32 } { dest "sum" } { a "a" } { b "b" } }
  92                     T{ ret }
  93                 } }
  94             }
  95             T{ ptx-func
  96                 { name "frob" }
  97                 { params {
  98                     T{ ptx-variable { align 8 } { storage-space .param } { type .u64 } { name "in" } }
  99                     T{ ptx-variable { align 8 } { storage-space .param } { type .u64 } { name "out" } }
 100                     T{ ptx-variable { align 8 } { storage-space .param } { type .u64 } { name "len" } }
 101                 } }
 102                 { body {
 103                     T{ ret }
 104                 } }
 105             }
 106             T{ ptx-func
 107                 { name "twib" }
 108                 { body {
 109                     T{ ret }
 110                 } }
 111             }
 112         } }
 113     } ptx>string
 114 ] unit-test
 115
 116 [ "a" ] [ [ "a" write-ptx-operand ] with-string-writer ] unit-test
 117 [ "2" ] [ [ 2 write-ptx-operand ] with-string-writer ] unit-test
 118 [ "0d4000000000000000" ] [ [ 2.0 write-ptx-operand ] with-string-writer ] unit-test
 119 [ "!a" ] [ [ T{ ptx-negation f "a" } write-ptx-operand ] with-string-writer ] unit-test
 120 [ "{a, b, c, d}" ] [ [ T{ ptx-vector f { "a" "b" "c" "d" } } write-ptx-operand ] with-string-writer ] unit-test
 121 [ "[a]" ] [ [ T{ ptx-indirect f "a" 0 } write-ptx-operand ] with-string-writer ] unit-test
 122 [ "[a+1]" ] [ [ T{ ptx-indirect f "a" 1 } write-ptx-operand ] with-string-writer ] unit-test
 123 [ "[a-1]" ] [ [ T{ ptx-indirect f "a" -1 } write-ptx-operand ] with-string-writer ] unit-test
 124 [ "a[1]" ] [ [ T{ ptx-element f "a" 1 } write-ptx-operand ] with-string-writer ] unit-test
 125 [ "{a, b[2], 3, 0d4000000000000000}" ] [ [ T{ ptx-vector f { "a" T{ ptx-element f "b" 2 } 3 2.0 } } write-ptx-operand ] with-string-writer ] unit-test
 126
 127 [ """   .version 2.0
 128         .target sm_20
 129         abs.s32 a, b;
 130         @p abs.s32 a, b;
 131         @!p abs.s32 a, b;
 132 foo:    abs.s32 a, b;
 133         abs.ftz.f32 a, b;
 134 """ ] [
 135     T{ ptx
 136         { version "2.0" }
 137         { target T{ ptx-target { arch sm_20 } } }
 138         { body {
 139             T{ abs { type .s32 } { dest "a" } { a "b" } }
 140             T{ abs
 141                 { predicate "p" }
 142                 { type .s32 } { dest "a" } { a "b" }
 143             }
 144             T{ abs
 145                 { predicate T{ ptx-negation f "p" } }
 146                 { type .s32 } { dest "a" } { a "b" }
 147             }
 148             T{ abs
 149                 { label "foo" }
 150                 { type .s32 } { dest "a" } { a "b" }
 151             }
 152             T{ abs { type .f32 } { dest "a" } { a "b" } { ftz? t } }
 153         } }
 154     } ptx>string
 155 ] unit-test
 156
 157 [ """   .version 2.0
 158         .target sm_20
 159         add.s32 a, b, c;
 160         add.cc.s32 a, b, c;
 161         add.sat.s32 a, b, c;
 162         add.ftz.f32 a, b, c;
 163         add.ftz.sat.f32 a, b, c;
 164         add.rz.sat.f32 a, b, c;
 165         add.rz.ftz.sat.f32 a, b, c;
 166 """ ] [
 167     T{ ptx
 168         { version "2.0" }
 169         { target T{ ptx-target { arch sm_20 } } }
 170         { body {
 171             T{ add { type .s32 } { dest "a" } { a "b" } { b "c" } }
 172             T{ add { cc? t } { type .s32 } { dest "a" } { a "b" } { b "c" } }
 173             T{ add { sat? t } { type .s32 } { dest "a" } { a "b" } { b "c" } }
 174             T{ add { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } }
 175             T{ add { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } }
 176             T{ add { round .rz } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } }
 177             T{ add { round .rz } { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } }
 178         } }
 179     } ptx>string
 180 ] unit-test
 181
 182 [ """   .version 2.0
 183         .target sm_20
 184         addc.s32 a, b, c;
 185         addc.cc.s32 a, b, c;
 186 """ ] [
 187     T{ ptx
 188         { version "2.0" }
 189         { target T{ ptx-target { arch sm_20 } } }
 190         { body {
 191             T{ addc { type .s32 } { dest "a" } { a "b" } { b "c" } }
 192             T{ addc { cc? t } { type .s32 } { dest "a" } { a "b" } { b "c" } }
 193         } }
 194     } ptx>string
 195 ] unit-test
 196
 197 [ """   .version 2.0
 198         .target sm_20
 199         and.b32 a, b, c;
 200 """ ] [
 201     T{ ptx
 202         { version "2.0" }
 203         { target T{ ptx-target { arch sm_20 } } }
 204         { body {
 205             T{ and { type .b32 } { dest "a" } { a "b" } { b "c" } }
 206         } }
 207     } ptx>string
 208 ] unit-test
 209
 210 [ """   .version 2.0
 211         .target sm_20
 212         atom.and.u32 a, [b], c;
 213         atom.global.or.u32 a, [b], c;
 214         atom.shared.cas.u32 a, [b], c, d;
 215 """ ] [
 216     T{ ptx
 217         { version "2.0" }
 218         { target T{ ptx-target { arch sm_20 } } }
 219         { body {
 220             T{ atom { op .and } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } { b "c" } }
 221             T{ atom { storage-space .global } { op .or } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } { b "c" } }
 222             T{ atom { storage-space .shared } { op .cas } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } { b "c" } { c "d" } }
 223
 224         } }
 225     } ptx>string
 226 ] unit-test
 227
 228 [ """   .version 2.0
 229         .target sm_20
 230         bar.arrive a, b;
 231         bar.red.popc.u32 a, b, d;
 232         bar.red.popc.u32 a, b, !d;
 233         bar.red.popc.u32 a, b, c, !d;
 234         bar.sync a;
 235         bar.sync a, b;
 236 """ ] [
 237     T{ ptx
 238         { version "2.0" }
 239         { target T{ ptx-target { arch sm_20 } } }
 240         { body {
 241             T{ bar.arrive { a "a" } { b "b" } }
 242             T{ bar.red { op .popc } { type .u32 } { dest "a" } { a "b" } { c "d" } }
 243             T{ bar.red { op .popc } { type .u32 } { dest "a" } { a "b" } { c T{ ptx-negation f "d" } } }
 244             T{ bar.red { op .popc } { type .u32 } { dest "a" } { a "b" } { b "c" } { c T{ ptx-negation f "d" } } }
 245             T{ bar.sync { a "a" } }
 246             T{ bar.sync { a "a" } { b "b" } }
 247         } }
 248     } ptx>string
 249 ] unit-test
 250
 251 [ """   .version 2.0
 252         .target sm_20
 253         bfe.u32 a, b, c, d;
 254 """ ] [
 255     T{ ptx
 256         { version "2.0" }
 257         { target T{ ptx-target { arch sm_20 } } }
 258         { body {
 259             T{ bfe { type .u32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 260         } }
 261     } ptx>string
 262 ] unit-test
 263
 264 [ """   .version 2.0
 265         .target sm_20
 266         bfi.u32 a, b, c, d, e;
 267 """ ] [
 268     T{ ptx
 269         { version "2.0" }
 270         { target T{ ptx-target { arch sm_20 } } }
 271         { body {
 272             T{ bfi { type .u32 } { dest "a" } { a "b" } { b "c" } { c "d" } { d "e" } }
 273         } }
 274     } ptx>string
 275 ] unit-test
 276
 277 [ """   .version 2.0
 278         .target sm_20
 279         bfind.u32 a, b;
 280         bfind.shiftamt.u32 a, b;
 281 """ ] [
 282     T{ ptx
 283         { version "2.0" }
 284         { target T{ ptx-target { arch sm_20 } } }
 285         { body {
 286             T{ bfind { type .u32 } { dest "a" } { a "b" } }
 287             T{ bfind { type .u32 } { shiftamt? t } { dest "a" } { a "b" } }
 288         } }
 289     } ptx>string
 290 ] unit-test
 291
 292 [ """   .version 2.0
 293         .target sm_20
 294         bra foo;
 295         bra.uni bar;
 296 """ ] [
 297     T{ ptx
 298         { version "2.0" }
 299         { target T{ ptx-target { arch sm_20 } } }
 300         { body {
 301             T{ bra { target "foo" } }
 302             T{ bra { uni? t } { target "bar" } }
 303         } }
 304     } ptx>string
 305 ] unit-test
 306
 307 [ """   .version 2.0
 308         .target sm_20
 309         brev.b32 a, b;
 310 """ ] [
 311     T{ ptx
 312         { version "2.0" }
 313         { target T{ ptx-target { arch sm_20 } } }
 314         { body {
 315             T{ brev { type .b32 } { dest "a" } { a "b" } }
 316         } }
 317     } ptx>string
 318 ] unit-test
 319
 320 [ """   .version 2.0
 321         .target sm_20
 322         brkpt;
 323 """ ] [
 324     T{ ptx
 325         { version "2.0" }
 326         { target T{ ptx-target { arch sm_20 } } }
 327         { body {
 328             T{ brkpt }
 329         } }
 330     } ptx>string
 331 ] unit-test
 332
 333 [ """   .version 2.0
 334         .target sm_20
 335         call foo;
 336         call.uni foo;
 337         call (a), foo;
 338         call (a), foo, (b);
 339         call (a), foo, (b, c);
 340         call (a), foo, (b, c, d);
 341         call (a[2]), foo, (b, c, d[3]);
 342         call foo, (b, c, d);
 343 """ ] [
 344     T{ ptx
 345         { version "2.0" }
 346         { target T{ ptx-target { arch sm_20 } } }
 347         { body {
 348             T{ call { target "foo" } }
 349             T{ call { uni? t } { target "foo" } }
 350             T{ call { return "a" } { target "foo" } }
 351             T{ call { return "a" } { target "foo" } { params { "b" } } }
 352             T{ call { return "a" } { target "foo" } { params { "b" "c" } } }
 353             T{ call { return "a" } { target "foo" } { params { "b" "c" "d" } } }
 354             T{ call { return T{ ptx-element f "a" 2 } } { target "foo" } { params { "b" "c" T{ ptx-element f "d" 3 } } } }
 355             T{ call { target "foo" } { params { "b" "c" "d" } } }
 356         } }
 357     } ptx>string
 358 ] unit-test
 359
 360 [ """   .version 2.0
 361         .target sm_20
 362         clz.b32 a, b;
 363 """ ] [
 364     T{ ptx
 365         { version "2.0" }
 366         { target T{ ptx-target { arch sm_20 } } }
 367         { body {
 368             T{ clz { type .b32 } { dest "a" } { a "b" } }
 369         } }
 370     } ptx>string
 371 ] unit-test
 372
 373 [ """   .version 2.0
 374         .target sm_20
 375         cnot.b32 a, b;
 376 """ ] [
 377     T{ ptx
 378         { version "2.0" }
 379         { target T{ ptx-target { arch sm_20 } } }
 380         { body {
 381             T{ cnot { type .b32 } { dest "a" } { a "b" } }
 382         } }
 383     } ptx>string
 384 ] unit-test
 385
 386 [ """   .version 2.0
 387         .target sm_20
 388         copysign.f64 a, b, c;
 389 """ ] [
 390     T{ ptx
 391         { version "2.0" }
 392         { target T{ ptx-target { arch sm_20 } } }
 393         { body {
 394             T{ copysign { type .f64 } { dest "a" } { a "b" } { b "c" } }
 395         } }
 396     } ptx>string
 397 ] unit-test
 398
 399 [ """   .version 2.0
 400         .target sm_20
 401         cos.approx.f32 a, b;
 402 """ ] [
 403     T{ ptx
 404         { version "2.0" }
 405         { target T{ ptx-target { arch sm_20 } } }
 406         { body {
 407             T{ cos { round .approx } { type .f32 } { dest "a" } { a "b" } }
 408         } }
 409     } ptx>string
 410 ] unit-test
 411
 412 [ """   .version 2.0
 413         .target sm_20
 414         cvt.f32.s32 a, b;
 415         cvt.s32.f32 a, b;
 416         cvt.rp.f32.f64 a, b;
 417         cvt.rpi.s32.f32 a, b;
 418         cvt.ftz.f32.f64 a, b;
 419         cvt.sat.f32.f64 a, b;
 420         cvt.ftz.sat.f32.f64 a, b;
 421         cvt.rp.ftz.sat.f32.f64 a, b;
 422 """ ] [
 423     T{ ptx
 424         { version "2.0" }
 425         { target T{ ptx-target { arch sm_20 } } }
 426         { body {
 427             T{ cvt { dest-type .f32 } { type .s32 } { dest "a" } { a "b" } }
 428             T{ cvt { dest-type .s32 } { type .f32 } { dest "a" } { a "b" } }
 429             T{ cvt { round .rp } { dest-type .f32 } { type .f64 } { dest "a" } { a "b" } }
 430             T{ cvt { round .rpi } { dest-type .s32 } { type .f32 } { dest "a" } { a "b" } }
 431             T{ cvt { ftz? t } { dest-type .f32 } { type .f64 } { dest "a" } { a "b" } }
 432             T{ cvt { sat? t } { dest-type .f32 } { type .f64 } { dest "a" } { a "b" } }
 433             T{ cvt { ftz? t } { sat? t } { dest-type .f32 } { type .f64 } { dest "a" } { a "b" } }
 434             T{ cvt { round .rp } { ftz? t } { sat? t } { dest-type .f32 } { type .f64 } { dest "a" } { a "b" } }
 435         } }
 436     } ptx>string
 437 ] unit-test
 438
 439 [ """   .version 2.0
 440         .target sm_20
 441         cvta.global.u64 a, b;
 442         cvta.shared.u64 a, b;
 443         cvta.to.shared.u64 a, b;
 444 """ ] [
 445     T{ ptx
 446         { version "2.0" }
 447         { target T{ ptx-target { arch sm_20 } } }
 448         { body {
 449             T{ cvta { storage-space .global } { type .u64 } { dest "a" } { a "b" } }
 450             T{ cvta { storage-space .shared } { type .u64 } { dest "a" } { a "b" } }
 451             T{ cvta { to? t } { storage-space .shared } { type .u64 } { dest "a" } { a "b" } }
 452         } }
 453     } ptx>string
 454 ] unit-test
 455
 456 [ """   .version 2.0
 457         .target sm_20
 458         div.u32 a, b, c;
 459         div.approx.f32 a, b, c;
 460         div.approx.ftz.f32 a, b, c;
 461         div.full.f32 a, b, c;
 462         div.full.ftz.f32 a, b, c;
 463         div.f32 a, b, c;
 464         div.rz.f32 a, b, c;
 465         div.ftz.f32 a, b, c;
 466         div.rz.ftz.f32 a, b, c;
 467         div.f64 a, b, c;
 468         div.rz.f64 a, b, c;
 469 """ ] [
 470     T{ ptx
 471         { version "2.0" }
 472         { target T{ ptx-target { arch sm_20 } } }
 473         { body {
 474             T{ div { type .u32 } { dest "a" } { a "b" } { b "c" } }
 475             T{ div { round .approx } { type .f32 } { dest "a" } { a "b" } { b "c" } }
 476             T{ div { round .approx } { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } }
 477             T{ div { round .full } { type .f32 } { dest "a" } { a "b" } { b "c" } }
 478             T{ div { round .full } { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } }
 479             T{ div { type .f32 } { dest "a" } { a "b" } { b "c" } }
 480             T{ div { round .rz } { type .f32 } { dest "a" } { a "b" } { b "c" } }
 481             T{ div { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } }
 482             T{ div { round .rz } { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } }
 483             T{ div { type .f64 } { dest "a" } { a "b" } { b "c" } }
 484             T{ div { round .rz } { type .f64 } { dest "a" } { a "b" } { b "c" } }
 485         } }
 486     } ptx>string
 487 ] unit-test
 488
 489 [ """   .version 2.0
 490         .target sm_20
 491         ex2.approx.f32 a, b;
 492 """ ] [
 493     T{ ptx
 494         { version "2.0" }
 495         { target T{ ptx-target { arch sm_20 } } }
 496         { body {
 497             T{ ex2 { round .approx } { type .f32 } { dest "a" } { a "b" } }
 498         } }
 499     } ptx>string
 500 ] unit-test
 501
 502 [ """   .version 2.0
 503         .target sm_20
 504         exit;
 505 """ ] [
 506     T{ ptx
 507         { version "2.0" }
 508         { target T{ ptx-target { arch sm_20 } } }
 509         { body {
 510             T{ exit }
 511         } }
 512     } ptx>string
 513 ] unit-test
 514
 515 [ """   .version 2.0
 516         .target sm_20
 517         fma.f32 a, b, c, d;
 518         fma.sat.f32 a, b, c, d;
 519         fma.ftz.f32 a, b, c, d;
 520         fma.ftz.sat.f32 a, b, c, d;
 521         fma.rz.sat.f32 a, b, c, d;
 522         fma.rz.ftz.sat.f32 a, b, c, d;
 523 """ ] [
 524     T{ ptx
 525         { version "2.0" }
 526         { target T{ ptx-target { arch sm_20 } } }
 527         { body {
 528             T{ fma { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 529             T{ fma { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 530             T{ fma { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 531             T{ fma { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 532             T{ fma { round .rz } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 533             T{ fma { round .rz } { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 534         } }
 535     } ptx>string
 536 ] unit-test
 537
 538 [ """   .version 2.0
 539         .target sm_20
 540         isspacep.shared a, b;
 541 """ ] [
 542     T{ ptx
 543         { version "2.0" }
 544         { target T{ ptx-target { arch sm_20 } } }
 545         { body {
 546             T{ isspacep { storage-space .shared } { dest "a" } { a "b" } }
 547         } }
 548     } ptx>string
 549 ] unit-test
 550
 551 [ """   .version 2.0
 552         .target sm_20
 553         ld.u32 a, [b];
 554         ld.v2.u32 a, [b];
 555         ld.v4.u32 a, [b];
 556         ld.v4.u32 {a, b, c, d}, [e];
 557         ld.lu.u32 a, [b];
 558         ld.const.lu.u32 a, [b];
 559         ld.volatile.const[5].u32 a, [b];
 560 """ ] [
 561     T{ ptx
 562         { version "2.0" }
 563         { target T{ ptx-target { arch sm_20 } } }
 564         { body {
 565             T{ ld { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } }
 566             T{ ld { type T{ .v2 { of .u32 } } } { dest "a" } { a T{ ptx-indirect f "b" } } }
 567             T{ ld { type T{ .v4 { of .u32 } } } { dest "a" } { a T{ ptx-indirect f "b" } } }
 568             T{ ld { type T{ .v4 { of .u32 } } } { dest T{ ptx-vector f { "a" "b" "c" "d" } } } { a "[e]" } }
 569             T{ ld { cache-op .lu } { type .u32 } { dest "a" } { a "[b]" } }
 570             T{ ld { storage-space T{ .const } } { cache-op .lu } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } }
 571             T{ ld { volatile? t } { storage-space T{ .const { bank 5 } } } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } }
 572         } }
 573     } ptx>string
 574 ] unit-test
 575
 576 [ """   .version 2.0
 577         .target sm_20
 578         ldu.u32 a, [b];
 579         ldu.v2.u32 a, [b];
 580         ldu.v4.u32 a, [b];
 581         ldu.v4.u32 {a, b, c, d}, [e];
 582         ldu.lu.u32 a, [b];
 583         ldu.const.lu.u32 a, [b];
 584         ldu.volatile.const[5].u32 a, [b];
 585 """ ] [
 586     T{ ptx
 587         { version "2.0" }
 588         { target T{ ptx-target { arch sm_20 } } }
 589         { body {
 590             T{ ldu { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } }
 591             T{ ldu { type T{ .v2 { of .u32 } } } { dest "a" } { a T{ ptx-indirect f "b" } } }
 592             T{ ldu { type T{ .v4 { of .u32 } } } { dest "a" } { a T{ ptx-indirect f "b" } } }
 593             T{ ldu { type T{ .v4 { of .u32 } } } { dest T{ ptx-vector f { "a" "b" "c" "d" } } } { a "[e]" } }
 594             T{ ldu { cache-op .lu } { type .u32 } { dest "a" } { a "[b]" } }
 595             T{ ldu { storage-space T{ .const } } { cache-op .lu } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } }
 596             T{ ldu { volatile? t } { storage-space T{ .const { bank 5 } } } { type .u32 } { dest "a" } { a T{ ptx-indirect f "b" } } }
 597         } }
 598     } ptx>string
 599 ] unit-test
 600
 601 [ """   .version 2.0
 602         .target sm_20
 603         lg2.approx.f32 a, b;
 604 """ ] [
 605     T{ ptx
 606         { version "2.0" }
 607         { target T{ ptx-target { arch sm_20 } } }
 608         { body {
 609             T{ lg2 { round .approx } { type .f32 } { dest "a" } { a "b" } }
 610         } }
 611     } ptx>string
 612 ] unit-test
 613
 614 [ """   .version 2.0
 615         .target sm_20
 616         mad.s32 a, b, c, d;
 617         mad.lo.s32 a, b, c, d;
 618         mad.sat.s32 a, b, c, d;
 619         mad.hi.sat.s32 a, b, c, d;
 620         mad.ftz.f32 a, b, c, d;
 621         mad.ftz.sat.f32 a, b, c, d;
 622         mad.rz.sat.f32 a, b, c, d;
 623         mad.rz.ftz.sat.f32 a, b, c, d;
 624 """ ] [
 625     T{ ptx
 626         { version "2.0" }
 627         { target T{ ptx-target { arch sm_20 } } }
 628         { body {
 629             T{ mad { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 630             T{ mad { mode .lo } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 631             T{ mad { sat? t } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 632             T{ mad { mode .hi } { sat? t } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 633             T{ mad { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 634             T{ mad { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 635             T{ mad { round .rz } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 636             T{ mad { round .rz } { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 637         } }
 638     } ptx>string
 639 ] unit-test
 640
 641 [ """   .version 2.0
 642         .target sm_20
 643         mad24.s32 a, b, c, d;
 644         mad24.lo.s32 a, b, c, d;
 645         mad24.sat.s32 a, b, c, d;
 646         mad24.hi.sat.s32 a, b, c, d;
 647 """ ] [
 648     T{ ptx
 649         { version "2.0" }
 650         { target T{ ptx-target { arch sm_20 } } }
 651         { body {
 652             T{ mad24 { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 653             T{ mad24 { mode .lo } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 654             T{ mad24 { sat? t } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 655             T{ mad24 { mode .hi } { sat? t } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 656         } }
 657     } ptx>string
 658 ] unit-test
 659
 660 [ """   .version 2.0
 661         .target sm_20
 662         neg.s32 a, b;
 663         neg.f32 a, b;
 664         neg.ftz.f32 a, b;
 665 """ ] [
 666     T{ ptx
 667         { version "2.0" }
 668         { target T{ ptx-target { arch sm_20 } } }
 669         { body {
 670             T{ neg { type .s32 } { dest "a" } { a "b" } }
 671             T{ neg { type .f32 } { dest "a" } { a "b" } }
 672             T{ neg { ftz? t } { type .f32 } { dest "a" } { a "b" } }
 673         } }
 674     } ptx>string
 675 ] unit-test
 676
 677 [ """   .version 2.0
 678         .target sm_20
 679         not.b32 a, b;
 680 """ ] [
 681     T{ ptx
 682         { version "2.0" }
 683         { target T{ ptx-target { arch sm_20 } } }
 684         { body {
 685             T{ not { type .b32 } { dest "a" } { a "b" } }
 686         } }
 687     } ptx>string
 688 ] unit-test
 689
 690 [ """   .version 2.0
 691         .target sm_20
 692         or.b32 a, b, c;
 693 """ ] [
 694     T{ ptx
 695         { version "2.0" }
 696         { target T{ ptx-target { arch sm_20 } } }
 697         { body {
 698             T{ or { type .b32 } { dest "a" } { a "b" } { b "c" } }
 699         } }
 700     } ptx>string
 701 ] unit-test
 702
 703 [ """   .version 2.0
 704         .target sm_20
 705         pmevent a;
 706 """ ] [
 707     T{ ptx
 708         { version "2.0" }
 709         { target T{ ptx-target { arch sm_20 } } }
 710         { body {
 711             T{ pmevent { a "a" } }
 712         } }
 713     } ptx>string
 714 ] unit-test
 715
 716 [ """   .version 2.0
 717         .target sm_20
 718         popc.b64 a, b;
 719 """ ] [
 720     T{ ptx
 721         { version "2.0" }
 722         { target T{ ptx-target { arch sm_20 } } }
 723         { body {
 724             T{ popc { type .b64 } { dest "a" } { a "b" } }
 725         } }
 726     } ptx>string
 727 ] unit-test
 728
 729 [ """   .version 2.0
 730         .target sm_20
 731         prefetch.L1 [a];
 732         prefetch.local.L2 [a];
 733         prefetchu.L1 [a];
 734 """ ] [
 735     T{ ptx
 736         { version "2.0" }
 737         { target T{ ptx-target { arch sm_20 } } }
 738         { body {
 739             T{ prefetch { level .L1 } { a T{ ptx-indirect f "a" } } }
 740             T{ prefetch { storage-space .local } { level .L2 } { a T{ ptx-indirect f "a" } } }
 741             T{ prefetchu { level .L1 } { a T{ ptx-indirect f "a" } } }
 742         } }
 743     } ptx>string
 744 ] unit-test
 745
 746 [ """   .version 2.0
 747         .target sm_20
 748         prmt.b32 a, b, c, d;
 749         prmt.b32.f4e a, b, c, d;
 750 """ ] [
 751     T{ ptx
 752         { version "2.0" }
 753         { target T{ ptx-target { arch sm_20 } } }
 754         { body {
 755             T{ prmt { type .b32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 756             T{ prmt { type .b32 } { mode .f4e } { dest "a" } { a "b" } { b "c" } { c "d" } }
 757         } }
 758     } ptx>string
 759 ] unit-test
 760
 761 [ """   .version 2.0
 762         .target sm_20
 763         rcp.approx.f32 a, b;
 764         rcp.approx.ftz.f32 a, b;
 765         rcp.f32 a, b;
 766         rcp.rz.f32 a, b;
 767         rcp.ftz.f32 a, b;
 768         rcp.rz.ftz.f32 a, b;
 769         rcp.f64 a, b;
 770         rcp.rz.f64 a, b;
 771 """ ] [
 772     T{ ptx
 773         { version "2.0" }
 774         { target T{ ptx-target { arch sm_20 } } }
 775         { body {
 776             T{ rcp { round .approx } { type .f32 } { dest "a" } { a "b" } }
 777             T{ rcp { round .approx } { ftz? t } { type .f32 } { dest "a" } { a "b" } }
 778             T{ rcp { type .f32 } { dest "a" } { a "b" } }
 779             T{ rcp { round .rz } { type .f32 } { dest "a" } { a "b" } }
 780             T{ rcp { ftz? t } { type .f32 } { dest "a" } { a "b" } }
 781             T{ rcp { round .rz } { ftz? t } { type .f32 } { dest "a" } { a "b" } }
 782             T{ rcp { type .f64 } { dest "a" } { a "b" } }
 783             T{ rcp { round .rz } { type .f64 } { dest "a" } { a "b" } }
 784         } }
 785     } ptx>string
 786 ] unit-test
 787
 788 [ """   .version 2.0
 789         .target sm_20
 790         red.and.u32 [a], b;
 791         red.global.and.u32 [a], b;
 792 """ ] [
 793     T{ ptx
 794         { version "2.0" }
 795         { target T{ ptx-target { arch sm_20 } } }
 796         { body {
 797             T{ red { op .and } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } }
 798             T{ red { storage-space .global } { op .and } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } }
 799         } }
 800     } ptx>string
 801 ] unit-test
 802
 803 [ """   .version 2.0
 804         .target sm_20
 805         rsqrt.approx.f32 a, b;
 806         rsqrt.approx.ftz.f32 a, b;
 807         rsqrt.approx.f64 a, b;
 808 """ ] [
 809     T{ ptx
 810         { version "2.0" }
 811         { target T{ ptx-target { arch sm_20 } } }
 812         { body {
 813             T{ rsqrt { round .approx } { type .f32 } { dest "a" } { a "b" } }
 814             T{ rsqrt { round .approx } { ftz? t } { type .f32 } { dest "a" } { a "b" } }
 815             T{ rsqrt { round .approx } { type .f64 } { dest "a" } { a "b" } }
 816         } }
 817     } ptx>string
 818 ] unit-test
 819
 820 [ """   .version 2.0
 821         .target sm_20
 822         rsqrt.approx.f32 a, b;
 823         rsqrt.approx.ftz.f32 a, b;
 824         rsqrt.approx.f64 a, b;
 825 """ ] [
 826     T{ ptx
 827         { version "2.0" }
 828         { target T{ ptx-target { arch sm_20 } } }
 829         { body {
 830             T{ rsqrt { round .approx } { type .f32 } { dest "a" } { a "b" } }
 831             T{ rsqrt { round .approx } { ftz? t } { type .f32 } { dest "a" } { a "b" } }
 832             T{ rsqrt { round .approx } { type .f64 } { dest "a" } { a "b" } }
 833         } }
 834     } ptx>string
 835 ] unit-test
 836
 837 [ """   .version 2.0
 838         .target sm_20
 839         sad.u32 a, b, c, d;
 840 """ ] [
 841     T{ ptx
 842         { version "2.0" }
 843         { target T{ ptx-target { arch sm_20 } } }
 844         { body {
 845             T{ sad { type .u32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 846         } }
 847     } ptx>string
 848 ] unit-test
 849
 850 [ """   .version 2.0
 851         .target sm_20
 852         selp.u32 a, b, c, d;
 853 """ ] [
 854     T{ ptx
 855         { version "2.0" }
 856         { target T{ ptx-target { arch sm_20 } } }
 857         { body {
 858             T{ selp { type .u32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 859         } }
 860     } ptx>string
 861 ] unit-test
 862
 863 [ """   .version 2.0
 864         .target sm_20
 865         set.gt.u32.s32 a, b, c;
 866         set.gt.ftz.u32.f32 a, b, c;
 867         set.gt.and.ftz.u32.f32 a, b, c, d;
 868         set.gt.and.ftz.u32.f32 a, b, c, !d;
 869 """ ] [
 870     T{ ptx
 871         { version "2.0" }
 872         { target T{ ptx-target { arch sm_20 } } }
 873         { body {
 874             T{ set { cmp-op .gt } { dest-type .u32 } { type .s32 } { dest "a" } { a "b" } { b "c" } }
 875             T{ set { cmp-op .gt } { ftz? t } { dest-type .u32 } { type .f32 } { dest "a" } { a "b" } { b "c" } }
 876             T{ set { cmp-op .gt } { bool-op .and } { ftz? t } { dest-type .u32 } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 877             T{ set { cmp-op .gt } { bool-op .and } { ftz? t } { dest-type .u32 } { type .f32 } { dest "a" } { a "b" } { b "c" } { c T{ ptx-negation f "d" } } }
 878         } }
 879     } ptx>string
 880 ] unit-test
 881
 882 [ """   .version 2.0
 883         .target sm_20
 884         setp.gt.s32 a, b, c;
 885         setp.gt.s32 a|z, b, c;
 886         setp.gt.ftz.f32 a, b, c;
 887         setp.gt.and.ftz.f32 a, b, c, d;
 888         setp.gt.and.ftz.f32 a, b, c, !d;
 889 """ ] [
 890     T{ ptx
 891         { version "2.0" }
 892         { target T{ ptx-target { arch sm_20 } } }
 893         { body {
 894             T{ setp { cmp-op .gt } { type .s32 } { dest "a" } { a "b" } { b "c" } }
 895             T{ setp { cmp-op .gt } { type .s32 } { dest "a" } { |dest "z" } { a "b" } { b "c" } }
 896             T{ setp { cmp-op .gt } { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } }
 897             T{ setp { cmp-op .gt } { bool-op .and } { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 898             T{ setp { cmp-op .gt } { bool-op .and } { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } { c "!d" } }
 899         } }
 900     } ptx>string
 901 ] unit-test
 902
 903 [ """   .version 2.0
 904         .target sm_20
 905         shl.b32 a, b, c;
 906 """ ] [
 907     T{ ptx
 908         { version "2.0" }
 909         { target T{ ptx-target { arch sm_20 } } }
 910         { body {
 911             T{ shl { type .b32 } { dest "a" } { a "b" } { b "c" } }
 912         } }
 913     } ptx>string
 914 ] unit-test
 915
 916 [ """   .version 2.0
 917         .target sm_20
 918         shr.b32 a, b, c;
 919 """ ] [
 920     T{ ptx
 921         { version "2.0" }
 922         { target T{ ptx-target { arch sm_20 } } }
 923         { body {
 924             T{ shr { type .b32 } { dest "a" } { a "b" } { b "c" } }
 925         } }
 926     } ptx>string
 927 ] unit-test
 928
 929 [ """   .version 2.0
 930         .target sm_20
 931         sin.approx.f32 a, b;
 932 """ ] [
 933     T{ ptx
 934         { version "2.0" }
 935         { target T{ ptx-target { arch sm_20 } } }
 936         { body {
 937             T{ sin { round .approx } { type .f32 } { dest "a" } { a "b" } }
 938         } }
 939     } ptx>string
 940 ] unit-test
 941
 942 [ """   .version 2.0
 943         .target sm_20
 944         slct.f32.s32 a, b, c, d;
 945         slct.ftz.f32.s32 a, b, c, d;
 946 """ ] [
 947     T{ ptx
 948         { version "2.0" }
 949         { target T{ ptx-target { arch sm_20 } } }
 950         { body {
 951             T{ slct { dest-type .f32 } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 952             T{ slct { ftz? t } { dest-type .f32 } { type .s32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
 953         } }
 954     } ptx>string
 955 ] unit-test
 956
 957 [ """   .version 2.0
 958         .target sm_20
 959         sqrt.approx.f32 a, b;
 960         sqrt.approx.ftz.f32 a, b;
 961         sqrt.f32 a, b;
 962         sqrt.rz.f32 a, b;
 963         sqrt.ftz.f32 a, b;
 964         sqrt.rz.ftz.f32 a, b;
 965         sqrt.f64 a, b;
 966         sqrt.rz.f64 a, b;
 967 """ ] [
 968     T{ ptx
 969         { version "2.0" }
 970         { target T{ ptx-target { arch sm_20 } } }
 971         { body {
 972             T{ sqrt { round .approx } { type .f32 } { dest "a" } { a "b" } }
 973             T{ sqrt { round .approx } { ftz? t } { type .f32 } { dest "a" } { a "b" } }
 974             T{ sqrt { type .f32 } { dest "a" } { a "b" } }
 975             T{ sqrt { round .rz } { type .f32 } { dest "a" } { a "b" } }
 976             T{ sqrt { ftz? t } { type .f32 } { dest "a" } { a "b" } }
 977             T{ sqrt { round .rz } { ftz? t } { type .f32 } { dest "a" } { a "b" } }
 978             T{ sqrt { type .f64 } { dest "a" } { a "b" } }
 979             T{ sqrt { round .rz } { type .f64 } { dest "a" } { a "b" } }
 980         } }
 981     } ptx>string
 982 ] unit-test
 983
 984 [ """   .version 2.0
 985         .target sm_20
 986         st.u32 [a], b;
 987         st.v2.u32 [a], b;
 988         st.v4.u32 [a], b;
 989         st.v4.u32 [a], {b, c, d, e};
 990         st.lu.u32 [a], b;
 991         st.local.lu.u32 [a], b;
 992         st.volatile.local.u32 [a], b;
 993 """ ] [
 994     T{ ptx
 995         { version "2.0" }
 996         { target T{ ptx-target { arch sm_20 } } }
 997         { body {
 998             T{ st { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } }
 999             T{ st { type T{ .v2 { of .u32 } } } { dest T{ ptx-indirect f "a" } } { a "b" } }
1000             T{ st { type T{ .v4 { of .u32 } } } { dest T{ ptx-indirect f "a" } } { a "b" } }
1001             T{ st { type T{ .v4 { of .u32 } } } { dest T{ ptx-indirect f "a" } } { a T{ ptx-vector f { "b" "c" "d" "e" } } } }
1002             T{ st { cache-op .lu } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } }
1003             T{ st { storage-space .local } { cache-op .lu } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } }
1004             T{ st { volatile? t } { storage-space .local } { type .u32 } { dest T{ ptx-indirect f "a" } } { a "b" } }
1005         } }
1006     } ptx>string
1007 ] unit-test
1008
1009 [ """   .version 2.0
1010         .target sm_20
1011         sub.s32 a, b, c;
1012         sub.cc.s32 a, b, c;
1013         sub.sat.s32 a, b, c;
1014         sub.ftz.f32 a, b, c;
1015         sub.ftz.sat.f32 a, b, c;
1016         sub.rz.sat.f32 a, b, c;
1017         sub.rz.ftz.sat.f32 a, b, c;
1018 """ ] [
1019     T{ ptx
1020         { version "2.0" }
1021         { target T{ ptx-target { arch sm_20 } } }
1022         { body {
1023             T{ sub { type .s32 } { dest "a" } { a "b" } { b "c" } }
1024             T{ sub { cc? t } { type .s32 } { dest "a" } { a "b" } { b "c" } }
1025             T{ sub { sat? t } { type .s32 } { dest "a" } { a "b" } { b "c" } }
1026             T{ sub { ftz? t } { type .f32 } { dest "a" } { a "b" } { b "c" } }
1027             T{ sub { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } }
1028             T{ sub { round .rz } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } }
1029             T{ sub { round .rz } { ftz? t } { sat? t } { type .f32 } { dest "a" } { a "b" } { b "c" } }
1030         } }
1031     } ptx>string
1032 ] unit-test
1033
1034 [ """   .version 2.0
1035         .target sm_20
1036         subc.s32 a, b, c;
1037         subc.cc.s32 a, b, c;
1038 """ ] [
1039     T{ ptx
1040         { version "2.0" }
1041         { target T{ ptx-target { arch sm_20 } } }
1042         { body {
1043             T{ subc { type .s32 } { dest "a" } { a "b" } { b "c" } }
1044             T{ subc { cc? t } { type .s32 } { dest "a" } { a "b" } { b "c" } }
1045         } }
1046     } ptx>string
1047 ] unit-test
1048
1049 [ """   .version 2.0
1050         .target sm_20
1051         testp.finite.f32 a, b;
1052 """ ] [
1053     T{ ptx
1054         { version "2.0" }
1055         { target T{ ptx-target { arch sm_20 } } }
1056         { body {
1057             T{ testp { op .finite } { type .f32 } { dest "a" } { a "b" } }
1058         } }
1059     } ptx>string
1060 ] unit-test
1061
1062 [ """   .version 2.0
1063         .target sm_20
1064         trap;
1065 """ ] [
1066     T{ ptx
1067         { version "2.0" }
1068         { target T{ ptx-target { arch sm_20 } } }
1069         { body {
1070             T{ trap }
1071         } }
1072     } ptx>string
1073 ] unit-test
1074
1075 [ """   .version 2.0
1076         .target sm_20
1077         vote.all.pred a, b;
1078         vote.all.pred a, !b;
1079         vote.ballot.b32 a, b;
1080 """ ] [
1081     T{ ptx
1082         { version "2.0" }
1083         { target T{ ptx-target { arch sm_20 } } }
1084         { body {
1085             T{ vote { mode .all } { type .pred } { dest "a" } { a "b" } }
1086             T{ vote { mode .all } { type .pred } { dest "a" } { a "!b" } }
1087             T{ vote { mode .ballot } { type .b32 } { dest "a" } { a "b" } }
1088         } }
1089     } ptx>string
1090 ] unit-test
1091
1092 [ """   .version 2.0
1093         .target sm_20
1094         xor.b32 a, b, c;
1095 """ ] [
1096     T{ ptx
1097         { version "2.0" }
1098         { target T{ ptx-target { arch sm_20 } } }
1099         { body {
1100             T{ xor { type .b32 } { dest "a" } { a "b" } { b "c" } }
1101         } }
1102     } ptx>string
1103 ] unit-test