1 USING: cuda.ptx io.streams.string tools.test ;
9 { target T{ ptx-target { arch sm_20 } } }
14 .target sm_20, .texmode_independent
18 { target T{ ptx-target { arch sm_20 } { texmode .texmode_independent } } }
23 .target sm_11, map_f64_to_f32
27 { target T{ ptx-target
35 .target sm_11, map_f64_to_f32, .texmode_independent
39 { target T{ ptx-target
42 { texmode .texmode_independent }
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)
56 .func frob (.align 8 .param .u64 in, .align 8 .param .u64 out, .align 8 .param .u64 len)
67 { target T{ ptx-target { arch sm_20 } } }
70 { storage-space .global }
78 { storage-space .shared }
79 { type T{ .v4 f .f32 } }
84 { return T{ ptx-variable { storage-space .reg } { type .f32 } { name "sum" } } }
87 T{ ptx-variable { storage-space .reg } { type .f32 } { name "a" } }
88 T{ ptx-variable { storage-space .reg } { type .f32 } { name "b" } }
91 T{ add { round .rn } { type .f32 } { dest "sum" } { a "a" } { b "b" } }
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" } }
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
137 { target T{ ptx-target { arch sm_20 } } }
139 T{ abs { type .s32 } { dest "a" } { a "b" } }
142 { type .s32 } { dest "a" } { a "b" }
145 { predicate T{ ptx-negation f "p" } }
146 { type .s32 } { dest "a" } { a "b" }
150 { type .s32 } { dest "a" } { a "b" }
152 T{ abs { type .f32 } { dest "a" } { a "b" } { ftz? t } }
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;
169 { target T{ ptx-target { arch sm_20 } } }
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" } }
189 { target T{ ptx-target { arch sm_20 } } }
191 T{ addc { type .s32 } { dest "a" } { a "b" } { b "c" } }
192 T{ addc { cc? t } { type .s32 } { dest "a" } { a "b" } { b "c" } }
203 { target T{ ptx-target { arch sm_20 } } }
205 T{ and { type .b32 } { dest "a" } { a "b" } { b "c" } }
212 atom.and.u32 a, [b], c;
213 atom.global.or.u32 a, [b], c;
214 atom.shared.cas.u32 a, [b], c, d;
218 { target T{ ptx-target { arch sm_20 } } }
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" } }
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;
239 { target T{ ptx-target { arch sm_20 } } }
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" } }
257 { target T{ ptx-target { arch sm_20 } } }
259 T{ bfe { type .u32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
266 bfi.u32 a, b, c, d, e;
270 { target T{ ptx-target { arch sm_20 } } }
272 T{ bfi { type .u32 } { dest "a" } { a "b" } { b "c" } { c "d" } { d "e" } }
280 bfind.shiftamt.u32 a, b;
284 { target T{ ptx-target { arch sm_20 } } }
286 T{ bfind { type .u32 } { dest "a" } { a "b" } }
287 T{ bfind { type .u32 } { shiftamt? t } { dest "a" } { a "b" } }
299 { target T{ ptx-target { arch sm_20 } } }
301 T{ bra { target "foo" } }
302 T{ bra { uni? t } { target "bar" } }
313 { target T{ ptx-target { arch sm_20 } } }
315 T{ brev { type .b32 } { dest "a" } { a "b" } }
326 { target T{ ptx-target { arch sm_20 } } }
339 call (a), foo, (b, c);
340 call (a), foo, (b, c, d);
341 call (a[2]), foo, (b, c, d[3]);
346 { target T{ ptx-target { arch sm_20 } } }
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" } } }
366 { target T{ ptx-target { arch sm_20 } } }
368 T{ clz { type .b32 } { dest "a" } { a "b" } }
379 { target T{ ptx-target { arch sm_20 } } }
381 T{ cnot { type .b32 } { dest "a" } { a "b" } }
388 copysign.f64 a, b, c;
392 { target T{ ptx-target { arch sm_20 } } }
394 T{ copysign { type .f64 } { dest "a" } { a "b" } { b "c" } }
405 { target T{ ptx-target { arch sm_20 } } }
407 T{ cos { round .approx } { type .f32 } { dest "a" } { 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;
425 { target T{ ptx-target { arch sm_20 } } }
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" } }
441 cvta.global.u64 a, b;
442 cvta.shared.u64 a, b;
443 cvta.to.shared.u64 a, b;
447 { target T{ ptx-target { arch sm_20 } } }
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" } }
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;
466 div.rz.ftz.f32 a, b, c;
472 { target T{ ptx-target { arch sm_20 } } }
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" } }
495 { target T{ ptx-target { arch sm_20 } } }
497 T{ ex2 { round .approx } { type .f32 } { dest "a" } { a "b" } }
508 { target T{ ptx-target { arch sm_20 } } }
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;
526 { target T{ ptx-target { arch sm_20 } } }
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" } }
540 isspacep.shared a, b;
544 { target T{ ptx-target { arch sm_20 } } }
546 T{ isspacep { storage-space .shared } { dest "a" } { a "b" } }
556 ld.v4.u32 {a, b, c, d}, [e];
558 ld.const.lu.u32 a, [b];
559 ld.volatile.const[5].u32 a, [b];
563 { target T{ ptx-target { arch sm_20 } } }
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" } } }
581 ldu.v4.u32 {a, b, c, d}, [e];
583 ldu.const.lu.u32 a, [b];
584 ldu.volatile.const[5].u32 a, [b];
588 { target T{ ptx-target { arch sm_20 } } }
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" } } }
607 { target T{ ptx-target { arch sm_20 } } }
609 T{ lg2 { round .approx } { type .f32 } { dest "a" } { a "b" } }
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;
627 { target T{ ptx-target { arch sm_20 } } }
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" } }
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;
650 { target T{ ptx-target { arch sm_20 } } }
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" } }
668 { target T{ ptx-target { arch sm_20 } } }
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" } }
683 { target T{ ptx-target { arch sm_20 } } }
685 T{ not { type .b32 } { dest "a" } { a "b" } }
696 { target T{ ptx-target { arch sm_20 } } }
698 T{ or { type .b32 } { dest "a" } { a "b" } { b "c" } }
709 { target T{ ptx-target { arch sm_20 } } }
711 T{ pmevent { a "a" } }
722 { target T{ ptx-target { arch sm_20 } } }
724 T{ popc { type .b64 } { dest "a" } { a "b" } }
732 prefetch.local.L2 [a];
737 { target T{ ptx-target { arch sm_20 } } }
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" } } }
749 prmt.b32.f4e a, b, c, d;
753 { target T{ ptx-target { arch sm_20 } } }
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" } }
764 rcp.approx.ftz.f32 a, b;
774 { target T{ ptx-target { arch sm_20 } } }
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" } }
791 red.global.and.u32 [a], b;
795 { target T{ ptx-target { arch sm_20 } } }
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" } }
805 rsqrt.approx.f32 a, b;
806 rsqrt.approx.ftz.f32 a, b;
807 rsqrt.approx.f64 a, b;
811 { target T{ ptx-target { arch sm_20 } } }
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" } }
822 rsqrt.approx.f32 a, b;
823 rsqrt.approx.ftz.f32 a, b;
824 rsqrt.approx.f64 a, b;
828 { target T{ ptx-target { arch sm_20 } } }
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" } }
843 { target T{ ptx-target { arch sm_20 } } }
845 T{ sad { type .u32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
856 { target T{ ptx-target { arch sm_20 } } }
858 T{ selp { type .u32 } { dest "a" } { a "b" } { b "c" } { c "d" } }
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;
872 { target T{ ptx-target { arch sm_20 } } }
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" } } }
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;
892 { target T{ ptx-target { arch sm_20 } } }
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" } }
909 { target T{ ptx-target { arch sm_20 } } }
911 T{ shl { type .b32 } { dest "a" } { a "b" } { b "c" } }
922 { target T{ ptx-target { arch sm_20 } } }
924 T{ shr { type .b32 } { dest "a" } { a "b" } { b "c" } }
935 { target T{ ptx-target { arch sm_20 } } }
937 T{ sin { round .approx } { type .f32 } { dest "a" } { a "b" } }
944 slct.f32.s32 a, b, c, d;
945 slct.ftz.f32.s32 a, b, c, d;
949 { target T{ ptx-target { arch sm_20 } } }
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" } }
959 sqrt.approx.f32 a, b;
960 sqrt.approx.ftz.f32 a, b;
964 sqrt.rz.ftz.f32 a, b;
970 { target T{ ptx-target { arch sm_20 } } }
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" } }
989 st.v4.u32 [a], {b, c, d, e};
991 st.local.lu.u32 [a], b;
992 st.volatile.local.u32 [a], b;
996 { target T{ ptx-target { arch sm_20 } } }
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" } }
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;
1021 { target T{ ptx-target { arch sm_20 } } }
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" } }
1037 subc.cc.s32 a, b, c;
1041 { target T{ ptx-target { arch sm_20 } } }
1043 T{ subc { type .s32 } { dest "a" } { a "b" } { b "c" } }
1044 T{ subc { cc? t } { type .s32 } { dest "a" } { a "b" } { b "c" } }
1051 testp.finite.f32 a, b;
1055 { target T{ ptx-target { arch sm_20 } } }
1057 T{ testp { op .finite } { type .f32 } { dest "a" } { a "b" } }
1068 { target T{ ptx-target { arch sm_20 } } }
1078 vote.all.pred a, !b;
1079 vote.ballot.b32 a, b;
1083 { target T{ ptx-target { arch sm_20 } } }
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" } }
1098 { target T{ ptx-target { arch sm_20 } } }
1100 T{ xor { type .b32 } { dest "a" } { a "b" } { b "c" } }