Lines Matching refs:b32
147 def INT_NVVM_CLZ_I : F_MATH_1<"clz.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
152 def INT_NVVM_POPC_I : F_MATH_1<"popc.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
157 def INT_NVVM_PRMT : F_MATH_3<"prmt.b32 \t$dst, $src0, $src1, $src2;", Int32Regs,
284 def INT_NVVM_BREV32 : F_MATH_1<"brev.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
664 !strconcat(".reg .b32 %temp; \n\t",
669 !strconcat(".reg .b32 %temp; \n\t",
811 def INT_NVVM_BITCAST_F2I : F_MATH_1<"mov.b32 \t$dst, $src0;", Int32Regs,
813 def INT_NVVM_BITCAST_I2F : F_MATH_1<"mov.b32 \t$dst, $src0;", Float32Regs,
1040 defm INT_PTX_ATOM_SWAP_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".exch",
1042 defm INT_PTX_ATOM_SWAP_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".exch",
1044 defm INT_PTX_ATOM_SWAP_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".exch",
1046 defm INT_PTX_ATOM_SWAP_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1224 defm INT_PTX_ATOM_AND_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".and",
1226 defm INT_PTX_ATOM_AND_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".and",
1228 defm INT_PTX_ATOM_AND_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".and",
1230 defm INT_PTX_ATOM_AND_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1256 defm INT_PTX_ATOM_OR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".or",
1258 defm INT_PTX_ATOM_OR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".or",
1260 defm INT_PTX_ATOM_OR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1262 defm INT_PTX_ATOM_OR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".or",
1288 defm INT_PTX_ATOM_XOR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".xor",
1290 defm INT_PTX_ATOM_XOR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".xor",
1292 defm INT_PTX_ATOM_XOR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".xor",
1294 defm INT_PTX_ATOM_XOR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1320 defm INT_PTX_ATOM_CAS_G_32 : F_ATOMIC_3<Int32Regs, ".global", ".b32", ".cas",
1322 defm INT_PTX_ATOM_CAS_S_32 : F_ATOMIC_3<Int32Regs, ".shared", ".b32", ".cas",
1324 defm INT_PTX_ATOM_CAS_GEN_32 : F_ATOMIC_3<Int32Regs, "", ".b32", ".cas",
1326 defm INT_PTX_ATOM_CAS_GEN_32_USE_G : F_ATOMIC_3<Int32Regs, ".global", ".b32",
1660 "mov.b32 \t$r, $s;",
1777 "mov.b32\t$d, $r;", []>;
1818 "shf.l.wrap.b32 \t$dst, $src, $src, $amt;",
1826 "shf.l.wrap.b32 \t$dst, $src, $src, $amt;",
1842 !strconcat(".reg .b32 %dummy;\n\t",
1850 !strconcat(".reg .b32 %dummy;\n\t",
1867 "shf.l.wrap.b32 \t$dst, $lo, $hi, $amt;",[]>,
1873 "shf.l.wrap.b32 \t$dst, $lo, $hi, $amt;",[]>,
1879 "shf.r.wrap.b32 \t$dst, $lo, $hi, $amt;",[]>,
1885 "shf.r.wrap.b32 \t$dst, $lo, $hi, $amt;",[]>,
3327 "suld.b.1d.b32.clamp \\{$r\\}, [$s, \\{$x\\}];",
3348 "suld.b.a1d.b32.clamp \\{$r\\}, [$s, \\{$l, $x\\}];",
3369 "suld.b.2d.b32.clamp \\{$r\\}, [$s, \\{$x, $y\\}];",
3390 "suld.b.a2d.b32.clamp \\{$r\\}, [$s, \\{$l, $x, $y, $y\\}];",
3411 "suld.b.3d.b32.clamp \\{$r\\}, [$s, \\{$x, $y, $z, $z\\}];",
3434 "suld.b.1d.v2.b32.clamp \\{$r, $g\\}, [$s, \\{$x\\}];",
3455 "suld.b.a1d.v2.b32.clamp \\{$r, $g\\}, [$s, \\{$l, $x\\}];",
3476 "suld.b.2d.v2.b32.clamp \\{$r, $g\\}, [$s, \\{$x, $y\\}];",
3499 "suld.b.a2d.v2.b32.clamp \\{$r, $g\\}, "
3522 "suld.b.3d.v2.b32.clamp \\{$r, $g\\}, [$s, \\{$x, $y, $z, $z\\}];",
3545 "suld.b.1d.v4.b32.clamp \\{$r, $g, $b, $a\\}, [$s, \\{$x\\}];",
3563 "suld.b.a1d.v4.b32.clamp \\{$r, $g, $b, $a\\}, "
3580 "suld.b.2d.v4.b32.clamp \\{$r, $g, $b, $a\\}, [$s, \\{$x, $y\\}];",
3598 "suld.b.a2d.v4.b32.clamp \\{$r, $g, $b, $a\\}, "
3618 "suld.b.3d.v4.b32.clamp \\{$r, $g, $b, $a\\}, "
3639 "suld.b.1d.b32.trap \\{$r\\}, [$s, \\{$x\\}];",
3660 "suld.b.a1d.b32.trap \\{$r\\}, [$s, \\{$l, $x\\}];",
3681 "suld.b.2d.b32.trap \\{$r\\}, [$s, \\{$x, $y\\}];",
3702 "suld.b.a2d.b32.trap \\{$r\\}, [$s, \\{$l, $x, $y, $y\\}];",
3723 "suld.b.3d.b32.trap \\{$r\\}, [$s, \\{$x, $y, $z, $z\\}];",
3746 "suld.b.1d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$x\\}];",
3767 "suld.b.a1d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$l, $x\\}];",
3788 "suld.b.2d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$x, $y\\}];",
3811 "suld.b.a2d.v2.b32.trap \\{$r, $g\\}, "
3834 "suld.b.3d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$x, $y, $z, $z\\}];",
3857 "suld.b.1d.v4.b32.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x\\}];",
3875 "suld.b.a1d.v4.b32.trap \\{$r, $g, $b, $a\\}, "
3892 "suld.b.2d.v4.b32.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x, $y\\}];",
3910 "suld.b.a2d.v4.b32.trap \\{$r, $g, $b, $a\\}, "
3930 "suld.b.3d.v4.b32.trap \\{$r, $g, $b, $a\\}, "
3950 "suld.b.1d.b32.zero \\{$r\\}, [$s, \\{$x\\}];",
3971 "suld.b.a1d.b32.zero \\{$r\\}, [$s, \\{$l, $x\\}];",
3992 "suld.b.2d.b32.zero \\{$r\\}, [$s, \\{$x, $y\\}];",
4013 "suld.b.a2d.b32.zero \\{$r\\}, [$s, \\{$l, $x, $y, $y\\}];",
4034 "suld.b.3d.b32.zero \\{$r\\}, [$s, \\{$x, $y, $z, $z\\}];",
4057 "suld.b.1d.v2.b32.zero \\{$r, $g\\}, [$s, \\{$x\\}];",
4078 "suld.b.a1d.v2.b32.zero \\{$r, $g\\}, [$s, \\{$l, $x\\}];",
4099 "suld.b.2d.v2.b32.zero \\{$r, $g\\}, [$s, \\{$x, $y\\}];",
4122 "suld.b.a2d.v2.b32.zero \\{$r, $g\\}, "
4145 "suld.b.3d.v2.b32.zero \\{$r, $g\\}, [$s, \\{$x, $y, $z, $z\\}];",
4168 "suld.b.1d.v4.b32.zero \\{$r, $g, $b, $a\\}, [$s, \\{$x\\}];",
4186 "suld.b.a1d.v4.b32.zero \\{$r, $g, $b, $a\\}, "
4203 "suld.b.2d.v4.b32.zero \\{$r, $g, $b, $a\\}, [$s, \\{$x, $y\\}];",
4221 "suld.b.a2d.v4.b32.zero \\{$r, $g, $b, $a\\}, "
4241 "suld.b.3d.v4.b32.zero \\{$r, $g, $b, $a\\}, "
4253 "txq.channel_order.b32 \t$d, [$a];",
4257 "txq.channel_data_type.b32 \t$d, [$a];",
4261 "txq.width.b32 \t$d, [$a];",
4265 "txq.height.b32 \t$d, [$a];",
4269 "txq.depth.b32 \t$d, [$a];",
4273 "txq.array_size.b32 \t$d, [$a];",
4277 "txq.num_samples.b32 \t$d, [$a];",
4281 "txq.num_mipmap_levels.b32 \t$d, [$a];",
4310 "suq.channel_order.b32 \t$d, [$a];",
4314 "suq.channel_data_type.b32 \t$d, [$a];",
4318 "suq.width.b32 \t$d, [$a];",
4322 "suq.height.b32 \t$d, [$a];",
4326 "suq.depth.b32 \t$d, [$a];",
4330 "suq.array_size.b32 \t$d, [$a];",
4382 "sust.b.1d.b32.clamp \t[$s, \\{$x\\}], \\{$r\\};",
4402 "sust.b.1d.v2.b32.clamp \t[$s, \\{$x\\}], \\{$r, $g\\};",
4425 "sust.b.1d.v4.b32.clamp \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
4442 "sust.b.a1d.b32.clamp \t[$s, \\{$idx, $x\\}], \\{$r\\};",
4465 "sust.b.a1d.v2.b32.clamp \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
4491 "sust.b.a1d.v4.b32.clamp \t[$s, \\{$idx, $x\\}], "
4509 "sust.b.2d.b32.clamp \t[$s, \\{$x, $y\\}], \\{$r\\};",
4532 "sust.b.2d.v2.b32.clamp \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
4558 "sust.b.2d.v4.b32.clamp \t[$s, \\{$x, $y\\}], "
4579 "sust.b.a2d.b32.clamp \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
4605 "sust.b.a2d.v2.b32.clamp \t[$s, \\{$idx, $x, $y, $y\\}], "
4633 "sust.b.a2d.v4.b32.clamp \t[$s, \\{$idx, $x, $y, $y\\}], "
4654 "sust.b.3d.b32.clamp \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
4680 "sust.b.3d.v2.b32.clamp \t[$s, \\{$x, $y, $z, $z\\}], "
4708 "sust.b.3d.v4.b32.clamp \t[$s, \\{$x, $y, $z, $z\\}], "
4727 "sust.b.1d.b32.trap \t[$s, \\{$x\\}], \\{$r\\};",
4747 "sust.b.1d.v2.b32.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
4770 "sust.b.1d.v4.b32.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
4787 "sust.b.a1d.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
4810 "sust.b.a1d.v2.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
4836 "sust.b.a1d.v4.b32.trap \t[$s, \\{$idx, $x\\}], "
4854 "sust.b.2d.b32.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
4877 "sust.b.2d.v2.b32.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
4903 "sust.b.2d.v4.b32.trap \t[$s, \\{$x, $y\\}], "
4924 "sust.b.a2d.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
4950 "sust.b.a2d.v2.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
4978 "sust.b.a2d.v4.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
4999 "sust.b.3d.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
5025 "sust.b.3d.v2.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
5053 "sust.b.3d.v4.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
5072 "sust.b.1d.b32.zero \t[$s, \\{$x\\}], \\{$r\\};",
5092 "sust.b.1d.v2.b32.zero \t[$s, \\{$x\\}], \\{$r, $g\\};",
5115 "sust.b.1d.v4.b32.zero \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
5132 "sust.b.a1d.b32.zero \t[$s, \\{$idx, $x\\}], \\{$r\\};",
5155 "sust.b.a1d.v2.b32.zero \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
5181 "sust.b.a1d.v4.b32.zero \t[$s, \\{$idx, $x\\}], "
5199 "sust.b.2d.b32.zero \t[$s, \\{$x, $y\\}], \\{$r\\};",
5222 "sust.b.2d.v2.b32.zero \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
5248 "sust.b.2d.v4.b32.zero \t[$s, \\{$x, $y\\}], "
5269 "sust.b.a2d.b32.zero \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
5295 "sust.b.a2d.v2.b32.zero \t[$s, \\{$idx, $x, $y, $y\\}], "
5323 "sust.b.a2d.v4.b32.zero \t[$s, \\{$idx, $x, $y, $y\\}], "
5344 "sust.b.3d.b32.zero \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
5370 "sust.b.3d.v2.b32.zero \t[$s, \\{$x, $y, $z, $z\\}], "
5398 "sust.b.3d.v4.b32.zero \t[$s, \\{$x, $y, $z, $z\\}], "
5419 "sust.p.1d.b32.trap \t[$s, \\{$x\\}], \\{$r\\};",
5434 "sust.p.1d.v2.b32.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
5452 "sust.p.1d.v4.b32.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
5469 "sust.p.a1d.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
5487 "sust.p.a1d.v2.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
5507 "sust.p.a1d.v4.b32.trap \t[$s, \\{$idx, $x\\}], "
5525 "sust.p.2d.b32.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
5543 "sust.p.2d.v2.b32.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
5563 "sust.p.2d.v4.b32.trap \t[$s, \\{$x, $y\\}], "
5584 "sust.p.a2d.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
5604 "sust.p.a2d.v2.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
5625 "sust.p.a2d.v4.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
5646 "sust.p.3d.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
5666 "sust.p.3d.v2.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
5687 "sust.p.3d.v4.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "