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

[Anastasia Stulova via Phabricator via cfe-commits]

Anastasia created this revision.
Anastasia added reviewers: svenvh, rjmccall.
Herald added subscribers: ebevhan, yaxunl.
Anastasia retitled this revision from "[OpenCL][PR42385]Improve addr space deduction for pointers/references to arrays" to "[OpenCL][PR42385] Improve 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 pointer/references to arrays we can have any number of parenthesis expressions as well as nested pointers.


https://reviews.llvm.org/D66137

Files:
  lib/Sema/SemaType.cpp
  test/SemaOpenCLCXX/address-space-deduction.cl


Index: test/SemaOpenCLCXX/address-space-deduction.cl
===================================================================
--- test/SemaOpenCLCXX/address-space-deduction.cl
+++ test/SemaOpenCLCXX/address-space-deduction.cl
@@ -78,3 +78,26 @@
   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);
+}
Index: lib/Sema/SemaType.cpp
===================================================================
--- lib/Sema/SemaType.cpp
+++ lib/Sema/SemaType.cpp
@@ -7390,8 +7390,22 @@
   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 parenthesis.
+  if (!IsPointee && ChunkIndex > 1) {
+    auto AdjustedCI = ChunkIndex-1;
+    if (D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Array)
+      AdjustedCI--;
+    // Skip over all parenthesis.
+    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;


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D66137.214806.patch
Type: text/x-patch
Size: 2343 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20190813/9533c670/attachment.bin>