diff options
| author | Matt Arsenault <Matthew.Arsenault@amd.com> | 2019-01-28 23:59:18 +0000 |
|---|---|---|
| committer | Matt Arsenault <Matthew.Arsenault@amd.com> | 2019-01-28 23:59:18 +0000 |
| commit | b72888647baee2093c2cbd7a0c3da6dc8b870f20 (patch) | |
| tree | 46b3c066b9e0f03ef065f854908f0d96ba14ef61 /clang/lib | |
| parent | 33f87b8aef87cd365ec59f0961ad8de8bd848bc0 (diff) | |
| download | bcm5719-llvm-b72888647baee2093c2cbd7a0c3da6dc8b870f20.tar.gz bcm5719-llvm-b72888647baee2093c2cbd7a0c3da6dc8b870f20.zip | |
AMDGPU: Add ds append/consume builtins
llvm-svn: 352443
Diffstat (limited to 'clang/lib')
| -rw-r--r-- | clang/lib/CodeGen/CGBuiltin.cpp | 8 |
1 files changed, 8 insertions, 0 deletions
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 195b0c078bd..fa4d5126433 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -12478,6 +12478,14 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, 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")); |
