aboutsummaryrefslogtreecommitdiff
path: root/clang/lib/CodeGen
diff options
context:
space:
mode:
Diffstat (limited to 'clang/lib/CodeGen')
-rw-r--r--clang/lib/CodeGen/CGCXXABI.h20
-rw-r--r--clang/lib/CodeGen/CGDebugInfo.cpp2
-rw-r--r--clang/lib/CodeGen/CGDecl.cpp15
-rw-r--r--clang/lib/CodeGen/CGExpr.cpp59
-rw-r--r--clang/lib/CodeGen/CGExprCXX.cpp20
-rw-r--r--clang/lib/CodeGen/CGHLSLRuntime.cpp37
-rw-r--r--clang/lib/CodeGen/CGHLSLRuntime.h4
-rw-r--r--clang/lib/CodeGen/CGObjCGNU.cpp2
-rw-r--r--clang/lib/CodeGen/CodeGenTBAA.cpp2
-rw-r--r--clang/lib/CodeGen/ItaniumCXXABI.cpp89
-rw-r--r--clang/lib/CodeGen/MicrosoftCXXABI.cpp8
-rw-r--r--clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp25
12 files changed, 206 insertions, 77 deletions
diff --git a/clang/lib/CodeGen/CGCXXABI.h b/clang/lib/CodeGen/CGCXXABI.h
index 96fe046..2dd320d 100644
--- a/clang/lib/CodeGen/CGCXXABI.h
+++ b/clang/lib/CodeGen/CGCXXABI.h
@@ -294,14 +294,22 @@ public:
Address Value,
QualType SrcRecordTy) = 0;
+ struct ExactDynamicCastInfo {
+ bool RequiresCastToPrimaryBase;
+ CharUnits Offset;
+ };
+
+ virtual std::optional<ExactDynamicCastInfo>
+ getExactDynamicCastInfo(QualType SrcRecordTy, QualType DestTy,
+ QualType DestRecordTy) = 0;
+
/// Emit a dynamic_cast from SrcRecordTy to DestRecordTy. The cast fails if
/// the dynamic type of Value is not exactly DestRecordTy.
- virtual llvm::Value *emitExactDynamicCast(CodeGenFunction &CGF, Address Value,
- QualType SrcRecordTy,
- QualType DestTy,
- QualType DestRecordTy,
- llvm::BasicBlock *CastSuccess,
- llvm::BasicBlock *CastFail) = 0;
+ virtual llvm::Value *emitExactDynamicCast(
+ CodeGenFunction &CGF, Address Value, QualType SrcRecordTy,
+ QualType DestTy, QualType DestRecordTy,
+ const ExactDynamicCastInfo &CastInfo, llvm::BasicBlock *CastSuccess,
+ llvm::BasicBlock *CastFail) = 0;
virtual bool EmitBadCastCall(CodeGenFunction &CGF) = 0;
diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp
index 5f50b1cc..2b469f2 100644
--- a/clang/lib/CodeGen/CGDebugInfo.cpp
+++ b/clang/lib/CodeGen/CGDebugInfo.cpp
@@ -2874,7 +2874,7 @@ static bool isDefinedInClangModule(const RecordDecl *RD) {
if (!Explicit && CXXDecl->getEnclosingNamespaceContext())
return false;
// This is a template, check the origin of the first member.
- if (CXXDecl->field_begin() == CXXDecl->field_end())
+ if (CXXDecl->fields().empty())
return TemplateKind == TSK_ExplicitInstantiationDeclaration;
if (!CXXDecl->field_begin()->isFromASTFile())
return false;
diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp
index 04f13c7..ff2dada 100644
--- a/clang/lib/CodeGen/CGDecl.cpp
+++ b/clang/lib/CodeGen/CGDecl.cpp
@@ -599,10 +599,11 @@ namespace {
llvm::Constant *CleanupFn;
const CGFunctionInfo &FnInfo;
const VarDecl &Var;
+ const CleanupAttr *Attribute;
CallCleanupFunction(llvm::Constant *CleanupFn, const CGFunctionInfo *Info,
- const VarDecl *Var)
- : CleanupFn(CleanupFn), FnInfo(*Info), Var(*Var) {}
+ const VarDecl *Var, const CleanupAttr *Attr)
+ : CleanupFn(CleanupFn), FnInfo(*Info), Var(*Var), Attribute(Attr) {}
void Emit(CodeGenFunction &CGF, Flags flags) override {
DeclRefExpr DRE(CGF.getContext(), const_cast<VarDecl *>(&Var), false,
@@ -624,8 +625,11 @@ namespace {
CallArgList Args;
Args.add(RValue::get(Arg),
CGF.getContext().getPointerType(Var.getType()));
- auto Callee = CGCallee::forDirect(CleanupFn);
- CGF.EmitCall(FnInfo, Callee, ReturnValueSlot(), Args);
+ GlobalDecl GD = GlobalDecl(Attribute->getFunctionDecl());
+ auto Callee = CGCallee::forDirect(CleanupFn, CGCalleeInfo(GD));
+ CGF.EmitCall(FnInfo, Callee, ReturnValueSlot(), Args,
+ /*callOrInvoke*/ nullptr, /*IsMustTail*/ false,
+ Attribute->getLoc());
}
};
} // end anonymous namespace
@@ -2231,7 +2235,8 @@ void CodeGenFunction::EmitAutoVarCleanups(const AutoVarEmission &emission) {
assert(F && "Could not find function!");
const CGFunctionInfo &Info = CGM.getTypes().arrangeFunctionDeclaration(FD);
- EHStack.pushCleanup<CallCleanupFunction>(NormalAndEHCleanup, F, &Info, &D);
+ EHStack.pushCleanup<CallCleanupFunction>(NormalAndEHCleanup, F, &Info, &D,
+ CA);
}
// If this is a block variable, call _Block_object_destroy
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index 5a3d4e4..ed35a05 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -3789,33 +3789,50 @@ void CodeGenFunction::EmitCheck(
Branch->setMetadata(llvm::LLVMContext::MD_prof, Node);
EmitBlock(Handlers);
+ // Clear arguments for the MinimalRuntime handler.
+ if (CGM.getCodeGenOpts().SanitizeMinimalRuntime) {
+ switch (CheckHandler) {
+ case SanitizerHandler::TypeMismatch:
+ // Pass value pointer only. It adds minimal overhead.
+ StaticArgs = {};
+ assert(DynamicArgs.size() == 1);
+ break;
+ default:
+ // No arguments for other checks.
+ StaticArgs = {};
+ DynamicArgs = {};
+ break;
+ }
+ }
+
// Handler functions take an i8* pointing to the (handler-specific) static
// information block, followed by a sequence of intptr_t arguments
// representing operand values.
SmallVector<llvm::Value *, 4> Args;
SmallVector<llvm::Type *, 4> ArgTypes;
- if (!CGM.getCodeGenOpts().SanitizeMinimalRuntime) {
- Args.reserve(DynamicArgs.size() + 1);
- ArgTypes.reserve(DynamicArgs.size() + 1);
-
- // Emit handler arguments and create handler function type.
- if (!StaticArgs.empty()) {
- llvm::Constant *Info = llvm::ConstantStruct::getAnon(StaticArgs);
- auto *InfoPtr = new llvm::GlobalVariable(
- CGM.getModule(), Info->getType(), false,
- llvm::GlobalVariable::PrivateLinkage, Info, "", nullptr,
- llvm::GlobalVariable::NotThreadLocal,
- CGM.getDataLayout().getDefaultGlobalsAddressSpace());
- InfoPtr->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
- CGM.getSanitizerMetadata()->disableSanitizerForGlobal(InfoPtr);
- Args.push_back(InfoPtr);
- ArgTypes.push_back(Args.back()->getType());
- }
- for (llvm::Value *DynamicArg : DynamicArgs) {
- Args.push_back(EmitCheckValue(DynamicArg));
- ArgTypes.push_back(IntPtrTy);
- }
+ Args.reserve(DynamicArgs.size() + 1);
+ ArgTypes.reserve(DynamicArgs.size() + 1);
+
+ // Emit handler arguments and create handler function type.
+ if (!StaticArgs.empty()) {
+ llvm::Constant *Info = llvm::ConstantStruct::getAnon(StaticArgs);
+ auto *InfoPtr = new llvm::GlobalVariable(
+ CGM.getModule(), Info->getType(),
+ // Non-constant global is used in a handler to deduplicate reports.
+ // TODO: change deduplication logic and make it constant.
+ /*isConstant=*/false, llvm::GlobalVariable::PrivateLinkage, Info, "",
+ nullptr, llvm::GlobalVariable::NotThreadLocal,
+ CGM.getDataLayout().getDefaultGlobalsAddressSpace());
+ InfoPtr->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
+ CGM.getSanitizerMetadata()->disableSanitizerForGlobal(InfoPtr);
+ Args.push_back(InfoPtr);
+ ArgTypes.push_back(Args.back()->getType());
+ }
+
+ for (llvm::Value *DynamicArg : DynamicArgs) {
+ Args.push_back(EmitCheckValue(DynamicArg));
+ ArgTypes.push_back(IntPtrTy);
}
llvm::FunctionType *FnType =
diff --git a/clang/lib/CodeGen/CGExprCXX.cpp b/clang/lib/CodeGen/CGExprCXX.cpp
index b8238a4..49d5d8a 100644
--- a/clang/lib/CodeGen/CGExprCXX.cpp
+++ b/clang/lib/CodeGen/CGExprCXX.cpp
@@ -373,7 +373,7 @@ RValue CodeGenFunction::EmitCXXMemberOrOperatorMemberCallExpr(
bool UseVirtualCall = CanUseVirtualCall && !DevirtualizedMethod;
if (const CXXDestructorDecl *Dtor = dyn_cast<CXXDestructorDecl>(CalleeDecl)) {
- assert(CE->arg_begin() == CE->arg_end() &&
+ assert(CE->arguments().empty() &&
"Destructor shouldn't have explicit parameters");
assert(ReturnValue.isNull() && "Destructor shouldn't have return value");
if (UseVirtualCall) {
@@ -2292,7 +2292,20 @@ llvm::Value *CodeGenFunction::EmitDynamicCast(Address ThisAddr,
bool IsExact = !IsDynamicCastToVoid &&
CGM.getCodeGenOpts().OptimizationLevel > 0 &&
DestRecordTy->getAsCXXRecordDecl()->isEffectivelyFinal() &&
- CGM.getCXXABI().shouldEmitExactDynamicCast(DestRecordTy);
+ CGM.getCXXABI().shouldEmitExactDynamicCast(DestRecordTy) &&
+ !getLangOpts().PointerAuthCalls;
+
+ std::optional<CGCXXABI::ExactDynamicCastInfo> ExactCastInfo;
+ if (IsExact) {
+ ExactCastInfo = CGM.getCXXABI().getExactDynamicCastInfo(SrcRecordTy, DestTy,
+ DestRecordTy);
+ if (!ExactCastInfo) {
+ llvm::Value *NullValue = EmitDynamicCastToNull(*this, DestTy);
+ if (!Builder.GetInsertBlock())
+ EmitBlock(createBasicBlock("dynamic_cast.unreachable"));
+ return NullValue;
+ }
+ }
// C++ [expr.dynamic.cast]p4:
// If the value of v is a null pointer value in the pointer case, the result
@@ -2321,7 +2334,8 @@ llvm::Value *CodeGenFunction::EmitDynamicCast(Address ThisAddr,
// If the destination type is effectively final, this pointer points to the
// right type if and only if its vptr has the right value.
Value = CGM.getCXXABI().emitExactDynamicCast(
- *this, ThisAddr, SrcRecordTy, DestTy, DestRecordTy, CastEnd, CastNull);
+ *this, ThisAddr, SrcRecordTy, DestTy, DestRecordTy, *ExactCastInfo,
+ CastEnd, CastNull);
} else {
assert(DestRecordTy->isRecordType() &&
"destination type must be a record type!");
diff --git a/clang/lib/CodeGen/CGHLSLRuntime.cpp b/clang/lib/CodeGen/CGHLSLRuntime.cpp
index a47d1cc..f64ac20 100644
--- a/clang/lib/CodeGen/CGHLSLRuntime.cpp
+++ b/clang/lib/CodeGen/CGHLSLRuntime.cpp
@@ -273,10 +273,14 @@ void CGHLSLRuntime::addBuffer(const HLSLBufferDecl *BufDecl) {
emitBufferGlobalsAndMetadata(BufDecl, BufGV);
// Initialize cbuffer from binding (implicit or explicit)
- HLSLResourceBindingAttr *RBA = BufDecl->getAttr<HLSLResourceBindingAttr>();
- assert(RBA &&
- "cbuffer/tbuffer should always have resource binding attribute");
- initializeBufferFromBinding(BufDecl, BufGV, RBA);
+ if (HLSLVkBindingAttr *VkBinding = BufDecl->getAttr<HLSLVkBindingAttr>()) {
+ initializeBufferFromBinding(BufDecl, BufGV, VkBinding);
+ } else {
+ HLSLResourceBindingAttr *RBA = BufDecl->getAttr<HLSLResourceBindingAttr>();
+ assert(RBA &&
+ "cbuffer/tbuffer should always have resource binding attribute");
+ initializeBufferFromBinding(BufDecl, BufGV, RBA);
+ }
}
llvm::TargetExtType *
@@ -593,6 +597,31 @@ static void initializeBuffer(CodeGenModule &CGM, llvm::GlobalVariable *GV,
CGM.AddCXXGlobalInit(InitResFunc);
}
+static Value *buildNameForResource(llvm::StringRef BaseName,
+ CodeGenModule &CGM) {
+ std::string Str(BaseName);
+ std::string GlobalName(Str + ".str");
+ return CGM.GetAddrOfConstantCString(Str, GlobalName.c_str()).getPointer();
+}
+
+void CGHLSLRuntime::initializeBufferFromBinding(const HLSLBufferDecl *BufDecl,
+ llvm::GlobalVariable *GV,
+ HLSLVkBindingAttr *VkBinding) {
+ assert(VkBinding && "expect a nonnull binding attribute");
+ llvm::Type *Int1Ty = llvm::Type::getInt1Ty(CGM.getLLVMContext());
+ auto *NonUniform = llvm::ConstantInt::get(Int1Ty, false);
+ auto *Index = llvm::ConstantInt::get(CGM.IntTy, 0);
+ auto *RangeSize = llvm::ConstantInt::get(CGM.IntTy, 1);
+ auto *Set = llvm::ConstantInt::get(CGM.IntTy, VkBinding->getSet());
+ auto *Binding = llvm::ConstantInt::get(CGM.IntTy, VkBinding->getBinding());
+ Value *Name = buildNameForResource(BufDecl->getName(), CGM);
+ llvm::Intrinsic::ID IntrinsicID =
+ CGM.getHLSLRuntime().getCreateHandleFromBindingIntrinsic();
+
+ SmallVector<Value *> Args{Set, Binding, RangeSize, Index, NonUniform, Name};
+ initializeBuffer(CGM, GV, IntrinsicID, Args);
+}
+
void CGHLSLRuntime::initializeBufferFromBinding(const HLSLBufferDecl *BufDecl,
llvm::GlobalVariable *GV,
HLSLResourceBindingAttr *RBA) {
diff --git a/clang/lib/CodeGen/CGHLSLRuntime.h b/clang/lib/CodeGen/CGHLSLRuntime.h
index 89d2aff8..31d1728 100644
--- a/clang/lib/CodeGen/CGHLSLRuntime.h
+++ b/clang/lib/CodeGen/CGHLSLRuntime.h
@@ -62,6 +62,7 @@ class VarDecl;
class ParmVarDecl;
class InitListExpr;
class HLSLBufferDecl;
+class HLSLVkBindingAttr;
class HLSLResourceBindingAttr;
class Type;
class RecordType;
@@ -168,6 +169,9 @@ private:
llvm::GlobalVariable *BufGV);
void initializeBufferFromBinding(const HLSLBufferDecl *BufDecl,
llvm::GlobalVariable *GV,
+ HLSLVkBindingAttr *VkBinding);
+ void initializeBufferFromBinding(const HLSLBufferDecl *BufDecl,
+ llvm::GlobalVariable *GV,
HLSLResourceBindingAttr *RBA);
llvm::Triple::ArchType getArch();
diff --git a/clang/lib/CodeGen/CGObjCGNU.cpp b/clang/lib/CodeGen/CGObjCGNU.cpp
index 8acf8d2..e4147de 100644
--- a/clang/lib/CodeGen/CGObjCGNU.cpp
+++ b/clang/lib/CodeGen/CGObjCGNU.cpp
@@ -1750,7 +1750,7 @@ class CGObjCGNUstep2 : public CGObjCGNUstep {
// struct objc_method_list *methods
// FIXME: Almost identical code is copied and pasted below for the
// class, but refactoring it cleanly requires C++14 generic lambdas.
- if (OID->classmeth_begin() == OID->classmeth_end())
+ if (OID->class_methods().empty())
metaclassFields.addNullPointer(PtrTy);
else {
SmallVector<ObjCMethodDecl*, 16> ClassMethods;
diff --git a/clang/lib/CodeGen/CodeGenTBAA.cpp b/clang/lib/CodeGen/CodeGenTBAA.cpp
index a02a009..90eafe2 100644
--- a/clang/lib/CodeGen/CodeGenTBAA.cpp
+++ b/clang/lib/CodeGen/CodeGenTBAA.cpp
@@ -439,7 +439,7 @@ CodeGenTBAA::CollectFields(uint64_t BaseOffset,
// TODO: Handle C++ base classes.
if (const CXXRecordDecl *Decl = dyn_cast<CXXRecordDecl>(RD))
- if (Decl->bases_begin() != Decl->bases_end())
+ if (!Decl->bases().empty())
return false;
const ASTRecordLayout &Layout = Context.getASTRecordLayout(RD);
diff --git a/clang/lib/CodeGen/ItaniumCXXABI.cpp b/clang/lib/CodeGen/ItaniumCXXABI.cpp
index f4a9946..5ffc1ed 100644
--- a/clang/lib/CodeGen/ItaniumCXXABI.cpp
+++ b/clang/lib/CodeGen/ItaniumCXXABI.cpp
@@ -226,6 +226,10 @@ public:
return hasUniqueVTablePointer(DestRecordTy);
}
+ std::optional<ExactDynamicCastInfo>
+ getExactDynamicCastInfo(QualType SrcRecordTy, QualType DestTy,
+ QualType DestRecordTy) override;
+
llvm::Value *emitDynamicCastCall(CodeGenFunction &CGF, Address Value,
QualType SrcRecordTy, QualType DestTy,
QualType DestRecordTy,
@@ -234,6 +238,7 @@ public:
llvm::Value *emitExactDynamicCast(CodeGenFunction &CGF, Address ThisAddr,
QualType SrcRecordTy, QualType DestTy,
QualType DestRecordTy,
+ const ExactDynamicCastInfo &CastInfo,
llvm::BasicBlock *CastSuccess,
llvm::BasicBlock *CastFail) override;
@@ -1681,10 +1686,11 @@ llvm::Value *ItaniumCXXABI::emitDynamicCastCall(
return Value;
}
-llvm::Value *ItaniumCXXABI::emitExactDynamicCast(
- CodeGenFunction &CGF, Address ThisAddr, QualType SrcRecordTy,
- QualType DestTy, QualType DestRecordTy, llvm::BasicBlock *CastSuccess,
- llvm::BasicBlock *CastFail) {
+std::optional<CGCXXABI::ExactDynamicCastInfo>
+ItaniumCXXABI::getExactDynamicCastInfo(QualType SrcRecordTy, QualType DestTy,
+ QualType DestRecordTy) {
+ assert(shouldEmitExactDynamicCast(DestRecordTy));
+
ASTContext &Context = getContext();
// Find all the inheritance paths.
@@ -1722,41 +1728,56 @@ llvm::Value *ItaniumCXXABI::emitExactDynamicCast(
if (!Offset)
Offset = PathOffset;
else if (Offset != PathOffset) {
- // Base appears in at least two different places. Find the most-derived
- // object and see if it's a DestDecl. Note that the most-derived object
- // must be at least as aligned as this base class subobject, and must
- // have a vptr at offset 0.
- ThisAddr = Address(emitDynamicCastToVoid(CGF, ThisAddr, SrcRecordTy),
- CGF.VoidPtrTy, ThisAddr.getAlignment());
- SrcDecl = DestDecl;
- Offset = CharUnits::Zero();
- break;
+ // Base appears in at least two different places.
+ return ExactDynamicCastInfo{/*RequiresCastToPrimaryBase=*/true,
+ CharUnits::Zero()};
}
}
+ if (!Offset)
+ return std::nullopt;
+ return ExactDynamicCastInfo{/*RequiresCastToPrimaryBase=*/false, *Offset};
+}
- if (!Offset) {
- // If there are no public inheritance paths, the cast always fails.
- CGF.EmitBranch(CastFail);
- return llvm::PoisonValue::get(CGF.VoidPtrTy);
- }
+llvm::Value *ItaniumCXXABI::emitExactDynamicCast(
+ CodeGenFunction &CGF, Address ThisAddr, QualType SrcRecordTy,
+ QualType DestTy, QualType DestRecordTy,
+ const ExactDynamicCastInfo &ExactCastInfo, llvm::BasicBlock *CastSuccess,
+ llvm::BasicBlock *CastFail) {
+ const CXXRecordDecl *SrcDecl = SrcRecordTy->getAsCXXRecordDecl();
+ const CXXRecordDecl *DestDecl = DestRecordTy->getAsCXXRecordDecl();
+
+ llvm::Value *VTable = nullptr;
+ if (ExactCastInfo.RequiresCastToPrimaryBase) {
+ // Base appears in at least two different places. Find the most-derived
+ // object and see if it's a DestDecl. Note that the most-derived object
+ // must be at least as aligned as this base class subobject, and must
+ // have a vptr at offset 0.
+ llvm::Value *PrimaryBase =
+ emitDynamicCastToVoid(CGF, ThisAddr, SrcRecordTy);
+ ThisAddr = Address(PrimaryBase, CGF.VoidPtrTy, ThisAddr.getAlignment());
+ SrcDecl = DestDecl;
+ Address VTablePtrPtr = ThisAddr.withElementType(CGF.VoidPtrPtrTy);
+ VTable = CGF.Builder.CreateLoad(VTablePtrPtr, "vtable");
+ } else
+ VTable = CGF.GetVTablePtr(ThisAddr, CGF.UnqualPtrTy, SrcDecl);
// Compare the vptr against the expected vptr for the destination type at
- // this offset. Note that we do not know what type ThisAddr points to in
- // the case where the derived class multiply inherits from the base class
- // so we can't use GetVTablePtr, so we load the vptr directly instead.
- llvm::Instruction *VPtr = CGF.Builder.CreateLoad(
- ThisAddr.withElementType(CGF.VoidPtrPtrTy), "vtable");
- CGM.DecorateInstructionWithTBAA(
- VPtr, CGM.getTBAAVTablePtrAccessInfo(CGF.VoidPtrPtrTy));
- llvm::Value *Success = CGF.Builder.CreateICmpEQ(
- VPtr, getVTableAddressPoint(BaseSubobject(SrcDecl, *Offset), DestDecl));
- llvm::Value *Result = ThisAddr.emitRawPointer(CGF);
- if (!Offset->isZero())
- Result = CGF.Builder.CreateInBoundsGEP(
- CGF.CharTy, Result,
- {llvm::ConstantInt::get(CGF.PtrDiffTy, -Offset->getQuantity())});
+ // this offset.
+ llvm::Constant *ExpectedVTable = getVTableAddressPoint(
+ BaseSubobject(SrcDecl, ExactCastInfo.Offset), DestDecl);
+ llvm::Value *Success = CGF.Builder.CreateICmpEQ(VTable, ExpectedVTable);
+ llvm::Value *AdjustedThisPtr = ThisAddr.emitRawPointer(CGF);
+
+ if (!ExactCastInfo.Offset.isZero()) {
+ CharUnits::QuantityType Offset = ExactCastInfo.Offset.getQuantity();
+ llvm::Constant *OffsetConstant =
+ llvm::ConstantInt::get(CGF.PtrDiffTy, -Offset);
+ AdjustedThisPtr = CGF.Builder.CreateInBoundsGEP(CGF.CharTy, AdjustedThisPtr,
+ OffsetConstant);
+ }
+
CGF.Builder.CreateCondBr(Success, CastSuccess, CastFail);
- return Result;
+ return AdjustedThisPtr;
}
llvm::Value *ItaniumCXXABI::emitDynamicCastToVoid(CodeGenFunction &CGF,
@@ -2262,7 +2283,7 @@ llvm::Value *ItaniumCXXABI::EmitVirtualDestructorCall(
auto *CE = dyn_cast<const CXXMemberCallExpr *>(E);
auto *D = dyn_cast<const CXXDeleteExpr *>(E);
assert((CE != nullptr) ^ (D != nullptr));
- assert(CE == nullptr || CE->arg_begin() == CE->arg_end());
+ assert(CE == nullptr || CE->arguments().empty());
assert(DtorType == Dtor_Deleting || DtorType == Dtor_Complete);
GlobalDecl GD(Dtor, DtorType);
diff --git a/clang/lib/CodeGen/MicrosoftCXXABI.cpp b/clang/lib/CodeGen/MicrosoftCXXABI.cpp
index a181559..e8d2451 100644
--- a/clang/lib/CodeGen/MicrosoftCXXABI.cpp
+++ b/clang/lib/CodeGen/MicrosoftCXXABI.cpp
@@ -158,9 +158,15 @@ public:
// TODO: Add support for exact dynamic_casts.
return false;
}
+ std::optional<ExactDynamicCastInfo>
+ getExactDynamicCastInfo(QualType SrcRecordTy, QualType DestTy,
+ QualType DestRecordTy) override {
+ llvm_unreachable("unsupported");
+ }
llvm::Value *emitExactDynamicCast(CodeGenFunction &CGF, Address Value,
QualType SrcRecordTy, QualType DestTy,
QualType DestRecordTy,
+ const ExactDynamicCastInfo &CastInfo,
llvm::BasicBlock *CastSuccess,
llvm::BasicBlock *CastFail) override {
llvm_unreachable("unsupported");
@@ -2000,7 +2006,7 @@ llvm::Value *MicrosoftCXXABI::EmitVirtualDestructorCall(
auto *CE = dyn_cast<const CXXMemberCallExpr *>(E);
auto *D = dyn_cast<const CXXDeleteExpr *>(E);
assert((CE != nullptr) ^ (D != nullptr));
- assert(CE == nullptr || CE->arg_begin() == CE->arg_end());
+ assert(CE == nullptr || CE->arguments().empty());
assert(DtorType == Dtor_Deleting || DtorType == Dtor_Complete);
// We have only one destructor in the vftable but can get both behaviors
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index d5b9299..dad1f95 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -894,6 +894,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
case AMDGPU::BI__builtin_amdgcn_wmma_f32_32x16x128_f4:
case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4:
case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4:
+ case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_32x16x128_f4:
+ case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_32x16x128_f4:
case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_f16:
case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_bf16:
case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x64_f16:
@@ -1172,6 +1174,14 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
ArgsForMatchingMatrixTypes = {3, 0, 1};
BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_32x16x128_f4;
break;
+ case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_32x16x128_f4:
+ ArgsForMatchingMatrixTypes = {3, 0, 1};
+ BuiltinWMMAOp = Intrinsic::amdgcn_wmma_scale_f32_32x16x128_f4;
+ break;
+ case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_32x16x128_f4:
+ ArgsForMatchingMatrixTypes = {3, 0, 1};
+ BuiltinWMMAOp = Intrinsic::amdgcn_wmma_scale16_f32_32x16x128_f4;
+ break;
case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_f16:
ArgsForMatchingMatrixTypes = {4, 1, 3, 5};
BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x64_f16;
@@ -1497,6 +1507,21 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
F, {EmitScalarExpr(E->getArg(0)), EmitScalarExpr(E->getArg(1)),
EmitScalarExpr(E->getArg(2)), EmitScalarExpr(E->getArg(3))});
}
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32:
+ return emitBuiltinWithOneOverloadedType<5>(
+ *this, E, Intrinsic::amdgcn_raw_ptr_buffer_atomic_add);
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32:
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16:
+ return emitBuiltinWithOneOverloadedType<5>(
+ *this, E, Intrinsic::amdgcn_raw_ptr_buffer_atomic_fadd);
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32:
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64:
+ return emitBuiltinWithOneOverloadedType<5>(
+ *this, E, Intrinsic::amdgcn_raw_ptr_buffer_atomic_fmin);
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32:
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64:
+ return emitBuiltinWithOneOverloadedType<5>(
+ *this, E, Intrinsic::amdgcn_raw_ptr_buffer_atomic_fmax);
case AMDGPU::BI__builtin_amdgcn_s_prefetch_data:
return emitBuiltinWithOneOverloadedType<2>(
*this, E, Intrinsic::amdgcn_s_prefetch_data);
generated by cgit v1.1 at 2025-08-07 14:27:34 +0000