Author: jvesely Date: Fri Feb 24 22:20:20 2017 New Revision: 296239 URL: http://llvm.org/viewvc/llvm-project?rev=296239&view=rev Log: AMDGPU: export s_waitcnt builtin
Differential Revision: https://reviews.llvm.org/D30359 Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=296239&r1=296238&r2=296239&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original) +++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Fri Feb 24 22:20:20 2017 @@ -36,6 +36,7 @@ BUILTIN(__builtin_amdgcn_workitem_id_z, // Instruction builtins. //===----------------------------------------------------------------------===// BUILTIN(__builtin_amdgcn_s_getreg, "UiIi", "n") +BUILTIN(__builtin_amdgcn_s_waitcnt, "vIi", "n") BUILTIN(__builtin_amdgcn_s_barrier, "v", "n") BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n") BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n") Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=296239&r1=296238&r2=296239&view=diff ============================================================================== --- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original) +++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Fri Feb 24 22:20:20 2017 @@ -263,6 +263,13 @@ void test_class_f64(global double* out, *out = __builtin_amdgcn_class(a, b); } +// CHECK-LABEL: @test_s_waitcnt +// CHECK: call void @llvm.amdgcn.s.waitcnt( +void test_s_waitcnt() +{ + __builtin_amdgcn_s_waitcnt(0); +} + // CHECK-LABEL: @test_s_barrier // CHECK: call void @llvm.amdgcn.s.barrier( void test_s_barrier() Modified: cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl?rev=296239&r1=296238&r2=296239&view=diff ============================================================================== --- cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl (original) +++ cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl Fri Feb 24 22:20:20 2017 @@ -18,6 +18,11 @@ void test_s_sleep(int x) __builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}} } +void test_s_waitcnt(int x) +{ + __builtin_amdgcn_s_waitcnt(x); // expected-error {{argument to '__builtin_amdgcn_s_waitcnt' must be a constant integer}} +} + void test_s_incperflevel(int x) { __builtin_amdgcn_s_incperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_incperflevel' must be a constant integer}} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
