4 Commits

Author SHA1 Message Date
Eli Friedman
1762e01cca
Fix codegen of consteval functions returning an empty class, and related issues (#93115)
Fix codegen of consteval functions returning an empty class, and related
issues

If a class is empty, don't store it to memory: the store might overwrite
useful data. Similarly, if a class has tail padding that might overlap
other fields, don't store the tail padding to memory.

The problem here turned out a bit more general than I initially thought:
basically all uses of EmitAggregateStore were broken. Call lowering had
a method that did mostly the right thing, though: CreateCoercedStore.
Adapt CreateCoercedStore so it always does the conservatively right
thing, and use it for both calls and ConstantExpr.

Also, along the way, fix the "overlap" bit in AggValueSlot: the bit was
set incorrectly for empty classes in some cases.

Fixes #93040.
2024-08-01 16:18:20 -07:00
Pravin Jagtap
1f21e49870
Revert "Revert "[AMDGPU] const-fold imm operands of (#71669)
amdgcn_update_dpp intrinsic (#71139)""

This reverts commit d1fb9307951319eea3e869d78470341d603c8363 and fixes
the lit test clang/test/CodeGenHIP/dpp-const-fold.hip

---------

Authored-by: Pravin Jagtap <Pravin.Jagtap@amd.com>
2023-11-09 10:09:22 +05:30
Mitch Phillips
d1fb930795 Revert "[AMDGPU] const-fold imm operands of amdgcn_update_dpp intrinsic (#71139)"
This reverts commit 32a3f2afe6ea7ffb02a6a188b123ded6f4c89f6c.

Reason: Broke the sanitizer buildbots. More details at
32a3f2afe6
2023-11-08 12:50:53 +01:00
Pravin Jagtap
32a3f2afe6
[AMDGPU] const-fold imm operands of amdgcn_update_dpp intrinsic (#71139)
Operands of `__builtin_amdgcn_update_dpp` need to evaluate to constant
to match the intrinsic requirements.

Fixes: SWDEV-426822, SWDEV-431138
---------

Authored-by: Pravin Jagtap <Pravin.Jagtap@amd.com>
2023-11-08 15:09:10 +05:30