177 lines · plain
1; RUN: opt %s -passes='print<uniformity>' -disable-output 2>&1 | FileCheck %s2 3target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"4target triple = "nvptx64-nvidia-cuda"5 6; return (n < 0 ? a + threadIdx.x : b + threadIdx.x)7define ptx_kernel i32 @no_diverge(i32 %n, i32 %a, i32 %b) {8; CHECK-LABEL: for function 'no_diverge'9entry:10 %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()11 %cond = icmp slt i32 %n, 012 br i1 %cond, label %then, label %else ; uniform13; CHECK-NOT: DIVERGENT: %cond =14; CHECK-NOT: DIVERGENT: br i1 %cond,15then:16 %a1 = add i32 %a, %tid17 br label %merge18else:19 %b2 = add i32 %b, %tid20 br label %merge21merge:22 %c = phi i32 [ %a1, %then ], [ %b2, %else ]23 ret i32 %c24}25 26; c = a;27; if (threadIdx.x < 5) // divergent: data dependent28; c = b;29; return c; // c is divergent: sync dependent30define ptx_kernel i32 @sync(i32 %a, i32 %b) {31; CHECK-LABEL: for function 'sync'32bb1:33 %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.y()34 %cond = icmp slt i32 %tid, 535 br i1 %cond, label %bb2, label %bb336; CHECK: DIVERGENT: %cond =37; CHECK: DIVERGENT: br i1 %cond,38bb2:39 br label %bb340bb3:41 %c = phi i32 [ %a, %bb1 ], [ %b, %bb2 ] ; sync dependent on tid42; CHECK: DIVERGENT: %c =43 ret i32 %c44}45 46; c = 0;47; if (threadIdx.x >= 5) { // divergent48; c = (n < 0 ? a : b); // c here is uniform because n is uniform49; }50; // c here is divergent because it is sync dependent on threadIdx.x >= 551; return c;52define ptx_kernel i32 @mixed(i32 %n, i32 %a, i32 %b) {53; CHECK-LABEL: for function 'mixed'54bb1:55 %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.z()56 %cond = icmp slt i32 %tid, 557 br i1 %cond, label %bb6, label %bb258; CHECK: DIVERGENT: %cond =59; CHECK: DIVERGENT: br i1 %cond,60bb2:61 %cond2 = icmp slt i32 %n, 062 br i1 %cond2, label %bb4, label %bb363bb3:64 br label %bb565bb4:66 br label %bb567bb5:68 %c = phi i32 [ %a, %bb3 ], [ %b, %bb4 ]69; CHECK-NOT: DIVERGENT: %c =70 br label %bb671bb6:72 %c2 = phi i32 [ 0, %bb1], [ %c, %bb5 ]73; CHECK: DIVERGENT: %c2 =74 ret i32 %c275}76 77; We conservatively treats all parameters of a __device__ function as divergent.78define i32 @device(i32 %n, i32 %a, i32 %b) {79; CHECK-LABEL: for function 'device'80; CHECK-DAG: DIVERGENT: i32 %n81; CHECK-DAG: DIVERGENT: i32 %a82; CHECK-DAG: DIVERGENT: i32 %b83entry:84 %cond = icmp slt i32 %n, 085 br i1 %cond, label %then, label %else86; CHECK: DIVERGENT: %cond =87; CHECK: DIVERGENT: br i1 %cond,88then:89 br label %merge90else:91 br label %merge92merge:93 %c = phi i32 [ %a, %then ], [ %b, %else ]94 ret i32 %c95}96 97; int i = 0;98; do {99; i++; // i here is uniform100; } while (i < laneid);101; return i == 10 ? 0 : 1; // i here is divergent102;103; The i defined in the loop is used outside.104define ptx_kernel i32 @loop() {105; CHECK-LABEL: for function 'loop'106entry:107 %laneid = call i32 @llvm.nvvm.read.ptx.sreg.laneid()108 br label %loop109loop:110 %i = phi i32 [ 0, %entry ], [ %i1, %loop ]111; CHECK-NOT: DIVERGENT: %i =112 %i1 = add i32 %i, 1113 %exit_cond = icmp sge i32 %i1, %laneid114 br i1 %exit_cond, label %loop_exit, label %loop115loop_exit:116 %cond = icmp eq i32 %i, 10117 br i1 %cond, label %then, label %else118; CHECK: DIVERGENT: %cond =119; CHECK: DIVERGENT: br i1 %cond,120then:121 ret i32 0122else:123 ret i32 1124}125 126; Same as @loop, but the loop is in the LCSSA form.127define i32 @lcssa() {128; CHECK-LABEL: for function 'lcssa'129entry:130 %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()131 br label %loop132loop:133 %i = phi i32 [ 0, %entry ], [ %i1, %loop ]134; CHECK-NOT: DIVERGENT: %i =135 %i1 = add i32 %i, 1136 %exit_cond = icmp sge i32 %i1, %tid137 br i1 %exit_cond, label %loop_exit, label %loop138loop_exit:139 %i.lcssa = phi i32 [ %i, %loop ]140; CHECK: DIVERGENT: %i.lcssa =141 %cond = icmp eq i32 %i.lcssa, 10142 br i1 %cond, label %then, label %else143; CHECK: DIVERGENT: %cond =144; CHECK: DIVERGENT: br i1 %cond,145then:146 ret i32 0147else:148 ret i32 1149}150 151; Verifies sync-dependence is computed correctly in the absense of loops.152define ptx_kernel i32 @sync_no_loop(i32 %arg) {153; CHECK-LABEL: for function 'sync_no_loop'154entry:155 %0 = add i32 %arg, 1156 %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()157 %1 = icmp sge i32 %tid, 10158 br i1 %1, label %bb1, label %bb2159 160bb1:161 br label %bb3162 163bb2:164 br label %bb3165 166bb3:167 %2 = add i32 %0, 2168 ; CHECK-NOT: DIVERGENT: %2169 ret i32 %2170}171 172declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()173declare i32 @llvm.nvvm.read.ptx.sreg.tid.y()174declare i32 @llvm.nvvm.read.ptx.sreg.tid.z()175declare i32 @llvm.nvvm.read.ptx.sreg.laneid()176 177