diff options
| author | Justin Lebar <jlebar@google.com> | 2016-02-17 17:46:54 +0000 |
|---|---|---|
| committer | Justin Lebar <jlebar@google.com> | 2016-02-17 17:46:54 +0000 |
| commit | f9b5add6adb7a178b8d5a385b734226f9bb4dea2 (patch) | |
| tree | 2e4708577246760332e8e39aefc1836e7313a86a /llvm/lib/Target | |
| parent | b5c7b1c00fc3bf4bdfb7fa7e7429594990c50609 (diff) | |
| download | bcm5719-llvm-f9b5add6adb7a178b8d5a385b734226f9bb4dea2.tar.gz bcm5719-llvm-f9b5add6adb7a178b8d5a385b734226f9bb4dea2.zip | |
[NVPTX] Annotate convergent intrinsics as convergent.
Summary:
Previously the machine instructions for bar.sync &co. were not marked as
convergent. This resulted in some MI passes (such as TailDuplication,
fixed in an upcoming patch) doing unsafe things to these instructions.
Reviewers: jingyue
Subscribers: llvm-commits, tra, jholewinski, hfinkel
Differential Revision: http://reviews.llvm.org/D17318
llvm-svn: 261115
Diffstat (limited to 'llvm/lib/Target')
| -rw-r--r-- | llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 2 |
1 files changed, 2 insertions, 0 deletions
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td index 14e51aa309e..208731f4738 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -32,6 +32,7 @@ def immDouble1 : PatLeaf<(fpimm), [{ //----------------------------------- // Synchronization Functions //----------------------------------- +let isConvergent = 1 in { def INT_CUDA_SYNCTHREADS : NVPTXInst<(outs), (ins), "bar.sync \t0;", [(int_cuda_syncthreads)]>; @@ -63,6 +64,7 @@ def INT_BARRIER0_OR : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred), !strconcat("selp.u32 \t$dst, 1, 0, %p2; \n\t", !strconcat("}}", ""))))))), [(set Int32Regs:$dst, (int_nvvm_barrier0_or Int32Regs:$pred))]>; +} // isConvergent = 1 //----------------------------------- |
