https://github.com/addmisol created
https://github.com/llvm/llvm-project/pull/185083
Fixes #184150
This PR removes the ABI lowering code that coerces 33-64 bit aggregate types
to [2 x i32] arrays for AMDGPU targets. Instead, structs are now passed
through using ABIArgInfo::getDirect() without a coercion type, which
preserves the original struct type and its field types.
Previously, a struct like { float, int } would be lowered to [2 x i32],
losing the floating-point type information. This prevents attaching
FP-specific attributes like nofpclass to the float component.
## Changes
- *clang/lib/CodeGen/Targets/AMDGPU.cpp*: Remove the special-case coercion
for 33-64 bit aggregates in both classifyReturnType and
classifyArgumentType. The existing numRegsForType path handles these
cases correctly while preserving type information.
- *clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl*: Update expected
output to reflect that structs now preserve their original types instead
of being coerced to [2 x i32] or integer types.
- *clang/test/CodeGen/amdgpu-abi-struct-coerce.c* (new): Add test coverage
for the specific case mentioned in the issue - structs containing mixed
float and int fields.
## Before/After
```c
typedef struct { float f; int i; } fp_int_pair;
fp_int_pair foo(fp_int_pair x) { return x; }
Before:
define [2 x i32] @foo([2 x i32] %x.coerce)
After:
define %struct.fp_int_pair @foo(float %x.coerce0, i32 %x.coerce1)
Test Plan
- Updated existing ABI tests in amdgpu-abi-struct-coerce.cl
- Added new test amdgpu-abi-struct-coerce.c for mixed FP/int structs
- Updated affected OpenMP complex math header tests
>From c5ffb2e73bcf69513f94d8e7b89e8372d0d280b2 Mon Sep 17 00:00:00 2001
From: addmisol <[email protected]>
Date: Fri, 6 Mar 2026 23:56:34 +0530
Subject: [PATCH 1/7] Create amdgpu-abi-struct-coerce.c
---
.../test/CodeGen/amdgpu-abi-struct-coerce.c | 71 +++++++++++++++++++
1 file changed, 71 insertions(+)
create mode 100644
clang/test/CodeGen/clang/test/CodeGen/amdgpu-abi-struct-coerce.c
diff --git a/clang/test/CodeGen/clang/test/CodeGen/amdgpu-abi-struct-coerce.c
b/clang/test/CodeGen/clang/test/CodeGen/amdgpu-abi-struct-coerce.c
new file mode 100644
index 0000000000000..2399630ff797b
--- /dev/null
+++ b/clang/test/CodeGen/clang/test/CodeGen/amdgpu-abi-struct-coerce.c
@@ -0,0 +1,71 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm -o - %s | FileCheck %s
+
+// Check that structs containing mixed float and int types are not coerced
+// to integer arrays. They should preserve the original struct type and
+// individual field types.
+
+typedef struct fp_int_pair {
+ float f;
+ int i;
+} fp_int_pair;
+
+// CHECK-LABEL: define{{.*}} %struct.fp_int_pair @return_fp_int_pair(float
%x.coerce0, i32 %x.coerce1)
+// CHECK: ret %struct.fp_int_pair
+fp_int_pair return_fp_int_pair(fp_int_pair x) {
+ return x;
+}
+
+typedef struct int_fp_pair {
+ int i;
+ float f;
+} int_fp_pair;
+
+// CHECK-LABEL: define{{.*}} %struct.int_fp_pair @return_int_fp_pair(i32
%x.coerce0, float %x.coerce1)
+// CHECK: ret %struct.int_fp_pair
+int_fp_pair return_int_fp_pair(int_fp_pair x) {
+ return x;
+}
+
+typedef struct two_floats {
+ float a;
+ float b;
+} two_floats;
+
+// CHECK-LABEL: define{{.*}} %struct.two_floats @return_two_floats(float
%x.coerce0, float %x.coerce1)
+// CHECK: ret %struct.two_floats
+two_floats return_two_floats(two_floats x) {
+ return x;
+}
+
+typedef struct two_ints {
+ int a;
+ int b;
+} two_ints;
+
+// CHECK-LABEL: define{{.*}} %struct.two_ints @return_two_ints(i32 %x.coerce0,
i32 %x.coerce1)
+// CHECK: ret %struct.two_ints
+two_ints return_two_ints(two_ints x) {
+ return x;
+}
+
+// Structs <= 32 bits should still be coerced to i32 for return value
+typedef struct small_struct {
+ short a;
+ short b;
+} small_struct;
+
+// CHECK-LABEL: define{{.*}} i32 @return_small_struct(i16 %x.coerce0, i16
%x.coerce1)
+small_struct return_small_struct(small_struct x) {
+ return x;
+}
+
+// Structs <= 16 bits should still be coerced to i16 for return value
+typedef struct tiny_struct {
+ char a;
+ char b;
+} tiny_struct;
+
+// CHECK-LABEL: define{{.*}} i16 @return_tiny_struct(i8 %x.coerce0, i8
%x.coerce1)
+tiny_struct return_tiny_struct(tiny_struct x) {
+ return x;
+}
>From 68c200f848058ab22b3d25ce810f1639eac50556 Mon Sep 17 00:00:00 2001
From: addmisol <[email protected]>
Date: Fri, 6 Mar 2026 23:57:11 +0530
Subject: [PATCH 2/7] Delete
clang/test/CodeGen/clang/test/CodeGen/amdgpu-abi-struct-coerce.c
---
.../test/CodeGen/amdgpu-abi-struct-coerce.c | 71 -------------------
1 file changed, 71 deletions(-)
delete mode 100644
clang/test/CodeGen/clang/test/CodeGen/amdgpu-abi-struct-coerce.c
diff --git a/clang/test/CodeGen/clang/test/CodeGen/amdgpu-abi-struct-coerce.c
b/clang/test/CodeGen/clang/test/CodeGen/amdgpu-abi-struct-coerce.c
deleted file mode 100644
index 2399630ff797b..0000000000000
--- a/clang/test/CodeGen/clang/test/CodeGen/amdgpu-abi-struct-coerce.c
+++ /dev/null
@@ -1,71 +0,0 @@
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm -o - %s | FileCheck %s
-
-// Check that structs containing mixed float and int types are not coerced
-// to integer arrays. They should preserve the original struct type and
-// individual field types.
-
-typedef struct fp_int_pair {
- float f;
- int i;
-} fp_int_pair;
-
-// CHECK-LABEL: define{{.*}} %struct.fp_int_pair @return_fp_int_pair(float
%x.coerce0, i32 %x.coerce1)
-// CHECK: ret %struct.fp_int_pair
-fp_int_pair return_fp_int_pair(fp_int_pair x) {
- return x;
-}
-
-typedef struct int_fp_pair {
- int i;
- float f;
-} int_fp_pair;
-
-// CHECK-LABEL: define{{.*}} %struct.int_fp_pair @return_int_fp_pair(i32
%x.coerce0, float %x.coerce1)
-// CHECK: ret %struct.int_fp_pair
-int_fp_pair return_int_fp_pair(int_fp_pair x) {
- return x;
-}
-
-typedef struct two_floats {
- float a;
- float b;
-} two_floats;
-
-// CHECK-LABEL: define{{.*}} %struct.two_floats @return_two_floats(float
%x.coerce0, float %x.coerce1)
-// CHECK: ret %struct.two_floats
-two_floats return_two_floats(two_floats x) {
- return x;
-}
-
-typedef struct two_ints {
- int a;
- int b;
-} two_ints;
-
-// CHECK-LABEL: define{{.*}} %struct.two_ints @return_two_ints(i32 %x.coerce0,
i32 %x.coerce1)
-// CHECK: ret %struct.two_ints
-two_ints return_two_ints(two_ints x) {
- return x;
-}
-
-// Structs <= 32 bits should still be coerced to i32 for return value
-typedef struct small_struct {
- short a;
- short b;
-} small_struct;
-
-// CHECK-LABEL: define{{.*}} i32 @return_small_struct(i16 %x.coerce0, i16
%x.coerce1)
-small_struct return_small_struct(small_struct x) {
- return x;
-}
-
-// Structs <= 16 bits should still be coerced to i16 for return value
-typedef struct tiny_struct {
- char a;
- char b;
-} tiny_struct;
-
-// CHECK-LABEL: define{{.*}} i16 @return_tiny_struct(i8 %x.coerce0, i8
%x.coerce1)
-tiny_struct return_tiny_struct(tiny_struct x) {
- return x;
-}
>From 3c5401a8e20cdac719d6817e198cc330dc0e4e80 Mon Sep 17 00:00:00 2001
From: addmisol <[email protected]>
Date: Fri, 6 Mar 2026 23:58:43 +0530
Subject: [PATCH 3/7] fix for clang abi lowering
---
clang/test/CodeGen/amdgpu-abi-struct-coerce.c | 71 +++++++++++++++++++
1 file changed, 71 insertions(+)
create mode 100644 clang/test/CodeGen/amdgpu-abi-struct-coerce.c
diff --git a/clang/test/CodeGen/amdgpu-abi-struct-coerce.c
b/clang/test/CodeGen/amdgpu-abi-struct-coerce.c
new file mode 100644
index 0000000000000..2399630ff797b
--- /dev/null
+++ b/clang/test/CodeGen/amdgpu-abi-struct-coerce.c
@@ -0,0 +1,71 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm -o - %s | FileCheck %s
+
+// Check that structs containing mixed float and int types are not coerced
+// to integer arrays. They should preserve the original struct type and
+// individual field types.
+
+typedef struct fp_int_pair {
+ float f;
+ int i;
+} fp_int_pair;
+
+// CHECK-LABEL: define{{.*}} %struct.fp_int_pair @return_fp_int_pair(float
%x.coerce0, i32 %x.coerce1)
+// CHECK: ret %struct.fp_int_pair
+fp_int_pair return_fp_int_pair(fp_int_pair x) {
+ return x;
+}
+
+typedef struct int_fp_pair {
+ int i;
+ float f;
+} int_fp_pair;
+
+// CHECK-LABEL: define{{.*}} %struct.int_fp_pair @return_int_fp_pair(i32
%x.coerce0, float %x.coerce1)
+// CHECK: ret %struct.int_fp_pair
+int_fp_pair return_int_fp_pair(int_fp_pair x) {
+ return x;
+}
+
+typedef struct two_floats {
+ float a;
+ float b;
+} two_floats;
+
+// CHECK-LABEL: define{{.*}} %struct.two_floats @return_two_floats(float
%x.coerce0, float %x.coerce1)
+// CHECK: ret %struct.two_floats
+two_floats return_two_floats(two_floats x) {
+ return x;
+}
+
+typedef struct two_ints {
+ int a;
+ int b;
+} two_ints;
+
+// CHECK-LABEL: define{{.*}} %struct.two_ints @return_two_ints(i32 %x.coerce0,
i32 %x.coerce1)
+// CHECK: ret %struct.two_ints
+two_ints return_two_ints(two_ints x) {
+ return x;
+}
+
+// Structs <= 32 bits should still be coerced to i32 for return value
+typedef struct small_struct {
+ short a;
+ short b;
+} small_struct;
+
+// CHECK-LABEL: define{{.*}} i32 @return_small_struct(i16 %x.coerce0, i16
%x.coerce1)
+small_struct return_small_struct(small_struct x) {
+ return x;
+}
+
+// Structs <= 16 bits should still be coerced to i16 for return value
+typedef struct tiny_struct {
+ char a;
+ char b;
+} tiny_struct;
+
+// CHECK-LABEL: define{{.*}} i16 @return_tiny_struct(i8 %x.coerce0, i8
%x.coerce1)
+tiny_struct return_tiny_struct(tiny_struct x) {
+ return x;
+}
>From 6cd1099ec2e06c33fd5d7092206e778a1e8ba58a Mon Sep 17 00:00:00 2001
From: addmisol <[email protected]>
Date: Sat, 7 Mar 2026 00:00:29 +0530
Subject: [PATCH 4/7] Update amdgcn-openmp-device-math-complex.c
---
clang/test/Headers/amdgcn-openmp-device-math-complex.c | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/clang/test/Headers/amdgcn-openmp-device-math-complex.c
b/clang/test/Headers/amdgcn-openmp-device-math-complex.c
index b347cf4716df2..34c05e2974a64 100644
--- a/clang/test/Headers/amdgcn-openmp-device-math-complex.c
+++ b/clang/test/Headers/amdgcn-openmp-device-math-complex.c
@@ -30,8 +30,8 @@ void test_complex_f32(float _Complex a) {
// CHECK-LABEL: define {{.*}}test_complex_f32
#pragma omp target
{
- // CHECK: call [2 x i32] @__divsc3
- // CHECK: call [2 x i32] @__mulsc3
+ // CHECK: call { float, float } @__divsc3
+ // CHECK: call { float, float } @__mulsc3
(void)(a * (a / a));
}
}
>From a67bcdb1baecf786c7714a07d05306b614634ce5 Mon Sep 17 00:00:00 2001
From: addmisol <[email protected]>
Date: Sat, 7 Mar 2026 00:11:07 +0530
Subject: [PATCH 5/7] Update amdgpu-abi-struct-coerce.cl
---
.../CodeGenOpenCL/amdgpu-abi-struct-coerce.cl | 16 +++++++++-------
1 file changed, 9 insertions(+), 7 deletions(-)
diff --git a/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl
b/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl
index 06d3cdb01deb2..a13f8e8bbe119 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl
@@ -288,16 +288,16 @@ void func_struct_arg(struct_arg_t arg1) { }
// CHECK: void @func_struct_padding_arg(i8 %arg1.coerce0, i64 %arg1.coerce1)
void func_struct_padding_arg(struct_padding_arg arg1) { }
-// CHECK: define{{.*}} void @func_struct_char_x8([2 x i32] %arg.coerce)
+// CHECK: define{{.*}} void @func_struct_char_x8(i8 %arg.coerce0, i8
%arg.coerce1, i8 %arg.coerce2, i8 %arg.coerce3, i8 %arg.coerce4, i8
%arg.coerce5, i8 %arg.coerce6, i8 %arg.coerce7)
void func_struct_char_x8(struct_char_x8 arg) { }
-// CHECK: define{{.*}} void @func_struct_char_x4(i32 %arg.coerce)
+// CHECK: define{{.*}} void @func_struct_char_x4(i8 %arg.coerce0, i8
%arg.coerce1, i8 %arg.coerce2, i8 %arg.coerce3)
void func_struct_char_x4(struct_char_x4 arg) { }
-// CHECK: define{{.*}} void @func_struct_char_x3(i32 %arg.coerce)
+// CHECK: define{{.*}} void @func_struct_char_x3(i8 %arg.coerce0, i8
%arg.coerce1, i8 %arg.coerce2)
void func_struct_char_x3(struct_char_x3 arg) { }
-// CHECK: define{{.*}} void @func_struct_char_x2(i16 %arg.coerce)
+// CHECK: define{{.*}} void @func_struct_char_x2(i8 %arg.coerce0, i8
%arg.coerce1)
void func_struct_char_x2(struct_char_x2 arg) { }
// CHECK: define{{.*}} void @func_struct_char_x1(i8 %arg.coerce)
@@ -363,8 +363,10 @@ struct_padding_arg func_struct_padding_ret()
return s;
}
-// CHECK: define{{.*}} [2 x i32] @func_struct_char_x8_ret()
-// CHECK: ret [2 x i32] zeroinitializer
+// CHECK: define{{.*}} %struct.struct_char_x8 @func_struct_char_x8_ret()
+// CHECK: ret %struct.struct_char_x8 zeroinitializer
+ struct_char_x8 func_struct_char_x8_ret()
+ {
struct_char_x8 func_struct_char_x8_ret()
{
struct_char_x8 s = { 0 };
@@ -525,5 +527,5 @@ void v2i8_reg_count(char2 arg0, char2 arg1, char2 arg2,
char2 arg3,
void v2i8_reg_count_over(char2 arg0, char2 arg1, char2 arg2, char2 arg3,
char2 arg4, char2 arg5, int arg6, struct_4regs arg7)
{ }
-// CHECK: define{{.*}} void @num_regs_left_64bit_aggregate(<4 x i32> noundef
%arg0, <4 x i32> noundef %arg1, <4 x i32> noundef %arg2, <3 x i32> noundef
%arg3, [2 x i32] %arg4.coerce, i32 noundef %arg5)
+// CHECK: define{{.}} void @num_regs_left_64bit_aggregate(<4 x i32> noundef
%arg0, <4 x i32> noundef %arg1, <4 x i32> noundef %arg2, <3 x i32> noundef
%arg3, ptr addrspace(5) noundef readnone byref(%struct.struct_char_x8) align 1
captures(none) %{{.}}, i32 noundef %arg5)
void num_regs_left_64bit_aggregate(int4 arg0, int4 arg1, int4 arg2, int3 arg3,
struct_char_x8 arg4, int arg5) { }
>From c299160a68b48335ff616aa586098403a9bb81b3 Mon Sep 17 00:00:00 2001
From: addmisol <[email protected]>
Date: Sat, 7 Mar 2026 00:13:09 +0530
Subject: [PATCH 6/7] Update AMDGPU.cpp
---
clang/lib/CodeGen/Targets/AMDGPU.cpp | 22 ----------------------
1 file changed, 22 deletions(-)
diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp
b/clang/lib/CodeGen/Targets/AMDGPU.cpp
index 4ac7f42289d6d..f3c4b5ad0837b 100644
--- a/clang/lib/CodeGen/Targets/AMDGPU.cpp
+++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp
@@ -163,11 +163,6 @@ ABIArgInfo AMDGPUABIInfo::classifyReturnType(QualType
RetTy) const {
if (Size <= 32)
return ABIArgInfo::getDirect(llvm::Type::getInt32Ty(getVMContext()));
- if (Size <= 64) {
- llvm::Type *I32Ty = llvm::Type::getInt32Ty(getVMContext());
- return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2));
- }
-
if (numRegsForType(RetTy) <= MaxNumRegsForArgsRet)
return ABIArgInfo::getDirect();
}
@@ -246,23 +241,6 @@ ABIArgInfo AMDGPUABIInfo::classifyArgumentType(QualType
Ty, bool Variadic,
RD && RD->hasFlexibleArrayMember())
return DefaultABIInfo::classifyArgumentType(Ty);
- // Pack aggregates <= 8 bytes into single VGPR or pair.
- uint64_t Size = getContext().getTypeSize(Ty);
- if (Size <= 64) {
- unsigned NumRegs = (Size + 31) / 32;
- NumRegsLeft -= std::min(NumRegsLeft, NumRegs);
-
- if (Size <= 16)
- return ABIArgInfo::getDirect(llvm::Type::getInt16Ty(getVMContext()));
-
- if (Size <= 32)
- return ABIArgInfo::getDirect(llvm::Type::getInt32Ty(getVMContext()));
-
- // XXX: Should this be i64 instead, and should the limit increase?
- llvm::Type *I32Ty = llvm::Type::getInt32Ty(getVMContext());
- return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2));
- }
-
if (NumRegsLeft > 0) {
uint64_t NumRegs = numRegsForType(Ty);
if (NumRegsLeft >= NumRegs) {
>From 3c87855bcfb0874e8abad1f3735350bb56e369c7 Mon Sep 17 00:00:00 2001
From: addmisol <[email protected]>
Date: Sat, 7 Mar 2026 00:31:52 +0530
Subject: [PATCH 7/7] Update amdgpu-abi-struct-coerce.cl
---
clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl | 2 --
1 file changed, 2 deletions(-)
diff --git a/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl
b/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl
index a13f8e8bbe119..fb5ba69c86c6d 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl
@@ -365,8 +365,6 @@ struct_padding_arg func_struct_padding_ret()
// CHECK: define{{.*}} %struct.struct_char_x8 @func_struct_char_x8_ret()
// CHECK: ret %struct.struct_char_x8 zeroinitializer
- struct_char_x8 func_struct_char_x8_ret()
- {
struct_char_x8 func_struct_char_x8_ret()
{
struct_char_x8 s = { 0 };
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
