From: Matt Arsenault Date: Fri, 22 Jan 2016 21:56:30 +0000 (+0000) Subject: AMDGPU: Add barrier builtin X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=ca57cfb3059b822634b40099055e8d377c73c934;p=clang 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 e26e72e11c..17d6547bb9 100644 --- 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 2f2ccb3447..6f9a8cf921 100644 --- 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