Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 4 commits into from
May 2, 2024
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
c5e7b2d
[OpenCL] Put constant initializer globals into constant addrspace
svenvh Apr 25, 2024
08936e5
Address code review comment: set ArrayQTy addrspace
svenvh Apr 26, 2024
ebefe4b
Address review comment: use GetGlobalConstantAddressSpace and extract…
svenvh Apr 29, 2024
99b7d42
Address review comment: override previous address space
svenvh Apr 30, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
11 changes: 7 additions & 4 deletions clang/lib/CodeGen/CGExprAgg.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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),
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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'

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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".

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@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.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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];
}

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

CGM.GetGlobalConstantAddressSpace());
LangAS AS = GVArrayQTy.getAddressSpace();
if (llvm::Constant *C =
Emitter.tryEmitForInitializer(ExprToVisit, AS, ArrayQTy)) {
Emitter.tryEmitForInitializer(ExprToVisit, AS, GVArrayQTy)) {
auto GV = new llvm::GlobalVariable(
CGM.getModule(), C->getType(),
/* isConstant= */ true, llvm::GlobalValue::PrivateLinkage, C,
"constinit",
/* InsertBefore= */ nullptr, llvm::GlobalVariable::NotThreadLocal,
CGM.getContext().getTargetAddressSpace(AS));
Emitter.finalize(GV);
CharUnits Align = CGM.getContext().getTypeAlignInChars(ArrayQTy);
CharUnits Align = CGM.getContext().getTypeAlignInChars(GVArrayQTy);
GV->setAlignment(Align.getAsAlign());
Address GVAddr(GV, GV->getValueType(), Align);
EmitFinalDestCopy(ArrayQTy, CGF.MakeAddrLValue(GVAddr, ArrayQTy));
EmitFinalDestCopy(ArrayQTy, CGF.MakeAddrLValue(GVAddr, GVArrayQTy));
return;
}
}
Expand Down
5 changes: 4 additions & 1 deletion clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) {}
Expand All @@ -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
Expand All @@ -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)
Expand Down