[llvm] r261115 - [NVPTX] Annotate convergent intrinsics as convergent.

[Justin Lebar via llvm-commits]

Author: jlebar
Date: Wed Feb 17 11:46:54 2016
New Revision: 261115

URL: http://llvm.org/viewvc/llvm-project?rev=261115&view=rev
Log:
[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

Modified:
    llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td

Modified: llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td?rev=261115&r1=261114&r2=261115&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td Wed Feb 17 11:46:54 2016
@@ -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 In
       !strconcat("selp.u32 \t$dst, 1, 0, %p2; \n\t",
         !strconcat("}}", ""))))))),
       [(set Int32Regs:$dst, (int_nvvm_barrier0_or Int32Regs:$pred))]>;
+} // isConvergent = 1
 
 
 //-----------------------------------