<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=iso-8859-1">
<style type="text/css" style="display:none;"><!-- P {margin-top:0;margin-bottom:0;} --></style>
</head>
<body dir="ltr">
<div id="divtagdefaultwrapper" style="font-size:12pt;color:#000000;font-family:Calibri,Helvetica,sans-serif;" dir="ltr">
<div id="divtagdefaultwrapper" style="font-size: 12pt; color: rgb(0, 0, 0); font-family: Calibri, Helvetica, sans-serif, "EmojiFont", "Apple Color Emoji", "Segoe UI Emoji", NotoColorEmoji, "Segoe UI Symbol", "Android Emoji", EmojiSymbols;" dir="ltr">
<p style="margin-top:0;margin-bottom:0"><br>
</p>
Hi Hans,</div>
<div style="font-size: 12pt; color: rgb(0, 0, 0); font-family: Calibri, Helvetica, sans-serif, "EmojiFont", "Apple Color Emoji", "Segoe UI Emoji", NotoColorEmoji, "Segoe UI Symbol", "Android Emoji", EmojiSymbols;" dir="ltr">
<br>
</div>
<div style="font-size: 12pt; color: rgb(0, 0, 0); font-family: Calibri, Helvetica, sans-serif, "EmojiFont", "Apple Color Emoji", "Segoe UI Emoji", NotoColorEmoji, "Segoe UI Symbol", "Android Emoji", EmojiSymbols;" dir="ltr">
Is it still possible to port this fix to the release branch?</div>
<div style="font-size: 12pt; color: rgb(0, 0, 0); font-family: Calibri, Helvetica, sans-serif, "EmojiFont", "Apple Color Emoji", "Segoe UI Emoji", NotoColorEmoji, "Segoe UI Symbol", "Android Emoji", EmojiSymbols;" dir="ltr">
<br>
</div>
<div style="font-size: 12pt; color: rgb(0, 0, 0); font-family: Calibri, Helvetica, sans-serif, "EmojiFont", "Apple Color Emoji", "Segoe UI Emoji", NotoColorEmoji, "Segoe UI Symbol", "Android Emoji", EmojiSymbols;" dir="ltr">
Thanks,</div>
<div style="font-size: 12pt; color: rgb(0, 0, 0); font-family: Calibri, Helvetica, sans-serif, "EmojiFont", "Apple Color Emoji", "Segoe UI Emoji", NotoColorEmoji, "Segoe UI Symbol", "Android Emoji", EmojiSymbols;" dir="ltr">
Anastasia<br>
<br>
<div style="color: rgb(0, 0, 0);">
<hr style="display:inline-block;width:98%" tabindex="-1">
<div id="divRplyFwdMsg" dir="ltr"><font style="font-size:11pt" face="Calibri, sans-serif" color="#000000"><b>From:</b> cfe-commits <cfe-commits-bounces@lists.llvm.org> on behalf of Anastasia Stulova via cfe-commits <cfe-commits@lists.llvm.org><br>
<b>Sent:</b> 19 August 2019 12:43<br>
<b>To:</b> cfe-commits@lists.llvm.org <cfe-commits@lists.llvm.org><br>
<b>Subject:</b> r369251 - [OpenCL] Fix addr space deduction for pointers/references to arrays.</font>
<div> </div>
</div>
<div class="BodyFragment"><font size="2"><span style="font-size:11pt;">
<div class="PlainText">Author: stulova<br>
Date: Mon Aug 19 04:43:16 2019<br>
New Revision: 369251<br>
<br>
URL: <a href="http://llvm.org/viewvc/llvm-project?rev=369251&view=rev">http://llvm.org/viewvc/llvm-project?rev=369251&view=rev</a><br>
Log:<br>
[OpenCL] Fix addr space deduction for pointers/references to arrays.<br>
<br>
Rewrite the logic for detecting if we are deducing addr space of<br>
a pointee type to take into account special logic for arrays. For<br>
pointers/references to arrays we can have any number of parentheses<br>
expressions as well as nested pointers.<br>
<br>
Differential Revision: <a href="https://reviews.llvm.org/D66137">https://reviews.llvm.org/D66137</a><br>
<br>
<br>
Modified:<br>
cfe/trunk/lib/Sema/SemaType.cpp<br>
cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl<br>
<br>
Modified: cfe/trunk/lib/Sema/SemaType.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=369251&r1=369250&r2=369251&view=diff">
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=369251&r1=369250&r2=369251&view=diff</a><br>
==============================================================================<br>
--- cfe/trunk/lib/Sema/SemaType.cpp (original)<br>
+++ cfe/trunk/lib/Sema/SemaType.cpp Mon Aug 19 04:43:16 2019<br>
@@ -7385,8 +7385,22 @@ static void deduceOpenCLImplicitAddrSpac<br>
bool IsPointee =<br>
ChunkIndex > 0 &&<br>
(D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Pointer ||<br>
- D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::BlockPointer ||<br>
- D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Reference);<br>
+ D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Reference ||<br>
+ D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::BlockPointer);<br>
+ // For pointers/references to arrays the next chunk is always an array<br>
+ // followed by any number of parentheses.<br>
+ if (!IsPointee && ChunkIndex > 1) {<br>
+ auto AdjustedCI = ChunkIndex - 1;<br>
+ if (D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Array)<br>
+ AdjustedCI--;<br>
+ // Skip over all parentheses.<br>
+ while (AdjustedCI > 0 &&<br>
+ D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Paren)<br>
+ AdjustedCI--;<br>
+ if (D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Pointer ||<br>
+ D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Reference)<br>
+ IsPointee = true;<br>
+ }<br>
bool IsFuncReturnType =<br>
ChunkIndex > 0 &&<br>
D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Function;<br>
<br>
Modified: cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl?rev=369251&r1=369250&r2=369251&view=diff">
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl?rev=369251&r1=369250&r2=369251&view=diff</a><br>
==============================================================================<br>
--- cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl (original)<br>
+++ cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl Mon Aug 19 04:43:16 2019<br>
@@ -78,3 +78,25 @@ __kernel void test() {<br>
int foo[10];<br>
xxx(&foo[0]);<br>
}<br>
+<br>
+// Addr space for pointer/reference to an array<br>
+//CHECK: FunctionDecl {{.*}} t1 'void (const __generic float (&)[2])'<br>
+void t1(const float (&fYZ)[2]);<br>
+//CHECK: FunctionDecl {{.*}} t2 'void (const __generic float (*)[2])'<br>
+void t2(const float (*fYZ)[2]);<br>
+//CHECK: FunctionDecl {{.*}} t3 'void (__generic float (((*)))[2])'<br>
+void t3(float(((*fYZ)))[2]);<br>
+//CHECK: FunctionDecl {{.*}} t4 'void (__generic float (((*__generic *)))[2])'<br>
+void t4(float(((**fYZ)))[2]);<br>
+//CHECK: FunctionDecl {{.*}} t5 'void (__generic float (*__generic (*))[2])'<br>
+void t5(float (*(*fYZ))[2]);<br>
+<br>
+__kernel void k() {<br>
+ __local float x[2];<br>
+ __local float(*p)[2];<br>
+ t1(x);<br>
+ t2(&x);<br>
+ t3(&x);<br>
+ t4(&p);<br>
+ t5(&p);<br>
+}<br>
<br>
<br>
_______________________________________________<br>
cfe-commits mailing list<br>
cfe-commits@lists.llvm.org<br>
<a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits">https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits</a><br>
</div>
</span></font></div>
</div>
</div>
</div>
</body>
</html>