[llvm-branch-commits] [cfe-branch] r369499 - Merging r369251:

[Hans Wennborg via llvm-branch-commits]

Author: hans
Date: Wed Aug 21 00:33:40 2019
New Revision: 369499

URL: http://llvm.org/viewvc/llvm-project?rev=369499&view=rev
Log:
Merging r369251:
------------------------------------------------------------------------
r369251 | stulova | 2019-08-19 13:43:16 +0200 (Mon, 19 Aug 2019) | 10 lines

[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/branches/release_90/   (props changed)
    cfe/branches/release_90/lib/Sema/SemaType.cpp
    cfe/branches/release_90/test/SemaOpenCLCXX/address-space-deduction.cl

Propchange: cfe/branches/release_90/
------------------------------------------------------------------------------
--- svn:mergeinfo (original)
+++ svn:mergeinfo Wed Aug 21 00:33:40 2019
@@ -1,4 +1,4 @@
 /cfe/branches/type-system-rewrite:134693-134817
-/cfe/trunk:366429,366448,366457,366474,366480,366483,366511,366670,366694,366699,366878,367008,367039,367055,367103,367134,367301,367305,367323,367387,367403,367520,367530,367661,367675,367802,367823,367906,368104,368202,368552,368561,368874,368940,369043
+/cfe/trunk:366429,366448,366457,366474,366480,366483,366511,366670,366694,366699,366878,367008,367039,367055,367103,367134,367301,367305,367323,367387,367403,367520,367530,367661,367675,367802,367823,367906,368104,368202,368552,368561,368874,368940,369043,369251
 /cfe/trunk/test:170344
 /cfe/trunk/test/SemaTemplate:126920

Modified: cfe/branches/release_90/lib/Sema/SemaType.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/branches/release_90/lib/Sema/SemaType.cpp?rev=369499&r1=369498&r2=369499&view=diff
==============================================================================
--- cfe/branches/release_90/lib/Sema/SemaType.cpp (original)
+++ cfe/branches/release_90/lib/Sema/SemaType.cpp Wed Aug 21 00:33:40 2019
@@ -7390,8 +7390,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/branches/release_90/test/SemaOpenCLCXX/address-space-deduction.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/branches/release_90/test/SemaOpenCLCXX/address-space-deduction.cl?rev=369499&r1=369498&r2=369499&view=diff
==============================================================================
--- cfe/branches/release_90/test/SemaOpenCLCXX/address-space-deduction.cl (original)
+++ cfe/branches/release_90/test/SemaOpenCLCXX/address-space-deduction.cl Wed Aug 21 00:33:40 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);
+}