From d2d2d3d5db3f639aab178f9ca9a20db2842d2b65 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 29 Oct 2024 14:20:44 +0000 Subject: [PATCH 01/27] `sret` args should always point to the `alloca` AS, so we can use that. --- clang/lib/CodeGen/CGCall.cpp | 15 ++++++++------- clang/test/CodeGen/partial-reinitialization2.c | 4 ++-- clang/test/CodeGen/sret.c | 11 +++++++++++ 3 files changed, 21 insertions(+), 9 deletions(-) diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 8f4f5d3ed8160..56acfae7ae9e5 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -1672,8 +1672,7 @@ CodeGenTypes::GetFunctionType(const CGFunctionInfo &FI) { // Add type for sret argument. if (IRFunctionArgs.hasSRetArg()) { - QualType Ret = FI.getReturnType(); - unsigned AddressSpace = CGM.getTypes().getTargetAddressSpace(Ret); + unsigned AddressSpace = CGM.getDataLayout().getAllocaAddrSpace(); ArgTypes[IRFunctionArgs.getSRetArgNo()] = llvm::PointerType::get(getLLVMContext(), AddressSpace); } @@ -5145,7 +5144,6 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo, // If the call returns a temporary with struct return, create a temporary // alloca to hold the result, unless one is given to us. Address SRetPtr = Address::invalid(); - RawAddress SRetAlloca = RawAddress::invalid(); llvm::Value *UnusedReturnSizePtr = nullptr; if (RetAI.isIndirect() || RetAI.isInAlloca() || RetAI.isCoerceAndExpand()) { // For virtual function pointer thunks and musttail calls, we must always @@ -5159,16 +5157,19 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo, } else if (!ReturnValue.isNull()) { SRetPtr = ReturnValue.getAddress(); } else { - SRetPtr = CreateMemTemp(RetTy, "tmp", &SRetAlloca); + SRetPtr = CreateMemTempWithoutCast(RetTy, "tmp"); if (HaveInsertPoint() && ReturnValue.isUnused()) { llvm::TypeSize size = CGM.getDataLayout().getTypeAllocSize(ConvertTypeForMem(RetTy)); - UnusedReturnSizePtr = EmitLifetimeStart(size, SRetAlloca.getPointer()); + UnusedReturnSizePtr = EmitLifetimeStart(size, SRetPtr.getBasePointer()); } } if (IRFunctionArgs.hasSRetArg()) { + // If the caller allocated the return slot, it is possible that the + // alloca was AS casted to the default as, so we ensure the cast is + // stripped before binding to the sret arg, which is in the allocaAS. IRCallArgs[IRFunctionArgs.getSRetArgNo()] = - getAsNaturalPointerTo(SRetPtr, RetTy); + getAsNaturalPointerTo(SRetPtr, RetTy)->stripPointerCasts(); } else if (RetAI.isInAlloca()) { Address Addr = Builder.CreateStructGEP(ArgMemory, RetAI.getInAllocaFieldIndex()); @@ -5740,7 +5741,7 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo, // pop this cleanup later on. Being eager about this is OK, since this // temporary is 'invisible' outside of the callee. if (UnusedReturnSizePtr) - pushFullExprCleanup(NormalEHLifetimeMarker, SRetAlloca, + pushFullExprCleanup(NormalEHLifetimeMarker, SRetPtr, UnusedReturnSizePtr); llvm::BasicBlock *InvokeDest = CannotThrow ? nullptr : getInvokeDest(); diff --git a/clang/test/CodeGen/partial-reinitialization2.c b/clang/test/CodeGen/partial-reinitialization2.c index e709c1d4ad1ee..7949a69555031 100644 --- a/clang/test/CodeGen/partial-reinitialization2.c +++ b/clang/test/CodeGen/partial-reinitialization2.c @@ -91,8 +91,8 @@ void test5(void) // CHECK-LABEL: test6 void test6(void) { - // CHECK: [[LP:%[a-z0-9]+]] = getelementptr{{.*}}%struct.LLP2P2, ptr{{.*}}, i32 0, i32 0 - // CHECK: call {{.*}}get456789(ptr {{.*}}[[LP]]) + // CHECK: [[VAR:%[a-z0-9]+]] = alloca + // CHECK: call {{.*}}get456789(ptr {{.*}}sret{{.*}} [[VAR]]) // CHECK: [[CALL:%[a-z0-9]+]] = call {{.*}}@get235() // CHECK: store{{.*}}[[CALL]], {{.*}}[[TMP0:%[a-z0-9.]+]] diff --git a/clang/test/CodeGen/sret.c b/clang/test/CodeGen/sret.c index 6d905e89b2c6f..3b4914f29d2bf 100644 --- a/clang/test/CodeGen/sret.c +++ b/clang/test/CodeGen/sret.c @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -Wno-strict-prototypes -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -Wno-strict-prototypes -triple amdgcn-amd-amdhsa -emit-llvm -o - | FileCheck --check-prefix=NONZEROALLOCAAS %s struct abc { long a; @@ -6,18 +7,28 @@ struct abc { long c; long d; long e; + long f; + long g; + long h; + long i; + long j; }; struct abc foo1(void); // CHECK-DAG: declare {{.*}} @foo1(ptr dead_on_unwind writable sret(%struct.abc) +// NONZEROALLOCAAS-DAG: declare {{.*}} @foo1(ptr addrspace(5) dead_on_unwind writable sret(%struct.abc) struct abc foo2(); // CHECK-DAG: declare {{.*}} @foo2(ptr dead_on_unwind writable sret(%struct.abc) +// NONZEROALLOCAAS-DAG: declare {{.*}} @foo2(ptr addrspace(5) dead_on_unwind writable sret(%struct.abc) struct abc foo3(void){} // CHECK-DAG: define {{.*}} @foo3(ptr dead_on_unwind noalias writable sret(%struct.abc) +// NONZEROALLOCAAS-DAG: define {{.*}} @foo3(ptr addrspace(5) dead_on_unwind noalias writable sret(%struct.abc) void bar(void) { struct abc dummy1 = foo1(); // CHECK-DAG: call {{.*}} @foo1(ptr dead_on_unwind writable sret(%struct.abc) + // NONZEROALLOCAAS-DAG: call {{.*}} @foo1(ptr addrspace(5) dead_on_unwind writable sret(%struct.abc) struct abc dummy2 = foo2(); // CHECK-DAG: call {{.*}} @foo2(ptr dead_on_unwind writable sret(%struct.abc) + // NONZEROALLOCAAS-DAG: call {{.*}} @foo2(ptr addrspace(5) dead_on_unwind writable sret(%struct.abc) } From b5a7df0a771cb70d60e58a8727a5d856219dacb3 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 29 Oct 2024 17:16:17 +0000 Subject: [PATCH 02/27] Fix broken tests. --- clang/test/CodeGenOpenCL/addr-space-struct-arg.cl | 4 ++-- clang/test/CodeGenOpenCL/amdgpu-abi-struct-arg-byref.cl | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/test/CodeGenOpenCL/addr-space-struct-arg.cl b/clang/test/CodeGenOpenCL/addr-space-struct-arg.cl index 57d056b0ff9d5..4a1db2c3564a5 100644 --- a/clang/test/CodeGenOpenCL/addr-space-struct-arg.cl +++ b/clang/test/CodeGenOpenCL/addr-space-struct-arg.cl @@ -250,7 +250,7 @@ kernel void ker(global Mat3X3 *in, global Mat4X4 *out) { // AMDGCN-NEXT: ret void // // AMDGCN20-LABEL: define dso_local void @foo_large( -// AMDGCN20-SAME: ptr dead_on_unwind noalias writable sret([[STRUCT_MAT64X64:%.*]]) align 4 [[AGG_RESULT:%.*]], ptr addrspace(5) noundef byref([[STRUCT_MAT32X32:%.*]]) align 4 [[TMP0:%.*]]) #[[ATTR0]] { +// AMDGCN20-SAME: ptr addrspace(5) dead_on_unwind noalias writable sret([[STRUCT_MAT64X64:%.*]]) align 4 [[AGG_RESULT:%.*]], ptr addrspace(5) noundef byref([[STRUCT_MAT32X32:%.*]]) align 4 [[TMP0:%.*]]) #[[ATTR0]] { // AMDGCN20-NEXT: [[ENTRY:.*:]] // AMDGCN20-NEXT: [[COERCE:%.*]] = alloca [[STRUCT_MAT32X32]], align 4, addrspace(5) // AMDGCN20-NEXT: [[IN:%.*]] = addrspacecast ptr addrspace(5) [[COERCE]] to ptr @@ -335,7 +335,7 @@ Mat64X64 __attribute__((noinline)) foo_large(Mat32X32 in) { // AMDGCN20-NEXT: [[TMP1:%.*]] = load ptr addrspace(1), ptr [[IN_ADDR_ASCAST]], align 8 // AMDGCN20-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [[STRUCT_MAT32X32]], ptr addrspace(1) [[TMP1]], i64 1 // AMDGCN20-NEXT: call void @llvm.memcpy.p5.p1.i64(ptr addrspace(5) align 4 [[BYVAL_TEMP]], ptr addrspace(1) align 4 [[ARRAYIDX1]], i64 4096, i1 false) -// AMDGCN20-NEXT: call void @foo_large(ptr dead_on_unwind writable sret([[STRUCT_MAT64X64]]) align 4 [[TMP_ASCAST]], ptr addrspace(5) noundef byref([[STRUCT_MAT32X32]]) align 4 [[BYVAL_TEMP]]) #[[ATTR3]] +// AMDGCN20-NEXT: call void @foo_large(ptr addrspace(5) dead_on_unwind writable sret([[STRUCT_MAT64X64]]) align 4 [[TMP]], ptr addrspace(5) noundef byref([[STRUCT_MAT32X32]]) align 4 [[BYVAL_TEMP]]) #[[ATTR3]] // AMDGCN20-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 4 [[ARRAYIDX]], ptr align 4 [[TMP_ASCAST]], i64 16384, i1 false) // AMDGCN20-NEXT: ret void // diff --git a/clang/test/CodeGenOpenCL/amdgpu-abi-struct-arg-byref.cl b/clang/test/CodeGenOpenCL/amdgpu-abi-struct-arg-byref.cl index 084281a8cada4..c2b2e00d15e13 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-abi-struct-arg-byref.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-abi-struct-arg-byref.cl @@ -91,7 +91,7 @@ kernel void ker(global Mat3X3 *in, global Mat4X4 *out) { } // AMDGCN-LABEL: define dso_local void @foo_large( -// AMDGCN-SAME: ptr dead_on_unwind noalias writable sret([[STRUCT_MAT64X64:%.*]]) align 4 [[AGG_RESULT:%.*]], ptr addrspace(5) noundef byref([[STRUCT_MAT32X32:%.*]]) align 4 [[TMP0:%.*]]) #[[ATTR0]] { +// AMDGCN-SAME: ptr addrspace(5) dead_on_unwind noalias writable sret([[STRUCT_MAT64X64:%.*]]) align 4 [[AGG_RESULT:%.*]], ptr addrspace(5) noundef byref([[STRUCT_MAT32X32:%.*]]) align 4 [[TMP0:%.*]]) #[[ATTR0]] { // AMDGCN-NEXT: [[ENTRY:.*:]] // AMDGCN-NEXT: [[COERCE:%.*]] = alloca [[STRUCT_MAT32X32]], align 4, addrspace(5) // AMDGCN-NEXT: [[IN:%.*]] = addrspacecast ptr addrspace(5) [[COERCE]] to ptr @@ -120,7 +120,7 @@ Mat64X64 __attribute__((noinline)) foo_large(Mat32X32 in) { // AMDGCN-NEXT: [[TMP1:%.*]] = load ptr addrspace(1), ptr [[IN_ADDR_ASCAST]], align 8 // AMDGCN-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [[STRUCT_MAT32X32]], ptr addrspace(1) [[TMP1]], i64 1 // AMDGCN-NEXT: call void @llvm.memcpy.p5.p1.i64(ptr addrspace(5) align 4 [[BYVAL_TEMP]], ptr addrspace(1) align 4 [[ARRAYIDX1]], i64 4096, i1 false) -// AMDGCN-NEXT: call void @foo_large(ptr dead_on_unwind writable sret([[STRUCT_MAT64X64]]) align 4 [[TMP_ASCAST]], ptr addrspace(5) noundef byref([[STRUCT_MAT32X32]]) align 4 [[BYVAL_TEMP]]) #[[ATTR3]] +// AMDGCN-NEXT: call void @foo_large(ptr addrspace(5) dead_on_unwind writable sret([[STRUCT_MAT64X64]]) align 4 [[TMP]], ptr addrspace(5) noundef byref([[STRUCT_MAT32X32]]) align 4 [[BYVAL_TEMP]]) #[[ATTR3]] // AMDGCN-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 4 [[ARRAYIDX]], ptr align 4 [[TMP_ASCAST]], i64 16384, i1 false) // AMDGCN-NEXT: ret void // From 2de33d4cfb210dc50a55b9ba87fa0d086d4b8d9f Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Wed, 30 Oct 2024 00:10:59 +0000 Subject: [PATCH 03/27] Handle passing an `alloca`ed `sret` arg directly to a callee that expects a pointer to the default AS. --- clang/lib/CodeGen/CGCall.cpp | 16 ++++++++++++---- clang/test/CodeGenCXX/no-elide-constructors.cpp | 4 ++++ 2 files changed, 16 insertions(+), 4 deletions(-) diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 56acfae7ae9e5..7171d85b0d0ab 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -5391,11 +5391,19 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo, V->getType()->isIntegerTy()) V = Builder.CreateZExt(V, ArgInfo.getCoerceToType()); - // If the argument doesn't match, perform a bitcast to coerce it. This - // can happen due to trivial type mismatches. + // If the argument doesn't match, we are either trying to pass an + // alloca-ed sret argument directly, and the alloca AS does not match + // the default AS, case in which we AS cast it, or we have a trivial + // type mismatch, and thus perform a bitcast to coerce it. if (FirstIRArg < IRFuncTy->getNumParams() && - V->getType() != IRFuncTy->getParamType(FirstIRArg)) - V = Builder.CreateBitCast(V, IRFuncTy->getParamType(FirstIRArg)); + V->getType() != IRFuncTy->getParamType(FirstIRArg)) { + auto IRTy = IRFuncTy->getParamType(FirstIRArg); + auto MaybeSRetArg = dyn_cast_or_null(V); + if (MaybeSRetArg && MaybeSRetArg->hasStructRetAttr()) + V = Builder.CreateAddrSpaceCast(V, IRTy); + else + V = Builder.CreateBitCast(V, IRTy); + } if (ArgHasMaybeUndefAttr) V = Builder.CreateFreeze(V); diff --git a/clang/test/CodeGenCXX/no-elide-constructors.cpp b/clang/test/CodeGenCXX/no-elide-constructors.cpp index 750392a43e05c..098163f957f75 100644 --- a/clang/test/CodeGenCXX/no-elide-constructors.cpp +++ b/clang/test/CodeGenCXX/no-elide-constructors.cpp @@ -1,7 +1,9 @@ // RUN: %clang_cc1 -std=c++98 -triple i386-unknown-unknown -fno-elide-constructors -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK --check-prefix=CHECK-CXX98 // RUN: %clang_cc1 -std=c++11 -triple i386-unknown-unknown -fno-elide-constructors -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK --check-prefix=CHECK-CXX11 +// RUN: %clang_cc1 -std=c++11 -triple amdgcn-amd-amdhsa -fno-elide-constructors -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK --check-prefix=CHECK-CXX11-NONZEROALLOCAAS // RUN: %clang_cc1 -std=c++98 -triple i386-unknown-unknown -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK --check-prefix=CHECK-CXX98-ELIDE // RUN: %clang_cc1 -std=c++11 -triple i386-unknown-unknown -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK --check-prefix=CHECK-CXX11-ELIDE +// RUN: %clang_cc1 -std=c++11 -triple amdgcn-amd-amdhsa -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK --check-prefix=CHECK-CXX11-NONZEROALLOCAAS-ELIDE // Reduced from PR12208 class X { @@ -23,8 +25,10 @@ X Test() // sret argument. // CHECK-CXX98: call void @_ZN1XC1ERKS_( // CHECK-CXX11: call void @_ZN1XC1EOS_( + // CHECK-CXX11-NONZEROALLOCAAS: call void @_ZN1XC1EOS_( // CHECK-CXX98-ELIDE-NOT: call void @_ZN1XC1ERKS_( // CHECK-CXX11-ELIDE-NOT: call void @_ZN1XC1EOS_( + // CHECK-CXX11-NONZEROALLOCAAS-ELIDE-NOT: call void @_ZN1XC1EOS_( // Make sure that the destructor for X is called. // FIXME: This call is present even in the -ELIDE runs, but is guarded by a From b209d6779cccaa9c2f272d839263cf7ca139b945 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Sat, 2 Nov 2024 00:57:17 +0000 Subject: [PATCH 04/27] Add query for a possible target specific indirect arg AS. --- clang/include/clang/Basic/TargetInfo.h | 8 ++++++++ clang/lib/CodeGen/CGCall.cpp | 6 ++++-- 2 files changed, 12 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h index 25eda907d20a7..fa5021baf667b 100644 --- a/clang/include/clang/Basic/TargetInfo.h +++ b/clang/include/clang/Basic/TargetInfo.h @@ -1780,6 +1780,14 @@ class TargetInfo : public TransferrableTargetInfo, return 0; } + /// \returns Target specific address space for indirect (e.g. sret) arguments. + /// If such an address space exists, it must be convertible to and from the + /// alloca address space. If it does not, std::nullopt is returned and the + /// alloca address space will be used. + virtual std::optional getIndirectArgAddressSpace() const { + return std::nullopt; + } + /// \returns If a target requires an address within a target specific address /// space \p AddressSpace to be converted in order to be used, then return the /// corresponding target specific DWARF address space. diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 7171d85b0d0ab..87e70df795a98 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -1672,9 +1672,11 @@ CodeGenTypes::GetFunctionType(const CGFunctionInfo &FI) { // Add type for sret argument. if (IRFunctionArgs.hasSRetArg()) { - unsigned AddressSpace = CGM.getDataLayout().getAllocaAddrSpace(); + auto AddressSpace = CGM.getTarget().getIndirectArgAddressSpace(); + if (!AddressSpace) + AddressSpace = getDataLayout().getAllocaAddrSpace(); ArgTypes[IRFunctionArgs.getSRetArgNo()] = - llvm::PointerType::get(getLLVMContext(), AddressSpace); + llvm::PointerType::get(getLLVMContext(), *AddressSpace); } // Add type for inalloca argument. From ac6367be734abec8f2c46f4fe8a13e950e13578f Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Sat, 2 Nov 2024 01:20:12 +0000 Subject: [PATCH 05/27] Add more context to test. --- clang/test/CodeGenCXX/no-elide-constructors.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/clang/test/CodeGenCXX/no-elide-constructors.cpp b/clang/test/CodeGenCXX/no-elide-constructors.cpp index 098163f957f75..994282debb0d0 100644 --- a/clang/test/CodeGenCXX/no-elide-constructors.cpp +++ b/clang/test/CodeGenCXX/no-elide-constructors.cpp @@ -17,6 +17,7 @@ class X { }; // CHECK-LABEL: define{{.*}} void @_Z4Testv( +// CHECK-SAME: ptr {{.*}}dead_on_unwind noalias writable sret([[CLASS_X:%.*]]) align 1 [[AGG_RESULT:%.*]]) X Test() { X x; @@ -25,7 +26,8 @@ X Test() // sret argument. // CHECK-CXX98: call void @_ZN1XC1ERKS_( // CHECK-CXX11: call void @_ZN1XC1EOS_( - // CHECK-CXX11-NONZEROALLOCAAS: call void @_ZN1XC1EOS_( + // CHECK-CXX11-NONZEROALLOCAAS: [[TMP0:%.*]] = addrspacecast ptr addrspace(5) [[AGG_RESULT]] to ptr + // CHECK-CXX11-NONZEROALLOCAAS-NEXT: call void @_ZN1XC1EOS_(ptr noundef nonnull align 1 dereferenceable(1) [[TMP0]] // CHECK-CXX98-ELIDE-NOT: call void @_ZN1XC1ERKS_( // CHECK-CXX11-ELIDE-NOT: call void @_ZN1XC1EOS_( // CHECK-CXX11-NONZEROALLOCAAS-ELIDE-NOT: call void @_ZN1XC1EOS_( From 9ff1d0dd16bbda206753348ab9671dcfe0b5eb7b Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Wed, 6 Nov 2024 13:16:03 +0200 Subject: [PATCH 06/27] Extend Indirect Args to carry an address space. --- clang/include/clang/CodeGen/CGFunctionInfo.h | 11 ++++++----- clang/lib/CodeGen/ABIInfo.cpp | 2 +- clang/lib/CodeGen/ItaniumCXXABI.cpp | 2 +- clang/lib/CodeGen/MicrosoftCXXABI.cpp | 2 +- clang/lib/CodeGen/SwiftCallingConv.cpp | 4 ++-- clang/lib/CodeGen/Targets/AMDGPU.cpp | 5 +++++ clang/lib/CodeGen/Targets/ARC.cpp | 2 +- clang/lib/CodeGen/Targets/ARM.cpp | 4 ++-- clang/lib/CodeGen/Targets/Lanai.cpp | 2 +- clang/lib/CodeGen/Targets/PPC.cpp | 4 ++-- clang/lib/CodeGen/Targets/X86.cpp | 16 ++++++++-------- 11 files changed, 30 insertions(+), 24 deletions(-) diff --git a/clang/include/clang/CodeGen/CGFunctionInfo.h b/clang/include/clang/CodeGen/CGFunctionInfo.h index 9d785d878b61d..4ca5d2b654812 100644 --- a/clang/include/clang/CodeGen/CGFunctionInfo.h +++ b/clang/include/clang/CodeGen/CGFunctionInfo.h @@ -206,8 +206,8 @@ class ABIArgInfo { static ABIArgInfo getIgnore() { return ABIArgInfo(Ignore); } - static ABIArgInfo getIndirect(CharUnits Alignment, bool ByVal = true, - bool Realign = false, + static ABIArgInfo getIndirect(CharUnits Alignment, unsigned AddrSpace = 0, + bool ByVal = true, bool Realign = false, llvm::Type *Padding = nullptr) { auto AI = ABIArgInfo(Indirect); AI.setIndirectAlign(Alignment); @@ -215,6 +215,7 @@ class ABIArgInfo { AI.setIndirectRealign(Realign); AI.setSRetAfterThis(false); AI.setPaddingType(Padding); + AI.setIndirectAddrSpace(AddrSpace); return AI; } @@ -232,7 +233,7 @@ class ABIArgInfo { static ABIArgInfo getIndirectInReg(CharUnits Alignment, bool ByVal = true, bool Realign = false) { - auto AI = getIndirect(Alignment, ByVal, Realign); + auto AI = getIndirect(Alignment, 0, ByVal, Realign); AI.setInReg(true); return AI; } @@ -422,12 +423,12 @@ class ABIArgInfo { } unsigned getIndirectAddrSpace() const { - assert(isIndirectAliased() && "Invalid kind!"); + assert((isIndirect() || isIndirectAliased()) && "Invalid kind!"); return IndirectAttr.AddrSpace; } void setIndirectAddrSpace(unsigned AddrSpace) { - assert(isIndirectAliased() && "Invalid kind!"); + assert((isIndirect() || isIndirectAliased()) && "Invalid kind!"); IndirectAttr.AddrSpace = AddrSpace; } diff --git a/clang/lib/CodeGen/ABIInfo.cpp b/clang/lib/CodeGen/ABIInfo.cpp index edd7146dc1ac7..7ab9f0aeb6099 100644 --- a/clang/lib/CodeGen/ABIInfo.cpp +++ b/clang/lib/CodeGen/ABIInfo.cpp @@ -174,7 +174,7 @@ bool ABIInfo::isPromotableIntegerTypeForABI(QualType Ty) const { ABIArgInfo ABIInfo::getNaturalAlignIndirect(QualType Ty, bool ByVal, bool Realign, llvm::Type *Padding) const { - return ABIArgInfo::getIndirect(getContext().getTypeAlignInChars(Ty), ByVal, + return ABIArgInfo::getIndirect(getContext().getTypeAlignInChars(Ty), 0, ByVal, Realign, Padding); } diff --git a/clang/lib/CodeGen/ItaniumCXXABI.cpp b/clang/lib/CodeGen/ItaniumCXXABI.cpp index 9b3c2f1b2af67..f5e2b096212f4 100644 --- a/clang/lib/CodeGen/ItaniumCXXABI.cpp +++ b/clang/lib/CodeGen/ItaniumCXXABI.cpp @@ -1351,7 +1351,7 @@ bool ItaniumCXXABI::classifyReturnType(CGFunctionInfo &FI) const { // If C++ prohibits us from making a copy, return by address. if (!RD->canPassInRegisters()) { auto Align = CGM.getContext().getTypeAlignInChars(FI.getReturnType()); - FI.getReturnInfo() = ABIArgInfo::getIndirect(Align, /*ByVal=*/false); + FI.getReturnInfo() = ABIArgInfo::getIndirect(Align, 0, /*ByVal=*/false); return true; } return false; diff --git a/clang/lib/CodeGen/MicrosoftCXXABI.cpp b/clang/lib/CodeGen/MicrosoftCXXABI.cpp index 3802dc8bcafc4..3b5b860a1b087 100644 --- a/clang/lib/CodeGen/MicrosoftCXXABI.cpp +++ b/clang/lib/CodeGen/MicrosoftCXXABI.cpp @@ -1171,7 +1171,7 @@ bool MicrosoftCXXABI::classifyReturnType(CGFunctionInfo &FI) const { if (isIndirectReturn) { CharUnits Align = CGM.getContext().getTypeAlignInChars(FI.getReturnType()); - FI.getReturnInfo() = ABIArgInfo::getIndirect(Align, /*ByVal=*/false); + FI.getReturnInfo() = ABIArgInfo::getIndirect(Align, 0, /*ByVal=*/false); // MSVC always passes `this` before the `sret` parameter. FI.getReturnInfo().setSRetAfterThis(FI.isInstanceMethod()); diff --git a/clang/lib/CodeGen/SwiftCallingConv.cpp b/clang/lib/CodeGen/SwiftCallingConv.cpp index ab2e2bd0b3064..e178c0fab5910 100644 --- a/clang/lib/CodeGen/SwiftCallingConv.cpp +++ b/clang/lib/CodeGen/SwiftCallingConv.cpp @@ -801,7 +801,7 @@ static ABIArgInfo classifyExpandedType(SwiftAggLowering &lowering, if (lowering.empty()) { return ABIArgInfo::getIgnore(); } else if (lowering.shouldPassIndirectly(forReturn)) { - return ABIArgInfo::getIndirect(alignmentForIndirect, /*byval*/ false); + return ABIArgInfo::getIndirect(alignmentForIndirect, 0, /*byval*/ false); } else { auto types = lowering.getCoerceAndExpandTypes(); return ABIArgInfo::getCoerceAndExpand(types.first, types.second); @@ -815,7 +815,7 @@ static ABIArgInfo classifyType(CodeGenModule &CGM, CanQualType type, auto &layout = CGM.getContext().getASTRecordLayout(record); if (mustPassRecordIndirectly(CGM, record)) - return ABIArgInfo::getIndirect(layout.getAlignment(), /*byval*/ false); + return ABIArgInfo::getIndirect(layout.getAlignment(), 0, /*byval*/ false); SwiftAggLowering lowering(CGM); lowering.addTypedData(recordType->getDecl(), CharUnits::Zero(), layout); diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 56ad0503a11ab..c45e7020de3f5 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -105,6 +105,11 @@ void AMDGPUABIInfo::computeInfo(CGFunctionInfo &FI) const { if (!getCXXABI().classifyReturnType(FI)) FI.getReturnInfo() = classifyReturnType(FI.getReturnType()); + // srets / indirect returns are unconditionally in the alloca AS. + if (FI.getReturnInfo().isIndirect()) + FI.getReturnInfo().setIndirectAddrSpace( + getDataLayout().getAllocaAddrSpace()); + unsigned ArgumentIndex = 0; const unsigned numFixedArguments = FI.getNumRequiredArgs(); diff --git a/clang/lib/CodeGen/Targets/ARC.cpp b/clang/lib/CodeGen/Targets/ARC.cpp index 1904e8fdb3888..ee0db9778bdcb 100644 --- a/clang/lib/CodeGen/Targets/ARC.cpp +++ b/clang/lib/CodeGen/Targets/ARC.cpp @@ -77,7 +77,7 @@ ABIArgInfo ARCABIInfo::getIndirectByValue(QualType Ty) const { // Compute the byval alignment. const unsigned MinABIStackAlignInBytes = 4; unsigned TypeAlign = getContext().getTypeAlign(Ty) / 8; - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(4), /*ByVal=*/true, + return ABIArgInfo::getIndirect(CharUnits::fromQuantity(4), 0, /*ByVal=*/true, TypeAlign > MinABIStackAlignInBytes); } diff --git a/clang/lib/CodeGen/Targets/ARM.cpp b/clang/lib/CodeGen/Targets/ARM.cpp index 2d858fa2f3c3a..d89a0bdff56a3 100644 --- a/clang/lib/CodeGen/Targets/ARM.cpp +++ b/clang/lib/CodeGen/Targets/ARM.cpp @@ -397,7 +397,7 @@ ABIArgInfo ARMABIInfo::classifyArgumentType(QualType Ty, bool isVariadic, // bigger than 128-bits, they get placed in space allocated by the caller, // and a pointer is passed. return ABIArgInfo::getIndirect( - CharUnits::fromQuantity(getContext().getTypeAlign(Ty) / 8), false); + CharUnits::fromQuantity(getContext().getTypeAlign(Ty) / 8), 0, false); } // Support byval for ARM. @@ -415,7 +415,7 @@ ABIArgInfo ARMABIInfo::classifyArgumentType(QualType Ty, bool isVariadic, } if (getContext().getTypeSizeInChars(Ty) > CharUnits::fromQuantity(64)) { assert(getABIKind() != ARMABIKind::AAPCS16_VFP && "unexpected byval"); - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(ABIAlign), + return ABIArgInfo::getIndirect(CharUnits::fromQuantity(ABIAlign), 0, /*ByVal=*/true, /*Realign=*/TyAlign > ABIAlign); } diff --git a/clang/lib/CodeGen/Targets/Lanai.cpp b/clang/lib/CodeGen/Targets/Lanai.cpp index 2578fc0291e76..ffacb0ccbea53 100644 --- a/clang/lib/CodeGen/Targets/Lanai.cpp +++ b/clang/lib/CodeGen/Targets/Lanai.cpp @@ -78,7 +78,7 @@ ABIArgInfo LanaiABIInfo::getIndirectResult(QualType Ty, bool ByVal, // Compute the byval alignment. const unsigned MinABIStackAlignInBytes = 4; unsigned TypeAlign = getContext().getTypeAlign(Ty) / 8; - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(4), /*ByVal=*/true, + return ABIArgInfo::getIndirect(CharUnits::fromQuantity(4), 0, /*ByVal=*/true, /*Realign=*/TypeAlign > MinABIStackAlignInBytes); } diff --git a/clang/lib/CodeGen/Targets/PPC.cpp b/clang/lib/CodeGen/Targets/PPC.cpp index 989e46f4b66a7..c8796036b214f 100644 --- a/clang/lib/CodeGen/Targets/PPC.cpp +++ b/clang/lib/CodeGen/Targets/PPC.cpp @@ -213,7 +213,7 @@ ABIArgInfo AIXABIInfo::classifyArgumentType(QualType Ty) const { CharUnits CCAlign = getParamTypeAlignment(Ty); CharUnits TyAlign = getContext().getTypeAlignInChars(Ty); - return ABIArgInfo::getIndirect(CCAlign, /*ByVal*/ true, + return ABIArgInfo::getIndirect(CCAlign, 0, /*ByVal*/ true, /*Realign*/ TyAlign > CCAlign); } @@ -887,7 +887,7 @@ PPC64_SVR4_ABIInfo::classifyArgumentType(QualType Ty) const { } // All other aggregates are passed ByVal. - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(ABIAlign), + return ABIArgInfo::getIndirect(CharUnits::fromQuantity(ABIAlign), 0, /*ByVal=*/true, /*Realign=*/TyAlign > ABIAlign); } diff --git a/clang/lib/CodeGen/Targets/X86.cpp b/clang/lib/CodeGen/Targets/X86.cpp index 7f73bf2a65266..f097c27bd8947 100644 --- a/clang/lib/CodeGen/Targets/X86.cpp +++ b/clang/lib/CodeGen/Targets/X86.cpp @@ -606,12 +606,12 @@ ABIArgInfo X86_32ABIInfo::getIndirectResult(QualType Ty, bool ByVal, unsigned TypeAlign = getContext().getTypeAlign(Ty) / 8; unsigned StackAlign = getTypeStackAlignInBytes(Ty, TypeAlign); if (StackAlign == 0) - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(4), /*ByVal=*/true); + return ABIArgInfo::getIndirect(CharUnits::fromQuantity(4), 0, /*ByVal=*/true); // If the stack alignment is less than the type alignment, realign the // argument. bool Realign = TypeAlign > StackAlign; - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(StackAlign), + return ABIArgInfo::getIndirect(CharUnits::fromQuantity(StackAlign), 0, /*ByVal=*/true, Realign); } @@ -2247,7 +2247,7 @@ ABIArgInfo X86_64ABIInfo::getIndirectResult(QualType Ty, Size)); } - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(Align)); + return ABIArgInfo::getIndirect(CharUnits::fromQuantity(Align), 0); } /// The ABI specifies that a value should be passed in a full vector XMM/YMM @@ -3304,7 +3304,7 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, return ABIArgInfo::getDirect(); return ABIArgInfo::getExpand(); } - return ABIArgInfo::getIndirect(Align, /*ByVal=*/false); + return ABIArgInfo::getIndirect(Align, 0, /*ByVal=*/false); } else if (IsVectorCall) { if (FreeSSERegs >= NumElts && (IsReturnType || Ty->isBuiltinType() || Ty->isVectorType())) { @@ -3314,7 +3314,7 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, return ABIArgInfo::getExpand(); } else if (!Ty->isBuiltinType() && !Ty->isVectorType()) { // HVAs are delayed and reclassified in the 2nd step. - return ABIArgInfo::getIndirect(Align, /*ByVal=*/false); + return ABIArgInfo::getIndirect(Align, 0, /*ByVal=*/false); } } } @@ -3350,7 +3350,7 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, if (IsMingw64) { const llvm::fltSemantics *LDF = &getTarget().getLongDoubleFormat(); if (LDF == &llvm::APFloat::x87DoubleExtended()) - return ABIArgInfo::getIndirect(Align, /*ByVal=*/false); + return ABIArgInfo::getIndirect(Align, 0, /*ByVal=*/false); } break; @@ -3360,7 +3360,7 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, // than 8 bytes are passed indirectly. GCC follows it. We follow it too, // even though it isn't particularly efficient. if (!IsReturnType) - return ABIArgInfo::getIndirect(Align, /*ByVal=*/false); + return ABIArgInfo::getIndirect(Align, 0, /*ByVal=*/false); // Mingw64 GCC returns i128 in XMM0. Coerce to v2i64 to handle that. // Clang matches them for compatibility. @@ -3380,7 +3380,7 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, // the power of 2. if (Width <= 64) return ABIArgInfo::getDirect(); - return ABIArgInfo::getIndirect(Align, /*ByVal=*/false); + return ABIArgInfo::getIndirect(Align, 0, /*ByVal=*/false); } return ABIArgInfo::getDirect(); From 1c3e67cdebf6025aacd1900c22f033504d8e7963 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Wed, 6 Nov 2024 13:21:51 +0200 Subject: [PATCH 07/27] Fix formatting. --- clang/lib/CodeGen/Targets/X86.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/Targets/X86.cpp b/clang/lib/CodeGen/Targets/X86.cpp index f097c27bd8947..6e5b46d5f91c8 100644 --- a/clang/lib/CodeGen/Targets/X86.cpp +++ b/clang/lib/CodeGen/Targets/X86.cpp @@ -606,7 +606,8 @@ ABIArgInfo X86_32ABIInfo::getIndirectResult(QualType Ty, bool ByVal, unsigned TypeAlign = getContext().getTypeAlign(Ty) / 8; unsigned StackAlign = getTypeStackAlignInBytes(Ty, TypeAlign); if (StackAlign == 0) - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(4), 0, /*ByVal=*/true); + return ABIArgInfo::getIndirect(CharUnits::fromQuantity(4), 0, + /*ByVal=*/true); // If the stack alignment is less than the type alignment, realign the // argument. From c9288fc9d38c603ef120714343b2a57611fda424 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Fri, 8 Nov 2024 01:13:11 +0200 Subject: [PATCH 08/27] Drop vestigial target hook. --- clang/include/clang/Basic/TargetInfo.h | 8 -------- clang/lib/CodeGen/CGCall.cpp | 7 ++----- 2 files changed, 2 insertions(+), 13 deletions(-) diff --git a/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h index fa5021baf667b..25eda907d20a7 100644 --- a/clang/include/clang/Basic/TargetInfo.h +++ b/clang/include/clang/Basic/TargetInfo.h @@ -1780,14 +1780,6 @@ class TargetInfo : public TransferrableTargetInfo, return 0; } - /// \returns Target specific address space for indirect (e.g. sret) arguments. - /// If such an address space exists, it must be convertible to and from the - /// alloca address space. If it does not, std::nullopt is returned and the - /// alloca address space will be used. - virtual std::optional getIndirectArgAddressSpace() const { - return std::nullopt; - } - /// \returns If a target requires an address within a target specific address /// space \p AddressSpace to be converted in order to be used, then return the /// corresponding target specific DWARF address space. diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 87e70df795a98..32200ada7cf7d 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -1672,11 +1672,8 @@ CodeGenTypes::GetFunctionType(const CGFunctionInfo &FI) { // Add type for sret argument. if (IRFunctionArgs.hasSRetArg()) { - auto AddressSpace = CGM.getTarget().getIndirectArgAddressSpace(); - if (!AddressSpace) - AddressSpace = getDataLayout().getAllocaAddrSpace(); - ArgTypes[IRFunctionArgs.getSRetArgNo()] = - llvm::PointerType::get(getLLVMContext(), *AddressSpace); + ArgTypes[IRFunctionArgs.getSRetArgNo()] = llvm::PointerType::get( + getLLVMContext(), FI.getReturnInfo().getIndirectAddrSpace()); } // Add type for inalloca argument. From 013790c1ad46cbcc143fb30fccc1ead25947da24 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Fri, 15 Nov 2024 23:04:53 +0000 Subject: [PATCH 09/27] Tweak handling potential AS mismatches. --- clang/lib/CodeGen/CGCall.cpp | 24 ++++++++++--------- clang/lib/CodeGen/CGExprAgg.cpp | 19 ++++++++++----- .../CodeGenOpenCL/addr-space-struct-arg.cl | 10 ++++---- .../amdgpu-abi-struct-arg-byref.cl | 10 ++++---- 4 files changed, 34 insertions(+), 29 deletions(-) diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 32200ada7cf7d..41105262b5f14 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -5168,7 +5168,7 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo, // alloca was AS casted to the default as, so we ensure the cast is // stripped before binding to the sret arg, which is in the allocaAS. IRCallArgs[IRFunctionArgs.getSRetArgNo()] = - getAsNaturalPointerTo(SRetPtr, RetTy)->stripPointerCasts(); + getAsNaturalPointerTo(SRetPtr, RetTy); } else if (RetAI.isInAlloca()) { Address Addr = Builder.CreateStructGEP(ArgMemory, RetAI.getInAllocaFieldIndex()); @@ -5390,18 +5390,20 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo, V->getType()->isIntegerTy()) V = Builder.CreateZExt(V, ArgInfo.getCoerceToType()); - // If the argument doesn't match, we are either trying to pass an - // alloca-ed sret argument directly, and the alloca AS does not match - // the default AS, case in which we AS cast it, or we have a trivial - // type mismatch, and thus perform a bitcast to coerce it. + // The only plausible mismatch here would be for pointer address spaces, + // which can happen e.g. when passing a sret arg that is in the AllocaAS + // to a function that takes a pointer to and argument in the DefaultAS. + // We assume that the target has a reasonable mapping for the DefaultAS + // (it can be casted to from incoming specific ASes), and insert an AS + // cast to address the mismatch. if (FirstIRArg < IRFuncTy->getNumParams() && V->getType() != IRFuncTy->getParamType(FirstIRArg)) { - auto IRTy = IRFuncTy->getParamType(FirstIRArg); - auto MaybeSRetArg = dyn_cast_or_null(V); - if (MaybeSRetArg && MaybeSRetArg->hasStructRetAttr()) - V = Builder.CreateAddrSpaceCast(V, IRTy); - else - V = Builder.CreateBitCast(V, IRTy); + assert(V->getType()->isPointerTy() && "Only pointers can mismatch!"); + auto FormalAS = + CallInfo.arguments()[ArgNo].type.getQualifiers().getAddressSpace(); + auto ActualAS = I->Ty.getAddressSpace(); + V = getTargetHooks().performAddrSpaceCast( + *this, V, ActualAS, FormalAS, IRFuncTy->getParamType(FirstIRArg)); } if (ArgHasMaybeUndefAttr) diff --git a/clang/lib/CodeGen/CGExprAgg.cpp b/clang/lib/CodeGen/CGExprAgg.cpp index 2ad6587089f10..f9c9c5df80163 100644 --- a/clang/lib/CodeGen/CGExprAgg.cpp +++ b/clang/lib/CodeGen/CGExprAgg.cpp @@ -296,18 +296,25 @@ void AggExprEmitter::withReturnValueSlot( (RequiresDestruction && Dest.isIgnored()); Address RetAddr = Address::invalid(); - RawAddress RetAllocaAddr = RawAddress::invalid(); EHScopeStack::stable_iterator LifetimeEndBlock; llvm::Value *LifetimeSizePtr = nullptr; llvm::IntrinsicInst *LifetimeStartInst = nullptr; if (!UseTemp) { - RetAddr = Dest.getAddress(); + // It is possible for the existing slot we are using directly to have been + // allocated in the correct AS for an indirect return, and then cast to + // the default AS (this is the behaviour of CreateMemTemp), however we know + // that the return address is expected to point to the uncasted AS, hence we + // strip possible pointer casts here. + if (Dest.getAddress().isValid()) + RetAddr = Dest.getAddress().withPointer( + Dest.getAddress().getBasePointer()->stripPointerCasts(), + Dest.getAddress().isKnownNonNull()); } else { - RetAddr = CGF.CreateMemTemp(RetTy, "tmp", &RetAllocaAddr); + RetAddr = CGF.CreateMemTempWithoutCast(RetTy, "tmp"); llvm::TypeSize Size = CGF.CGM.getDataLayout().getTypeAllocSize(CGF.ConvertTypeForMem(RetTy)); - LifetimeSizePtr = CGF.EmitLifetimeStart(Size, RetAllocaAddr.getPointer()); + LifetimeSizePtr = CGF.EmitLifetimeStart(Size, RetAddr.getBasePointer()); if (LifetimeSizePtr) { LifetimeStartInst = cast(std::prev(Builder.GetInsertPoint())); @@ -316,7 +323,7 @@ void AggExprEmitter::withReturnValueSlot( "Last insertion wasn't a lifetime.start?"); CGF.pushFullExprCleanup( - NormalEHLifetimeMarker, RetAllocaAddr, LifetimeSizePtr); + NormalEHLifetimeMarker, RetAddr, LifetimeSizePtr); LifetimeEndBlock = CGF.EHStack.stable_begin(); } } @@ -337,7 +344,7 @@ void AggExprEmitter::withReturnValueSlot( // Since we're not guaranteed to be in an ExprWithCleanups, clean up // eagerly. CGF.DeactivateCleanupBlock(LifetimeEndBlock, LifetimeStartInst); - CGF.EmitLifetimeEnd(LifetimeSizePtr, RetAllocaAddr.getPointer()); + CGF.EmitLifetimeEnd(LifetimeSizePtr, RetAddr.getBasePointer()); } } diff --git a/clang/test/CodeGenOpenCL/addr-space-struct-arg.cl b/clang/test/CodeGenOpenCL/addr-space-struct-arg.cl index 4a1db2c3564a5..effdeb9546800 100644 --- a/clang/test/CodeGenOpenCL/addr-space-struct-arg.cl +++ b/clang/test/CodeGenOpenCL/addr-space-struct-arg.cl @@ -154,7 +154,6 @@ Mat4X4 __attribute__((noinline)) foo(Mat3X3 in) { // AMDGCN20-NEXT: [[TMP:%.*]] = alloca [[STRUCT_MAT4X4:%.*]], align 4, addrspace(5) // AMDGCN20-NEXT: [[IN_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[IN_ADDR]] to ptr // AMDGCN20-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr -// AMDGCN20-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr // AMDGCN20-NEXT: store ptr addrspace(1) [[IN]], ptr [[IN_ADDR_ASCAST]], align 8 // AMDGCN20-NEXT: store ptr addrspace(1) [[OUT]], ptr [[OUT_ADDR_ASCAST]], align 8 // AMDGCN20-NEXT: [[TMP0:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8 @@ -164,10 +163,10 @@ Mat4X4 __attribute__((noinline)) foo(Mat3X3 in) { // AMDGCN20-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT_MAT3X3]], ptr addrspace(1) [[ARRAYIDX1]], i32 0, i32 0 // AMDGCN20-NEXT: [[TMP3:%.*]] = load [9 x i32], ptr addrspace(1) [[TMP2]], align 4 // AMDGCN20-NEXT: [[CALL:%.*]] = call [[STRUCT_MAT4X4]] @[[FOO:[a-zA-Z0-9_$\"\\.-]*[a-zA-Z_$\"\\.-][a-zA-Z0-9_$\"\\.-]*]]([9 x i32] [[TMP3]]) #[[ATTR3:[0-9]+]] -// AMDGCN20-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT_MAT4X4]], ptr [[TMP_ASCAST]], i32 0, i32 0 +// AMDGCN20-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT_MAT4X4]], ptr addrspace(5) [[TMP]], i32 0, i32 0 // AMDGCN20-NEXT: [[TMP5:%.*]] = extractvalue [[STRUCT_MAT4X4]] [[CALL]], 0 -// AMDGCN20-NEXT: store [16 x i32] [[TMP5]], ptr [[TMP4]], align 4 -// AMDGCN20-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 4 [[ARRAYIDX]], ptr align 4 [[TMP_ASCAST]], i64 64, i1 false) +// AMDGCN20-NEXT: store [16 x i32] [[TMP5]], ptr addrspace(5) [[TMP4]], align 4 +// AMDGCN20-NEXT: call void @llvm.memcpy.p1.p5.i64(ptr addrspace(1) align 4 [[ARRAYIDX]], ptr addrspace(5) align 4 [[TMP]], i64 64, i1 false) // AMDGCN20-NEXT: ret void // // SPIR-LABEL: define dso_local spir_kernel void @ker( @@ -327,7 +326,6 @@ Mat64X64 __attribute__((noinline)) foo_large(Mat32X32 in) { // AMDGCN20-NEXT: [[BYVAL_TEMP:%.*]] = alloca [[STRUCT_MAT32X32:%.*]], align 4, addrspace(5) // AMDGCN20-NEXT: [[IN_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[IN_ADDR]] to ptr // AMDGCN20-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr -// AMDGCN20-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr // AMDGCN20-NEXT: store ptr addrspace(1) [[IN]], ptr [[IN_ADDR_ASCAST]], align 8 // AMDGCN20-NEXT: store ptr addrspace(1) [[OUT]], ptr [[OUT_ADDR_ASCAST]], align 8 // AMDGCN20-NEXT: [[TMP0:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8 @@ -336,7 +334,7 @@ Mat64X64 __attribute__((noinline)) foo_large(Mat32X32 in) { // AMDGCN20-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [[STRUCT_MAT32X32]], ptr addrspace(1) [[TMP1]], i64 1 // AMDGCN20-NEXT: call void @llvm.memcpy.p5.p1.i64(ptr addrspace(5) align 4 [[BYVAL_TEMP]], ptr addrspace(1) align 4 [[ARRAYIDX1]], i64 4096, i1 false) // AMDGCN20-NEXT: call void @foo_large(ptr addrspace(5) dead_on_unwind writable sret([[STRUCT_MAT64X64]]) align 4 [[TMP]], ptr addrspace(5) noundef byref([[STRUCT_MAT32X32]]) align 4 [[BYVAL_TEMP]]) #[[ATTR3]] -// AMDGCN20-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 4 [[ARRAYIDX]], ptr align 4 [[TMP_ASCAST]], i64 16384, i1 false) +// AMDGCN20-NEXT: call void @llvm.memcpy.p1.p5.i64(ptr addrspace(1) align 4 [[ARRAYIDX]], ptr addrspace(5) align 4 [[TMP]], i64 16384, i1 false) // AMDGCN20-NEXT: ret void // // SPIR-LABEL: define dso_local spir_kernel void @ker_large( diff --git a/clang/test/CodeGenOpenCL/amdgpu-abi-struct-arg-byref.cl b/clang/test/CodeGenOpenCL/amdgpu-abi-struct-arg-byref.cl index c2b2e00d15e13..2f8ba99a3e416 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-abi-struct-arg-byref.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-abi-struct-arg-byref.cl @@ -70,7 +70,6 @@ Mat4X4 __attribute__((noinline)) foo(Mat3X3 in) { // AMDGCN-NEXT: [[TMP:%.*]] = alloca [[STRUCT_MAT4X4:%.*]], align 4, addrspace(5) // AMDGCN-NEXT: [[IN_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[IN_ADDR]] to ptr // AMDGCN-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr -// AMDGCN-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr // AMDGCN-NEXT: store ptr addrspace(1) [[IN]], ptr [[IN_ADDR_ASCAST]], align 8 // AMDGCN-NEXT: store ptr addrspace(1) [[OUT]], ptr [[OUT_ADDR_ASCAST]], align 8 // AMDGCN-NEXT: [[TMP0:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8 @@ -80,10 +79,10 @@ Mat4X4 __attribute__((noinline)) foo(Mat3X3 in) { // AMDGCN-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT_MAT3X3]], ptr addrspace(1) [[ARRAYIDX1]], i32 0, i32 0 // AMDGCN-NEXT: [[TMP3:%.*]] = load [9 x i32], ptr addrspace(1) [[TMP2]], align 4 // AMDGCN-NEXT: [[CALL:%.*]] = call [[STRUCT_MAT4X4]] @[[FOO:[a-zA-Z0-9_$\"\\.-]*[a-zA-Z_$\"\\.-][a-zA-Z0-9_$\"\\.-]*]]([9 x i32] [[TMP3]]) #[[ATTR3:[0-9]+]] -// AMDGCN-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT_MAT4X4]], ptr [[TMP_ASCAST]], i32 0, i32 0 +// AMDGCN-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[STRUCT_MAT4X4]], ptr addrspace(5) [[TMP]], i32 0, i32 0 // AMDGCN-NEXT: [[TMP5:%.*]] = extractvalue [[STRUCT_MAT4X4]] [[CALL]], 0 -// AMDGCN-NEXT: store [16 x i32] [[TMP5]], ptr [[TMP4]], align 4 -// AMDGCN-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 4 [[ARRAYIDX]], ptr align 4 [[TMP_ASCAST]], i64 64, i1 false) +// AMDGCN-NEXT: store [16 x i32] [[TMP5]], ptr addrspace(5) [[TMP4]], align 4 +// AMDGCN-NEXT: call void @llvm.memcpy.p1.p5.i64(ptr addrspace(1) align 4 [[ARRAYIDX]], ptr addrspace(5) align 4 [[TMP]], i64 64, i1 false) // AMDGCN-NEXT: ret void // kernel void ker(global Mat3X3 *in, global Mat4X4 *out) { @@ -112,7 +111,6 @@ Mat64X64 __attribute__((noinline)) foo_large(Mat32X32 in) { // AMDGCN-NEXT: [[BYVAL_TEMP:%.*]] = alloca [[STRUCT_MAT32X32:%.*]], align 4, addrspace(5) // AMDGCN-NEXT: [[IN_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[IN_ADDR]] to ptr // AMDGCN-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr -// AMDGCN-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr // AMDGCN-NEXT: store ptr addrspace(1) [[IN]], ptr [[IN_ADDR_ASCAST]], align 8 // AMDGCN-NEXT: store ptr addrspace(1) [[OUT]], ptr [[OUT_ADDR_ASCAST]], align 8 // AMDGCN-NEXT: [[TMP0:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8 @@ -121,7 +119,7 @@ Mat64X64 __attribute__((noinline)) foo_large(Mat32X32 in) { // AMDGCN-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [[STRUCT_MAT32X32]], ptr addrspace(1) [[TMP1]], i64 1 // AMDGCN-NEXT: call void @llvm.memcpy.p5.p1.i64(ptr addrspace(5) align 4 [[BYVAL_TEMP]], ptr addrspace(1) align 4 [[ARRAYIDX1]], i64 4096, i1 false) // AMDGCN-NEXT: call void @foo_large(ptr addrspace(5) dead_on_unwind writable sret([[STRUCT_MAT64X64]]) align 4 [[TMP]], ptr addrspace(5) noundef byref([[STRUCT_MAT32X32]]) align 4 [[BYVAL_TEMP]]) #[[ATTR3]] -// AMDGCN-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 4 [[ARRAYIDX]], ptr align 4 [[TMP_ASCAST]], i64 16384, i1 false) +// AMDGCN-NEXT: call void @llvm.memcpy.p1.p5.i64(ptr addrspace(1) align 4 [[ARRAYIDX]], ptr addrspace(5) align 4 [[TMP]], i64 16384, i1 false) // AMDGCN-NEXT: ret void // kernel void ker_large(global Mat32X32 *in, global Mat64X64 *out) { From c4bdeabaddd3b8fd6fbd269644ce6dfb8bd49739 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Fri, 15 Nov 2024 23:12:40 +0000 Subject: [PATCH 10/27] Fix formatting. --- clang/lib/CodeGen/CGCall.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 41105262b5f14..40c41e0895d66 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -5399,8 +5399,9 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo, if (FirstIRArg < IRFuncTy->getNumParams() && V->getType() != IRFuncTy->getParamType(FirstIRArg)) { assert(V->getType()->isPointerTy() && "Only pointers can mismatch!"); - auto FormalAS = - CallInfo.arguments()[ArgNo].type.getQualifiers().getAddressSpace(); + auto FormalAS = CallInfo.arguments()[ArgNo] + .type.getQualifiers() + .getAddressSpace(); auto ActualAS = I->Ty.getAddressSpace(); V = getTargetHooks().performAddrSpaceCast( *this, V, ActualAS, FormalAS, IRFuncTy->getParamType(FirstIRArg)); From eeb54e4648ec72217399c13451687385c5ad5b16 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Sun, 24 Nov 2024 21:54:00 +0000 Subject: [PATCH 11/27] Remove lie. --- clang/lib/CodeGen/CGCall.cpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 4f2ea9d18fc16..8198669f058d3 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -5163,9 +5163,6 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo, } } if (IRFunctionArgs.hasSRetArg()) { - // If the caller allocated the return slot, it is possible that the - // alloca was AS casted to the default as, so we ensure the cast is - // stripped before binding to the sret arg, which is in the allocaAS. IRCallArgs[IRFunctionArgs.getSRetArgNo()] = getAsNaturalPointerTo(SRetPtr, RetTy); } else if (RetAI.isInAlloca()) { From f16d1d922e038fcf5a9fe8254cafc1bc3804c332 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 5 Dec 2024 03:13:10 +0000 Subject: [PATCH 12/27] Generalise placing `sret`/returns in the alloca AS; remove risky defaulted arg. --- clang/include/clang/CodeGen/CGFunctionInfo.h | 2 +- clang/lib/CodeGen/ABIInfo.cpp | 8 +-- clang/lib/CodeGen/ABIInfo.h | 3 +- clang/lib/CodeGen/ABIInfoImpl.cpp | 19 ++++-- clang/lib/CodeGen/ItaniumCXXABI.cpp | 4 +- clang/lib/CodeGen/MicrosoftCXXABI.cpp | 4 +- clang/lib/CodeGen/SwiftCallingConv.cpp | 10 ++- clang/lib/CodeGen/Targets/AArch64.cpp | 26 +++++--- clang/lib/CodeGen/Targets/AMDGPU.cpp | 10 ++- clang/lib/CodeGen/Targets/ARC.cpp | 12 ++-- clang/lib/CodeGen/Targets/ARM.cpp | 34 ++++++---- clang/lib/CodeGen/Targets/AVR.cpp | 2 +- clang/lib/CodeGen/Targets/BPF.cpp | 12 ++-- clang/lib/CodeGen/Targets/CSKY.cpp | 8 ++- clang/lib/CodeGen/Targets/Hexagon.cpp | 21 +++++-- clang/lib/CodeGen/Targets/Lanai.cpp | 14 +++-- clang/lib/CodeGen/Targets/LoongArch.cpp | 14 +++-- clang/lib/CodeGen/Targets/Mips.cpp | 13 ++-- clang/lib/CodeGen/Targets/NVPTX.cpp | 9 ++- clang/lib/CodeGen/Targets/PNaCl.cpp | 16 +++-- clang/lib/CodeGen/Targets/PPC.cpp | 41 ++++++++---- clang/lib/CodeGen/Targets/RISCV.cpp | 14 +++-- clang/lib/CodeGen/Targets/SPIR.cpp | 9 ++- clang/lib/CodeGen/Targets/Sparc.cpp | 8 ++- clang/lib/CodeGen/Targets/SystemZ.cpp | 18 ++++-- clang/lib/CodeGen/Targets/WebAssembly.cpp | 4 +- clang/lib/CodeGen/Targets/X86.cpp | 66 ++++++++++++++------ 27 files changed, 276 insertions(+), 125 deletions(-) diff --git a/clang/include/clang/CodeGen/CGFunctionInfo.h b/clang/include/clang/CodeGen/CGFunctionInfo.h index 4ca5d2b654812..040ee025afaa8 100644 --- a/clang/include/clang/CodeGen/CGFunctionInfo.h +++ b/clang/include/clang/CodeGen/CGFunctionInfo.h @@ -206,7 +206,7 @@ class ABIArgInfo { static ABIArgInfo getIgnore() { return ABIArgInfo(Ignore); } - static ABIArgInfo getIndirect(CharUnits Alignment, unsigned AddrSpace = 0, + static ABIArgInfo getIndirect(CharUnits Alignment, unsigned AddrSpace, bool ByVal = true, bool Realign = false, llvm::Type *Padding = nullptr) { auto AI = ABIArgInfo(Indirect); diff --git a/clang/lib/CodeGen/ABIInfo.cpp b/clang/lib/CodeGen/ABIInfo.cpp index 2d6280f8fc508..19d8c66b000cb 100644 --- a/clang/lib/CodeGen/ABIInfo.cpp +++ b/clang/lib/CodeGen/ABIInfo.cpp @@ -171,11 +171,11 @@ bool ABIInfo::isPromotableIntegerTypeForABI(QualType Ty) const { return false; } -ABIArgInfo ABIInfo::getNaturalAlignIndirect(QualType Ty, bool ByVal, - bool Realign, +ABIArgInfo ABIInfo::getNaturalAlignIndirect(QualType Ty, unsigned AddrSpace, + bool ByVal, bool Realign, llvm::Type *Padding) const { - return ABIArgInfo::getIndirect(getContext().getTypeAlignInChars(Ty), 0, ByVal, - Realign, Padding); + return ABIArgInfo::getIndirect(getContext().getTypeAlignInChars(Ty), + AddrSpace, ByVal, Realign, Padding); } ABIArgInfo ABIInfo::getNaturalAlignIndirectInReg(QualType Ty, diff --git a/clang/lib/CodeGen/ABIInfo.h b/clang/lib/CodeGen/ABIInfo.h index b8a8de57e5b97..57bcb3dd0a852 100644 --- a/clang/lib/CodeGen/ABIInfo.h +++ b/clang/lib/CodeGen/ABIInfo.h @@ -109,7 +109,8 @@ class ABIInfo { /// A convenience method to return an indirect ABIArgInfo with an /// expected alignment equal to the ABI alignment of the given type. CodeGen::ABIArgInfo - getNaturalAlignIndirect(QualType Ty, bool ByVal = true, bool Realign = false, + getNaturalAlignIndirect(QualType Ty, unsigned AddrSpace, bool ByVal = true, + bool Realign = false, llvm::Type *Padding = nullptr) const; CodeGen::ABIArgInfo getNaturalAlignIndirectInReg(QualType Ty, diff --git a/clang/lib/CodeGen/ABIInfoImpl.cpp b/clang/lib/CodeGen/ABIInfoImpl.cpp index 79300df15d0e2..d55fcc9e125f0 100644 --- a/clang/lib/CodeGen/ABIInfoImpl.cpp +++ b/clang/lib/CodeGen/ABIInfoImpl.cpp @@ -21,9 +21,13 @@ ABIArgInfo DefaultABIInfo::classifyArgumentType(QualType Ty) const { // Records with non-trivial destructors/copy-constructors should not be // passed by value. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, + getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); - return getNaturalAlignIndirect(Ty); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default)); } // Treat an enum type as its underlying type. @@ -36,7 +40,8 @@ ABIArgInfo DefaultABIInfo::classifyArgumentType(QualType Ty) const { Context.getTypeSize(Context.getTargetInfo().hasInt128Type() ? Context.Int128Ty : Context.LongLongTy)) - return getNaturalAlignIndirect(Ty); + return getNaturalAlignIndirect( + Ty, Context.getTargetAddressSpace(LangAS::Default)); return (isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty, CGT.ConvertType(Ty)) @@ -48,7 +53,7 @@ ABIArgInfo DefaultABIInfo::classifyReturnType(QualType RetTy) const { return ABIArgInfo::getIgnore(); if (isAggregateTypeForABI(RetTy)) - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect(RetTy, getDataLayout().getAllocaAddrSpace()); // Treat an enum type as its underlying type. if (const EnumType *EnumTy = RetTy->getAs()) @@ -59,7 +64,8 @@ ABIArgInfo DefaultABIInfo::classifyReturnType(QualType RetTy) const { getContext().getTypeSize(getContext().getTargetInfo().hasInt128Type() ? getContext().Int128Ty : getContext().LongLongTy)) - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect( + RetTy, getDataLayout().getAllocaAddrSpace()); return (isPromotableIntegerTypeForABI(RetTy) ? ABIArgInfo::getExtend(RetTy) : ABIArgInfo::getDirect()); @@ -126,7 +132,8 @@ bool CodeGen::classifyReturnType(const CGCXXABI &CXXABI, CGFunctionInfo &FI, if (const auto *RT = Ty->getAs()) if (!isa(RT->getDecl()) && !RT->getDecl()->canPassInRegisters()) { - FI.getReturnInfo() = Info.getNaturalAlignIndirect(Ty); + FI.getReturnInfo() = Info.getNaturalAlignIndirect( + Ty, Info.getDataLayout().getAllocaAddrSpace()); return true; } diff --git a/clang/lib/CodeGen/ItaniumCXXABI.cpp b/clang/lib/CodeGen/ItaniumCXXABI.cpp index 0687020015349..c0e11d2269e18 100644 --- a/clang/lib/CodeGen/ItaniumCXXABI.cpp +++ b/clang/lib/CodeGen/ItaniumCXXABI.cpp @@ -1350,7 +1350,9 @@ bool ItaniumCXXABI::classifyReturnType(CGFunctionInfo &FI) const { // If C++ prohibits us from making a copy, return by address. if (!RD->canPassInRegisters()) { auto Align = CGM.getContext().getTypeAlignInChars(FI.getReturnType()); - FI.getReturnInfo() = ABIArgInfo::getIndirect(Align, 0, /*ByVal=*/false); + FI.getReturnInfo() = ABIArgInfo::getIndirect( + Align, /*AddrSpace=*/CGM.getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); return true; } return false; diff --git a/clang/lib/CodeGen/MicrosoftCXXABI.cpp b/clang/lib/CodeGen/MicrosoftCXXABI.cpp index b6eb02a394f32..bc5bb24a00344 100644 --- a/clang/lib/CodeGen/MicrosoftCXXABI.cpp +++ b/clang/lib/CodeGen/MicrosoftCXXABI.cpp @@ -1172,7 +1172,9 @@ bool MicrosoftCXXABI::classifyReturnType(CGFunctionInfo &FI) const { if (isIndirectReturn) { CharUnits Align = CGM.getContext().getTypeAlignInChars(FI.getReturnType()); - FI.getReturnInfo() = ABIArgInfo::getIndirect(Align, 0, /*ByVal=*/false); + FI.getReturnInfo() = ABIArgInfo::getIndirect( + Align, /*AddrSpace=*/CGM.getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); // MSVC always passes `this` before the `sret` parameter. FI.getReturnInfo().setSRetAfterThis(FI.isInstanceMethod()); diff --git a/clang/lib/CodeGen/SwiftCallingConv.cpp b/clang/lib/CodeGen/SwiftCallingConv.cpp index ef619e9613d83..01a6402b42c6a 100644 --- a/clang/lib/CodeGen/SwiftCallingConv.cpp +++ b/clang/lib/CodeGen/SwiftCallingConv.cpp @@ -800,7 +800,10 @@ static ABIArgInfo classifyExpandedType(SwiftAggLowering &lowering, if (lowering.empty()) { return ABIArgInfo::getIgnore(); } else if (lowering.shouldPassIndirectly(forReturn)) { - return ABIArgInfo::getIndirect(alignmentForIndirect, 0, /*byval*/ false); + return ABIArgInfo::getIndirect( + alignmentForIndirect, + /*AddrSpace*/ 0, + /*byval*/ false); } else { auto types = lowering.getCoerceAndExpandTypes(); return ABIArgInfo::getCoerceAndExpand(types.first, types.second); @@ -814,7 +817,10 @@ static ABIArgInfo classifyType(CodeGenModule &CGM, CanQualType type, auto &layout = CGM.getContext().getASTRecordLayout(record); if (mustPassRecordIndirectly(CGM, record)) - return ABIArgInfo::getIndirect(layout.getAlignment(), 0, /*byval*/ false); + return ABIArgInfo::getIndirect( + layout.getAlignment(), + /*AddrSpace*/ CGM.getContext().getTargetAddressSpace(LangAS::Default), + /*byval*/ false); SwiftAggLowering lowering(CGM); lowering.addTypedData(recordType->getDecl(), CharUnits::Zero(), layout); diff --git a/clang/lib/CodeGen/Targets/AArch64.cpp b/clang/lib/CodeGen/Targets/AArch64.cpp index be33e26f04784..c224b29a243a8 100644 --- a/clang/lib/CodeGen/Targets/AArch64.cpp +++ b/clang/lib/CodeGen/Targets/AArch64.cpp @@ -325,7 +325,8 @@ ABIArgInfo AArch64ABIInfo::coerceIllegalVector(QualType Ty, unsigned &NSRN, return ABIArgInfo::getDirect(ResType); } - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), /*ByVal=*/false); } ABIArgInfo AArch64ABIInfo::coerceAndExpandPureScalableAggregate( @@ -333,7 +334,9 @@ ABIArgInfo AArch64ABIInfo::coerceAndExpandPureScalableAggregate( const SmallVectorImpl &UnpaddedCoerceToSeq, unsigned &NSRN, unsigned &NPRN) const { if (!IsNamedArg || NSRN + NVec > 8 || NPRN + NPred > 4) - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); NSRN += NVec; NPRN += NPred; @@ -369,7 +372,8 @@ ABIArgInfo AArch64ABIInfo::classifyArgumentType(QualType Ty, bool IsVariadicFn, if (const auto *EIT = Ty->getAs()) if (EIT->getNumBits() > 128) - return getNaturalAlignIndirect(Ty, false); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), false); if (Ty->isVectorType()) NSRN = std::min(NSRN + 1, 8u); @@ -409,8 +413,9 @@ ABIArgInfo AArch64ABIInfo::classifyArgumentType(QualType Ty, bool IsVariadicFn, // Structures with either a non-trivial destructor or a non-trivial // copy constructor are always indirect. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { - return getNaturalAlignIndirect(Ty, /*ByVal=*/RAA == - CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } // Empty records are always ignored on Darwin, but actually passed in C++ mode @@ -484,7 +489,9 @@ ABIArgInfo AArch64ABIInfo::classifyArgumentType(QualType Ty, bool IsVariadicFn, : llvm::ArrayType::get(BaseTy, Size / Alignment)); } - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, /*AddrSpace*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); } ABIArgInfo AArch64ABIInfo::classifyReturnType(QualType RetTy, @@ -502,7 +509,7 @@ ABIArgInfo AArch64ABIInfo::classifyReturnType(QualType RetTy, // Large vector types should be returned via memory. if (RetTy->isVectorType() && getContext().getTypeSize(RetTy) > 128) - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect(RetTy, getDataLayout().getAllocaAddrSpace()); if (!isAggregateTypeForABI(RetTy)) { // Treat an enum type as its underlying type. @@ -511,7 +518,8 @@ ABIArgInfo AArch64ABIInfo::classifyReturnType(QualType RetTy, if (const auto *EIT = RetTy->getAs()) if (EIT->getNumBits() > 128) - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect(RetTy, + getDataLayout().getAllocaAddrSpace()); return (isPromotableIntegerTypeForABI(RetTy) && isDarwinPCS() ? ABIArgInfo::getExtend(RetTy) @@ -569,7 +577,7 @@ ABIArgInfo AArch64ABIInfo::classifyReturnType(QualType RetTy, return ABIArgInfo::getDirect(llvm::IntegerType::get(getVMContext(), Size)); } - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect(RetTy, getDataLayout().getAllocaAddrSpace()); } /// isIllegalVectorType - check whether the vector type is legal for AArch64. diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index c45e7020de3f5..b180b1b8fa00c 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -105,11 +105,6 @@ void AMDGPUABIInfo::computeInfo(CGFunctionInfo &FI) const { if (!getCXXABI().classifyReturnType(FI)) FI.getReturnInfo() = classifyReturnType(FI.getReturnType()); - // srets / indirect returns are unconditionally in the alloca AS. - if (FI.getReturnInfo().isIndirect()) - FI.getReturnInfo().setIndirectAddrSpace( - getDataLayout().getAllocaAddrSpace()); - unsigned ArgumentIndex = 0; const unsigned numFixedArguments = FI.getNumRequiredArgs(); @@ -230,7 +225,10 @@ ABIArgInfo AMDGPUABIInfo::classifyArgumentType(QualType Ty, bool Variadic, // Records with non-trivial destructors/copy-constructors should not be // passed by value. if (auto RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, + getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); // Ignore empty structs/unions. if (isEmptyRecord(getContext(), Ty, true)) diff --git a/clang/lib/CodeGen/Targets/ARC.cpp b/clang/lib/CodeGen/Targets/ARC.cpp index ee0db9778bdcb..a73b668c30ae1 100644 --- a/clang/lib/CodeGen/Targets/ARC.cpp +++ b/clang/lib/CodeGen/Targets/ARC.cpp @@ -69,16 +69,20 @@ class ARCTargetCodeGenInfo : public TargetCodeGenInfo { ABIArgInfo ARCABIInfo::getIndirectByRef(QualType Ty, bool HasFreeRegs) const { - return HasFreeRegs ? getNaturalAlignIndirectInReg(Ty) : - getNaturalAlignIndirect(Ty, false); + return HasFreeRegs + ? getNaturalAlignIndirectInReg(Ty) + : getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), false); } ABIArgInfo ARCABIInfo::getIndirectByValue(QualType Ty) const { // Compute the byval alignment. const unsigned MinABIStackAlignInBytes = 4; unsigned TypeAlign = getContext().getTypeAlign(Ty) / 8; - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(4), 0, /*ByVal=*/true, - TypeAlign > MinABIStackAlignInBytes); + return ABIArgInfo::getIndirect( + CharUnits::fromQuantity(4), + /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/true, TypeAlign > MinABIStackAlignInBytes); } RValue ARCABIInfo::EmitVAArg(CodeGenFunction &CGF, Address VAListAddr, diff --git a/clang/lib/CodeGen/Targets/ARM.cpp b/clang/lib/CodeGen/Targets/ARM.cpp index d89a0bdff56a3..c663d02ebb88b 100644 --- a/clang/lib/CodeGen/Targets/ARM.cpp +++ b/clang/lib/CodeGen/Targets/ARM.cpp @@ -298,7 +298,9 @@ ABIArgInfo ARMABIInfo::coerceIllegalVector(QualType Ty) const { llvm::Type::getInt32Ty(getVMContext()), Size / 32); return ABIArgInfo::getDirect(ResType); } - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); } ABIArgInfo ARMABIInfo::classifyHomogeneousAggregate(QualType Ty, @@ -355,7 +357,10 @@ ABIArgInfo ARMABIInfo::classifyArgumentType(QualType Ty, bool isVariadic, if (const auto *EIT = Ty->getAs()) if (EIT->getNumBits() > 64) - return getNaturalAlignIndirect(Ty, /*ByVal=*/true); + return getNaturalAlignIndirect( + Ty, + /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/true); return (isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty, CGT.ConvertType(Ty)) @@ -363,7 +368,9 @@ ABIArgInfo ARMABIInfo::classifyArgumentType(QualType Ty, bool isVariadic, } if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { - return getNaturalAlignIndirect(Ty, RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); } // Ignore empty records. @@ -397,7 +404,8 @@ ABIArgInfo ARMABIInfo::classifyArgumentType(QualType Ty, bool isVariadic, // bigger than 128-bits, they get placed in space allocated by the caller, // and a pointer is passed. return ABIArgInfo::getIndirect( - CharUnits::fromQuantity(getContext().getTypeAlign(Ty) / 8), 0, false); + CharUnits::fromQuantity(getContext().getTypeAlign(Ty) / 8), + getContext().getTargetAddressSpace(LangAS::Default), false); } // Support byval for ARM. @@ -415,9 +423,10 @@ ABIArgInfo ARMABIInfo::classifyArgumentType(QualType Ty, bool isVariadic, } if (getContext().getTypeSizeInChars(Ty) > CharUnits::fromQuantity(64)) { assert(getABIKind() != ARMABIKind::AAPCS16_VFP && "unexpected byval"); - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(ABIAlign), 0, - /*ByVal=*/true, - /*Realign=*/TyAlign > ABIAlign); + return ABIArgInfo::getIndirect( + CharUnits::fromQuantity(ABIAlign), + /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/true, /*Realign=*/TyAlign > ABIAlign); } // Otherwise, pass by coercing to a structure of the appropriate size. @@ -534,7 +543,8 @@ ABIArgInfo ARMABIInfo::classifyReturnType(QualType RetTy, bool isVariadic, if (const VectorType *VT = RetTy->getAs()) { // Large vector types should be returned via memory. if (getContext().getTypeSize(RetTy) > 128) - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect( + RetTy, getDataLayout().getAllocaAddrSpace()); // TODO: FP16/BF16 vectors should be converted to integer vectors // This check is similar to isIllegalVectorType - refactor? if ((!getTarget().hasLegalHalfType() && @@ -552,7 +562,9 @@ ABIArgInfo ARMABIInfo::classifyReturnType(QualType RetTy, bool isVariadic, if (const auto *EIT = RetTy->getAs()) if (EIT->getNumBits() > 64) - return getNaturalAlignIndirect(RetTy, /*ByVal=*/false); + return getNaturalAlignIndirect( + RetTy, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); return isPromotableIntegerTypeForABI(RetTy) ? ABIArgInfo::getExtend(RetTy) : ABIArgInfo::getDirect(); @@ -583,7 +595,7 @@ ABIArgInfo ARMABIInfo::classifyReturnType(QualType RetTy, bool isVariadic, } // Otherwise return in memory. - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect(RetTy, getDataLayout().getAllocaAddrSpace()); } // Otherwise this is an AAPCS variant. @@ -620,7 +632,7 @@ ABIArgInfo ARMABIInfo::classifyReturnType(QualType RetTy, bool isVariadic, return ABIArgInfo::getDirect(CoerceTy); } - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect(RetTy, getDataLayout().getAllocaAddrSpace()); } /// isIllegalVector - check whether Ty is an illegal vector type. diff --git a/clang/lib/CodeGen/Targets/AVR.cpp b/clang/lib/CodeGen/Targets/AVR.cpp index 50547dd6dec5e..26e2a22f14d1e 100644 --- a/clang/lib/CodeGen/Targets/AVR.cpp +++ b/clang/lib/CodeGen/Targets/AVR.cpp @@ -45,7 +45,7 @@ class AVRABIInfo : public DefaultABIInfo { // stack slot, along with a pointer as the function's implicit argument. if (getContext().getTypeSize(Ty) > RetRegs * 8) { LargeRet = true; - return getNaturalAlignIndirect(Ty); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace()); } // An i8 return value should not be extended to i16, since AVR has 8-bit // registers. diff --git a/clang/lib/CodeGen/Targets/BPF.cpp b/clang/lib/CodeGen/Targets/BPF.cpp index 2849222f7a186..612c61d83a00a 100644 --- a/clang/lib/CodeGen/Targets/BPF.cpp +++ b/clang/lib/CodeGen/Targets/BPF.cpp @@ -42,7 +42,8 @@ class BPFABIInfo : public DefaultABIInfo { } return ABIArgInfo::getDirect(CoerceTy); } else { - return getNaturalAlignIndirect(Ty); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default)); } } @@ -52,7 +53,8 @@ class BPFABIInfo : public DefaultABIInfo { ASTContext &Context = getContext(); if (const auto *EIT = Ty->getAs()) if (EIT->getNumBits() > Context.getTypeSize(Context.Int128Ty)) - return getNaturalAlignIndirect(Ty); + return getNaturalAlignIndirect( + Ty, Context.getTargetAddressSpace(LangAS::Default)); return (isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty) : ABIArgInfo::getDirect()); @@ -63,7 +65,8 @@ class BPFABIInfo : public DefaultABIInfo { return ABIArgInfo::getIgnore(); if (isAggregateTypeForABI(RetTy)) - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect( + RetTy, getDataLayout().getAllocaAddrSpace()); // Treat an enum type as its underlying type. if (const EnumType *EnumTy = RetTy->getAs()) @@ -72,7 +75,8 @@ class BPFABIInfo : public DefaultABIInfo { ASTContext &Context = getContext(); if (const auto *EIT = RetTy->getAs()) if (EIT->getNumBits() > Context.getTypeSize(Context.Int128Ty)) - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect( + RetTy, getDataLayout().getAllocaAddrSpace()); // Caller will do necessary sign/zero extension. return ABIArgInfo::getDirect(); diff --git a/clang/lib/CodeGen/Targets/CSKY.cpp b/clang/lib/CodeGen/Targets/CSKY.cpp index d8720afd1a713..41dd921ae309e 100644 --- a/clang/lib/CodeGen/Targets/CSKY.cpp +++ b/clang/lib/CodeGen/Targets/CSKY.cpp @@ -82,8 +82,9 @@ ABIArgInfo CSKYABIInfo::classifyArgumentType(QualType Ty, int &ArgGPRsLeft, if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { if (ArgGPRsLeft) ArgGPRsLeft -= 1; - return getNaturalAlignIndirect(Ty, /*ByVal=*/RAA == - CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } // Ignore empty structs/unions. @@ -144,7 +145,8 @@ ABIArgInfo CSKYABIInfo::classifyArgumentType(QualType Ty, int &ArgGPRsLeft, llvm::IntegerType::get(getVMContext(), XLen), (Size + 31) / XLen)); } } - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), /*ByVal=*/false); } ABIArgInfo CSKYABIInfo::classifyReturnType(QualType RetTy) const { diff --git a/clang/lib/CodeGen/Targets/Hexagon.cpp b/clang/lib/CodeGen/Targets/Hexagon.cpp index 8fd2a81494d99..151e269a10d38 100644 --- a/clang/lib/CodeGen/Targets/Hexagon.cpp +++ b/clang/lib/CodeGen/Targets/Hexagon.cpp @@ -105,14 +105,18 @@ ABIArgInfo HexagonABIInfo::classifyArgumentType(QualType Ty, HexagonAdjustRegsLeft(Size, RegsLeft); if (Size > 64 && Ty->isBitIntType()) - return getNaturalAlignIndirect(Ty, /*ByVal=*/true); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/true); return isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty) : ABIArgInfo::getDirect(); } if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); // Ignore empty records. if (isEmptyRecord(getContext(), Ty, true)) @@ -122,7 +126,9 @@ ABIArgInfo HexagonABIInfo::classifyArgumentType(QualType Ty, unsigned Align = getContext().getTypeAlign(Ty); if (Size > 64) - return getNaturalAlignIndirect(Ty, /*ByVal=*/true); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/true); if (HexagonAdjustRegsLeft(Size, RegsLeft)) Align = Size <= 32 ? 32 : 64; @@ -151,7 +157,8 @@ ABIArgInfo HexagonABIInfo::classifyReturnType(QualType RetTy) const { } // Large vector types should be returned via memory. if (Size > 64) - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect( + RetTy, getDataLayout().getAllocaAddrSpace()); } if (!isAggregateTypeForABI(RetTy)) { @@ -160,7 +167,8 @@ ABIArgInfo HexagonABIInfo::classifyReturnType(QualType RetTy) const { RetTy = EnumTy->getDecl()->getIntegerType(); if (Size > 64 && RetTy->isBitIntType()) - return getNaturalAlignIndirect(RetTy, /*ByVal=*/false); + return getNaturalAlignIndirect( + RetTy, getDataLayout().getAllocaAddrSpace(), /*ByVal=*/false); return isPromotableIntegerTypeForABI(RetTy) ? ABIArgInfo::getExtend(RetTy) : ABIArgInfo::getDirect(); @@ -176,7 +184,8 @@ ABIArgInfo HexagonABIInfo::classifyReturnType(QualType RetTy) const { Size = llvm::bit_ceil(Size); return ABIArgInfo::getDirect(llvm::Type::getIntNTy(getVMContext(), Size)); } - return getNaturalAlignIndirect(RetTy, /*ByVal=*/true); + return getNaturalAlignIndirect( + RetTy, getDataLayout().getAllocaAddrSpace(), /*ByVal=*/true); } Address HexagonABIInfo::EmitVAArgFromMemory(CodeGenFunction &CGF, diff --git a/clang/lib/CodeGen/Targets/Lanai.cpp b/clang/lib/CodeGen/Targets/Lanai.cpp index ffacb0ccbea53..9b40655fb5289 100644 --- a/clang/lib/CodeGen/Targets/Lanai.cpp +++ b/clang/lib/CodeGen/Targets/Lanai.cpp @@ -72,15 +72,17 @@ ABIArgInfo LanaiABIInfo::getIndirectResult(QualType Ty, bool ByVal, --State.FreeRegs; // Non-byval indirects just use one pointer. return getNaturalAlignIndirectInReg(Ty); } - return getNaturalAlignIndirect(Ty, false); + return getNaturalAlignIndirect( + Ty, getDataLayout().getAllocaAddrSpace(), false); } // Compute the byval alignment. const unsigned MinABIStackAlignInBytes = 4; unsigned TypeAlign = getContext().getTypeAlign(Ty) / 8; - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(4), 0, /*ByVal=*/true, - /*Realign=*/TypeAlign > - MinABIStackAlignInBytes); + return ABIArgInfo::getIndirect( + CharUnits::fromQuantity(4), + /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), /*ByVal=*/true, + /*Realign=*/TypeAlign > MinABIStackAlignInBytes); } ABIArgInfo LanaiABIInfo::classifyArgumentType(QualType Ty, @@ -92,7 +94,9 @@ ABIArgInfo LanaiABIInfo::classifyArgumentType(QualType Ty, if (RAA == CGCXXABI::RAA_Indirect) { return getIndirectResult(Ty, /*ByVal=*/false, State); } else if (RAA == CGCXXABI::RAA_DirectInMemory) { - return getNaturalAlignIndirect(Ty, /*ByVal=*/true); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/true); } } diff --git a/clang/lib/CodeGen/Targets/LoongArch.cpp b/clang/lib/CodeGen/Targets/LoongArch.cpp index 6af9375461f09..b89130f93720d 100644 --- a/clang/lib/CodeGen/Targets/LoongArch.cpp +++ b/clang/lib/CodeGen/Targets/LoongArch.cpp @@ -305,8 +305,10 @@ ABIArgInfo LoongArchABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { if (GARsLeft) GARsLeft -= 1; - return getNaturalAlignIndirect(Ty, /*ByVal=*/RAA == - CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, + /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } uint64_t Size = getContext().getTypeSize(Ty); @@ -381,7 +383,9 @@ ABIArgInfo LoongArchABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (EIT->getNumBits() > 128 || (!getContext().getTargetInfo().hasInt128Type() && EIT->getNumBits() > 64)) - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); } return ABIArgInfo::getDirect(); @@ -404,7 +408,9 @@ ABIArgInfo LoongArchABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, return ABIArgInfo::getDirect( llvm::ArrayType::get(llvm::IntegerType::get(getVMContext(), GRLen), 2)); } - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty,/*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); } ABIArgInfo LoongArchABIInfo::classifyReturnType(QualType RetTy) const { diff --git a/clang/lib/CodeGen/Targets/Mips.cpp b/clang/lib/CodeGen/Targets/Mips.cpp index 06d9b6d4a5761..067ffa2c2cd81 100644 --- a/clang/lib/CodeGen/Targets/Mips.cpp +++ b/clang/lib/CodeGen/Targets/Mips.cpp @@ -209,7 +209,10 @@ MipsABIInfo::classifyArgumentType(QualType Ty, uint64_t &Offset) const { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { Offset = OrigOffset + MinABIStackAlignInBytes; - return getNaturalAlignIndirect(Ty, RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, + getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); } // If we have reached here, aggregates are passed directly by coercing to @@ -231,7 +234,8 @@ MipsABIInfo::classifyArgumentType(QualType Ty, uint64_t &Offset) const { if (EIT->getNumBits() > 128 || (EIT->getNumBits() > 64 && !getContext().getTargetInfo().hasInt128Type())) - return getNaturalAlignIndirect(Ty); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default)); // All integral types are promoted to the GPR width. if (Ty->isIntegralOrEnumerationType()) @@ -310,7 +314,7 @@ ABIArgInfo MipsABIInfo::classifyReturnType(QualType RetTy) const { } } - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect(RetTy, getDataLayout().getAllocaAddrSpace()); } // Treat an enum type as its underlying type. @@ -322,7 +326,8 @@ ABIArgInfo MipsABIInfo::classifyReturnType(QualType RetTy) const { if (EIT->getNumBits() > 128 || (EIT->getNumBits() > 64 && !getContext().getTargetInfo().hasInt128Type())) - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect( + RetTy, getDataLayout().getAllocaAddrSpace()); if (isPromotableIntegerTypeForABI(RetTy)) return ABIArgInfo::getExtend(RetTy); diff --git a/clang/lib/CodeGen/Targets/NVPTX.cpp b/clang/lib/CodeGen/Targets/NVPTX.cpp index 0431d2cc4ddc3..ce59f63fcf6f1 100644 --- a/clang/lib/CodeGen/Targets/NVPTX.cpp +++ b/clang/lib/CodeGen/Targets/NVPTX.cpp @@ -193,14 +193,19 @@ ABIArgInfo NVPTXABIInfo::classifyArgumentType(QualType Ty) const { return ABIArgInfo::getDirect( CGInfo.getCUDADeviceBuiltinTextureDeviceType()); } - return getNaturalAlignIndirect(Ty, /* byval */ true); + return getNaturalAlignIndirect( + Ty, /* AddrSpace */ getContext().getTargetAddressSpace(LangAS::Default), + /* byval */ true); } if (const auto *EIT = Ty->getAs()) { if ((EIT->getNumBits() > 128) || (!getContext().getTargetInfo().hasInt128Type() && EIT->getNumBits() > 64)) - return getNaturalAlignIndirect(Ty, /* byval */ true); + return getNaturalAlignIndirect( + Ty, + /* AddrSpace */ getContext().getTargetAddressSpace(LangAS::Default), + /* byval */ true); } return (isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty) diff --git a/clang/lib/CodeGen/Targets/PNaCl.cpp b/clang/lib/CodeGen/Targets/PNaCl.cpp index 9b7d757df3a39..ac0bc598ff84a 100644 --- a/clang/lib/CodeGen/Targets/PNaCl.cpp +++ b/clang/lib/CodeGen/Targets/PNaCl.cpp @@ -63,8 +63,11 @@ RValue PNaClABIInfo::EmitVAArg(CodeGenFunction &CGF, Address VAListAddr, ABIArgInfo PNaClABIInfo::classifyArgumentType(QualType Ty) const { if (isAggregateTypeForABI(Ty)) { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, RAA == CGCXXABI::RAA_DirectInMemory); - return getNaturalAlignIndirect(Ty); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default)); } else if (const EnumType *EnumTy = Ty->getAs()) { // Treat an enum type as its underlying type. Ty = EnumTy->getDecl()->getIntegerType(); @@ -75,7 +78,8 @@ ABIArgInfo PNaClABIInfo::classifyArgumentType(QualType Ty) const { // Treat bit-precise integers as integers if <= 64, otherwise pass // indirectly. if (EIT->getNumBits() > 64) - return getNaturalAlignIndirect(Ty); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default)); return ABIArgInfo::getDirect(); } @@ -89,12 +93,14 @@ ABIArgInfo PNaClABIInfo::classifyReturnType(QualType RetTy) const { // In the PNaCl ABI we always return records/structures on the stack. if (isAggregateTypeForABI(RetTy)) - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect( + RetTy, getDataLayout().getAllocaAddrSpace()); // Treat bit-precise integers as integers if <= 64, otherwise pass indirectly. if (const auto *EIT = RetTy->getAs()) { if (EIT->getNumBits() > 64) - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect( + RetTy, getDataLayout().getAllocaAddrSpace()); return ABIArgInfo::getDirect(); } diff --git a/clang/lib/CodeGen/Targets/PPC.cpp b/clang/lib/CodeGen/Targets/PPC.cpp index c8796036b214f..5204bfddf1ab1 100644 --- a/clang/lib/CodeGen/Targets/PPC.cpp +++ b/clang/lib/CodeGen/Targets/PPC.cpp @@ -189,7 +189,7 @@ ABIArgInfo AIXABIInfo::classifyReturnType(QualType RetTy) const { return ABIArgInfo::getIgnore(); if (isAggregateTypeForABI(RetTy)) - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect(RetTy, getDataLayout().getAllocaAddrSpace()); return (isPromotableTypeForABI(RetTy) ? ABIArgInfo::getExtend(RetTy) : ABIArgInfo::getDirect()); @@ -208,13 +208,18 @@ ABIArgInfo AIXABIInfo::classifyArgumentType(QualType Ty) const { // Records with non-trivial destructors/copy-constructors should not be // passed by value. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, + getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); CharUnits CCAlign = getParamTypeAlignment(Ty); CharUnits TyAlign = getContext().getTypeAlignInChars(Ty); - return ABIArgInfo::getIndirect(CCAlign, 0, /*ByVal*/ true, - /*Realign*/ TyAlign > CCAlign); + return ABIArgInfo::getIndirect( + CCAlign, + /*AddrSpace*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal*/ true, /*Realign*/ TyAlign > CCAlign); } return (isPromotableTypeForABI(Ty) @@ -833,7 +838,9 @@ PPC64_SVR4_ABIInfo::classifyArgumentType(QualType Ty) const { if (Ty->isVectorType()) { uint64_t Size = getContext().getTypeSize(Ty); if (Size > 128) - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); else if (Size < 128) { llvm::Type *CoerceTy = llvm::IntegerType::get(getVMContext(), Size); return ABIArgInfo::getDirect(CoerceTy); @@ -842,11 +849,16 @@ PPC64_SVR4_ABIInfo::classifyArgumentType(QualType Ty) const { if (const auto *EIT = Ty->getAs()) if (EIT->getNumBits() > 128) - return getNaturalAlignIndirect(Ty, /*ByVal=*/true); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/true); if (isAggregateTypeForABI(Ty)) { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, + getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); uint64_t ABIAlign = getParamTypeAlignment(Ty).getQuantity(); uint64_t TyAlign = getContext().getTypeAlignInChars(Ty).getQuantity(); @@ -887,9 +899,10 @@ PPC64_SVR4_ABIInfo::classifyArgumentType(QualType Ty) const { } // All other aggregates are passed ByVal. - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(ABIAlign), 0, - /*ByVal=*/true, - /*Realign=*/TyAlign > ABIAlign); + return ABIArgInfo::getIndirect( + CharUnits::fromQuantity(ABIAlign), + /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/true, /*Realign=*/TyAlign > ABIAlign); } return (isPromotableTypeForABI(Ty) @@ -910,7 +923,8 @@ PPC64_SVR4_ABIInfo::classifyReturnType(QualType RetTy) const { if (RetTy->isVectorType()) { uint64_t Size = getContext().getTypeSize(RetTy); if (Size > 128) - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect(RetTy, + getDataLayout().getAllocaAddrSpace()); else if (Size < 128) { llvm::Type *CoerceTy = llvm::IntegerType::get(getVMContext(), Size); return ABIArgInfo::getDirect(CoerceTy); @@ -919,7 +933,8 @@ PPC64_SVR4_ABIInfo::classifyReturnType(QualType RetTy) const { if (const auto *EIT = RetTy->getAs()) if (EIT->getNumBits() > 128) - return getNaturalAlignIndirect(RetTy, /*ByVal=*/false); + return getNaturalAlignIndirect( + RetTy, getDataLayout().getAllocaAddrSpace(), /*ByVal=*/false); if (isAggregateTypeForABI(RetTy)) { // ELFv2 homogeneous aggregates are returned as array types. @@ -949,7 +964,7 @@ PPC64_SVR4_ABIInfo::classifyReturnType(QualType RetTy) const { } // All other aggregates are returned indirectly. - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect(RetTy, getDataLayout().getAllocaAddrSpace()); } return (isPromotableTypeForABI(RetTy) ? ABIArgInfo::getExtend(RetTy) diff --git a/clang/lib/CodeGen/Targets/RISCV.cpp b/clang/lib/CodeGen/Targets/RISCV.cpp index b04e436c665f5..dc55924f747f4 100644 --- a/clang/lib/CodeGen/Targets/RISCV.cpp +++ b/clang/lib/CodeGen/Targets/RISCV.cpp @@ -410,8 +410,9 @@ ABIArgInfo RISCVABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { if (ArgGPRsLeft) ArgGPRsLeft -= 1; - return getNaturalAlignIndirect(Ty, /*ByVal=*/RAA == - CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } uint64_t Size = getContext().getTypeSize(Ty); @@ -492,7 +493,10 @@ ABIArgInfo RISCVABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (EIT->getNumBits() > 128 || (!getContext().getTargetInfo().hasInt128Type() && EIT->getNumBits() > 64)) - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, + /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); } ABIArgInfo Info = ABIArgInfo::getDirect(); @@ -530,7 +534,9 @@ ABIArgInfo RISCVABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, llvm::IntegerType::get(getVMContext(), XLen), 2)); } } - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); } ABIArgInfo RISCVABIInfo::classifyReturnType(QualType RetTy) const { diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp index a48fe9d5f1ee9..920d2151b622d 100644 --- a/clang/lib/CodeGen/Targets/SPIR.cpp +++ b/clang/lib/CodeGen/Targets/SPIR.cpp @@ -155,7 +155,9 @@ ABIArgInfo SPIRVABIInfo::classifyKernelArgumentType(QualType Ty) const { // This behavior follows the CUDA spec // https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-function-argument-processing, // and matches the NVPTX implementation. - return getNaturalAlignIndirect(Ty, /* byval */ true); + return getNaturalAlignIndirect( + Ty, /* AddrSpace */ getContext().getTargetAddressSpace(LangAS::Default), + /* byval */ true); } } return classifyArgumentType(Ty); @@ -170,7 +172,10 @@ ABIArgInfo SPIRVABIInfo::classifyArgumentType(QualType Ty) const { // Records with non-trivial destructors/copy-constructors should not be // passed by value. if (auto RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, + getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); if (const RecordType *RT = Ty->getAs()) { const RecordDecl *RD = RT->getDecl(); diff --git a/clang/lib/CodeGen/Targets/Sparc.cpp b/clang/lib/CodeGen/Targets/Sparc.cpp index da8c7219be263..56069a7d818ee 100644 --- a/clang/lib/CodeGen/Targets/Sparc.cpp +++ b/clang/lib/CodeGen/Targets/Sparc.cpp @@ -232,7 +232,9 @@ SparcV9ABIInfo::classifyType(QualType Ty, unsigned SizeLimit) const { // Anything too big to fit in registers is passed with an explicit indirect // pointer / sret pointer. if (Size > SizeLimit) - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); // Treat an enum type as its underlying type. if (const EnumType *EnumTy = Ty->getAs()) @@ -253,7 +255,9 @@ SparcV9ABIInfo::classifyType(QualType Ty, unsigned SizeLimit) const { // If a C++ object has either a non-trivial copy constructor or a non-trivial // destructor, it is passed with an explicit indirect pointer / sret pointer. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); // This is a small aggregate type that should be passed in registers. // Build a coercion type from the LLVM struct type. diff --git a/clang/lib/CodeGen/Targets/SystemZ.cpp b/clang/lib/CodeGen/Targets/SystemZ.cpp index 23c96fa5cf98c..2dfd590af5b45 100644 --- a/clang/lib/CodeGen/Targets/SystemZ.cpp +++ b/clang/lib/CodeGen/Targets/SystemZ.cpp @@ -406,7 +406,7 @@ ABIArgInfo SystemZABIInfo::classifyReturnType(QualType RetTy) const { if (isVectorArgumentType(RetTy)) return ABIArgInfo::getDirect(); if (isCompoundType(RetTy) || getContext().getTypeSize(RetTy) > 64) - return getNaturalAlignIndirect(RetTy); + return getNaturalAlignIndirect(RetTy, getDataLayout().getAllocaAddrSpace()); return (isPromotableIntegerTypeForABI(RetTy) ? ABIArgInfo::getExtend(RetTy) : ABIArgInfo::getDirect()); } @@ -417,7 +417,9 @@ ABIArgInfo SystemZABIInfo::classifyArgumentType(QualType Ty) const { // Handle the generic C++ ABI. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); // Integers and enums are extended to full register width. if (isPromotableIntegerTypeForABI(Ty)) @@ -434,7 +436,9 @@ ABIArgInfo SystemZABIInfo::classifyArgumentType(QualType Ty) const { // Values that are not 1, 2, 4 or 8 bytes in size are passed indirectly. if (Size != 8 && Size != 16 && Size != 32 && Size != 64) - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); // Handle small structures. if (const RecordType *RT = Ty->getAs()) { @@ -442,7 +446,9 @@ ABIArgInfo SystemZABIInfo::classifyArgumentType(QualType Ty) const { // fail the size test above. const RecordDecl *RD = RT->getDecl(); if (RD->hasFlexibleArrayMember()) - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); // The structure is passed as an unextended integer, a float, or a double. if (isFPArgumentType(SingleElementTy)) { @@ -459,7 +465,9 @@ ABIArgInfo SystemZABIInfo::classifyArgumentType(QualType Ty) const { // Non-structure compounds are passed indirectly. if (isCompoundType(Ty)) - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); return ABIArgInfo::getDirect(nullptr); } diff --git a/clang/lib/CodeGen/Targets/WebAssembly.cpp b/clang/lib/CodeGen/Targets/WebAssembly.cpp index 70a968fe93ca7..dc45062c345ad 100644 --- a/clang/lib/CodeGen/Targets/WebAssembly.cpp +++ b/clang/lib/CodeGen/Targets/WebAssembly.cpp @@ -103,7 +103,9 @@ ABIArgInfo WebAssemblyABIInfo::classifyArgumentType(QualType Ty) const { // Records with non-trivial destructors/copy-constructors should not be // passed by value. if (auto RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); // Ignore empty structs/unions. if (isEmptyRecord(getContext(), Ty, true)) return ABIArgInfo::getIgnore(); diff --git a/clang/lib/CodeGen/Targets/X86.cpp b/clang/lib/CodeGen/Targets/X86.cpp index 6e5b46d5f91c8..61132a0c9ce30 100644 --- a/clang/lib/CodeGen/Targets/X86.cpp +++ b/clang/lib/CodeGen/Targets/X86.cpp @@ -462,7 +462,9 @@ ABIArgInfo X86_32ABIInfo::getIndirectReturnResult(QualType RetTy, CCState &State if (!IsMCUABI) return getNaturalAlignIndirectInReg(RetTy); } - return getNaturalAlignIndirect(RetTy, /*ByVal=*/false); + return getNaturalAlignIndirect( + RetTy, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); } ABIArgInfo X86_32ABIInfo::classifyReturnType(QualType RetTy, @@ -599,21 +601,26 @@ ABIArgInfo X86_32ABIInfo::getIndirectResult(QualType Ty, bool ByVal, if (!IsMCUABI) return getNaturalAlignIndirectInReg(Ty); } - return getNaturalAlignIndirect(Ty, false); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + false); } // Compute the byval alignment. unsigned TypeAlign = getContext().getTypeAlign(Ty) / 8; unsigned StackAlign = getTypeStackAlignInBytes(Ty, TypeAlign); if (StackAlign == 0) - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(4), 0, - /*ByVal=*/true); + return ABIArgInfo::getIndirect( + CharUnits::fromQuantity(4), + /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/true); // If the stack alignment is less than the type alignment, realign the // argument. bool Realign = TypeAlign > StackAlign; - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(StackAlign), 0, - /*ByVal=*/true, Realign); + return ABIArgInfo::getIndirect( + CharUnits::fromQuantity(StackAlign), + /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), /*ByVal=*/true, + Realign); } X86_32ABIInfo::Class X86_32ABIInfo::classify(QualType Ty) const { @@ -2165,13 +2172,13 @@ ABIArgInfo X86_64ABIInfo::getIndirectReturnResult(QualType Ty) const { Ty = EnumTy->getDecl()->getIntegerType(); if (Ty->isBitIntType()) - return getNaturalAlignIndirect(Ty); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace()); return (isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty) : ABIArgInfo::getDirect()); } - return getNaturalAlignIndirect(Ty); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace()); } bool X86_64ABIInfo::IsIllegalVectorType(QualType Ty) const { @@ -2211,7 +2218,8 @@ ABIArgInfo X86_64ABIInfo::getIndirectResult(QualType Ty, } if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + RAA == CGCXXABI::RAA_DirectInMemory); // Compute the byval alignment. We specify the alignment of the byval in all // cases so that the mid-level optimizer knows the alignment of the byval. @@ -2248,7 +2256,8 @@ ABIArgInfo X86_64ABIInfo::getIndirectResult(QualType Ty, Size)); } - return ABIArgInfo::getIndirect(CharUnits::fromQuantity(Align), 0); + return ABIArgInfo::getIndirect(CharUnits::fromQuantity(Align), + getDataLayout().getAllocaAddrSpace()); } /// The ABI specifies that a value should be passed in a full vector XMM/YMM @@ -3284,11 +3293,15 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, if (RT) { if (!IsReturnType) { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(RT, getCXXABI())) - return getNaturalAlignIndirect(Ty, RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); } if (RT->getDecl()->hasFlexibleArrayMember()) - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); } @@ -3305,7 +3318,10 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, return ABIArgInfo::getDirect(); return ABIArgInfo::getExpand(); } - return ABIArgInfo::getIndirect(Align, 0, /*ByVal=*/false); + return ABIArgInfo::getIndirect( + Align, + /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); } else if (IsVectorCall) { if (FreeSSERegs >= NumElts && (IsReturnType || Ty->isBuiltinType() || Ty->isVectorType())) { @@ -3315,7 +3331,10 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, return ABIArgInfo::getExpand(); } else if (!Ty->isBuiltinType() && !Ty->isVectorType()) { // HVAs are delayed and reclassified in the 2nd step. - return ABIArgInfo::getIndirect(Align, 0, /*ByVal=*/false); + return ABIArgInfo::getIndirect( + Align, + /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); } } } @@ -3332,7 +3351,9 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, // MS x64 ABI requirement: "Any argument that doesn't fit in 8 bytes, or is // not 1, 2, 4, or 8 bytes, must be passed by reference." if (Width > 64 || !llvm::isPowerOf2_64(Width)) - return getNaturalAlignIndirect(Ty, /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); // Otherwise, coerce it to a small integer. return ABIArgInfo::getDirect(llvm::IntegerType::get(getVMContext(), Width)); @@ -3351,7 +3372,10 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, if (IsMingw64) { const llvm::fltSemantics *LDF = &getTarget().getLongDoubleFormat(); if (LDF == &llvm::APFloat::x87DoubleExtended()) - return ABIArgInfo::getIndirect(Align, 0, /*ByVal=*/false); + return ABIArgInfo::getIndirect( + Align, + /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); } break; @@ -3361,7 +3385,10 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, // than 8 bytes are passed indirectly. GCC follows it. We follow it too, // even though it isn't particularly efficient. if (!IsReturnType) - return ABIArgInfo::getIndirect(Align, 0, /*ByVal=*/false); + return ABIArgInfo::getIndirect( + Align, + /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); // Mingw64 GCC returns i128 in XMM0. Coerce to v2i64 to handle that. // Clang matches them for compatibility. @@ -3381,7 +3408,10 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, // the power of 2. if (Width <= 64) return ABIArgInfo::getDirect(); - return ABIArgInfo::getIndirect(Align, 0, /*ByVal=*/false); + return ABIArgInfo::getIndirect( + Align, + /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); } return ABIArgInfo::getDirect(); From 02775164c6a1cd483f7780ea0f2fecd36e6cb730 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 5 Dec 2024 03:21:02 +0000 Subject: [PATCH 13/27] Fix formatting. --- clang/lib/CodeGen/ABIInfoImpl.cpp | 9 ++++----- clang/lib/CodeGen/SwiftCallingConv.cpp | 7 +++---- clang/lib/CodeGen/Targets/AArch64.cpp | 6 +++--- clang/lib/CodeGen/Targets/AMDGPU.cpp | 3 +-- clang/lib/CodeGen/Targets/ARC.cpp | 7 ++++--- clang/lib/CodeGen/Targets/ARM.cpp | 8 ++++---- clang/lib/CodeGen/Targets/BPF.cpp | 8 ++++---- clang/lib/CodeGen/Targets/Hexagon.cpp | 12 ++++++------ clang/lib/CodeGen/Targets/Lanai.cpp | 4 ++-- clang/lib/CodeGen/Targets/LoongArch.cpp | 15 ++++++++------- clang/lib/CodeGen/Targets/Mips.cpp | 7 +++---- clang/lib/CodeGen/Targets/PNaCl.cpp | 7 +++---- clang/lib/CodeGen/Targets/PPC.cpp | 12 +++++------- clang/lib/CodeGen/Targets/RISCV.cpp | 4 ++-- clang/lib/CodeGen/Targets/SPIR.cpp | 8 ++++---- clang/lib/CodeGen/Targets/Sparc.cpp | 4 ++-- clang/lib/CodeGen/Targets/WebAssembly.cpp | 4 ++-- clang/lib/CodeGen/Targets/X86.cpp | 1 - 18 files changed, 60 insertions(+), 66 deletions(-) diff --git a/clang/lib/CodeGen/ABIInfoImpl.cpp b/clang/lib/CodeGen/ABIInfoImpl.cpp index d55fcc9e125f0..65a186542692d 100644 --- a/clang/lib/CodeGen/ABIInfoImpl.cpp +++ b/clang/lib/CodeGen/ABIInfoImpl.cpp @@ -22,9 +22,8 @@ ABIArgInfo DefaultABIInfo::classifyArgumentType(QualType Ty) const { // passed by value. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) return getNaturalAlignIndirect( - Ty, - getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + Ty, getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); return getNaturalAlignIndirect( Ty, getContext().getTargetAddressSpace(LangAS::Default)); @@ -64,8 +63,8 @@ ABIArgInfo DefaultABIInfo::classifyReturnType(QualType RetTy) const { getContext().getTypeSize(getContext().getTargetInfo().hasInt128Type() ? getContext().Int128Ty : getContext().LongLongTy)) - return getNaturalAlignIndirect( - RetTy, getDataLayout().getAllocaAddrSpace()); + return getNaturalAlignIndirect(RetTy, + getDataLayout().getAllocaAddrSpace()); return (isPromotableIntegerTypeForABI(RetTy) ? ABIArgInfo::getExtend(RetTy) : ABIArgInfo::getDirect()); diff --git a/clang/lib/CodeGen/SwiftCallingConv.cpp b/clang/lib/CodeGen/SwiftCallingConv.cpp index 01a6402b42c6a..a3a8de8028e71 100644 --- a/clang/lib/CodeGen/SwiftCallingConv.cpp +++ b/clang/lib/CodeGen/SwiftCallingConv.cpp @@ -800,10 +800,9 @@ static ABIArgInfo classifyExpandedType(SwiftAggLowering &lowering, if (lowering.empty()) { return ABIArgInfo::getIgnore(); } else if (lowering.shouldPassIndirectly(forReturn)) { - return ABIArgInfo::getIndirect( - alignmentForIndirect, - /*AddrSpace*/ 0, - /*byval*/ false); + return ABIArgInfo::getIndirect(alignmentForIndirect, + /*AddrSpace*/ 0, + /*byval*/ false); } else { auto types = lowering.getCoerceAndExpandTypes(); return ABIArgInfo::getCoerceAndExpand(types.first, types.second); diff --git a/clang/lib/CodeGen/Targets/AArch64.cpp b/clang/lib/CodeGen/Targets/AArch64.cpp index c224b29a243a8..996bc7fa87b07 100644 --- a/clang/lib/CodeGen/Targets/AArch64.cpp +++ b/clang/lib/CodeGen/Targets/AArch64.cpp @@ -414,8 +414,8 @@ ABIArgInfo AArch64ABIInfo::classifyArgumentType(QualType Ty, bool IsVariadicFn, // copy constructor are always indirect. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); + Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } // Empty records are always ignored on Darwin, but actually passed in C++ mode @@ -490,7 +490,7 @@ ABIArgInfo AArch64ABIInfo::classifyArgumentType(QualType Ty, bool IsVariadicFn, } return getNaturalAlignIndirect( - Ty, /*AddrSpace*/getContext().getTargetAddressSpace(LangAS::Default), + Ty, /*AddrSpace*/ getContext().getTargetAddressSpace(LangAS::Default), /*ByVal=*/false); } diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index b180b1b8fa00c..56b12c18eb4f6 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -226,8 +226,7 @@ ABIArgInfo AMDGPUABIInfo::classifyArgumentType(QualType Ty, bool Variadic, // passed by value. if (auto RAA = getRecordArgABI(Ty, getCXXABI())) return getNaturalAlignIndirect( - Ty, - getContext().getTargetAddressSpace(LangAS::Default), + Ty, getContext().getTargetAddressSpace(LangAS::Default), RAA == CGCXXABI::RAA_DirectInMemory); // Ignore empty structs/unions. diff --git a/clang/lib/CodeGen/Targets/ARC.cpp b/clang/lib/CodeGen/Targets/ARC.cpp index a73b668c30ae1..4e2c869fd4b2f 100644 --- a/clang/lib/CodeGen/Targets/ARC.cpp +++ b/clang/lib/CodeGen/Targets/ARC.cpp @@ -70,9 +70,10 @@ class ARCTargetCodeGenInfo : public TargetCodeGenInfo { ABIArgInfo ARCABIInfo::getIndirectByRef(QualType Ty, bool HasFreeRegs) const { return HasFreeRegs - ? getNaturalAlignIndirectInReg(Ty) - : getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), false); + ? getNaturalAlignIndirectInReg(Ty) + : getNaturalAlignIndirect( + Ty, getContext().getTargetAddressSpace(LangAS::Default), + false); } ABIArgInfo ARCABIInfo::getIndirectByValue(QualType Ty) const { diff --git a/clang/lib/CodeGen/Targets/ARM.cpp b/clang/lib/CodeGen/Targets/ARM.cpp index c663d02ebb88b..6796b4074c30f 100644 --- a/clang/lib/CodeGen/Targets/ARM.cpp +++ b/clang/lib/CodeGen/Targets/ARM.cpp @@ -299,8 +299,8 @@ ABIArgInfo ARMABIInfo::coerceIllegalVector(QualType Ty) const { return ABIArgInfo::getDirect(ResType); } return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); } ABIArgInfo ARMABIInfo::classifyHomogeneousAggregate(QualType Ty, @@ -543,8 +543,8 @@ ABIArgInfo ARMABIInfo::classifyReturnType(QualType RetTy, bool isVariadic, if (const VectorType *VT = RetTy->getAs()) { // Large vector types should be returned via memory. if (getContext().getTypeSize(RetTy) > 128) - return getNaturalAlignIndirect( - RetTy, getDataLayout().getAllocaAddrSpace()); + return getNaturalAlignIndirect(RetTy, + getDataLayout().getAllocaAddrSpace()); // TODO: FP16/BF16 vectors should be converted to integer vectors // This check is similar to isIllegalVectorType - refactor? if ((!getTarget().hasLegalHalfType() && diff --git a/clang/lib/CodeGen/Targets/BPF.cpp b/clang/lib/CodeGen/Targets/BPF.cpp index 612c61d83a00a..ec3402b19b728 100644 --- a/clang/lib/CodeGen/Targets/BPF.cpp +++ b/clang/lib/CodeGen/Targets/BPF.cpp @@ -65,8 +65,8 @@ class BPFABIInfo : public DefaultABIInfo { return ABIArgInfo::getIgnore(); if (isAggregateTypeForABI(RetTy)) - return getNaturalAlignIndirect( - RetTy, getDataLayout().getAllocaAddrSpace()); + return getNaturalAlignIndirect(RetTy, + getDataLayout().getAllocaAddrSpace()); // Treat an enum type as its underlying type. if (const EnumType *EnumTy = RetTy->getAs()) @@ -75,8 +75,8 @@ class BPFABIInfo : public DefaultABIInfo { ASTContext &Context = getContext(); if (const auto *EIT = RetTy->getAs()) if (EIT->getNumBits() > Context.getTypeSize(Context.Int128Ty)) - return getNaturalAlignIndirect( - RetTy, getDataLayout().getAllocaAddrSpace()); + return getNaturalAlignIndirect(RetTy, + getDataLayout().getAllocaAddrSpace()); // Caller will do necessary sign/zero extension. return ABIArgInfo::getDirect(); diff --git a/clang/lib/CodeGen/Targets/Hexagon.cpp b/clang/lib/CodeGen/Targets/Hexagon.cpp index 151e269a10d38..6f8a5ad650b1a 100644 --- a/clang/lib/CodeGen/Targets/Hexagon.cpp +++ b/clang/lib/CodeGen/Targets/Hexagon.cpp @@ -106,8 +106,8 @@ ABIArgInfo HexagonABIInfo::classifyArgumentType(QualType Ty, if (Size > 64 && Ty->isBitIntType()) return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/true); + Ty, getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/true); return isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty) : ABIArgInfo::getDirect(); @@ -157,8 +157,8 @@ ABIArgInfo HexagonABIInfo::classifyReturnType(QualType RetTy) const { } // Large vector types should be returned via memory. if (Size > 64) - return getNaturalAlignIndirect( - RetTy, getDataLayout().getAllocaAddrSpace()); + return getNaturalAlignIndirect(RetTy, + getDataLayout().getAllocaAddrSpace()); } if (!isAggregateTypeForABI(RetTy)) { @@ -184,8 +184,8 @@ ABIArgInfo HexagonABIInfo::classifyReturnType(QualType RetTy) const { Size = llvm::bit_ceil(Size); return ABIArgInfo::getDirect(llvm::Type::getIntNTy(getVMContext(), Size)); } - return getNaturalAlignIndirect( - RetTy, getDataLayout().getAllocaAddrSpace(), /*ByVal=*/true); + return getNaturalAlignIndirect(RetTy, getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/true); } Address HexagonABIInfo::EmitVAArgFromMemory(CodeGenFunction &CGF, diff --git a/clang/lib/CodeGen/Targets/Lanai.cpp b/clang/lib/CodeGen/Targets/Lanai.cpp index 9b40655fb5289..b3d8a81a56aac 100644 --- a/clang/lib/CodeGen/Targets/Lanai.cpp +++ b/clang/lib/CodeGen/Targets/Lanai.cpp @@ -72,8 +72,8 @@ ABIArgInfo LanaiABIInfo::getIndirectResult(QualType Ty, bool ByVal, --State.FreeRegs; // Non-byval indirects just use one pointer. return getNaturalAlignIndirectInReg(Ty); } - return getNaturalAlignIndirect( - Ty, getDataLayout().getAllocaAddrSpace(), false); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + false); } // Compute the byval alignment. diff --git a/clang/lib/CodeGen/Targets/LoongArch.cpp b/clang/lib/CodeGen/Targets/LoongArch.cpp index b89130f93720d..4dc05fd0cb8eb 100644 --- a/clang/lib/CodeGen/Targets/LoongArch.cpp +++ b/clang/lib/CodeGen/Targets/LoongArch.cpp @@ -306,9 +306,9 @@ ABIArgInfo LoongArchABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (GARsLeft) GARsLeft -= 1; return getNaturalAlignIndirect( - Ty, - /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); + Ty, + /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } uint64_t Size = getContext().getTypeSize(Ty); @@ -384,8 +384,9 @@ ABIArgInfo LoongArchABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, (!getContext().getTargetInfo().hasInt128Type() && EIT->getNumBits() > 64)) return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + Ty, + /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); } return ABIArgInfo::getDirect(); @@ -409,8 +410,8 @@ ABIArgInfo LoongArchABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, llvm::ArrayType::get(llvm::IntegerType::get(getVMContext(), GRLen), 2)); } return getNaturalAlignIndirect( - Ty,/*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); } ABIArgInfo LoongArchABIInfo::classifyReturnType(QualType RetTy) const { diff --git a/clang/lib/CodeGen/Targets/Mips.cpp b/clang/lib/CodeGen/Targets/Mips.cpp index 067ffa2c2cd81..ae0125c48a0c3 100644 --- a/clang/lib/CodeGen/Targets/Mips.cpp +++ b/clang/lib/CodeGen/Targets/Mips.cpp @@ -210,8 +210,7 @@ MipsABIInfo::classifyArgumentType(QualType Ty, uint64_t &Offset) const { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { Offset = OrigOffset + MinABIStackAlignInBytes; return getNaturalAlignIndirect( - Ty, - getContext().getTargetAddressSpace(LangAS::Default), + Ty, getContext().getTargetAddressSpace(LangAS::Default), RAA == CGCXXABI::RAA_DirectInMemory); } @@ -326,8 +325,8 @@ ABIArgInfo MipsABIInfo::classifyReturnType(QualType RetTy) const { if (EIT->getNumBits() > 128 || (EIT->getNumBits() > 64 && !getContext().getTargetInfo().hasInt128Type())) - return getNaturalAlignIndirect( - RetTy, getDataLayout().getAllocaAddrSpace()); + return getNaturalAlignIndirect(RetTy, + getDataLayout().getAllocaAddrSpace()); if (isPromotableIntegerTypeForABI(RetTy)) return ABIArgInfo::getExtend(RetTy); diff --git a/clang/lib/CodeGen/Targets/PNaCl.cpp b/clang/lib/CodeGen/Targets/PNaCl.cpp index ac0bc598ff84a..fc5cc90a9e995 100644 --- a/clang/lib/CodeGen/Targets/PNaCl.cpp +++ b/clang/lib/CodeGen/Targets/PNaCl.cpp @@ -93,14 +93,13 @@ ABIArgInfo PNaClABIInfo::classifyReturnType(QualType RetTy) const { // In the PNaCl ABI we always return records/structures on the stack. if (isAggregateTypeForABI(RetTy)) - return getNaturalAlignIndirect( - RetTy, getDataLayout().getAllocaAddrSpace()); + return getNaturalAlignIndirect(RetTy, getDataLayout().getAllocaAddrSpace()); // Treat bit-precise integers as integers if <= 64, otherwise pass indirectly. if (const auto *EIT = RetTy->getAs()) { if (EIT->getNumBits() > 64) - return getNaturalAlignIndirect( - RetTy, getDataLayout().getAllocaAddrSpace()); + return getNaturalAlignIndirect(RetTy, + getDataLayout().getAllocaAddrSpace()); return ABIArgInfo::getDirect(); } diff --git a/clang/lib/CodeGen/Targets/PPC.cpp b/clang/lib/CodeGen/Targets/PPC.cpp index 5204bfddf1ab1..2a5d454c74d6a 100644 --- a/clang/lib/CodeGen/Targets/PPC.cpp +++ b/clang/lib/CodeGen/Targets/PPC.cpp @@ -209,8 +209,7 @@ ABIArgInfo AIXABIInfo::classifyArgumentType(QualType Ty) const { // passed by value. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) return getNaturalAlignIndirect( - Ty, - getContext().getTargetAddressSpace(LangAS::Default), + Ty, getContext().getTargetAddressSpace(LangAS::Default), RAA == CGCXXABI::RAA_DirectInMemory); CharUnits CCAlign = getParamTypeAlignment(Ty); @@ -218,7 +217,7 @@ ABIArgInfo AIXABIInfo::classifyArgumentType(QualType Ty) const { return ABIArgInfo::getIndirect( CCAlign, - /*AddrSpace*/getContext().getTargetAddressSpace(LangAS::Default), + /*AddrSpace*/ getContext().getTargetAddressSpace(LangAS::Default), /*ByVal*/ true, /*Realign*/ TyAlign > CCAlign); } @@ -850,14 +849,13 @@ PPC64_SVR4_ABIInfo::classifyArgumentType(QualType Ty) const { if (const auto *EIT = Ty->getAs()) if (EIT->getNumBits() > 128) return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/true); + Ty, getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/true); if (isAggregateTypeForABI(Ty)) { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) return getNaturalAlignIndirect( - Ty, - getContext().getTargetAddressSpace(LangAS::Default), + Ty, getContext().getTargetAddressSpace(LangAS::Default), RAA == CGCXXABI::RAA_DirectInMemory); uint64_t ABIAlign = getParamTypeAlignment(Ty).getQuantity(); diff --git a/clang/lib/CodeGen/Targets/RISCV.cpp b/clang/lib/CodeGen/Targets/RISCV.cpp index dc55924f747f4..03289d958f142 100644 --- a/clang/lib/CodeGen/Targets/RISCV.cpp +++ b/clang/lib/CodeGen/Targets/RISCV.cpp @@ -411,8 +411,8 @@ ABIArgInfo RISCVABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (ArgGPRsLeft) ArgGPRsLeft -= 1; return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); + Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } uint64_t Size = getContext().getTypeSize(Ty); diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp index 920d2151b622d..ef30db8cfb221 100644 --- a/clang/lib/CodeGen/Targets/SPIR.cpp +++ b/clang/lib/CodeGen/Targets/SPIR.cpp @@ -156,8 +156,9 @@ ABIArgInfo SPIRVABIInfo::classifyKernelArgumentType(QualType Ty) const { // https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-function-argument-processing, // and matches the NVPTX implementation. return getNaturalAlignIndirect( - Ty, /* AddrSpace */ getContext().getTargetAddressSpace(LangAS::Default), - /* byval */ true); + Ty, + /* AddrSpace */ getContext().getTargetAddressSpace(LangAS::Default), + /* byval */ true); } } return classifyArgumentType(Ty); @@ -173,8 +174,7 @@ ABIArgInfo SPIRVABIInfo::classifyArgumentType(QualType Ty) const { // passed by value. if (auto RAA = getRecordArgABI(Ty, getCXXABI())) return getNaturalAlignIndirect( - Ty, - getContext().getTargetAddressSpace(LangAS::Default), + Ty, getContext().getTargetAddressSpace(LangAS::Default), RAA == CGCXXABI::RAA_DirectInMemory); if (const RecordType *RT = Ty->getAs()) { diff --git a/clang/lib/CodeGen/Targets/Sparc.cpp b/clang/lib/CodeGen/Targets/Sparc.cpp index 56069a7d818ee..38107a75f9c97 100644 --- a/clang/lib/CodeGen/Targets/Sparc.cpp +++ b/clang/lib/CodeGen/Targets/Sparc.cpp @@ -233,8 +233,8 @@ SparcV9ABIInfo::classifyType(QualType Ty, unsigned SizeLimit) const { // pointer / sret pointer. if (Size > SizeLimit) return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + /*ByVal=*/false); // Treat an enum type as its underlying type. if (const EnumType *EnumTy = Ty->getAs()) diff --git a/clang/lib/CodeGen/Targets/WebAssembly.cpp b/clang/lib/CodeGen/Targets/WebAssembly.cpp index dc45062c345ad..502f5260fbabf 100644 --- a/clang/lib/CodeGen/Targets/WebAssembly.cpp +++ b/clang/lib/CodeGen/Targets/WebAssembly.cpp @@ -104,8 +104,8 @@ ABIArgInfo WebAssemblyABIInfo::classifyArgumentType(QualType Ty) const { // passed by value. if (auto RAA = getRecordArgABI(Ty, getCXXABI())) return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + Ty, getContext().getTargetAddressSpace(LangAS::Default), + RAA == CGCXXABI::RAA_DirectInMemory); // Ignore empty structs/unions. if (isEmptyRecord(getContext(), Ty, true)) return ABIArgInfo::getIgnore(); diff --git a/clang/lib/CodeGen/Targets/X86.cpp b/clang/lib/CodeGen/Targets/X86.cpp index 61132a0c9ce30..af6d62fc05999 100644 --- a/clang/lib/CodeGen/Targets/X86.cpp +++ b/clang/lib/CodeGen/Targets/X86.cpp @@ -3302,7 +3302,6 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, return getNaturalAlignIndirect( Ty, getContext().getTargetAddressSpace(LangAS::Default), /*ByVal=*/false); - } const Type *Base = nullptr; From f6c8e01551296301e9f2a5f6187c1871655c24a9 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Sun, 5 Jan 2025 18:02:43 +0200 Subject: [PATCH 14/27] Add helper accessor for `LangAS::Default -> TargetAS` queries. --- clang/lib/CodeGen/ABIInfo.cpp | 4 +++ clang/lib/CodeGen/ABIInfo.h | 2 ++ clang/lib/CodeGen/Targets/AArch64.cpp | 16 ++++------- clang/lib/CodeGen/Targets/ARC.cpp | 7 ++--- clang/lib/CodeGen/Targets/ARM.cpp | 21 ++++++--------- clang/lib/CodeGen/Targets/BPF.cpp | 5 ++-- clang/lib/CodeGen/Targets/CSKY.cpp | 5 ++-- clang/lib/CodeGen/Targets/Hexagon.cpp | 13 +++------ clang/lib/CodeGen/Targets/Lanai.cpp | 5 ++-- clang/lib/CodeGen/Targets/LoongArch.cpp | 12 +++------ clang/lib/CodeGen/Targets/Mips.cpp | 8 +++--- clang/lib/CodeGen/Targets/NVPTX.cpp | 11 +++----- clang/lib/CodeGen/Targets/PNaCl.cpp | 8 +++--- clang/lib/CodeGen/Targets/PPC.cpp | 28 +++++++------------ clang/lib/CodeGen/Targets/RISCV.cpp | 13 ++++----- clang/lib/CodeGen/Targets/SPIR.cpp | 11 +++----- clang/lib/CodeGen/Targets/Sparc.cpp | 10 +++---- clang/lib/CodeGen/Targets/SystemZ.cpp | 17 ++++-------- clang/lib/CodeGen/Targets/WebAssembly.cpp | 5 ++-- clang/lib/CodeGen/Targets/X86.cpp | 33 +++++++---------------- 20 files changed, 85 insertions(+), 149 deletions(-) diff --git a/clang/lib/CodeGen/ABIInfo.cpp b/clang/lib/CodeGen/ABIInfo.cpp index 19d8c66b000cb..b7f355eb335f4 100644 --- a/clang/lib/CodeGen/ABIInfo.cpp +++ b/clang/lib/CodeGen/ABIInfo.cpp @@ -12,6 +12,10 @@ using namespace clang; using namespace clang::CodeGen; +unsigned ABIInfo::getTargetDefaultAS() const { + return getContext().getTargetAddressSpace(LangAS::Default); +} + // Pin the vtable to this file. ABIInfo::~ABIInfo() = default; diff --git a/clang/lib/CodeGen/ABIInfo.h b/clang/lib/CodeGen/ABIInfo.h index 57bcb3dd0a852..538b4ed26b113 100644 --- a/clang/lib/CodeGen/ABIInfo.h +++ b/clang/lib/CodeGen/ABIInfo.h @@ -49,6 +49,8 @@ class ABIInfo { CodeGen::CodeGenTypes &CGT; llvm::CallingConv::ID RuntimeCC; + unsigned getTargetDefaultAS() const; + public: ABIInfo(CodeGen::CodeGenTypes &cgt) : CGT(cgt), RuntimeCC(llvm::CallingConv::C) {} diff --git a/clang/lib/CodeGen/Targets/AArch64.cpp b/clang/lib/CodeGen/Targets/AArch64.cpp index b010ea8407b06..5f23b99adaa0c 100644 --- a/clang/lib/CodeGen/Targets/AArch64.cpp +++ b/clang/lib/CodeGen/Targets/AArch64.cpp @@ -326,8 +326,7 @@ ABIArgInfo AArch64ABIInfo::coerceIllegalVector(QualType Ty, unsigned &NSRN, return ABIArgInfo::getDirect(ResType); } - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); } ABIArgInfo AArch64ABIInfo::coerceAndExpandPureScalableAggregate( @@ -335,9 +334,7 @@ ABIArgInfo AArch64ABIInfo::coerceAndExpandPureScalableAggregate( const SmallVectorImpl &UnpaddedCoerceToSeq, unsigned &NSRN, unsigned &NPRN) const { if (!IsNamedArg || NSRN + NVec > 8 || NPRN + NPred > 4) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); NSRN += NVec; NPRN += NPred; @@ -377,8 +374,7 @@ ABIArgInfo AArch64ABIInfo::classifyArgumentType(QualType Ty, bool IsVariadicFn, if (const auto *EIT = Ty->getAs()) if (EIT->getNumBits() > 128) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), false); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), false); if (Ty->isVectorType()) NSRN = std::min(NSRN + 1, 8u); @@ -419,7 +415,7 @@ ABIArgInfo AArch64ABIInfo::classifyArgumentType(QualType Ty, bool IsVariadicFn, // copy constructor are always indirect. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + Ty, /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } @@ -494,9 +490,7 @@ ABIArgInfo AArch64ABIInfo::classifyArgumentType(QualType Ty, bool IsVariadicFn, : llvm::ArrayType::get(BaseTy, Size / Alignment)); } - return getNaturalAlignIndirect( - Ty, /*AddrSpace*/ getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); } ABIArgInfo AArch64ABIInfo::classifyReturnType(QualType RetTy, diff --git a/clang/lib/CodeGen/Targets/ARC.cpp b/clang/lib/CodeGen/Targets/ARC.cpp index 4e2c869fd4b2f..1e5516e1b4a0e 100644 --- a/clang/lib/CodeGen/Targets/ARC.cpp +++ b/clang/lib/CodeGen/Targets/ARC.cpp @@ -71,9 +71,7 @@ class ARCTargetCodeGenInfo : public TargetCodeGenInfo { ABIArgInfo ARCABIInfo::getIndirectByRef(QualType Ty, bool HasFreeRegs) const { return HasFreeRegs ? getNaturalAlignIndirectInReg(Ty) - : getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - false); + : getNaturalAlignIndirect(Ty, getTargetDefaultAS(), false); } ABIArgInfo ARCABIInfo::getIndirectByValue(QualType Ty) const { @@ -81,8 +79,7 @@ ABIArgInfo ARCABIInfo::getIndirectByValue(QualType Ty) const { const unsigned MinABIStackAlignInBytes = 4; unsigned TypeAlign = getContext().getTypeAlign(Ty) / 8; return ABIArgInfo::getIndirect( - CharUnits::fromQuantity(4), - /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + CharUnits::fromQuantity(4), /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/true, TypeAlign > MinABIStackAlignInBytes); } diff --git a/clang/lib/CodeGen/Targets/ARM.cpp b/clang/lib/CodeGen/Targets/ARM.cpp index 6796b4074c30f..e4628a7e9731a 100644 --- a/clang/lib/CodeGen/Targets/ARM.cpp +++ b/clang/lib/CodeGen/Targets/ARM.cpp @@ -298,9 +298,8 @@ ABIArgInfo ARMABIInfo::coerceIllegalVector(QualType Ty) const { llvm::Type::getInt32Ty(getVMContext()), Size / 32); return ABIArgInfo::getDirect(ResType); } - return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), + /*ByVal=*/false); } ABIArgInfo ARMABIInfo::classifyHomogeneousAggregate(QualType Ty, @@ -357,10 +356,8 @@ ABIArgInfo ARMABIInfo::classifyArgumentType(QualType Ty, bool isVariadic, if (const auto *EIT = Ty->getAs()) if (EIT->getNumBits() > 64) - return getNaturalAlignIndirect( - Ty, - /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/true); + return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), + /*ByVal=*/true); return (isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty, CGT.ConvertType(Ty)) @@ -368,9 +365,8 @@ ABIArgInfo ARMABIInfo::classifyArgumentType(QualType Ty, bool isVariadic, } if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + RAA == CGCXXABI::RAA_DirectInMemory); } // Ignore empty records. @@ -405,7 +401,7 @@ ABIArgInfo ARMABIInfo::classifyArgumentType(QualType Ty, bool isVariadic, // and a pointer is passed. return ABIArgInfo::getIndirect( CharUnits::fromQuantity(getContext().getTypeAlign(Ty) / 8), - getContext().getTargetAddressSpace(LangAS::Default), false); + getTargetDefaultAS(), false); } // Support byval for ARM. @@ -424,8 +420,7 @@ ABIArgInfo ARMABIInfo::classifyArgumentType(QualType Ty, bool isVariadic, if (getContext().getTypeSizeInChars(Ty) > CharUnits::fromQuantity(64)) { assert(getABIKind() != ARMABIKind::AAPCS16_VFP && "unexpected byval"); return ABIArgInfo::getIndirect( - CharUnits::fromQuantity(ABIAlign), - /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + CharUnits::fromQuantity(ABIAlign), /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/true, /*Realign=*/TyAlign > ABIAlign); } diff --git a/clang/lib/CodeGen/Targets/BPF.cpp b/clang/lib/CodeGen/Targets/BPF.cpp index ec3402b19b728..31ade54f7ef71 100644 --- a/clang/lib/CodeGen/Targets/BPF.cpp +++ b/clang/lib/CodeGen/Targets/BPF.cpp @@ -43,7 +43,7 @@ class BPFABIInfo : public DefaultABIInfo { return ABIArgInfo::getDirect(CoerceTy); } else { return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default)); + Ty, getTargetDefaultAS()); } } @@ -53,8 +53,7 @@ class BPFABIInfo : public DefaultABIInfo { ASTContext &Context = getContext(); if (const auto *EIT = Ty->getAs()) if (EIT->getNumBits() > Context.getTypeSize(Context.Int128Ty)) - return getNaturalAlignIndirect( - Ty, Context.getTargetAddressSpace(LangAS::Default)); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS()); return (isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty) : ABIArgInfo::getDirect()); diff --git a/clang/lib/CodeGen/Targets/CSKY.cpp b/clang/lib/CodeGen/Targets/CSKY.cpp index 41dd921ae309e..29b239b24e56b 100644 --- a/clang/lib/CodeGen/Targets/CSKY.cpp +++ b/clang/lib/CodeGen/Targets/CSKY.cpp @@ -83,7 +83,7 @@ ABIArgInfo CSKYABIInfo::classifyArgumentType(QualType Ty, int &ArgGPRsLeft, if (ArgGPRsLeft) ArgGPRsLeft -= 1; return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + Ty, /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } @@ -145,8 +145,7 @@ ABIArgInfo CSKYABIInfo::classifyArgumentType(QualType Ty, int &ArgGPRsLeft, llvm::IntegerType::get(getVMContext(), XLen), (Size + 31) / XLen)); } } - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); } ABIArgInfo CSKYABIInfo::classifyReturnType(QualType RetTy) const { diff --git a/clang/lib/CodeGen/Targets/Hexagon.cpp b/clang/lib/CodeGen/Targets/Hexagon.cpp index 6f8a5ad650b1a..38d1c0232d6e1 100644 --- a/clang/lib/CodeGen/Targets/Hexagon.cpp +++ b/clang/lib/CodeGen/Targets/Hexagon.cpp @@ -105,18 +105,15 @@ ABIArgInfo HexagonABIInfo::classifyArgumentType(QualType Ty, HexagonAdjustRegsLeft(Size, RegsLeft); if (Size > 64 && Ty->isBitIntType()) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/true); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/true); return isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty) : ABIArgInfo::getDirect(); } if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + RAA == CGCXXABI::RAA_DirectInMemory); // Ignore empty records. if (isEmptyRecord(getContext(), Ty, true)) @@ -126,9 +123,7 @@ ABIArgInfo HexagonABIInfo::classifyArgumentType(QualType Ty, unsigned Align = getContext().getTypeAlign(Ty); if (Size > 64) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/true); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/true); if (HexagonAdjustRegsLeft(Size, RegsLeft)) Align = Size <= 32 ? 32 : 64; diff --git a/clang/lib/CodeGen/Targets/Lanai.cpp b/clang/lib/CodeGen/Targets/Lanai.cpp index b3d8a81a56aac..4ea078a9e48af 100644 --- a/clang/lib/CodeGen/Targets/Lanai.cpp +++ b/clang/lib/CodeGen/Targets/Lanai.cpp @@ -94,9 +94,8 @@ ABIArgInfo LanaiABIInfo::classifyArgumentType(QualType Ty, if (RAA == CGCXXABI::RAA_Indirect) { return getIndirectResult(Ty, /*ByVal=*/false, State); } else if (RAA == CGCXXABI::RAA_DirectInMemory) { - return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/true); + return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), + /*ByVal=*/true); } } diff --git a/clang/lib/CodeGen/Targets/LoongArch.cpp b/clang/lib/CodeGen/Targets/LoongArch.cpp index 4dc05fd0cb8eb..8edff7b59fee9 100644 --- a/clang/lib/CodeGen/Targets/LoongArch.cpp +++ b/clang/lib/CodeGen/Targets/LoongArch.cpp @@ -306,8 +306,7 @@ ABIArgInfo LoongArchABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (GARsLeft) GARsLeft -= 1; return getNaturalAlignIndirect( - Ty, - /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + Ty, /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } @@ -384,9 +383,7 @@ ABIArgInfo LoongArchABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, (!getContext().getTargetInfo().hasInt128Type() && EIT->getNumBits() > 64)) return getNaturalAlignIndirect( - Ty, - /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + Ty, /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/false); } return ABIArgInfo::getDirect(); @@ -409,9 +406,8 @@ ABIArgInfo LoongArchABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, return ABIArgInfo::getDirect( llvm::ArrayType::get(llvm::IntegerType::get(getVMContext(), GRLen), 2)); } - return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), + /*ByVal=*/false); } ABIArgInfo LoongArchABIInfo::classifyReturnType(QualType RetTy) const { diff --git a/clang/lib/CodeGen/Targets/Mips.cpp b/clang/lib/CodeGen/Targets/Mips.cpp index ae0125c48a0c3..49613c3341b1a 100644 --- a/clang/lib/CodeGen/Targets/Mips.cpp +++ b/clang/lib/CodeGen/Targets/Mips.cpp @@ -209,9 +209,8 @@ MipsABIInfo::classifyArgumentType(QualType Ty, uint64_t &Offset) const { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { Offset = OrigOffset + MinABIStackAlignInBytes; - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + RAA == CGCXXABI::RAA_DirectInMemory); } // If we have reached here, aggregates are passed directly by coercing to @@ -233,8 +232,7 @@ MipsABIInfo::classifyArgumentType(QualType Ty, uint64_t &Offset) const { if (EIT->getNumBits() > 128 || (EIT->getNumBits() > 64 && !getContext().getTargetInfo().hasInt128Type())) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default)); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS()); // All integral types are promoted to the GPR width. if (Ty->isIntegralOrEnumerationType()) diff --git a/clang/lib/CodeGen/Targets/NVPTX.cpp b/clang/lib/CodeGen/Targets/NVPTX.cpp index ce59f63fcf6f1..5aab9702c467c 100644 --- a/clang/lib/CodeGen/Targets/NVPTX.cpp +++ b/clang/lib/CodeGen/Targets/NVPTX.cpp @@ -193,19 +193,16 @@ ABIArgInfo NVPTXABIInfo::classifyArgumentType(QualType Ty) const { return ABIArgInfo::getDirect( CGInfo.getCUDADeviceBuiltinTextureDeviceType()); } - return getNaturalAlignIndirect( - Ty, /* AddrSpace */ getContext().getTargetAddressSpace(LangAS::Default), - /* byval */ true); + return getNaturalAlignIndirect(Ty, /* AddrSpace */ getTargetDefaultAS(), + /* byval */ true); } if (const auto *EIT = Ty->getAs()) { if ((EIT->getNumBits() > 128) || (!getContext().getTargetInfo().hasInt128Type() && EIT->getNumBits() > 64)) - return getNaturalAlignIndirect( - Ty, - /* AddrSpace */ getContext().getTargetAddressSpace(LangAS::Default), - /* byval */ true); + return getNaturalAlignIndirect(Ty, /* AddrSpace */ getTargetDefaultAS(), + /* byval */ true); } return (isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty) diff --git a/clang/lib/CodeGen/Targets/PNaCl.cpp b/clang/lib/CodeGen/Targets/PNaCl.cpp index fc5cc90a9e995..85c2743a16985 100644 --- a/clang/lib/CodeGen/Targets/PNaCl.cpp +++ b/clang/lib/CodeGen/Targets/PNaCl.cpp @@ -63,9 +63,8 @@ RValue PNaClABIInfo::EmitVAArg(CodeGenFunction &CGF, Address VAListAddr, ABIArgInfo PNaClABIInfo::classifyArgumentType(QualType Ty) const { if (isAggregateTypeForABI(Ty)) { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + RAA == CGCXXABI::RAA_DirectInMemory); return getNaturalAlignIndirect( Ty, getContext().getTargetAddressSpace(LangAS::Default)); } else if (const EnumType *EnumTy = Ty->getAs()) { @@ -78,8 +77,7 @@ ABIArgInfo PNaClABIInfo::classifyArgumentType(QualType Ty) const { // Treat bit-precise integers as integers if <= 64, otherwise pass // indirectly. if (EIT->getNumBits() > 64) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default)); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS()); return ABIArgInfo::getDirect(); } diff --git a/clang/lib/CodeGen/Targets/PPC.cpp b/clang/lib/CodeGen/Targets/PPC.cpp index 2a5d454c74d6a..2acaf181677af 100644 --- a/clang/lib/CodeGen/Targets/PPC.cpp +++ b/clang/lib/CodeGen/Targets/PPC.cpp @@ -208,17 +208,15 @@ ABIArgInfo AIXABIInfo::classifyArgumentType(QualType Ty) const { // Records with non-trivial destructors/copy-constructors should not be // passed by value. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + RAA == CGCXXABI::RAA_DirectInMemory); CharUnits CCAlign = getParamTypeAlignment(Ty); CharUnits TyAlign = getContext().getTypeAlignInChars(Ty); - return ABIArgInfo::getIndirect( - CCAlign, - /*AddrSpace*/ getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal*/ true, /*Realign*/ TyAlign > CCAlign); + return ABIArgInfo::getIndirect(CCAlign, /*AddrSpace*/ getTargetDefaultAS(), + /*ByVal*/ true, + /*Realign*/ TyAlign > CCAlign); } return (isPromotableTypeForABI(Ty) @@ -837,9 +835,7 @@ PPC64_SVR4_ABIInfo::classifyArgumentType(QualType Ty) const { if (Ty->isVectorType()) { uint64_t Size = getContext().getTypeSize(Ty); if (Size > 128) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); else if (Size < 128) { llvm::Type *CoerceTy = llvm::IntegerType::get(getVMContext(), Size); return ABIArgInfo::getDirect(CoerceTy); @@ -848,15 +844,12 @@ PPC64_SVR4_ABIInfo::classifyArgumentType(QualType Ty) const { if (const auto *EIT = Ty->getAs()) if (EIT->getNumBits() > 128) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/true); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/true); if (isAggregateTypeForABI(Ty)) { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + RAA == CGCXXABI::RAA_DirectInMemory); uint64_t ABIAlign = getParamTypeAlignment(Ty).getQuantity(); uint64_t TyAlign = getContext().getTypeAlignInChars(Ty).getQuantity(); @@ -898,8 +891,7 @@ PPC64_SVR4_ABIInfo::classifyArgumentType(QualType Ty) const { // All other aggregates are passed ByVal. return ABIArgInfo::getIndirect( - CharUnits::fromQuantity(ABIAlign), - /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + CharUnits::fromQuantity(ABIAlign), /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/true, /*Realign=*/TyAlign > ABIAlign); } diff --git a/clang/lib/CodeGen/Targets/RISCV.cpp b/clang/lib/CodeGen/Targets/RISCV.cpp index 60dd48ba5532c..7f8df7458fcb3 100644 --- a/clang/lib/CodeGen/Targets/RISCV.cpp +++ b/clang/lib/CodeGen/Targets/RISCV.cpp @@ -411,7 +411,7 @@ ABIArgInfo RISCVABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (ArgGPRsLeft) ArgGPRsLeft -= 1; return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), + Ty, /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } @@ -493,10 +493,8 @@ ABIArgInfo RISCVABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (EIT->getNumBits() > 128 || (!getContext().getTargetInfo().hasInt128Type() && EIT->getNumBits() > 64)) - return getNaturalAlignIndirect( - Ty, - /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), + /*ByVal=*/false); } return ABIArgInfo::getDirect(); @@ -528,9 +526,8 @@ ABIArgInfo RISCVABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, llvm::IntegerType::get(getVMContext(), XLen), 2)); } } - return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), + /*ByVal=*/false); } ABIArgInfo RISCVABIInfo::classifyReturnType(QualType RetTy) const { diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp index ef30db8cfb221..1a66be1eff552 100644 --- a/clang/lib/CodeGen/Targets/SPIR.cpp +++ b/clang/lib/CodeGen/Targets/SPIR.cpp @@ -155,10 +155,8 @@ ABIArgInfo SPIRVABIInfo::classifyKernelArgumentType(QualType Ty) const { // This behavior follows the CUDA spec // https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-function-argument-processing, // and matches the NVPTX implementation. - return getNaturalAlignIndirect( - Ty, - /* AddrSpace */ getContext().getTargetAddressSpace(LangAS::Default), - /* byval */ true); + return getNaturalAlignIndirect(Ty, /* AddrSpace */ getTargetDefaultAS(), + /* byval */ true); } } return classifyArgumentType(Ty); @@ -173,9 +171,8 @@ ABIArgInfo SPIRVABIInfo::classifyArgumentType(QualType Ty) const { // Records with non-trivial destructors/copy-constructors should not be // passed by value. if (auto RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + RAA == CGCXXABI::RAA_DirectInMemory); if (const RecordType *RT = Ty->getAs()) { const RecordDecl *RD = RT->getDecl(); diff --git a/clang/lib/CodeGen/Targets/Sparc.cpp b/clang/lib/CodeGen/Targets/Sparc.cpp index 38107a75f9c97..d0b40aa9ceab1 100644 --- a/clang/lib/CodeGen/Targets/Sparc.cpp +++ b/clang/lib/CodeGen/Targets/Sparc.cpp @@ -232,9 +232,8 @@ SparcV9ABIInfo::classifyType(QualType Ty, unsigned SizeLimit) const { // Anything too big to fit in registers is passed with an explicit indirect // pointer / sret pointer. if (Size > SizeLimit) - return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), + /*ByVal=*/false); // Treat an enum type as its underlying type. if (const EnumType *EnumTy = Ty->getAs()) @@ -255,9 +254,8 @@ SparcV9ABIInfo::classifyType(QualType Ty, unsigned SizeLimit) const { // If a C++ object has either a non-trivial copy constructor or a non-trivial // destructor, it is passed with an explicit indirect pointer / sret pointer. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + RAA == CGCXXABI::RAA_DirectInMemory); // This is a small aggregate type that should be passed in registers. // Build a coercion type from the LLVM struct type. diff --git a/clang/lib/CodeGen/Targets/SystemZ.cpp b/clang/lib/CodeGen/Targets/SystemZ.cpp index 2dfd590af5b45..4fd141c694c8b 100644 --- a/clang/lib/CodeGen/Targets/SystemZ.cpp +++ b/clang/lib/CodeGen/Targets/SystemZ.cpp @@ -417,9 +417,8 @@ ABIArgInfo SystemZABIInfo::classifyArgumentType(QualType Ty) const { // Handle the generic C++ ABI. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + RAA == CGCXXABI::RAA_DirectInMemory); // Integers and enums are extended to full register width. if (isPromotableIntegerTypeForABI(Ty)) @@ -436,9 +435,7 @@ ABIArgInfo SystemZABIInfo::classifyArgumentType(QualType Ty) const { // Values that are not 1, 2, 4 or 8 bytes in size are passed indirectly. if (Size != 8 && Size != 16 && Size != 32 && Size != 64) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); // Handle small structures. if (const RecordType *RT = Ty->getAs()) { @@ -446,9 +443,7 @@ ABIArgInfo SystemZABIInfo::classifyArgumentType(QualType Ty) const { // fail the size test above. const RecordDecl *RD = RT->getDecl(); if (RD->hasFlexibleArrayMember()) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); // The structure is passed as an unextended integer, a float, or a double. if (isFPArgumentType(SingleElementTy)) { @@ -465,9 +460,7 @@ ABIArgInfo SystemZABIInfo::classifyArgumentType(QualType Ty) const { // Non-structure compounds are passed indirectly. if (isCompoundType(Ty)) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); return ABIArgInfo::getDirect(nullptr); } diff --git a/clang/lib/CodeGen/Targets/WebAssembly.cpp b/clang/lib/CodeGen/Targets/WebAssembly.cpp index 502f5260fbabf..0601213483864 100644 --- a/clang/lib/CodeGen/Targets/WebAssembly.cpp +++ b/clang/lib/CodeGen/Targets/WebAssembly.cpp @@ -103,9 +103,8 @@ ABIArgInfo WebAssemblyABIInfo::classifyArgumentType(QualType Ty) const { // Records with non-trivial destructors/copy-constructors should not be // passed by value. if (auto RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + RAA == CGCXXABI::RAA_DirectInMemory); // Ignore empty structs/unions. if (isEmptyRecord(getContext(), Ty, true)) return ABIArgInfo::getIgnore(); diff --git a/clang/lib/CodeGen/Targets/X86.cpp b/clang/lib/CodeGen/Targets/X86.cpp index af6d62fc05999..53e1df25522db 100644 --- a/clang/lib/CodeGen/Targets/X86.cpp +++ b/clang/lib/CodeGen/Targets/X86.cpp @@ -3293,9 +3293,8 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, if (RT) { if (!IsReturnType) { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(RT, getCXXABI())) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + RAA == CGCXXABI::RAA_DirectInMemory); } if (RT->getDecl()->hasFlexibleArrayMember()) @@ -3317,10 +3316,8 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, return ABIArgInfo::getDirect(); return ABIArgInfo::getExpand(); } - return ABIArgInfo::getIndirect( - Align, - /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + return ABIArgInfo::getIndirect(Align, /*AddrSpace=*/getTargetDefaultAS(), + /*ByVal=*/false); } else if (IsVectorCall) { if (FreeSSERegs >= NumElts && (IsReturnType || Ty->isBuiltinType() || Ty->isVectorType())) { @@ -3331,9 +3328,7 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, } else if (!Ty->isBuiltinType() && !Ty->isVectorType()) { // HVAs are delayed and reclassified in the 2nd step. return ABIArgInfo::getIndirect( - Align, - /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + Align, /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/false); } } } @@ -3350,9 +3345,7 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, // MS x64 ABI requirement: "Any argument that doesn't fit in 8 bytes, or is // not 1, 2, 4, or 8 bytes, must be passed by reference." if (Width > 64 || !llvm::isPowerOf2_64(Width)) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); // Otherwise, coerce it to a small integer. return ABIArgInfo::getDirect(llvm::IntegerType::get(getVMContext(), Width)); @@ -3372,9 +3365,7 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, const llvm::fltSemantics *LDF = &getTarget().getLongDoubleFormat(); if (LDF == &llvm::APFloat::x87DoubleExtended()) return ABIArgInfo::getIndirect( - Align, - /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + Align, /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/false); } break; @@ -3385,9 +3376,7 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, // even though it isn't particularly efficient. if (!IsReturnType) return ABIArgInfo::getIndirect( - Align, - /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + Align, /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/false); // Mingw64 GCC returns i128 in XMM0. Coerce to v2i64 to handle that. // Clang matches them for compatibility. @@ -3407,10 +3396,8 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, // the power of 2. if (Width <= 64) return ABIArgInfo::getDirect(); - return ABIArgInfo::getIndirect( - Align, - /*AddrSpace=*/getContext().getTargetAddressSpace(LangAS::Default), - /*ByVal=*/false); + return ABIArgInfo::getIndirect(Align, /*AddrSpace=*/getTargetDefaultAS(), + /*ByVal=*/false); } return ABIArgInfo::getDirect(); From 0f724f8725eb921db844b9037914a10757472ab5 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Sun, 5 Jan 2025 18:13:13 +0200 Subject: [PATCH 15/27] Align AMDGPU argument classification. --- clang/lib/CodeGen/Targets/AMDGPU.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 56b12c18eb4f6..10ef72dfa55d9 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -225,9 +225,8 @@ ABIArgInfo AMDGPUABIInfo::classifyArgumentType(QualType Ty, bool Variadic, // Records with non-trivial destructors/copy-constructors should not be // passed by value. if (auto RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + RAA == CGCXXABI::RAA_DirectInMemory); // Ignore empty structs/unions. if (isEmptyRecord(getContext(), Ty, true)) From 8f472f377c9aba3df9b9930018c9f3e4e39628da Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 7 Jan 2025 13:50:57 +0200 Subject: [PATCH 16/27] Tweak Swift's use of AS aware `getIndirect`. --- clang/lib/CodeGen/SwiftCallingConv.cpp | 17 ++++++++++------- 1 file changed, 10 insertions(+), 7 deletions(-) diff --git a/clang/lib/CodeGen/SwiftCallingConv.cpp b/clang/lib/CodeGen/SwiftCallingConv.cpp index a3a8de8028e71..299a379977704 100644 --- a/clang/lib/CodeGen/SwiftCallingConv.cpp +++ b/clang/lib/CodeGen/SwiftCallingConv.cpp @@ -796,12 +796,13 @@ bool swiftcall::mustPassRecordIndirectly(CodeGenModule &CGM, static ABIArgInfo classifyExpandedType(SwiftAggLowering &lowering, bool forReturn, - CharUnits alignmentForIndirect) { + CharUnits alignmentForIndirect, + unsigned IndirectAS) { if (lowering.empty()) { return ABIArgInfo::getIgnore(); } else if (lowering.shouldPassIndirectly(forReturn)) { return ABIArgInfo::getIndirect(alignmentForIndirect, - /*AddrSpace*/ 0, + /*AddrSpace*/ IndirectAS, /*byval*/ false); } else { auto types = lowering.getCoerceAndExpandTypes(); @@ -811,21 +812,23 @@ static ABIArgInfo classifyExpandedType(SwiftAggLowering &lowering, static ABIArgInfo classifyType(CodeGenModule &CGM, CanQualType type, bool forReturn) { + unsigned IndirectAS = forReturn + ? CGM.getDataLayout().getAllocaAddrSpace() + : CGM.getContext().getTargetAddressSpace(LangAS::Default); if (auto recordType = dyn_cast(type)) { auto record = recordType->getDecl(); auto &layout = CGM.getContext().getASTRecordLayout(record); if (mustPassRecordIndirectly(CGM, record)) return ABIArgInfo::getIndirect( - layout.getAlignment(), - /*AddrSpace*/ CGM.getContext().getTargetAddressSpace(LangAS::Default), - /*byval*/ false); + layout.getAlignment(), /*AddrSpace=*/ IndirectAS, /*byval=*/ false); SwiftAggLowering lowering(CGM); lowering.addTypedData(recordType->getDecl(), CharUnits::Zero(), layout); lowering.finish(); - return classifyExpandedType(lowering, forReturn, layout.getAlignment()); + return classifyExpandedType(lowering, forReturn, layout.getAlignment(), + IndirectAS); } // Just assume that all of our target ABIs can support returning at least @@ -841,7 +844,7 @@ static ABIArgInfo classifyType(CodeGenModule &CGM, CanQualType type, lowering.finish(); CharUnits alignment = CGM.getContext().getTypeAlignInChars(type); - return classifyExpandedType(lowering, forReturn, alignment); + return classifyExpandedType(lowering, forReturn, alignment, IndirectAS); } // Member pointer types need to be expanded, but it's a simple form of From 2bdb085b0debd87651d0c0b81fecff181ee0e541 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 7 Jan 2025 14:17:29 +0200 Subject: [PATCH 17/27] Fix formatting. --- clang/lib/CodeGen/SwiftCallingConv.cpp | 10 +++++----- clang/lib/CodeGen/Targets/AArch64.cpp | 6 +++--- clang/lib/CodeGen/Targets/ARC.cpp | 5 ++--- clang/lib/CodeGen/Targets/BPF.cpp | 3 +-- clang/lib/CodeGen/Targets/CSKY.cpp | 6 +++--- clang/lib/CodeGen/Targets/LoongArch.cpp | 10 +++++----- clang/lib/CodeGen/Targets/RISCV.cpp | 6 +++--- 7 files changed, 22 insertions(+), 24 deletions(-) diff --git a/clang/lib/CodeGen/SwiftCallingConv.cpp b/clang/lib/CodeGen/SwiftCallingConv.cpp index 299a379977704..b91a35ef0e65e 100644 --- a/clang/lib/CodeGen/SwiftCallingConv.cpp +++ b/clang/lib/CodeGen/SwiftCallingConv.cpp @@ -812,16 +812,16 @@ static ABIArgInfo classifyExpandedType(SwiftAggLowering &lowering, static ABIArgInfo classifyType(CodeGenModule &CGM, CanQualType type, bool forReturn) { - unsigned IndirectAS = forReturn - ? CGM.getDataLayout().getAllocaAddrSpace() - : CGM.getContext().getTargetAddressSpace(LangAS::Default); + unsigned IndirectAS = + forReturn ? CGM.getDataLayout().getAllocaAddrSpace() + : CGM.getContext().getTargetAddressSpace(LangAS::Default); if (auto recordType = dyn_cast(type)) { auto record = recordType->getDecl(); auto &layout = CGM.getContext().getASTRecordLayout(record); if (mustPassRecordIndirectly(CGM, record)) - return ABIArgInfo::getIndirect( - layout.getAlignment(), /*AddrSpace=*/ IndirectAS, /*byval=*/ false); + return ABIArgInfo::getIndirect(layout.getAlignment(), + /*AddrSpace=*/IndirectAS, /*byval=*/false); SwiftAggLowering lowering(CGM); lowering.addTypedData(recordType->getDecl(), CharUnits::Zero(), layout); diff --git a/clang/lib/CodeGen/Targets/AArch64.cpp b/clang/lib/CodeGen/Targets/AArch64.cpp index 19520e0229b64..8241c8e1bef50 100644 --- a/clang/lib/CodeGen/Targets/AArch64.cpp +++ b/clang/lib/CodeGen/Targets/AArch64.cpp @@ -414,9 +414,9 @@ ABIArgInfo AArch64ABIInfo::classifyArgumentType(QualType Ty, bool IsVariadicFn, // Structures with either a non-trivial destructor or a non-trivial // copy constructor are always indirect. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { - return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), + /*ByVal=*/RAA == + CGCXXABI::RAA_DirectInMemory); } // Empty records are always ignored on Darwin, but actually passed in C++ mode diff --git a/clang/lib/CodeGen/Targets/ARC.cpp b/clang/lib/CodeGen/Targets/ARC.cpp index 1e5516e1b4a0e..218250484aab5 100644 --- a/clang/lib/CodeGen/Targets/ARC.cpp +++ b/clang/lib/CodeGen/Targets/ARC.cpp @@ -69,9 +69,8 @@ class ARCTargetCodeGenInfo : public TargetCodeGenInfo { ABIArgInfo ARCABIInfo::getIndirectByRef(QualType Ty, bool HasFreeRegs) const { - return HasFreeRegs - ? getNaturalAlignIndirectInReg(Ty) - : getNaturalAlignIndirect(Ty, getTargetDefaultAS(), false); + return HasFreeRegs ? getNaturalAlignIndirectInReg(Ty) + : getNaturalAlignIndirect(Ty, getTargetDefaultAS(), false); } ABIArgInfo ARCABIInfo::getIndirectByValue(QualType Ty) const { diff --git a/clang/lib/CodeGen/Targets/BPF.cpp b/clang/lib/CodeGen/Targets/BPF.cpp index 31ade54f7ef71..1d575f464e1e4 100644 --- a/clang/lib/CodeGen/Targets/BPF.cpp +++ b/clang/lib/CodeGen/Targets/BPF.cpp @@ -42,8 +42,7 @@ class BPFABIInfo : public DefaultABIInfo { } return ABIArgInfo::getDirect(CoerceTy); } else { - return getNaturalAlignIndirect( - Ty, getTargetDefaultAS()); + return getNaturalAlignIndirect(Ty, getTargetDefaultAS()); } } diff --git a/clang/lib/CodeGen/Targets/CSKY.cpp b/clang/lib/CodeGen/Targets/CSKY.cpp index 29b239b24e56b..ce15b068141d0 100644 --- a/clang/lib/CodeGen/Targets/CSKY.cpp +++ b/clang/lib/CodeGen/Targets/CSKY.cpp @@ -82,9 +82,9 @@ ABIArgInfo CSKYABIInfo::classifyArgumentType(QualType Ty, int &ArgGPRsLeft, if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { if (ArgGPRsLeft) ArgGPRsLeft -= 1; - return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), + /*ByVal=*/RAA == + CGCXXABI::RAA_DirectInMemory); } // Ignore empty structs/unions. diff --git a/clang/lib/CodeGen/Targets/LoongArch.cpp b/clang/lib/CodeGen/Targets/LoongArch.cpp index 8edff7b59fee9..691333b2b6f9a 100644 --- a/clang/lib/CodeGen/Targets/LoongArch.cpp +++ b/clang/lib/CodeGen/Targets/LoongArch.cpp @@ -305,9 +305,9 @@ ABIArgInfo LoongArchABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { if (GARsLeft) GARsLeft -= 1; - return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), + /*ByVal=*/RAA == + CGCXXABI::RAA_DirectInMemory); } uint64_t Size = getContext().getTypeSize(Ty); @@ -382,8 +382,8 @@ ABIArgInfo LoongArchABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (EIT->getNumBits() > 128 || (!getContext().getTargetInfo().hasInt128Type() && EIT->getNumBits() > 64)) - return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), + /*ByVal=*/false); } return ABIArgInfo::getDirect(); diff --git a/clang/lib/CodeGen/Targets/RISCV.cpp b/clang/lib/CodeGen/Targets/RISCV.cpp index 7f8df7458fcb3..50802a29da1a4 100644 --- a/clang/lib/CodeGen/Targets/RISCV.cpp +++ b/clang/lib/CodeGen/Targets/RISCV.cpp @@ -410,9 +410,9 @@ ABIArgInfo RISCVABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { if (ArgGPRsLeft) ArgGPRsLeft -= 1; - return getNaturalAlignIndirect( - Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), + /*ByVal=*/RAA == + CGCXXABI::RAA_DirectInMemory); } uint64_t Size = getContext().getTypeSize(Ty); From 4b47cd79cae2d8ffdd63fbe5137224b0455e5526 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Wed, 8 Jan 2025 23:14:02 +0200 Subject: [PATCH 18/27] Remove helper, switch to using the AllocaAS for all indirects. --- clang/lib/CodeGen/ABIInfo.cpp | 4 ---- clang/lib/CodeGen/ABIInfo.h | 2 -- clang/lib/CodeGen/ABIInfoImpl.cpp | 11 ++++------- clang/lib/CodeGen/Targets/AArch64.cpp | 18 ++++++++++------- clang/lib/CodeGen/Targets/ARC.cpp | 6 ++++-- clang/lib/CodeGen/Targets/ARM.cpp | 17 +++++++++------- clang/lib/CodeGen/Targets/BPF.cpp | 6 ++++-- clang/lib/CodeGen/Targets/CSKY.cpp | 9 +++++---- clang/lib/CodeGen/Targets/Hexagon.cpp | 8 +++++--- clang/lib/CodeGen/Targets/Lanai.cpp | 5 +++-- clang/lib/CodeGen/Targets/LoongArch.cpp | 16 ++++++++------- clang/lib/CodeGen/Targets/Mips.cpp | 4 ++-- clang/lib/CodeGen/Targets/NVPTX.cpp | 10 ++++++---- clang/lib/CodeGen/Targets/PNaCl.cpp | 4 ++-- clang/lib/CodeGen/Targets/PPC.cpp | 20 +++++++++++-------- clang/lib/CodeGen/Targets/RISCV.cpp | 16 ++++++++------- clang/lib/CodeGen/Targets/SPIR.cpp | 9 +++++---- clang/lib/CodeGen/Targets/Sparc.cpp | 7 ++++--- clang/lib/CodeGen/Targets/SystemZ.cpp | 11 +++++++---- clang/lib/CodeGen/Targets/WebAssembly.cpp | 2 +- clang/lib/CodeGen/Targets/X86.cpp | 24 ++++++++++++++--------- 21 files changed, 118 insertions(+), 91 deletions(-) diff --git a/clang/lib/CodeGen/ABIInfo.cpp b/clang/lib/CodeGen/ABIInfo.cpp index b7f355eb335f4..19d8c66b000cb 100644 --- a/clang/lib/CodeGen/ABIInfo.cpp +++ b/clang/lib/CodeGen/ABIInfo.cpp @@ -12,10 +12,6 @@ using namespace clang; using namespace clang::CodeGen; -unsigned ABIInfo::getTargetDefaultAS() const { - return getContext().getTargetAddressSpace(LangAS::Default); -} - // Pin the vtable to this file. ABIInfo::~ABIInfo() = default; diff --git a/clang/lib/CodeGen/ABIInfo.h b/clang/lib/CodeGen/ABIInfo.h index 538b4ed26b113..57bcb3dd0a852 100644 --- a/clang/lib/CodeGen/ABIInfo.h +++ b/clang/lib/CodeGen/ABIInfo.h @@ -49,8 +49,6 @@ class ABIInfo { CodeGen::CodeGenTypes &CGT; llvm::CallingConv::ID RuntimeCC; - unsigned getTargetDefaultAS() const; - public: ABIInfo(CodeGen::CodeGenTypes &cgt) : CGT(cgt), RuntimeCC(llvm::CallingConv::C) {} diff --git a/clang/lib/CodeGen/ABIInfoImpl.cpp b/clang/lib/CodeGen/ABIInfoImpl.cpp index 65a186542692d..982563cd7fa9f 100644 --- a/clang/lib/CodeGen/ABIInfoImpl.cpp +++ b/clang/lib/CodeGen/ABIInfoImpl.cpp @@ -21,12 +21,10 @@ ABIArgInfo DefaultABIInfo::classifyArgumentType(QualType Ty) const { // Records with non-trivial destructors/copy-constructors should not be // passed by value. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), - RAA == CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + RAA == CGCXXABI::RAA_DirectInMemory); - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default)); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace()); } // Treat an enum type as its underlying type. @@ -39,8 +37,7 @@ ABIArgInfo DefaultABIInfo::classifyArgumentType(QualType Ty) const { Context.getTypeSize(Context.getTargetInfo().hasInt128Type() ? Context.Int128Ty : Context.LongLongTy)) - return getNaturalAlignIndirect( - Ty, Context.getTargetAddressSpace(LangAS::Default)); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace()); return (isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty, CGT.ConvertType(Ty)) diff --git a/clang/lib/CodeGen/Targets/AArch64.cpp b/clang/lib/CodeGen/Targets/AArch64.cpp index 8241c8e1bef50..cd61aa57d4c37 100644 --- a/clang/lib/CodeGen/Targets/AArch64.cpp +++ b/clang/lib/CodeGen/Targets/AArch64.cpp @@ -326,7 +326,8 @@ ABIArgInfo AArch64ABIInfo::coerceIllegalVector(QualType Ty, unsigned &NSRN, return ABIArgInfo::getDirect(ResType); } - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); } ABIArgInfo AArch64ABIInfo::coerceAndExpandPureScalableAggregate( @@ -334,7 +335,8 @@ ABIArgInfo AArch64ABIInfo::coerceAndExpandPureScalableAggregate( const SmallVectorImpl &UnpaddedCoerceToSeq, unsigned &NSRN, unsigned &NPRN) const { if (!IsNamedArg || NSRN + NVec > 8 || NPRN + NPred > 4) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); NSRN += NVec; NPRN += NPred; @@ -374,7 +376,8 @@ ABIArgInfo AArch64ABIInfo::classifyArgumentType(QualType Ty, bool IsVariadicFn, if (const auto *EIT = Ty->getAs()) if (EIT->getNumBits() > 128) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), false); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + false); if (Ty->isVectorType()) NSRN = std::min(NSRN + 1, 8u); @@ -414,9 +417,9 @@ ABIArgInfo AArch64ABIInfo::classifyArgumentType(QualType Ty, bool IsVariadicFn, // Structures with either a non-trivial destructor or a non-trivial // copy constructor are always indirect. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { - return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/RAA == - CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } // Empty records are always ignored on Darwin, but actually passed in C++ mode @@ -490,7 +493,8 @@ ABIArgInfo AArch64ABIInfo::classifyArgumentType(QualType Ty, bool IsVariadicFn, : llvm::ArrayType::get(BaseTy, Size / Alignment)); } - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); } ABIArgInfo AArch64ABIInfo::classifyReturnType(QualType RetTy, diff --git a/clang/lib/CodeGen/Targets/ARC.cpp b/clang/lib/CodeGen/Targets/ARC.cpp index 218250484aab5..c8db7e8f9706e 100644 --- a/clang/lib/CodeGen/Targets/ARC.cpp +++ b/clang/lib/CodeGen/Targets/ARC.cpp @@ -70,7 +70,8 @@ class ARCTargetCodeGenInfo : public TargetCodeGenInfo { ABIArgInfo ARCABIInfo::getIndirectByRef(QualType Ty, bool HasFreeRegs) const { return HasFreeRegs ? getNaturalAlignIndirectInReg(Ty) - : getNaturalAlignIndirect(Ty, getTargetDefaultAS(), false); + : getNaturalAlignIndirect( + Ty, getDataLayout().getAllocaAddrSpace(), false); } ABIArgInfo ARCABIInfo::getIndirectByValue(QualType Ty) const { @@ -78,7 +79,8 @@ ABIArgInfo ARCABIInfo::getIndirectByValue(QualType Ty) const { const unsigned MinABIStackAlignInBytes = 4; unsigned TypeAlign = getContext().getTypeAlign(Ty) / 8; return ABIArgInfo::getIndirect( - CharUnits::fromQuantity(4), /*AddrSpace=*/getTargetDefaultAS(), + CharUnits::fromQuantity(4), + /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), /*ByVal=*/true, TypeAlign > MinABIStackAlignInBytes); } diff --git a/clang/lib/CodeGen/Targets/ARM.cpp b/clang/lib/CodeGen/Targets/ARM.cpp index e4628a7e9731a..de11c1fd1fd78 100644 --- a/clang/lib/CodeGen/Targets/ARM.cpp +++ b/clang/lib/CodeGen/Targets/ARM.cpp @@ -298,8 +298,9 @@ ABIArgInfo ARMABIInfo::coerceIllegalVector(QualType Ty) const { llvm::Type::getInt32Ty(getVMContext()), Size / 32); return ABIArgInfo::getDirect(ResType); } - return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); } ABIArgInfo ARMABIInfo::classifyHomogeneousAggregate(QualType Ty, @@ -356,8 +357,9 @@ ABIArgInfo ARMABIInfo::classifyArgumentType(QualType Ty, bool isVariadic, if (const auto *EIT = Ty->getAs()) if (EIT->getNumBits() > 64) - return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/true); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/true); return (isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty, CGT.ConvertType(Ty)) @@ -365,7 +367,7 @@ ABIArgInfo ARMABIInfo::classifyArgumentType(QualType Ty, bool isVariadic, } if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), RAA == CGCXXABI::RAA_DirectInMemory); } @@ -401,7 +403,7 @@ ABIArgInfo ARMABIInfo::classifyArgumentType(QualType Ty, bool isVariadic, // and a pointer is passed. return ABIArgInfo::getIndirect( CharUnits::fromQuantity(getContext().getTypeAlign(Ty) / 8), - getTargetDefaultAS(), false); + getDataLayout().getAllocaAddrSpace(), false); } // Support byval for ARM. @@ -420,7 +422,8 @@ ABIArgInfo ARMABIInfo::classifyArgumentType(QualType Ty, bool isVariadic, if (getContext().getTypeSizeInChars(Ty) > CharUnits::fromQuantity(64)) { assert(getABIKind() != ARMABIKind::AAPCS16_VFP && "unexpected byval"); return ABIArgInfo::getIndirect( - CharUnits::fromQuantity(ABIAlign), /*AddrSpace=*/getTargetDefaultAS(), + CharUnits::fromQuantity(ABIAlign), + /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), /*ByVal=*/true, /*Realign=*/TyAlign > ABIAlign); } diff --git a/clang/lib/CodeGen/Targets/BPF.cpp b/clang/lib/CodeGen/Targets/BPF.cpp index 1d575f464e1e4..880a891083c3a 100644 --- a/clang/lib/CodeGen/Targets/BPF.cpp +++ b/clang/lib/CodeGen/Targets/BPF.cpp @@ -42,7 +42,8 @@ class BPFABIInfo : public DefaultABIInfo { } return ABIArgInfo::getDirect(CoerceTy); } else { - return getNaturalAlignIndirect(Ty, getTargetDefaultAS()); + return getNaturalAlignIndirect(Ty, + getDataLayout().getAllocaAddrSpace()); } } @@ -52,7 +53,8 @@ class BPFABIInfo : public DefaultABIInfo { ASTContext &Context = getContext(); if (const auto *EIT = Ty->getAs()) if (EIT->getNumBits() > Context.getTypeSize(Context.Int128Ty)) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS()); + return getNaturalAlignIndirect(Ty, + getDataLayout().getAllocaAddrSpace()); return (isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty) : ABIArgInfo::getDirect()); diff --git a/clang/lib/CodeGen/Targets/CSKY.cpp b/clang/lib/CodeGen/Targets/CSKY.cpp index ce15b068141d0..ef26d483a180a 100644 --- a/clang/lib/CodeGen/Targets/CSKY.cpp +++ b/clang/lib/CodeGen/Targets/CSKY.cpp @@ -82,9 +82,9 @@ ABIArgInfo CSKYABIInfo::classifyArgumentType(QualType Ty, int &ArgGPRsLeft, if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { if (ArgGPRsLeft) ArgGPRsLeft -= 1; - return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/RAA == - CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } // Ignore empty structs/unions. @@ -145,7 +145,8 @@ ABIArgInfo CSKYABIInfo::classifyArgumentType(QualType Ty, int &ArgGPRsLeft, llvm::IntegerType::get(getVMContext(), XLen), (Size + 31) / XLen)); } } - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); } ABIArgInfo CSKYABIInfo::classifyReturnType(QualType RetTy) const { diff --git a/clang/lib/CodeGen/Targets/Hexagon.cpp b/clang/lib/CodeGen/Targets/Hexagon.cpp index 38d1c0232d6e1..667599d2d9a66 100644 --- a/clang/lib/CodeGen/Targets/Hexagon.cpp +++ b/clang/lib/CodeGen/Targets/Hexagon.cpp @@ -105,14 +105,15 @@ ABIArgInfo HexagonABIInfo::classifyArgumentType(QualType Ty, HexagonAdjustRegsLeft(Size, RegsLeft); if (Size > 64 && Ty->isBitIntType()) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/true); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/true); return isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty) : ABIArgInfo::getDirect(); } if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), RAA == CGCXXABI::RAA_DirectInMemory); // Ignore empty records. @@ -123,7 +124,8 @@ ABIArgInfo HexagonABIInfo::classifyArgumentType(QualType Ty, unsigned Align = getContext().getTypeAlign(Ty); if (Size > 64) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/true); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/true); if (HexagonAdjustRegsLeft(Size, RegsLeft)) Align = Size <= 32 ? 32 : 64; diff --git a/clang/lib/CodeGen/Targets/Lanai.cpp b/clang/lib/CodeGen/Targets/Lanai.cpp index 4ea078a9e48af..6f75bd54a8ef2 100644 --- a/clang/lib/CodeGen/Targets/Lanai.cpp +++ b/clang/lib/CodeGen/Targets/Lanai.cpp @@ -94,8 +94,9 @@ ABIArgInfo LanaiABIInfo::classifyArgumentType(QualType Ty, if (RAA == CGCXXABI::RAA_Indirect) { return getIndirectResult(Ty, /*ByVal=*/false, State); } else if (RAA == CGCXXABI::RAA_DirectInMemory) { - return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/true); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/true); } } diff --git a/clang/lib/CodeGen/Targets/LoongArch.cpp b/clang/lib/CodeGen/Targets/LoongArch.cpp index 691333b2b6f9a..f1d972b3d4ffc 100644 --- a/clang/lib/CodeGen/Targets/LoongArch.cpp +++ b/clang/lib/CodeGen/Targets/LoongArch.cpp @@ -305,9 +305,9 @@ ABIArgInfo LoongArchABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { if (GARsLeft) GARsLeft -= 1; - return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/RAA == - CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } uint64_t Size = getContext().getTypeSize(Ty); @@ -382,8 +382,9 @@ ABIArgInfo LoongArchABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (EIT->getNumBits() > 128 || (!getContext().getTargetInfo().hasInt128Type() && EIT->getNumBits() > 64)) - return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); } return ABIArgInfo::getDirect(); @@ -406,8 +407,9 @@ ABIArgInfo LoongArchABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, return ABIArgInfo::getDirect( llvm::ArrayType::get(llvm::IntegerType::get(getVMContext(), GRLen), 2)); } - return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); } ABIArgInfo LoongArchABIInfo::classifyReturnType(QualType RetTy) const { diff --git a/clang/lib/CodeGen/Targets/Mips.cpp b/clang/lib/CodeGen/Targets/Mips.cpp index 49613c3341b1a..2ba0b36330619 100644 --- a/clang/lib/CodeGen/Targets/Mips.cpp +++ b/clang/lib/CodeGen/Targets/Mips.cpp @@ -209,7 +209,7 @@ MipsABIInfo::classifyArgumentType(QualType Ty, uint64_t &Offset) const { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { Offset = OrigOffset + MinABIStackAlignInBytes; - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), RAA == CGCXXABI::RAA_DirectInMemory); } @@ -232,7 +232,7 @@ MipsABIInfo::classifyArgumentType(QualType Ty, uint64_t &Offset) const { if (EIT->getNumBits() > 128 || (EIT->getNumBits() > 64 && !getContext().getTargetInfo().hasInt128Type())) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS()); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace()); // All integral types are promoted to the GPR width. if (Ty->isIntegralOrEnumerationType()) diff --git a/clang/lib/CodeGen/Targets/NVPTX.cpp b/clang/lib/CodeGen/Targets/NVPTX.cpp index 1920a2ff4aecf..c236de8db70e4 100644 --- a/clang/lib/CodeGen/Targets/NVPTX.cpp +++ b/clang/lib/CodeGen/Targets/NVPTX.cpp @@ -192,16 +192,18 @@ ABIArgInfo NVPTXABIInfo::classifyArgumentType(QualType Ty) const { return ABIArgInfo::getDirect( CGInfo.getCUDADeviceBuiltinTextureDeviceType()); } - return getNaturalAlignIndirect(Ty, /* AddrSpace */ getTargetDefaultAS(), - /* byval */ true); + return getNaturalAlignIndirect( + Ty, /* AddrSpace */ getDataLayout().getAllocaAddrSpace(), + /* byval */ true); } if (const auto *EIT = Ty->getAs()) { if ((EIT->getNumBits() > 128) || (!getContext().getTargetInfo().hasInt128Type() && EIT->getNumBits() > 64)) - return getNaturalAlignIndirect(Ty, /* AddrSpace */ getTargetDefaultAS(), - /* byval */ true); + return getNaturalAlignIndirect( + Ty, /* AddrSpace */ getDataLayout().getAllocaAddrSpace(), + /* byval */ true); } return (isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty) diff --git a/clang/lib/CodeGen/Targets/PNaCl.cpp b/clang/lib/CodeGen/Targets/PNaCl.cpp index 85c2743a16985..abe95c01cda70 100644 --- a/clang/lib/CodeGen/Targets/PNaCl.cpp +++ b/clang/lib/CodeGen/Targets/PNaCl.cpp @@ -63,7 +63,7 @@ RValue PNaClABIInfo::EmitVAArg(CodeGenFunction &CGF, Address VAListAddr, ABIArgInfo PNaClABIInfo::classifyArgumentType(QualType Ty) const { if (isAggregateTypeForABI(Ty)) { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), RAA == CGCXXABI::RAA_DirectInMemory); return getNaturalAlignIndirect( Ty, getContext().getTargetAddressSpace(LangAS::Default)); @@ -77,7 +77,7 @@ ABIArgInfo PNaClABIInfo::classifyArgumentType(QualType Ty) const { // Treat bit-precise integers as integers if <= 64, otherwise pass // indirectly. if (EIT->getNumBits() > 64) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS()); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace()); return ABIArgInfo::getDirect(); } diff --git a/clang/lib/CodeGen/Targets/PPC.cpp b/clang/lib/CodeGen/Targets/PPC.cpp index 2acaf181677af..b03eb8c9d5035 100644 --- a/clang/lib/CodeGen/Targets/PPC.cpp +++ b/clang/lib/CodeGen/Targets/PPC.cpp @@ -208,15 +208,16 @@ ABIArgInfo AIXABIInfo::classifyArgumentType(QualType Ty) const { // Records with non-trivial destructors/copy-constructors should not be // passed by value. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), RAA == CGCXXABI::RAA_DirectInMemory); CharUnits CCAlign = getParamTypeAlignment(Ty); CharUnits TyAlign = getContext().getTypeAlignInChars(Ty); - return ABIArgInfo::getIndirect(CCAlign, /*AddrSpace*/ getTargetDefaultAS(), - /*ByVal*/ true, - /*Realign*/ TyAlign > CCAlign); + return ABIArgInfo::getIndirect( + CCAlign, /*AddrSpace*/ getDataLayout().getAllocaAddrSpace(), + /*ByVal*/ true, + /*Realign*/ TyAlign > CCAlign); } return (isPromotableTypeForABI(Ty) @@ -835,7 +836,8 @@ PPC64_SVR4_ABIInfo::classifyArgumentType(QualType Ty) const { if (Ty->isVectorType()) { uint64_t Size = getContext().getTypeSize(Ty); if (Size > 128) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); else if (Size < 128) { llvm::Type *CoerceTy = llvm::IntegerType::get(getVMContext(), Size); return ABIArgInfo::getDirect(CoerceTy); @@ -844,11 +846,12 @@ PPC64_SVR4_ABIInfo::classifyArgumentType(QualType Ty) const { if (const auto *EIT = Ty->getAs()) if (EIT->getNumBits() > 128) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/true); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/true); if (isAggregateTypeForABI(Ty)) { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), RAA == CGCXXABI::RAA_DirectInMemory); uint64_t ABIAlign = getParamTypeAlignment(Ty).getQuantity(); @@ -891,7 +894,8 @@ PPC64_SVR4_ABIInfo::classifyArgumentType(QualType Ty) const { // All other aggregates are passed ByVal. return ABIArgInfo::getIndirect( - CharUnits::fromQuantity(ABIAlign), /*AddrSpace=*/getTargetDefaultAS(), + CharUnits::fromQuantity(ABIAlign), + /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), /*ByVal=*/true, /*Realign=*/TyAlign > ABIAlign); } diff --git a/clang/lib/CodeGen/Targets/RISCV.cpp b/clang/lib/CodeGen/Targets/RISCV.cpp index 50802a29da1a4..cdf9cf4b8f3dc 100644 --- a/clang/lib/CodeGen/Targets/RISCV.cpp +++ b/clang/lib/CodeGen/Targets/RISCV.cpp @@ -410,9 +410,9 @@ ABIArgInfo RISCVABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) { if (ArgGPRsLeft) ArgGPRsLeft -= 1; - return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/RAA == - CGCXXABI::RAA_DirectInMemory); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/RAA == CGCXXABI::RAA_DirectInMemory); } uint64_t Size = getContext().getTypeSize(Ty); @@ -493,8 +493,9 @@ ABIArgInfo RISCVABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, if (EIT->getNumBits() > 128 || (!getContext().getTargetInfo().hasInt128Type() && EIT->getNumBits() > 64)) - return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); } return ABIArgInfo::getDirect(); @@ -526,8 +527,9 @@ ABIArgInfo RISCVABIInfo::classifyArgumentType(QualType Ty, bool IsFixed, llvm::IntegerType::get(getVMContext(), XLen), 2)); } } - return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); } ABIArgInfo RISCVABIInfo::classifyReturnType(QualType RetTy) const { diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp index 2bc7dfc62ae44..56e59306f2382 100644 --- a/clang/lib/CodeGen/Targets/SPIR.cpp +++ b/clang/lib/CodeGen/Targets/SPIR.cpp @@ -156,9 +156,10 @@ ABIArgInfo SPIRVABIInfo::classifyKernelArgumentType(QualType Ty) const { // copied to be valid on the device. // This behavior follows the CUDA spec // https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-function-argument-processing, - // and matches the NVPTX implementation. - return getNaturalAlignIndirect(Ty, /* AddrSpace */ getTargetDefaultAS(), - /* byval */ true); + // and matches the NVPTX implementation. TODO: hardcoding to 0 should be + // revisited if HIPSPV / byval starts making use of the AS of an indirect + // arg. + return getNaturalAlignIndirect(Ty, /* AddrSpace */ 0, /* byval */ true); } } return classifyArgumentType(Ty); @@ -173,7 +174,7 @@ ABIArgInfo SPIRVABIInfo::classifyArgumentType(QualType Ty) const { // Records with non-trivial destructors/copy-constructors should not be // passed by value. if (auto RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), RAA == CGCXXABI::RAA_DirectInMemory); if (const RecordType *RT = Ty->getAs()) { diff --git a/clang/lib/CodeGen/Targets/Sparc.cpp b/clang/lib/CodeGen/Targets/Sparc.cpp index d0b40aa9ceab1..9642196b78c63 100644 --- a/clang/lib/CodeGen/Targets/Sparc.cpp +++ b/clang/lib/CodeGen/Targets/Sparc.cpp @@ -232,8 +232,9 @@ SparcV9ABIInfo::classifyType(QualType Ty, unsigned SizeLimit) const { // Anything too big to fit in registers is passed with an explicit indirect // pointer / sret pointer. if (Size > SizeLimit) - return getNaturalAlignIndirect(Ty, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/false); + return getNaturalAlignIndirect( + Ty, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); // Treat an enum type as its underlying type. if (const EnumType *EnumTy = Ty->getAs()) @@ -254,7 +255,7 @@ SparcV9ABIInfo::classifyType(QualType Ty, unsigned SizeLimit) const { // If a C++ object has either a non-trivial copy constructor or a non-trivial // destructor, it is passed with an explicit indirect pointer / sret pointer. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), RAA == CGCXXABI::RAA_DirectInMemory); // This is a small aggregate type that should be passed in registers. diff --git a/clang/lib/CodeGen/Targets/SystemZ.cpp b/clang/lib/CodeGen/Targets/SystemZ.cpp index 4fd141c694c8b..c928d3c029caa 100644 --- a/clang/lib/CodeGen/Targets/SystemZ.cpp +++ b/clang/lib/CodeGen/Targets/SystemZ.cpp @@ -417,7 +417,7 @@ ABIArgInfo SystemZABIInfo::classifyArgumentType(QualType Ty) const { // Handle the generic C++ ABI. if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), RAA == CGCXXABI::RAA_DirectInMemory); // Integers and enums are extended to full register width. @@ -435,7 +435,8 @@ ABIArgInfo SystemZABIInfo::classifyArgumentType(QualType Ty) const { // Values that are not 1, 2, 4 or 8 bytes in size are passed indirectly. if (Size != 8 && Size != 16 && Size != 32 && Size != 64) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); // Handle small structures. if (const RecordType *RT = Ty->getAs()) { @@ -443,7 +444,8 @@ ABIArgInfo SystemZABIInfo::classifyArgumentType(QualType Ty) const { // fail the size test above. const RecordDecl *RD = RT->getDecl(); if (RD->hasFlexibleArrayMember()) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); // The structure is passed as an unextended integer, a float, or a double. if (isFPArgumentType(SingleElementTy)) { @@ -460,7 +462,8 @@ ABIArgInfo SystemZABIInfo::classifyArgumentType(QualType Ty) const { // Non-structure compounds are passed indirectly. if (isCompoundType(Ty)) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); return ABIArgInfo::getDirect(nullptr); } diff --git a/clang/lib/CodeGen/Targets/WebAssembly.cpp b/clang/lib/CodeGen/Targets/WebAssembly.cpp index 0601213483864..9217c78a540a3 100644 --- a/clang/lib/CodeGen/Targets/WebAssembly.cpp +++ b/clang/lib/CodeGen/Targets/WebAssembly.cpp @@ -103,7 +103,7 @@ ABIArgInfo WebAssemblyABIInfo::classifyArgumentType(QualType Ty) const { // Records with non-trivial destructors/copy-constructors should not be // passed by value. if (auto RAA = getRecordArgABI(Ty, getCXXABI())) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), RAA == CGCXXABI::RAA_DirectInMemory); // Ignore empty structs/unions. if (isEmptyRecord(getContext(), Ty, true)) diff --git a/clang/lib/CodeGen/Targets/X86.cpp b/clang/lib/CodeGen/Targets/X86.cpp index 53e1df25522db..329055706d3a8 100644 --- a/clang/lib/CodeGen/Targets/X86.cpp +++ b/clang/lib/CodeGen/Targets/X86.cpp @@ -3293,7 +3293,7 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, if (RT) { if (!IsReturnType) { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(RT, getCXXABI())) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), RAA == CGCXXABI::RAA_DirectInMemory); } @@ -3316,8 +3316,9 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, return ABIArgInfo::getDirect(); return ABIArgInfo::getExpand(); } - return ABIArgInfo::getIndirect(Align, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/false); + return ABIArgInfo::getIndirect( + Align, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); } else if (IsVectorCall) { if (FreeSSERegs >= NumElts && (IsReturnType || Ty->isBuiltinType() || Ty->isVectorType())) { @@ -3328,7 +3329,8 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, } else if (!Ty->isBuiltinType() && !Ty->isVectorType()) { // HVAs are delayed and reclassified in the 2nd step. return ABIArgInfo::getIndirect( - Align, /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/false); + Align, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); } } } @@ -3345,7 +3347,8 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, // MS x64 ABI requirement: "Any argument that doesn't fit in 8 bytes, or is // not 1, 2, 4, or 8 bytes, must be passed by reference." if (Width > 64 || !llvm::isPowerOf2_64(Width)) - return getNaturalAlignIndirect(Ty, getTargetDefaultAS(), /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); // Otherwise, coerce it to a small integer. return ABIArgInfo::getDirect(llvm::IntegerType::get(getVMContext(), Width)); @@ -3365,7 +3368,8 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, const llvm::fltSemantics *LDF = &getTarget().getLongDoubleFormat(); if (LDF == &llvm::APFloat::x87DoubleExtended()) return ABIArgInfo::getIndirect( - Align, /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/false); + Align, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); } break; @@ -3376,7 +3380,8 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, // even though it isn't particularly efficient. if (!IsReturnType) return ABIArgInfo::getIndirect( - Align, /*AddrSpace=*/getTargetDefaultAS(), /*ByVal=*/false); + Align, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); // Mingw64 GCC returns i128 in XMM0. Coerce to v2i64 to handle that. // Clang matches them for compatibility. @@ -3396,8 +3401,9 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, // the power of 2. if (Width <= 64) return ABIArgInfo::getDirect(); - return ABIArgInfo::getIndirect(Align, /*AddrSpace=*/getTargetDefaultAS(), - /*ByVal=*/false); + return ABIArgInfo::getIndirect( + Align, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/false); } return ABIArgInfo::getDirect(); From d1032557e27e0750977e2a78f3ffaa77d7cb80ad Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Wed, 8 Jan 2025 23:36:53 +0200 Subject: [PATCH 19/27] Fix Swift mismatch. --- clang/lib/CodeGen/SwiftCallingConv.cpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/SwiftCallingConv.cpp b/clang/lib/CodeGen/SwiftCallingConv.cpp index b91a35ef0e65e..6f10ff96d3a94 100644 --- a/clang/lib/CodeGen/SwiftCallingConv.cpp +++ b/clang/lib/CodeGen/SwiftCallingConv.cpp @@ -812,9 +812,7 @@ static ABIArgInfo classifyExpandedType(SwiftAggLowering &lowering, static ABIArgInfo classifyType(CodeGenModule &CGM, CanQualType type, bool forReturn) { - unsigned IndirectAS = - forReturn ? CGM.getDataLayout().getAllocaAddrSpace() - : CGM.getContext().getTargetAddressSpace(LangAS::Default); + unsigned IndirectAS = CGM.getDataLayout().getAllocaAddrSpace(); if (auto recordType = dyn_cast(type)) { auto record = recordType->getDecl(); auto &layout = CGM.getContext().getASTRecordLayout(record); From 5227aefb2208672b2ecbaa5d703b8c705bce9351 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 23 Jan 2025 01:12:19 +0000 Subject: [PATCH 20/27] Fix leftover LangAS::Default. --- clang/lib/CodeGen/Targets/X86.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/Targets/X86.cpp b/clang/lib/CodeGen/Targets/X86.cpp index c0c0f028da271..0bef50117e8cf 100644 --- a/clang/lib/CodeGen/Targets/X86.cpp +++ b/clang/lib/CodeGen/Targets/X86.cpp @@ -3299,7 +3299,7 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, if (RT->getDecl()->hasFlexibleArrayMember()) return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default), + Ty, getDataLayout().getAllocaAddressSpace(), /*ByVal=*/false); } From 94b51d565f217816f927be2a952b633ec00573fa Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 23 Jan 2025 01:16:48 +0000 Subject: [PATCH 21/27] Fix leftover use of LangAS::Default. --- clang/lib/CodeGen/Targets/PNaCl.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/Targets/PNaCl.cpp b/clang/lib/CodeGen/Targets/PNaCl.cpp index abe95c01cda70..358010785850e 100644 --- a/clang/lib/CodeGen/Targets/PNaCl.cpp +++ b/clang/lib/CodeGen/Targets/PNaCl.cpp @@ -65,8 +65,7 @@ ABIArgInfo PNaClABIInfo::classifyArgumentType(QualType Ty) const { if (CGCXXABI::RecordArgABI RAA = getRecordArgABI(Ty, getCXXABI())) return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), RAA == CGCXXABI::RAA_DirectInMemory); - return getNaturalAlignIndirect( - Ty, getContext().getTargetAddressSpace(LangAS::Default)); + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace()); } else if (const EnumType *EnumTy = Ty->getAs()) { // Treat an enum type as its underlying type. Ty = EnumTy->getDecl()->getIntegerType(); From 53d8462f4f6c43e9005847ec9b15444a73fe7e02 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 23 Jan 2025 01:18:29 +0000 Subject: [PATCH 22/27] Apply formatting suggestions. --- clang/lib/CodeGen/SwiftCallingConv.cpp | 4 ++-- clang/lib/CodeGen/Targets/PPC.cpp | 6 +++--- clang/lib/CodeGen/Targets/SPIR.cpp | 2 +- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/clang/lib/CodeGen/SwiftCallingConv.cpp b/clang/lib/CodeGen/SwiftCallingConv.cpp index 4e656d921e9a5..dd98053fa9330 100644 --- a/clang/lib/CodeGen/SwiftCallingConv.cpp +++ b/clang/lib/CodeGen/SwiftCallingConv.cpp @@ -802,8 +802,8 @@ static ABIArgInfo classifyExpandedType(SwiftAggLowering &lowering, return ABIArgInfo::getIgnore(); } else if (lowering.shouldPassIndirectly(forReturn)) { return ABIArgInfo::getIndirect(alignmentForIndirect, - /*AddrSpace*/ IndirectAS, - /*byval*/ false); + /*AddrSpace=*/ IndirectAS, + /*byval=*/ false); } else { auto types = lowering.getCoerceAndExpandTypes(); return ABIArgInfo::getCoerceAndExpand(types.first, types.second); diff --git a/clang/lib/CodeGen/Targets/PPC.cpp b/clang/lib/CodeGen/Targets/PPC.cpp index b03eb8c9d5035..f14f356813765 100644 --- a/clang/lib/CodeGen/Targets/PPC.cpp +++ b/clang/lib/CodeGen/Targets/PPC.cpp @@ -215,9 +215,9 @@ ABIArgInfo AIXABIInfo::classifyArgumentType(QualType Ty) const { CharUnits TyAlign = getContext().getTypeAlignInChars(Ty); return ABIArgInfo::getIndirect( - CCAlign, /*AddrSpace*/ getDataLayout().getAllocaAddrSpace(), - /*ByVal*/ true, - /*Realign*/ TyAlign > CCAlign); + CCAlign, /*AddrSpace=*/ getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/ true, + /*Realign=*/ TyAlign > CCAlign); } return (isPromotableTypeForABI(Ty) diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp index 56e59306f2382..d475cf56c19e0 100644 --- a/clang/lib/CodeGen/Targets/SPIR.cpp +++ b/clang/lib/CodeGen/Targets/SPIR.cpp @@ -159,7 +159,7 @@ ABIArgInfo SPIRVABIInfo::classifyKernelArgumentType(QualType Ty) const { // and matches the NVPTX implementation. TODO: hardcoding to 0 should be // revisited if HIPSPV / byval starts making use of the AS of an indirect // arg. - return getNaturalAlignIndirect(Ty, /* AddrSpace */ 0, /* byval */ true); + return getNaturalAlignIndirect(Ty, /*AddrSpace=*/ 0, /*byval=*/ true); } } return classifyArgumentType(Ty); From 4d2b9f7904b255094ec89fd821c30c7b9e9d7546 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 23 Jan 2025 01:30:52 +0000 Subject: [PATCH 23/27] Fix formatting. --- clang/lib/CodeGen/SwiftCallingConv.cpp | 4 ++-- clang/lib/CodeGen/Targets/PPC.cpp | 6 +++--- clang/lib/CodeGen/Targets/SPIR.cpp | 2 +- clang/lib/CodeGen/Targets/X86.cpp | 6 +++--- 4 files changed, 9 insertions(+), 9 deletions(-) diff --git a/clang/lib/CodeGen/SwiftCallingConv.cpp b/clang/lib/CodeGen/SwiftCallingConv.cpp index dd98053fa9330..10f9f20bca313 100644 --- a/clang/lib/CodeGen/SwiftCallingConv.cpp +++ b/clang/lib/CodeGen/SwiftCallingConv.cpp @@ -802,8 +802,8 @@ static ABIArgInfo classifyExpandedType(SwiftAggLowering &lowering, return ABIArgInfo::getIgnore(); } else if (lowering.shouldPassIndirectly(forReturn)) { return ABIArgInfo::getIndirect(alignmentForIndirect, - /*AddrSpace=*/ IndirectAS, - /*byval=*/ false); + /*AddrSpace=*/IndirectAS, + /*byval=*/false); } else { auto types = lowering.getCoerceAndExpandTypes(); return ABIArgInfo::getCoerceAndExpand(types.first, types.second); diff --git a/clang/lib/CodeGen/Targets/PPC.cpp b/clang/lib/CodeGen/Targets/PPC.cpp index f14f356813765..4df4c9f3c3475 100644 --- a/clang/lib/CodeGen/Targets/PPC.cpp +++ b/clang/lib/CodeGen/Targets/PPC.cpp @@ -215,9 +215,9 @@ ABIArgInfo AIXABIInfo::classifyArgumentType(QualType Ty) const { CharUnits TyAlign = getContext().getTypeAlignInChars(Ty); return ABIArgInfo::getIndirect( - CCAlign, /*AddrSpace=*/ getDataLayout().getAllocaAddrSpace(), - /*ByVal=*/ true, - /*Realign=*/ TyAlign > CCAlign); + CCAlign, /*AddrSpace=*/getDataLayout().getAllocaAddrSpace(), + /*ByVal=*/true, + /*Realign=*/TyAlign > CCAlign); } return (isPromotableTypeForABI(Ty) diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp index d475cf56c19e0..b81ed29a5159b 100644 --- a/clang/lib/CodeGen/Targets/SPIR.cpp +++ b/clang/lib/CodeGen/Targets/SPIR.cpp @@ -159,7 +159,7 @@ ABIArgInfo SPIRVABIInfo::classifyKernelArgumentType(QualType Ty) const { // and matches the NVPTX implementation. TODO: hardcoding to 0 should be // revisited if HIPSPV / byval starts making use of the AS of an indirect // arg. - return getNaturalAlignIndirect(Ty, /*AddrSpace=*/ 0, /*byval=*/ true); + return getNaturalAlignIndirect(Ty, /*AddrSpace=*/0, /*byval=*/true); } } return classifyArgumentType(Ty); diff --git a/clang/lib/CodeGen/Targets/X86.cpp b/clang/lib/CodeGen/Targets/X86.cpp index 0bef50117e8cf..f2a71184c283b 100644 --- a/clang/lib/CodeGen/Targets/X86.cpp +++ b/clang/lib/CodeGen/Targets/X86.cpp @@ -3298,9 +3298,9 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, } if (RT->getDecl()->hasFlexibleArrayMember()) - return getNaturalAlignIndirect( - Ty, getDataLayout().getAllocaAddressSpace(), - /*ByVal=*/false); + return getNaturalAlignIndirect(Ty, + getDataLayout().getAllocaAddressSpace(), + /*ByVal=*/false); } const Type *Base = nullptr; From 3acc4ffbf8c609e86e3f996ddc4bcef6c93d5a3c Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 23 Jan 2025 17:51:54 +0000 Subject: [PATCH 24/27] Fix typo. --- clang/lib/CodeGen/Targets/X86.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/Targets/X86.cpp b/clang/lib/CodeGen/Targets/X86.cpp index f2a71184c283b..704810a21209a 100644 --- a/clang/lib/CodeGen/Targets/X86.cpp +++ b/clang/lib/CodeGen/Targets/X86.cpp @@ -3299,7 +3299,7 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, if (RT->getDecl()->hasFlexibleArrayMember()) return getNaturalAlignIndirect(Ty, - getDataLayout().getAllocaAddressSpace(), + getDataLayout().getAllocaAddrSpace(), /*ByVal=*/false); } From 69b7937f7f04a2180bc15f4df3a24adb271c2721 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 23 Jan 2025 18:25:05 +0000 Subject: [PATCH 25/27] Add test. --- ...plicit-addrspacecast-function-parameter.cl | 68 +++++++++++++++++++ 1 file changed, 68 insertions(+) create mode 100644 clang/test/CodeGenOpenCL/implicit-addrspacecast-function-parameter.cl diff --git a/clang/test/CodeGenOpenCL/implicit-addrspacecast-function-parameter.cl b/clang/test/CodeGenOpenCL/implicit-addrspacecast-function-parameter.cl new file mode 100644 index 0000000000000..997c8a4a5e5cd --- /dev/null +++ b/clang/test/CodeGenOpenCL/implicit-addrspacecast-function-parameter.cl @@ -0,0 +1,68 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5 +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s + +// Check there's no assertion when passing a pointer to an address space +// qualified argument. + +extern void private_ptr(__private int *); +extern void local_ptr(__local int *); +extern void generic_ptr(__generic int *); + +// CHECK-LABEL: define dso_local void @use_of_private_var( +// CHECK-SAME: ) #[[ATTR0:[0-9]+]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[X:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr +// CHECK-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[X]]) #[[ATTR4:[0-9]+]] +// CHECK-NEXT: store i32 0, ptr [[X_ASCAST]], align 4, !tbaa [[TBAA4:![0-9]+]] +// CHECK-NEXT: [[TMP0:%.*]] = addrspacecast ptr [[X_ASCAST]] to ptr addrspace(5) +// CHECK-NEXT: call void @private_ptr(ptr addrspace(5) noundef [[TMP0]]) #[[ATTR5:[0-9]+]] +// CHECK-NEXT: call void @generic_ptr(ptr noundef [[X_ASCAST]]) #[[ATTR5]] +// CHECK-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[X]]) #[[ATTR4]] +// CHECK-NEXT: ret void +// +void use_of_private_var() +{ + int x = 0 ; + private_ptr(&x); + generic_ptr(&x); +} + +// CHECK-LABEL: define dso_local void @addr_of_arg( +// CHECK-SAME: i32 noundef [[X:%.*]]) #[[ATTR0]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[X_ADDR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr +// CHECK-NEXT: store i32 [[X]], ptr [[X_ADDR_ASCAST]], align 4, !tbaa [[TBAA4]] +// CHECK-NEXT: [[TMP0:%.*]] = addrspacecast ptr [[X_ADDR_ASCAST]] to ptr addrspace(5) +// CHECK-NEXT: call void @private_ptr(ptr addrspace(5) noundef [[TMP0]]) #[[ATTR5]] +// CHECK-NEXT: call void @generic_ptr(ptr noundef [[X_ADDR_ASCAST]]) #[[ATTR5]] +// CHECK-NEXT: ret void +// +void addr_of_arg(int x) +{ + private_ptr(&x); + generic_ptr(&x); +} + +// CHECK-LABEL: define dso_local amdgpu_kernel void @use_of_local_var( +// CHECK-SAME: ) #[[ATTR3:[0-9]+]] !kernel_arg_addr_space [[META8:![0-9]+]] !kernel_arg_access_qual [[META8]] !kernel_arg_type [[META8]] !kernel_arg_base_type [[META8]] !kernel_arg_type_qual [[META8]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: call void @local_ptr(ptr addrspace(3) noundef @use_of_local_var.x) #[[ATTR5]] +// CHECK-NEXT: call void @generic_ptr(ptr noundef addrspacecast (ptr addrspace(3) @use_of_local_var.x to ptr)) #[[ATTR5]] +// CHECK-NEXT: ret void +// +__kernel void use_of_local_var() +{ + __local int x; + local_ptr(&x); + generic_ptr(&x); +} + +//. +// CHECK: [[TBAA4]] = !{[[META5:![0-9]+]], [[META5]], i64 0} +// CHECK: [[META5]] = !{!"int", [[META6:![0-9]+]], i64 0} +// CHECK: [[META6]] = !{!"omnipotent char", [[META7:![0-9]+]], i64 0} +// CHECK: [[META7]] = !{!"Simple C/C++ TBAA"} +// CHECK: [[META8]] = !{} +//. \ No newline at end of file From ddaccb83ce828fae83835f89d611cb82c39e14c0 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 23 Jan 2025 18:33:49 +0000 Subject: [PATCH 26/27] Fix formatting (again). --- clang/lib/CodeGen/Targets/X86.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/Targets/X86.cpp b/clang/lib/CodeGen/Targets/X86.cpp index 704810a21209a..b3664513c119d 100644 --- a/clang/lib/CodeGen/Targets/X86.cpp +++ b/clang/lib/CodeGen/Targets/X86.cpp @@ -3298,8 +3298,7 @@ ABIArgInfo WinX86_64ABIInfo::classify(QualType Ty, unsigned &FreeSSERegs, } if (RT->getDecl()->hasFlexibleArrayMember()) - return getNaturalAlignIndirect(Ty, - getDataLayout().getAllocaAddrSpace(), + return getNaturalAlignIndirect(Ty, getDataLayout().getAllocaAddrSpace(), /*ByVal=*/false); } From 3e10da36c12606dc6b56524ca745caf001f6c65d Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Thu, 13 Feb 2025 12:58:50 +0700 Subject: [PATCH 27/27] Update clang/test/CodeGenOpenCL/implicit-addrspacecast-function-parameter.cl --- .../CodeGenOpenCL/implicit-addrspacecast-function-parameter.cl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/CodeGenOpenCL/implicit-addrspacecast-function-parameter.cl b/clang/test/CodeGenOpenCL/implicit-addrspacecast-function-parameter.cl index 997c8a4a5e5cd..4a7bb8227c339 100644 --- a/clang/test/CodeGenOpenCL/implicit-addrspacecast-function-parameter.cl +++ b/clang/test/CodeGenOpenCL/implicit-addrspacecast-function-parameter.cl @@ -65,4 +65,4 @@ __kernel void use_of_local_var() // CHECK: [[META6]] = !{!"omnipotent char", [[META7:![0-9]+]], i64 0} // CHECK: [[META7]] = !{!"Simple C/C++ TBAA"} // CHECK: [[META8]] = !{} -//. \ No newline at end of file +//.