diff options
| author | Yaxun Liu <Yaxun.Liu@amd.com> | 2018-10-17 02:32:26 +0000 |
|---|---|---|
| committer | Yaxun Liu <Yaxun.Liu@amd.com> | 2018-10-17 02:32:26 +0000 |
| commit | aae1e87f4b80118eb0f4106fbf8fadd315f915f2 (patch) | |
| tree | 59188adfbff16dd38ba38ef3be0b53924b5d5fea /clang/lib/CodeGen/CGBuiltin.cpp | |
| parent | d2198a1bd01747557828e929b583a5916c5f5c64 (diff) | |
| download | bcm5719-llvm-aae1e87f4b80118eb0f4106fbf8fadd315f915f2.tar.gz bcm5719-llvm-aae1e87f4b80118eb0f4106fbf8fadd315f915f2.zip | |
AMDGPU: add __builtin_amdgcn_update_dpp
Emit llvm.amdgcn.update.dpp for both __builtin_amdgcn_mov_dpp and
__builtin_amdgcn_update_dpp. The first argument to
llvm.amdgcn.update.dpp will be undef for __builtin_amdgcn_mov_dpp.
Differential Revision: https://reviews.llvm.org/D52320
llvm-svn: 344665
Diffstat (limited to 'clang/lib/CodeGen/CGBuiltin.cpp')
| -rw-r--r-- | clang/lib/CodeGen/CGBuiltin.cpp | 14 |
1 files changed, 9 insertions, 5 deletions
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 70c50d37823..54d9fbba61d 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -11347,12 +11347,16 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, case AMDGPU::BI__builtin_amdgcn_ds_swizzle: return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_ds_swizzle); - case AMDGPU::BI__builtin_amdgcn_mov_dpp: { - llvm::SmallVector<llvm::Value *, 5> Args; - for (unsigned I = 0; I != 5; ++I) + case AMDGPU::BI__builtin_amdgcn_mov_dpp: + case AMDGPU::BI__builtin_amdgcn_update_dpp: { + llvm::SmallVector<llvm::Value *, 6> Args; + for (unsigned I = 0; I != E->getNumArgs(); ++I) Args.push_back(EmitScalarExpr(E->getArg(I))); - Value *F = CGM.getIntrinsic(Intrinsic::amdgcn_mov_dpp, - Args[0]->getType()); + assert(Args.size() == 5 || Args.size() == 6); + if (Args.size() == 5) + Args.insert(Args.begin(), llvm::UndefValue::get(Args[0]->getType())); + Value *F = + CGM.getIntrinsic(Intrinsic::amdgcn_update_dpp, Args[0]->getType()); return Builder.CreateCall(F, Args); } case AMDGPU::BI__builtin_amdgcn_div_fixup: |
