From: Matt Arsenault <Matthew.Arsenault@amd.com>
Date: Fri, 20 Jan 2017 19:24:22 +0000 (+0000)
Subject: AMDGPU: Add builtin for getreg intrinsic
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=f1da5d517624133ea5c6fe34589cc9873c74a8c6;p=clang

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 f0f63fa73a..309dc91cd7 100644
--- 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 ce566b6e87..69f6ab304b 100644
--- 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 83ccbefddc..d8576e2bb5 100644
--- 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}}
+}