commit f12e3b848afe8aec6b7bbabb5e26875e8b5636fc
author Matt Arsenault <Matthew.Arsenault@amd.com> Mon Oct 09 20:06:37 2017 +0000
committer Matt Arsenault <Matthew.Arsenault@amd.com> Mon Oct 09 20:06:37 2017 +0000
tree 491b9e77ff73ee4000dc605b9d2e677f0486f9ab
parent c1d5955684dba9a31cb0ff3b4b61fe2a84e392e8

AMDGPU: Add read_exec_lo/hi builtins

llvm-svn: 315238

clang/lib/CodeGen/CGBuiltin.cpp

diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 7ae558f..a4e6452 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/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: