Lines Matching refs:b32
784 defm SELP_b32 : SELP_PATTERN<"b32", i32, Int32Regs, i32imm, imm>;
798 // defm SELP_f16x2 : SELP_PATTERN<"b32", v2f16, Int32Regs, v2f16imm, imm>;
1523 // .b32, .b64, and .pred (predicate registers -- i.e., i1) versions of OpcStr.
1543 !strconcat(OpcStr, ".b32 \t$dst, $a, $b;"),
1547 !strconcat(OpcStr, ".b32 \t$dst, $a, $b;"),
1573 // Lower logical v2i16/v4i8 ops as bitwise ops on b32.
1599 "not.b32 \t$dst, $src;",
1648 "brev.b32 \t$dst, $a;",
1664 "shf.l.wrap.b32 \t$dst, $src, $src, $amt;",
1670 "shf.l.wrap.b32 \t$dst, $src, $src, $amt;",
1679 "shf.r.wrap.b32 \t$dst, $src, $src, $amt;",
1685 "shf.r.wrap.b32 \t$dst, $src, $src, $amt;",
1694 ".reg .b32 %lhs;\n\t"
1695 ".reg .b32 %rhs;\n\t"
1696 "shl.b32 \t%lhs, $src, $amt1;\n\t"
1697 "shr.b32 \t%rhs, $src, $amt2;\n\t"
1717 ".reg .b32 %lhs;\n\t"
1718 ".reg .b32 %rhs;\n\t"
1719 ".reg .b32 %amt2;\n\t"
1720 "shl.b32 \t%lhs, $src, $amt;\n\t"
1722 "shr.b32 \t%rhs, $src, %amt2;\n\t"
1732 ".reg .b32 %lhs;\n\t"
1733 ".reg .b32 %rhs;\n\t"
1734 ".reg .b32 %amt2;\n\t"
1735 "shr.b32 \t%lhs, $src, $amt;\n\t"
1737 "shl.b32 \t%rhs, $src, %amt2;\n\t"
1772 "and.b32 \t%amt2, $amt, 63;\n\t"
1786 "and.b32 \t%amt2, $amt, 63;\n\t"
1806 "shf.l.clamp.b32 \t$dst, $lo, $hi, $amt;",
1813 "shf.r.clamp.b32 \t$dst, $lo, $hi, $amt;",
1897 !strconcat("prmt.b32${mode}", " \t$d, $a, $b, $c;"),
1902 !strconcat("prmt.b32${mode}", " \t$d, $a, $b, $c;"),
1907 !strconcat("prmt.b32${mode}", " \t$d, $a, $b, $c;"),
1921 defm BFI_B32 : BFI<"bfi.b32", i32, Int32Regs, i32imm>;
1986 defm SETP_b32 : SETP<"b32", Int32Regs, i32imm>;
2041 defm SET_b32 : SET<"b32", Int32Regs, i32imm>;
2123 "mov.b32 \t$dst, $sss;", []>;
2152 "mov.b32 \t$dst, $src;", []>;
2762 def LoadParamMemI32 : LoadParamMemInst<Int32Regs, ".b32">;
2766 def LoadParamMemV2I32 : LoadParamV2MemInst<Int32Regs, ".b32">;
2769 def LoadParamMemV4I32 : LoadParamV4MemInst<Int32Regs, ".b32">;
2779 defm StoreParamI32 : StoreParamInst<Int32Regs, i32imm, ".b32">;
2787 defm StoreParamV2I32 : StoreParamV2Inst<Int32Regs, i32imm, ".b32">;
2791 defm StoreParamV4I32 : StoreParamV4Inst<Int32Regs, i32imm, ".b32">;
2804 def StoreRetvalI32 : StoreRetvalInst<Int32Regs, ".b32">;
2810 def StoreRetvalV2I32 : StoreRetvalV2Inst<Int32Regs, ".b32">;
2813 def StoreRetvalV4I32 : StoreRetvalV4Inst<Int32Regs, ".b32">;
2912 def MoveParamI32 : MoveParamInst<i32, Int32Regs, ".b32">;
2915 def MoveParamSymbolI32 : MoveParamSymbolInst<Int32Regs, i32imm, i32, ".b32">;
2942 def ProxyRegI32 : ProxyRegInst<"b32", i32, Int32Regs>;
3563 "mov.b32 \t$d, {{$s1, $s2}};", []>;
3581 "mov.b32 \t{{$d1, $d2}}, $s;", []>;
3594 "{{ .reg .b16 tmp; mov.b32 {tmp, $high}, $s; }}",
3598 "{{ .reg .b16 tmp; mov.b32 {$low, tmp}, $s; }}",
3602 "{{ .reg .b32 tmp; mov.b64 {tmp, $high}, $s; }}",
3606 "{{ .reg .b32 tmp; mov.b64 {$low, tmp}, $s; }}",
3644 "clz.b32 \t$d, $a;", []>;
3663 // TODO: NVPTX has a mov.b32 b32reg, {imm, b16reg} instruction, which we could
3666 // mov.b32 $tmp, {0xffff, $a}
3667 // ctlz.b32 $result, $tmp
3682 "popc.b32 \t$d, $a;", []>;