Lines Matching refs:shfl
3 declare {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32, i32, i32, i32)
4 declare {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32, float, i32, i32)
5 declare {i32, i1} @llvm.nvvm.shfl.sync.up.i32p(i32, i32, i32, i32)
6 declare {float, i1} @llvm.nvvm.shfl.sync.up.f32p(i32, float, i32, i32)
7 declare {i32, i1} @llvm.nvvm.shfl.sync.bfly.i32p(i32, i32, i32, i32)
8 declare {float, i1} @llvm.nvvm.shfl.sync.bfly.f32p(i32, float, i32, i32)
9 declare {i32, i1} @llvm.nvvm.shfl.sync.idx.i32p(i32, i32, i32, i32)
10 declare {float, i1} @llvm.nvvm.shfl.sync.idx.f32p(i32, float, i32, i32)
12 ; CHECK-LABEL: .func{{.*}}shfl.sync.i32.rrr
13 define {i32, i1} @shfl.sync.i32.rrr(i32 %mask, i32 %a, i32 %b, i32 %c) {
18 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]], [[MASK]];
20 %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 %mask, i32 %a, i32 %b, i32 %c)
24 ; CHECK-LABEL: .func{{.*}}shfl.sync.i32.irr
25 define {i32, i1} @shfl.sync.i32.irr(i32 %a, i32 %b, i32 %c) {
29 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]], 1;
31 %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 1, i32 %a, i32 %b, i32 %c)
35 ; CHECK-LABEL: .func{{.*}}shfl.sync.i32.rri
36 define {i32, i1} @shfl.sync.i32.rri(i32 %mask, i32 %a, i32 %b) {
40 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 1, [[MASK]];
42 %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 %mask, i32 %a, i32 %b, i32 1)
46 ; CHECK-LABEL: .func{{.*}}shfl.sync.i32.iri
47 define {i32, i1} @shfl.sync.i32.iri(i32 %a, i32 %b) {
50 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 2, 1;
52 %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 1, i32 %a, i32 %b, i32 2)
56 ; CHECK-LABEL: .func{{.*}}shfl.sync.i32.rir
57 define {i32, i1} @shfl.sync.i32.rir(i32 %mask, i32 %a, i32 %c) {
61 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, [[C]], [[MASK]];
63 %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 %mask, i32 %a, i32 1, i32 %c)
67 ; CHECK-LABEL: .func{{.*}}shfl.sync.i32.iir
68 define {i32, i1} @shfl.sync.i32.iir(i32 %a, i32 %c) {
71 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, [[C]], 1;
73 %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 1, i32 %a, i32 2, i32 %c)
77 ; CHECK-LABEL: .func{{.*}}shfl.sync.i32.rii
78 define {i32, i1} @shfl.sync.i32.rii(i32 %mask, i32 %a) {
81 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, 2, [[MASK]];
83 %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 %mask, i32 %a, i32 1, i32 2)
87 ; CHECK-LABEL: .func{{.*}}shfl.sync.i32.iii
88 define {i32, i1} @shfl.sync.i32.iii(i32 %a, i32 %b) {
90 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, 3, 1;
92 %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 1, i32 %a, i32 2, i32 3)
98 ; CHECK-LABEL: .func{{.*}}shfl.sync.f32.rrr
99 define {float, i1} @shfl.sync.f32.rrr(i32 %mask, float %a, i32 %b, i32 %c) {
104 ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]], [[MASK]];
106 %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 %mask, float %a, i32 %b, i32 %c)
110 ; CHECK-LABEL: .func{{.*}}shfl.sync.f32.irr
111 define {float, i1} @shfl.sync.f32.irr(float %a, i32 %b, i32 %c) {
115 ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]], 1;
117 %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 1, float %a, i32 %b, i32 %c)
121 ; CHECK-LABEL: .func{{.*}}shfl.sync.f32.rri
122 define {float, i1} @shfl.sync.f32.rri(i32 %mask, float %a, i32 %b) {
126 ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 1, [[MASK]];
128 %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 %mask, float %a, i32 %b, i32 1)
132 ; CHECK-LABEL: .func{{.*}}shfl.sync.f32.iri
133 define {float, i1} @shfl.sync.f32.iri(float %a, i32 %b) {
136 ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 2, 1;
138 %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 1, float %a, i32 %b, i32 2)
142 ; CHECK-LABEL: .func{{.*}}shfl.sync.f32.rir
143 define {float, i1} @shfl.sync.f32.rir(i32 %mask, float %a, i32 %c) {
147 ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, [[C]], [[MASK]];
149 %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 %mask, float %a, i32 1, i32 %c)
153 ; CHECK-LABEL: .func{{.*}}shfl.sync.f32.iir
154 define {float, i1} @shfl.sync.f32.iir(float %a, i32 %c) {
157 ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, [[C]], 1;
159 %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 1, float %a, i32 2, i32 %c)
163 ; CHECK-LABEL: .func{{.*}}shfl.sync.f32.rii
164 define {float, i1} @shfl.sync.f32.rii(i32 %mask, float %a) {
167 ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, 2, [[MASK]];
169 %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 %mask, float %a, i32 1, i32 2)
173 ; CHECK-LABEL: .func{{.*}}shfl.sync.f32.iii
174 define {float, i1} @shfl.sync.f32.iii(float %a, i32 %b) {
176 ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, 3, 1;
178 %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 1, float %a, i32 2, i32 3)