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 <[email protected]> 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 <[email protected]> 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) { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
