AMDGPU: Add barrier builtin

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

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index e26e72e11c9c1a7ebd671f22dd7b118a3fdaa5a9..17d6547bb967db81d602f6894772f3804375cfdd 100644 (file)
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/include/clang/Basic/BuiltinsAMDGPU.def
@@ -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")

diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl
index 2f2ccb34472d9196aeac3295c2ba919ae95cc421..6f9a8cf921699c2f168a10ab66308cd1f671a1da 100644 (file)
--- a/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -141,6 +141,12 @@ void test_class_f64(global double* out, double a, int b)
   *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