diff options
Diffstat (limited to 'clang/lib/CodeGen')
-rw-r--r-- | clang/lib/CodeGen/ABIInfo.cpp | 4 | ||||
-rw-r--r-- | clang/lib/CodeGen/CGBuiltin.cpp | 16 | ||||
-rw-r--r-- | clang/lib/CodeGen/CGCall.cpp | 38 | ||||
-rw-r--r-- | clang/lib/CodeGen/CGCoroutine.cpp | 8 | ||||
-rw-r--r-- | clang/lib/CodeGen/CGDebugInfo.cpp | 31 | ||||
-rw-r--r-- | clang/lib/CodeGen/CGExprCXX.cpp | 27 | ||||
-rw-r--r-- | clang/lib/CodeGen/CGObjCMac.cpp | 2 | ||||
-rw-r--r-- | clang/lib/CodeGen/CGOpenMPRuntime.cpp | 157 | ||||
-rw-r--r-- | clang/lib/CodeGen/CGStmt.cpp | 11 | ||||
-rw-r--r-- | clang/lib/CodeGen/CodeGenFunction.cpp | 3 | ||||
-rw-r--r-- | clang/lib/CodeGen/CodeGenModule.cpp | 7 | ||||
-rw-r--r-- | clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 42 | ||||
-rw-r--r-- | clang/lib/CodeGen/TargetBuiltins/ARM.cpp | 4 | ||||
-rw-r--r-- | clang/lib/CodeGen/TargetBuiltins/WebAssembly.cpp | 61 |
14 files changed, 290 insertions, 121 deletions
diff --git a/clang/lib/CodeGen/ABIInfo.cpp b/clang/lib/CodeGen/ABIInfo.cpp index d981d69..3ef430e1 100644 --- a/clang/lib/CodeGen/ABIInfo.cpp +++ b/clang/lib/CodeGen/ABIInfo.cpp @@ -218,8 +218,8 @@ void ABIInfo::appendAttributeMangling(StringRef AttrStr, // only have "+" prefixes here. assert(LHS.starts_with("+") && RHS.starts_with("+") && "Features should always have a prefix."); - return TI.getFMVPriority({LHS.substr(1)}) > - TI.getFMVPriority({RHS.substr(1)}); + return TI.getFMVPriority({LHS.substr(1)}) + .ugt(TI.getFMVPriority({RHS.substr(1)})); }); bool IsFirst = true; diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 5f2eb76..3f784fc 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -4108,6 +4108,22 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID, return RValue::get(Result); } + case Builtin::BI__builtin_elementwise_maximumnum: { + Value *Op0 = EmitScalarExpr(E->getArg(0)); + Value *Op1 = EmitScalarExpr(E->getArg(1)); + Value *Result = Builder.CreateBinaryIntrinsic( + Intrinsic::maximumnum, Op0, Op1, nullptr, "elt.maximumnum"); + return RValue::get(Result); + } + + case Builtin::BI__builtin_elementwise_minimumnum: { + Value *Op0 = EmitScalarExpr(E->getArg(0)); + Value *Op1 = EmitScalarExpr(E->getArg(1)); + Value *Result = Builder.CreateBinaryIntrinsic( + Intrinsic::minimumnum, Op0, Op1, nullptr, "elt.minimumnum"); + return RValue::get(Result); + } + case Builtin::BI__builtin_reduce_max: { auto GetIntrinsicID = [this](QualType QT) { if (auto *VecTy = QT->getAs<VectorType>()) diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 19d8ba2..d9bd443 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -214,7 +214,7 @@ static void appendParameterTypes( for (unsigned I = 0, E = FPT->getNumParams(); I != E; ++I) { prefix.push_back(FPT->getParamType(I)); if (ExtInfos[I].hasPassObjectSize()) - prefix.push_back(CGT.getContext().getSizeType()); + prefix.push_back(CGT.getContext().getCanonicalSizeType()); } addExtParameterInfosForCall(paramInfos, FPT.getTypePtr(), PrefixSize, @@ -2852,20 +2852,28 @@ void CodeGenModule::ConstructAttributeList(StringRef Name, if (AI.getInReg()) Attrs.addAttribute(llvm::Attribute::InReg); - // Depending on the ABI, this may be either a byval or a dead_on_return - // argument. - if (AI.getIndirectByVal()) { - Attrs.addByValAttr(getTypes().ConvertTypeForMem(ParamType)); - } else { - // Add dead_on_return when the object's lifetime ends in the callee. - // This includes trivially-destructible objects, as well as objects - // whose destruction / clean-up is carried out within the callee (e.g., - // Obj-C ARC-managed structs, MSVC callee-destroyed objects). - if (!ParamType.isDestructedType() || !ParamType->isRecordType() || - ParamType->castAs<RecordType>() - ->getDecl() - ->isParamDestroyedInCallee()) - Attrs.addAttribute(llvm::Attribute::DeadOnReturn); + // HLSL out and inout parameters must not be marked with ByVal or + // DeadOnReturn attributes because stores to these parameters by the + // callee are visible to the caller. + if (auto ParamABI = FI.getExtParameterInfo(ArgNo).getABI(); + ParamABI != ParameterABI::HLSLOut && + ParamABI != ParameterABI::HLSLInOut) { + + // Depending on the ABI, this may be either a byval or a dead_on_return + // argument. + if (AI.getIndirectByVal()) { + Attrs.addByValAttr(getTypes().ConvertTypeForMem(ParamType)); + } else { + // Add dead_on_return when the object's lifetime ends in the callee. + // This includes trivially-destructible objects, as well as objects + // whose destruction / clean-up is carried out within the callee + // (e.g., Obj-C ARC-managed structs, MSVC callee-destroyed objects). + if (!ParamType.isDestructedType() || !ParamType->isRecordType() || + ParamType->castAs<RecordType>() + ->getDecl() + ->isParamDestroyedInCallee()) + Attrs.addAttribute(llvm::Attribute::DeadOnReturn); + } } auto *Decl = ParamType->getAsRecordDecl(); diff --git a/clang/lib/CodeGen/CGCoroutine.cpp b/clang/lib/CodeGen/CGCoroutine.cpp index 117ef3d..5ee9089 100644 --- a/clang/lib/CodeGen/CGCoroutine.cpp +++ b/clang/lib/CodeGen/CGCoroutine.cpp @@ -1006,15 +1006,15 @@ RValue CodeGenFunction::EmitCoroutineIntrinsic(const CallExpr *E, } case llvm::Intrinsic::coro_size: { auto &Context = getContext(); - CanQualType SizeTy = Context.getSizeType(); - llvm::IntegerType *T = Builder.getIntNTy(Context.getTypeSize(SizeTy)); + llvm::IntegerType *T = + Builder.getIntNTy(Context.getTypeSize(Context.getSizeType())); llvm::Function *F = CGM.getIntrinsic(llvm::Intrinsic::coro_size, T); return RValue::get(Builder.CreateCall(F)); } case llvm::Intrinsic::coro_align: { auto &Context = getContext(); - CanQualType SizeTy = Context.getSizeType(); - llvm::IntegerType *T = Builder.getIntNTy(Context.getTypeSize(SizeTy)); + llvm::IntegerType *T = + Builder.getIntNTy(Context.getTypeSize(Context.getSizeType())); llvm::Function *F = CGM.getIntrinsic(llvm::Intrinsic::coro_align, T); return RValue::get(Builder.CreateCall(F)); } diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp index 446cf8d..a371b67 100644 --- a/clang/lib/CodeGen/CGDebugInfo.cpp +++ b/clang/lib/CodeGen/CGDebugInfo.cpp @@ -58,13 +58,6 @@ using namespace clang; using namespace clang::CodeGen; -// TODO: consider deprecating ClArrayBoundsPseudoFn; functionality is subsumed -// by -fsanitize-annotate-debug-info -static llvm::cl::opt<bool> ClArrayBoundsPseudoFn( - "array-bounds-pseudofn", llvm::cl::Hidden, llvm::cl::Optional, - llvm::cl::desc("Emit debug info that places array-bounds instrumentation " - "in an inline function called __ubsan_check_array_bounds.")); - static uint32_t getTypeAlignIfRequired(const Type *Ty, const ASTContext &Ctx) { auto TI = Ctx.getTypeInfo(Ty); if (TI.isAlignRequired()) @@ -4052,7 +4045,8 @@ llvm::DIType *CGDebugInfo::CreateTypeNode(QualType Ty, llvm::DIFile *Unit) { return CreateType(cast<HLSLAttributedResourceType>(Ty), Unit); case Type::HLSLInlineSpirv: return CreateType(cast<HLSLInlineSpirvType>(Ty), Unit); - + case Type::PredefinedSugar: + return getOrCreateType(cast<PredefinedSugarType>(Ty)->desugar(), Unit); case Type::CountAttributed: case Type::Auto: case Type::Attributed: @@ -6068,11 +6062,10 @@ void CGDebugInfo::EmitPseudoVariable(CGBuilderTy &Builder, // ptr, in this case its debug info may not match the actual type of object // being used as in the next instruction, so we will need to emit a pseudo // variable for type-casted value. - auto DeclareTypeMatches = [&](auto *DbgDeclare) { + auto DeclareTypeMatches = [&](llvm::DbgVariableRecord *DbgDeclare) { return DbgDeclare->getVariable()->getType() == Type; }; - if (any_of(llvm::findDbgDeclares(Var), DeclareTypeMatches) || - any_of(llvm::findDVRDeclares(Var), DeclareTypeMatches)) + if (any_of(llvm::findDVRDeclares(Var), DeclareTypeMatches)) return; } @@ -6482,7 +6475,11 @@ llvm::DILocation *CodeGenFunction::SanitizerAnnotateDebugInfo( SanitizerHandler Handler) { llvm::DILocation *CheckDebugLoc = Builder.getCurrentDebugLocation(); auto *DI = getDebugInfo(); - if (!DI) + if (!DI || !CheckDebugLoc) + return CheckDebugLoc; + const auto &AnnotateDebugInfo = + CGM.getCodeGenOpts().SanitizeAnnotateDebugInfo; + if (AnnotateDebugInfo.empty()) return CheckDebugLoc; std::string Label; @@ -6491,14 +6488,8 @@ llvm::DILocation *CodeGenFunction::SanitizerAnnotateDebugInfo( else Label = SanitizerHandlerToCheckLabel(Handler); - for (auto Ord : Ordinals) { - // TODO: deprecate ClArrayBoundsPseudoFn - if (((ClArrayBoundsPseudoFn && Ord == SanitizerKind::SO_ArrayBounds) || - CGM.getCodeGenOpts().SanitizeAnnotateDebugInfo.has(Ord)) && - CheckDebugLoc) { - return DI->CreateSyntheticInlineAt(CheckDebugLoc, Label); - } - } + if (any_of(Ordinals, [&](auto Ord) { return AnnotateDebugInfo.has(Ord); })) + return DI->CreateSyntheticInlineAt(CheckDebugLoc, Label); return CheckDebugLoc; } diff --git a/clang/lib/CodeGen/CGExprCXX.cpp b/clang/lib/CodeGen/CGExprCXX.cpp index 359e30c..b8238a4 100644 --- a/clang/lib/CodeGen/CGExprCXX.cpp +++ b/clang/lib/CodeGen/CGExprCXX.cpp @@ -2146,30 +2146,9 @@ void CodeGenFunction::EmitCXXDeleteExpr(const CXXDeleteExpr *E) { return; } - // We might be deleting a pointer to array. If so, GEP down to the - // first non-array element. - // (this assumes that A(*)[3][7] is converted to [3 x [7 x %A]]*) - if (DeleteTy->isConstantArrayType()) { - llvm::Value *Zero = Builder.getInt32(0); - SmallVector<llvm::Value*,8> GEP; - - GEP.push_back(Zero); // point at the outermost array - - // For each layer of array type we're pointing at: - while (const ConstantArrayType *Arr - = getContext().getAsConstantArrayType(DeleteTy)) { - // 1. Unpeel the array type. - DeleteTy = Arr->getElementType(); - - // 2. GEP to the first element of the array. - GEP.push_back(Zero); - } - - Ptr = Builder.CreateInBoundsGEP(Ptr, GEP, ConvertTypeForMem(DeleteTy), - Ptr.getAlignment(), "del.first"); - } - - assert(ConvertTypeForMem(DeleteTy) == Ptr.getElementType()); + // We might be deleting a pointer to array. + DeleteTy = getContext().getBaseElementType(DeleteTy); + Ptr = Ptr.withElementType(ConvertTypeForMem(DeleteTy)); if (E->isArrayForm()) { EmitArrayDelete(*this, E, Ptr, DeleteTy); diff --git a/clang/lib/CodeGen/CGObjCMac.cpp b/clang/lib/CodeGen/CGObjCMac.cpp index 8e71a57..8c66176 100644 --- a/clang/lib/CodeGen/CGObjCMac.cpp +++ b/clang/lib/CodeGen/CGObjCMac.cpp @@ -285,7 +285,7 @@ public: SmallVector<CanQualType, 5> Params; Params.push_back(Ctx.VoidPtrTy); Params.push_back(Ctx.VoidPtrTy); - Params.push_back(Ctx.getSizeType()); + Params.push_back(Ctx.getCanonicalSizeType()); Params.push_back(Ctx.BoolTy); Params.push_back(Ctx.BoolTy); llvm::FunctionType *FTy = Types.GetFunctionType( diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index ce2dd4d..91237cf 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7080,6 +7080,110 @@ private: return ConstLength.getSExtValue() != 1; } + /// A helper class to copy structures with overlapped elements, i.e. those + /// which have mappings of both "s" and "s.mem". Consecutive elements that + /// are not explicitly copied have mapping nodes synthesized for them, + /// taking care to avoid generating zero-sized copies. + class CopyOverlappedEntryGaps { + CodeGenFunction &CGF; + MapCombinedInfoTy &CombinedInfo; + OpenMPOffloadMappingFlags Flags = OpenMPOffloadMappingFlags::OMP_MAP_NONE; + const ValueDecl *MapDecl = nullptr; + const Expr *MapExpr = nullptr; + Address BP = Address::invalid(); + bool IsNonContiguous = false; + uint64_t DimSize = 0; + // These elements track the position as the struct is iterated over + // (in order of increasing element address). + const RecordDecl *LastParent = nullptr; + uint64_t Cursor = 0; + unsigned LastIndex = -1u; + Address LB = Address::invalid(); + + public: + CopyOverlappedEntryGaps(CodeGenFunction &CGF, + MapCombinedInfoTy &CombinedInfo, + OpenMPOffloadMappingFlags Flags, + const ValueDecl *MapDecl, const Expr *MapExpr, + Address BP, Address LB, bool IsNonContiguous, + uint64_t DimSize) + : CGF(CGF), CombinedInfo(CombinedInfo), Flags(Flags), MapDecl(MapDecl), + MapExpr(MapExpr), BP(BP), IsNonContiguous(IsNonContiguous), + DimSize(DimSize), LB(LB) {} + + void processField( + const OMPClauseMappableExprCommon::MappableComponent &MC, + const FieldDecl *FD, + llvm::function_ref<LValue(CodeGenFunction &, const MemberExpr *)> + EmitMemberExprBase) { + const RecordDecl *RD = FD->getParent(); + const ASTRecordLayout &RL = CGF.getContext().getASTRecordLayout(RD); + uint64_t FieldOffset = RL.getFieldOffset(FD->getFieldIndex()); + uint64_t FieldSize = + CGF.getContext().getTypeSize(FD->getType().getCanonicalType()); + Address ComponentLB = Address::invalid(); + + if (FD->getType()->isLValueReferenceType()) { + const auto *ME = cast<MemberExpr>(MC.getAssociatedExpression()); + LValue BaseLVal = EmitMemberExprBase(CGF, ME); + ComponentLB = + CGF.EmitLValueForFieldInitialization(BaseLVal, FD).getAddress(); + } else { + ComponentLB = + CGF.EmitOMPSharedLValue(MC.getAssociatedExpression()).getAddress(); + } + + if (!LastParent) + LastParent = RD; + if (FD->getParent() == LastParent) { + if (FD->getFieldIndex() != LastIndex + 1) + copyUntilField(FD, ComponentLB); + } else { + LastParent = FD->getParent(); + if (((int64_t)FieldOffset - (int64_t)Cursor) > 0) + copyUntilField(FD, ComponentLB); + } + Cursor = FieldOffset + FieldSize; + LastIndex = FD->getFieldIndex(); + LB = CGF.Builder.CreateConstGEP(ComponentLB, 1); + } + + void copyUntilField(const FieldDecl *FD, Address ComponentLB) { + llvm::Value *ComponentLBPtr = ComponentLB.emitRawPointer(CGF); + llvm::Value *LBPtr = LB.emitRawPointer(CGF); + llvm::Value *Size = + CGF.Builder.CreatePtrDiff(CGF.Int8Ty, ComponentLBPtr, LBPtr); + copySizedChunk(LBPtr, Size); + } + + void copyUntilEnd(Address HB) { + if (LastParent) { + const ASTRecordLayout &RL = + CGF.getContext().getASTRecordLayout(LastParent); + if ((uint64_t)CGF.getContext().toBits(RL.getSize()) <= Cursor) + return; + } + llvm::Value *LBPtr = LB.emitRawPointer(CGF); + llvm::Value *Size = CGF.Builder.CreatePtrDiff( + CGF.Int8Ty, CGF.Builder.CreateConstGEP(HB, 1).emitRawPointer(CGF), + LBPtr); + copySizedChunk(LBPtr, Size); + } + + void copySizedChunk(llvm::Value *Base, llvm::Value *Size) { + CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr); + CombinedInfo.BasePointers.push_back(BP.emitRawPointer(CGF)); + CombinedInfo.DevicePtrDecls.push_back(nullptr); + CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None); + CombinedInfo.Pointers.push_back(Base); + CombinedInfo.Sizes.push_back( + CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true)); + CombinedInfo.Types.push_back(Flags); + CombinedInfo.Mappers.push_back(nullptr); + CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize : 1); + } + }; + /// Generate the base pointers, section pointers, sizes, map type bits, and /// user-defined mappers (all included in \a CombinedInfo) for the provided /// map type, map or motion modifiers, and expression components. @@ -7570,63 +7674,22 @@ private: getMapTypeBits(MapType, MapModifiers, MotionModifiers, IsImplicit, /*AddPtrFlag=*/false, /*AddIsTargetParamFlag=*/false, IsNonContiguous); - llvm::Value *Size = nullptr; + CopyOverlappedEntryGaps CopyGaps(CGF, CombinedInfo, Flags, MapDecl, + MapExpr, BP, LB, IsNonContiguous, + DimSize); // Do bitcopy of all non-overlapped structure elements. for (OMPClauseMappableExprCommon::MappableExprComponentListRef Component : OverlappedElements) { - Address ComponentLB = Address::invalid(); for (const OMPClauseMappableExprCommon::MappableComponent &MC : Component) { if (const ValueDecl *VD = MC.getAssociatedDeclaration()) { - const auto *FD = dyn_cast<FieldDecl>(VD); - if (FD && FD->getType()->isLValueReferenceType()) { - const auto *ME = - cast<MemberExpr>(MC.getAssociatedExpression()); - LValue BaseLVal = EmitMemberExprBase(CGF, ME); - ComponentLB = - CGF.EmitLValueForFieldInitialization(BaseLVal, FD) - .getAddress(); - } else { - ComponentLB = - CGF.EmitOMPSharedLValue(MC.getAssociatedExpression()) - .getAddress(); + if (const auto *FD = dyn_cast<FieldDecl>(VD)) { + CopyGaps.processField(MC, FD, EmitMemberExprBase); } - llvm::Value *ComponentLBPtr = ComponentLB.emitRawPointer(CGF); - llvm::Value *LBPtr = LB.emitRawPointer(CGF); - Size = CGF.Builder.CreatePtrDiff(CGF.Int8Ty, ComponentLBPtr, - LBPtr); - break; } } - assert(Size && "Failed to determine structure size"); - CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr); - CombinedInfo.BasePointers.push_back(BP.emitRawPointer(CGF)); - CombinedInfo.DevicePtrDecls.push_back(nullptr); - CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None); - CombinedInfo.Pointers.push_back(LB.emitRawPointer(CGF)); - CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast( - Size, CGF.Int64Ty, /*isSigned=*/true)); - CombinedInfo.Types.push_back(Flags); - CombinedInfo.Mappers.push_back(nullptr); - CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize - : 1); - LB = CGF.Builder.CreateConstGEP(ComponentLB, 1); } - CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr); - CombinedInfo.BasePointers.push_back(BP.emitRawPointer(CGF)); - CombinedInfo.DevicePtrDecls.push_back(nullptr); - CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None); - CombinedInfo.Pointers.push_back(LB.emitRawPointer(CGF)); - llvm::Value *LBPtr = LB.emitRawPointer(CGF); - Size = CGF.Builder.CreatePtrDiff( - CGF.Int8Ty, CGF.Builder.CreateConstGEP(HB, 1).emitRawPointer(CGF), - LBPtr); - CombinedInfo.Sizes.push_back( - CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true)); - CombinedInfo.Types.push_back(Flags); - CombinedInfo.Mappers.push_back(nullptr); - CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize - : 1); + CopyGaps.copyUntilEnd(HB); break; } llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression()); diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp index e065006..1a8c6f0 100644 --- a/clang/lib/CodeGen/CGStmt.cpp +++ b/clang/lib/CodeGen/CGStmt.cpp @@ -846,11 +846,13 @@ void CodeGenFunction::EmitGotoStmt(const GotoStmt &S) { if (HaveInsertPoint()) EmitStopPoint(&S); + ApplyAtomGroup Grp(getDebugInfo()); EmitBranchThroughCleanup(getJumpDestForLabel(S.getLabel())); } void CodeGenFunction::EmitIndirectGotoStmt(const IndirectGotoStmt &S) { + ApplyAtomGroup Grp(getDebugInfo()); if (const LabelDecl *Target = S.getConstantTarget()) { EmitBranchThroughCleanup(getJumpDestForLabel(Target)); return; @@ -869,6 +871,8 @@ void CodeGenFunction::EmitIndirectGotoStmt(const IndirectGotoStmt &S) { cast<llvm::PHINode>(IndGotoBB->begin())->addIncoming(V, CurBB); EmitBranch(IndGotoBB); + if (CurBB && CurBB->getTerminator()) + addInstToCurrentSourceAtom(CurBB->getTerminator(), nullptr); } void CodeGenFunction::EmitIfStmt(const IfStmt &S) { @@ -2672,6 +2676,9 @@ static void UpdateAsmCallInst(llvm::CallBase &Result, bool HasSideEffect, llvm::ConstantAsMetadata::get(Loc))); } + // Make inline-asm calls Key for the debug info feature Key Instructions. + CGF.addInstToNewSourceAtom(&Result, nullptr); + if (!NoConvergent && CGF.getLangOpts().assumeFunctionsAreConvergent()) // Conservatively, mark all inline asm blocks in CUDA or OpenCL as // convergent (meaning, they may call an intrinsically convergent op, such @@ -2750,6 +2757,7 @@ EmitAsmStores(CodeGenFunction &CGF, const AsmStmt &S, } } + ApplyAtomGroup Grp(CGF.getDebugInfo()); LValue Dest = ResultRegDests[i]; // ResultTypeRequiresCast elements correspond to the first // ResultTypeRequiresCast.size() elements of RegResults. @@ -2757,7 +2765,8 @@ EmitAsmStores(CodeGenFunction &CGF, const AsmStmt &S, unsigned Size = CGF.getContext().getTypeSize(ResultRegQualTys[i]); Address A = Dest.getAddress().withElementType(ResultRegTypes[i]); if (CGF.getTargetHooks().isScalarizableAsmOperand(CGF, TruncTy)) { - Builder.CreateStore(Tmp, A); + llvm::StoreInst *S = Builder.CreateStore(Tmp, A); + CGF.addInstToCurrentSourceAtom(S, S->getValueOperand()); continue; } diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 0fda31c..ab345a5 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -720,7 +720,7 @@ static bool matchesStlAllocatorFn(const Decl *D, const ASTContext &Ctx) { (MD->getNumParams() != 1 && MD->getNumParams() != 2)) return false; - if (MD->parameters()[0]->getType().getCanonicalType() != Ctx.getSizeType()) + if (!Ctx.hasSameType(MD->parameters()[0]->getType(), Ctx.getSizeType())) return false; if (MD->getNumParams() == 2) { @@ -2491,6 +2491,7 @@ void CodeGenFunction::EmitVariablyModifiedType(QualType type) { case Type::ObjCObjectPointer: case Type::BitInt: case Type::HLSLInlineSpirv: + case Type::PredefinedSugar: llvm_unreachable("type class is never variably-modified!"); case Type::Elaborated: diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 236cc3d..834b1c0 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4418,8 +4418,9 @@ void CodeGenModule::EmitGlobalDefinition(GlobalDecl GD, llvm::GlobalValue *GV) { static void ReplaceUsesOfNonProtoTypeWithRealFunction(llvm::GlobalValue *Old, llvm::Function *NewFn); -static uint64_t getFMVPriority(const TargetInfo &TI, - const CodeGenFunction::FMVResolverOption &RO) { +static llvm::APInt +getFMVPriority(const TargetInfo &TI, + const CodeGenFunction::FMVResolverOption &RO) { llvm::SmallVector<StringRef, 8> Features{RO.Features}; if (RO.Architecture) Features.push_back(*RO.Architecture); @@ -4544,7 +4545,7 @@ void CodeGenModule::emitMultiVersionFunctions() { llvm::stable_sort( Options, [&TI](const CodeGenFunction::FMVResolverOption &LHS, const CodeGenFunction::FMVResolverOption &RHS) { - return getFMVPriority(TI, LHS) > getFMVPriority(TI, RHS); + return getFMVPriority(TI, LHS).ugt(getFMVPriority(TI, RHS)); }); CodeGenFunction CGF(*this); CGF.EmitMultiVersionResolver(ResolverFunc, Options); diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index ee736a2..70f510a 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -274,7 +274,7 @@ llvm::Value *CodeGenFunction::EmitScalarOrConstFoldImmArg(unsigned ICEArguments, void CodeGenFunction::AddAMDGPUFenceAddressSpaceMMRA(llvm::Instruction *Inst, const CallExpr *E) { - constexpr const char *Tag = "amdgpu-as"; + constexpr const char *Tag = "amdgpu-synchronize-as"; LLVMContext &Ctx = Inst->getContext(); SmallVector<MMRAMetadata::TagT, 3> MMRAs; @@ -633,6 +633,41 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy}); return Builder.CreateCall(F, {Addr}); } + case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32: + case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64: + case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128: + case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32: + case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64: + case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128: { + + Intrinsic::ID IID; + switch (BuiltinID) { + case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32: + IID = Intrinsic::amdgcn_global_load_monitor_b32; + break; + case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64: + IID = Intrinsic::amdgcn_global_load_monitor_b64; + break; + case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128: + IID = Intrinsic::amdgcn_global_load_monitor_b128; + break; + case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32: + IID = Intrinsic::amdgcn_flat_load_monitor_b32; + break; + case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64: + IID = Intrinsic::amdgcn_flat_load_monitor_b64; + break; + case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128: + IID = Intrinsic::amdgcn_flat_load_monitor_b128; + break; + } + + llvm::Type *LoadTy = ConvertType(E->getType()); + llvm::Value *Addr = EmitScalarExpr(E->getArg(0)); + llvm::Value *Val = EmitScalarExpr(E->getArg(1)); + llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy}); + return Builder.CreateCall(F, {Addr, Val}); + } case AMDGPU::BI__builtin_amdgcn_load_to_lds: { // Should this have asan instrumentation? return emitBuiltinWithOneOverloadedType<5>(*this, E, @@ -855,6 +890,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8: case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_bf8: case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8: + case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_f8f6f4: case AMDGPU::BI__builtin_amdgcn_wmma_f32_32x16x128_f4: case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_f16: case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_bf16: @@ -1118,6 +1154,10 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, ArgsForMatchingMatrixTypes = {4, 1}; BuiltinWMMAOp = Intrinsic::amdgcn_wmma_i32_16x16x64_iu8; break; + case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_f8f6f4: + ArgsForMatchingMatrixTypes = {5, 1, 3}; + BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x128_f8f6f4; + break; case AMDGPU::BI__builtin_amdgcn_wmma_f32_32x16x128_f4: ArgsForMatchingMatrixTypes = {3, 0, 1}; BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_32x16x128_f4; diff --git a/clang/lib/CodeGen/TargetBuiltins/ARM.cpp b/clang/lib/CodeGen/TargetBuiltins/ARM.cpp index 7e6a47f..2e6b4b3 100644 --- a/clang/lib/CodeGen/TargetBuiltins/ARM.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/ARM.cpp @@ -8112,7 +8112,7 @@ Value *CodeGenFunction::EmitAArch64CpuSupports(const CallExpr *E) { llvm::Value * CodeGenFunction::EmitAArch64CpuSupports(ArrayRef<StringRef> FeaturesStrs) { - uint64_t FeaturesMask = llvm::AArch64::getCpuSupportsMask(FeaturesStrs); + llvm::APInt FeaturesMask = llvm::AArch64::getCpuSupportsMask(FeaturesStrs); Value *Result = Builder.getTrue(); if (FeaturesMask != 0) { // Get features from structure in runtime library @@ -8128,7 +8128,7 @@ CodeGenFunction::EmitAArch64CpuSupports(ArrayRef<StringRef> FeaturesStrs) { {ConstantInt::get(Int32Ty, 0), ConstantInt::get(Int32Ty, 0)}); Value *Features = Builder.CreateAlignedLoad(Int64Ty, CpuFeatures, CharUnits::fromQuantity(8)); - Value *Mask = Builder.getInt64(FeaturesMask); + Value *Mask = Builder.getInt(FeaturesMask.trunc(64)); Value *Bitset = Builder.CreateAnd(Features, Mask); Value *Cmp = Builder.CreateICmpEQ(Bitset, Mask); Result = Builder.CreateAnd(Result, Cmp); diff --git a/clang/lib/CodeGen/TargetBuiltins/WebAssembly.cpp b/clang/lib/CodeGen/TargetBuiltins/WebAssembly.cpp index b7fd70e..33a8d8f 100644 --- a/clang/lib/CodeGen/TargetBuiltins/WebAssembly.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/WebAssembly.cpp @@ -12,7 +12,10 @@ #include "CGBuiltin.h" #include "clang/Basic/TargetBuiltins.h" +#include "llvm/ADT/APInt.h" +#include "llvm/IR/Constants.h" #include "llvm/IR/IntrinsicsWebAssembly.h" +#include "llvm/Support/ErrorHandling.h" using namespace clang; using namespace CodeGen; @@ -218,6 +221,64 @@ Value *CodeGenFunction::EmitWebAssemblyBuiltinExpr(unsigned BuiltinID, Function *Callee = CGM.getIntrinsic(Intrinsic::wasm_ref_null_func); return Builder.CreateCall(Callee); } + case WebAssembly::BI__builtin_wasm_test_function_pointer_signature: { + Value *FuncRef = EmitScalarExpr(E->getArg(0)); + + // Get the function type from the argument's static type + QualType ArgType = E->getArg(0)->getType(); + const PointerType *PtrTy = ArgType->getAs<PointerType>(); + assert(PtrTy && "Sema should have ensured this is a function pointer"); + + const FunctionType *FuncTy = PtrTy->getPointeeType()->getAs<FunctionType>(); + assert(FuncTy && "Sema should have ensured this is a function pointer"); + + // In the llvm IR, we won't have access any more to the type of the function + // pointer so we need to insert this type information somehow. The + // @llvm.wasm.ref.test.func takes varargs arguments whose values are unused + // to indicate the type of the function to test for. See the test here: + // llvm/test/CodeGen/WebAssembly/ref-test-func.ll + // + // The format is: first we include the return types (since this is a C + // function pointer, there will be 0 or one of these) then a token type to + // indicate the boundary between return types and param types, then the + // param types. + + llvm::FunctionType *LLVMFuncTy = + cast<llvm::FunctionType>(ConvertType(QualType(FuncTy, 0))); + + unsigned NParams = LLVMFuncTy->getNumParams(); + std::vector<Value *> Args; + Args.reserve(NParams + 3); + // The only real argument is the FuncRef + Args.push_back(FuncRef); + + // Add the type information + auto addType = [this, &Args](llvm::Type *T) { + if (T->isVoidTy()) { + // Do nothing + } else if (T->isFloatingPointTy()) { + Args.push_back(ConstantFP::get(T, 0)); + } else if (T->isIntegerTy()) { + Args.push_back(ConstantInt::get(T, 0)); + } else if (T->isPointerTy()) { + Args.push_back(ConstantPointerNull::get(llvm::PointerType::get( + getLLVMContext(), T->getPointerAddressSpace()))); + } else { + // TODO: Handle reference types. For now, we reject them in Sema. + llvm_unreachable("Unhandled type"); + } + }; + + addType(LLVMFuncTy->getReturnType()); + // The token type indicates the boundary between return types and param + // types. + Args.push_back(PoisonValue::get(llvm::Type::getTokenTy(getLLVMContext()))); + for (unsigned i = 0; i < NParams; i++) { + addType(LLVMFuncTy->getParamType(i)); + } + Function *Callee = CGM.getIntrinsic(Intrinsic::wasm_ref_test_func); + return Builder.CreateCall(Callee, Args); + } case WebAssembly::BI__builtin_wasm_swizzle_i8x16: { Value *Src = EmitScalarExpr(E->getArg(0)); Value *Indices = EmitScalarExpr(E->getArg(1)); |