[clang] [clang][SPIR-V][AMDGPU] Fix variadic aggregate ABI classification for AMDGCN (PR #216326)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Wed Aug 26 23:10:41 PDT 2026
https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/216326
>From 0d90147fbc852cbbb6df7c6b49ed65429ec5bffa Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Fri, 14 Aug 2026 16:26:39 +0200
Subject: [PATCH 1/2] [clang][SPIR-V] Fix variadic aggregate ABI classification
for AMDGCN
Mirror AMDGPUABIInfo fixed/variadic split
AMDGCNSPIRVABIInfo::classifyArgumentType did not distinguish fixed from variadic arguments, so aggregates passed through `...` were misclassified as indirect byref instead of direct
---
clang/lib/CodeGen/Targets/SPIR.cpp | 24 +++-
.../amdgcnspirv-uses-amdgpu-abi.cpp | 115 ++++++++++++++++++
2 files changed, 133 insertions(+), 6 deletions(-)
diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp
index 177d166d5b85b..f11aaf201a7ef 100644
--- a/clang/lib/CodeGen/Targets/SPIR.cpp
+++ b/clang/lib/CodeGen/Targets/SPIR.cpp
@@ -73,7 +73,7 @@ class AMDGCNSPIRVABIInfo : public SPIRVABIInfo {
ABIArgInfo classifyReturnType(QualType RetTy) const;
ABIArgInfo classifyKernelArgumentType(QualType Ty) const;
- ABIArgInfo classifyArgumentType(QualType Ty) const;
+ ABIArgInfo classifyArgumentType(QualType Ty, bool Variadic) const;
public:
AMDGCNSPIRVABIInfo(CodeGenTypes &CGT) : SPIRVABIInfo(CGT) {}
@@ -320,12 +320,19 @@ ABIArgInfo AMDGCNSPIRVABIInfo::classifyKernelArgumentType(QualType Ty) const {
return ABIArgInfo::getDirect(LTy, 0, nullptr, false);
}
-ABIArgInfo AMDGCNSPIRVABIInfo::classifyArgumentType(QualType Ty) const {
+ABIArgInfo AMDGCNSPIRVABIInfo::classifyArgumentType(QualType Ty,
+ bool Variadic) const {
assert(NumRegsLeft <= MaxNumRegsForArgsRet && "register estimate underflow");
Ty = useFirstFieldIfTransparentUnion(Ty);
- // TODO: support for variadics.
+ if (Variadic) {
+ return ABIArgInfo::getDirect(/*T=*/nullptr,
+ /*Offset=*/0,
+ /*Padding=*/nullptr,
+ /*CanBeFlattened=*/false,
+ /*Align=*/0);
+ }
if (!isAggregateTypeForABI(Ty)) {
ABIArgInfo ArgInfo = DefaultABIInfo::classifyArgumentType(Ty);
@@ -396,12 +403,17 @@ void AMDGCNSPIRVABIInfo::computeInfo(CGFunctionInfo &FI) const {
if (!getCXXABI().classifyReturnType(FI))
FI.getReturnInfo() = classifyReturnType(FI.getReturnType());
+ unsigned ArgumentIndex = 0;
+ const unsigned NumRequiredArgs = FI.getNumRequiredArgs();
+
NumRegsLeft = MaxNumRegsForArgsRet;
for (auto &I : FI.arguments()) {
- if (CC == llvm::CallingConv::SPIR_KERNEL)
+ if (CC == llvm::CallingConv::SPIR_KERNEL) {
I.info = classifyKernelArgumentType(I.type);
- else
- I.info = classifyArgumentType(I.type);
+ } else {
+ bool FixedArgument = ArgumentIndex++ < NumRequiredArgs;
+ I.info = classifyArgumentType(I.type, !FixedArgument);
+ }
}
}
diff --git a/clang/test/CodeGenHIP/amdgcnspirv-uses-amdgpu-abi.cpp b/clang/test/CodeGenHIP/amdgcnspirv-uses-amdgpu-abi.cpp
index 0a146ce6485d6..9741c58762ac6 100644
--- a/clang/test/CodeGenHIP/amdgcnspirv-uses-amdgpu-abi.cpp
+++ b/clang/test/CodeGenHIP/amdgcnspirv-uses-amdgpu-abi.cpp
@@ -316,6 +316,121 @@ __device__ V3 f15() { return {}; }
// AMDGPU-NEXT: ret <4 x i32> zeroinitializer
//
__device__ V4 f16() { return {}; }
+
+extern "C" __device__ void variadic(int, ...);
+// AMDGCNSPIRV-LABEL: define spir_func void @_Z3f175ByRef(
+// AMDGCNSPIRV-SAME: ptr nofree noundef readonly byref([[STRUCT_BYREF:%.*]]) align 4 captures(none) [[TMP0:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3:[0-9]+]] {
+// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
+// AMDGCNSPIRV-NEXT: [[B_SROA_0_0_COPYLOAD:%.*]] = load i32, ptr [[TMP0]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_2_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_2_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_2_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_3_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 8
+// AMDGCNSPIRV-NEXT: [[B_SROA_3_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_3_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_4_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 12
+// AMDGCNSPIRV-NEXT: [[B_SROA_4_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_4_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_5_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 16
+// AMDGCNSPIRV-NEXT: [[B_SROA_5_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_5_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_6_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 20
+// AMDGCNSPIRV-NEXT: [[B_SROA_6_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_6_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_7_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 24
+// AMDGCNSPIRV-NEXT: [[B_SROA_7_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_7_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_8_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 28
+// AMDGCNSPIRV-NEXT: [[B_SROA_8_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_8_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_9_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 32
+// AMDGCNSPIRV-NEXT: [[B_SROA_9_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_9_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_10_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 36
+// AMDGCNSPIRV-NEXT: [[B_SROA_10_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_10_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_11_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 40
+// AMDGCNSPIRV-NEXT: [[B_SROA_11_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_11_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_12_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 44
+// AMDGCNSPIRV-NEXT: [[B_SROA_12_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_12_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_13_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 48
+// AMDGCNSPIRV-NEXT: [[B_SROA_13_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_13_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_14_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 52
+// AMDGCNSPIRV-NEXT: [[B_SROA_14_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_14_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_15_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 56
+// AMDGCNSPIRV-NEXT: [[B_SROA_15_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_15_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_16_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 60
+// AMDGCNSPIRV-NEXT: [[B_SROA_16_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_16_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[B_SROA_17_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP0]], i64 64
+// AMDGCNSPIRV-NEXT: [[B_SROA_17_0_COPYLOAD:%.*]] = load i32, ptr [[B_SROA_17_0__SROA_IDX]], align 4
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_0_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] poison, i32 [[B_SROA_0_0_COPYLOAD]], 0, 0
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_1_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_0_INSERT]], i32 [[B_SROA_2_0_COPYLOAD]], 0, 1
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_2_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_1_INSERT]], i32 [[B_SROA_3_0_COPYLOAD]], 0, 2
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_3_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_2_INSERT]], i32 [[B_SROA_4_0_COPYLOAD]], 0, 3
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_4_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_3_INSERT]], i32 [[B_SROA_5_0_COPYLOAD]], 0, 4
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_5_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_4_INSERT]], i32 [[B_SROA_6_0_COPYLOAD]], 0, 5
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_6_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_5_INSERT]], i32 [[B_SROA_7_0_COPYLOAD]], 0, 6
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_7_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_6_INSERT]], i32 [[B_SROA_8_0_COPYLOAD]], 0, 7
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_8_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_7_INSERT]], i32 [[B_SROA_9_0_COPYLOAD]], 0, 8
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_9_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_8_INSERT]], i32 [[B_SROA_10_0_COPYLOAD]], 0, 9
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_10_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_9_INSERT]], i32 [[B_SROA_11_0_COPYLOAD]], 0, 10
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_11_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_10_INSERT]], i32 [[B_SROA_12_0_COPYLOAD]], 0, 11
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_12_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_11_INSERT]], i32 [[B_SROA_13_0_COPYLOAD]], 0, 12
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_13_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_12_INSERT]], i32 [[B_SROA_14_0_COPYLOAD]], 0, 13
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_14_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_13_INSERT]], i32 [[B_SROA_15_0_COPYLOAD]], 0, 14
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_15_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_14_INSERT]], i32 [[B_SROA_16_0_COPYLOAD]], 0, 15
+// AMDGCNSPIRV-NEXT: [[DOTFCA_0_16_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_15_INSERT]], i32 [[B_SROA_17_0_COPYLOAD]], 0, 16
+// AMDGCNSPIRV-NEXT: tail call spir_func addrspace(4) void (i32, ...) @variadic(i32 noundef 1, [[STRUCT_BYREF]] [[DOTFCA_0_16_INSERT]]) #[[ATTR5:[0-9]+]]
+// AMDGCNSPIRV-NEXT: ret void
+//
+// AMDGPU-LABEL: define dso_local void @_Z3f175ByRef(
+// AMDGPU-SAME: ptr addrspace(5) nofree noundef readonly byref([[STRUCT_BYREF:%.*]]) align 4 captures(none) [[TMP0:%.*]]) local_unnamed_addr #[[ATTR4:[0-9]+]] {
+// AMDGPU-NEXT: [[ENTRY:.*:]]
+// AMDGPU-NEXT: [[B_SROA_0_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[TMP0]], align 4
+// AMDGPU-NEXT: [[B_SROA_2_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 4
+// AMDGPU-NEXT: [[B_SROA_2_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_2_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_3_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 8
+// AMDGPU-NEXT: [[B_SROA_3_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_3_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_4_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 12
+// AMDGPU-NEXT: [[B_SROA_4_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_4_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_5_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 16
+// AMDGPU-NEXT: [[B_SROA_5_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_5_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_6_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 20
+// AMDGPU-NEXT: [[B_SROA_6_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_6_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_7_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 24
+// AMDGPU-NEXT: [[B_SROA_7_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_7_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_8_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 28
+// AMDGPU-NEXT: [[B_SROA_8_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_8_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_9_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 32
+// AMDGPU-NEXT: [[B_SROA_9_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_9_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_10_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 36
+// AMDGPU-NEXT: [[B_SROA_10_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_10_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_11_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 40
+// AMDGPU-NEXT: [[B_SROA_11_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_11_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_12_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 44
+// AMDGPU-NEXT: [[B_SROA_12_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_12_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_13_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 48
+// AMDGPU-NEXT: [[B_SROA_13_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_13_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_14_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 52
+// AMDGPU-NEXT: [[B_SROA_14_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_14_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_15_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 56
+// AMDGPU-NEXT: [[B_SROA_15_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_15_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_16_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 60
+// AMDGPU-NEXT: [[B_SROA_16_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_16_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[B_SROA_17_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(5) [[TMP0]], i32 64
+// AMDGPU-NEXT: [[B_SROA_17_0_COPYLOAD:%.*]] = load i32, ptr addrspace(5) [[B_SROA_17_0__SROA_IDX]], align 4
+// AMDGPU-NEXT: [[DOTFCA_0_0_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] poison, i32 [[B_SROA_0_0_COPYLOAD]], 0, 0
+// AMDGPU-NEXT: [[DOTFCA_0_1_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_0_INSERT]], i32 [[B_SROA_2_0_COPYLOAD]], 0, 1
+// AMDGPU-NEXT: [[DOTFCA_0_2_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_1_INSERT]], i32 [[B_SROA_3_0_COPYLOAD]], 0, 2
+// AMDGPU-NEXT: [[DOTFCA_0_3_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_2_INSERT]], i32 [[B_SROA_4_0_COPYLOAD]], 0, 3
+// AMDGPU-NEXT: [[DOTFCA_0_4_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_3_INSERT]], i32 [[B_SROA_5_0_COPYLOAD]], 0, 4
+// AMDGPU-NEXT: [[DOTFCA_0_5_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_4_INSERT]], i32 [[B_SROA_6_0_COPYLOAD]], 0, 5
+// AMDGPU-NEXT: [[DOTFCA_0_6_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_5_INSERT]], i32 [[B_SROA_7_0_COPYLOAD]], 0, 6
+// AMDGPU-NEXT: [[DOTFCA_0_7_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_6_INSERT]], i32 [[B_SROA_8_0_COPYLOAD]], 0, 7
+// AMDGPU-NEXT: [[DOTFCA_0_8_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_7_INSERT]], i32 [[B_SROA_9_0_COPYLOAD]], 0, 8
+// AMDGPU-NEXT: [[DOTFCA_0_9_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_8_INSERT]], i32 [[B_SROA_10_0_COPYLOAD]], 0, 9
+// AMDGPU-NEXT: [[DOTFCA_0_10_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_9_INSERT]], i32 [[B_SROA_11_0_COPYLOAD]], 0, 10
+// AMDGPU-NEXT: [[DOTFCA_0_11_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_10_INSERT]], i32 [[B_SROA_12_0_COPYLOAD]], 0, 11
+// AMDGPU-NEXT: [[DOTFCA_0_12_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_11_INSERT]], i32 [[B_SROA_13_0_COPYLOAD]], 0, 12
+// AMDGPU-NEXT: [[DOTFCA_0_13_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_12_INSERT]], i32 [[B_SROA_14_0_COPYLOAD]], 0, 13
+// AMDGPU-NEXT: [[DOTFCA_0_14_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_13_INSERT]], i32 [[B_SROA_15_0_COPYLOAD]], 0, 14
+// AMDGPU-NEXT: [[DOTFCA_0_15_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_14_INSERT]], i32 [[B_SROA_16_0_COPYLOAD]], 0, 15
+// AMDGPU-NEXT: [[DOTFCA_0_16_INSERT:%.*]] = insertvalue [[STRUCT_BYREF]] [[DOTFCA_0_15_INSERT]], i32 [[B_SROA_17_0_COPYLOAD]], 0, 16
+// AMDGPU-NEXT: tail call void (i32, ...) @variadic(i32 noundef 1, [[STRUCT_BYREF]] [[DOTFCA_0_16_INSERT]]) #[[ATTR6:[0-9]+]]
+// AMDGPU-NEXT: ret void
+//
+__device__ void f17(ByRef b) { variadic(1, b); }
//.
// AMDGCNSPIRV: [[META8]] = !{i32 1024, i32 1, i32 1}
//.
>From 368a3a21ecc6d4c2e8baf544703dc2d2d7e26118 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 2/2] move to ABI lib
---
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 f11aaf201a7ef..d721d8761ce1e 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) {
More information about the cfe-commits
mailing list