r292636 - AMDGPU: Add builtin for getreg intrinsic

[Matt Arsenault via cfe-commits]

Author: arsenm
Date: Fri Jan 20 13:24:22 2017
New Revision: 292636

URL: http://llvm.org/viewvc/llvm-project?rev=292636&view=rev
Log:
AMDGPU: Add builtin for getreg intrinsic

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=292636&r1=292635&r2=292636&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Fri Jan 20 13:24:22 2017
@@ -35,6 +35,7 @@ BUILTIN(__builtin_amdgcn_workitem_id_z,
 //===----------------------------------------------------------------------===//
 // 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")

Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=292636&r1=292635&r2=292636&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Fri Jan 20 13:24:22 2017
@@ -371,6 +371,17 @@ void test_get_group_id(int d, global int
 	}
 }
 
+// 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]]

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=292636&r1=292635&r2=292636&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl (original)
+++ cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl Fri Jan 20 13:24:22 2017
@@ -62,3 +62,8 @@ void test_ds_swizzle(global int* out, in
 {
   *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}}
+}