1*9880d681SAndroid Build Coastguard Worker; RUN: llc -O2 -tail-dup-size=100 -enable-tail-merge=0 < %s | FileCheck %s 2*9880d681SAndroid Build Coastguard Workertarget triple = "nvptx64-nvidia-cuda" 3*9880d681SAndroid Build Coastguard Worker 4*9880d681SAndroid Build Coastguard Workerdeclare void @foo() 5*9880d681SAndroid Build Coastguard Workerdeclare void @llvm.nvvm.barrier0() 6*9880d681SAndroid Build Coastguard Worker 7*9880d681SAndroid Build Coastguard Worker; syncthreads shouldn't be duplicated. 8*9880d681SAndroid Build Coastguard Worker; CHECK: .func call_syncthreads 9*9880d681SAndroid Build Coastguard Worker; CHECK: bar.sync 10*9880d681SAndroid Build Coastguard Worker; CHECK-NOT: bar.sync 11*9880d681SAndroid Build Coastguard Workerdefine void @call_syncthreads(i32* %a, i32* %b, i1 %cond, i1 %cond2) nounwind { 12*9880d681SAndroid Build Coastguard Worker br i1 %cond, label %L1, label %L2 13*9880d681SAndroid Build Coastguard Worker br i1 %cond2, label %Ret, label %L1 14*9880d681SAndroid Build Coastguard WorkerRet: 15*9880d681SAndroid Build Coastguard Worker ret void 16*9880d681SAndroid Build Coastguard WorkerL1: 17*9880d681SAndroid Build Coastguard Worker store i32 0, i32* %a 18*9880d681SAndroid Build Coastguard Worker br label %L42 19*9880d681SAndroid Build Coastguard WorkerL2: 20*9880d681SAndroid Build Coastguard Worker store i32 1, i32* %a 21*9880d681SAndroid Build Coastguard Worker br label %L42 22*9880d681SAndroid Build Coastguard WorkerL42: 23*9880d681SAndroid Build Coastguard Worker call void @llvm.nvvm.barrier0() 24*9880d681SAndroid Build Coastguard Worker br label %Ret 25*9880d681SAndroid Build Coastguard Worker} 26*9880d681SAndroid Build Coastguard Worker 27*9880d681SAndroid Build Coastguard Worker; Check that call_syncthreads really does trigger tail duplication. 28*9880d681SAndroid Build Coastguard Worker; CHECK: .func call_foo 29*9880d681SAndroid Build Coastguard Worker; CHECK: call 30*9880d681SAndroid Build Coastguard Worker; CHECK: call 31*9880d681SAndroid Build Coastguard Workerdefine void @call_foo(i32* %a, i32* %b, i1 %cond, i1 %cond2) nounwind { 32*9880d681SAndroid Build Coastguard Worker br i1 %cond, label %L1, label %L2 33*9880d681SAndroid Build Coastguard Worker br i1 %cond2, label %Ret, label %L1 34*9880d681SAndroid Build Coastguard WorkerRet: 35*9880d681SAndroid Build Coastguard Worker ret void 36*9880d681SAndroid Build Coastguard WorkerL1: 37*9880d681SAndroid Build Coastguard Worker store i32 0, i32* %a 38*9880d681SAndroid Build Coastguard Worker br label %L42 39*9880d681SAndroid Build Coastguard WorkerL2: 40*9880d681SAndroid Build Coastguard Worker store i32 1, i32* %a 41*9880d681SAndroid Build Coastguard Worker br label %L42 42*9880d681SAndroid Build Coastguard WorkerL42: 43*9880d681SAndroid Build Coastguard Worker call void @foo() 44*9880d681SAndroid Build Coastguard Worker br label %Ret 45*9880d681SAndroid Build Coastguard Worker} 46