[clang] [clang][CodeGen] `sret` args should always point to the `alloca` AS, so use that (PR #114062)
Alex Voicu via cfe-commits
cfe-commits at lists.llvm.org
Thu Jan 23 10:27:17 PST 2025
https://github.com/AlexVlx updated https://github.com/llvm/llvm-project/pull/114062
>From d2d2d3d5db3f639aab178f9ca9a20db2842d2b65 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu at amd.com>
Date: Tue, 29 Oct 2024 14:20:44 +0000
Subject: [PATCH 01/25] `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 8f4f5d3ed81601..56acfae7ae9e51 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<CallLifetimeEnd>(NormalEHLifetimeMarker, SRetAlloca,
+ pushFullExprCleanup<CallLifetimeEnd>(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 e709c1d4ad1ee1..7949a69555031e 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 6d905e89b2c6fd..3b4914f29d2bfe 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 <alexandru.voicu at amd.com>
Date: Tue, 29 Oct 2024 17:16:17 +0000
Subject: [PATCH 02/25] 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 57d056b0ff9d51..4a1db2c3564a57 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 084281a8cada46..c2b2e00d15e13f 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 <alexandru.voicu at amd.com>
Date: Wed, 30 Oct 2024 00:10:59 +0000
Subject: [PATCH 03/25] 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 56acfae7ae9e51..7171d85b0d0ab0 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<llvm::Argument>(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 750392a43e05cc..098163f957f759 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 <alexandru.voicu at amd.com>
Date: Sat, 2 Nov 2024 00:57:17 +0000
Subject: [PATCH 04/25] 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 25eda907d20a7b..fa5021baf667b5 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<unsigned> 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 7171d85b0d0ab0..87e70df795a986 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 <alexandru.voicu at amd.com>
Date: Sat, 2 Nov 2024 01:20:12 +0000
Subject: [PATCH 05/25] 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 098163f957f759..994282debb0d08 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 <alexandru.voicu at amd.com>
Date: Wed, 6 Nov 2024 13:16:03 +0200
Subject: [PATCH 06/25] 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 9d785d878b61dc..4ca5d2b6548124 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 edd7146dc1ac76..7ab9f0aeb60993 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 9b3c2f1b2af677..f5e2b096212f4d 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 3802dc8bcafc49..3b5b860a1b087f 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 ab2e2bd0b30646..e178c0fab5910d 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 56ad0503a11ab2..c45e7020de3f52 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 1904e8fdb3888a..ee0db9778bdcb0 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 2d858fa2f3c3a3..d89a0bdff56a35 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 2578fc0291e760..ffacb0ccbea53f 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 989e46f4b66a7d..c8796036b214f5 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 7f73bf2a65266e..f097c27bd89478 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 <alexandru.voicu at amd.com>
Date: Wed, 6 Nov 2024 13:21:51 +0200
Subject: [PATCH 07/25] 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 f097c27bd89478..6e5b46d5f91c8a 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 <alexandru.voicu at amd.com>
Date: Fri, 8 Nov 2024 01:13:11 +0200
Subject: [PATCH 08/25] 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 fa5021baf667b5..25eda907d20a7b 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<unsigned> 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 87e70df795a986..32200ada7cf7de 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 <alexandru.voicu at amd.com>
Date: Fri, 15 Nov 2024 23:04:53 +0000
Subject: [PATCH 09/25] 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 32200ada7cf7de..41105262b5f14d 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<llvm::Argument>(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 2ad6587089f101..f9c9c5df801631 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<llvm::IntrinsicInst>(std::prev(Builder.GetInsertPoint()));
@@ -316,7 +323,7 @@ void AggExprEmitter::withReturnValueSlot(
"Last insertion wasn't a lifetime.start?");
CGF.pushFullExprCleanup<CodeGenFunction::CallLifetimeEnd>(
- 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 4a1db2c3564a57..effdeb9546800d 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 c2b2e00d15e13f..2f8ba99a3e4160 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 <alexandru.voicu at amd.com>
Date: Fri, 15 Nov 2024 23:12:40 +0000
Subject: [PATCH 10/25] 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 41105262b5f14d..40c41e0895d66f 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 <alexandru.voicu at amd.com>
Date: Sun, 24 Nov 2024 21:54:00 +0000
Subject: [PATCH 11/25] 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 4f2ea9d18fc161..8198669f058d3d 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 <alexandru.voicu at amd.com>
Date: Thu, 5 Dec 2024 03:13:10 +0000
Subject: [PATCH 12/25] 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 4ca5d2b6548124..040ee025afaa86 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 2d6280f8fc508c..19d8c66b000cb0 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 b8a8de57e5b971..57bcb3dd0a852f 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 79300df15d0e29..d55fcc9e125f00 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<EnumType>())
@@ -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<RecordType>())
if (!isa<CXXRecordDecl>(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 0687020015349c..c0e11d2269e18e 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 b6eb02a394f321..bc5bb24a003448 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 ef619e9613d837..01a6402b42c6a0 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 be33e26f047841..c224b29a243a84 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<llvm::Type *> &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<BitIntType>())
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<BitIntType>())
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 c45e7020de3f52..b180b1b8fa00c9 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 ee0db9778bdcb0..a73b668c30ae1f 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 d89a0bdff56a35..c663d02ebb88b7 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<BitIntType>())
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<VectorType>()) {
// 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<BitIntType>())
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 50547dd6dec5e7..26e2a22f14d1e2 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 2849222f7a1869..612c61d83a00ac 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<BitIntType>())
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<EnumType>())
@@ -72,7 +75,8 @@ class BPFABIInfo : public DefaultABIInfo {
ASTContext &Context = getContext();
if (const auto *EIT = RetTy->getAs<BitIntType>())
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 d8720afd1a7132..41dd921ae309ef 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 8fd2a81494d998..151e269a10d380 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 ffacb0ccbea53f..9b40655fb52890 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 6af9375461f095..b89130f93720db 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 06d9b6d4a57615..067ffa2c2cd81c 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 0431d2cc4ddc39..ce59f63fcf6f19 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<BitIntType>()) {
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 9b7d757df3a392..ac0bc598ff84a2 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<EnumType>()) {
// 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<BitIntType>()) {
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 c8796036b214f5..5204bfddf1ab11 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<BitIntType>())
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<BitIntType>())
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 b04e436c665f52..dc55924f747f48 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 a48fe9d5f1ee9c..920d2151b622d7 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<RecordType>()) {
const RecordDecl *RD = RT->getDecl();
diff --git a/clang/lib/CodeGen/Targets/Sparc.cpp b/clang/lib/CodeGen/Targets/Sparc.cpp
index da8c7219be2639..56069a7d818eeb 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<EnumType>())
@@ -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 23c96fa5cf98cb..2dfd590af5b453 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<RecordType>()) {
@@ -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 70a968fe93ca76..dc45062c345ad4 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 6e5b46d5f91c8a..61132a0c9ce30e 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 <alexandru.voicu at amd.com>
Date: Thu, 5 Dec 2024 03:21:02 +0000
Subject: [PATCH 13/25] 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 d55fcc9e125f00..65a186542692d1 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 01a6402b42c6a0..a3a8de8028e710 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 c224b29a243a84..996bc7fa87b07b 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 b180b1b8fa00c9..56b12c18eb4f61 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 a73b668c30ae1f..4e2c869fd4b2fc 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 c663d02ebb88b7..6796b4074c30fa 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<VectorType>()) {
// 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 612c61d83a00ac..ec3402b19b7281 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<EnumType>())
@@ -75,8 +75,8 @@ class BPFABIInfo : public DefaultABIInfo {
ASTContext &Context = getContext();
if (const auto *EIT = RetTy->getAs<BitIntType>())
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 151e269a10d380..6f8a5ad650b1a0 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 9b40655fb52890..b3d8a81a56aacd 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 b89130f93720db..4dc05fd0cb8eb1 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 067ffa2c2cd81c..ae0125c48a0c35 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 ac0bc598ff84a2..fc5cc90a9e995e 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<BitIntType>()) {
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 5204bfddf1ab11..2a5d454c74d6a7 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<BitIntType>())
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 dc55924f747f48..03289d958f1428 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 920d2151b622d7..ef30db8cfb221f 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<RecordType>()) {
diff --git a/clang/lib/CodeGen/Targets/Sparc.cpp b/clang/lib/CodeGen/Targets/Sparc.cpp
index 56069a7d818eeb..38107a75f9c974 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<EnumType>())
diff --git a/clang/lib/CodeGen/Targets/WebAssembly.cpp b/clang/lib/CodeGen/Targets/WebAssembly.cpp
index dc45062c345ad4..502f5260fbabf2 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 61132a0c9ce30e..af6d62fc059991 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 <alexandru.voicu at amd.com>
Date: Sun, 5 Jan 2025 18:02:43 +0200
Subject: [PATCH 14/25] 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 19d8c66b000cb0..b7f355eb335f42 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 57bcb3dd0a852f..538b4ed26b113a 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 b010ea8407b06f..5f23b99adaa0ca 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<llvm::Type *> &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<BitIntType>())
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 4e2c869fd4b2fc..1e5516e1b4a0e1 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 6796b4074c30fa..e4628a7e9731a8 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<BitIntType>())
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 ec3402b19b7281..31ade54f7ef71f 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<BitIntType>())
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 41dd921ae309ef..29b239b24e56bf 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 6f8a5ad650b1a0..38d1c0232d6e10 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 b3d8a81a56aacd..4ea078a9e48af8 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 4dc05fd0cb8eb1..8edff7b59fee9b 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 ae0125c48a0c35..49613c3341b1ac 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 ce59f63fcf6f19..5aab9702c467c2 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<BitIntType>()) {
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 fc5cc90a9e995e..85c2743a16985f 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<EnumType>()) {
@@ -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 2a5d454c74d6a7..2acaf181677afe 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<BitIntType>())
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 60dd48ba5532c7..7f8df7458fcb3b 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 ef30db8cfb221f..1a66be1eff552b 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<RecordType>()) {
const RecordDecl *RD = RT->getDecl();
diff --git a/clang/lib/CodeGen/Targets/Sparc.cpp b/clang/lib/CodeGen/Targets/Sparc.cpp
index 38107a75f9c974..d0b40aa9ceab1c 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<EnumType>())
@@ -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 2dfd590af5b453..4fd141c694c8bd 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<RecordType>()) {
@@ -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 502f5260fbabf2..06012134838647 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 af6d62fc059991..53e1df25522dbf 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 <alexandru.voicu at amd.com>
Date: Sun, 5 Jan 2025 18:13:13 +0200
Subject: [PATCH 15/25] 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 56b12c18eb4f61..10ef72dfa55d99 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 <alexandru.voicu at amd.com>
Date: Tue, 7 Jan 2025 13:50:57 +0200
Subject: [PATCH 16/25] 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 a3a8de8028e710..299a3799777044 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<RecordType>(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 <alexandru.voicu at amd.com>
Date: Tue, 7 Jan 2025 14:17:29 +0200
Subject: [PATCH 17/25] 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 299a3799777044..b91a35ef0e65e0 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<RecordType>(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 19520e0229b64e..8241c8e1bef501 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 1e5516e1b4a0e1..218250484aab5b 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 31ade54f7ef71f..1d575f464e1e45 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 29b239b24e56bf..ce15b068141d02 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 8edff7b59fee9b..691333b2b6f9a6 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 7f8df7458fcb3b..50802a29da1a4a 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 <alexandru.voicu at amd.com>
Date: Wed, 8 Jan 2025 23:14:02 +0200
Subject: [PATCH 18/25] 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 b7f355eb335f42..19d8c66b000cb0 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 538b4ed26b113a..57bcb3dd0a852f 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 65a186542692d1..982563cd7fa9f6 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 8241c8e1bef501..cd61aa57d4c372 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<llvm::Type *> &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<BitIntType>())
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 218250484aab5b..c8db7e8f9706ed 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 e4628a7e9731a8..de11c1fd1fd78d 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<BitIntType>())
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 1d575f464e1e45..880a891083c3a8 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<BitIntType>())
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 ce15b068141d02..ef26d483a180a4 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 38d1c0232d6e10..667599d2d9a669 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 4ea078a9e48af8..6f75bd54a8ef25 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 691333b2b6f9a6..f1d972b3d4ffc6 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 49613c3341b1ac..2ba0b363306190 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 1920a2ff4aecff..c236de8db70e48 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<BitIntType>()) {
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 85c2743a16985f..abe95c01cda703 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 2acaf181677afe..b03eb8c9d50354 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<BitIntType>())
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 50802a29da1a4a..cdf9cf4b8f3dc1 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 2bc7dfc62ae441..56e59306f23820 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<RecordType>()) {
diff --git a/clang/lib/CodeGen/Targets/Sparc.cpp b/clang/lib/CodeGen/Targets/Sparc.cpp
index d0b40aa9ceab1c..9642196b78c63a 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<EnumType>())
@@ -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 4fd141c694c8bd..c928d3c029caa4 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<RecordType>()) {
@@ -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 06012134838647..9217c78a540a38 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 53e1df25522dbf..329055706d3a8b 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 <alexandru.voicu at amd.com>
Date: Wed, 8 Jan 2025 23:36:53 +0200
Subject: [PATCH 19/25] 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 b91a35ef0e65e0..6f10ff96d3a948 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<RecordType>(type)) {
auto record = recordType->getDecl();
auto &layout = CGM.getContext().getASTRecordLayout(record);
>From 5227aefb2208672b2ecbaa5d703b8c705bce9351 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu at amd.com>
Date: Thu, 23 Jan 2025 01:12:19 +0000
Subject: [PATCH 20/25] 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 c0c0f028da2712..0bef50117e8cf3 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 <alexandru.voicu at amd.com>
Date: Thu, 23 Jan 2025 01:16:48 +0000
Subject: [PATCH 21/25] 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 abe95c01cda703..358010785850e7 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<EnumType>()) {
// Treat an enum type as its underlying type.
Ty = EnumTy->getDecl()->getIntegerType();
>From 53d8462f4f6c43e9005847ec9b15444a73fe7e02 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu at amd.com>
Date: Thu, 23 Jan 2025 01:18:29 +0000
Subject: [PATCH 22/25] 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 4e656d921e9a59..dd98053fa9330a 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 b03eb8c9d50354..f14f3568137651 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 56e59306f23820..d475cf56c19e08 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 <alexandru.voicu at amd.com>
Date: Thu, 23 Jan 2025 01:30:52 +0000
Subject: [PATCH 23/25] 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 dd98053fa9330a..10f9f20bca3137 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 f14f3568137651..4df4c9f3c34754 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 d475cf56c19e08..b81ed29a5159b1 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 0bef50117e8cf3..f2a71184c283b0 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 <alexandru.voicu at amd.com>
Date: Thu, 23 Jan 2025 17:51:54 +0000
Subject: [PATCH 24/25] 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 f2a71184c283b0..704810a21209a8 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 <alexandru.voicu at amd.com>
Date: Thu, 23 Jan 2025 18:25:05 +0000
Subject: [PATCH 25/25] 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 00000000000000..997c8a4a5e5cd7
--- /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
More information about the cfe-commits
mailing list