96 lines · plain
1; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | FileCheck %s2; RUN: %if ptxas-isa-6.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | %ptxas-verify %}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_rrr14define i32 @shfl_sync_rrr(i32 %mask, i32 %a, i32 %b, i32 %c) {15 ; CHECK: ld.param.b32 [[MASK:%r[0-9]+]]16 ; CHECK: ld.param.b32 [[A:%r[0-9]+]]17 ; CHECK: ld.param.b32 [[B:%r[0-9]+]]18 ; CHECK: ld.param.b32 [[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 %val23}24 25; CHECK-LABEL: .func{{.*}}shfl_sync_irr26define i32 @shfl_sync_irr(i32 %a, i32 %b, i32 %c) {27 ; CHECK: ld.param.b32 [[A:%r[0-9]+]]28 ; CHECK: ld.param.b32 [[B:%r[0-9]+]]29 ; CHECK: ld.param.b32 [[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 %val34}35 36; CHECK-LABEL: .func{{.*}}shfl_sync_rri37define i32 @shfl_sync_rri(i32 %mask, i32 %a, i32 %b) {38 ; CHECK: ld.param.b32 [[MASK:%r[0-9]+]]39 ; CHECK: ld.param.b32 [[A:%r[0-9]+]]40 ; CHECK: ld.param.b32 [[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 %val45}46 47; CHECK-LABEL: .func{{.*}}shfl_sync_iri48define i32 @shfl_sync_iri(i32 %a, i32 %b) {49 ; CHECK: ld.param.b32 [[A:%r[0-9]+]]50 ; CHECK: ld.param.b32 [[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 %val55}56 57; CHECK-LABEL: .func{{.*}}shfl_sync_rir58define i32 @shfl_sync_rir(i32 %mask, i32 %a, i32 %c) {59 ; CHECK: ld.param.b32 [[MASK:%r[0-9]+]]60 ; CHECK: ld.param.b32 [[A:%r[0-9]+]]61 ; CHECK: ld.param.b32 [[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 %val66}67 68; CHECK-LABEL: .func{{.*}}shfl_sync_iir69define i32 @shfl_sync_iir(i32 %a, i32 %c) {70 ; CHECK: ld.param.b32 [[A:%r[0-9]+]]71 ; CHECK: ld.param.b32 [[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 %val76}77 78; CHECK-LABEL: .func{{.*}}shfl_sync_rii79define i32 @shfl_sync_rii(i32 %mask, i32 %a) {80 ; CHECK: ld.param.b32 [[MASK:%r[0-9]+]]81 ; CHECK: ld.param.b32 [[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 %val86}87 88; CHECK-LABEL: .func{{.*}}shfl_sync_iii89define i32 @shfl_sync_iii(i32 %a, i32 %b) {90 ; CHECK: ld.param.b32 [[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 %val95}96