diff options
Diffstat (limited to 'clang/lib')
| -rw-r--r-- | clang/lib/Basic/Targets.cpp | 1 | ||||
| -rw-r--r-- | clang/lib/CodeGen/CGBuiltin.cpp | 8 |
2 files changed, 9 insertions, 0 deletions
diff --git a/clang/lib/Basic/Targets.cpp b/clang/lib/Basic/Targets.cpp index 6b0b9100b3e..e23ff3cc380 100644 --- a/clang/lib/Basic/Targets.cpp +++ b/clang/lib/Basic/Targets.cpp @@ -2387,6 +2387,7 @@ bool AMDGPUTargetInfo::initFeatureMap( case GK_GFX8: Features["s-memrealtime"] = true; Features["16-bit-insts"] = true; + Features["dpp"] = true; break; case GK_NONE: 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: |
