r258564 - AMDGPU: Add barrier builtin

[Matt Arsenault via cfe-commits]

Author: arsenm
Date: Fri Jan 22 15:56:30 2016
New Revision: 258564

URL: http://llvm.org/viewvc/llvm-project?rev=258564&view=rev
Log:
AMDGPU: Add barrier builtin

Modified:
    cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
    cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl

Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=258564&r1=258563&r2=258564&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Fri Jan 22 15:56:30 2016
@@ -14,6 +14,7 @@
 
 // The format of this database matches clang/Basic/Builtins.def.
 
+BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
 BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n")
 BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc")

Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=258564&r1=258563&r2=258564&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Fri Jan 22 15:56:30 2016
@@ -141,6 +141,12 @@ void test_class_f64(global double* out,
   *out = __builtin_amdgcn_class(a, b);
 }
 
+// CHECK-LABEL: @test_s_barrier
+// CHECK: call void @llvm.amdgcn.s.barrier(
+void test_s_barrier()
+{
+  __builtin_amdgcn_s_barrier();
+}
 
 // Legacy intrinsics with AMDGPU prefix