[PATCH] D17318: [NVPTX] Annotate convergent intrinsics as convergent.

[Justin Lebar via llvm-commits]

jlebar created this revision.
jlebar added reviewers: jingyue, hfinkel.
jlebar added subscribers: tra, llvm-commits.
Herald added a subscriber: jholewinski.

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.

http://reviews.llvm.org/D17318

Files:
  lib/Target/NVPTX/NVPTXIntrinsics.td

Index: lib/Target/NVPTX/NVPTXIntrinsics.td
===================================================================
--- lib/Target/NVPTX/NVPTXIntrinsics.td
+++ lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -32,6 +32,7 @@
 //-----------------------------------
 // Synchronization Functions
 //-----------------------------------
+let isConvergent = 1 in {
 def INT_CUDA_SYNCTHREADS : NVPTXInst<(outs), (ins),
                   "bar.sync \t0;",
       [(int_cuda_syncthreads)]>;
@@ -63,6 +64,7 @@
       !strconcat("selp.u32 \t$dst, 1, 0, %p2; \n\t",
         !strconcat("}}", ""))))))),
       [(set Int32Regs:$dst, (int_nvvm_barrier0_or Int32Regs:$pred))]>;
+} // isConvergent = 1
 
 
 //-----------------------------------


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D17318.48133.patch
Type: text/x-patch
Size: 714 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160217/92329834/attachment.bin>