llvm/test/CodeGen/NVPTX/shfl-sync.ll

   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 %}
   3
   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)
  12
  13 ; CHECK-LABEL: .func{{.*}}shfl_sync_rrr
  14 define 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
  26 define 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
  37 define 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
  48 define 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
  58 define 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
  69 define 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
  79 define 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
  89 define 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 }