[clang] [OpenCL] Put constant initializer globals into constant addrspace (PR #90048)

Changpeng Fang via cfe-commits cfe-commits at lists.llvm.org
Fri May 17 11:25:07 PDT 2024


================
@@ -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),
----------------
changpeng wrote:

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

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


More information about the cfe-commits mailing list