Lines Matching refs:shfl

3 declare i32 @llvm.nvvm.shfl.sync.down.i32(i32, i32, i32, i32)
4 declare float @llvm.nvvm.shfl.sync.down.f32(float, i32, i32, i32)
5 declare i32 @llvm.nvvm.shfl.sync.up.i32(i32, i32, i32, i32)
6 declare float @llvm.nvvm.shfl.sync.up.f32(float, i32, i32, i32)
7 declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32)
8 declare float @llvm.nvvm.shfl.sync.bfly.f32(float, i32, i32, i32)
9 declare i32 @llvm.nvvm.shfl.sync.idx.i32(i32, i32, i32, i32)
10 declare float @llvm.nvvm.shfl.sync.idx.f32(float, i32, i32, i32)
12 ; CHECK-LABEL: .func{{.*}}shfl.sync.rrr
13 define i32 @shfl.sync.rrr(i32 %mask, i32 %a, i32 %b, i32 %c) {
18   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], [[C]], [[MASK]];
20   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 %b, i32 %c)
24 ; CHECK-LABEL: .func{{.*}}shfl.sync.irr
25 define i32 @shfl.sync.irr(i32 %a, i32 %b, i32 %c) {
29   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], [[C]], 1;
31   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 %b, i32 %c)
35 ; CHECK-LABEL: .func{{.*}}shfl.sync.rri
36 define i32 @shfl.sync.rri(i32 %mask, i32 %a, i32 %b) {
40   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 1, [[MASK]];
42   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 %b, i32 1)
46 ; CHECK-LABEL: .func{{.*}}shfl.sync.iri
47 define i32 @shfl.sync.iri(i32 %a, i32 %b) {
50   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 2, 1;
52   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 %b, i32 2)
56 ; CHECK-LABEL: .func{{.*}}shfl.sync.rir
57 define i32 @shfl.sync.rir(i32 %mask, i32 %a, i32 %c) {
61   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 1, [[C]], [[MASK]];
63   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 1, i32 %c)
67 ; CHECK-LABEL: .func{{.*}}shfl.sync.iir
68 define i32 @shfl.sync.iir(i32 %a, i32 %c) {
71   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, [[C]], 1;
73   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 2, i32 %c)
77 ; CHECK-LABEL: .func{{.*}}shfl.sync.rii
78 define i32 @shfl.sync.rii(i32 %mask, i32 %a) {
81   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 1, 2, [[MASK]];
83   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 1, i32 2)
87 ; CHECK-LABEL: .func{{.*}}shfl.sync.iii
88 define i32 @shfl.sync.iii(i32 %a, i32 %b) {
90   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, 3, 1;
92   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 2, i32 3)