[OpenCL] Fix an infinite loop in builidng AddrSpaceQualType #92612
Conversation
In building AddrSpaceQualType (llvm#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.
llvmbot
added
clang
|
@llvm/pr-subscribers-clang-codegen @llvm/pr-subscribers-clang Author: Changpeng Fang (changpeng) ChangesIn building AddrSpaceQualType (#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 instead, which strips the qualifier off the element type, then reconstruct the array type. Full diff: https://github.com/llvm/llvm-project/pull/92612.diff 2 Files Affected:
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"}
+//.
|
clang/lib/CodeGen/CGExprAgg.cpp
Outdated
| QualType GVArrayQTy = CGM.getContext().getAddrSpaceQualType( | ||
| CGM.getContext().removeAddrSpaceQualType(ArrayQTy), | ||
| CGM.getContext().getUnqualifiedArrayType(ArrayQTy, Quals), |
There was a problem hiding this comment.
I see two issues with this fix:
- It leaves removeAddrSpaceQualType broken for anyone else who tries to use it. Ideally, removeAddrSpaceQualType should do what it says.
- This drops other qualifiers on the floor; they might be relevant in some cases.
There was a problem hiding this comment.
Do you mean we should actually fix removeAddrSpaceQualType? Somewhere inside removeAddrSpaceQualType, we
should use getUnqualifiedArrayType if it is an arrayType, and getSingleStepDesugaredType othereise?
I have to admit that I have no experience in this field, so I am relying on you and @svenvh to move on for a reasonable fix. Thanks.
There was a problem hiding this comment.
Something like that, yes.
There was a problem hiding this comment.
Something like that, yes.
Thanks. Updated. Should be closer!
Fix ASTContext::removeAddrSpaceQualType()
llvmbot
added
the
clang:frontend
| @@ -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 | |||
There was a problem hiding this comment.
codegen tests need an explicit target
There was a problem hiding this comment.
add "triple spir", same as other tests in the same directory
clang/lib/AST/ASTContext.cpp
Outdated
| // array type | ||
| if (T.getTypePtr()->isArrayType()) { | ||
| Qualifiers Qualfs; | ||
| return getUnqualifiedArrayType(T, Qualfs); |
There was a problem hiding this comment.
You need to preserve the non-address-space qualifiers, the same way the code after the loop does. Actually, I'd just recommend restructuring to use the existing code, something like:
if (T.getTypePtr()->isArrayType()) {
getUnqualifiedArrayType()
} else {
while (T.hasAddressSpace()) {
[...]
}
}
There was a problem hiding this comment.
Thanks. Can I do as the following?
Note that I am passing QualifierCollector to getUnqualifiedArrayType, which has Qualifiers as the second argument.
Also, TypeNode = T.getTypePtr(); after I is unqualified.
QualifierCollector Quals;
const Type *TypeNode;
if (T.getTypePtr()->isArrayType()) {
T = getUnqualifiedArrayType(T, Quals);
TypeNode = T.getTypePtr();
} else {
while (T.hasAddressSpace()) {
....
In building AddrSpaceQualType (#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 instead, which strips the qualifier off the element type, then reconstruct the array type.