Lines Matching refs:sync
4 declare i32 @llvm.nvvm.shfl.sync.down.i32(i32, i32, i32, i32)
5 declare float @llvm.nvvm.shfl.sync.down.f32(float, i32, i32, i32)
6 declare i32 @llvm.nvvm.shfl.sync.up.i32(i32, i32, i32, i32)
7 declare float @llvm.nvvm.shfl.sync.up.f32(float, i32, i32, i32)
8 declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32)
9 declare float @llvm.nvvm.shfl.sync.bfly.f32(float, i32, i32, i32)
10 declare i32 @llvm.nvvm.shfl.sync.idx.i32(i32, i32, i32, i32)
11 declare float @llvm.nvvm.shfl.sync.idx.f32(float, i32, i32, i32)
19 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], [[C]], [[MASK]];
21 %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 %b, i32 %c)
30 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], [[C]], 1;
32 %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 %b, i32 %c)
41 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 1, [[MASK]];
43 %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 %b, i32 1)
51 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 2, 1;
53 %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 %b, i32 2)
62 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 1, [[C]], [[MASK]];
64 %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 1, i32 %c)
72 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, [[C]], 1;
74 %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 2, i32 %c)
82 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 1, 2, [[MASK]];
84 %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 1, i32 2)
91 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, 3, 1;
93 %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 2, i32 3)