[PATCH] D17318: [NVPTX] Annotate convergent intrinsics as convergent.
Justin Lebar via llvm-commits
llvm-commits at lists.llvm.org
Tue Feb 16 17:03:11 PST 2016
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>
More information about the llvm-commits
mailing list