https://github.com/Pierre-vh updated
https://github.com/llvm/llvm-project/pull/141053
>From 68db9fad42369be31d935257a2d80962a4018892 Mon Sep 17 00:00:00 2001
From: pvanhout <pierre.vanhoutr...@amd.com>
Date: Thu, 22 May 2025 14:13:54 +0200
Subject: [PATCH 1/5] [clang][CodeGen] Fix crash on non-natural type in
CheckAtomicAlignment
In some specific scenarios, `Ptr.getElementType()` won't be a primitive
type or a vector of primitive types, and thus `getScalarSizeInBits()` returns
zero.
Use the datalayout to get the proper size of the type instead of making an
implicit
assumption that the type is a simple primitive type.
Solves SWDEV-534184
---
clang/lib/CodeGen/CGBuiltin.cpp | 3 +-
.../CodeGenOpenCL/check-atomic-alignment.cl | 45 +++++++++++++++++++
2 files changed, 47 insertions(+), 1 deletion(-)
create mode 100644 clang/test/CodeGenOpenCL/check-atomic-alignment.cl
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 04a0d9ba2bbce..749f716acfbe8 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -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.getTypeSizeInBits(Ptr.getElementType()) / 8;
unsigned Align = Ptr.getAlignment().getQuantity();
if (Align % Bytes != 0) {
DiagnosticsEngine &Diags = CGF.CGM.getDiags();
diff --git a/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
b/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
new file mode 100644
index 0000000000000..6aa4bc26633f8
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
@@ -0,0 +1,45 @@
+// 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 - | 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 noundef [[ADDR:%.*]], <2 x half> noundef [[VAL:%.*]])
#[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[RETVAL:%.*]] = alloca <2 x half>, align 4, addrspace(5)
+// CHECK-NEXT: [[ADDR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[VAL_ADDR:%.*]] = alloca <2 x half>, align 4, addrspace(5)
+// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[RETVAL]] to ptr
+// CHECK-NEXT: [[ADDR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[ADDR_ADDR]] to ptr
+// CHECK-NEXT: [[VAL_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[VAL_ADDR]] to ptr
+// CHECK-NEXT: store ptr [[ADDR]], ptr [[ADDR_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store <2 x half> [[VAL]], ptr [[VAL_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ADDR_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load <2 x half>, ptr [[VAL_ADDR_ASCAST]],
align 4
+// CHECK-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[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(__generic short2 *addr, half2 val) {
+ return __builtin_amdgcn_flat_atomic_fadd_v2f16((struct __half2*)addr, val);
+}
+//.
+// CHECK: [[META4]] = !{}
+//.
>From 11ddc4d93d36d1605390f9b4ffcf3233c342c90c Mon Sep 17 00:00:00 2001
From: pvanhout <pierre.vanhoutr...@amd.com>
Date: Thu, 22 May 2025 14:46:37 +0200
Subject: [PATCH 2/5] comments
---
clang/lib/CodeGen/CGBuiltin.cpp | 2 +-
.../CodeGenOpenCL/check-atomic-alignment.cl | 26 +++++++------------
2 files changed, 11 insertions(+), 17 deletions(-)
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 749f716acfbe8..809ffe549be88 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -277,7 +277,7 @@ Address CheckAtomicAlignment(CodeGenFunction &CGF, const
CallExpr *E) {
const llvm::DataLayout &DL = CGF.CGM.getDataLayout();
unsigned Bytes = Ptr.getElementType()->isPointerTy()
? Ctx.getTypeSizeInChars(Ctx.VoidPtrTy).getQuantity()
- : DL.getTypeSizeInBits(Ptr.getElementType()) / 8;
+ : DL.getTypeStoreSize(Ptr.getElementType());
unsigned Align = Ptr.getAlignment().getQuantity();
if (Align % Bytes != 0) {
DiagnosticsEngine &Diags = CGF.CGM.getDiags();
diff --git a/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
b/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
index 6aa4bc26633f8..a536d284f8016 100644
--- a/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
+++ b/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
@@ -1,8 +1,6 @@
// 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 - | FileCheck %s
-
-// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -O0 -cl-std=CL1.2 -triple amdgcn-amd-amdhsa -target-cpu
gfx942 \
+// RUN: %s -emit-llvm -o - -disable-llvm-passes | FileCheck %s
// `Ptr.getElementType()` in `CheckAtomicAlignment` returns
// %struct.__half2 = type { %union.anon }
@@ -22,22 +20,18 @@ struct __half2 {
};
// CHECK-LABEL: define dso_local <2 x half> @test_flat_add_2f16(
-// CHECK-SAME: ptr noundef [[ADDR:%.*]], <2 x half> noundef [[VAL:%.*]])
#[[ATTR0:[0-9]+]] {
+// CHECK-SAME: ptr addrspace(5) noundef [[ADDR:%.*]], <2 x half> noundef
[[VAL:%.*]]) #[[ATTR0:[0-9]+]] {
// CHECK-NEXT: [[ENTRY:.*:]]
-// CHECK-NEXT: [[RETVAL:%.*]] = alloca <2 x half>, align 4, addrspace(5)
-// CHECK-NEXT: [[ADDR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// 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: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[RETVAL]] to ptr
-// CHECK-NEXT: [[ADDR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[ADDR_ADDR]] to ptr
-// CHECK-NEXT: [[VAL_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[VAL_ADDR]] to ptr
-// CHECK-NEXT: store ptr [[ADDR]], ptr [[ADDR_ADDR_ASCAST]], align 8
-// CHECK-NEXT: store <2 x half> [[VAL]], ptr [[VAL_ADDR_ASCAST]], align 4
-// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ADDR_ADDR_ASCAST]], align 8
-// CHECK-NEXT: [[TMP1:%.*]] = load <2 x half>, ptr [[VAL_ADDR_ASCAST]],
align 4
-// CHECK-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], <2 x half>
[[TMP1]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory
[[META4:![0-9]+]]
+// 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(__generic short2 *addr, half2 val) {
+half2 test_flat_add_2f16(short2 *addr, half2 val) {
return __builtin_amdgcn_flat_atomic_fadd_v2f16((struct __half2*)addr, val);
}
//.
>From 5062f00419562d0b7253d0855cce9d5ab52c4cde Mon Sep 17 00:00:00 2001
From: pvanhout <pierre.vanhoutr...@amd.com>
Date: Thu, 22 May 2025 14:47:21 +0200
Subject: [PATCH 3/5] restore amdgpu-registered-target
---
clang/test/CodeGenOpenCL/check-atomic-alignment.cl | 2 ++
1 file changed, 2 insertions(+)
diff --git a/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
b/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
index a536d284f8016..3b2d56f83b345 100644
--- a/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
+++ b/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
@@ -2,6 +2,8 @@
// RUN: %clang_cc1 -O0 -cl-std=CL1.2 -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.
>From bfee7a5d02fa36de9cb1bfb4a12f6b735dff4961 Mon Sep 17 00:00:00 2001
From: pvanhout <pierre.vanhoutr...@amd.com>
Date: Thu, 22 May 2025 14:51:40 +0200
Subject: [PATCH 4/5] Use CL2.0
---
clang/test/CodeGenOpenCL/check-atomic-alignment.cl | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
b/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
index 3b2d56f83b345..aaaa60a35c374 100644
--- a/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
+++ b/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
@@ -1,5 +1,5 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
UTC_ARGS: --version 5
-// RUN: %clang_cc1 -O0 -cl-std=CL1.2 -triple amdgcn-amd-amdhsa -target-cpu
gfx942 \
+// 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
>From 50ceb4fd027e7c21960250a2415813a69c07bd13 Mon Sep 17 00:00:00 2001
From: pvanhout <pierre.vanhoutr...@amd.com>
Date: Thu, 22 May 2025 16:09:43 +0200
Subject: [PATCH 5/5] fix test
---
.../CodeGenOpenCL/check-atomic-alignment.cl | 18 +++++++++++-------
1 file changed, 11 insertions(+), 7 deletions(-)
diff --git a/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
b/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
index aaaa60a35c374..89771d20ad090 100644
--- a/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
+++ b/clang/test/CodeGenOpenCL/check-atomic-alignment.cl
@@ -22,15 +22,19 @@ struct __half2 {
};
// 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-SAME: ptr 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: [[RETVAL:%.*]] = alloca <2 x half>, align 4, addrspace(5)
+// CHECK-NEXT: [[ADDR_ADDR:%.*]] = alloca ptr, align 8, 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: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[RETVAL]] to ptr
+// CHECK-NEXT: [[ADDR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[ADDR_ADDR]] to ptr
+// CHECK-NEXT: [[VAL_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[VAL_ADDR]] to ptr
+// CHECK-NEXT: store ptr [[ADDR]], ptr [[ADDR_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store <2 x half> [[VAL]], ptr [[VAL_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[ADDR_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load <2 x half>, ptr [[VAL_ADDR_ASCAST]],
align 4
+// CHECK-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[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) {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
