[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

llvm-svn: 369251
This commit is contained in:
Anastasia Stulova 2019-08-19 11:43:16 +00:00
parent b38bac3699
commit eb801abd58
2 changed files with 38 additions and 2 deletions

View File

@ -7385,8 +7385,22 @@ static void deduceOpenCLImplicitAddrSpace(TypeProcessingState &State,
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;

View File

@ -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);
}