From: Valery Pykhtin <Valery.Pykhtin@amd.com>
Date: Fri, 19 Aug 2016 12:54:31 +0000 (+0000)
Subject: [AMDGPU] add s_incperflevel/s_decperflevel builtins
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=d16a5514753a18cdf8af8bb3c0b1f0cbd471acc6;p=clang

[AMDGPU] add s_incperflevel/s_decperflevel builtins

Differential revision: https://reviews.llvm.org/D23668

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

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index 0e54128080..9e80ae7b71 100644
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/include/clang/Basic/BuiltinsAMDGPU.def
@@ -70,6 +70,8 @@ BUILTIN(__builtin_amdgcn_cubetc, "ffff", "nc")
 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")
diff --git a/test/CodeGenOpenCL/builtins-amdgcn-error.cl b/test/CodeGenOpenCL/builtins-amdgcn-error.cl
index 5513c5a958..97b834f519 100644
--- a/test/CodeGenOpenCL/builtins-amdgcn-error.cl
+++ b/test/CodeGenOpenCL/builtins-amdgcn-error.cl
@@ -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}}
diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl
index 3280f19076..a715d98d13 100644
--- a/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -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) {