Author: stulova Date: Mon Aug 19 04:43:16 2019 New Revision: 369251 URL: http://llvm.org/viewvc/llvm-project?rev=369251&view=rev Log: [OpenCL] Fix addr space deduction for pointers/references to arrays.
Rewrite the logic for detecting if we are deducing addr space of a pointee type to take into account special logic for arrays. For pointers/references to arrays we can have any number of parentheses expressions as well as nested pointers. Differential Revision: https://reviews.llvm.org/D66137 Modified: cfe/trunk/lib/Sema/SemaType.cpp cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl Modified: cfe/trunk/lib/Sema/SemaType.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=369251&r1=369250&r2=369251&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaType.cpp (original) +++ cfe/trunk/lib/Sema/SemaType.cpp Mon Aug 19 04:43:16 2019 @@ -7385,8 +7385,22 @@ static void deduceOpenCLImplicitAddrSpac bool IsPointee = ChunkIndex > 0 && (D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Pointer || - D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::BlockPointer || - D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Reference); + D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Reference || + D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::BlockPointer); + // For pointers/references to arrays the next chunk is always an array + // followed by any number of parentheses. + if (!IsPointee && ChunkIndex > 1) { + auto AdjustedCI = ChunkIndex - 1; + if (D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Array) + AdjustedCI--; + // Skip over all parentheses. + while (AdjustedCI > 0 && + D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Paren) + AdjustedCI--; + if (D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Pointer || + D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Reference) + IsPointee = true; + } bool IsFuncReturnType = ChunkIndex > 0 && D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Function; Modified: cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl?rev=369251&r1=369250&r2=369251&view=diff ============================================================================== --- cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl (original) +++ cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl Mon Aug 19 04:43:16 2019 @@ -78,3 +78,25 @@ __kernel void test() { int foo[10]; xxx(&foo[0]); } + +// Addr space for pointer/reference to an array +//CHECK: FunctionDecl {{.*}} t1 'void (const __generic float (&)[2])' +void t1(const float (&fYZ)[2]); +//CHECK: FunctionDecl {{.*}} t2 'void (const __generic float (*)[2])' +void t2(const float (*fYZ)[2]); +//CHECK: FunctionDecl {{.*}} t3 'void (__generic float (((*)))[2])' +void t3(float(((*fYZ)))[2]); +//CHECK: FunctionDecl {{.*}} t4 'void (__generic float (((*__generic *)))[2])' +void t4(float(((**fYZ)))[2]); +//CHECK: FunctionDecl {{.*}} t5 'void (__generic float (*__generic (*))[2])' +void t5(float (*(*fYZ))[2]); + +__kernel void k() { + __local float x[2]; + __local float(*p)[2]; + t1(x); + t2(&x); + t3(&x); + t4(&p); + t5(&p); +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
