AMDGPU: Add builtin for getreg intrinsic

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

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index f0f63fa73a79914e418a81a87612773642804599..309dc91cd74fe894ac818b71c1783cd6c1297ecb 100644 (file)
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/include/clang/Basic/BuiltinsAMDGPU.def
@@ -35,6 +35,7 @@ BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc")
 //===----------------------------------------------------------------------===//
 // Instruction builtins.
 //===----------------------------------------------------------------------===//
+BUILTIN(__builtin_amdgcn_s_getreg, "UiIi", "n")
 BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")

diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl
index ce566b6e876a1d3db6fd0b7d22a79006eb073939..69f6ab304b7a9e98d060c6f248031b3e9a79fcbb 100644 (file)
--- a/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -371,6 +371,17 @@ void test_get_group_id(int d, global int *out)
        }
 }
 
+// CHECK-LABEL: @test_s_getreg(
+// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 0)
+// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 1)
+// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 65535)
+void test_s_getreg(volatile global uint *out)
+{
+  *out = __builtin_amdgcn_s_getreg(0);
+  *out = __builtin_amdgcn_s_getreg(1);
+  *out = __builtin_amdgcn_s_getreg(65535);
+}
+
 // CHECK-LABEL: @test_get_local_id(
 // CHECK: tail call i32 @llvm.amdgcn.workitem.id.x(), !range [[WI_RANGE:![0-9]*]]
 // CHECK: tail call i32 @llvm.amdgcn.workitem.id.y(), !range [[WI_RANGE]]

diff --git a/test/SemaOpenCL/builtins-amdgcn-error.cl b/test/SemaOpenCL/builtins-amdgcn-error.cl
index 83ccbefddc6f742bdb9b8f4145608f0c8672ff28..d8576e2bb53d4512a234d31ac0e558b29c989b7b 100644 (file)
--- a/test/SemaOpenCL/builtins-amdgcn-error.cl
+++ b/test/SemaOpenCL/builtins-amdgcn-error.cl
@@ -62,3 +62,8 @@ void test_ds_swizzle(global int* out, int a, int b)
 {
   *out = __builtin_amdgcn_ds_swizzle(a, b); // expected-error {{argument to '__builtin_amdgcn_ds_swizzle' must be a constant integer}}
 }
+
+void test_s_getreg(global int* out, int a)
+{
+  *out = __builtin_amdgcn_s_getreg(a); // expected-error {{argument to '__builtin_amdgcn_s_getreg' must be a constant integer}}
+}