Author: arsenm
Date: Mon Oct 9 13:06:37 2017
New Revision: 315238
URL: http://llvm.org/viewvc/llvm-project?rev=315238&view=rev
Log:
AMDGPU: Add read_exec_lo/hi builtins
Modified:
cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
cfe/trunk/lib/CodeGen/CGBuiltin.cpp
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=315238&r1=315237&r2=315238&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Mon Oct 9 13:06:37 2017
@@ -121,6 +121,8 @@ TARGET_BUILTIN(__builtin_amdgcn_fmed3h,
// Special builtins.
//===----------------------------------------------------------------------===//
BUILTIN(__builtin_amdgcn_read_exec, "LUi", "nc")
+BUILTIN(__builtin_amdgcn_read_exec_lo, "Ui", "nc")
+BUILTIN(__builtin_amdgcn_read_exec_hi, "Ui", "nc")
//===----------------------------------------------------------------------===//
// R600-NI only builtins.
Modified: cfe/trunk/lib/CodeGen/CGBuiltin.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGBuiltin.cpp?rev=315238&r1=315237&r2=315238&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGBuiltin.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp Mon Oct 9 13:06:37 2017
@@ -9103,6 +9103,15 @@ Value *CodeGenFunction::EmitAMDGPUBuilti
CI->setConvergent();
return CI;
}
+ case AMDGPU::BI__builtin_amdgcn_read_exec_lo:
+ case AMDGPU::BI__builtin_amdgcn_read_exec_hi: {
+ StringRef RegName = BuiltinID == AMDGPU::BI__builtin_amdgcn_read_exec_lo ?
+ "exec_lo" : "exec_hi";
+ CallInst *CI = cast<CallInst>(
+ EmitSpecialRegisterBuiltin(*this, E, Int32Ty, Int32Ty, true, RegName));
+ CI->setConvergent();
+ return CI;
+ }
// amdgcn workitem
case AMDGPU::BI__builtin_amdgcn_workitem_id_x:
Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=315238&r1=315237&r2=315238&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Mon Oct 9 13:06:37 2017
@@ -421,6 +421,18 @@ void test_read_exec(global ulong* out) {
// CHECK: declare i64 @llvm.read_register.i64(metadata)
#[[NOUNWIND_READONLY:[0-9]+]]
+// CHECK-LABEL: @test_read_exec_lo(
+// CHECK: call i32 @llvm.read_register.i32(metadata ![[EXEC_LO:[0-9]+]])
#[[READ_EXEC_ATTRS]]
+void test_read_exec_lo(global uint* out) {
+ *out = __builtin_amdgcn_read_exec_lo();
+}
+
+// CHECK-LABEL: @test_read_exec_hi(
+// CHECK: call i32 @llvm.read_register.i32(metadata ![[EXEC_HI:[0-9]+]])
#[[READ_EXEC_ATTRS]]
+void test_read_exec_hi(global uint* out) {
+ *out = __builtin_amdgcn_read_exec_hi();
+}
+
// CHECK-LABEL: @test_dispatch_ptr
// CHECK: call i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr()
void test_dispatch_ptr(__attribute__((address_space(2))) unsigned char ** out)
@@ -499,3 +511,5 @@ void test_s_getpc(global ulong* out)
// CHECK-DAG: attributes #[[NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly }
// CHECK-DAG: attributes #[[READ_EXEC_ATTRS]] = { convergent }
// CHECK-DAG: ![[EXEC]] = !{!"exec"}
+// CHECK-DAG: ![[EXEC_LO]] = !{!"exec_lo"}
+// CHECK-DAG: ![[EXEC_HI]] = !{!"exec_hi"}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
