1; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | FileCheck %s
2; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | %ptxas-verify %if !ptxas-11.0 %{-arch=sm_30%} %}
3
4declare i32 @llvm.nvvm.shfl.sync.down.i32(i32, i32, i32, i32)
5declare float @llvm.nvvm.shfl.sync.down.f32(float, i32, i32, i32)
6declare i32 @llvm.nvvm.shfl.sync.up.i32(i32, i32, i32, i32)
7declare float @llvm.nvvm.shfl.sync.up.f32(float, i32, i32, i32)
8declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32)
9declare float @llvm.nvvm.shfl.sync.bfly.f32(float, i32, i32, i32)
10declare i32 @llvm.nvvm.shfl.sync.idx.i32(i32, i32, i32, i32)
11declare float @llvm.nvvm.shfl.sync.idx.f32(float, i32, i32, i32)
12
13; CHECK-LABEL: .func{{.*}}shfl_sync_rrr
14define i32 @shfl_sync_rrr(i32 %mask, i32 %a, i32 %b, i32 %c) {
15  ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
16  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
17  ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
18  ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
19  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], [[C]], [[MASK]];
20  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
21  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 %b, i32 %c)
22  ret i32 %val
23}
24
25; CHECK-LABEL: .func{{.*}}shfl_sync_irr
26define i32 @shfl_sync_irr(i32 %a, i32 %b, i32 %c) {
27  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
28  ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
29  ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
30  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], [[C]], 1;
31  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
32  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 %b, i32 %c)
33  ret i32 %val
34}
35
36; CHECK-LABEL: .func{{.*}}shfl_sync_rri
37define i32 @shfl_sync_rri(i32 %mask, i32 %a, i32 %b) {
38  ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
39  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
40  ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
41  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 1, [[MASK]];
42  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
43  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 %b, i32 1)
44  ret i32 %val
45}
46
47; CHECK-LABEL: .func{{.*}}shfl_sync_iri
48define i32 @shfl_sync_iri(i32 %a, i32 %b) {
49  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
50  ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
51  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 2, 1;
52  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
53  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 %b, i32 2)
54  ret i32 %val
55}
56
57; CHECK-LABEL: .func{{.*}}shfl_sync_rir
58define i32 @shfl_sync_rir(i32 %mask, i32 %a, i32 %c) {
59  ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
60  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
61  ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
62  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 1, [[C]], [[MASK]];
63  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
64  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 1, i32 %c)
65  ret i32 %val
66}
67
68; CHECK-LABEL: .func{{.*}}shfl_sync_iir
69define i32 @shfl_sync_iir(i32 %a, i32 %c) {
70  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
71  ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
72  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, [[C]], 1;
73  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
74  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 2, i32 %c)
75  ret i32 %val
76}
77
78; CHECK-LABEL: .func{{.*}}shfl_sync_rii
79define i32 @shfl_sync_rii(i32 %mask, i32 %a) {
80  ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
81  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
82  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 1, 2, [[MASK]];
83  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
84  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 1, i32 2)
85  ret i32 %val
86}
87
88; CHECK-LABEL: .func{{.*}}shfl_sync_iii
89define i32 @shfl_sync_iii(i32 %a, i32 %b) {
90  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
91  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, 3, 1;
92  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
93  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 2, i32 3)
94  ret i32 %val
95}
96