1 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 | FileCheck %s
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) {
14 ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
15 ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
16 ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
17 ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
18 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], [[C]], [[MASK]];
19 ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
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) {
26 ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
27 ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
28 ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
29 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], [[C]], 1;
30 ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
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) {
37 ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
38 ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
39 ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
40 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 1, [[MASK]];
41 ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
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) {
48 ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
49 ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
50 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 2, 1;
51 ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
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) {
58 ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
59 ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
60 ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
61 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 1, [[C]], [[MASK]];
62 ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
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) {
69 ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
70 ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
71 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, [[C]], 1;
72 ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
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) {
79 ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
80 ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
81 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 1, 2, [[MASK]];
82 ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
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) {
89 ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
90 ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, 3, 1;
91 ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
92 %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 2, i32 3)