[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;



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;

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits