[OpenCL] Put constant initializer globals into constant addrspace #90048

svenvh · 2024-04-25T13:14:51Z

Place constant initializer globals into the constant address space. Clang generates such globals for e.g. larger array member initializers of classes and then emits copy operations from the global to the object(s). The globals are never written so they ought to be in the constant address space.

llvmbot · 2024-04-25T13:15:21Z

@llvm/pr-subscribers-clang

@llvm/pr-subscribers-clang-codegen

Author: Sven van Haastregt (svenvh)

Changes

Place constant initializer globals into the constant address space. Clang generates such globals for e.g. larger array member initializers of classes and then emits copy operations from the global to the object(s). The globals are never written so they ought to be in the constant address space.

Full diff: https://github.com/llvm/llvm-project/pull/90048.diff

2 Files Affected:

(modified) clang/lib/CodeGen/CGExprAgg.cpp (+2)
(modified) clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp (+4-1)

diff --git a/clang/lib/CodeGen/CGExprAgg.cpp b/clang/lib/CodeGen/CGExprAgg.cpp
index 355fec42be4489..30cde245cc837c 100644
--- a/clang/lib/CodeGen/CGExprAgg.cpp
+++ b/clang/lib/CodeGen/CGExprAgg.cpp
@@ -536,6 +536,8 @@ void AggExprEmitter::EmitArrayInit(Address DestPtr, llvm::ArrayType *AType,
     CodeGen::CodeGenModule &CGM = CGF.CGM;
     ConstantEmitter Emitter(CGF);
     LangAS AS = ArrayQTy.getAddressSpace();
+    if (CGF.getLangOpts().OpenCL)
+      AS = LangAS::opencl_constant;
     if (llvm::Constant *C =
             Emitter.tryEmitForInitializer(ExprToVisit, AS, ArrayQTy)) {
       auto GV = new llvm::GlobalVariable(
diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp
index 18d97a989a4364..a0ed03b25535c8 100644
--- a/clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp
@@ -5,7 +5,7 @@
 // for constructors, member functions and destructors.
 // See also atexit.cl and global_init.cl for other specific tests.
 
-// CHECK: %struct.MyType = type { i32 }
+// CHECK: %struct.MyType = type { i32, [5 x i32] }
 struct MyType {
   MyType(int i) : i(i) {}
   MyType(int i) __constant : i(i) {}
@@ -14,6 +14,7 @@ struct MyType {
   int bar() { return i + 2; }
   int bar() __constant { return i + 1; }
   int i;
+  int a[5] = {42, 43, 44, 45, 46};
 };
 
 // CHECK: @const1 ={{.*}} addrspace(2) global %struct.MyType zeroinitializer
@@ -23,6 +24,8 @@ __constant MyType const2(2);
 // CHECK: @glob ={{.*}} addrspace(1) global %struct.MyType zeroinitializer
 MyType glob(1);
 
+// CHECK: @constinit ={{.*}} addrspace(2) constant [5 x i32] [i32 42, i32 43, i32 44, i32 45, i32 46]
+
 // CHECK: call spir_func void @_ZNU3AS26MyTypeC1Ei(ptr addrspace(2) {{[^,]*}} @const1, i32 noundef 1)
 // CHECK: call spir_func void @_ZNU3AS26MyTypeC1Ei(ptr addrspace(2) {{[^,]*}} @const2, i32 noundef 2)
 // CHECK: call spir_func void @_ZNU3AS46MyTypeC1Ei(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(1) @glob to ptr addrspace(4)), i32 noundef 1)

clang/lib/CodeGen/CGExprAgg.cpp

… GVArrayQTy

clang/lib/CodeGen/CGExprAgg.cpp

efriedma-quic

LGTM

changpeng · 2024-05-16T22:04:48Z

clang/lib/CodeGen/CGExprAgg.cpp

@@ -535,20 +535,23 @@ void AggExprEmitter::EmitArrayInit(Address DestPtr, llvm::ArrayType *AType,
      elementType.isTriviallyCopyableType(CGF.getContext())) {
    CodeGen::CodeGenModule &CGM = CGF.CGM;
    ConstantEmitter Emitter(CGF);
-    LangAS AS = ArrayQTy.getAddressSpace();
+    QualType GVArrayQTy = CGM.getContext().getAddrSpaceQualType(
+        CGM.getContext().removeAddrSpaceQualType(ArrayQTy),


We saw a regression caused by this PR. It is a soft hang in CGM.getContext().removeAddrSpaceQualType.
Specifically it is in the following while loop:
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);

}
We found that "T == T.getSingleStepDesugaredType(*this);" and this it could not proceed.

I am not sure whether we should break out this loop when "T == T.getSingleStepDesugaredType(*this)"
or something else is wrong that we should never see such case.

Here is the dump of T:
ConstantArrayType 0x555565b40640 '__private ulong[16]' 16
-QualType 0x555565b403f8 '__private ulong' __private -ElaboratedType 0x555565b3ff40 'ulong' sugar imported
-TypedefType 0x555565b3fef0 'ulong' sugar imported |-Typedef 0x555565b3fe80 'ulong' -BuiltinType 0x55556583f430 'unsigned long'

I think it's a bug in removeAddrSpaceQualType(): it needs to special-case 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. So it needs to strip the qualifier off the element type, then reconstruct the array type. Maybe it can use ASTContext::getUnqualifiedArrayType.

@changpeng would you be able to provide an input source that demonstrates the issue?

I think it's a bug in removeAddrSpaceQualType(): it needs to special-case 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. So it needs to strip the qualifier off the element type, then reconstruct the array type. Maybe it can use ASTC
getSingleStepDesugaredType
Yes, the issue is in removeAddrSpaceQualType(ArrayQTy), And getSingleStepDesugaredType can not remove "Sugar".

@changpeng would you be able to provide an input source that demonstrates the issue?

Hi, @svenvh : I attached test.cl.txt here which is the dumped opencl source file. Unfortunately I do not know exactly how to reproduce the infinite loop offline with this source. I extracted out the following simplified kernel which can reproduce the hang with

clang -c -Xclang -emit-llvm -O0 test.clcpp

__kernel void nonceGrind(__global ulong *headerIn, __global ulong *nonceOut) {
ulong m[16] = { headerIn[0], headerIn[1],
headerIn[2], headerIn[3],
0, headerIn[5],
headerIn[6], headerIn[7],
headerIn[8], headerIn[9], 0, 0, 0, 0, 0, 0 };
*nonceOut = m[4];
}

However, I am afraid it may not fully represent the original issue. This is because after I break out the loop in
ASTContext::removeAddrSpaceQualType, I am seeing the following assert:

clang: /home/chfang/llvm-project/clang/include/clang/AST/Type.h:677: void clang::Qualifiers::addConsistentQualifiers(Qualifiers): Assertion `getAddressSpace() == qs.getAddressSpace() || !hasAddressSpace() || !qs.hasAddressSpace()' failed.

Hopefully the information is useful, and you are able to help. Thanks.

Reduced further:

clang -c -Xclang -emit-llvm -O0 test.clcpp

__kernel void test(__global ulong *In, __global ulong *Out) {
ulong m[4] = { In[0], In[1], 0, 0 };
*Out = m[1];
}

I think it's a bug in removeAddrSpaceQualType(): it needs to special-case 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. So it needs to strip the qualifier off the element type, then reconstruct the array type. Maybe it can use ASTContext::getUnqualifiedArrayType.

Thanks for the suggestion. I drafted a fix:
#92612

changpeng · 2024-05-17T18:25:10Z

test.cl.txt

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.

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.

svenvh added the OpenCL label Apr 25, 2024

svenvh requested a review from AnastasiaStulova April 25, 2024 13:14

llvmbot added clang Clang issues not falling into any other category clang:codegen IR generation bugs: mangling, exceptions, etc. labels Apr 25, 2024

svenvh mentioned this pull request Apr 25, 2024

InvalidInstruction: Can't translate llvm instruction with LLVM 18 (working with 17) KhronosGroup/SPIRV-LLVM-Translator#2531

Closed

davidrohr mentioned this pull request Apr 25, 2024

OpenCL (CLC++2021) compilation to SPIR-V fails #68305

Open

efriedma-quic reviewed Apr 25, 2024

View reviewed changes

clang/lib/CodeGen/CGExprAgg.cpp Outdated Show resolved Hide resolved

clang/lib/CodeGen/CGExprAgg.cpp Outdated Show resolved Hide resolved

clang/lib/CodeGen/CGExprAgg.cpp Show resolved Hide resolved

svenvh added 2 commits April 26, 2024 13:56

Address code review comment: set ArrayQTy addrspace

08936e5

Address review comment: use GetGlobalConstantAddressSpace and extract…

ebefe4b

… GVArrayQTy

efriedma-quic reviewed Apr 29, 2024

View reviewed changes

clang/lib/CodeGen/CGExprAgg.cpp Outdated Show resolved Hide resolved

Address review comment: override previous address space

99b7d42

efriedma-quic approved these changes Apr 30, 2024

View reviewed changes

svenvh merged commit b6328db into llvm:main May 2, 2024

svenvh deleted the constinit-constant-addrspace branch May 2, 2024 09:46

changpeng reviewed May 16, 2024

View reviewed changes

changpeng mentioned this pull request May 17, 2024

[OpenCL] Fix an infinite loop in builidng AddrSpaceQualType #92612

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[OpenCL] Put constant initializer globals into constant addrspace #90048

[OpenCL] Put constant initializer globals into constant addrspace #90048

Uh oh!

svenvh commented Apr 25, 2024

Uh oh!

llvmbot commented Apr 25, 2024 •

edited

Loading

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

efriedma-quic left a comment

Uh oh!

changpeng May 16, 2024

Uh oh!

efriedma-quic May 16, 2024

Uh oh!

svenvh May 17, 2024

Uh oh!

changpeng May 17, 2024

Uh oh!

changpeng May 17, 2024

Uh oh!

changpeng May 17, 2024

Uh oh!

changpeng May 17, 2024

Uh oh!

changpeng commented May 17, 2024

Uh oh!

Uh oh!

[OpenCL] Put constant initializer globals into constant addrspace #90048

[OpenCL] Put constant initializer globals into constant addrspace #90048

Uh oh!

Conversation

svenvh commented Apr 25, 2024

Uh oh!

llvmbot commented Apr 25, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

efriedma-quic left a comment

Choose a reason for hiding this comment

Uh oh!

changpeng May 16, 2024

Choose a reason for hiding this comment

Uh oh!

efriedma-quic May 16, 2024

Choose a reason for hiding this comment

Uh oh!

svenvh May 17, 2024

Choose a reason for hiding this comment

Uh oh!

changpeng May 17, 2024

Choose a reason for hiding this comment

Uh oh!

changpeng May 17, 2024

Choose a reason for hiding this comment

Uh oh!

changpeng May 17, 2024

Choose a reason for hiding this comment

Uh oh!

changpeng May 17, 2024

Choose a reason for hiding this comment

Uh oh!

changpeng commented May 17, 2024

Uh oh!

Uh oh!

llvmbot commented Apr 25, 2024 •

edited

Loading