[OpenCL][PR42385] Improve addr space deduction for pointers/references to arrays

[Anastasia]

Summary

Rewrite the logic for detecting if we are deducing addr space of a pointee type to take into account special logic for arrays.
For pointer/references to arrays we can have any number of parenthesis expressions as well as nested pointers.

[svenvh]

LGTM, but please clang-format your patch before committing (esp. spacing around binary operators).

lib/Sema/SemaType.cpp

7396 ↗

(On Diff #214806)

parentheses (also below).

cfe/trunk/lib/Sema/SemaType.cpp

static void deduceOpenCLImplicitAddrSpace(TypeProcessingState &State,
   // space, except when it is used as a pointee type.
   //
   // Since LLVM assumes function type is in default address space, it should not
   // have address space.
   auto ChunkIndex = State.getCurrentChunkIndex();
   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;
   bool IsFuncType =
       ChunkIndex < D.getNumTypeObjects() &&
       D.getTypeObject(ChunkIndex).Kind == DeclaratorChunk::Function;
   if ( // Do not deduce addr space for function return type and function type,
        // otherwise it will fail some sema check.

cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl

T xxx(T *in) {
   T ii;
   return *i;
 }
 
 __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);
+}