https://github.com/yqtian-se created
https://github.com/llvm/llvm-project/pull/237655
A union is lowered to the LLVM type of one member, and NVPTX returned it
directly, so bytes that are padding in that member were dropped. For
`union U { unsigned long a; struct { unsigned short x; unsigned long y; } s; }`,
returning a `U` with `a` active gave the caller only the low 16 bits of `a`.
This patch returns types that are or contain a union as an integer array,
with elements no wider than the alignment of the previous return type, so the
PTX return parameter is unchanged.
This issue was found by a non-AI automated testing tool. AI assisted with the
analysis and development of the fix.
Fixes #237641
>From 6fa229603efdff27242ccf327b8b13c005750ae1 Mon Sep 17 00:00:00 2001
From: Yongqiang Tian <[email protected]>
Date: Sat, 31 Oct 2026 10:21:41 +0000
Subject: [PATCH] [Clang][NVPTX] Return all bytes of unions from device
functions
A union is lowered to the LLVM type of one member, and NVPTX returned it
directly, so bytes that are padding in that member were dropped. For
`union U { unsigned long a; struct { unsigned short x; unsigned long y; } s; }`,
returning a `U` with `a` active gave the caller only the low 16 bits of `a`.
Return types that are or contain a union are now returned as an integer
array. Its elements are no wider than the alignment of the previous return
type, so the PTX return parameter is unchanged.
Fixes #237641
Assisted-by: Claude Opus 5.5, co-reviewed by Yongqiang Tian.
---
clang/lib/CodeGen/Targets/NVPTX.cpp | 32 +++++++
clang/test/CodeGen/nvptx-abi.c | 94 ++++++++++++++++++-
.../CodeGenCUDA/nvptx-union-return-padding.cu | 20 ++++
clang/test/CodeGenCXX/nvptx-union-return.cpp | 23 +++++
4 files changed, 168 insertions(+), 1 deletion(-)
create mode 100644 clang/test/CodeGenCUDA/nvptx-union-return-padding.cu
create mode 100644 clang/test/CodeGenCXX/nvptx-union-return.cpp
diff --git a/clang/lib/CodeGen/Targets/NVPTX.cpp
b/clang/lib/CodeGen/Targets/NVPTX.cpp
index ed70b01fcf48bb..c4fbf0145213eb 100644
--- a/clang/lib/CodeGen/Targets/NVPTX.cpp
+++ b/clang/lib/CodeGen/Targets/NVPTX.cpp
@@ -156,6 +156,28 @@ bool NVPTXABIInfo::isUnsupportedType(QualType T) const {
return false;
}
+/// Whether \p Ty is or contains a union. Unions are lowered to the LLVM type
of
+/// one member, so bytes that are padding in that member, but hold data of
+/// another, are not part of the LLVM type and are not copied by an aggregate
+/// load or store.
+static bool isOrContainsUnion(const ASTContext &Ctx, QualType Ty) {
+ if (const ConstantArrayType *AT = Ctx.getAsConstantArrayType(Ty))
+ return isOrContainsUnion(Ctx, AT->getElementType());
+ const RecordDecl *RD = Ty->getAsRecordDecl();
+ if (!RD)
+ return false;
+ if (RD->isUnion())
+ return true;
+ if (const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD))
+ for (const CXXBaseSpecifier &Base : CXXRD->bases())
+ if (isOrContainsUnion(Ctx, Base.getType()))
+ return true;
+ for (const FieldDecl *FD : RD->fields())
+ if (isOrContainsUnion(Ctx, FD->getType()))
+ return true;
+ return false;
+}
+
/// Coerce the given type into an array with maximum allowed size of elements.
ABIArgInfo NVPTXABIInfo::coerceToIntArrayWithLimit(QualType Ty,
unsigned MaxSize) const {
@@ -177,6 +199,16 @@ ABIArgInfo NVPTXABIInfo::classifyReturnType(QualType
RetTy) const {
isUnsupportedType(RetTy))
return coerceToIntArrayWithLimit(RetTy, 64);
+ // Return an aggregate that contains a union as an integer array of the same
+ // size. Use the original LLVM type's ABI alignment for the array elements
+ // so the PTX return parameter keeps its alignment while all bytes of the
+ // active union member are returned.
+ if (isOrContainsUnion(getContext(), RetTy)) {
+ unsigned Align =
+ getDataLayout().getABITypeAlign(CGT.ConvertType(RetTy)).value() * 8;
+ return coerceToIntArrayWithLimit(RetTy, Align);
+ }
+
// note: this is different from default ABI
if (!RetTy->isScalarType())
return ABIArgInfo::getDirect();
diff --git a/clang/test/CodeGen/nvptx-abi.c b/clang/test/CodeGen/nvptx-abi.c
index f9c1de5713f220..a6aa7ac2a5a4e7 100644
--- a/clang/test/CodeGen/nvptx-abi.c
+++ b/clang/test/CodeGen/nvptx-abi.c
@@ -1,5 +1,5 @@
// RUN: %clang_cc1 -triple nvptx-unknown-unknown -o - %s -emit-llvm |
FileCheck %s
-// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -o - %s -emit-llvm |
FileCheck %s
+// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -o - %s -emit-llvm |
FileCheck %s --check-prefixes=CHECK,NVPTX64
typedef struct float4_s {
float x, y, z, w;
@@ -41,3 +41,95 @@ void baz(nested_t x) {
// CHECK-LABEL: @baz
// CHECK: ptr noundef byval(%struct.nested_s) align 8 %x)
}
+
+// A union is lowered to the LLVM type of one member, here pad_t, whose padding
+// holds data of the other member. Aggregates containing a union are returned
+// as an integer array so that those bytes are not dropped.
+typedef struct pad_s {
+ short s;
+ unsigned long long l;
+} pad_t;
+
+typedef union pad_u {
+ unsigned long long l;
+ pad_t p;
+} pad_u_t;
+
+pad_u_t ret_union(void) {
+// CHECK-LABEL: define{{.*}} [2 x i64] @ret_union
+// CHECK: load [2 x i64], ptr
+// CHECK: ret [2 x i64]
+ pad_u_t u = {0x1122334455667788ULL};
+ return u;
+}
+
+unsigned long long call_ret_union(void) {
+// CHECK-LABEL: @call_ret_union
+// CHECK: call [2 x i64] @ret_union
+// CHECK: store [2 x i64]
+ return ret_union().l;
+}
+
+typedef struct with_union_s {
+ int i;
+ pad_u_t u;
+} with_union_t;
+
+with_union_t ret_with_union(void) {
+// CHECK-LABEL: define{{.*}} [3 x i64] @ret_with_union
+ with_union_t t;
+ return t;
+}
+
+typedef struct with_union_array_s {
+ pad_u_t u[2];
+} with_union_array_t;
+
+with_union_array_t ret_with_union_array(void) {
+// CHECK-LABEL: define{{.*}} [4 x i64] @ret_with_union_array
+ with_union_array_t t;
+ return t;
+}
+
+// Keep the ABI alignment of the original LLVM return type, which can differ
+// from the source type's explicit alignment.
+typedef struct __attribute__((aligned(16))) {
+ int x;
+} aligned16_t;
+typedef union {
+ aligned16_t s;
+ unsigned char bytes[16];
+} union_align16_t;
+
+union_align16_t ret_align16(void) {
+// CHECK-LABEL: define{{.*}} [4 x i32] @ret_align16
+ union_align16_t u = {0};
+ return u;
+}
+
+typedef struct __attribute__((aligned(8))) {
+ short x;
+} aligned8_t;
+typedef union {
+ aligned8_t s;
+ unsigned char bytes[8];
+} union_align8_t;
+
+union_align8_t ret_align8(void) {
+// CHECK-LABEL: define{{.*}} [4 x i16] @ret_align8
+ union_align8_t u = {0};
+ return u;
+}
+
+#if __SIZEOF_POINTER__ == 8
+typedef union {
+ __int128 x;
+ unsigned char bytes[16];
+} union_i128_t;
+
+union_i128_t ret_i128(void) {
+// NVPTX64-LABEL: define{{.*}} [1 x i128] @ret_i128
+ union_i128_t u = {0};
+ return u;
+}
+#endif
diff --git a/clang/test/CodeGenCUDA/nvptx-union-return-padding.cu
b/clang/test/CodeGenCUDA/nvptx-union-return-padding.cu
new file mode 100644
index 00000000000000..552b80855f03b2
--- /dev/null
+++ b/clang/test/CodeGenCUDA/nvptx-union-return-padding.cu
@@ -0,0 +1,20 @@
+// REQUIRES: nvptx-registered-target
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -O0 -S -o - %s
| FileCheck %s
+
+#define __device__ __attribute__((device))
+
+union U {
+ unsigned long a;
+ struct {
+ unsigned short x;
+ unsigned long y;
+ } s;
+};
+
+// CHECK-LABEL: .visible .func (.param .align 8 .b8 func_retval0[16]) _Z1fv()
+// CHECK: st.param{{.*}}.b64{{.*}}[func_retval0],
+// CHECK: st.param{{.*}}.b64{{.*}}[func_retval0+8],
+__device__ __attribute__((noinline)) U f() {
+ U u = {0x1122334455667788UL};
+ return u;
+}
diff --git a/clang/test/CodeGenCXX/nvptx-union-return.cpp
b/clang/test/CodeGenCXX/nvptx-union-return.cpp
new file mode 100644
index 00000000000000..81e885710eb6bd
--- /dev/null
+++ b/clang/test/CodeGenCXX/nvptx-union-return.cpp
@@ -0,0 +1,23 @@
+// RUN: %clang_cc1 -triple nvptx64 -emit-llvm -o - %s | FileCheck %s
+
+// A union in a base class is returned as an integer array, like a union
+// member, so that bytes of the active member are not dropped.
+union U {
+ unsigned long long l;
+ struct {
+ short s;
+ unsigned long long l;
+ } p;
+};
+
+struct Base {
+ U u;
+};
+
+struct Derived : Base {
+ int i;
+};
+
+// The original packed LLVM return type has byte alignment.
+// CHECK-LABEL: define{{.*}} [24 x i8] @_Z11ret_derivedv
+Derived ret_derived() { return Derived(); }
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
