https://github.com/michaelselehov created https://github.com/llvm/llvm-project/pull/199981
…ger arrays (#185083)" This reverts commit 50b859cca1ccf7d174ee61a8a130ae14220209e4. >From 411bb66fe7477a22e426a7c1e414244fe2b3adf7 Mon Sep 17 00:00:00 2001 From: mselehov <[email protected]> Date: Wed, 27 May 2026 07:19:43 -0500 Subject: [PATCH] Revert "[AMDGPU] Stop coercing structs with FP and int fields to integer arrays (#185083)" This reverts commit 50b859cca1ccf7d174ee61a8a130ae14220209e4. --- clang/lib/CodeGen/Targets/AMDGPU.cpp | 111 +-- clang/test/CodeGen/amdgpu-abi-struct-coerce.c | 702 ------------------ clang/test/CodeGen/amdgpu-variadic-call.c | 10 +- .../CodeGenOpenCL/amdgpu-abi-struct-coerce.cl | 4 +- .../amdgcn-openmp-device-math-complex.c | 4 +- 5 files changed, 29 insertions(+), 802 deletions(-) delete mode 100644 clang/test/CodeGen/amdgpu-abi-struct-coerce.c diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 0d36f166328c7..a3a596bb9d822 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -78,62 +78,6 @@ bool AMDGPUABIInfo::isHomogeneousAggregateSmallEnough( return Members * NumRegs <= MaxNumRegsForArgsRet; } -/// Check if all fields in an aggregate type contain only sub-32-bit integer -/// types. Such aggregates should be packed into i32 registers rather than -/// passed as individual elements. Aggregates containing floats or full-sized -/// integer types (i32, i64) should preserve their original types. -static bool containsOnlyPackableIntegerTypes(const RecordDecl *RD, - const ASTContext &Context) { - for (const FieldDecl *Field : RD->fields()) { - QualType FieldTy = Field->getType(); - - // For bitfields, they are always integer types so they're always packable. - // A bitfield like "unsigned a : 4" should be packable even though - // 'unsigned' is 32 bits. Similarly, larger bitfields that fill into - // wider ints (like i64) should also be packed. - if (Field->isBitField()) { - continue; - } - - // Recursively check nested structs - if (const RecordDecl *NestedRD = FieldTy->getAsRecordDecl()) { - if (!containsOnlyPackableIntegerTypes(NestedRD, Context)) - return false; - continue; - } - - // Arrays - check the element type - if (const ConstantArrayType *AT = Context.getAsConstantArrayType(FieldTy)) { - QualType EltTy = AT->getElementType(); - if (const RecordDecl *NestedRD = EltTy->getAsRecordDecl()) { - if (!containsOnlyPackableIntegerTypes(NestedRD, Context)) - return false; - continue; - } - // For non-struct array elements, check if they're packable integers - if (!EltTy->isIntegerType()) - return false; - uint64_t EltSize = Context.getTypeSize(EltTy); - if (EltSize >= 32) - return false; - continue; - } - - // Floating point types should not be packed into integers - if (FieldTy->isFloatingType()) - return false; - - // Only integer types that are smaller than 32 bits should be packed - if (!FieldTy->isIntegerType()) - return false; - - uint64_t FieldSize = Context.getTypeSize(FieldTy); - if (FieldSize >= 32) - return false; - } - return true; -} - /// Estimate number of registers the type will use when passed in registers. uint64_t AMDGPUABIInfo::numRegsForType(QualType Ty) const { uint64_t NumRegs = 0; @@ -212,27 +156,17 @@ ABIArgInfo AMDGPUABIInfo::classifyReturnType(QualType RetTy) const { RD && RD->hasFlexibleArrayMember()) return DefaultABIInfo::classifyReturnType(RetTy); - // Pack aggregates <= 8 bytes into single VGPR or pair, but only if they - // contain sub-32-bit integer types. Aggregates with floats or full-sized - // integers should preserve their original types. + // Pack aggregates <= 4 bytes into single VGPR or pair. uint64_t Size = getContext().getTypeSize(RetTy); + if (Size <= 16) + return ABIArgInfo::getDirect(llvm::Type::getInt16Ty(getVMContext())); + + if (Size <= 32) + return ABIArgInfo::getDirect(llvm::Type::getInt32Ty(getVMContext())); + if (Size <= 64) { - const RecordDecl *RD = RetTy->getAsRecordDecl(); - bool ShouldPackToInt = - RD && containsOnlyPackableIntegerTypes(RD, getContext()); - - if (ShouldPackToInt) { - if (Size <= 16) - return ABIArgInfo::getDirect( - llvm::Type::getInt16Ty(getVMContext())); - - if (Size <= 32) - return ABIArgInfo::getDirect( - llvm::Type::getInt32Ty(getVMContext())); - - llvm::Type *I32Ty = llvm::Type::getInt32Ty(getVMContext()); - return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); - } + llvm::Type *I32Ty = llvm::Type::getInt32Ty(getVMContext()); + return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); } if (numRegsForType(RetTy) <= MaxNumRegsForArgsRet) @@ -313,28 +247,21 @@ ABIArgInfo AMDGPUABIInfo::classifyArgumentType(QualType Ty, bool Variadic, RD && RD->hasFlexibleArrayMember()) return DefaultABIInfo::classifyArgumentType(Ty); - // Pack aggregates <= 8 bytes into single VGPR or pair, but only if they - // contain sub-32-bit integer types. Aggregates with floats or full-sized - // integers (i32, i64) should preserve their original types. + // Pack aggregates <= 8 bytes into single VGPR or pair. uint64_t Size = getContext().getTypeSize(Ty); if (Size <= 64) { - const RecordDecl *RD = Ty->getAsRecordDecl(); - bool ShouldPackToInt = - RD && containsOnlyPackableIntegerTypes(RD, getContext()); - - if (ShouldPackToInt) { - unsigned NumRegs = (Size + 31) / 32; - NumRegsLeft -= std::min(NumRegsLeft, NumRegs); + unsigned NumRegs = (Size + 31) / 32; + NumRegsLeft -= std::min(NumRegsLeft, NumRegs); - if (Size <= 16) - return ABIArgInfo::getDirect(llvm::Type::getInt16Ty(getVMContext())); + if (Size <= 16) + return ABIArgInfo::getDirect(llvm::Type::getInt16Ty(getVMContext())); - if (Size <= 32) - return ABIArgInfo::getDirect(llvm::Type::getInt32Ty(getVMContext())); + if (Size <= 32) + return ABIArgInfo::getDirect(llvm::Type::getInt32Ty(getVMContext())); - llvm::Type *I32Ty = llvm::Type::getInt32Ty(getVMContext()); - return ABIArgInfo::getDirect(llvm::ArrayType::get(I32Ty, 2)); - } + // 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) { diff --git a/clang/test/CodeGen/amdgpu-abi-struct-coerce.c b/clang/test/CodeGen/amdgpu-abi-struct-coerce.c deleted file mode 100644 index 2a1ebf0437f61..0000000000000 --- a/clang/test/CodeGen/amdgpu-abi-struct-coerce.c +++ /dev/null @@ -1,702 +0,0 @@ -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm -o - %s | FileCheck %s - -// Test AMDGPU ABI struct coercion behavior: -// - Structs containing ONLY sub-32-bit integers (char, short) should be packed into i32 registers -// - Structs containing floats or full-sized integers (i32, i64) should preserve their original types -// -// This tests the fix for the issue where structs like {float, int} were incorrectly -// coerced to [2 x i32], losing float type information. - -// ============================================================================ -// SECTION 1: Structs with floats - should NOT be coerced to integers -// ============================================================================ - -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; -} - -// Double precision floats -typedef struct double_struct { - double d; -} double_struct; - -// CHECK-LABEL: define{{.*}} double @return_double_struct(double %x.coerce) -double_struct return_double_struct(double_struct x) { - return x; -} - -// ============================================================================ -// SECTION 2: Structs with full-sized integers - should NOT be coerced -// ============================================================================ - -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; -} - -typedef struct single_int { - int a; -} single_int; - -// CHECK-LABEL: define{{.*}} i32 @return_single_int(i32 %x.coerce) -single_int return_single_int(single_int x) { - return x; -} - -typedef struct int64_struct { - long long a; -} int64_struct; - -// CHECK-LABEL: define{{.*}} i64 @return_int64_struct(i64 %x.coerce) -int64_struct return_int64_struct(int64_struct x) { - return x; -} - -// ============================================================================ -// SECTION 3: Structs with ONLY sub-32-bit integers - SHOULD be coerced -// ============================================================================ - -// Structs of small integers <= 32 bits should be coerced to i32 -typedef struct small_struct { - short a; - short b; -} small_struct; - -// CHECK-LABEL: define{{.*}} i32 @return_small_struct(i32 %x.coerce) -small_struct return_small_struct(small_struct x) { - return x; -} - -// Structs of small integers <= 16 bits should be coerced to i16 -typedef struct tiny_struct { - char a; - char b; -} tiny_struct; - -// CHECK-LABEL: define{{.*}} i16 @return_tiny_struct(i16 %x.coerce) -tiny_struct return_tiny_struct(tiny_struct x) { - return x; -} - -// Struct of 8 chars (64 bits) should be coerced to [2 x i32] -typedef struct eight_chars { - char a, b, c, d, e, f, g, h; -} eight_chars; - -// CHECK-LABEL: define{{.*}} [2 x i32] @return_eight_chars([2 x i32] %x.coerce) -eight_chars return_eight_chars(eight_chars x) { - return x; -} - -// Struct of 4 chars (32 bits) should be coerced to i32 -typedef struct four_chars { - char a, b, c, d; -} four_chars; - -// CHECK-LABEL: define{{.*}} i32 @return_four_chars(i32 %x.coerce) -four_chars return_four_chars(four_chars x) { - return x; -} - -// Struct of 4 shorts (64 bits) should be coerced to [2 x i32] -typedef struct four_shorts { - short a, b, c, d; -} four_shorts; - -// CHECK-LABEL: define{{.*}} [2 x i32] @return_four_shorts([2 x i32] %x.coerce) -four_shorts return_four_shorts(four_shorts x) { - return x; -} - -// ============================================================================ -// SECTION 4: Mixed types - floats prevent coercion even with small integers -// ============================================================================ - -typedef struct char_and_float { - char c; - float f; -} char_and_float; - -// CHECK-LABEL: define{{.*}} %struct.char_and_float @return_char_and_float(i8 %x.coerce0, float %x.coerce1) -// CHECK: ret %struct.char_and_float -char_and_float return_char_and_float(char_and_float x) { - return x; -} - -typedef struct short_and_float { - short s; - float f; -} short_and_float; - -// CHECK-LABEL: define{{.*}} %struct.short_and_float @return_short_and_float(i16 %x.coerce0, float %x.coerce1) -// CHECK: ret %struct.short_and_float -short_and_float return_short_and_float(short_and_float x) { - return x; -} - -// Small int + full-sized int should NOT be coerced -typedef struct char_and_int { - char c; - int i; -} char_and_int; - -// CHECK-LABEL: define{{.*}} %struct.char_and_int @return_char_and_int(i8 %x.coerce0, i32 %x.coerce1) -// CHECK: ret %struct.char_and_int -char_and_int return_char_and_int(char_and_int x) { - return x; -} - -// ============================================================================ -// SECTION 5: Exotic/Complex aggregates (per reviewer request) -// ============================================================================ - -// --- Nested structs --- - -typedef struct inner_chars { - char a, b; -} inner_chars; - -typedef struct outer_with_inner_chars { - inner_chars inner; - char c, d; -} outer_with_inner_chars; - -// All chars, 32 bits total - should be coerced to i32 -// CHECK-LABEL: define{{.*}} i32 @return_nested_chars(i32 %x.coerce) -outer_with_inner_chars return_nested_chars(outer_with_inner_chars x) { - return x; -} - -typedef struct inner_with_float { - char c; - float f; -} inner_with_float; - -typedef struct outer_with_float_inner { - inner_with_float inner; -} outer_with_float_inner; - -// Nested struct contains float - should NOT be coerced -// CHECK-LABEL: define{{.*}} %struct.outer_with_float_inner @return_nested_with_float(%struct.inner_with_float %x.coerce) -// CHECK: ret %struct.outer_with_float_inner -outer_with_float_inner return_nested_with_float(outer_with_float_inner x) { - return x; -} - -// --- Arrays within structs --- - -typedef struct char_array_struct { - char arr[4]; -} char_array_struct; - -// Array of 4 chars = 32 bits, all small ints - should be coerced to i32 -// CHECK-LABEL: define{{.*}} i32 @return_char_array(i32 %x.coerce) -char_array_struct return_char_array(char_array_struct x) { - return x; -} - -typedef struct short_array_struct { - short arr[2]; -} short_array_struct; - -// Array of 2 shorts = 32 bits, all small ints - should be coerced to i32 -// CHECK-LABEL: define{{.*}} i32 @return_short_array(i32 %x.coerce) -short_array_struct return_short_array(short_array_struct x) { - return x; -} - -typedef struct int_array_struct { - int arr[2]; -} int_array_struct; - -// Array of 2 ints = 64 bits, but ints are full-sized - should NOT be coerced -// CHECK-LABEL: define{{.*}} %struct.int_array_struct @return_int_array([2 x i32] %x.coerce) -// CHECK: ret %struct.int_array_struct -int_array_struct return_int_array(int_array_struct x) { - return x; -} - -typedef struct float_array_struct { - float arr[2]; -} float_array_struct; - -// Array of 2 floats - should NOT be coerced -// CHECK-LABEL: define{{.*}} %struct.float_array_struct @return_float_array([2 x float] %x.coerce) -// CHECK: ret %struct.float_array_struct -float_array_struct return_float_array(float_array_struct x) { - return x; -} - -// --- Complex combinations --- - -typedef struct mixed_nested { - struct { - char a; - char b; - } inner; - short s; -} mixed_nested; - -// All small integers (nested anonymous struct + short) = 32 bits - should be coerced -// CHECK-LABEL: define{{.*}} i32 @return_mixed_nested(i32 %x.coerce) -mixed_nested return_mixed_nested(mixed_nested x) { - return x; -} - -typedef struct deeply_nested_chars { - struct { - struct { - char a, b; - } level2; - char c, d; - } level1; -} deeply_nested_chars; - -// Deeply nested, but all chars = 32 bits - should be coerced -// CHECK-LABEL: define{{.*}} i32 @return_deeply_nested(i32 %x.coerce) -deeply_nested_chars return_deeply_nested(deeply_nested_chars x) { - return x; -} - -typedef struct deeply_nested_with_float { - struct { - struct { - char a; - float f; // Float buried deep - } level2; - } level1; -} deeply_nested_with_float; - -// Float buried in nested struct - should NOT be coerced -// CHECK-LABEL: define{{.*}} %struct.deeply_nested_with_float @return_deeply_nested_float -// CHECK: ret %struct.deeply_nested_with_float -deeply_nested_with_float return_deeply_nested_float(deeply_nested_with_float x) { - return x; -} - -// --- Edge cases --- - -// Single char -typedef struct single_char { - char c; -} single_char; - -// CHECK-LABEL: define{{.*}} i8 @return_single_char(i8 %x.coerce) -single_char return_single_char(single_char x) { - return x; -} - -// Three chars (24 bits, rounds up to 32) -typedef struct three_chars { - char a, b, c; -} three_chars; - -// CHECK-LABEL: define{{.*}} i32 @return_three_chars(i32 %x.coerce) -three_chars return_three_chars(three_chars x) { - return x; -} - -// Five chars (40 bits, rounds up to 64) -typedef struct five_chars { - char a, b, c, d, e; -} five_chars; - -// CHECK-LABEL: define{{.*}} [2 x i32] @return_five_chars([2 x i32] %x.coerce) -five_chars return_five_chars(five_chars x) { - return x; -} - -// --- Union tests --- - -typedef union char_int_union { - char c; - int i; -} char_int_union; - -// Union with int - preserves union type -// CHECK-LABEL: define{{.*}} %union.char_int_union @return_char_int_union(i32 %x.coerce) -char_int_union return_char_int_union(char_int_union x) { - return x; -} - -typedef union float_int_union { - float f; - int i; -} float_int_union; - -// Union with float - preserves union type -// CHECK-LABEL: define{{.*}} %union.float_int_union @return_float_int_union(float %x.coerce) -float_int_union return_float_int_union(float_int_union x) { - return x; -} - -// --- Padding scenarios --- - -typedef struct char_with_padding { - char c; - // 3 bytes padding - int i; -} char_with_padding; - -// Has int, should NOT be coerced even though small + padding -// CHECK-LABEL: define{{.*}} %struct.char_with_padding @return_char_with_padding(i8 %x.coerce0, i32 %x.coerce1) -// CHECK: ret %struct.char_with_padding -char_with_padding return_char_with_padding(char_with_padding x) { - return x; -} - -// ============================================================================ -// SECTION 6: Additional exotic aggregates -// ============================================================================ - -// --- Bitfields --- - -typedef struct bitfield_small { - unsigned a : 4; - unsigned b : 4; - unsigned c : 8; -} bitfield_small; - -// Bitfields with small bit-widths should be coerced to i32 -// Even though backing type is 'unsigned' (32 bits), the actual bit-widths are 4+4+8=16 bits -// CHECK-LABEL: define{{.*}} i32 @return_bitfield_small(i32 %x.coerce) -bitfield_small return_bitfield_small(bitfield_small x) { - return x; -} - -typedef struct bitfield_chars { - char a : 4; - char b : 4; -} bitfield_chars; - -// Bitfields with char backing type (8-bit) - should be coerced to i16 -// CHECK-LABEL: define{{.*}} i16 @return_bitfield_chars(i16 %x.coerce) -bitfield_chars return_bitfield_chars(bitfield_chars x) { - return x; -} - -typedef struct bitfield_with_int { - unsigned a : 4; - unsigned b : 4; - int i; -} bitfield_with_int; - -// Bitfields + full int - should NOT be coerced -// Bitfield packs into i8, then padding, then i32 -// CHECK-LABEL: define{{.*}} %struct.bitfield_with_int @return_bitfield_with_int(i8 %x.coerce0, i32 %x.coerce1) -// CHECK: ret %struct.bitfield_with_int -bitfield_with_int return_bitfield_with_int(bitfield_with_int x) { - return x; -} - -typedef struct bitfield_with_float { - unsigned a : 16; - float f; -} bitfield_with_float; - -// Bitfield + float - should NOT be coerced -// CHECK-LABEL: define{{.*}} %struct.bitfield_with_float @return_bitfield_with_float(i16 %x.coerce0, float %x.coerce1) -// CHECK: ret %struct.bitfield_with_float -bitfield_with_float return_bitfield_with_float(bitfield_with_float x) { - return x; -} - -// Bitfields that fill wider ints (up to i64) should also be packed -typedef struct bitfield_large { - unsigned long long a : 40; - unsigned long long b : 20; -} bitfield_large; - -// 40 + 20 = 60 bits, fits in 64-bit storage - should be coerced to [2 x i32] -// CHECK-LABEL: define{{.*}} [2 x i32] @return_bitfield_large([2 x i32] %x.coerce) -bitfield_large return_bitfield_large(bitfield_large x) { - return x; -} - -typedef struct bitfield_exactly_32 { - unsigned a : 16; - unsigned b : 16; -} bitfield_exactly_32; - -// 16 + 16 = 32 bits exactly - should be coerced to i32 -// CHECK-LABEL: define{{.*}} i32 @return_bitfield_exactly_32(i32 %x.coerce) -bitfield_exactly_32 return_bitfield_exactly_32(bitfield_exactly_32 x) { - return x; -} - -typedef struct bitfield_48 { - unsigned long long a : 32; - unsigned long long b : 16; -} bitfield_48; - -// 32 + 16 = 48 bits, stored in 64-bit - should be coerced to [2 x i32] -// CHECK-LABEL: define{{.*}} [2 x i32] @return_bitfield_48([2 x i32] %x.coerce) -bitfield_48 return_bitfield_48(bitfield_48 x) { - return x; -} - -// --- _Bool fields --- - -typedef struct bool_struct { - _Bool a; - _Bool b; - _Bool c; - _Bool d; -} bool_struct; - -// 4 bools = 32 bits, all sub-32-bit - should be coerced to i32 -// CHECK-LABEL: define{{.*}} i32 @return_bool_struct(i32 %x.coerce) -bool_struct return_bool_struct(bool_struct x) { - return x; -} - -typedef struct bool_and_float { - _Bool b; - float f; -} bool_and_float; - -// Bool + float - should NOT be coerced -// CHECK-LABEL: define{{.*}} %struct.bool_and_float @return_bool_and_float(i8 %x.coerce0, float %x.coerce1) -// CHECK: ret %struct.bool_and_float -bool_and_float return_bool_and_float(bool_and_float x) { - return x; -} - -typedef struct bool_and_int { - _Bool b; - int i; -} bool_and_int; - -// Bool + int - should NOT be coerced (int is full-sized) -// CHECK-LABEL: define{{.*}} %struct.bool_and_int @return_bool_and_int(i8 %x.coerce0, i32 %x.coerce1) -// CHECK: ret %struct.bool_and_int -bool_and_int return_bool_and_int(bool_and_int x) { - return x; -} - -// --- Half-precision floats --- - -typedef struct half_struct { - __fp16 a; - __fp16 b; -} half_struct; - -// Two halfs = 32 bits, but floats - should NOT be coerced -// CHECK-LABEL: define{{.*}} %struct.half_struct @return_half_struct(half %x.coerce0, half %x.coerce1) -// CHECK: ret %struct.half_struct -half_struct return_half_struct(half_struct x) { - return x; -} - -typedef struct half_and_char { - __fp16 h; - char c; -} half_and_char; - -// Half + char - should NOT be coerced (half is float type) -// CHECK-LABEL: define{{.*}} %struct.half_and_char @return_half_and_char(half %x.coerce0, i8 %x.coerce1) -// CHECK: ret %struct.half_and_char -half_and_char return_half_and_char(half_and_char x) { - return x; -} - -typedef struct four_halfs { - __fp16 a, b, c, d; -} four_halfs; - -// Four halfs = 64 bits - should NOT be coerced -// CHECK-LABEL: define{{.*}} %struct.four_halfs @return_four_halfs(half %x.coerce0, half %x.coerce1, half %x.coerce2, half %x.coerce3) -// CHECK: ret %struct.four_halfs -four_halfs return_four_halfs(four_halfs x) { - return x; -} - -// --- Bfloat16 tests --- - -typedef struct bfloat_struct { - __bf16 a; - __bf16 b; -} bfloat_struct; - -// Two bfloats = 32 bits, but floats - should NOT be coerced -// CHECK-LABEL: define{{.*}} %struct.bfloat_struct @return_bfloat_struct(bfloat %x.coerce0, bfloat %x.coerce1) -// CHECK: ret %struct.bfloat_struct -bfloat_struct return_bfloat_struct(bfloat_struct x) { - return x; -} - -typedef struct bfloat_and_char { - __bf16 b; - char c; -} bfloat_and_char; - -// Bfloat + char - should NOT be coerced (bfloat is float type) -// CHECK-LABEL: define{{.*}} %struct.bfloat_and_char @return_bfloat_and_char(bfloat %x.coerce0, i8 %x.coerce1) -// CHECK: ret %struct.bfloat_and_char -bfloat_and_char return_bfloat_and_char(bfloat_and_char x) { - return x; -} - -typedef struct four_bfloats { - __bf16 a, b, c, d; -} four_bfloats; - -// Four bfloats = 64 bits - should NOT be coerced -// CHECK-LABEL: define{{.*}} %struct.four_bfloats @return_four_bfloats(bfloat %x.coerce0, bfloat %x.coerce1, bfloat %x.coerce2, bfloat %x.coerce3) -// CHECK: ret %struct.four_bfloats -four_bfloats return_four_bfloats(four_bfloats x) { - return x; -} - -// --- Mixed half and bfloat --- - -typedef struct mixed_half_bfloat { - __fp16 h; - __bf16 b; -} mixed_half_bfloat; - -// Mixed half + bfloat - should NOT be coerced -// CHECK-LABEL: define{{.*}} %struct.mixed_half_bfloat @return_mixed_half_bfloat(half %x.coerce0, bfloat %x.coerce1) -// CHECK: ret %struct.mixed_half_bfloat -mixed_half_bfloat return_mixed_half_bfloat(mixed_half_bfloat x) { - return x; -} - -typedef struct bfloat_and_float { - __bf16 b; - float f; -} bfloat_and_float; - -// Bfloat + float - should NOT be coerced -// CHECK-LABEL: define{{.*}} %struct.bfloat_and_float @return_bfloat_and_float(bfloat %x.coerce0, float %x.coerce1) -// CHECK: ret %struct.bfloat_and_float -bfloat_and_float return_bfloat_and_float(bfloat_and_float x) { - return x; -} - -// --- Vectors inside structs --- - -typedef int int2 __attribute__((ext_vector_type(2))); -typedef float float2 __attribute__((ext_vector_type(2))); -typedef char char4 __attribute__((ext_vector_type(4))); - -typedef struct vec_int2_struct { - int2 v; -} vec_int2_struct; - -// Single-element vector struct - unwrapped to vector type -// CHECK-LABEL: define{{.*}} <2 x i32> @return_vec_int2(<2 x i32> %x.coerce) -vec_int2_struct return_vec_int2(vec_int2_struct x) { - return x; -} - -typedef struct vec_float2_struct { - float2 v; -} vec_float2_struct; - -// Single-element vector struct - unwrapped to vector type -// CHECK-LABEL: define{{.*}} <2 x float> @return_vec_float2(<2 x float> %x.coerce) -vec_float2_struct return_vec_float2(vec_float2_struct x) { - return x; -} - -typedef struct vec_char4_struct { - char4 v; -} vec_char4_struct; - -// Single-element vector struct - unwrapped to vector type -// CHECK-LABEL: define{{.*}} <4 x i8> @return_vec_char4(<4 x i8> %x.coerce) -vec_char4_struct return_vec_char4(vec_char4_struct x) { - return x; -} - -typedef struct vec_and_scalar { - char4 v; - int i; -} vec_and_scalar; - -// Vector + scalar - should NOT be coerced (vector is not a packable integer type) -// CHECK-LABEL: define{{.*}} %struct.vec_and_scalar @return_vec_and_scalar(<4 x i8> %x.coerce0, i32 %x.coerce1) -// CHECK: ret %struct.vec_and_scalar -vec_and_scalar return_vec_and_scalar(vec_and_scalar x) { - return x; -} - -// --- Arrays of nested structs --- - -typedef struct inner_two_chars { - char a, b; -} inner_two_chars; - -typedef struct array_of_nested_chars { - inner_two_chars arr[2]; -} array_of_nested_chars; - -// Array of 2 nested structs, each with 2 chars = 32 bits total - should be coerced -// CHECK-LABEL: define{{.*}} i32 @return_array_of_nested_chars(i32 %x.coerce) -array_of_nested_chars return_array_of_nested_chars(array_of_nested_chars x) { - return x; -} - -typedef struct inner_char_float { - char c; - float f; -} inner_char_float; - -typedef struct array_of_nested_floats { - inner_char_float arr[1]; -} array_of_nested_floats; - -// Array of nested struct containing float - should NOT be coerced -// CHECK-LABEL: define{{.*}} %struct.array_of_nested_floats @return_array_of_nested_floats([1 x %struct.inner_char_float] %x.coerce) -// CHECK: ret %struct.array_of_nested_floats -array_of_nested_floats return_array_of_nested_floats(array_of_nested_floats x) { - return x; -} - -typedef struct nested_array_of_shorts { - struct { - short arr[2]; - } inner; -} nested_array_of_shorts; - -// Nested struct with array of shorts = 32 bits - should be coerced -// CHECK-LABEL: define{{.*}} i32 @return_nested_array_of_shorts(i32 %x.coerce) -nested_array_of_shorts return_nested_array_of_shorts(nested_array_of_shorts x) { - return x; -} diff --git a/clang/test/CodeGen/amdgpu-variadic-call.c b/clang/test/CodeGen/amdgpu-variadic-call.c index 22402118d862f..17eda215211a2 100644 --- a/clang/test/CodeGen/amdgpu-variadic-call.c +++ b/clang/test/CodeGen/amdgpu-variadic-call.c @@ -217,9 +217,10 @@ typedef union } union_f32_i32; // CHECK-LABEL: define {{[^@]+}}@one_pair_union_f32_i32 -// CHECK-SAME: (i32 noundef [[F0:%.*]], double noundef [[F1:%.*]], float [[V0_COERCE:%.*]]) local_unnamed_addr #[[ATTR0]] { +// CHECK-SAME: (i32 noundef [[F0:%.*]], double noundef [[F1:%.*]], i32 [[V0_COERCE:%.*]]) local_unnamed_addr #[[ATTR0]] { // CHECK-NEXT: entry: -// CHECK-NEXT: [[DOTFCA_0_INSERT:%.*]] = insertvalue [[UNION_UNION_F32_I32:%.*]] poison, float [[V0_COERCE]], 0 +// CHECK-NEXT: [[TMP0:%.*]] = bitcast i32 [[V0_COERCE]] to float +// CHECK-NEXT: [[DOTFCA_0_INSERT:%.*]] = insertvalue [[UNION_UNION_F32_I32:%.*]] poison, float [[TMP0]], 0 // CHECK-NEXT: tail call void (...) @sink_0([[UNION_UNION_F32_I32]] [[DOTFCA_0_INSERT]]) #[[ATTR2]] // CHECK-NEXT: tail call void (i32, ...) @sink_1(i32 noundef [[F0]], [[UNION_UNION_F32_I32]] [[DOTFCA_0_INSERT]]) #[[ATTR2]] // CHECK-NEXT: tail call void (double, i32, ...) @sink_2(double noundef [[F1]], i32 noundef [[F0]], [[UNION_UNION_F32_I32]] [[DOTFCA_0_INSERT]]) #[[ATTR2]] @@ -272,12 +273,13 @@ void multiple_one(int f0, double f1, int v0, double v1) } // CHECK-LABEL: define {{[^@]+}}@multiple_two -// CHECK-SAME: (i32 noundef [[F0:%.*]], double noundef [[F1:%.*]], double [[V0_COERCE0:%.*]], double [[V0_COERCE1:%.*]], float noundef [[V1:%.*]], float [[V2_COERCE:%.*]], i32 noundef [[V3:%.*]]) local_unnamed_addr #[[ATTR0]] { +// CHECK-SAME: (i32 noundef [[F0:%.*]], double noundef [[F1:%.*]], double [[V0_COERCE0:%.*]], double [[V0_COERCE1:%.*]], float noundef [[V1:%.*]], i32 [[V2_COERCE:%.*]], i32 noundef [[V3:%.*]]) local_unnamed_addr #[[ATTR0]] { // CHECK-NEXT: entry: +// CHECK-NEXT: [[TMP0:%.*]] = bitcast i32 [[V2_COERCE]] to float // CHECK-NEXT: [[CONV:%.*]] = fpext float [[V1]] to double // CHECK-NEXT: [[DOTFCA_0_INSERT16:%.*]] = insertvalue [[STRUCT_PAIR_F64:%.*]] poison, double [[V0_COERCE0]], 0 // CHECK-NEXT: [[DOTFCA_1_INSERT:%.*]] = insertvalue [[STRUCT_PAIR_F64]] [[DOTFCA_0_INSERT16]], double [[V0_COERCE1]], 1 -// CHECK-NEXT: [[DOTFCA_0_INSERT:%.*]] = insertvalue [[UNION_UNION_F32_I32:%.*]] poison, float [[V2_COERCE]], 0 +// CHECK-NEXT: [[DOTFCA_0_INSERT:%.*]] = insertvalue [[UNION_UNION_F32_I32:%.*]] poison, float [[TMP0]], 0 // CHECK-NEXT: tail call void (...) @sink_0([[STRUCT_PAIR_F64]] [[DOTFCA_1_INSERT]], double noundef [[CONV]], [[UNION_UNION_F32_I32]] [[DOTFCA_0_INSERT]], i32 noundef [[V3]]) #[[ATTR2]] // CHECK-NEXT: tail call void (i32, ...) @sink_1(i32 noundef [[F0]], [[STRUCT_PAIR_F64]] [[DOTFCA_1_INSERT]], double noundef [[CONV]], [[UNION_UNION_F32_I32]] [[DOTFCA_0_INSERT]], i32 noundef [[V3]]) #[[ATTR2]] // CHECK-NEXT: tail call void (double, i32, ...) @sink_2(double noundef [[F1]], i32 noundef [[F0]], [[STRUCT_PAIR_F64]] [[DOTFCA_1_INSERT]], double noundef [[CONV]], [[UNION_UNION_F32_I32]] [[DOTFCA_0_INSERT]], i32 noundef [[V3]]) #[[ATTR2]] diff --git a/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl b/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl index e9cdb7f5da32a..06d3cdb01deb2 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl @@ -431,8 +431,8 @@ struct_char_arr32 func_ret_struct_char_arr32() return s; } -// CHECK: define{{.*}} %union.transparent_u @func_transparent_union_ret() local_unnamed_addr #[[ATTR1:[0-9]+]] { -// CHECK: ret %union.transparent_u zeroinitializer +// CHECK: define{{.*}} i32 @func_transparent_union_ret() local_unnamed_addr #[[ATTR1:[0-9]+]] { +// CHECK: ret i32 0 transparent_u func_transparent_union_ret() { transparent_u u = { 0 }; diff --git a/clang/test/Headers/amdgcn-openmp-device-math-complex.c b/clang/test/Headers/amdgcn-openmp-device-math-complex.c index 34c05e2974a64..b347cf4716df2 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 { float, float } @__divsc3 - // CHECK: call { float, float } @__mulsc3 + // CHECK: call [2 x i32] @__divsc3 + // CHECK: call [2 x i32] @__mulsc3 (void)(a * (a / a)); } } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
