[CUDA] Rename the __nvvm_bar0 builtin back to __syncthreads.

The builtin was renamed in r274770.  But __syncthreads is part of our
user-facing API, so we need to keep the name as-is.

Patch by Justin Bogner.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@274780 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/clang/Basic/BuiltinsNVPTX.def b/include/clang/Basic/BuiltinsNVPTX.def
index acfed606d5280c11069ae6e49f126eae8834fa9a..456d0001a12d4b4f80d1da148a8a35befec3a273 100644 (file)
--- a/include/clang/Basic/BuiltinsNVPTX.def
+++ b/include/clang/Basic/BuiltinsNVPTX.def
@@ -393,7 +393,7 @@ BUILTIN(__nvvm_bitcast_d2ll, "LLid", "")
 
 // Sync
 
-BUILTIN(__nvvm_bar0, "v", "")
+BUILTIN(__syncthreads, "v", "")
 BUILTIN(__nvvm_bar0_popc, "ii", "")
 BUILTIN(__nvvm_bar0_and, "ii", "")
 BUILTIN(__nvvm_bar0_or, "ii", "")

diff --git a/test/CodeGen/builtins-nvptx.c b/test/CodeGen/builtins-nvptx.c
index 2cb0ff3e37ff7e16a9bcb2b6fed1c163747f2440..cd21361140bcddb628fa40f898ca49f1d905a592 100644 (file)
--- a/test/CodeGen/builtins-nvptx.c
+++ b/test/CodeGen/builtins-nvptx.c
@@ -179,7 +179,7 @@ __device__ void nvvm_math(float f1, float f2, double d1, double d2) {
 // CHECK: call void @llvm.nvvm.membar.sys()
   __nvvm_membar_sys();
 // CHECK: call void @llvm.nvvm.barrier0()
-  __nvvm_bar0();
+  __syncthreads();
 }
 
 __device__ int di;