Skip to content

[clang][CodeGen] Fix crash on non-natural type in CheckAtomicAlignment #141053

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 5 commits into from
May 22, 2025
Merged

[clang][CodeGen] Fix crash on non-natural type in CheckAtomicAlignment #141053

Show file tree
Hide file tree
Changes from 4 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
68db9fa
[clang][CodeGen] Fix crash on non-natural type in CheckAtomicAlignment
Pierre-vh May 22, 2025
11ddc4d
comments
Pierre-vh May 22, 2025
5062f00
restore amdgpu-registered-target
Pierre-vh May 22, 2025
bfee7a5
Use CL2.0
Pierre-vh May 22, 2025
50ceb4f
fix test
Pierre-vh May 22, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion clang/lib/CodeGen/CGBuiltin.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -274,9 +274,10 @@ Value *EmitFromInt(CodeGenFunction &CGF, llvm::Value *V,
Address CheckAtomicAlignment(CodeGenFunction &CGF, const CallExpr *E) {
ASTContext &Ctx = CGF.getContext();
Address Ptr = CGF.EmitPointerWithAlignment(E->getArg(0));
const llvm::DataLayout &DL = CGF.CGM.getDataLayout();
unsigned Bytes = Ptr.getElementType()->isPointerTy()
? Ctx.getTypeSizeInChars(Ctx.VoidPtrTy).getQuantity()
: Ptr.getElementType()->getScalarSizeInBits() / 8;
: DL.getTypeStoreSize(Ptr.getElementType());
unsigned Align = Ptr.getAlignment().getQuantity();
if (Align % Bytes != 0) {
DiagnosticsEngine &Diags = CGF.CGM.getDiags();
Expand Down
41 changes: 41 additions & 0 deletions clang/test/CodeGenOpenCL/check-atomic-alignment.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,41 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx942 \
// RUN: %s -emit-llvm -o - -disable-llvm-passes | FileCheck %s

// REQUIRES: amdgpu-registered-target

// `Ptr.getElementType()` in `CheckAtomicAlignment` returns
// %struct.__half2 = type { %union.anon }
// Check we do not crash when handling that.

typedef half __attribute__((ext_vector_type(2))) half2;
typedef short __attribute__((ext_vector_type(2))) short2;

struct __half2 {
union {
struct {
half x;
half y;
};
half2 data;
};
};

// CHECK-LABEL: define dso_local <2 x half> @test_flat_add_2f16(
// CHECK-SAME: ptr addrspace(5) noundef [[ADDR:%.*]], <2 x half> noundef [[VAL:%.*]]) #[[ATTR0:[0-9]+]] {
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[ADDR_ADDR:%.*]] = alloca ptr addrspace(5), align 4, addrspace(5)
// CHECK-NEXT: [[VAL_ADDR:%.*]] = alloca <2 x half>, align 4, addrspace(5)
// CHECK-NEXT: store ptr addrspace(5) [[ADDR]], ptr addrspace(5) [[ADDR_ADDR]], align 4
// CHECK-NEXT: store <2 x half> [[VAL]], ptr addrspace(5) [[VAL_ADDR]], align 4
// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(5), ptr addrspace(5) [[ADDR_ADDR]], align 4
// CHECK-NEXT: [[TMP1:%.*]] = load <2 x half>, ptr addrspace(5) [[VAL_ADDR]], align 4
// CHECK-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(5) [[TMP0]], <2 x half> [[TMP1]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4:![0-9]+]]
// CHECK-NEXT: ret <2 x half> [[TMP2]]
//
half2 test_flat_add_2f16(short2 *addr, half2 val) {
return __builtin_amdgcn_flat_atomic_fadd_v2f16((struct __half2*)addr, val);
}
//.
// CHECK: [[META4]] = !{}
//.
Loading