[NVPTX] Annotate convergent intrinsics as convergent.

[jlebar]

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.

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

def immDouble1 : PatLeaf<(fpimm), [{
     return (d==1.0);
 }]>;
 
 
 
 //-----------------------------------
 // Synchronization Functions
 //-----------------------------------
+let isConvergent = 1 in {
 def INT_CUDA_SYNCTHREADS : NVPTXInst<(outs), (ins),
                   "bar.sync \t0;",
       [(int_cuda_syncthreads)]>;
 def INT_BARRIER0 : NVPTXInst<(outs), (ins),
                   "bar.sync \t0;",
       [(int_nvvm_barrier0)]>;
 def INT_BARRIER0_POPC : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
   !strconcat("{{ \n\t",
def INT_BARRIER0_OR : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
   !strconcat("{{ \n\t",
       !strconcat(".reg .pred \t%p1; \n\t",
       !strconcat(".reg .pred \t%p2; \n\t",
       !strconcat("setp.ne.u32 \t%p1, $pred, 0; \n\t",
       !strconcat("bar.red.or.pred \t%p2, 0, %p1; \n\t",
       !strconcat("selp.u32 \t$dst, 1, 0, %p2; \n\t",
         !strconcat("}}", ""))))))),
       [(set Int32Regs:$dst, (int_nvvm_barrier0_or Int32Regs:$pred))]>;
+} // isConvergent = 1
 
 
 //-----------------------------------
 // Explicit Memory Fence Functions
 //-----------------------------------
 class MEMBAR<string StrOp, Intrinsic IntOP> :
               NVPTXInst<(outs), (ins),
             StrOp, [(IntOP)]>;