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}}
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
