[clang] [llvm] [clang][SPIR-V][AMDGPU] Factor shared ABI classification into ABIInfoImpl (PR #223259)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Mon Sep 14 00:38:24 PDT 2026
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 <arseniy.obolenskiy at amd.com>
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 <arseniy.obolenskiy at amd.com>
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",
More information about the cfe-commits
mailing list