https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/223259
>From 14382ba8798ff994ac3de6de3f5d65c4c2ef9eb6 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Thu, 27 Aug 2026 08:10:29 +0200 Subject: [PATCH 1/2] [clang][SPIR-V][AMDGPU] Factor shared ABI classification into ABIInfoImpl Deduplicate the near-identical argument classification logic between AMDGPUABIInfo and AMDGCNSPIRVABIInfo into a shared AMDGPUABIInfoCommon template in ABIInfoImpl.h --- clang/lib/CodeGen/ABIInfoImpl.h | 193 +++++++++++++++++++++++++ clang/lib/CodeGen/Targets/AMDGPU.cpp | 203 +-------------------------- clang/lib/CodeGen/Targets/SPIR.cpp | 195 +------------------------ 3 files changed, 200 insertions(+), 391 deletions(-) diff --git a/clang/lib/CodeGen/ABIInfoImpl.h b/clang/lib/CodeGen/ABIInfoImpl.h index d9d79c6a55ddb..4649060c8e712 100644 --- a/clang/lib/CodeGen/ABIInfoImpl.h +++ b/clang/lib/CodeGen/ABIInfoImpl.h @@ -11,6 +11,7 @@ #include "ABIInfo.h" #include "CGCXXABI.h" +#include "llvm/IR/DerivedTypes.h" namespace clang::CodeGen { @@ -140,6 +141,198 @@ bool isEmptyRecordForLayout(const ASTContext &Context, QualType T); /// it exists. const Type *isSingleElementStruct(QualType T, ASTContext &Context); +/// Shared classification rules for AMDGPU and AMDGCN-SPIR-V, with \p Base as +/// the fallback ABIInfo for non-register-packed cases. +template <typename Base> class AMDGPUABIInfoCommon : public Base { +protected: + static constexpr unsigned MaxNumRegsForArgsRet = 16; // 16 32-bit registers + mutable unsigned NumRegsLeft = 0; + + using Base::Base; + + /// Estimate number of registers the type will use when passed in registers. + uint64_t numRegsForType(QualType Ty) const { + uint64_t NumRegs = 0; + + if (const VectorType *VT = Ty->template getAs<VectorType>()) { + // Compute from the number of elements. The reported size is based on + // the in-memory size, which includes the padding 4th element for + // 3-vectors. + QualType EltTy = VT->getElementType(); + uint64_t EltSize = this->getContext().getTypeSize(EltTy); + + // 16-bit element vectors should be passed as packed. + if (EltSize == 16) + return (VT->getNumElements() + 1) / 2; + + uint64_t EltNumRegs = (EltSize + 31) / 32; + return EltNumRegs * VT->getNumElements(); + } + + if (const auto *RD = Ty->getAsRecordDecl()) { + assert(!RD->hasFlexibleArrayMember()); + + for (const FieldDecl *Field : RD->fields()) + NumRegs += numRegsForType(Field->getType()); + + return NumRegs; + } + + return (this->getContext().getTypeSize(Ty) + 31) / 32; + } + + bool isHomogeneousAggregateBaseType(QualType Ty) const override { + return true; + } + + bool isHomogeneousAggregateSmallEnough(const Type *T, + uint64_t Members) const override { + uint32_t NumRegs = (this->getContext().getTypeSize(T) + 31) / 32; + + // Homogeneous Aggregates may occupy at most 16 registers. + return Members * NumRegs <= MaxNumRegsForArgsRet; + } + + // Coerce scalar pointer arguments from generic pointers to a fixed AS. + llvm::Type *coerceKernelArgumentType(llvm::Type *Ty, unsigned FromAS, + unsigned ToAS) const { + // Single value types. + auto *PtrTy = llvm::dyn_cast<llvm::PointerType>(Ty); + if (PtrTy && PtrTy->getAddressSpace() == FromAS) + return llvm::PointerType::get(Ty->getContext(), ToAS); + return Ty; + } + + ABIArgInfo classifyReturnType(QualType RetTy) const { + if (!isAggregateTypeForABI(RetTy) || + getRecordArgABI(RetTy, this->getCXXABI())) + return Base::classifyReturnType(RetTy); + + // Ignore empty structs/unions. + if (isEmptyRecord(this->getContext(), RetTy, true)) + return ABIArgInfo::getIgnore(); + + // Lower single-element structs to just return a regular value. + if (const Type *SeltTy = isSingleElementStruct(RetTy, this->getContext())) + return ABIArgInfo::getDirect(this->CGT.ConvertType(QualType(SeltTy, 0))); + + if (const auto *RD = RetTy->getAsRecordDecl(); + RD && RD->hasFlexibleArrayMember()) + return Base::classifyReturnType(RetTy); + + // Pack aggregates <= 4 bytes into single VGPR or pair. + uint64_t Size = this->getContext().getTypeSize(RetTy); + if (Size <= 16) + return ABIArgInfo::getDirect( + llvm::Type::getInt16Ty(this->getVMContext())); + + if (Size <= 32) + return ABIArgInfo::getDirect( + llvm::Type::getInt32Ty(this->getVMContext())); + + if (Size <= 64) { + llvm::Type *I32Ty = llvm::Type::getInt32Ty(this->getVMContext()); + return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); + } + + if (numRegsForType(RetTy) <= MaxNumRegsForArgsRet) + return ABIArgInfo::getDirect(); + + return Base::classifyReturnType(RetTy); + } + + ABIArgInfo classifyArgumentType(QualType Ty, bool Variadic) const { + assert(NumRegsLeft <= MaxNumRegsForArgsRet && + "register estimate underflow"); + + Ty = useFirstFieldIfTransparentUnion(Ty); + + if (Variadic) { + return ABIArgInfo::getDirect(/*T=*/nullptr, + /*Offset=*/0, + /*Padding=*/nullptr, + /*CanBeFlattened=*/false, + /*Align=*/0); + } + + if (!isAggregateTypeForABI(Ty)) { + ABIArgInfo ArgInfo = Base::classifyArgumentType(Ty); + if (!ArgInfo.isIndirect()) { + uint64_t NumRegs = numRegsForType(Ty); + NumRegsLeft -= std::min(NumRegs, uint64_t{NumRegsLeft}); + } + + return ArgInfo; + } + + // Records with non-trivial destructors/copy-constructors should not be + // passed by value. + if (auto RAA = getRecordArgABI(Ty, this->getCXXABI())) + return this->getNaturalAlignIndirect( + Ty, this->getDataLayout().getAllocaAddrSpace(), + RAA == CGCXXABI::RAA_DirectInMemory); + + // Ignore empty structs/unions. + if (isEmptyRecord(this->getContext(), Ty, true)) + return ABIArgInfo::getIgnore(); + + // Lower single-element structs to just pass a regular value. TODO: We + // could do reasonable-size multiple-element structs too, using + // getExpand(), though watch out for things like bitfields. + if (const Type *SeltTy = isSingleElementStruct(Ty, this->getContext())) + return ABIArgInfo::getDirect(this->CGT.ConvertType(QualType(SeltTy, 0))); + + if (const auto *RD = Ty->getAsRecordDecl(); + RD && RD->hasFlexibleArrayMember()) + return Base::classifyArgumentType(Ty); + + // Pack aggregates <= 8 bytes into single VGPR or pair. + uint64_t Size = this->getContext().getTypeSize(Ty); + if (Size <= 64) { + unsigned NumRegs = (Size + 31) / 32; + NumRegsLeft -= std::min(NumRegsLeft, NumRegs); + + if (Size <= 16) + return ABIArgInfo::getDirect( + llvm::Type::getInt16Ty(this->getVMContext())); + + if (Size <= 32) + return ABIArgInfo::getDirect( + llvm::Type::getInt32Ty(this->getVMContext())); + + // XXX: Should this be i64 instead, and should the limit increase? + llvm::Type *I32Ty = llvm::Type::getInt32Ty(this->getVMContext()); + return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); + } + + if (NumRegsLeft > 0) { + uint64_t NumRegs = numRegsForType(Ty); + if (NumRegsLeft >= NumRegs) { + NumRegsLeft -= NumRegs; + return ABIArgInfo::getDirect(); + } + } + + // Use pass-by-reference instead of pass-by-value for struct arguments in + // function ABI. + return ABIArgInfo::getIndirectAliased( + this->getContext().getTypeAlignInChars(Ty), + this->getContext().getTargetAddressSpace(LangAS::opencl_private)); + } + + llvm::FixedVectorType * + getOptimalVectorMemoryType(llvm::FixedVectorType *Ty, + const LangOptions &LangOpt) const override { + // We have legal instructions for 96-bit so 3x32 can be supported. + // FIXME: This check should be a subtarget feature as technically SI + // doesn't support it. + if (Ty->getNumElements() == 3 && + this->getDataLayout().getTypeSizeInBits(Ty) == 96) + return Ty; + return Base::getOptimalVectorMemoryType(Ty, LangOpt); + } +}; + Address EmitVAArgInstr(CodeGenFunction &CGF, Address VAListAddr, QualType Ty, const ABIArgInfo &AI); diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 07e2eac39305d..230742255073e 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -22,95 +22,18 @@ using namespace clang::CodeGen; namespace { -class AMDGPUABIInfo final : public DefaultABIInfo { -private: - static const unsigned MaxNumRegsForArgsRet = 16; - - uint64_t numRegsForType(QualType Ty) const; - - bool isHomogeneousAggregateBaseType(QualType Ty) const override; - bool isHomogeneousAggregateSmallEnough(const Type *Base, - uint64_t Members) const override; - - // Coerce HIP scalar pointer arguments from generic pointers to global ones. - llvm::Type *coerceKernelArgumentType(llvm::Type *Ty, unsigned FromAS, - unsigned ToAS) const { - // Single value types. - auto *PtrTy = llvm::dyn_cast<llvm::PointerType>(Ty); - if (PtrTy && PtrTy->getAddressSpace() == FromAS) - return llvm::PointerType::get(Ty->getContext(), ToAS); - return Ty; - } - +class AMDGPUABIInfo final : public AMDGPUABIInfoCommon<DefaultABIInfo> { public: - explicit AMDGPUABIInfo(CodeGen::CodeGenTypes &CGT) : - DefaultABIInfo(CGT) {} + explicit AMDGPUABIInfo(CodeGen::CodeGenTypes &CGT) + : AMDGPUABIInfoCommon(CGT) {} - ABIArgInfo classifyReturnType(QualType RetTy) const; ABIArgInfo classifyKernelArgumentType(QualType Ty) const; - ABIArgInfo classifyArgumentType(QualType Ty, bool Variadic, - unsigned &NumRegsLeft) const; void computeInfo(CGFunctionInfo &FI) const override; RValue EmitVAArg(CodeGenFunction &CGF, Address VAListAddr, QualType Ty, AggValueSlot Slot) const override; - - llvm::FixedVectorType * - getOptimalVectorMemoryType(llvm::FixedVectorType *T, - const LangOptions &Opt) const override { - // We have legal instructions for 96-bit so 3x32 can be supported. - // FIXME: This check should be a subtarget feature as technically SI doesn't - // support it. - if (T->getNumElements() == 3 && getDataLayout().getTypeSizeInBits(T) == 96) - return T; - return DefaultABIInfo::getOptimalVectorMemoryType(T, Opt); - } }; -bool AMDGPUABIInfo::isHomogeneousAggregateBaseType(QualType Ty) const { - return true; -} - -bool AMDGPUABIInfo::isHomogeneousAggregateSmallEnough( - const Type *Base, uint64_t Members) const { - uint32_t NumRegs = (getContext().getTypeSize(Base) + 31) / 32; - - // Homogeneous Aggregates may occupy at most 16 registers. - return Members * NumRegs <= MaxNumRegsForArgsRet; -} - -/// Estimate number of registers the type will use when passed in registers. -uint64_t AMDGPUABIInfo::numRegsForType(QualType Ty) const { - uint64_t NumRegs = 0; - - if (const VectorType *VT = Ty->getAs<VectorType>()) { - // Compute from the number of elements. The reported size is based on the - // in-memory size, which includes the padding 4th element for 3-vectors. - QualType EltTy = VT->getElementType(); - uint64_t EltSize = getContext().getTypeSize(EltTy); - - // 16-bit element vectors should be passed as packed. - if (EltSize == 16) - return (VT->getNumElements() + 1) / 2; - - uint64_t EltNumRegs = (EltSize + 31) / 32; - return EltNumRegs * VT->getNumElements(); - } - - if (const auto *RD = Ty->getAsRecordDecl()) { - assert(!RD->hasFlexibleArrayMember()); - - for (const FieldDecl *Field : RD->fields()) { - QualType FieldTy = Field->getType(); - NumRegs += numRegsForType(FieldTy); - } - - return NumRegs; - } - - return (getContext().getTypeSize(Ty) + 31) / 32; -} - void AMDGPUABIInfo::computeInfo(CGFunctionInfo &FI) const { llvm::CallingConv::ID CC = FI.getCallingConvention(); @@ -120,13 +43,13 @@ void AMDGPUABIInfo::computeInfo(CGFunctionInfo &FI) const { unsigned ArgumentIndex = 0; const unsigned numFixedArguments = FI.getNumRequiredArgs(); - unsigned NumRegsLeft = MaxNumRegsForArgsRet; + NumRegsLeft = MaxNumRegsForArgsRet; for (auto &Arg : FI.arguments()) { if (CC == llvm::CallingConv::AMDGPU_KERNEL) { Arg.info = classifyKernelArgumentType(Arg.type); } else { bool FixedArgument = ArgumentIndex++ < numFixedArguments; - Arg.info = classifyArgumentType(Arg.type, !FixedArgument, NumRegsLeft); + Arg.info = classifyArgumentType(Arg.type, !FixedArgument); } } } @@ -140,45 +63,6 @@ RValue AMDGPUABIInfo::EmitVAArg(CodeGenFunction &CGF, Address VAListAddr, CharUnits::fromQuantity(4), AllowHigherAlign, Slot); } -ABIArgInfo AMDGPUABIInfo::classifyReturnType(QualType RetTy) const { - if (isAggregateTypeForABI(RetTy)) { - // Records with non-trivial destructors/copy-constructors should not be - // returned by value. - if (!getRecordArgABI(RetTy, getCXXABI())) { - // Ignore empty structs/unions. - if (isEmptyRecord(getContext(), RetTy, true)) - return ABIArgInfo::getIgnore(); - - // Lower single-element structs to just return a regular value. - if (const Type *SeltTy = isSingleElementStruct(RetTy, getContext())) - return ABIArgInfo::getDirect(CGT.ConvertType(QualType(SeltTy, 0))); - - if (const auto *RD = RetTy->getAsRecordDecl(); - RD && RD->hasFlexibleArrayMember()) - return DefaultABIInfo::classifyReturnType(RetTy); - - // Pack aggregates <= 4 bytes into single VGPR or pair. - uint64_t Size = getContext().getTypeSize(RetTy); - if (Size <= 16) - return ABIArgInfo::getDirect(llvm::Type::getInt16Ty(getVMContext())); - - if (Size <= 32) - return ABIArgInfo::getDirect(llvm::Type::getInt32Ty(getVMContext())); - - if (Size <= 64) { - llvm::Type *I32Ty = llvm::Type::getInt32Ty(getVMContext()); - return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); - } - - if (numRegsForType(RetTy) <= MaxNumRegsForArgsRet) - return ABIArgInfo::getDirect(); - } - } - - // Otherwise just do the default thing. - return DefaultABIInfo::classifyReturnType(RetTy); -} - /// For kernels all parameters are really passed in a special buffer. It doesn't /// make sense to pass anything byval, so everything must be direct. ABIArgInfo AMDGPUABIInfo::classifyKernelArgumentType(QualType Ty) const { @@ -213,83 +97,6 @@ ABIArgInfo AMDGPUABIInfo::classifyKernelArgumentType(QualType Ty) const { return ABIArgInfo::getDirect(LTy, 0, nullptr, false); } -ABIArgInfo AMDGPUABIInfo::classifyArgumentType(QualType Ty, bool Variadic, - unsigned &NumRegsLeft) const { - assert(NumRegsLeft <= MaxNumRegsForArgsRet && "register estimate underflow"); - - Ty = useFirstFieldIfTransparentUnion(Ty); - - if (Variadic) { - return ABIArgInfo::getDirect(/*T=*/nullptr, - /*Offset=*/0, - /*Padding=*/nullptr, - /*CanBeFlattened=*/false, - /*Align=*/0); - } - - if (isAggregateTypeForABI(Ty)) { - // Records with non-trivial destructors/copy-constructors should not be - // passed by value. - if (auto RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), - RAA == CGCXXABI::RAA_DirectInMemory); - - // Ignore empty structs/unions. - if (isEmptyRecord(getContext(), Ty, true)) - return ABIArgInfo::getIgnore(); - - // Lower single-element structs to just pass a regular value. TODO: We - // could do reasonable-size multiple-element structs too, using getExpand(), - // though watch out for things like bitfields. - if (const Type *SeltTy = isSingleElementStruct(Ty, getContext())) - return ABIArgInfo::getDirect(CGT.ConvertType(QualType(SeltTy, 0))); - - if (const auto *RD = Ty->getAsRecordDecl(); - RD && RD->hasFlexibleArrayMember()) - return DefaultABIInfo::classifyArgumentType(Ty); - - // Pack aggregates <= 8 bytes into single VGPR or pair. - uint64_t Size = getContext().getTypeSize(Ty); - if (Size <= 64) { - unsigned NumRegs = (Size + 31) / 32; - NumRegsLeft -= std::min(NumRegsLeft, NumRegs); - - if (Size <= 16) - return ABIArgInfo::getDirect(llvm::Type::getInt16Ty(getVMContext())); - - if (Size <= 32) - return ABIArgInfo::getDirect(llvm::Type::getInt32Ty(getVMContext())); - - // XXX: Should this be i64 instead, and should the limit increase? - llvm::Type *I32Ty = llvm::Type::getInt32Ty(getVMContext()); - return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); - } - - if (NumRegsLeft > 0) { - uint64_t NumRegs = numRegsForType(Ty); - if (NumRegsLeft >= NumRegs) { - NumRegsLeft -= NumRegs; - return ABIArgInfo::getDirect(); - } - } - - // Use pass-by-reference in stead of pass-by-value for struct arguments in - // function ABI. - return ABIArgInfo::getIndirectAliased( - getContext().getTypeAlignInChars(Ty), - getContext().getTargetAddressSpace(LangAS::opencl_private)); - } - - // Otherwise just do the default thing. - ABIArgInfo ArgInfo = DefaultABIInfo::classifyArgumentType(Ty); - if (!ArgInfo.isIndirect()) { - uint64_t NumRegs = numRegsForType(Ty); - NumRegsLeft -= std::min(NumRegs, uint64_t{NumRegsLeft}); - } - - return ArgInfo; -} - class AMDGPUTargetCodeGenInfo : public TargetCodeGenInfo { public: AMDGPUTargetCodeGenInfo(CodeGenTypes &CGT) diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp index e8148d1566f85..7dbcf5e439095 100644 --- a/clang/lib/CodeGen/Targets/SPIR.cpp +++ b/clang/lib/CodeGen/Targets/SPIR.cpp @@ -48,40 +48,12 @@ class SPIRVABIInfo : public CommonSPIRABIInfo { ABIArgInfo classifyKernelArgumentType(QualType Ty) const; }; -class AMDGCNSPIRVABIInfo : public SPIRVABIInfo { - // TODO: this should be unified / shared with AMDGPU, ideally we'd like to - // re-use AMDGPUABIInfo eventually, rather than duplicate. - static constexpr unsigned MaxNumRegsForArgsRet = 16; // 16 32-bit registers - mutable unsigned NumRegsLeft = 0; - - uint64_t numRegsForType(QualType Ty) const; - - bool isHomogeneousAggregateBaseType(QualType Ty) const override { - return true; - } - bool isHomogeneousAggregateSmallEnough(const Type *Base, - uint64_t Members) const override { - uint32_t NumRegs = (getContext().getTypeSize(Base) + 31) / 32; - - // Homogeneous Aggregates may occupy at most 16 registers. - return Members * NumRegs <= MaxNumRegsForArgsRet; - } - - // Coerce HIP scalar pointer arguments from generic pointers to global ones. - llvm::Type *coerceKernelArgumentType(llvm::Type *Ty, unsigned FromAS, - unsigned ToAS) const; - - ABIArgInfo classifyReturnType(QualType RetTy) const; +class AMDGCNSPIRVABIInfo : public AMDGPUABIInfoCommon<SPIRVABIInfo> { ABIArgInfo classifyKernelArgumentType(QualType Ty) const; - ABIArgInfo classifyArgumentType(QualType Ty, bool Variadic) const; public: - AMDGCNSPIRVABIInfo(CodeGenTypes &CGT) : SPIRVABIInfo(CGT) {} + AMDGCNSPIRVABIInfo(CodeGenTypes &CGT) : AMDGPUABIInfoCommon(CGT) {} void computeInfo(CGFunctionInfo &FI) const override; - - llvm::FixedVectorType * - getOptimalVectorMemoryType(llvm::FixedVectorType *Ty, - const LangOptions &LangOpt) const override; }; } // end anonymous namespace namespace { @@ -209,84 +181,6 @@ RValue SPIRVABIInfo::EmitVAArg(CodeGenFunction &CGF, Address VAListAddr, /*AllowHigherAlign=*/true, Slot); } -uint64_t AMDGCNSPIRVABIInfo::numRegsForType(QualType Ty) const { - // This duplicates the AMDGPUABI computation. - uint64_t NumRegs = 0; - - if (const VectorType *VT = Ty->getAs<VectorType>()) { - // Compute from the number of elements. The reported size is based on the - // in-memory size, which includes the padding 4th element for 3-vectors. - QualType EltTy = VT->getElementType(); - uint64_t EltSize = getContext().getTypeSize(EltTy); - - // 16-bit element vectors should be passed as packed. - if (EltSize == 16) - return (VT->getNumElements() + 1) / 2; - - uint64_t EltNumRegs = (EltSize + 31) / 32; - return EltNumRegs * VT->getNumElements(); - } - - if (const auto *RD = Ty->getAsRecordDecl()) { - assert(!RD->hasFlexibleArrayMember()); - - for (const FieldDecl *Field : RD->fields()) { - QualType FieldTy = Field->getType(); - NumRegs += numRegsForType(FieldTy); - } - - return NumRegs; - } - - return (getContext().getTypeSize(Ty) + 31) / 32; -} - -llvm::Type *AMDGCNSPIRVABIInfo::coerceKernelArgumentType(llvm::Type *Ty, - unsigned FromAS, - unsigned ToAS) const { - // Single value types. - auto *PtrTy = llvm::dyn_cast<llvm::PointerType>(Ty); - if (PtrTy && PtrTy->getAddressSpace() == FromAS) - return llvm::PointerType::get(Ty->getContext(), ToAS); - return Ty; -} - -ABIArgInfo AMDGCNSPIRVABIInfo::classifyReturnType(QualType RetTy) const { - if (!isAggregateTypeForABI(RetTy) || getRecordArgABI(RetTy, getCXXABI())) - return DefaultABIInfo::classifyReturnType(RetTy); - - // Ignore empty structs/unions. - if (isEmptyRecord(getContext(), RetTy, true)) - return ABIArgInfo::getIgnore(); - - // Lower single-element structs to just return a regular value. - if (const Type *SeltTy = isSingleElementStruct(RetTy, getContext())) - return ABIArgInfo::getDirect(CGT.ConvertType(QualType(SeltTy, 0))); - - if (const auto *RD = RetTy->getAsRecordDecl(); - RD && RD->hasFlexibleArrayMember()) - return DefaultABIInfo::classifyReturnType(RetTy); - - // Pack aggregates <= 4 bytes into single VGPR or pair. - uint64_t Size = getContext().getTypeSize(RetTy); - if (Size <= 16) - return ABIArgInfo::getDirect(llvm::Type::getInt16Ty(getVMContext())); - - if (Size <= 32) - return ABIArgInfo::getDirect(llvm::Type::getInt32Ty(getVMContext())); - - // TODO: This carried over from AMDGPU oddity, we retain it to - // ensure consistency, but it might be reasonable to return Int64. - if (Size <= 64) { - llvm::Type *I32Ty = llvm::Type::getInt32Ty(getVMContext()); - return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); - } - - if (numRegsForType(RetTy) <= MaxNumRegsForArgsRet) - return ABIArgInfo::getDirect(); - return DefaultABIInfo::classifyReturnType(RetTy); -} - /// For kernels all parameters are really passed in a special buffer. It doesn't /// make sense to pass anything byval, so everything must be direct. ABIArgInfo AMDGCNSPIRVABIInfo::classifyKernelArgumentType(QualType Ty) const { @@ -320,83 +214,6 @@ ABIArgInfo AMDGCNSPIRVABIInfo::classifyKernelArgumentType(QualType Ty) const { return ABIArgInfo::getDirect(LTy, 0, nullptr, false); } -ABIArgInfo AMDGCNSPIRVABIInfo::classifyArgumentType(QualType Ty, - bool Variadic) const { - assert(NumRegsLeft <= MaxNumRegsForArgsRet && "register estimate underflow"); - - Ty = useFirstFieldIfTransparentUnion(Ty); - - if (Variadic) { - return ABIArgInfo::getDirect(/*T=*/nullptr, - /*Offset=*/0, - /*Padding=*/nullptr, - /*CanBeFlattened=*/false, - /*Align=*/0); - } - - if (!isAggregateTypeForABI(Ty)) { - ABIArgInfo ArgInfo = DefaultABIInfo::classifyArgumentType(Ty); - if (!ArgInfo.isIndirect()) { - uint64_t NumRegs = numRegsForType(Ty); - NumRegsLeft -= std::min(NumRegs, uint64_t{NumRegsLeft}); - } - - return ArgInfo; - } - - // Records with non-trivial destructors/copy-constructors should not be - // passed by value. - if (auto RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), - RAA == CGCXXABI::RAA_DirectInMemory); - - // Ignore empty structs/unions. - if (isEmptyRecord(getContext(), Ty, true)) - return ABIArgInfo::getIgnore(); - - // Lower single-element structs to just pass a regular value. TODO: We - // could do reasonable-size multiple-element structs too, using getExpand(), - // though watch out for things like bitfields. - if (const Type *SeltTy = isSingleElementStruct(Ty, getContext())) - return ABIArgInfo::getDirect(CGT.ConvertType(QualType(SeltTy, 0))); - - if (const auto *RD = Ty->getAsRecordDecl(); - RD && RD->hasFlexibleArrayMember()) - return DefaultABIInfo::classifyArgumentType(Ty); - - uint64_t Size = getContext().getTypeSize(Ty); - if (Size <= 64) { - // Pack aggregates <= 8 bytes into single VGPR or pair. - unsigned NumRegs = (Size + 31) / 32; - NumRegsLeft -= std::min(NumRegsLeft, NumRegs); - - if (Size <= 16) - return ABIArgInfo::getDirect(llvm::Type::getInt16Ty(getVMContext())); - - if (Size <= 32) - return ABIArgInfo::getDirect(llvm::Type::getInt32Ty(getVMContext())); - - // TODO: This is an AMDGPU oddity, and might be vestigial, we retain it to - // ensure consistency, but it should be revisited. - llvm::Type *I32Ty = llvm::Type::getInt32Ty(getVMContext()); - return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); - } - - if (NumRegsLeft > 0) { - uint64_t NumRegs = numRegsForType(Ty); - if (NumRegsLeft >= NumRegs) { - NumRegsLeft -= NumRegs; - return ABIArgInfo::getDirect(); - } - } - - // Use pass-by-reference in stead of pass-by-value for struct arguments in - // function ABI. - return ABIArgInfo::getIndirectAliased( - getContext().getTypeAlignInChars(Ty), - getContext().getTargetAddressSpace(LangAS::opencl_private)); -} - void AMDGCNSPIRVABIInfo::computeInfo(CGFunctionInfo &FI) const { llvm::CallingConv::ID CC = FI.getCallingConvention(); @@ -428,14 +245,6 @@ SPIRVABIInfo::getOptimalVectorMemoryType(llvm::FixedVectorType *Ty, return DefaultABIInfo::getOptimalVectorMemoryType(Ty, LangOpt); } -llvm::FixedVectorType *AMDGCNSPIRVABIInfo::getOptimalVectorMemoryType( - llvm::FixedVectorType *Ty, const LangOptions &LangOpt) const { - // AMDGPU has legal instructions for 96-bit so 3x32 can be supported. - if (Ty->getNumElements() == 3 && getDataLayout().getTypeSizeInBits(Ty) == 96) - return Ty; - return DefaultABIInfo::getOptimalVectorMemoryType(Ty, LangOpt); -} - namespace clang { namespace CodeGen { void computeSPIRKernelABIInfo(CodeGenModule &CGM, CGFunctionInfo &FI) { >From 72634dbd5bc9e9ed60dc3dda32ae39e6b5c18855 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Mon, 14 Sep 2026 09:38:10 +0200 Subject: [PATCH 2/2] move to llvm --- clang/lib/CodeGen/ABIInfoImpl.h | 193 ------------ clang/lib/CodeGen/CGCall.cpp | 14 +- clang/lib/CodeGen/CodeGenModule.cpp | 29 ++ clang/lib/CodeGen/QualTypeMapper.cpp | 30 +- clang/lib/CodeGen/QualTypeMapper.h | 4 + clang/lib/CodeGen/Targets/AMDGPU.cpp | 203 ++++++++++++- clang/lib/CodeGen/Targets/SPIR.cpp | 195 +++++++++++- .../amdgcnspirv-uses-amdgpu-abi.cpp | 4 + .../CodeGenOpenCL/amdgpu-abi-struct-coerce.cl | 1 + clang/test/CodeGenOpenCL/opencl_types.cl | 1 + llvm/include/llvm/ABI/FunctionInfo.h | 36 ++- llvm/include/llvm/ABI/TargetInfo.h | 36 ++- llvm/include/llvm/ABI/Types.h | 4 + llvm/lib/ABI/CMakeLists.txt | 1 + llvm/lib/ABI/TargetInfo.cpp | 71 ++++- llvm/lib/ABI/Targets/AMDGPU.cpp | 279 ++++++++++++++++++ llvm/lib/ABI/Targets/X86.cpp | 55 ---- llvm/utils/gn/secondary/llvm/lib/ABI/BUILD.gn | 1 + 18 files changed, 881 insertions(+), 276 deletions(-) create mode 100644 llvm/lib/ABI/Targets/AMDGPU.cpp diff --git a/clang/lib/CodeGen/ABIInfoImpl.h b/clang/lib/CodeGen/ABIInfoImpl.h index 4649060c8e712..d9d79c6a55ddb 100644 --- a/clang/lib/CodeGen/ABIInfoImpl.h +++ b/clang/lib/CodeGen/ABIInfoImpl.h @@ -11,7 +11,6 @@ #include "ABIInfo.h" #include "CGCXXABI.h" -#include "llvm/IR/DerivedTypes.h" namespace clang::CodeGen { @@ -141,198 +140,6 @@ bool isEmptyRecordForLayout(const ASTContext &Context, QualType T); /// it exists. const Type *isSingleElementStruct(QualType T, ASTContext &Context); -/// Shared classification rules for AMDGPU and AMDGCN-SPIR-V, with \p Base as -/// the fallback ABIInfo for non-register-packed cases. -template <typename Base> class AMDGPUABIInfoCommon : public Base { -protected: - static constexpr unsigned MaxNumRegsForArgsRet = 16; // 16 32-bit registers - mutable unsigned NumRegsLeft = 0; - - using Base::Base; - - /// Estimate number of registers the type will use when passed in registers. - uint64_t numRegsForType(QualType Ty) const { - uint64_t NumRegs = 0; - - if (const VectorType *VT = Ty->template getAs<VectorType>()) { - // Compute from the number of elements. The reported size is based on - // the in-memory size, which includes the padding 4th element for - // 3-vectors. - QualType EltTy = VT->getElementType(); - uint64_t EltSize = this->getContext().getTypeSize(EltTy); - - // 16-bit element vectors should be passed as packed. - if (EltSize == 16) - return (VT->getNumElements() + 1) / 2; - - uint64_t EltNumRegs = (EltSize + 31) / 32; - return EltNumRegs * VT->getNumElements(); - } - - if (const auto *RD = Ty->getAsRecordDecl()) { - assert(!RD->hasFlexibleArrayMember()); - - for (const FieldDecl *Field : RD->fields()) - NumRegs += numRegsForType(Field->getType()); - - return NumRegs; - } - - return (this->getContext().getTypeSize(Ty) + 31) / 32; - } - - bool isHomogeneousAggregateBaseType(QualType Ty) const override { - return true; - } - - bool isHomogeneousAggregateSmallEnough(const Type *T, - uint64_t Members) const override { - uint32_t NumRegs = (this->getContext().getTypeSize(T) + 31) / 32; - - // Homogeneous Aggregates may occupy at most 16 registers. - return Members * NumRegs <= MaxNumRegsForArgsRet; - } - - // Coerce scalar pointer arguments from generic pointers to a fixed AS. - llvm::Type *coerceKernelArgumentType(llvm::Type *Ty, unsigned FromAS, - unsigned ToAS) const { - // Single value types. - auto *PtrTy = llvm::dyn_cast<llvm::PointerType>(Ty); - if (PtrTy && PtrTy->getAddressSpace() == FromAS) - return llvm::PointerType::get(Ty->getContext(), ToAS); - return Ty; - } - - ABIArgInfo classifyReturnType(QualType RetTy) const { - if (!isAggregateTypeForABI(RetTy) || - getRecordArgABI(RetTy, this->getCXXABI())) - return Base::classifyReturnType(RetTy); - - // Ignore empty structs/unions. - if (isEmptyRecord(this->getContext(), RetTy, true)) - return ABIArgInfo::getIgnore(); - - // Lower single-element structs to just return a regular value. - if (const Type *SeltTy = isSingleElementStruct(RetTy, this->getContext())) - return ABIArgInfo::getDirect(this->CGT.ConvertType(QualType(SeltTy, 0))); - - if (const auto *RD = RetTy->getAsRecordDecl(); - RD && RD->hasFlexibleArrayMember()) - return Base::classifyReturnType(RetTy); - - // Pack aggregates <= 4 bytes into single VGPR or pair. - uint64_t Size = this->getContext().getTypeSize(RetTy); - if (Size <= 16) - return ABIArgInfo::getDirect( - llvm::Type::getInt16Ty(this->getVMContext())); - - if (Size <= 32) - return ABIArgInfo::getDirect( - llvm::Type::getInt32Ty(this->getVMContext())); - - if (Size <= 64) { - llvm::Type *I32Ty = llvm::Type::getInt32Ty(this->getVMContext()); - return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); - } - - if (numRegsForType(RetTy) <= MaxNumRegsForArgsRet) - return ABIArgInfo::getDirect(); - - return Base::classifyReturnType(RetTy); - } - - ABIArgInfo classifyArgumentType(QualType Ty, bool Variadic) const { - assert(NumRegsLeft <= MaxNumRegsForArgsRet && - "register estimate underflow"); - - Ty = useFirstFieldIfTransparentUnion(Ty); - - if (Variadic) { - return ABIArgInfo::getDirect(/*T=*/nullptr, - /*Offset=*/0, - /*Padding=*/nullptr, - /*CanBeFlattened=*/false, - /*Align=*/0); - } - - if (!isAggregateTypeForABI(Ty)) { - ABIArgInfo ArgInfo = Base::classifyArgumentType(Ty); - if (!ArgInfo.isIndirect()) { - uint64_t NumRegs = numRegsForType(Ty); - NumRegsLeft -= std::min(NumRegs, uint64_t{NumRegsLeft}); - } - - return ArgInfo; - } - - // Records with non-trivial destructors/copy-constructors should not be - // passed by value. - if (auto RAA = getRecordArgABI(Ty, this->getCXXABI())) - return this->getNaturalAlignIndirect( - Ty, this->getDataLayout().getAllocaAddrSpace(), - RAA == CGCXXABI::RAA_DirectInMemory); - - // Ignore empty structs/unions. - if (isEmptyRecord(this->getContext(), Ty, true)) - return ABIArgInfo::getIgnore(); - - // Lower single-element structs to just pass a regular value. TODO: We - // could do reasonable-size multiple-element structs too, using - // getExpand(), though watch out for things like bitfields. - if (const Type *SeltTy = isSingleElementStruct(Ty, this->getContext())) - return ABIArgInfo::getDirect(this->CGT.ConvertType(QualType(SeltTy, 0))); - - if (const auto *RD = Ty->getAsRecordDecl(); - RD && RD->hasFlexibleArrayMember()) - return Base::classifyArgumentType(Ty); - - // Pack aggregates <= 8 bytes into single VGPR or pair. - uint64_t Size = this->getContext().getTypeSize(Ty); - if (Size <= 64) { - unsigned NumRegs = (Size + 31) / 32; - NumRegsLeft -= std::min(NumRegsLeft, NumRegs); - - if (Size <= 16) - return ABIArgInfo::getDirect( - llvm::Type::getInt16Ty(this->getVMContext())); - - if (Size <= 32) - return ABIArgInfo::getDirect( - llvm::Type::getInt32Ty(this->getVMContext())); - - // XXX: Should this be i64 instead, and should the limit increase? - llvm::Type *I32Ty = llvm::Type::getInt32Ty(this->getVMContext()); - return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); - } - - if (NumRegsLeft > 0) { - uint64_t NumRegs = numRegsForType(Ty); - if (NumRegsLeft >= NumRegs) { - NumRegsLeft -= NumRegs; - return ABIArgInfo::getDirect(); - } - } - - // Use pass-by-reference instead of pass-by-value for struct arguments in - // function ABI. - return ABIArgInfo::getIndirectAliased( - this->getContext().getTypeAlignInChars(Ty), - this->getContext().getTargetAddressSpace(LangAS::opencl_private)); - } - - llvm::FixedVectorType * - getOptimalVectorMemoryType(llvm::FixedVectorType *Ty, - const LangOptions &LangOpt) const override { - // We have legal instructions for 96-bit so 3x32 can be supported. - // FIXME: This check should be a subtarget feature as technically SI - // doesn't support it. - if (Ty->getNumElements() == 3 && - this->getDataLayout().getTypeSizeInBits(Ty) == 96) - return Ty; - return Base::getOptimalVectorMemoryType(Ty, LangOpt); - } -}; - Address EmitVAArgInstr(CodeGenFunction &CGF, Address VAListAddr, QualType Ty, const ABIArgInfo &AI); diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 1221829871b9f..9b63fcc303ac6 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -976,6 +976,9 @@ void CodeGenModule::computeABIInfoUsingLib(CGFunctionInfo &FI) { CheckSimple(Target.getDirectAlign(), Res.getDirectAlign(), "DirectAlign"); CheckSimple(Target.getDirectOffset(), Res.getDirectOffset(), "DirectOffset"); + if (Res.isDirect()) + CheckSimple(Target.getCanBeFlattened(), Res.getCanBeFlattened(), + "CanBeFlattened"); break; case ABIArgInfo::Indirect: CheckSimple(Target.getIndirectByVal(), Res.getIndirectByVal(), @@ -1023,7 +1026,9 @@ ABIArgInfo CodeGenModule::convertABIArgInfo(const llvm::abi::ArgInfo &AbiInfo, CoercedType = AbiReverseMapper->convertType(AbiInfo.getCoerceToType()); if (!CoercedType) CoercedType = getTypes().ConvertType(Type); - return ABIArgInfo::getDirect(CoercedType, AbiInfo.getDirectOffset()); + return ABIArgInfo::getDirect(CoercedType, AbiInfo.getDirectOffset(), + /*Padding=*/nullptr, + AbiInfo.getCanBeFlattened()); } case llvm::abi::ArgInfo::Extend: { llvm::Type *CoercedType = nullptr; @@ -1049,6 +1054,13 @@ ABIArgInfo CodeGenModule::convertABIArgInfo(const llvm::abi::ArgInfo &AbiInfo, AbiInfo.getIndirectByVal(), AbiInfo.getIndirectRealign()); } + case llvm::abi::ArgInfo::IndirectAliased: { + CharUnits Alignment = + CharUnits::fromQuantity(AbiInfo.getIndirectAlign().value()); + return ABIArgInfo::getIndirectAliased(Alignment, + AbiInfo.getIndirectAddrSpace(), + AbiInfo.getIndirectRealign()); + } case llvm::abi::ArgInfo::Ignore: return ABIArgInfo::getIgnore(); } diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 40feb4b30f98c..43b69416b7f72 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -346,6 +346,12 @@ const TargetCodeGenInfo &CodeGenModule::getTargetCodeGenInfo() { return *TheTargetCodeGenInfo; } +/// AMDGCN, and the AMDGCN-flavoured SPIR-V that lowers to it, share one +/// classifier. +static bool usesAMDGPUABI(const llvm::Triple &T) { + return T.isAMDGCN() || (T.isSPIRV() && T.getVendor() == llvm::Triple::AMD); +} + bool CodeGenModule::shouldUseLLVMABILowering(unsigned CallingConv) const { if (!CodeGenOpts.ExperimentalABILowering) return false; @@ -359,6 +365,9 @@ bool CodeGenModule::shouldUseLLVMABILowering(unsigned CallingConv) const { T.getArch() == llvm::Triple::aarch64_be) return true; + if (usesAMDGPUABI(T)) + return true; + if (T.getArch() == llvm::Triple::x86_64 && !T.isOSWindows() && !T.isUEFI() && !T.isOSDarwin() && !T.isOSCygMing()) { switch (CallingConv) { @@ -390,6 +399,26 @@ CodeGenModule::getLLVMABITargetInfo(llvm::abi::TypeBuilder &TB) { const llvm::Triple &T = getTriple(); + if (usesAMDGPUABI(T)) { + ASTContext &Ctx = getContext(); + const bool IsSPIRV = T.isSPIRV(); + llvm::abi::AMDGPUABIOptions Opts; + Opts.KernelCC = IsSPIRV ? llvm::CallingConv::SPIR_KERNEL + : llvm::CallingConv::AMDGPU_KERNEL; + Opts.AllocaAddrSpace = getDataLayout().getAllocaAddrSpace(); + Opts.PrivateAddrSpace = Ctx.getTargetAddressSpace(LangAS::opencl_private); + Opts.ConstantAddrSpace = Ctx.getTargetAddressSpace(LangAS::opencl_constant); + Opts.GenericAddrSpace = Ctx.getTargetAddressSpace(LangAS::Default); + // Pre-existing divergences between the two classifiers, kept deliberately. + Opts.KernelArgAddrSpace = Ctx.getTargetAddressSpace( + IsSPIRV ? LangAS::opencl_global : LangAS::cuda_device); + Opts.CoerceKernelPointerArgs = + IsSPIRV ? getLangOpts().isTargetDevice() : getLangOpts().HIP; + Opts.HasInt128 = Ctx.getTargetInfo().hasInt128Type(); + TheLLVMABITargetInfo = llvm::abi::createAMDGPUTargetInfo(TB, Opts); + return *TheLLVMABITargetInfo; + } + switch (T.getArch()) { default: llvm_unreachable("LLVMABI lowering requested for an unsupported target"); diff --git a/clang/lib/CodeGen/QualTypeMapper.cpp b/clang/lib/CodeGen/QualTypeMapper.cpp index a6e96c63446f6..850ae406ab3a1 100644 --- a/clang/lib/CodeGen/QualTypeMapper.cpp +++ b/clang/lib/CodeGen/QualTypeMapper.cpp @@ -128,8 +128,9 @@ const llvm::abi::Type *QualTypeMapper::convertTypeImpl(QualType QT) { ASTCtx.getTypeSize(QT), getTypeAlign(QT)); } case Type::BlockPointer: - case Type::Pipe: return createPointerTypeForPointee(ASTCtx.VoidPtrTy); + case Type::Pipe: + return createOpenCLOpaqueType(QT.getTypePtr()); case Type::ConstantMatrix: { const auto *MT = cast<ConstantMatrixType>(QT); return Builder.getArrayType(convertType(MT->getElementType()), @@ -262,7 +263,7 @@ QualTypeMapper::convertBuiltinType(const BuiltinType *BT) { case BuiltinType::OCLClkEvent: case BuiltinType::OCLQueue: case BuiltinType::OCLReserveID: - return createPointerTypeForPointee(QT); + return createOpenCLOpaqueType(BT); // Objective-C builtin types are represented as opaque pointers. case BuiltinType::ObjCId: @@ -604,12 +605,26 @@ llvm::Align QualTypeMapper::getTypeAlign(QualType QT) const { return llvm::Align(ASTCtx.getTypeAlignInChars(QT).getQuantity()); } +const llvm::abi::Type * +QualTypeMapper::createPointerType(LangAS AddrSpace, + std::optional<unsigned> TargetAddrSpace) { + const clang::TargetInfo &TI = ASTCtx.getTargetInfo(); + return Builder.getPointerType( + TI.getPointerWidth(AddrSpace), + llvm::Align(TI.getPointerAlign(AddrSpace) / 8), + TargetAddrSpace.value_or(TI.getTargetAddressSpace(AddrSpace))); +} + +const llvm::abi::Type * +QualTypeMapper::createOpenCLOpaqueType(const clang::Type *T) { + // Mirrors CGOpenCLRuntime::getPointerType: the address space comes from the + // target hook, not from a qualifier on the type. + return createPointerType(ASTCtx.getOpenCLTypeAddrSpace(T)); +} + const llvm::abi::Type * QualTypeMapper::createPointerTypeForPointee(QualType PointeeType) { - auto AddrSpace = PointeeType.getAddressSpace(); - auto PointerSize = ASTCtx.getTargetInfo().getPointerWidth(AddrSpace); - llvm::Align Alignment = - llvm::Align(ASTCtx.getTargetInfo().getPointerAlign(AddrSpace)); + LangAS AddrSpace = PointeeType.getAddressSpace(); // Function types without an explicit address space qualifier use the program // address space, which may differ from the default data address space on // targets like AMDGPU. @@ -617,8 +632,7 @@ QualTypeMapper::createPointerTypeForPointee(QualType PointeeType) { PointeeType->isFunctionType() && !PointeeType.hasAddressSpace() ? DL.getProgramAddressSpace() : ASTCtx.getTargetInfo().getTargetAddressSpace(AddrSpace); - return Builder.getPointerType(PointerSize, llvm::Align(Alignment.value() / 8), - TargetAddrSpace); + return createPointerType(AddrSpace, TargetAddrSpace); } /// Processes the fields of a record (struct/class/union) and populates diff --git a/clang/lib/CodeGen/QualTypeMapper.h b/clang/lib/CodeGen/QualTypeMapper.h index f6bb43482c705..08c7214d85459 100644 --- a/clang/lib/CodeGen/QualTypeMapper.h +++ b/clang/lib/CodeGen/QualTypeMapper.h @@ -51,7 +51,11 @@ class QualTypeMapper { const llvm::abi::RecordType *convertStructType(const clang::RecordDecl *RD); const llvm::abi::RecordType *convertUnionType(const clang::RecordDecl *RD); + const llvm::abi::Type * + createPointerType(LangAS AddrSpace, + std::optional<unsigned> TargetAddrSpace = std::nullopt); const llvm::abi::Type *createPointerTypeForPointee(QualType PointeeType); + const llvm::abi::Type *createOpenCLOpaqueType(const clang::Type *T); const llvm::abi::RecordType *convertCXXRecordType(const CXXRecordDecl *RD); void computeFieldInfo(const clang::RecordDecl *RD, diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 230742255073e..07e2eac39305d 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -22,18 +22,95 @@ using namespace clang::CodeGen; namespace { -class AMDGPUABIInfo final : public AMDGPUABIInfoCommon<DefaultABIInfo> { +class AMDGPUABIInfo final : public DefaultABIInfo { +private: + static const unsigned MaxNumRegsForArgsRet = 16; + + uint64_t numRegsForType(QualType Ty) const; + + bool isHomogeneousAggregateBaseType(QualType Ty) const override; + bool isHomogeneousAggregateSmallEnough(const Type *Base, + uint64_t Members) const override; + + // Coerce HIP scalar pointer arguments from generic pointers to global ones. + llvm::Type *coerceKernelArgumentType(llvm::Type *Ty, unsigned FromAS, + unsigned ToAS) const { + // Single value types. + auto *PtrTy = llvm::dyn_cast<llvm::PointerType>(Ty); + if (PtrTy && PtrTy->getAddressSpace() == FromAS) + return llvm::PointerType::get(Ty->getContext(), ToAS); + return Ty; + } + public: - explicit AMDGPUABIInfo(CodeGen::CodeGenTypes &CGT) - : AMDGPUABIInfoCommon(CGT) {} + explicit AMDGPUABIInfo(CodeGen::CodeGenTypes &CGT) : + DefaultABIInfo(CGT) {} + ABIArgInfo classifyReturnType(QualType RetTy) const; ABIArgInfo classifyKernelArgumentType(QualType Ty) const; + ABIArgInfo classifyArgumentType(QualType Ty, bool Variadic, + unsigned &NumRegsLeft) const; void computeInfo(CGFunctionInfo &FI) const override; RValue EmitVAArg(CodeGenFunction &CGF, Address VAListAddr, QualType Ty, AggValueSlot Slot) const override; + + llvm::FixedVectorType * + getOptimalVectorMemoryType(llvm::FixedVectorType *T, + const LangOptions &Opt) const override { + // We have legal instructions for 96-bit so 3x32 can be supported. + // FIXME: This check should be a subtarget feature as technically SI doesn't + // support it. + if (T->getNumElements() == 3 && getDataLayout().getTypeSizeInBits(T) == 96) + return T; + return DefaultABIInfo::getOptimalVectorMemoryType(T, Opt); + } }; +bool AMDGPUABIInfo::isHomogeneousAggregateBaseType(QualType Ty) const { + return true; +} + +bool AMDGPUABIInfo::isHomogeneousAggregateSmallEnough( + const Type *Base, uint64_t Members) const { + uint32_t NumRegs = (getContext().getTypeSize(Base) + 31) / 32; + + // Homogeneous Aggregates may occupy at most 16 registers. + return Members * NumRegs <= MaxNumRegsForArgsRet; +} + +/// Estimate number of registers the type will use when passed in registers. +uint64_t AMDGPUABIInfo::numRegsForType(QualType Ty) const { + uint64_t NumRegs = 0; + + if (const VectorType *VT = Ty->getAs<VectorType>()) { + // Compute from the number of elements. The reported size is based on the + // in-memory size, which includes the padding 4th element for 3-vectors. + QualType EltTy = VT->getElementType(); + uint64_t EltSize = getContext().getTypeSize(EltTy); + + // 16-bit element vectors should be passed as packed. + if (EltSize == 16) + return (VT->getNumElements() + 1) / 2; + + uint64_t EltNumRegs = (EltSize + 31) / 32; + return EltNumRegs * VT->getNumElements(); + } + + if (const auto *RD = Ty->getAsRecordDecl()) { + assert(!RD->hasFlexibleArrayMember()); + + for (const FieldDecl *Field : RD->fields()) { + QualType FieldTy = Field->getType(); + NumRegs += numRegsForType(FieldTy); + } + + return NumRegs; + } + + return (getContext().getTypeSize(Ty) + 31) / 32; +} + void AMDGPUABIInfo::computeInfo(CGFunctionInfo &FI) const { llvm::CallingConv::ID CC = FI.getCallingConvention(); @@ -43,13 +120,13 @@ void AMDGPUABIInfo::computeInfo(CGFunctionInfo &FI) const { unsigned ArgumentIndex = 0; const unsigned numFixedArguments = FI.getNumRequiredArgs(); - NumRegsLeft = MaxNumRegsForArgsRet; + unsigned NumRegsLeft = MaxNumRegsForArgsRet; for (auto &Arg : FI.arguments()) { if (CC == llvm::CallingConv::AMDGPU_KERNEL) { Arg.info = classifyKernelArgumentType(Arg.type); } else { bool FixedArgument = ArgumentIndex++ < numFixedArguments; - Arg.info = classifyArgumentType(Arg.type, !FixedArgument); + Arg.info = classifyArgumentType(Arg.type, !FixedArgument, NumRegsLeft); } } } @@ -63,6 +140,45 @@ RValue AMDGPUABIInfo::EmitVAArg(CodeGenFunction &CGF, Address VAListAddr, CharUnits::fromQuantity(4), AllowHigherAlign, Slot); } +ABIArgInfo AMDGPUABIInfo::classifyReturnType(QualType RetTy) const { + if (isAggregateTypeForABI(RetTy)) { + // Records with non-trivial destructors/copy-constructors should not be + // returned by value. + if (!getRecordArgABI(RetTy, getCXXABI())) { + // Ignore empty structs/unions. + if (isEmptyRecord(getContext(), RetTy, true)) + return ABIArgInfo::getIgnore(); + + // Lower single-element structs to just return a regular value. + if (const Type *SeltTy = isSingleElementStruct(RetTy, getContext())) + return ABIArgInfo::getDirect(CGT.ConvertType(QualType(SeltTy, 0))); + + if (const auto *RD = RetTy->getAsRecordDecl(); + RD && RD->hasFlexibleArrayMember()) + return DefaultABIInfo::classifyReturnType(RetTy); + + // Pack aggregates <= 4 bytes into single VGPR or pair. + uint64_t Size = getContext().getTypeSize(RetTy); + if (Size <= 16) + return ABIArgInfo::getDirect(llvm::Type::getInt16Ty(getVMContext())); + + if (Size <= 32) + return ABIArgInfo::getDirect(llvm::Type::getInt32Ty(getVMContext())); + + if (Size <= 64) { + llvm::Type *I32Ty = llvm::Type::getInt32Ty(getVMContext()); + return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); + } + + if (numRegsForType(RetTy) <= MaxNumRegsForArgsRet) + return ABIArgInfo::getDirect(); + } + } + + // Otherwise just do the default thing. + return DefaultABIInfo::classifyReturnType(RetTy); +} + /// For kernels all parameters are really passed in a special buffer. It doesn't /// make sense to pass anything byval, so everything must be direct. ABIArgInfo AMDGPUABIInfo::classifyKernelArgumentType(QualType Ty) const { @@ -97,6 +213,83 @@ ABIArgInfo AMDGPUABIInfo::classifyKernelArgumentType(QualType Ty) const { return ABIArgInfo::getDirect(LTy, 0, nullptr, false); } +ABIArgInfo AMDGPUABIInfo::classifyArgumentType(QualType Ty, bool Variadic, + unsigned &NumRegsLeft) const { + assert(NumRegsLeft <= MaxNumRegsForArgsRet && "register estimate underflow"); + + Ty = useFirstFieldIfTransparentUnion(Ty); + + if (Variadic) { + return ABIArgInfo::getDirect(/*T=*/nullptr, + /*Offset=*/0, + /*Padding=*/nullptr, + /*CanBeFlattened=*/false, + /*Align=*/0); + } + + if (isAggregateTypeForABI(Ty)) { + // Records with non-trivial destructors/copy-constructors should not be + // passed by value. + if (auto RAA = getRecordArgABI(Ty, getCXXABI())) + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + RAA == CGCXXABI::RAA_DirectInMemory); + + // Ignore empty structs/unions. + if (isEmptyRecord(getContext(), Ty, true)) + return ABIArgInfo::getIgnore(); + + // Lower single-element structs to just pass a regular value. TODO: We + // could do reasonable-size multiple-element structs too, using getExpand(), + // though watch out for things like bitfields. + if (const Type *SeltTy = isSingleElementStruct(Ty, getContext())) + return ABIArgInfo::getDirect(CGT.ConvertType(QualType(SeltTy, 0))); + + if (const auto *RD = Ty->getAsRecordDecl(); + RD && RD->hasFlexibleArrayMember()) + return DefaultABIInfo::classifyArgumentType(Ty); + + // Pack aggregates <= 8 bytes into single VGPR or pair. + uint64_t Size = getContext().getTypeSize(Ty); + if (Size <= 64) { + unsigned NumRegs = (Size + 31) / 32; + NumRegsLeft -= std::min(NumRegsLeft, NumRegs); + + if (Size <= 16) + return ABIArgInfo::getDirect(llvm::Type::getInt16Ty(getVMContext())); + + if (Size <= 32) + return ABIArgInfo::getDirect(llvm::Type::getInt32Ty(getVMContext())); + + // XXX: Should this be i64 instead, and should the limit increase? + llvm::Type *I32Ty = llvm::Type::getInt32Ty(getVMContext()); + return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); + } + + if (NumRegsLeft > 0) { + uint64_t NumRegs = numRegsForType(Ty); + if (NumRegsLeft >= NumRegs) { + NumRegsLeft -= NumRegs; + return ABIArgInfo::getDirect(); + } + } + + // Use pass-by-reference in stead of pass-by-value for struct arguments in + // function ABI. + return ABIArgInfo::getIndirectAliased( + getContext().getTypeAlignInChars(Ty), + getContext().getTargetAddressSpace(LangAS::opencl_private)); + } + + // Otherwise just do the default thing. + ABIArgInfo ArgInfo = DefaultABIInfo::classifyArgumentType(Ty); + if (!ArgInfo.isIndirect()) { + uint64_t NumRegs = numRegsForType(Ty); + NumRegsLeft -= std::min(NumRegs, uint64_t{NumRegsLeft}); + } + + return ArgInfo; +} + class AMDGPUTargetCodeGenInfo : public TargetCodeGenInfo { public: AMDGPUTargetCodeGenInfo(CodeGenTypes &CGT) diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp index 7dbcf5e439095..e8148d1566f85 100644 --- a/clang/lib/CodeGen/Targets/SPIR.cpp +++ b/clang/lib/CodeGen/Targets/SPIR.cpp @@ -48,12 +48,40 @@ class SPIRVABIInfo : public CommonSPIRABIInfo { ABIArgInfo classifyKernelArgumentType(QualType Ty) const; }; -class AMDGCNSPIRVABIInfo : public AMDGPUABIInfoCommon<SPIRVABIInfo> { +class AMDGCNSPIRVABIInfo : public SPIRVABIInfo { + // TODO: this should be unified / shared with AMDGPU, ideally we'd like to + // re-use AMDGPUABIInfo eventually, rather than duplicate. + static constexpr unsigned MaxNumRegsForArgsRet = 16; // 16 32-bit registers + mutable unsigned NumRegsLeft = 0; + + uint64_t numRegsForType(QualType Ty) const; + + bool isHomogeneousAggregateBaseType(QualType Ty) const override { + return true; + } + bool isHomogeneousAggregateSmallEnough(const Type *Base, + uint64_t Members) const override { + uint32_t NumRegs = (getContext().getTypeSize(Base) + 31) / 32; + + // Homogeneous Aggregates may occupy at most 16 registers. + return Members * NumRegs <= MaxNumRegsForArgsRet; + } + + // Coerce HIP scalar pointer arguments from generic pointers to global ones. + llvm::Type *coerceKernelArgumentType(llvm::Type *Ty, unsigned FromAS, + unsigned ToAS) const; + + ABIArgInfo classifyReturnType(QualType RetTy) const; ABIArgInfo classifyKernelArgumentType(QualType Ty) const; + ABIArgInfo classifyArgumentType(QualType Ty, bool Variadic) const; public: - AMDGCNSPIRVABIInfo(CodeGenTypes &CGT) : AMDGPUABIInfoCommon(CGT) {} + AMDGCNSPIRVABIInfo(CodeGenTypes &CGT) : SPIRVABIInfo(CGT) {} void computeInfo(CGFunctionInfo &FI) const override; + + llvm::FixedVectorType * + getOptimalVectorMemoryType(llvm::FixedVectorType *Ty, + const LangOptions &LangOpt) const override; }; } // end anonymous namespace namespace { @@ -181,6 +209,84 @@ RValue SPIRVABIInfo::EmitVAArg(CodeGenFunction &CGF, Address VAListAddr, /*AllowHigherAlign=*/true, Slot); } +uint64_t AMDGCNSPIRVABIInfo::numRegsForType(QualType Ty) const { + // This duplicates the AMDGPUABI computation. + uint64_t NumRegs = 0; + + if (const VectorType *VT = Ty->getAs<VectorType>()) { + // Compute from the number of elements. The reported size is based on the + // in-memory size, which includes the padding 4th element for 3-vectors. + QualType EltTy = VT->getElementType(); + uint64_t EltSize = getContext().getTypeSize(EltTy); + + // 16-bit element vectors should be passed as packed. + if (EltSize == 16) + return (VT->getNumElements() + 1) / 2; + + uint64_t EltNumRegs = (EltSize + 31) / 32; + return EltNumRegs * VT->getNumElements(); + } + + if (const auto *RD = Ty->getAsRecordDecl()) { + assert(!RD->hasFlexibleArrayMember()); + + for (const FieldDecl *Field : RD->fields()) { + QualType FieldTy = Field->getType(); + NumRegs += numRegsForType(FieldTy); + } + + return NumRegs; + } + + return (getContext().getTypeSize(Ty) + 31) / 32; +} + +llvm::Type *AMDGCNSPIRVABIInfo::coerceKernelArgumentType(llvm::Type *Ty, + unsigned FromAS, + unsigned ToAS) const { + // Single value types. + auto *PtrTy = llvm::dyn_cast<llvm::PointerType>(Ty); + if (PtrTy && PtrTy->getAddressSpace() == FromAS) + return llvm::PointerType::get(Ty->getContext(), ToAS); + return Ty; +} + +ABIArgInfo AMDGCNSPIRVABIInfo::classifyReturnType(QualType RetTy) const { + if (!isAggregateTypeForABI(RetTy) || getRecordArgABI(RetTy, getCXXABI())) + return DefaultABIInfo::classifyReturnType(RetTy); + + // Ignore empty structs/unions. + if (isEmptyRecord(getContext(), RetTy, true)) + return ABIArgInfo::getIgnore(); + + // Lower single-element structs to just return a regular value. + if (const Type *SeltTy = isSingleElementStruct(RetTy, getContext())) + return ABIArgInfo::getDirect(CGT.ConvertType(QualType(SeltTy, 0))); + + if (const auto *RD = RetTy->getAsRecordDecl(); + RD && RD->hasFlexibleArrayMember()) + return DefaultABIInfo::classifyReturnType(RetTy); + + // Pack aggregates <= 4 bytes into single VGPR or pair. + uint64_t Size = getContext().getTypeSize(RetTy); + if (Size <= 16) + return ABIArgInfo::getDirect(llvm::Type::getInt16Ty(getVMContext())); + + if (Size <= 32) + return ABIArgInfo::getDirect(llvm::Type::getInt32Ty(getVMContext())); + + // TODO: This carried over from AMDGPU oddity, we retain it to + // ensure consistency, but it might be reasonable to return Int64. + if (Size <= 64) { + llvm::Type *I32Ty = llvm::Type::getInt32Ty(getVMContext()); + return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); + } + + if (numRegsForType(RetTy) <= MaxNumRegsForArgsRet) + return ABIArgInfo::getDirect(); + return DefaultABIInfo::classifyReturnType(RetTy); +} + /// For kernels all parameters are really passed in a special buffer. It doesn't /// make sense to pass anything byval, so everything must be direct. ABIArgInfo AMDGCNSPIRVABIInfo::classifyKernelArgumentType(QualType Ty) const { @@ -214,6 +320,83 @@ ABIArgInfo AMDGCNSPIRVABIInfo::classifyKernelArgumentType(QualType Ty) const { return ABIArgInfo::getDirect(LTy, 0, nullptr, false); } +ABIArgInfo AMDGCNSPIRVABIInfo::classifyArgumentType(QualType Ty, + bool Variadic) const { + assert(NumRegsLeft <= MaxNumRegsForArgsRet && "register estimate underflow"); + + Ty = useFirstFieldIfTransparentUnion(Ty); + + if (Variadic) { + return ABIArgInfo::getDirect(/*T=*/nullptr, + /*Offset=*/0, + /*Padding=*/nullptr, + /*CanBeFlattened=*/false, + /*Align=*/0); + } + + if (!isAggregateTypeForABI(Ty)) { + ABIArgInfo ArgInfo = DefaultABIInfo::classifyArgumentType(Ty); + if (!ArgInfo.isIndirect()) { + uint64_t NumRegs = numRegsForType(Ty); + NumRegsLeft -= std::min(NumRegs, uint64_t{NumRegsLeft}); + } + + return ArgInfo; + } + + // Records with non-trivial destructors/copy-constructors should not be + // passed by value. + if (auto RAA = getRecordArgABI(Ty, getCXXABI())) + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + RAA == CGCXXABI::RAA_DirectInMemory); + + // Ignore empty structs/unions. + if (isEmptyRecord(getContext(), Ty, true)) + return ABIArgInfo::getIgnore(); + + // Lower single-element structs to just pass a regular value. TODO: We + // could do reasonable-size multiple-element structs too, using getExpand(), + // though watch out for things like bitfields. + if (const Type *SeltTy = isSingleElementStruct(Ty, getContext())) + return ABIArgInfo::getDirect(CGT.ConvertType(QualType(SeltTy, 0))); + + if (const auto *RD = Ty->getAsRecordDecl(); + RD && RD->hasFlexibleArrayMember()) + return DefaultABIInfo::classifyArgumentType(Ty); + + uint64_t Size = getContext().getTypeSize(Ty); + if (Size <= 64) { + // Pack aggregates <= 8 bytes into single VGPR or pair. + unsigned NumRegs = (Size + 31) / 32; + NumRegsLeft -= std::min(NumRegsLeft, NumRegs); + + if (Size <= 16) + return ABIArgInfo::getDirect(llvm::Type::getInt16Ty(getVMContext())); + + if (Size <= 32) + return ABIArgInfo::getDirect(llvm::Type::getInt32Ty(getVMContext())); + + // TODO: This is an AMDGPU oddity, and might be vestigial, we retain it to + // ensure consistency, but it should be revisited. + llvm::Type *I32Ty = llvm::Type::getInt32Ty(getVMContext()); + return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); + } + + if (NumRegsLeft > 0) { + uint64_t NumRegs = numRegsForType(Ty); + if (NumRegsLeft >= NumRegs) { + NumRegsLeft -= NumRegs; + return ABIArgInfo::getDirect(); + } + } + + // Use pass-by-reference in stead of pass-by-value for struct arguments in + // function ABI. + return ABIArgInfo::getIndirectAliased( + getContext().getTypeAlignInChars(Ty), + getContext().getTargetAddressSpace(LangAS::opencl_private)); +} + void AMDGCNSPIRVABIInfo::computeInfo(CGFunctionInfo &FI) const { llvm::CallingConv::ID CC = FI.getCallingConvention(); @@ -245,6 +428,14 @@ SPIRVABIInfo::getOptimalVectorMemoryType(llvm::FixedVectorType *Ty, return DefaultABIInfo::getOptimalVectorMemoryType(Ty, LangOpt); } +llvm::FixedVectorType *AMDGCNSPIRVABIInfo::getOptimalVectorMemoryType( + llvm::FixedVectorType *Ty, const LangOptions &LangOpt) const { + // AMDGPU has legal instructions for 96-bit so 3x32 can be supported. + if (Ty->getNumElements() == 3 && getDataLayout().getTypeSizeInBits(Ty) == 96) + return Ty; + return DefaultABIInfo::getOptimalVectorMemoryType(Ty, LangOpt); +} + namespace clang { namespace CodeGen { void computeSPIRKernelABIInfo(CodeGenModule &CGM, CGFunctionInfo &FI) { diff --git a/clang/test/CodeGenHIP/amdgcnspirv-uses-amdgpu-abi.cpp b/clang/test/CodeGenHIP/amdgcnspirv-uses-amdgpu-abi.cpp index 9741c58762ac6..1cc49b0505a5f 100644 --- a/clang/test/CodeGenHIP/amdgcnspirv-uses-amdgpu-abi.cpp +++ b/clang/test/CodeGenHIP/amdgcnspirv-uses-amdgpu-abi.cpp @@ -3,6 +3,10 @@ // RUN: -o - %s | FileCheck --check-prefix=AMDGCNSPIRV %s // RUN: %clang_cc1 -triple amdgpu9.06-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -O3 \ // RUN: -o - %s | FileCheck --check-prefix=AMDGPU %s +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -O3 \ +// RUN: -fexperimental-abi-lowering -o - %s | FileCheck --check-prefix=AMDGCNSPIRV %s +// RUN: %clang_cc1 -triple amdgpu9.06-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -O3 \ +// RUN: -fexperimental-abi-lowering -o - %s | FileCheck --check-prefix=AMDGPU %s #define __global__ __attribute__((global)) #define __device__ __attribute__((device)) diff --git a/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl b/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl index c153fbf1a7d52..0383e5eaee403 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl @@ -1,6 +1,7 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple amdgpu-unknown-unknown -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple r600-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu-unknown-unknown -fexperimental-abi-lowering -emit-llvm -o - %s | FileCheck %s typedef __attribute__(( ext_vector_type(2) )) char char2; typedef __attribute__(( ext_vector_type(3) )) char char3; diff --git a/clang/test/CodeGenOpenCL/opencl_types.cl b/clang/test/CodeGenOpenCL/opencl_types.cl index aac3492b7a9e8..f0a6ac3c24f36 100644 --- a/clang/test/CodeGenOpenCL/opencl_types.cl +++ b/clang/test/CodeGenOpenCL/opencl_types.cl @@ -1,5 +1,6 @@ // RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "spir-unknown-unknown" -emit-llvm -o - -O0 | FileCheck %s --check-prefix=CHECK-SPIR // RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "amdgcn--amdhsa" -emit-llvm -o - -O0 | FileCheck %s --check-prefix=CHECK-AMDGCN +// RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "amdgcn--amdhsa" -fexperimental-abi-lowering -emit-llvm -o - -O0 | FileCheck %s --check-prefix=CHECK-AMDGCN #define CLK_ADDRESS_CLAMP_TO_EDGE 2 #define CLK_NORMALIZED_COORDS_TRUE 1 diff --git a/llvm/include/llvm/ABI/FunctionInfo.h b/llvm/include/llvm/ABI/FunctionInfo.h index caedafcbe9d22..248f2936dea13 100644 --- a/llvm/include/llvm/ABI/FunctionInfo.h +++ b/llvm/include/llvm/ABI/FunctionInfo.h @@ -39,6 +39,11 @@ class ArgInfo { /// Pass the argument indirectly via a hidden pointer with the specified /// alignment and address space. Indirect, + /// Like Indirect, but the object may be referenced elsewhere. Nothing + /// modifies it through another reference during the call, and the callee + /// must not modify it either, so a callee that cannot prove otherwise has + /// to copy it locally first. + IndirectAliased, /// Ignore the argument (treat as void). Useful for void and empty structs. Ignore, }; @@ -71,10 +76,11 @@ class ArgInfo { bool ZeroExt : 1; bool IndirectByVal : 1; bool IndirectRealign : 1; + bool CanBeFlattened : 1; ArgInfo(Kind K = Direct) : TheKind(K), SignExt(false), ZeroExt(false), IndirectByVal(false), - IndirectRealign(false) {} + IndirectRealign(false), CanBeFlattened(false) {} public: /// \param T The type to coerce to. If null, the argument's original type is @@ -85,12 +91,16 @@ class ArgInfo { /// return value on x86-64). /// \param Align Override for the argument's alignment. If absent, the /// default alignment for \p T is used. + /// \param CanBeFlattened Whether a record may be passed as its individual + /// elements rather than as one value. static ArgInfo getDirect(const Type *T = nullptr, unsigned Offset = 0, - MaybeAlign Align = std::nullopt) { + MaybeAlign Align = std::nullopt, + bool CanBeFlattened = true) { ArgInfo AI(Direct); AI.CoercionType = T; AI.Alignment = Align; AI.DirectAttr.Offset = Offset; + AI.CanBeFlattened = CanBeFlattened; return AI; } @@ -124,6 +134,16 @@ class ArgInfo { return AI; } + /// \p AddrSpace is the address space the object lives in. + static ArgInfo getIndirectAliased(Align Align, unsigned AddrSpace, + bool Realign = false) { + ArgInfo AI(IndirectAliased); + AI.Alignment = Align; + AI.IndirectAttr.AddrSpace = AddrSpace; + AI.IndirectRealign = Realign; + return AI; + } + static ArgInfo getIgnore() { return ArgInfo(Ignore); } ArgInfo &setSignExt(bool SignExtend = true) { @@ -143,6 +163,7 @@ class ArgInfo { Kind getKind() const { return TheKind; } bool isDirect() const { return TheKind == Direct; } bool isIndirect() const { return TheKind == Indirect; } + bool isIndirectAliased() const { return TheKind == IndirectAliased; } bool isIgnore() const { return TheKind == Ignore; } bool isExtend() const { return TheKind == Extend; } @@ -156,15 +177,20 @@ class ArgInfo { return Alignment; } + bool getCanBeFlattened() const { + assert(isDirect() && "Invalid Kind!"); + return CanBeFlattened; + } + Align getIndirectAlign() const { - assert(isIndirect() && "Invalid Kind!"); + assert((isIndirect() || isIndirectAliased()) && "Invalid Kind!"); assert(Alignment.has_value() && "Indirect arguments must have an alignment"); return *Alignment; } unsigned getIndirectAddrSpace() const { - assert(isIndirect() && "Invalid Kind!"); + assert((isIndirect() || isIndirectAliased()) && "Invalid Kind!"); return IndirectAttr.AddrSpace; } @@ -174,7 +200,7 @@ class ArgInfo { } bool getIndirectRealign() const { - assert(isIndirect() && "Invalid Kind!"); + assert((isIndirect() || isIndirectAliased()) && "Invalid Kind!"); return IndirectRealign; } diff --git a/llvm/include/llvm/ABI/TargetInfo.h b/llvm/include/llvm/ABI/TargetInfo.h index ea621597e5b22..8a9c5b4985247 100644 --- a/llvm/include/llvm/ABI/TargetInfo.h +++ b/llvm/include/llvm/ABI/TargetInfo.h @@ -80,16 +80,23 @@ class TargetInfo { LLVM_ABI RecordArgABI getRecordArgABI(const RecordType *RT) const; LLVM_ABI RecordArgABI getRecordArgABI(const Type *Ty) const; LLVM_ABI bool isPromotableInteger(const IntegerType *IT) const; - LLVM_ABI ArgInfo getNaturalAlignIndirect(const Type *Ty, - bool ByVal = true) const; + LLVM_ABI ArgInfo getNaturalAlignIndirect(const Type *Ty, bool ByVal = true, + unsigned AddrSpace = 0) const; LLVM_ABI bool isAggregateTypeForABI(const Type *Ty) const; + /// Returns the element type if \p Ty is a struct wrapping exactly one + /// non-empty element with no padding beyond it, else nullptr. Single-element + /// arrays are looked through. + LLVM_ABI const Type *isSingleElementStruct(const Type *Ty) const; + /// If Ty is a transparent union, return its first field type; otherwise /// return Ty unchanged. LLVM_ABI const Type *useFirstFieldIfTransparentUnion(const Type *Ty) const; /// Apply rules for classifying return types that are common to all targets. - LLVM_ABI bool maybeCommonClassifyReturnType(FunctionInfo &FI) const; + /// AddrSpace is the address space of the sret pointer. + LLVM_ABI bool maybeCommonClassifyReturnType(FunctionInfo &FI, + unsigned AddrSpace = 0) const; }; LLVM_ABI std::unique_ptr<TargetInfo> createBPFTargetInfo(TypeBuilder &TB); @@ -116,6 +123,29 @@ enum class AArch64ABIKind { LLVM_ABI std::unique_ptr<TargetInfo> createAArch64TargetInfo(TypeBuilder &TB, AArch64ABIKind Kind); +/// The parts of the AMDGPU ABI that differ between amdgcn and the +/// AMDGCN-flavoured SPIR-V that lowers to it. Address spaces are target +/// address space numbers, not language address spaces. +struct AMDGPUABIOptions { + CallingConv::ID KernelCC = CallingConv::C; + unsigned AllocaAddrSpace = 0; + /// Aggregate arguments that do not fit in registers are passed here. + unsigned PrivateAddrSpace = 0; + /// Indirect kernel arguments are passed here. + unsigned ConstantAddrSpace = 0; + /// Kernel pointer arguments are coerced from this space to + /// KernelArgAddrSpace. + unsigned GenericAddrSpace = 0; + unsigned KernelArgAddrSpace = 0; + /// False outside of device compilation, where no coercion applies. + bool CoerceKernelPointerArgs = false; + /// Sets the width above which a _BitInt is passed indirectly. + bool HasInt128 = false; +}; + +LLVM_ABI std::unique_ptr<TargetInfo> +createAMDGPUTargetInfo(TypeBuilder &TB, const AMDGPUABIOptions &Opts); + } // namespace abi } // namespace llvm diff --git a/llvm/include/llvm/ABI/Types.h b/llvm/include/llvm/ABI/Types.h index 5145c201ff1b9..fe690b2160bda 100644 --- a/llvm/include/llvm/ABI/Types.h +++ b/llvm/include/llvm/ABI/Types.h @@ -74,6 +74,10 @@ class Type { return alignTo(getTypeStoreSize(), getAlignment().value()); } + /// Unlike getSizeInBits this includes trailing padding, matching Clang + /// ASTContext::getTypeSize. + TypeSize getTypeAllocSizeInBits() const { return getTypeAllocSize() * 8; } + bool isVoid() const { return Kind == TypeKind::Void; } bool isAtomic() const { return Kind == TypeKind::Atomic; } bool isInteger() const { return Kind == TypeKind::Integer; } diff --git a/llvm/lib/ABI/CMakeLists.txt b/llvm/lib/ABI/CMakeLists.txt index 39e725ca9fd3b..54307bae8e07d 100644 --- a/llvm/lib/ABI/CMakeLists.txt +++ b/llvm/lib/ABI/CMakeLists.txt @@ -4,6 +4,7 @@ add_llvm_component_library(LLVMABI TargetInfo.cpp IRTypeMapper.cpp Targets/AArch64.cpp + Targets/AMDGPU.cpp Targets/BPF.cpp Targets/X86.cpp diff --git a/llvm/lib/ABI/TargetInfo.cpp b/llvm/lib/ABI/TargetInfo.cpp index 507e3eb5bc120..3a0506dd3601a 100644 --- a/llvm/lib/ABI/TargetInfo.cpp +++ b/llvm/lib/ABI/TargetInfo.cpp @@ -19,6 +19,10 @@ bool TargetInfo::isAggregateTypeForABI(const Type *Ty) const { if (Ty->isInteger() || Ty->isFloat() || Ty->isPointer() || Ty->isVector()) return false; + // Data member pointers have scalar evaluation kind. + if (const auto *MPT = dyn_cast<MemberPointerType>(Ty)) + return MPT->isFunctionPointer(); + // A matrix type is modeled as an array but lowers to a single flattened // vector and has scalar evaluation kind in classic CodeGen, so it is not an // aggregate for ABI purposes. @@ -37,8 +41,66 @@ bool TargetInfo::isPromotableInteger(const IntegerType *IT) const { return BitWidth < 32; } -ArgInfo TargetInfo::getNaturalAlignIndirect(const Type *Ty, bool ByVal) const { - return ArgInfo::getIndirect(Ty->getAlignment(), ByVal); +ArgInfo TargetInfo::getNaturalAlignIndirect(const Type *Ty, bool ByVal, + unsigned AddrSpace) const { + return ArgInfo::getIndirect(Ty->getAlignment(), ByVal, AddrSpace); +} + +const Type *TargetInfo::isSingleElementStruct(const Type *Ty) const { + const auto *RT = dyn_cast<RecordType>(Ty); + if (!RT) + return nullptr; + + if (RT->hasFlexibleArrayMember()) + return nullptr; + + const Type *Found = nullptr; + + for (const FieldInfo &Base : RT->getBaseClasses()) { + const auto *BaseRT = dyn_cast<RecordType>(Base.FieldType); + if (!BaseRT || BaseRT->isEmpty()) + continue; + + if (Found) + return nullptr; + + Found = isSingleElementStruct(Base.FieldType); + if (!Found) + return nullptr; + } + + for (const FieldInfo &Field : RT->getFields()) { + if (Field.isEmpty()) + continue; + + if (Found) + return nullptr; + + const Type *FieldTy = Field.FieldType; + + // Treat single element arrays as the element. + while (const auto *AT = dyn_cast<ArrayType>(FieldTy)) { + if (AT->getNumElements() != 1) + break; + FieldTy = AT->getElementType(); + } + + if (!isAggregateTypeForABI(FieldTy)) { + Found = FieldTy; + } else { + Found = isSingleElementStruct(FieldTy); + if (!Found) + return nullptr; + } + } + + // Padding beyond the element disqualifies the struct. Compare in-memory + // sizes, not raw bit widths, so an element with trailing padding of its own + // still matches the record wrapping it. + if (Found && Found->getTypeAllocSize() != Ty->getTypeAllocSize()) + return nullptr; + + return Found; } RecordArgABI TargetInfo::getRecordArgABI(const RecordType *RT) const { @@ -67,7 +129,8 @@ const Type *TargetInfo::useFirstFieldIfTransparentUnion(const Type *Ty) const { return Ty; } -bool TargetInfo::maybeCommonClassifyReturnType(FunctionInfo &FI) const { +bool TargetInfo::maybeCommonClassifyReturnType(FunctionInfo &FI, + unsigned AddrSpace) const { const abi::Type *Ty = FI.getReturnType(); // TODO: When Microsoft ABI is supported, CXX records may need different @@ -79,7 +142,7 @@ bool TargetInfo::maybeCommonClassifyReturnType(FunctionInfo &FI) const { // distinct from getIndirectReturnResult (plain aggregates), which uses // ByVal=true. FI.getReturnInfo() = - ArgInfo::getIndirect(RT->getAlignment(), /*ByVal=*/false); + ArgInfo::getIndirect(RT->getAlignment(), /*ByVal=*/false, AddrSpace); return true; } } diff --git a/llvm/lib/ABI/Targets/AMDGPU.cpp b/llvm/lib/ABI/Targets/AMDGPU.cpp new file mode 100644 index 0000000000000..9acf091d76059 --- /dev/null +++ b/llvm/lib/ABI/Targets/AMDGPU.cpp @@ -0,0 +1,279 @@ +//===- AMDGPU.cpp - AMDGPU ABI Implementation -----------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// Shared by amdgcn and the AMDGCN-flavoured SPIR-V that lowers to it. The two +// differ only in the values carried by AMDGPUABIOptions. +// +//===----------------------------------------------------------------------===// + +#include "llvm/ABI/FunctionInfo.h" +#include "llvm/ABI/TargetInfo.h" +#include "llvm/ABI/Types.h" +#include "llvm/ADT/STLExtras.h" +#include "llvm/Support/Casting.h" +#include "llvm/Support/MathExtras.h" +#include <algorithm> + +namespace llvm::abi { + +class AMDGPUTargetInfo : public TargetInfo { +private: + /// Registers available for arguments and return values, in 32-bit units. + static constexpr unsigned MaxNumRegsForArgsRet = 16; + + TypeBuilder &TB; + AMDGPUABIOptions Opts; + const IntegerType *Int16Ty; + const IntegerType *Int32Ty; + const ArrayType *Int32PairTy; + + uint64_t numRegsForType(const Type *Ty) const { + if (const auto *VT = dyn_cast<VectorType>(Ty)) { + // Compute from the number of elements. The reported size is based on the + // in-memory size, which includes the padding 4th element for 3-vectors. + uint64_t EltSize = + VT->getElementType()->getTypeAllocSizeInBits().getFixedValue(); + uint64_t NumElts = VT->getNumElements().getFixedValue(); + + // 16-bit element vectors should be passed as packed. + if (EltSize == 16) + return (NumElts + 1) / 2; + + return divideCeil(EltSize, 32) * NumElts; + } + + if (const auto *RT = dyn_cast<RecordType>(Ty)) { + assert(!RT->hasFlexibleArrayMember()); + + // Bases are deliberately not counted, matching RecordDecl::fields() in + // the classic classifier. + uint64_t NumRegs = 0; + for (const FieldInfo &Field : RT->getFields()) + NumRegs += numRegsForType(Field.FieldType); + + return NumRegs; + } + + return divideCeil(Ty->getTypeAllocSizeInBits().getFixedValue(), 32); + } + + /// Coerce a scalar pointer argument from the generic address space to the + /// one kernel arguments must use. + const Type *coerceKernelArgumentType(const Type *Ty) const { + const auto *PtrTy = dyn_cast<PointerType>(Ty); + if (PtrTy && PtrTy->getAddrSpace() == Opts.GenericAddrSpace) + return TB.getPointerType(PtrTy->getSizeInBits().getFixedValue(), + PtrTy->getAlignment(), Opts.KernelArgAddrSpace); + return Ty; + } + + /// The non-aggregate tail shared by the two Clang DefaultABIInfo rules. + ArgInfo defaultClassifyScalar(const Type *Ty) const { + if (const auto *IntTy = dyn_cast<IntegerType>(Ty)) { + if (IntTy->isBitInt() && + IntTy->getSizeInBits().getFixedValue() > getMaxDirectBitIntWidth()) + return getNaturalAlignIndirect(Ty, /*ByVal=*/true, + Opts.AllocaAddrSpace); + + if (isPromotableInteger(IntTy)) + return ArgInfo::getExtend(Ty); + } + + return ArgInfo::getDirect(); + } + + /// The Clang DefaultABIInfo rules, the fallback for anything AMDGPU does not + /// pack into registers itself. + ArgInfo defaultClassifyArgumentType(const Type *Ty) const { + Ty = useFirstFieldIfTransparentUnion(Ty); + + if (isAggregateTypeForABI(Ty)) + return getNaturalAlignIndirect(Ty, getRecordArgABI(Ty) != RAA_Indirect, + Opts.AllocaAddrSpace); + + return defaultClassifyScalar(Ty); + } + + ArgInfo defaultClassifyReturnType(const Type *RetTy) const { + if (RetTy->isVoid()) + return ArgInfo::getIgnore(); + + if (isAggregateTypeForABI(RetTy)) + return getNaturalAlignIndirect(RetTy, /*ByVal=*/true, + Opts.AllocaAddrSpace); + + return defaultClassifyScalar(RetTy); + } + + ArgInfo classifyReturnType(const Type *RetTy) const { + const auto *RT = dyn_cast<RecordType>(RetTy); + + if (RetTy->isVoid() || !isAggregateTypeForABI(RetTy) || getRecordArgABI(RT)) + return defaultClassifyReturnType(RetTy); + + // Ignore empty structs/unions. + if (RT && RT->isEmpty()) + return ArgInfo::getIgnore(); + + // Lower single-element structs to just return a regular value. + if (const Type *SeltTy = isSingleElementStruct(RetTy)) + return ArgInfo::getDirect(SeltTy); + + if (RT && RT->hasFlexibleArrayMember()) + return defaultClassifyReturnType(RetTy); + + // Pack aggregates <= 4 bytes into single VGPR or pair. + uint64_t Size = RetTy->getTypeAllocSizeInBits().getFixedValue(); + if (Size <= 16) + return ArgInfo::getDirect(Int16Ty); + + if (Size <= 32) + return ArgInfo::getDirect(Int32Ty); + + if (Size <= 64) + return ArgInfo::getDirect(Int32PairTy); + + if (numRegsForType(RetTy) <= MaxNumRegsForArgsRet) + return ArgInfo::getDirect(); + + return defaultClassifyReturnType(RetTy); + } + + ArgInfo classifyArgumentType(const Type *Ty, bool Variadic, + unsigned &NumRegsLeft) const { + assert(NumRegsLeft <= MaxNumRegsForArgsRet && + "register estimate underflow"); + + Ty = useFirstFieldIfTransparentUnion(Ty); + + if (Variadic) + return ArgInfo::getDirect(/*T=*/nullptr, /*Offset=*/0, + /*Align=*/std::nullopt, + /*CanBeFlattened=*/false); + + if (!isAggregateTypeForABI(Ty)) { + ArgInfo Info = defaultClassifyScalar(Ty); + if (!Info.isIndirect()) + NumRegsLeft -= std::min<uint64_t>(numRegsForType(Ty), NumRegsLeft); + + return Info; + } + + const auto *RT = dyn_cast<RecordType>(Ty); + + // Records with non-trivial destructors/copy-constructors should not be + // passed by value. + if (RecordArgABI RAA = getRecordArgABI(RT)) + return getNaturalAlignIndirect(Ty, RAA == RAA_DirectInMemory, + Opts.AllocaAddrSpace); + + // Ignore empty structs/unions. + if (RT && RT->isEmpty()) + return ArgInfo::getIgnore(); + + // Lower single-element structs to just pass a regular value. TODO: We + // could do reasonable-size multiple-element structs too, using getExpand(), + // though watch out for things like bitfields. + if (const Type *SeltTy = isSingleElementStruct(Ty)) + return ArgInfo::getDirect(SeltTy); + + if (RT && RT->hasFlexibleArrayMember()) + return defaultClassifyArgumentType(Ty); + + // Pack aggregates <= 8 bytes into single VGPR or pair. + uint64_t Size = Ty->getTypeAllocSizeInBits().getFixedValue(); + if (Size <= 64) { + NumRegsLeft -= std::min<uint64_t>(NumRegsLeft, divideCeil(Size, 32)); + + if (Size <= 16) + return ArgInfo::getDirect(Int16Ty); + + if (Size <= 32) + return ArgInfo::getDirect(Int32Ty); + + // XXX: Should this be i64 instead, and should the limit increase? + return ArgInfo::getDirect(Int32PairTy); + } + + if (NumRegsLeft > 0) { + uint64_t NumRegs = numRegsForType(Ty); + if (NumRegsLeft >= NumRegs) { + NumRegsLeft -= NumRegs; + return ArgInfo::getDirect(); + } + } + + // Use pass-by-reference instead of pass-by-value for struct arguments in + // function ABI. + return ArgInfo::getIndirectAliased(Ty->getAlignment(), + Opts.PrivateAddrSpace); + } + + /// For kernels all parameters are really passed in a special buffer. It + /// doesn't make sense to pass anything byval, so everything must be direct. + ArgInfo classifyKernelArgumentType(const Type *Ty) const { + Ty = useFirstFieldIfTransparentUnion(Ty); + + // TODO: Can we omit empty structs? + + if (const Type *SeltTy = isSingleElementStruct(Ty)) + Ty = SeltTy; + + // FIXME: This doesn't apply the optimization of coercing pointers in + // structs to global address space when using byref. This would require + // implementing a new kind of coercion of the in-memory type when for + // indirect arguments. + if (isAggregateTypeForABI(Ty)) + return ArgInfo::getIndirectAliased(Ty->getAlignment(), + Opts.ConstantAddrSpace); + + const Type *CoercedTy = Ty; + if (Opts.CoerceKernelPointerArgs) + CoercedTy = coerceKernelArgumentType(Ty); + + // If we set CanBeFlattened to true, CodeGen will expand the struct to its + // individual elements, which confuses the Clover OpenCL backend; therefore + // we have to set it to false here. + return ArgInfo::getDirect(CoercedTy, /*Offset=*/0, /*Align=*/std::nullopt, + /*CanBeFlattened=*/false); + } + + uint64_t getMaxDirectBitIntWidth() const { return Opts.HasInt128 ? 128 : 64; } + +public: + AMDGPUTargetInfo(TypeBuilder &TB, const AMDGPUABIOptions &Opts) + : TB(TB), Opts(Opts), + Int16Ty(TB.getIntegerType(16, Align(2), /*Signed=*/false)), + Int32Ty(TB.getIntegerType(32, Align(4), /*Signed=*/false)), + Int32PairTy(TB.getArrayType(Int32Ty, 2, 64)) {} + + void computeInfo(FunctionInfo &FI) const override { + // A record that cannot be copied is constructed in place, so the sret + // pointer uses the generic address space rather than the alloca one. + if (!maybeCommonClassifyReturnType(FI, Opts.GenericAddrSpace)) + FI.getReturnInfo() = classifyReturnType(FI.getReturnType()); + + const bool IsKernel = FI.getCallingConvention() == Opts.KernelCC; + const unsigned NumRequiredArgs = FI.getNumRequiredArgs(); + unsigned NumRegsLeft = MaxNumRegsForArgsRet; + + for (auto [Index, Arg] : enumerate(FI.arguments())) { + Arg.Info = + IsKernel ? classifyKernelArgumentType(Arg.ABIType) + : classifyArgumentType(Arg.ABIType, Index >= NumRequiredArgs, + NumRegsLeft); + } + } +}; + +std::unique_ptr<TargetInfo> +createAMDGPUTargetInfo(TypeBuilder &TB, const AMDGPUABIOptions &Opts) { + return std::make_unique<AMDGPUTargetInfo>(TB, Opts); +} + +} // namespace llvm::abi diff --git a/llvm/lib/ABI/Targets/X86.cpp b/llvm/lib/ABI/Targets/X86.cpp index fd7a5e15b3548..c60bb751cd47f 100644 --- a/llvm/lib/ABI/Targets/X86.cpp +++ b/llvm/lib/ABI/Targets/X86.cpp @@ -100,7 +100,6 @@ class X86_64TargetInfo : public TargetInfo { ArgInfo getIndirectReturnResult(const Type *Ty) const; const Type *getFPTypeAtOffset(const Type *Ty, unsigned Offset) const; - const Type *isSingleElementStruct(const Type *Ty) const; const Type *getByteVectorType(const Type *Ty) const; const Type *createPairType(const Type *Lo, const Type *Hi) const; @@ -1243,60 +1242,6 @@ const Type *X86_64TargetInfo::getByteVectorType(const Type *Ty) const { ElementCount::getFixed(Size / 64), Align(Size / 8)); } -// Returns the single element if this is a single-element struct wrapper -const Type *X86_64TargetInfo::isSingleElementStruct(const Type *Ty) const { - const auto *RT = dyn_cast<RecordType>(Ty); - if (!RT) - return nullptr; - - if (RT->hasFlexibleArrayMember()) - return nullptr; - - const Type *Found = nullptr; - - for (const auto &Base : RT->getBaseClasses()) { - const Type *BaseTy = Base.FieldType; - auto *BaseRT = dyn_cast<RecordType>(BaseTy); - - if (!BaseRT || BaseRT->isEmpty()) - continue; - - const Type *Elem = isSingleElementStruct(BaseTy); - if (!Elem || Found) - return nullptr; - Found = Elem; - } - - for (const auto &FI : RT->getFields()) { - if (FI.isEmpty()) - continue; - - const Type *FTy = FI.FieldType; - - while (auto *AT = dyn_cast<ArrayType>(FTy)) { - if (AT->getNumElements() != 1) - break; - FTy = AT->getElementType(); - } - - const Type *Elem; - if (auto *InnerRT = dyn_cast<RecordType>(FTy)) - Elem = isSingleElementStruct(InnerRT); - else - Elem = FTy; - if (!Elem || Found) - return nullptr; - Found = Elem; - } - - if (!Found) - return nullptr; - if (Found->getSizeInBits() != Ty->getSizeInBits()) - return nullptr; - - return Found; -} - bool X86_64TargetInfo::isIllegalVectorType(const Type *Ty) const { if (const auto *VecTy = dyn_cast<VectorType>(Ty)) { uint64_t Size = VecTy->getSizeInBits().getFixedValue(); diff --git a/llvm/utils/gn/secondary/llvm/lib/ABI/BUILD.gn b/llvm/utils/gn/secondary/llvm/lib/ABI/BUILD.gn index aaf5642b53d15..74cb27466ea3f 100644 --- a/llvm/utils/gn/secondary/llvm/lib/ABI/BUILD.gn +++ b/llvm/utils/gn/secondary/llvm/lib/ABI/BUILD.gn @@ -10,6 +10,7 @@ static_library("ABI") { "IRTypeMapper.cpp", "TargetInfo.cpp", "Targets/AArch64.cpp", + "Targets/AMDGPU.cpp", "Targets/BPF.cpp", "Targets/X86.cpp", "Types.cpp", _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
