[llvm-branch-commits] [cfe-branch] r369499 - Merging r369251:
Hans Wennborg via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Wed Aug 21 00:33:41 PDT 2019
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);
+}
More information about the llvm-branch-commits
mailing list