Home
last modified time | relevance | path

Searched refs:b32 (Results 1 – 25 of 141) sorted by relevance

123456

/llvm-project-15.0.7/llvm/test/CodeGen/NVPTX/
H A Df16x2-instructions.ll35 ; CHECK: mov.b32 [[R:%hh[0-9+]]], [[T]];
337 ; CHECK-DAG: st.param.b32 [param0+0], [[A]];
338 ; CHECK-DAG: st.param.b32 [param1+0], [[B]];
358 ; CHECK-DAG: st.param.b32 [param0+0], [[B]];
359 ; CHECK-DAG: st.param.b32 [param1+0], [[A]];
1023 ; CHECK: mov.b32 [[R:%hh[0-9]+]], [[AF1]];
1033 ; CHECK: mov.b32 [[AF1:%f[0-9]+]], [[R]];
1473 ; CHECK: mov.b32 {%h1, %h2}, %hh1;
1474 ; CHECK: mov.b32 %hh2, {%h2, %h1};
1481 ; CHECK: mov.b32 {%h2, %tmp_hi}, %hh1;
[all …]
H A Dparam-load-store.ll32 ; CHECK: .param .b32 param0;
34 ; CHECK: .param .b32 retval0;
55 ; CHECK: .param .b32 param0;
57 ; CHECK: .param .b32 retval0;
139 ; CHECK: .param .b32 param0;
141 ; CHECK: .param .b32 retval0;
156 ; CHECK: .param .b32 param0;
158 ; CHECK: .param .b32 retval0;
176 ; CHECK: .param .b32 param0;
196 ; CHECK: .param .b32 param0;
[all …]
H A Denvreg.ll41 ; CHECK: mov.b32 %r{{[0-9]+}}, %envreg0
43 ; CHECK: mov.b32 %r{{[0-9]+}}, %envreg1
45 ; CHECK: mov.b32 %r{{[0-9]+}}, %envreg2
47 ; CHECK: mov.b32 %r{{[0-9]+}}, %envreg3
49 ; CHECK: mov.b32 %r{{[0-9]+}}, %envreg4
51 ; CHECK: mov.b32 %r{{[0-9]+}}, %envreg5
53 ; CHECK: mov.b32 %r{{[0-9]+}}, %envreg6
55 ; CHECK: mov.b32 %r{{[0-9]+}}, %envreg7
57 ; CHECK: mov.b32 %r{{[0-9]+}}, %envreg8
59 ; CHECK: mov.b32 %r{{[0-9]+}}, %envreg9
[all …]
H A Dproxy-reg-erasure-ptx.ll15 ; PTX-DAG: ld.param.b32 [[LD:%r[0-9]+]], [retval0+0];
18 ; PTX-WITHOUT-DAG: mov.b32 [[PROXY:%r[0-9]+]], [[LD]];
19 ; PTX-WITHOUT-DAG: and.b32 [[RES:%r[0-9]+]], [[PROXY]], 1;
20 ; PTX-WITH-DAG: and.b32 [[RES:%r[0-9]+]], [[LD]], 1;
22 ; PTX-DAG: st.param.b32 [func_retval0+0], [[RES]];
32 ; PTX-DAG: ld.param.b32 [[LD:%r[0-9]+]], [retval0+0];
35 ; PTX-WITHOUT-DAG: mov.b32 [[PROXY:%r[0-9]+]], [[LD]];
39 ; PTX-DAG: st.param.b32 [func_retval0+0], [[RES]];
49 ; PTX-DAG: ld.param.b32 [[LD:%r[0-9]+]], [retval0+0];
52 ; PTX-WITHOUT-DAG: mov.b32 [[PROXY:%r[0-9]+]], [[LD]];
[all …]
H A Drotate.ll7 declare i32 @llvm.nvvm.rotate.b32(i32, i32)
14 ; SM20: shl.b32
16 ; SM20: shr.b32
18 ; SM35: shf.l.wrap.b32
30 ; SM35: shf.l.wrap.b32
31 ; SM35: shf.l.wrap.b32
43 ; SM35: shf.r.wrap.b32
44 ; SM35: shf.r.wrap.b32
52 ; SM20: shl.b32
53 ; SM20: shr.b32
[all …]
H A Datomicrmw-expand.ll26 ; ALL: atom.and.b32
28 ; ALL: atom.or.b32
30 ; ALL: atom.xor.b32
32 ; ALL: atom.exch.b32
90 ; ALL: atom.and.b32
92 ; ALL: atom.or.b32
94 ; ALL: atom.xor.b32
96 ; ALL: atom.cas.b32
104 ; ALL: atom.cas.b32
106 ; ALL: atom.cas.b32
[all …]
H A Dtexsurf-queries.ll21 ; SM20: txq.width.b32
22 ; SM30: txq.width.b32
32 ; SM20: txq.width.b32 %r{{[0-9]+}}, [tex0]
42 ; SM20: txq.height.b32
43 ; SM30: txq.height.b32
53 ; SM20: txq.height.b32 %r{{[0-9]+}}, [tex0]
63 ; SM20: suq.width.b32
64 ; SM30: suq.width.b32
74 ; SM20: suq.width.b32 %r{{[0-9]+}}, [surf0]
84 ; SM20: suq.height.b32
[all …]
H A Dstore-retval.ll27 ; CHECK-NOT: st.param.b32 [func_retval0+0], %r{{[0-9]+}};
28 ; CHECK-NOT: st.param.b32 [func_retval0+4], %r{{[0-9]+}};
29 ; CHECK-NOT: st.param.b32 [func_retval0+8], %r{{[0-9]+}};
30 ; CHECK-NOT: st.param.b32 [func_retval0+12], %r{{[0-9]+}};
31 ; CHECK: st.param.b32 [func_retval0+16], [[R1]];
44 ; CHECK-NOT: st.param.b32 [func_retval0+0], %r{{[0-9]+}};
45 ; CHECK-NOT: st.param.b32 [func_retval0+4], %r{{[0-9]+}};
46 ; CHECK-NOT: st.param.b32 [func_retval0+8], %r{{[0-9]+}};
47 ; CHECK-NOT: st.param.b32 [func_retval0+12], %r{{[0-9]+}};
48 ; CHECK: st.param.b32 [func_retval0+16], [[R1]];
[all …]
H A Dfns.ll12 ; CHECK: fns.b32 {{%r[0-9]+}}, [[MASK]], [[BASE]], [[OFFSET]];
14 ; CHECK: fns.b32 {{%r[0-9]+}}, [[MASK]], [[BASE]], 0;
17 ; CHECK: fns.b32 {{%r[0-9]+}}, [[MASK]], 1, [[OFFSET]];
19 ; CHECK: fns.b32 {{%r[0-9]+}}, [[MASK]], 1, 0;
23 ; CHECK: fns.b32 {{%r[0-9]+}}, 2, [[BASE]], [[OFFSET]];
25 ; CHECK: fns.b32 {{%r[0-9]+}}, 2, [[BASE]], 0;
28 ; CHECK: fns.b32 {{%r[0-9]+}}, 2, 1, [[OFFSET]];
30 ; CHECK: fns.b32 {{%r[0-9]+}}, 2, 1, 0;
H A Dparam-vectorize-device.ll84 ; CHECK: .param .b32 caller_St4x1_param_1
86 ; CHECK: .param .b32 param0;
87 ; CHECK: st.param.b32 [param0+0], {{%r[0-9]+}};
94 ; CHECK: ld.param.b32 {{%r[0-9]+}}, [retval0+0];
105 ; CHECK-NEXT: .param .b32 callee_St4x1_param_0
107 ; CHECK: st.param.b32 [func_retval0+0], [[R1]];
116 ; CHECK: .param .b32 caller_St4x2_param_1
157 ; CHECK: .param .b32 caller_St4x3_param_1
161 ; CHECK: st.param.b32 [param0+8], {{%r[0-9]+}};
205 ; CHECK: .param .b32 caller_St4x4_param_1
[all …]
H A Dshfl-p.ll18 ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]];
39 ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 1;
49 ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 2;
59 ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, [[C]];
69 ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, [[C]];
78 ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, 2;
87 ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, 3;
121 ; CHECK: shfl.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 1;
131 ; CHECK: shfl.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 2;
160 ; CHECK: shfl.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, 2;
[all …]
H A Dshfl-sync-p.ll30 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]], 1;
41 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 1, [[MASK]];
51 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 2, 1;
72 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, [[C]], 1;
82 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, 2, [[MASK]];
91 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, 3, 1;
116 ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]], 1;
137 ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 2, 1;
158 ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, [[C]], 1;
168 ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, 2, [[MASK]];
[all …]
H A Dmodule-inline-asm.ll7 ; CHECK-NOT: .global .b32 val;
12 ; CHECK: .global .b32 val;
13 module asm ".global .b32 val;"
21 ; CHECK-NOT: .global .b32 val;
H A Dshfl.ll19 ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]], [[IN]], 1, 2;
57 ; CHECK: shfl.down.b32 [[OUT:%f[0-9]+]], [[IN]], 5, 6;
66 ; CHECK: shfl.up.b32 %r{{[0-9]+}}, %r{{[0-9]+}}, 1, 2;
70 ; CHECK: shfl.up.b32 %f{{[0-9]+}}, %f{{[0-9]+}}, 3, 4;
74 ; CHECK: shfl.bfly.b32 %r{{[0-9]+}}, %r{{[0-9]+}}, 5, 6;
78 ; CHECK: shfl.bfly.b32 %f{{[0-9]+}}, %f{{[0-9]+}}, 7, 8;
82 ; CHECK: shfl.idx.b32 %r{{[0-9]+}}, %r{{[0-9]+}}, 9, 10;
86 ; CHECK: shfl.idx.b32 %f{{[0-9]+}}, %f{{[0-9]+}}, 11, 12;
H A Dshfl-sync.ll19 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], [[C]], [[MASK]];
30 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], [[C]], 1;
41 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 1, [[MASK]];
51 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 2, 1;
62 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 1, [[C]], [[MASK]];
72 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, [[C]], 1;
82 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 1, 2, [[MASK]];
91 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, 3, 1;
H A Dcttz.ll11 ; CHECK: popc.b32
17 ; CHECK: popc.b32
30 ; CHECK: popc.b32
36 ; CHECK: popc.b32
H A Di8-param.ll6 ; CHECK: .visible .func (.param .b32 func_retval0) callee
10 ; CHECK: st.param.b32
19 ; CHECK: ld.param.b32
H A Dmatch.ll12 ; CHECK: match.any.sync.b32 [[V0:%r[0-9]+]], [[VALUE]], [[MASK]];
14 ; CHECK: match.any.sync.b32 [[V1:%r[0-9]+]], [[VALUE]], 1;
16 ; CHECK: match.any.sync.b32 [[V2:%r[0-9]+]], 2, [[MASK]];
18 ; CHECK: match.any.sync.b32 [[V3:%r[0-9]+]], 4, 3;
53 ; CHECK: match.all.sync.b32 {{%r[0-9]+\|%p[0-9]+}}, [[VALUE]], [[MASK]];
58 ; CHECK: match.all.sync.b32 {{%r[0-9]+\|%p[0-9]+}}, 1, [[MASK]];
63 ; CHECK: match.all.sync.b32 {{%r[0-9]+\|%p[0-9]+}}, [[VALUE]], 2;
68 ; CHECK: match.all.sync.b32 {{%r[0-9]+\|%p[0-9]+}}, 4, 3;
/llvm-project-15.0.7/llvm/test/DebugInfo/NVPTX/
H A Ddebug-info.ll24 ; CHECK-DAG: .reg .b32 %r<6>;
1613 ; CHECK-NEXT:.b32 1508 // DW_AT_type
1616 ; CHECK-NEXT:.b32 1508 // DW_AT_type
1653 ; CHECK-NEXT:.b32 1554 // DW_AT_type
1656 ; CHECK-NEXT:.b32 1554 // DW_AT_type
1687 ; CHECK-NEXT:.b32 1554 // DW_AT_type
1690 ; CHECK-NEXT:.b32 1554 // DW_AT_type
1710 ; CHECK-NEXT:.b32 1554 // DW_AT_type
1713 ; CHECK-NEXT:.b32 1554 // DW_AT_type
1735 ; CHECK-NEXT:.b32 1554 // DW_AT_type
[all …]
H A Dcu-range-hole.ll6 ; CHECK: .visible .func (.param .b32 func_retval0) b(
7 ; CHECK: .param .b32 b_param_0
18 ; CHECK: .visible .func (.param .b32 func_retval0) a(
19 ; CHECK: .param .b32 a_param_0
28 ; CHECK: .visible .func (.param .b32 func_retval0) d(
29 ; CHECK: .param .b32 d_param_0
148 ; CHECK-NEXT: .b32 183 // Length of Unit
236 ; CHECK-NEXT: .b32 179 // DW_AT_type
243 ; CHECK-NEXT: .b32 179 // DW_AT_type
255 ; CHECK-NEXT: .b32 179 // DW_AT_type
[all …]
/llvm-project-15.0.7/llvm/test/CodeGen/Hexagon/
H A Dopt-addr-mode-subreg-use.ll91 br i1 %v31, label %b6, label %b32
97 br i1 undef, label %b7, label %b32
100 br i1 undef, label %b8, label %b32
103 br i1 undef, label %b9, label %b32
106 br i1 undef, label %b10, label %b32
110 br i1 %v32, label %b11, label %b32
120 br i1 %v33, label %b14, label %b32
136 br i1 undef, label %b19, label %b32
155 br i1 %v35, label %b25, label %b32
173 br i1 %v38, label %b30, label %b32
[all …]
H A Dmulti-cycle.ll46 br label %b32
48 b32: ; preds = %b32, %b18
49 %v33 = phi i32 [ 0, %b18 ], [ %v63, %b32 ]
50 %v34 = phi <16 x i32>* [ %v31, %b18 ], [ %v62, %b32 ]
51 %v35 = phi <16 x i32>* [ %v28, %b18 ], [ %v46, %b32 ]
52 %v36 = phi <16 x i32>* [ %v29, %b18 ], [ %v44, %b32 ]
54 %v38 = phi <16 x i32> [ %v15, %b18 ], [ %v39, %b32 ]
55 %v39 = phi <16 x i32> [ %v26, %b18 ], [ %v56, %b32 ]
56 %v40 = phi <16 x i32> [ %v27, %b18 ], [ %v51, %b32 ]
57 %v41 = phi <16 x i32> [ %v15, %b18 ], [ %v40, %b32 ]
[all …]
/llvm-project-15.0.7/llvm/test/CodeGen/AMDGPU/
H A Dimmv216.ll10 ; GCN: buffer_store_{{dword|b32}} [[REG]]
18 ; GCN: buffer_store_{{dword|b32}} [[REG]]
26 ; GCN: buffer_store_{{dword|b32}} [[REG]]
34 ; GCN: buffer_store_{{dword|b32}} [[REG]]
42 ; GCN: buffer_store_{{dword|b32}} [[REG]]
50 ; GCN: buffer_store_{{dword|b32}} [[REG]]
58 ; GCN: buffer_store_{{dword|b32}} [[REG]]
66 ; GCN: buffer_store_{{dword|b32}} [[REG]]
74 ; GCN: buffer_store_{{dword|b32}} [[REG]]
82 ; GCN: buffer_store_{{dword|b32}} [[REG]]
[all …]
H A Dllvm.amdgcn.workitem.id.ll22 ; ALL: {{buffer|flat|global}}_store_{{dword|b32}} {{.*}}v0
41 ; PACKED-TID: {{buffer|flat|global}}_store_{{dword|b32}} {{.*}}[[ID]]
59 ; PACKED-TID: {{buffer|flat|global}}_store_{{dword|b32}} {{.*}}[[ID]]
77 ; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
78 ; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
93 ; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
100 ; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
115 ; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
116 ; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
/llvm-project-15.0.7/llvm/test/CodeGen/AMDGPU/GlobalISel/
H A Dllvm.amdgcn.workitem.id.ll22 ; ALL: {{buffer|flat|global}}_store_{{dword|b32}} {{.*}}v0
69 ; ALL: {{flat|global}}_store_{{dword|b32}} v{{.*}}, v0
71 ; ALL: {{flat|global}}_store_{{dword|b32}} v{{.*}}, v0
83 ; ALL: {{flat|global}}_store_{{dword|b32}}
85 ; ALL: {{flat|global}}_store_{{dword|b32}} v{{.*}}, v0
140 ; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
141 ; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
156 ; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
163 ; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
178 ; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
[all …]

123456