https://github.com/changpeng updated https://github.com/llvm/llvm-project/pull/92612
>From 2468a85a47499d90a99610846c632332eb7307b8 Mon Sep 17 00:00:00 2001 From: Changpeng Fang <changpeng.f...@amd.com> Date: Fri, 17 May 2024 15:13:07 -0700 Subject: [PATCH 1/3] [OpenCL] Fix an infinite loop in builidng AddrSpaceQualType In building AddrSpaceQualType (https://github.com/llvm/llvm-project/pull/90048), there is a bug in removeAddrSpaceQualType() for arrays. Arrays are weird because qualifiers on the element type also count as qualifiers on the type, so getSingleStepDesugaredType() can't remove the sugar on arrays. This results in an infinite loop in removeAddrSpaceQualType. To fix the issue, we use ASTContext::getUnqualifiedArrayType, which strips the qualifier off the element type, then reconstruct the array type. --- clang/lib/CodeGen/CGExprAgg.cpp | 3 ++- .../array-type-infinite-loop.clcpp | 25 +++++++++++++++++++ 2 files changed, 27 insertions(+), 1 deletion(-) create mode 100644 clang/test/CodeGenOpenCLCXX/array-type-infinite-loop.clcpp diff --git a/clang/lib/CodeGen/CGExprAgg.cpp b/clang/lib/CodeGen/CGExprAgg.cpp index 6172eb9cdc1bb..53ce133e8cbc6 100644 --- a/clang/lib/CodeGen/CGExprAgg.cpp +++ b/clang/lib/CodeGen/CGExprAgg.cpp @@ -537,8 +537,9 @@ void AggExprEmitter::EmitArrayInit(Address DestPtr, llvm::ArrayType *AType, elementType.isTriviallyCopyableType(CGF.getContext())) { CodeGen::CodeGenModule &CGM = CGF.CGM; ConstantEmitter Emitter(CGF); + Qualifiers Quals; QualType GVArrayQTy = CGM.getContext().getAddrSpaceQualType( - CGM.getContext().removeAddrSpaceQualType(ArrayQTy), + CGM.getContext().getUnqualifiedArrayType(ArrayQTy, Quals), CGM.GetGlobalConstantAddressSpace()); LangAS AS = GVArrayQTy.getAddressSpace(); if (llvm::Constant *C = diff --git a/clang/test/CodeGenOpenCLCXX/array-type-infinite-loop.clcpp b/clang/test/CodeGenOpenCLCXX/array-type-infinite-loop.clcpp new file mode 100644 index 0000000000000..5a5b104e892f7 --- /dev/null +++ b/clang/test/CodeGenOpenCLCXX/array-type-infinite-loop.clcpp @@ -0,0 +1,25 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 4 +//RUN: %clang_cc1 %s -emit-llvm -O1 -o - | FileCheck %s + +// CHECK-LABEL: define dso_local spir_kernel void @test( +// CHECK-SAME: ptr nocapture noundef readonly align 8 [[IN:%.*]], ptr nocapture noundef writeonly align 8 [[OUT:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !kernel_arg_addr_space [[META3:![0-9]+]] !kernel_arg_access_qual [[META4:![0-9]+]] !kernel_arg_type [[META5:![0-9]+]] !kernel_arg_base_type [[META5]] !kernel_arg_type_qual [[META6:![0-9]+]] { +// CHECK-NEXT: entry: +// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds i8, ptr [[IN]], i64 8 +// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[ARRAYIDX1]], align 8, !tbaa [[TBAA7:![0-9]+]] +// CHECK-NEXT: store i64 [[TMP0]], ptr [[OUT]], align 8, !tbaa [[TBAA7]] +// CHECK-NEXT: ret void +// +__kernel void test(__global long *In, __global long *Out) { + long m[4] = { In[0], In[1], 0, 0 }; + *Out = m[1]; +} +//. +// CHECK: [[META3]] = !{i32 1, i32 1} +// CHECK: [[META4]] = !{!"none", !"none"} +// CHECK: [[META5]] = !{!"long*", !"long*"} +// CHECK: [[META6]] = !{!"", !""} +// CHECK: [[TBAA7]] = !{[[META8:![0-9]+]], [[META8]], i64 0} +// CHECK: [[META8]] = !{!"long", [[META9:![0-9]+]], i64 0} +// CHECK: [[META9]] = !{!"omnipotent char", [[META10:![0-9]+]], i64 0} +// CHECK: [[META10]] = !{!"Simple C++ TBAA"} +//. >From 17ac766cdcbf22af685b89b9a054a22afb42f46e Mon Sep 17 00:00:00 2001 From: Changpeng Fang <changpeng.f...@amd.com> Date: Fri, 17 May 2024 18:20:06 -0700 Subject: [PATCH 2/3] [OpenCL] Fix an infinite loop in builidng AddrSpaceQualType Fix ASTContext::removeAddrSpaceQualType() --- clang/include/clang/AST/ASTContext.h | 2 +- clang/lib/AST/ASTContext.cpp | 9 ++++++++- clang/lib/CodeGen/CGExprAgg.cpp | 3 +-- 3 files changed, 10 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h index e03b112194786..2ce2b810d3636 100644 --- a/clang/include/clang/AST/ASTContext.h +++ b/clang/include/clang/AST/ASTContext.h @@ -2611,7 +2611,7 @@ class ASTContext : public RefCountedBase<ASTContext> { /// /// \returns if this is an array type, the completely unqualified array type /// that corresponds to it. Otherwise, returns T.getUnqualifiedType(). - QualType getUnqualifiedArrayType(QualType T, Qualifiers &Quals); + QualType getUnqualifiedArrayType(QualType T, Qualifiers &Quals) const; /// Determine whether the given types are equivalent after /// cvr-qualifiers have been removed. diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 8fc2bb8c401c2..388233c554d46 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -3054,6 +3054,13 @@ QualType ASTContext::removeAddrSpaceQualType(QualType T) const { if (!T.hasAddressSpace()) return T; + // For arrays, strip the qualifier off the element type, then reconstruct the + // array type + if (T.getTypePtr()->isArrayType()) { + Qualifiers Qualfs; + return getUnqualifiedArrayType(T, Qualfs); + } + // If we are composing extended qualifiers together, merge together // into one ExtQuals node. QualifierCollector Quals; @@ -6093,7 +6100,7 @@ CanQualType ASTContext::getCanonicalParamType(QualType T) const { } QualType ASTContext::getUnqualifiedArrayType(QualType type, - Qualifiers &quals) { + Qualifiers &quals) const { SplitQualType splitType = type.getSplitUnqualifiedType(); // FIXME: getSplitUnqualifiedType() actually walks all the way to diff --git a/clang/lib/CodeGen/CGExprAgg.cpp b/clang/lib/CodeGen/CGExprAgg.cpp index 53ce133e8cbc6..6172eb9cdc1bb 100644 --- a/clang/lib/CodeGen/CGExprAgg.cpp +++ b/clang/lib/CodeGen/CGExprAgg.cpp @@ -537,9 +537,8 @@ void AggExprEmitter::EmitArrayInit(Address DestPtr, llvm::ArrayType *AType, elementType.isTriviallyCopyableType(CGF.getContext())) { CodeGen::CodeGenModule &CGM = CGF.CGM; ConstantEmitter Emitter(CGF); - Qualifiers Quals; QualType GVArrayQTy = CGM.getContext().getAddrSpaceQualType( - CGM.getContext().getUnqualifiedArrayType(ArrayQTy, Quals), + CGM.getContext().removeAddrSpaceQualType(ArrayQTy), CGM.GetGlobalConstantAddressSpace()); LangAS AS = GVArrayQTy.getAddressSpace(); if (llvm::Constant *C = >From dec36a5150455995e3e0e07d4916854906c30501 Mon Sep 17 00:00:00 2001 From: Changpeng Fang <changpeng.f...@amd.com> Date: Sat, 18 May 2024 15:07:39 -0700 Subject: [PATCH 3/3] [OpenCL] Fix an infinite loop in builidng AddrSpaceQualType --- clang/lib/AST/ASTContext.cpp | 35 +++++++++---------- .../array-type-infinite-loop.clcpp | 10 +++--- 2 files changed, 22 insertions(+), 23 deletions(-) diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 388233c554d46..52eab5feb062b 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -3054,28 +3054,27 @@ QualType ASTContext::removeAddrSpaceQualType(QualType T) const { if (!T.hasAddressSpace()) return T; + QualifierCollector Quals; + const Type *TypeNode; // For arrays, strip the qualifier off the element type, then reconstruct the // array type if (T.getTypePtr()->isArrayType()) { - Qualifiers Qualfs; - return getUnqualifiedArrayType(T, Qualfs); - } - - // If we are composing extended qualifiers together, merge together - // into one ExtQuals node. - QualifierCollector Quals; - const Type *TypeNode; - - while (T.hasAddressSpace()) { - TypeNode = Quals.strip(T); - - // If the type no longer has an address space after stripping qualifiers, - // jump out. - if (!QualType(TypeNode, 0).hasAddressSpace()) - break; + T = getUnqualifiedArrayType(T, Quals); + TypeNode = T.getTypePtr(); + } else { + // If we are composing extended qualifiers together, merge together + // into one ExtQuals node. + while (T.hasAddressSpace()) { + TypeNode = Quals.strip(T); + + // If the type no longer has an address space after stripping qualifiers, + // jump out. + if (!QualType(TypeNode, 0).hasAddressSpace()) + break; - // There might be sugar in the way. Strip it and try again. - T = T.getSingleStepDesugaredType(*this); + // There might be sugar in the way. Strip it and try again. + T = T.getSingleStepDesugaredType(*this); + } } Quals.removeAddressSpace(); diff --git a/clang/test/CodeGenOpenCLCXX/array-type-infinite-loop.clcpp b/clang/test/CodeGenOpenCLCXX/array-type-infinite-loop.clcpp index 5a5b104e892f7..db9d7eb3281fc 100644 --- a/clang/test/CodeGenOpenCLCXX/array-type-infinite-loop.clcpp +++ b/clang/test/CodeGenOpenCLCXX/array-type-infinite-loop.clcpp @@ -1,12 +1,12 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 4 -//RUN: %clang_cc1 %s -emit-llvm -O1 -o - | FileCheck %s +//RUN: %clang_cc1 %s -triple spir -emit-llvm -O1 -o - | FileCheck %s // CHECK-LABEL: define dso_local spir_kernel void @test( -// CHECK-SAME: ptr nocapture noundef readonly align 8 [[IN:%.*]], ptr nocapture noundef writeonly align 8 [[OUT:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !kernel_arg_addr_space [[META3:![0-9]+]] !kernel_arg_access_qual [[META4:![0-9]+]] !kernel_arg_type [[META5:![0-9]+]] !kernel_arg_base_type [[META5]] !kernel_arg_type_qual [[META6:![0-9]+]] { +// CHECK-SAME: ptr addrspace(1) nocapture noundef readonly align 8 [[IN:%.*]], ptr addrspace(1) nocapture noundef writeonly align 8 [[OUT:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !kernel_arg_addr_space [[META3:![0-9]+]] !kernel_arg_access_qual [[META4:![0-9]+]] !kernel_arg_type [[META5:![0-9]+]] !kernel_arg_base_type [[META5]] !kernel_arg_type_qual [[META6:![0-9]+]] { // CHECK-NEXT: entry: -// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds i8, ptr [[IN]], i64 8 -// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[ARRAYIDX1]], align 8, !tbaa [[TBAA7:![0-9]+]] -// CHECK-NEXT: store i64 [[TMP0]], ptr [[OUT]], align 8, !tbaa [[TBAA7]] +// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[IN]], i32 8 +// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr addrspace(1) [[ARRAYIDX1]], align 8, !tbaa [[TBAA7:![0-9]+]] +// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(1) [[OUT]], align 8, !tbaa [[TBAA7]] // CHECK-NEXT: ret void // __kernel void test(__global long *In, __global long *Out) { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
