r369251 - [OpenCL] Fix addr space deduction for pointers/references to arrays.

[Anastasia Stulova via cfe-commits]

Hi Hans,

Is it still possible to port this fix to the release branch?

Thanks,
Anastasia

________________________________
From: cfe-commits <cfe-commits-bounces at lists.llvm.org> on behalf of Anastasia Stulova via cfe-commits <cfe-commits at lists.llvm.org>
Sent: 19 August 2019 12:43
To: cfe-commits at lists.llvm.org <cfe-commits at lists.llvm.org>
Subject: r369251 - [OpenCL] Fix addr space deduction for pointers/references to arrays.

Author: stulova
Date: Mon Aug 19 04:43:16 2019
New Revision: 369251

URL: http://llvm.org/viewvc/llvm-project?rev=369251&view=rev
Log:
[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/trunk/lib/Sema/SemaType.cpp
    cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl

Modified: cfe/trunk/lib/Sema/SemaType.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=369251&r1=369250&r2=369251&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaType.cpp (original)
+++ cfe/trunk/lib/Sema/SemaType.cpp Mon Aug 19 04:43:16 2019
@@ -7385,8 +7385,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/trunk/test/SemaOpenCLCXX/address-space-deduction.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl?rev=369251&r1=369250&r2=369251&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl (original)
+++ cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl Mon Aug 19 04:43:16 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);
+}


_______________________________________________
cfe-commits mailing list
cfe-commits at lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20190820/667cd59f/attachment-0001.html>