commit b72888647baee2093c2cbd7a0c3da6dc8b870f20
author Matt Arsenault <Matthew.Arsenault@amd.com> Mon Jan 28 23:59:18 2019 +0000
committer Matt Arsenault <Matthew.Arsenault@amd.com> Mon Jan 28 23:59:18 2019 +0000
tree 46b3c066b9e0f03ef065f854908f0d96ba14ef61
parent 33f87b8aef87cd365ec59f0961ad8de8bd848bc0

AMDGPU: Add ds append/consume builtins

llvm-svn: 352443

clang/lib/CodeGen/CGBuiltin.cpp

diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 195b0c0..fa4d512 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -12478,6 +12478,14 @@
   case AMDGPU::BI__builtin_amdgcn_fmed3f:
   case AMDGPU::BI__builtin_amdgcn_fmed3h:
     return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_fmed3);
+  case AMDGPU::BI__builtin_amdgcn_ds_append:
+  case AMDGPU::BI__builtin_amdgcn_ds_consume: {
+    Intrinsic::ID Intrin = BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_append ?
+      Intrinsic::amdgcn_ds_append : Intrinsic::amdgcn_ds_consume;
+    Value *Src0 = EmitScalarExpr(E->getArg(0));
+    Function *F = CGM.getIntrinsic(Intrin, { Src0->getType() });
+    return Builder.CreateCall(F, { Src0, Builder.getFalse() });
+  }
   case AMDGPU::BI__builtin_amdgcn_read_exec: {
     CallInst *CI = cast<CallInst>(
       EmitSpecialRegisterBuiltin(*this, E, Int64Ty, Int64Ty, true, "exec"));