llvm/test/Analysis/UniformityAnalysis/NVPTX/diverge.ll

   1 ; RUN: opt %s -passes='print<uniformity>' -disable-output 2>&1 | FileCheck %s
   2
   3 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
   4 target triple = "nvptx64-nvidia-cuda"
   5
   6 ; return (n < 0 ? a + threadIdx.x : b + threadIdx.x)
   7 define i32 @no_diverge(i32 %n, i32 %a, i32 %b) {
   8 ; CHECK-LABEL: for function 'no_diverge'
   9 entry:
  10   %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
  11   %cond = icmp slt i32 %n, 0
  12   br i1 %cond, label %then, label %else ; uniform
  13 ; CHECK-NOT: DIVERGENT: %cond =
  14 ; CHECK-NOT: DIVERGENT: br i1 %cond,
  15 then:
  16   %a1 = add i32 %a, %tid
  17   br label %merge
  18 else:
  19   %b2 = add i32 %b, %tid
  20   br label %merge
  21 merge:
  22   %c = phi i32 [ %a1, %then ], [ %b2, %else ]
  23   ret i32 %c
  24 }
  25
  26 ; c = a;
  27 ; if (threadIdx.x < 5)    // divergent: data dependent
  28 ;   c = b;
  29 ; return c;               // c is divergent: sync dependent
  30 define i32 @sync(i32 %a, i32 %b) {
  31 ; CHECK-LABEL: for function 'sync'
  32 bb1:
  33   %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.y()
  34   %cond = icmp slt i32 %tid, 5
  35   br i1 %cond, label %bb2, label %bb3
  36 ; CHECK:  DIVERGENT: %cond =
  37 ; CHECK: DIVERGENT: br i1 %cond,
  38 bb2:
  39   br label %bb3
  40 bb3:
  41   %c = phi i32 [ %a, %bb1 ], [ %b, %bb2 ] ; sync dependent on tid
  42 ; CHECK: DIVERGENT: %c =
  43   ret i32 %c
  44 }
  45
  46 ; c = 0;
  47 ; if (threadIdx.x >= 5) {  // divergent
  48 ;   c = (n < 0 ? a : b);  // c here is uniform because n is uniform
  49 ; }
  50 ; // c here is divergent because it is sync dependent on threadIdx.x >= 5
  51 ; return c;
  52 define i32 @mixed(i32 %n, i32 %a, i32 %b) {
  53 ; CHECK-LABEL: for function 'mixed'
  54 bb1:
  55   %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.z()
  56   %cond = icmp slt i32 %tid, 5
  57   br i1 %cond, label %bb6, label %bb2
  58 ; CHECK:  DIVERGENT: %cond =
  59 ; CHECK: DIVERGENT: br i1 %cond,
  60 bb2:
  61   %cond2 = icmp slt i32 %n, 0
  62   br i1 %cond2, label %bb4, label %bb3
  63 bb3:
  64   br label %bb5
  65 bb4:
  66   br label %bb5
  67 bb5:
  68   %c = phi i32 [ %a, %bb3 ], [ %b, %bb4 ]
  69 ; CHECK-NOT: DIVERGENT: %c =
  70   br label %bb6
  71 bb6:
  72   %c2 = phi i32 [ 0, %bb1], [ %c, %bb5 ]
  73 ; CHECK: DIVERGENT: %c2 =
  74   ret i32 %c2
  75 }
  76
  77 ; We conservatively treats all parameters of a __device__ function as divergent.
  78 define i32 @device(i32 %n, i32 %a, i32 %b) {
  79 ; CHECK-LABEL: for function 'device'
  80 ; CHECK-DAG: DIVERGENT: i32 %n
  81 ; CHECK-DAG: DIVERGENT: i32 %a
  82 ; CHECK-DAG: DIVERGENT: i32 %b
  83 entry:
  84   %cond = icmp slt i32 %n, 0
  85   br i1 %cond, label %then, label %else
  86 ; CHECK:  DIVERGENT: %cond =
  87 ; CHECK: DIVERGENT: br i1 %cond,
  88 then:
  89   br label %merge
  90 else:
  91   br label %merge
  92 merge:
  93   %c = phi i32 [ %a, %then ], [ %b, %else ]
  94   ret i32 %c
  95 }
  96
  97 ; int i = 0;
  98 ; do {
  99 ;   i++;                  // i here is uniform
 100 ; } while (i < laneid);
 101 ; return i == 10 ? 0 : 1; // i here is divergent
 102 ;
 103 ; The i defined in the loop is used outside.
 104 define i32 @loop() {
 105 ; CHECK-LABEL: for function 'loop'
 106 entry:
 107   %laneid = call i32 @llvm.nvvm.read.ptx.sreg.laneid()
 108   br label %loop
 109 loop:
 110   %i = phi i32 [ 0, %entry ], [ %i1, %loop ]
 111 ; CHECK-NOT: DIVERGENT: %i =
 112   %i1 = add i32 %i, 1
 113   %exit_cond = icmp sge i32 %i1, %laneid
 114   br i1 %exit_cond, label %loop_exit, label %loop
 115 loop_exit:
 116   %cond = icmp eq i32 %i, 10
 117   br i1 %cond, label %then, label %else
 118 ; CHECK:  DIVERGENT: %cond =
 119 ; CHECK: DIVERGENT: br i1 %cond,
 120 then:
 121   ret i32 0
 122 else:
 123   ret i32 1
 124 }
 125
 126 ; Same as @loop, but the loop is in the LCSSA form.
 127 define i32 @lcssa() {
 128 ; CHECK-LABEL: for function 'lcssa'
 129 entry:
 130   %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
 131   br label %loop
 132 loop:
 133   %i = phi i32 [ 0, %entry ], [ %i1, %loop ]
 134 ; CHECK-NOT: DIVERGENT: %i =
 135   %i1 = add i32 %i, 1
 136   %exit_cond = icmp sge i32 %i1, %tid
 137   br i1 %exit_cond, label %loop_exit, label %loop
 138 loop_exit:
 139   %i.lcssa = phi i32 [ %i, %loop ]
 140 ; CHECK: DIVERGENT: %i.lcssa =
 141   %cond = icmp eq i32 %i.lcssa, 10
 142   br i1 %cond, label %then, label %else
 143 ; CHECK:  DIVERGENT: %cond =
 144 ; CHECK: DIVERGENT: br i1 %cond,
 145 then:
 146   ret i32 0
 147 else:
 148   ret i32 1
 149 }
 150
 151 ; Verifies sync-dependence is computed correctly in the absense of loops.
 152 define i32 @sync_no_loop(i32 %arg) {
 153 ; CHECK-LABEL: for function 'sync_no_loop'
 154 entry:
 155   %0 = add i32 %arg, 1
 156   %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
 157   %1 = icmp sge i32 %tid, 10
 158   br i1 %1, label %bb1, label %bb2
 159
 160 bb1:
 161   br label %bb3
 162
 163 bb2:
 164   br label %bb3
 165
 166 bb3:
 167   %2 = add i32 %0, 2
 168   ; CHECK-NOT: DIVERGENT: %2
 169   ret i32 %2
 170 }
 171
 172 declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
 173 declare i32 @llvm.nvvm.read.ptx.sreg.tid.y()
 174 declare i32 @llvm.nvvm.read.ptx.sreg.tid.z()
 175 declare i32 @llvm.nvvm.read.ptx.sreg.laneid()
 176
 177 !nvvm.annotations = !{!0, !1, !2, !3, !4}
 178 !0 = !{ptr @no_diverge, !"kernel", i32 1}
 179 !1 = !{ptr @sync, !"kernel", i32 1}
 180 !2 = !{ptr @mixed, !"kernel", i32 1}
 181 !3 = !{ptr @loop, !"kernel", i32 1}
 182 !4 = !{ptr @sync_no_loop, !"kernel", i32 1}