mirror of
https://github.com/RPCS3/llvm-mirror.git
synced 2024-10-21 20:12:56 +02:00
3251647c03
Summary: Convergent instrs shouldn't be made control-dependent on other values, but this is basically the whole point of tail duplication. So just bail if we see a convergent instruction. Reviewers: iteratee Subscribers: jholewinski, jhen, hfinkel, tra, jingyue, llvm-commits Differential Revision: http://reviews.llvm.org/D17320 llvm-svn: 261540
46 lines
1011 B
LLVM
46 lines
1011 B
LLVM
; RUN: llc -O2 -tail-dup-size=100 -enable-tail-merge=0 < %s | FileCheck %s
|
|
target triple = "nvptx64-nvidia-cuda"
|
|
|
|
declare void @foo()
|
|
declare void @llvm.cuda.syncthreads()
|
|
|
|
; syncthreads shouldn't be duplicated.
|
|
; CHECK: .func call_syncthreads
|
|
; CHECK: bar.sync
|
|
; CHECK-NOT: bar.sync
|
|
define void @call_syncthreads(i32* %a, i32* %b, i1 %cond, i1 %cond2) nounwind {
|
|
br i1 %cond, label %L1, label %L2
|
|
br i1 %cond2, label %Ret, label %L1
|
|
Ret:
|
|
ret void
|
|
L1:
|
|
store i32 0, i32* %a
|
|
br label %L42
|
|
L2:
|
|
store i32 1, i32* %a
|
|
br label %L42
|
|
L42:
|
|
call void @llvm.cuda.syncthreads()
|
|
br label %Ret
|
|
}
|
|
|
|
; Check that call_syncthreads really does trigger tail duplication.
|
|
; CHECK: .func call_foo
|
|
; CHECK: call
|
|
; CHECK: call
|
|
define void @call_foo(i32* %a, i32* %b, i1 %cond, i1 %cond2) nounwind {
|
|
br i1 %cond, label %L1, label %L2
|
|
br i1 %cond2, label %Ret, label %L1
|
|
Ret:
|
|
ret void
|
|
L1:
|
|
store i32 0, i32* %a
|
|
br label %L42
|
|
L2:
|
|
store i32 1, i32* %a
|
|
br label %L42
|
|
L42:
|
|
call void @foo()
|
|
br label %Ret
|
|
}
|