llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang-codegen @llvm/pr-subscribers-clang Author: Pierre van Houtryve (Pierre-vh) <details> <summary>Changes</summary> 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 --- Full diff: https://github.com/llvm/llvm-project/pull/141053.diff 2 Files Affected: - (modified) clang/lib/CodeGen/CGBuiltin.cpp (+2-1) - (added) clang/test/CodeGenOpenCL/check-atomic-alignment.cl (+45) ``````````diff 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]] = !{} +//. `````````` </details> https://github.com/llvm/llvm-project/pull/141053 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
