diff options
| author | Yaxun Liu <Yaxun.Liu@amd.com> | 2017-03-10 01:30:46 +0000 |
|---|---|---|
| committer | Yaxun Liu <Yaxun.Liu@amd.com> | 2017-03-10 01:30:46 +0000 |
| commit | 4d86799219c1027ff179824b96a213140b12d448 (patch) | |
| tree | ee95af6e47de2f51885556f92c4630faa73b408a /clang/lib/CodeGen/CGBuiltin.cpp | |
| parent | 65e2e6805a08f419271f916d01175d1684960be5 (diff) | |
| download | bcm5719-llvm-4d86799219c1027ff179824b96a213140b12d448.tar.gz bcm5719-llvm-4d86799219c1027ff179824b96a213140b12d448.zip | |
[AMDGPU] Add builtin functions readlane ds_permute mov_dpp
Differential Revision: https://reviews.llvm.org/D30551
llvm-svn: 297436
Diffstat (limited to 'clang/lib/CodeGen/CGBuiltin.cpp')
| -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 1c7ee2e6d49..85ed2b86777 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -8401,6 +8401,14 @@ 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) + Args.push_back(EmitScalarExpr(E->getArg(I))); + Value *F = CGM.getIntrinsic(Intrinsic::amdgcn_mov_dpp, + Args[0]->getType()); + return Builder.CreateCall(F, Args); + } case AMDGPU::BI__builtin_amdgcn_div_fixup: case AMDGPU::BI__builtin_amdgcn_div_fixupf: case AMDGPU::BI__builtin_amdgcn_div_fixuph: |
