[clang] [OpenCL] Fix an infinite loop in builidng AddrSpaceQualType (PR #92612)

[Changpeng Fang via cfe-commits]

https://github.com/changpeng created https://github.com/llvm/llvm-project/pull/92612

 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 instead, which strips the qualifier off the element type, then reconstruct the array type.

>From 2468a85a47499d90a99610846c632332eb7307b8 Mon Sep 17 00:00:00 2001
From: Changpeng Fang <changpeng.fang at amd.com>
Date: Fri, 17 May 2024 15:13:07 -0700
Subject: [PATCH] [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"}
+//.