diff options
Diffstat (limited to 'contrib/llvm-project/clang/lib/CodeGen/CGBuiltin.cpp')
| -rw-r--r-- | contrib/llvm-project/clang/lib/CodeGen/CGBuiltin.cpp | 45 |
1 files changed, 42 insertions, 3 deletions
diff --git a/contrib/llvm-project/clang/lib/CodeGen/CGBuiltin.cpp b/contrib/llvm-project/clang/lib/CodeGen/CGBuiltin.cpp index 998fcc3af581..7ef764b8e1ac 100644 --- a/contrib/llvm-project/clang/lib/CodeGen/CGBuiltin.cpp +++ b/contrib/llvm-project/clang/lib/CodeGen/CGBuiltin.cpp @@ -1019,7 +1019,7 @@ CodeGenFunction::emitFlexibleArrayMemberSize(const Expr *E, unsigned Type, FAMSize = Builder.CreateIntCast(FAMSize, ResType, IsSigned); Value *Res = FAMSize; - if (const auto *DRE = dyn_cast<DeclRefExpr>(Base)) { + if (isa<DeclRefExpr>(Base)) { // The whole struct is specificed in the __bdos. const ASTRecordLayout &Layout = Ctx.getASTRecordLayout(OuterRD); @@ -10056,7 +10056,7 @@ CodeGenFunction::getSVEOverloadTypes(const SVETypeFlags &TypeFlags, llvm::Type *DefaultType = getSVEType(TypeFlags); - if (TypeFlags.isOverloadWhile()) + if (TypeFlags.isOverloadWhileOrMultiVecCvt()) return {DefaultType, Ops[1]->getType()}; if (TypeFlags.isOverloadWhileRW()) @@ -13287,7 +13287,7 @@ Value *CodeGenFunction::EmitBPFBuiltinExpr(unsigned BuiltinID, const auto *DR = cast<DeclRefExpr>(CE->getSubExpr()); const auto *Enumerator = cast<EnumConstantDecl>(DR->getDecl()); - auto &InitVal = Enumerator->getInitVal(); + auto InitVal = Enumerator->getInitVal(); std::string InitValStr; if (InitVal.isNegative() || InitVal > uint64_t(INT64_MAX)) InitValStr = std::to_string(InitVal.getSExtValue()); @@ -18178,6 +18178,45 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, llvm::Function *F = CGM.getIntrinsic(IID, {ArgTy}); return Builder.CreateCall(F, {Addr, Val, ZeroI32, ZeroI32, ZeroI1}); } + case AMDGPU::BI__builtin_amdgcn_global_load_tr_i32: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v2i32: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v4f16: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v4i16: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v8f16: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v8i16: { + + llvm::Type *ArgTy; + switch (BuiltinID) { + case AMDGPU::BI__builtin_amdgcn_global_load_tr_i32: + ArgTy = llvm::Type::getInt32Ty(getLLVMContext()); + break; + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v2i32: + ArgTy = llvm::FixedVectorType::get( + llvm::Type::getInt32Ty(getLLVMContext()), 2); + break; + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v4f16: + ArgTy = llvm::FixedVectorType::get( + llvm::Type::getHalfTy(getLLVMContext()), 4); + break; + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v4i16: + ArgTy = llvm::FixedVectorType::get( + llvm::Type::getInt16Ty(getLLVMContext()), 4); + break; + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v8f16: + ArgTy = llvm::FixedVectorType::get( + llvm::Type::getHalfTy(getLLVMContext()), 8); + break; + case AMDGPU::BI__builtin_amdgcn_global_load_tr_v8i16: + ArgTy = llvm::FixedVectorType::get( + llvm::Type::getInt16Ty(getLLVMContext()), 8); + break; + } + + llvm::Value *Addr = EmitScalarExpr(E->getArg(0)); + llvm::Function *F = + CGM.getIntrinsic(Intrinsic::amdgcn_global_load_tr, {ArgTy}); + return Builder.CreateCall(F, {Addr}); + } case AMDGPU::BI__builtin_amdgcn_read_exec: return EmitAMDGCNBallotForExec(*this, E, Int64Ty, Int64Ty, false); case AMDGPU::BI__builtin_amdgcn_read_exec_lo: |