[PATCH] D38698: AMDGPU: Add read_exec_lo/hi builtins

[Matt Arsenault via Phabricator via cfe-commits]

arsenm created this revision.
Herald added subscribers: t-tye, tpr, dstuttard, nhaehnle, wdng, kzhuravl.

https://reviews.llvm.org/D38698

Files:
  include/clang/Basic/BuiltinsAMDGPU.def
  lib/CodeGen/CGBuiltin.cpp
  test/CodeGenOpenCL/builtins-amdgcn.cl


Index: test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn.cl
+++ test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -421,6 +421,18 @@
 
 // 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 @@
 // 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"}
Index: lib/CodeGen/CGBuiltin.cpp
===================================================================
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -9103,6 +9103,15 @@
     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:
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -121,6 +121,8 @@
 // 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.



Index: test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn.cl
+++ test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -421,6 +421,18 @@
 
 // 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 @@
 // 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"}
Index: lib/CodeGen/CGBuiltin.cpp
===================================================================
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -9103,6 +9103,15 @@
     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:
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -121,6 +121,8 @@
 // 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.

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits