[PATCH] D21161: [NVPTX] Mark bar.sync intrinsic as convergent.

Justin Lebar via llvm-commits llvm-commits at lists.llvm.org
Wed Jun 8 17:22:23 PDT 2016


jlebar created this revision.
jlebar added a reviewer: tra.
jlebar added subscribers: jholewinski, llvm-commits.

__syncthreads, which corresponds to bar.sync 0, is already convergent.
This makes the more general bar.sync n likewise convergent.

http://reviews.llvm.org/D21161

Files:
  include/llvm/IR/IntrinsicsNVVM.td

Index: include/llvm/IR/IntrinsicsNVVM.td
===================================================================
--- include/llvm/IR/IntrinsicsNVVM.td
+++ include/llvm/IR/IntrinsicsNVVM.td
@@ -3744,7 +3744,7 @@
 def int_ptx_read_pm3         : PTXReadSpecialRegisterIntrinsic_r32
                                <"__builtin_ptx_read_pm3">;
 
-def int_ptx_bar_sync : Intrinsic<[], [llvm_i32_ty], []>,
+def int_ptx_bar_sync : Intrinsic<[], [llvm_i32_ty], [IntrConvergent]>,
                        GCCBuiltin<"__builtin_ptx_bar_sync">;
 
 //


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D21161.60124.patch
Type: text/x-patch
Size: 535 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160609/73274fab/attachment.bin>


More information about the llvm-commits mailing list