[llvm] r274779 - [NVVM] Rename __nvvm_bar0 builtin back to __syncthreads.

Justin Lebar via llvm-commits llvm-commits at lists.llvm.org
Thu Jul 7 11:14:55 PDT 2016


Author: jlebar
Date: Thu Jul  7 13:14:55 2016
New Revision: 274779

URL: http://llvm.org/viewvc/llvm-project?rev=274779&view=rev
Log:
[NVVM] Rename __nvvm_bar0 builtin back to __syncthreads.

__syncthreads was renamed to __nvvm_bar0 in r274664.  But __syncthreads
is part of our user-facing API, so we need to keep the name.

This will momentarily break clang; we need a matching patch there.

Patch by Justin Bogner.

Modified:
    llvm/trunk/include/llvm/IR/IntrinsicsNVVM.td

Modified: llvm/trunk/include/llvm/IR/IntrinsicsNVVM.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsNVVM.td?rev=274779&r1=274778&r2=274779&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsNVVM.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsNVVM.td Thu Jul  7 13:14:55 2016
@@ -729,7 +729,10 @@ def llvm_anyi64ptr_ty     : LLVMAnyPoint
                                       [IntrArgMemOnly, NoCapture<0>]>;
 
 // Bar.Sync
-  def int_nvvm_barrier0 : GCCBuiltin<"__nvvm_bar0">,
+
+  // The builtin for "bar.sync 0" is called __syncthreads.  Unlike most of the
+  // intrinsics in this file, this one is a user-facing API.
+  def int_nvvm_barrier0 : GCCBuiltin<"__syncthreads">,
       Intrinsic<[], [], [IntrConvergent]>;
   def int_nvvm_barrier0_popc : GCCBuiltin<"__nvvm_bar0_popc">,
       Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrConvergent]>;




More information about the llvm-commits mailing list