[clang] [llvm] [LLVMABI][AMDGPU] Add an AMDGPU target to the ABI library (PR #223259)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Thu Sep 17 00:56:21 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/3] [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/3] 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",
>From a1f32a916b35a988b8116365396caf3347f3d61d Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Thu, 17 Sep 2026 09:56:02 +0200
Subject: [PATCH 3/3] address comments
---
clang/lib/CodeGen/CodeGenModule.cpp | 14 +-
clang/lib/CodeGen/QualTypeMapper.cpp | 26 +--
clang/lib/CodeGen/QualTypeMapper.h | 3 -
.../CodeGen/AMDGPU/abi-classify-arg-types.c | 128 ++++++++++++++
.../AMDGPU/abi-classify-cxx-records.cpp | 47 +++++
.../AMDGPU/abi-classify-kernel-args.hip | 37 ++++
.../AMDGPU/abi-classify-return-types.c | 79 +++++++++
llvm/include/llvm/ABI/TargetInfo.h | 31 ++--
llvm/lib/ABI/TargetInfo.cpp | 64 +------
llvm/lib/ABI/Targets/AMDGPU.cpp | 166 +++++++++++++-----
llvm/lib/ABI/Targets/X86.cpp | 55 ++++++
11 files changed, 499 insertions(+), 151 deletions(-)
create mode 100644 clang/test/CodeGen/AMDGPU/abi-classify-arg-types.c
create mode 100644 clang/test/CodeGen/AMDGPU/abi-classify-cxx-records.cpp
create mode 100644 clang/test/CodeGen/AMDGPU/abi-classify-kernel-args.hip
create mode 100644 clang/test/CodeGen/AMDGPU/abi-classify-return-types.c
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index f364da24bd0d2..e1b5f3df2fb9a 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -401,21 +401,13 @@ CodeGenModule::getLLVMABITargetInfo(llvm::abi::TypeBuilder &TB) {
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);
+ T.isSPIRV() ? getLangOpts().isTargetDevice() : getLangOpts().HIP;
+ TheLLVMABITargetInfo =
+ llvm::abi::createAMDGPUTargetInfo(TB, getDataLayout(), Opts);
return *TheLLVMABITargetInfo;
}
diff --git a/clang/lib/CodeGen/QualTypeMapper.cpp b/clang/lib/CodeGen/QualTypeMapper.cpp
index 850ae406ab3a1..175880dd5a858 100644
--- a/clang/lib/CodeGen/QualTypeMapper.cpp
+++ b/clang/lib/CodeGen/QualTypeMapper.cpp
@@ -595,6 +595,8 @@ QualTypeMapper::convertUnionType(const clang::RecordDecl *RD) {
RecFlags |= llvm::abi::RecordFlags::CanPassInRegisters;
if (isa<CXXRecordDecl>(RD))
RecFlags |= llvm::abi::RecordFlags::IsCXXRecord;
+ if (RD->hasFlexibleArrayMember())
+ RecFlags |= llvm::abi::RecordFlags::HasFlexibleArrayMember;
return Builder.getUnionType(AllFields, Size, Alignment,
llvm::abi::StructPacking::Default, RecFlags);
@@ -605,26 +607,23 @@ 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));
+ LangAS AddrSpace = ASTCtx.getOpenCLTypeAddrSpace(T);
+ const clang::TargetInfo &TI = ASTCtx.getTargetInfo();
+ return Builder.getPointerType(TI.getPointerWidth(AddrSpace),
+ llvm::Align(TI.getPointerAlign(AddrSpace) / 8),
+ TI.getTargetAddressSpace(AddrSpace));
}
const llvm::abi::Type *
QualTypeMapper::createPointerTypeForPointee(QualType PointeeType) {
- LangAS AddrSpace = PointeeType.getAddressSpace();
+ auto AddrSpace = PointeeType.getAddressSpace();
+ auto PointerSize = ASTCtx.getTargetInfo().getPointerWidth(AddrSpace);
+ llvm::Align Alignment =
+ llvm::Align(ASTCtx.getTargetInfo().getPointerAlign(AddrSpace));
// 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.
@@ -632,7 +631,8 @@ QualTypeMapper::createPointerTypeForPointee(QualType PointeeType) {
PointeeType->isFunctionType() && !PointeeType.hasAddressSpace()
? DL.getProgramAddressSpace()
: ASTCtx.getTargetInfo().getTargetAddressSpace(AddrSpace);
- return createPointerType(AddrSpace, TargetAddrSpace);
+ return Builder.getPointerType(PointerSize, llvm::Align(Alignment.value() / 8),
+ 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 08c7214d85459..79480d12e98c4 100644
--- a/clang/lib/CodeGen/QualTypeMapper.h
+++ b/clang/lib/CodeGen/QualTypeMapper.h
@@ -51,9 +51,6 @@ 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);
diff --git a/clang/test/CodeGen/AMDGPU/abi-classify-arg-types.c b/clang/test/CodeGen/AMDGPU/abi-classify-arg-types.c
new file mode 100644
index 0000000000000..15be7ba3ec2aa
--- /dev/null
+++ b/clang/test/CodeGen/AMDGPU/abi-classify-arg-types.c
@@ -0,0 +1,128 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fexperimental-max-bitint-width=1024 -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,GCN
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fexperimental-max-bitint-width=1024 -fexperimental-abi-lowering -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,GCN
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fexperimental-max-bitint-width=1024 -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,SPIRV
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fexperimental-max-bitint-width=1024 -fexperimental-abi-lowering -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,SPIRV
+
+// Check that the ABI library classifies arguments the same way Clang does.
+
+typedef struct {} empty_t;
+typedef struct { char c; } i8_wrapper_t;
+typedef struct { int i; } i32_wrapper_t;
+typedef struct { float f; } f32_wrapper_t;
+typedef struct { int a[4]; } array_wrapper_t;
+typedef struct { char a, b, c; } small_t;
+typedef struct { int a, b, c, d; } pair16_t;
+typedef struct { int a[64]; } large_t;
+typedef struct { int i; float f; double d; } mixed_t;
+typedef union __attribute__((transparent_union)) { int i; } transparent_u;
+typedef __attribute__((ext_vector_type(3))) char char3;
+typedef __attribute__((ext_vector_type(4))) short short4;
+typedef struct { _Bool b; } bool_wrapper_t;
+typedef struct { _BitInt(24) i; } bitint24_wrapper_t;
+typedef struct { char3 v; } char3_wrapper_t;
+typedef union { long long a; char c[]; } fam_union_t;
+
+void arg_void(void) {}
+// CHECK: define{{.*}} void @arg_void()
+
+void arg_bool(_Bool b) {}
+// CHECK: define{{.*}} void @arg_bool(i1 noundef zeroext %{{.*}})
+
+void arg_char(char c) {}
+// CHECK: define{{.*}} void @arg_char(i8 noundef signext %{{.*}})
+
+void arg_short(short s) {}
+// CHECK: define{{.*}} void @arg_short(i16 noundef signext %{{.*}})
+
+void arg_int(int i) {}
+// CHECK: define{{.*}} void @arg_int(i32 noundef %{{.*}})
+
+void arg_long(long l) {}
+// CHECK: define{{.*}} void @arg_long(i64 noundef %{{.*}})
+
+void arg_float(float f) {}
+// CHECK: define{{.*}} void @arg_float(float noundef %{{.*}})
+
+void arg_double(double d) {}
+// CHECK: define{{.*}} void @arg_double(double noundef %{{.*}})
+
+void arg_ptr(int *p) {}
+// GCN: define{{.*}} void @arg_ptr(ptr noundef %{{.*}})
+// SPIRV: define{{.*}} void @arg_ptr(ptr addrspace(4) noundef %{{.*}})
+
+void arg_bitint65(_BitInt(65) b) {}
+// CHECK: define{{.*}} void @arg_bitint65(i65 noundef %{{.*}})
+
+void arg_bitint129(_BitInt(129) b) {}
+// GCN: define{{.*}} void @arg_bitint129(ptr addrspace(5) noundef byval(i192) align 8 %{{.*}})
+// SPIRV: define{{.*}} void @arg_bitint129(ptr noundef byval(i192) align 8 %{{.*}})
+
+void arg_empty(empty_t e) {}
+// CHECK: define{{.*}} void @arg_empty()
+
+void arg_i8_wrapper(i8_wrapper_t s) {}
+// CHECK: define{{.*}} void @arg_i8_wrapper(i8 %{{.*}})
+
+void arg_i32_wrapper(i32_wrapper_t s) {}
+// CHECK: define{{.*}} void @arg_i32_wrapper(i32 %{{.*}})
+
+void arg_f32_wrapper(f32_wrapper_t s) {}
+// CHECK: define{{.*}} void @arg_f32_wrapper(float %{{.*}})
+
+// Compared at in-memory size, so an element narrower than its storage unwraps.
+void arg_bool_wrapper(bool_wrapper_t s) {}
+// CHECK: define{{.*}} void @arg_bool_wrapper(i1 %{{.*}})
+
+void arg_bitint24_wrapper(bitint24_wrapper_t s) {}
+// CHECK: define{{.*}} void @arg_bitint24_wrapper(i24 %{{.*}})
+
+void arg_char3_wrapper(char3_wrapper_t s) {}
+// CHECK: define{{.*}} void @arg_char3_wrapper(<3 x i8> %{{.*}})
+
+void arg_fam_union(fam_union_t u) {}
+// GCN: define{{.*}} void @arg_fam_union(ptr addrspace(5) noundef byval(%union.fam_union_t) align 8 %{{.*}})
+// SPIRV: define{{.*}} void @arg_fam_union(ptr noundef byval(%union.fam_union_t) align 8 %{{.*}})
+
+// Not a single element, so this falls through to register packing.
+void arg_array_wrapper(array_wrapper_t s) {}
+// CHECK: define{{.*}} void @arg_array_wrapper([4 x i32] %{{.*}})
+
+void arg_small(small_t s) {}
+// CHECK: define{{.*}} void @arg_small(i32 %{{.*}})
+
+void arg_pair16(pair16_t s) {}
+// CHECK: define{{.*}} void @arg_pair16(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
+
+void arg_large(large_t s) {}
+// GCN: define{{.*}} void @arg_large(ptr addrspace(5) noundef byref(%struct.large_t) align 4 %{{.*}})
+// SPIRV: define{{.*}} void @arg_large(ptr noundef byref(%struct.large_t) align 4 %{{.*}})
+
+void arg_mixed(mixed_t s) {}
+// CHECK: define{{.*}} void @arg_mixed(i32 %{{.*}}, float %{{.*}}, double %{{.*}})
+
+void arg_transparent_union(transparent_u u) {}
+// CHECK: define{{.*}} void @arg_transparent_union(i32 %{{.*}})
+
+void arg_char3(char3 v) {}
+// CHECK: define{{.*}} void @arg_char3(<3 x i8> noundef %{{.*}})
+
+void arg_short4(short4 v) {}
+// CHECK: define{{.*}} void @arg_short4(<4 x i16> noundef %{{.*}})
+
+void arg_variadic(int n, ...) {}
+// CHECK: define{{.*}} void @arg_variadic(i32 noundef %{{.*}}, ...)
+
+// A variadic argument is passed as-is, not packed like a fixed one.
+void call_variadic(small_t s) { arg_variadic(1, s); }
+// CHECK: define{{.*}} void @call_variadic(i32 %{{.*}})
+// CHECK: call{{.*}} void (i32, ...){{.*}} @arg_variadic(i32 noundef 1, %struct.small_t %{{.*}})
+
+// The 16-register budget is shared, so the fifth aggregate goes by reference.
+void arg_reg_budget(pair16_t a, pair16_t b, pair16_t c, pair16_t d, pair16_t e) {}
+// GCN: define{{.*}} void @arg_reg_budget(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, ptr addrspace(5) noundef byref(%struct.pair16_t) align 4 %{{.*}})
+// SPIRV: define{{.*}} void @arg_reg_budget(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, ptr noundef byref(%struct.pair16_t) align 4 %{{.*}})
+
+// A 16-bit element vector packs two per register, so six short4 cost 12 of 16.
+void arg_packed_short4_budget(short4 a, short4 b, short4 c, short4 d, short4 e,
+ short4 f, pair16_t g) {}
+// CHECK: define{{.*}} void @arg_packed_short4_budget(<4 x i16> noundef %{{.*}}, <4 x i16> noundef %{{.*}}, <4 x i16> noundef %{{.*}}, <4 x i16> noundef %{{.*}}, <4 x i16> noundef %{{.*}}, <4 x i16> noundef %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
diff --git a/clang/test/CodeGen/AMDGPU/abi-classify-cxx-records.cpp b/clang/test/CodeGen/AMDGPU/abi-classify-cxx-records.cpp
new file mode 100644
index 0000000000000..cd045fe918207
--- /dev/null
+++ b/clang/test/CodeGen/AMDGPU/abi-classify-cxx-records.cpp
@@ -0,0 +1,47 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,GCN
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fexperimental-abi-lowering -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,GCN
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,SPIRV
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fexperimental-abi-lowering -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,SPIRV
+
+// Check that the ABI library classifies C++ records the same way Clang does.
+
+struct NonTrivialDtor { int i; ~NonTrivialDtor(); };
+struct NonCopyable { int i; NonCopyable(const NonCopyable &) = delete; NonCopyable(); };
+struct TrivialBase { int i; };
+struct DerivedSingle : TrivialBase {};
+struct EmptyBase {};
+struct DerivedWithEmptyBase : EmptyBase { float f; };
+struct MemberPtrHolder { int TrivialBase::*p; };
+
+// Constructed in place, so the sret pointer uses the generic address space.
+NonTrivialDtor ret_non_trivial_dtor() { return NonTrivialDtor(); }
+// GCN: define{{.*}} void @_Z20ret_non_trivial_dtorv(ptr dead_on_unwind noalias writable sret(%struct.NonTrivialDtor) align 4 %{{.*}})
+// SPIRV: define{{.*}} void @_Z20ret_non_trivial_dtorv(ptr addrspace(4) dead_on_unwind noalias writable sret(%struct.NonTrivialDtor) align 4 %{{.*}})
+
+NonCopyable ret_non_copyable() { return NonCopyable(); }
+// GCN: define{{.*}} void @_Z16ret_non_copyablev(ptr dead_on_unwind noalias writable sret(%struct.NonCopyable) align 4 %{{.*}})
+// SPIRV: define{{.*}} void @_Z16ret_non_copyablev(ptr addrspace(4) dead_on_unwind noalias writable sret(%struct.NonCopyable) align 4 %{{.*}})
+
+DerivedSingle ret_derived_single() { return DerivedSingle(); }
+// CHECK: define{{.*}} i32 @_Z18ret_derived_singlev()
+
+DerivedWithEmptyBase ret_derived_empty_base() { return DerivedWithEmptyBase(); }
+// CHECK: define{{.*}} float @_Z22ret_derived_empty_basev()
+
+// A data member pointer is a scalar, so the wrapper is a single-element struct.
+MemberPtrHolder ret_member_ptr() { return MemberPtrHolder(); }
+// CHECK: define{{.*}} i64 @_Z14ret_member_ptrv()
+
+// By contrast an argument passed by address uses the alloca address space.
+void arg_non_trivial_dtor(NonTrivialDtor s) {}
+// GCN: define{{.*}} void @_Z20arg_non_trivial_dtor14NonTrivialDtor(ptr addrspace(5) nofreeobj noundef align 4 dereferenceable(4) %{{.*}})
+// SPIRV: define{{.*}} void @_Z20arg_non_trivial_dtor14NonTrivialDtor(ptr nofreeobj noundef align 4 dereferenceable(4) %{{.*}})
+
+void arg_derived_single(DerivedSingle s) {}
+// CHECK: define{{.*}} void @_Z18arg_derived_single13DerivedSingle(i32 %{{.*}})
+
+void arg_derived_empty_base(DerivedWithEmptyBase s) {}
+// CHECK: define{{.*}} void @_Z22arg_derived_empty_base20DerivedWithEmptyBase(float %{{.*}})
+
+void arg_member_ptr(MemberPtrHolder s) {}
+// CHECK: define{{.*}} void @_Z14arg_member_ptr15MemberPtrHolder(i64 %{{.*}})
diff --git a/clang/test/CodeGen/AMDGPU/abi-classify-kernel-args.hip b/clang/test/CodeGen/AMDGPU/abi-classify-kernel-args.hip
new file mode 100644
index 0000000000000..c4ef23a7391fc
--- /dev/null
+++ b/clang/test/CodeGen/AMDGPU/abi-classify-kernel-args.hip
@@ -0,0 +1,37 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -emit-llvm -o - %s | FileCheck %s --check-prefix=GCN
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fexperimental-abi-lowering -emit-llvm -o - %s | FileCheck %s --check-prefix=GCN
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -emit-llvm -o - %s | FileCheck %s --check-prefix=SPIRV
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -fexperimental-abi-lowering -emit-llvm -o - %s | FileCheck %s --check-prefix=SPIRV
+
+// Kernel arguments are never byval, so aggregates go byref in the constant
+// address space, which differs between amdgcn and SPIR-V.
+
+#define __global__ __attribute__((global))
+
+struct pair_t { int a, b; };
+struct large_t { int a[64]; };
+struct i32_wrapper_t { int i; };
+
+__global__ void kern_scalar(int i, float f) {}
+// GCN: define{{.*}} amdgpu_kernel void @_Z11kern_scalarif(i32 noundef %{{.*}}, float noundef %{{.*}})
+// SPIRV: define{{.*}} spir_kernel void @_Z11kern_scalarif(i32 noundef %{{.*}}, float noundef %{{.*}})
+
+__global__ void kern_ptr(int *p) {}
+// GCN: define{{.*}} amdgpu_kernel void @_Z8kern_ptrPi(ptr addrspace(1) noundef %{{.*}})
+// SPIRV: define{{.*}} spir_kernel void @_Z8kern_ptrPi(ptr addrspace(1) noundef %{{.*}})
+
+__global__ void kern_atomic_ptr(_Atomic(int *) p) {}
+// GCN: define{{.*}} amdgpu_kernel void @_Z15kern_atomic_ptrU7_AtomicPi(ptr addrspace(1) %{{.*}})
+// SPIRV: define{{.*}} spir_kernel void @_Z15kern_atomic_ptrU7_AtomicPi(ptr addrspace(1) %{{.*}})
+
+__global__ void kern_small(pair_t s) {}
+// GCN: define{{.*}} amdgpu_kernel void @_Z10kern_small6pair_t(ptr addrspace(4) noundef byref(%struct.pair_t) align 4 %{{.*}})
+// SPIRV: define{{.*}} spir_kernel void @_Z10kern_small6pair_t(ptr addrspace(2) noundef byref(%struct.pair_t) align 4 %{{.*}})
+
+__global__ void kern_large(large_t s) {}
+// GCN: define{{.*}} amdgpu_kernel void @_Z10kern_large7large_t(ptr addrspace(4) noundef byref(%struct.large_t) align 4 %{{.*}})
+// SPIRV: define{{.*}} spir_kernel void @_Z10kern_large7large_t(ptr addrspace(2) noundef byref(%struct.large_t) align 4 %{{.*}})
+
+__global__ void kern_wrapper(i32_wrapper_t s) {}
+// GCN: define{{.*}} amdgpu_kernel void @_Z12kern_wrapper13i32_wrapper_t(i32 %{{.*}})
+// SPIRV: define{{.*}} spir_kernel void @_Z12kern_wrapper13i32_wrapper_t(i32 %{{.*}})
diff --git a/clang/test/CodeGen/AMDGPU/abi-classify-return-types.c b/clang/test/CodeGen/AMDGPU/abi-classify-return-types.c
new file mode 100644
index 0000000000000..d9b7d08be6aaf
--- /dev/null
+++ b/clang/test/CodeGen/AMDGPU/abi-classify-return-types.c
@@ -0,0 +1,79 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fexperimental-max-bitint-width=1024 -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,GCN
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fexperimental-max-bitint-width=1024 -fexperimental-abi-lowering -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,GCN
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fexperimental-max-bitint-width=1024 -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,SPIRV
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fexperimental-max-bitint-width=1024 -fexperimental-abi-lowering -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,SPIRV
+
+// Check that the ABI library classifies return types the same way Clang does.
+
+typedef struct {} empty_t;
+typedef struct { char c; } i8_wrapper_t;
+typedef struct { float f; } f32_wrapper_t;
+typedef struct { char a, b, c; } small_t;
+typedef struct { int a, b; } pair_t;
+typedef struct { int a, b, c, d; } quad_t;
+typedef struct { int a[4]; } array_wrapper_t;
+typedef struct { int a[64]; } large_t;
+typedef __attribute__((ext_vector_type(3))) char char3;
+typedef struct { _Bool b; } bool_wrapper_t;
+typedef struct { _BitInt(24) i; } bitint24_wrapper_t;
+
+void ret_void(void) {}
+// CHECK: define{{.*}} void @ret_void()
+
+_Bool ret_bool(void) { return 0; }
+// CHECK: define{{.*}} zeroext i1 @ret_bool()
+
+char ret_char(void) { return 0; }
+// CHECK: define{{.*}} signext i8 @ret_char()
+
+int ret_int(void) { return 0; }
+// CHECK: define{{.*}} i32 @ret_int()
+
+long ret_long(void) { return 0; }
+// CHECK: define{{.*}} i64 @ret_long()
+
+float ret_float(void) { return 0; }
+// CHECK: define{{.*}} float @ret_float()
+
+_BitInt(65) ret_bitint65(void) { return 0; }
+// CHECK: define{{.*}} i65 @ret_bitint65()
+
+_BitInt(129) ret_bitint129(void) { return 0; }
+// GCN: define{{.*}} void @ret_bitint129(ptr addrspace(5) dead_on_unwind noalias writable sret(i192) align 8 %{{.*}})
+// SPIRV: define{{.*}} void @ret_bitint129(ptr dead_on_unwind noalias writable sret(i192) align 8 %{{.*}})
+
+empty_t ret_empty(void) { empty_t e; return e; }
+// CHECK: define{{.*}} void @ret_empty()
+
+i8_wrapper_t ret_i8_wrapper(void) { i8_wrapper_t s; return s; }
+// CHECK: define{{.*}} i8 @ret_i8_wrapper()
+
+f32_wrapper_t ret_f32_wrapper(void) { f32_wrapper_t s; return s; }
+// CHECK: define{{.*}} float @ret_f32_wrapper()
+
+// Compared at in-memory size, so an element narrower than its storage unwraps.
+bool_wrapper_t ret_bool_wrapper(void) { bool_wrapper_t s; return s; }
+// CHECK: define{{.*}} i1 @ret_bool_wrapper()
+
+bitint24_wrapper_t ret_bitint24_wrapper(void) { bitint24_wrapper_t s; return s; }
+// CHECK: define{{.*}} i24 @ret_bitint24_wrapper()
+
+small_t ret_small(void) { small_t s; return s; }
+// CHECK: define{{.*}} i32 @ret_small()
+
+pair_t ret_pair(void) { pair_t s; return s; }
+// CHECK: define{{.*}} [2 x i32] @ret_pair()
+
+// Over 8 bytes but within the register budget, so still returned directly.
+quad_t ret_quad(void) { quad_t s; return s; }
+// CHECK: define{{.*}} %struct.quad_t @ret_quad()
+
+array_wrapper_t ret_array_wrapper(void) { array_wrapper_t s; return s; }
+// CHECK: define{{.*}} %struct.array_wrapper_t @ret_array_wrapper()
+
+large_t ret_large(void) { large_t s; return s; }
+// GCN: define{{.*}} void @ret_large(ptr addrspace(5) dead_on_unwind noalias writable sret(%struct.large_t) align 4 %{{.*}})
+// SPIRV: define{{.*}} void @ret_large(ptr dead_on_unwind noalias writable sret(%struct.large_t) align 4 %{{.*}})
+
+char3 ret_char3(void) { char3 v; return v; }
+// CHECK: define{{.*}} <3 x i8> @ret_char3()
diff --git a/llvm/include/llvm/ABI/TargetInfo.h b/llvm/include/llvm/ABI/TargetInfo.h
index 87e21b966b5bf..745594fcb1a7b 100644
--- a/llvm/include/llvm/ABI/TargetInfo.h
+++ b/llvm/include/llvm/ABI/TargetInfo.h
@@ -21,6 +21,9 @@
#include <memory>
namespace llvm {
+
+class DataLayout;
+
namespace abi {
enum RecordArgABI {
@@ -88,19 +91,15 @@ class TargetInfo {
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;
+ /// Address space of the sret pointer for \p RT, a record returned in memory.
+ virtual unsigned getSRetAddrSpace(const RecordType *RT) const { return 0; }
+
/// Apply rules for classifying return types that are common to all targets.
- /// AddrSpace is the address space of the sret pointer.
- LLVM_ABI bool maybeCommonClassifyReturnType(FunctionInfo &FI,
- unsigned AddrSpace = 0) const;
+ LLVM_ABI bool maybeCommonClassifyReturnType(FunctionInfo &FI) const;
/// Return true if \p Ty is a valid base type for a homogeneous aggregate.
virtual bool isHomogeneousAggregateBaseType(const Type *Ty) const {
@@ -168,28 +167,18 @@ struct AArch64ABIOptions {
LLVM_ABI std::unique_ptr<TargetInfo>
createAArch64TargetInfo(TypeBuilder &TB, const AArch64ABIOptions &Opts);
-/// 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.
+ /// The language default address space.
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);
+createAMDGPUTargetInfo(TypeBuilder &TB, const DataLayout &DL,
+ const AMDGPUABIOptions &Opts);
} // namespace abi
} // namespace llvm
diff --git a/llvm/lib/ABI/TargetInfo.cpp b/llvm/lib/ABI/TargetInfo.cpp
index 9225bb0b6a538..6a2ae8747e069 100644
--- a/llvm/lib/ABI/TargetInfo.cpp
+++ b/llvm/lib/ABI/TargetInfo.cpp
@@ -50,63 +50,6 @@ ArgInfo TargetInfo::getNaturalAlignIndirect(const Type *Ty, bool ByVal,
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 {
if (RT && !RT->canPassInRegisters())
return RAA_Indirect;
@@ -133,8 +76,7 @@ const Type *TargetInfo::useFirstFieldIfTransparentUnion(const Type *Ty) const {
return Ty;
}
-bool TargetInfo::maybeCommonClassifyReturnType(FunctionInfo &FI,
- unsigned AddrSpace) const {
+bool TargetInfo::maybeCommonClassifyReturnType(FunctionInfo &FI) const {
const abi::Type *Ty = FI.getReturnType();
// TODO: When Microsoft ABI is supported, CXX records may need different
@@ -145,8 +87,8 @@ bool TargetInfo::maybeCommonClassifyReturnType(FunctionInfo &FI,
// is returned indirectly with ByVal=false. This is the RAA path and is
// distinct from getIndirectReturnResult (plain aggregates), which uses
// ByVal=true.
- FI.getReturnInfo() =
- ArgInfo::getIndirect(RT->getAlignment(), /*ByVal=*/false, AddrSpace);
+ FI.getReturnInfo() = ArgInfo::getIndirect(
+ RT->getAlignment(), /*ByVal=*/false, getSRetAddrSpace(RT));
return true;
}
}
diff --git a/llvm/lib/ABI/Targets/AMDGPU.cpp b/llvm/lib/ABI/Targets/AMDGPU.cpp
index 166975e170ec3..d5bcf68a089c8 100644
--- a/llvm/lib/ABI/Targets/AMDGPU.cpp
+++ b/llvm/lib/ABI/Targets/AMDGPU.cpp
@@ -6,8 +6,7 @@
//
//===----------------------------------------------------------------------===//
//
-// Shared by amdgcn and the AMDGCN-flavoured SPIR-V that lowers to it. The two
-// differ only in the values carried by AMDGPUABIOptions.
+// Shared by amdgcn and the AMDGCN-flavoured SPIR-V that lowers to it.
//
//===----------------------------------------------------------------------===//
@@ -15,21 +14,33 @@
#include "llvm/ABI/TargetInfo.h"
#include "llvm/ABI/Types.h"
#include "llvm/ADT/STLExtras.h"
+#include "llvm/IR/DataLayout.h"
#include "llvm/Support/Casting.h"
#include "llvm/Support/MathExtras.h"
#include <algorithm>
namespace llvm::abi {
+static const PointerType *getGlobalsPointerType(TypeBuilder &TB,
+ const DataLayout &DL) {
+ const unsigned AS = DL.getDefaultGlobalsAddressSpace();
+ return TB.getPointerType(DL.getPointerSizeInBits(AS),
+ DL.getPointerABIAlignment(AS), AS);
+}
+
class AMDGPUTargetInfo : public TargetInfo {
private:
/// Registers available for arguments and return values, in 32-bit units.
static constexpr unsigned MaxNumRegsForArgsRet = 16;
+ static constexpr uint64_t MaxDirectBitIntWidth = 128;
+
+ const DataLayout &DL;
AMDGPUABIOptions Opts;
const IntegerType *Int16Ty;
const IntegerType *Int32Ty;
const ArrayType *Int32PairTy;
+ const PointerType *KernelArgPtrTy;
uint64_t numRegsForType(const Type *Ty) const {
if (const auto *VT = dyn_cast<VectorType>(Ty)) {
@@ -41,7 +52,7 @@ class AMDGPUTargetInfo : public TargetInfo {
// 16-bit element vectors should be passed as packed.
if (EltSize == 16)
- return (NumElts + 1) / 2;
+ return divideCeil(NumElts, 2);
return divideCeil(EltSize, 32) * NumElts;
}
@@ -61,23 +72,100 @@ class AMDGPUTargetInfo : public TargetInfo {
return divideCeil(Ty->getTypeAllocSizeInBits().getFixedValue(), 32);
}
+ /// 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.
+ // TODO: X86 has its own copy of this. Hoist both into TargetInfo.
+ const Type *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;
+ }
+
/// 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 (!Opts.CoerceKernelPointerArgs)
+ return Ty;
+
+ // Classic CodeGen coerces the lowered type, which has no atomic wrapper.
+ const Type *Unwrapped = Ty;
+ if (const auto *AT = dyn_cast<AtomicType>(Ty))
+ Unwrapped = AT->getValueType();
+
+ const auto *PtrTy = dyn_cast<PointerType>(Unwrapped);
if (PtrTy && PtrTy->getAddrSpace() == Opts.GenericAddrSpace)
- return TB.getPointerType(PtrTy->getSizeInBits().getFixedValue(),
- PtrTy->getAlignment(), Opts.KernelArgAddrSpace);
+ return KernelArgPtrTy;
return Ty;
}
+ /// Pack an aggregate of \p Size bits into a VGPR or a VGPR pair.
+ const Type *getRegisterCoerceType(uint64_t Size) const {
+ if (Size <= 16)
+ return Int16Ty;
+ if (Size <= 32)
+ return Int32Ty;
+ return Int32PairTy;
+ }
+
/// 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())
+ IntTy->getSizeInBits().getFixedValue() > MaxDirectBitIntWidth)
return getNaturalAlignIndirect(Ty, /*ByVal=*/true,
- Opts.AllocaAddrSpace);
+ DL.getAllocaAddrSpace());
if (isPromotableInteger(IntTy))
return ArgInfo::getExtend(Ty);
@@ -93,7 +181,7 @@ class AMDGPUTargetInfo : public TargetInfo {
if (isAggregateTypeForABI(Ty))
return getNaturalAlignIndirect(Ty, getRecordArgABI(Ty) != RAA_Indirect,
- Opts.AllocaAddrSpace);
+ DL.getAllocaAddrSpace());
return defaultClassifyScalar(Ty);
}
@@ -104,7 +192,7 @@ class AMDGPUTargetInfo : public TargetInfo {
if (isAggregateTypeForABI(RetTy))
return getNaturalAlignIndirect(RetTy, /*ByVal=*/true,
- Opts.AllocaAddrSpace);
+ DL.getAllocaAddrSpace());
return defaultClassifyScalar(RetTy);
}
@@ -128,14 +216,8 @@ class AMDGPUTargetInfo : public TargetInfo {
// 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);
+ return ArgInfo::getDirect(getRegisterCoerceType(Size));
if (numRegsForType(RetTy) <= MaxNumRegsForArgsRet)
return ArgInfo::getDirect();
@@ -169,7 +251,7 @@ class AMDGPUTargetInfo : public TargetInfo {
// passed by value.
if (RecordArgABI RAA = getRecordArgABI(RT))
return getNaturalAlignIndirect(Ty, RAA == RAA_DirectInMemory,
- Opts.AllocaAddrSpace);
+ DL.getAllocaAddrSpace());
// Ignore empty structs/unions.
if (RT && RT->isEmpty())
@@ -189,14 +271,9 @@ class AMDGPUTargetInfo : public TargetInfo {
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);
+ // XXX: Should the 64-bit case be i64 instead, and should the limit
+ // increase?
+ return ArgInfo::getDirect(getRegisterCoerceType(Size));
}
if (NumRegsLeft > 0) {
@@ -210,7 +287,7 @@ class AMDGPUTargetInfo : public TargetInfo {
// Use pass-by-reference instead of pass-by-value for struct arguments in
// function ABI.
return ArgInfo::getIndirectAliased(Ty->getAlignment(),
- Opts.PrivateAddrSpace);
+ DL.getAllocaAddrSpace());
}
/// For kernels all parameters are really passed in a special buffer. It
@@ -231,33 +308,37 @@ class AMDGPUTargetInfo : public TargetInfo {
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,
+ return ArgInfo::getDirect(coerceKernelArgumentType(Ty), /*Offset=*/0,
+ /*Align=*/std::nullopt,
/*CanBeFlattened=*/false);
}
- uint64_t getMaxDirectBitIntWidth() const { return Opts.HasInt128 ? 128 : 64; }
+protected:
+ /// A record that cannot be copied is constructed in place, so the sret
+ /// pointer uses the generic address space rather than the alloca one.
+ unsigned getSRetAddrSpace(const RecordType *RT) const override {
+ return Opts.GenericAddrSpace;
+ }
public:
- AMDGPUTargetInfo(TypeBuilder &TB, const AMDGPUABIOptions &Opts)
- : TargetInfo(TB), Opts(Opts),
+ AMDGPUTargetInfo(TypeBuilder &TB, const DataLayout &DL,
+ const AMDGPUABIOptions &Opts)
+ : TargetInfo(TB), DL(DL), Opts(Opts),
Int16Ty(TB.getIntegerType(16, Align(2), /*Signed=*/false)),
Int32Ty(TB.getIntegerType(32, Align(4), /*Signed=*/false)),
- Int32PairTy(TB.getArrayType(Int32Ty, 2, 64)) {}
+ Int32PairTy(TB.getArrayType(Int32Ty, 2, 64)),
+ KernelArgPtrTy(getGlobalsPointerType(TB, DL)) {}
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))
+ if (!maybeCommonClassifyReturnType(FI))
FI.getReturnInfo() = classifyReturnType(FI.getReturnType());
- const bool IsKernel = FI.getCallingConvention() == Opts.KernelCC;
+ const CallingConv::ID CC = FI.getCallingConvention();
+ const bool IsKernel =
+ CC == CallingConv::AMDGPU_KERNEL || CC == CallingConv::SPIR_KERNEL;
const unsigned NumRequiredArgs = FI.getNumRequiredArgs();
unsigned NumRegsLeft = MaxNumRegsForArgsRet;
@@ -271,8 +352,9 @@ class AMDGPUTargetInfo : public TargetInfo {
};
std::unique_ptr<TargetInfo>
-createAMDGPUTargetInfo(TypeBuilder &TB, const AMDGPUABIOptions &Opts) {
- return std::make_unique<AMDGPUTargetInfo>(TB, Opts);
+createAMDGPUTargetInfo(TypeBuilder &TB, const DataLayout &DL,
+ const AMDGPUABIOptions &Opts) {
+ return std::make_unique<AMDGPUTargetInfo>(TB, DL, Opts);
}
} // namespace llvm::abi
diff --git a/llvm/lib/ABI/Targets/X86.cpp b/llvm/lib/ABI/Targets/X86.cpp
index 515be6158d106..08a5e20cb9070 100644
--- a/llvm/lib/ABI/Targets/X86.cpp
+++ b/llvm/lib/ABI/Targets/X86.cpp
@@ -97,6 +97,7 @@ 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;
@@ -1252,6 +1253,60 @@ 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();
More information about the cfe-commits
mailing list