https://github.com/yqtian-se updated https://github.com/llvm/llvm-project/pull/237655
>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 1/2] [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(); } >From 240edbbf50c6d60dc2ca7012f2ec1be5af5310be Mon Sep 17 00:00:00 2001 From: Yongqiang Tian <[email protected]> Date: Sun, 1 Nov 2026 02:11:24 +0000 Subject: [PATCH 2/2] [Clang][NVPTX] Return all non-empty records as integer arrays Returning a record as its LLVM struct type copies it field by field and drops bytes that are padding in that type, not only for unions. Return every non-empty record as an integer array, as most targets do, and update the tests whose expected return types change. Assisted-by: Claude Opus 5.5, co-reviewed by Yongqiang Tian. --- clang/lib/CodeGen/Targets/NVPTX.cpp | 33 ++++--------------- clang/test/CodeGen/nvptx-abi.c | 8 ++--- clang/test/CodeGenCXX/nvptx-union-return.cpp | 4 +-- clang/test/OpenMP/nvptx_NRVO_variable.cpp | 4 +-- .../OpenMP/nvptx_unsupported_type_codegen.cpp | 8 ++--- 5 files changed, 18 insertions(+), 39 deletions(-) diff --git a/clang/lib/CodeGen/Targets/NVPTX.cpp b/clang/lib/CodeGen/Targets/NVPTX.cpp index c4fbf0145213eb..0fded559ed9f45 100644 --- a/clang/lib/CodeGen/Targets/NVPTX.cpp +++ b/clang/lib/CodeGen/Targets/NVPTX.cpp @@ -156,28 +156,6 @@ 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 { @@ -199,11 +177,12 @@ 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)) { + // Return a record as an integer array of the same size. Returning its LLVM + // struct type would copy it field by field and drop bytes that are padding + // in that type, which can hold data of an active union member. Elements are + // no wider than the struct type's ABI alignment, so the PTX return parameter + // keeps its alignment. + if (RetTy->isRecordType() && !isEmptyRecord(getContext(), RetTy, true)) { unsigned Align = getDataLayout().getABITypeAlign(CGT.ConvertType(RetTy)).value() * 8; return coerceToIntArrayWithLimit(RetTy, Align); diff --git a/clang/test/CodeGen/nvptx-abi.c b/clang/test/CodeGen/nvptx-abi.c index a6aa7ac2a5a4e7..c467b87e9b5f5d 100644 --- a/clang/test/CodeGen/nvptx-abi.c +++ b/clang/test/CodeGen/nvptx-abi.c @@ -6,7 +6,7 @@ typedef struct float4_s { } float4_t; float4_t my_function(void) { -// CHECK-LABEL: define{{.*}} %struct.float4_s @my_function +// CHECK-LABEL: define{{.*}} [4 x i32] @my_function float4_t t; return t; }; @@ -14,7 +14,7 @@ float4_t my_function(void) { float bar(void) { float4_t ret; // CHECK-LABEL: @bar -// CHECK: call %struct.float4_s @my_function +// CHECK: call [4 x i32] @my_function ret = my_function(); return ret.x; } @@ -43,8 +43,8 @@ void baz(nested_t 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. +// holds data of the other member. Records are returned as an integer array so +// that those bytes are not dropped. typedef struct pad_s { short s; unsigned long long l; diff --git a/clang/test/CodeGenCXX/nvptx-union-return.cpp b/clang/test/CodeGenCXX/nvptx-union-return.cpp index 81e885710eb6bd..d15c96f9325080 100644 --- a/clang/test/CodeGenCXX/nvptx-union-return.cpp +++ b/clang/test/CodeGenCXX/nvptx-union-return.cpp @@ -1,7 +1,7 @@ // 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. +// A union in a base class is returned as an integer array, so that bytes of +// the active member are not dropped. union U { unsigned long long l; struct { diff --git a/clang/test/OpenMP/nvptx_NRVO_variable.cpp b/clang/test/OpenMP/nvptx_NRVO_variable.cpp index 010b43c2f5dab3..281d4cc5cf066c 100644 --- a/clang/test/OpenMP/nvptx_NRVO_variable.cpp +++ b/clang/test/OpenMP/nvptx_NRVO_variable.cpp @@ -19,8 +19,8 @@ S foo() { // CHECK: call void @{{.+}}bar{{.+}}(ptr {{.*}}[[S_REF:%.+]]) bar(s); // CHECK: call void @llvm.memcpy.p0.p0.i64(ptr {{.*}}[[RETVAL]], ptr {{.*}}[[S_REF]], i64 4, i1 false) - // CHECK: [[VAL:%.+]] = load %struct.S, ptr [[RETVAL]], - // CHECK: ret %struct.S [[VAL]] + // CHECK: [[VAL:%.+]] = load [1 x i32], ptr + // CHECK: ret [1 x i32] [[VAL]] return s; } #pragma omp end declare target diff --git a/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp b/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp index d1d087ee887f32..ff36760f28310f 100644 --- a/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp +++ b/clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp @@ -57,15 +57,15 @@ T1 f1 = a1; void foo1(T1 a = T1()) { return; } -// CHECK: define{{ hidden | }}[[T1]] @{{.+}}bar1{{.+}}() +// CHECK: define{{ hidden | }}[64 x i8] @{{.+}}bar1{{.+}}() T1 bar1() { -// CHECK: load [[T1]], ptr -// CHECK-NEXT: ret [[T1]] +// CHECK: load [64 x i8], ptr +// CHECK-NEXT: ret [64 x i8] return T1(); } // CHECK: define{{ hidden | }}void @{{.+}}baz1{{.+}}() void baz1() { -// CHECK: call [[T1]] @{{.+}}bar1{{.+}}() +// CHECK: call [64 x i8] @{{.+}}bar1{{.+}}() T1 t = bar1(); } #pragma omp end declare target _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
