Author: vpykhtin Date: Fri Aug 19 07:54:31 2016 New Revision: 279235 URL: http://llvm.org/viewvc/llvm-project?rev=279235&view=rev Log: [AMDGPU] add s_incperflevel/s_decperflevel builtins
Differential revision: https://reviews.llvm.org/D23668 Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl 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=279235&r1=279234&r2=279235&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original) +++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Fri Aug 19 07:54:31 2016 @@ -70,6 +70,8 @@ BUILTIN(__builtin_amdgcn_cubetc, "ffff", BUILTIN(__builtin_amdgcn_cubema, "ffff", "nc") BUILTIN(__builtin_amdgcn_s_memtime, "LUi", "n") BUILTIN(__builtin_amdgcn_s_sleep, "vIi", "n") +BUILTIN(__builtin_amdgcn_s_incperflevel, "vIi", "n") +BUILTIN(__builtin_amdgcn_s_decperflevel, "vIi", "n") BUILTIN(__builtin_amdgcn_uicmp, "LUiUiUiIi", "nc") BUILTIN(__builtin_amdgcn_uicmpl, "LUiLUiLUiIi", "nc") BUILTIN(__builtin_amdgcn_sicmp, "LUiiiIi", "nc") Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl?rev=279235&r1=279234&r2=279235&view=diff ============================================================================== --- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl (original) +++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl Fri Aug 19 07:54:31 2016 @@ -18,6 +18,16 @@ 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_incperflevel(int x) +{ + __builtin_amdgcn_s_incperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_incperflevel' must be a constant integer}} +} + +void test_s_decperflevel(int x) +{ + __builtin_amdgcn_s_decperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_decperflevel' must be a constant integer}} +} + void test_sicmp_i32(global ulong* out, int a, int b, uint c) { *out = __builtin_amdgcn_sicmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_sicmp' must be a constant integer}} Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=279235&r1=279234&r2=279235&view=diff ============================================================================== --- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original) +++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Fri Aug 19 07:54:31 2016 @@ -286,6 +286,24 @@ void test_s_sleep() __builtin_amdgcn_s_sleep(15); } +// CHECK-LABEL: @test_s_incperflevel +// CHECK: call void @llvm.amdgcn.s.incperflevel(i32 1) +// CHECK: call void @llvm.amdgcn.s.incperflevel(i32 15) +void test_s_incperflevel() +{ + __builtin_amdgcn_s_incperflevel(1); + __builtin_amdgcn_s_incperflevel(15); +} + +// CHECK-LABEL: @test_s_decperflevel +// CHECK: call void @llvm.amdgcn.s.decperflevel(i32 1) +// CHECK: call void @llvm.amdgcn.s.decperflevel(i32 15) +void test_s_decperflevel() +{ + __builtin_amdgcn_s_decperflevel(1); + __builtin_amdgcn_s_decperflevel(15); +} + // CHECK-LABEL: @test_cubeid( // CHECK: call float @llvm.amdgcn.cubeid(float %a, float %b, float %c) void test_cubeid(global float* out, float a, float b, float c) { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
