<html xmlns:v="urn:schemas-microsoft-com:vml" xmlns:o="urn:schemas-microsoft-com:office:office" xmlns:w="urn:schemas-microsoft-com:office:word" xmlns:m="http://schemas.microsoft.com/office/2004/12/omml" xmlns="http://www.w3.org/TR/REC-html40">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8">
<meta name="Generator" content="Microsoft Word 12 (filtered medium)">
<style><!--
/* Font Definitions */
@font-face
{font-family:SimSun;
panose-1:2 1 6 0 3 1 1 1 1 1;}
@font-face
{font-family:"Cambria Math";
panose-1:2 4 5 3 5 4 6 3 2 4;}
@font-face
{font-family:Calibri;
panose-1:2 15 5 2 2 2 4 3 2 4;}
@font-face
{font-family:Tahoma;
panose-1:2 11 6 4 3 5 4 4 2 4;}
@font-face
{font-family:"\@SimSun";
panose-1:2 1 6 0 3 1 1 1 1 1;}
/* Style Definitions */
p.MsoNormal, li.MsoNormal, div.MsoNormal
{margin:0in;
margin-bottom:.0001pt;
font-size:12.0pt;
font-family:"Times New Roman","serif";}
a:link, span.MsoHyperlink
{mso-style-priority:99;
color:blue;
text-decoration:underline;}
a:visited, span.MsoHyperlinkFollowed
{mso-style-priority:99;
color:purple;
text-decoration:underline;}
span.EmailStyle17
{mso-style-type:personal-reply;
font-family:"Calibri","sans-serif";
color:#1F497D;}
.MsoChpDefault
{mso-style-type:export-only;}
@page WordSection1
{size:8.5in 11.0in;
margin:1.0in 1.0in 1.0in 1.0in;}
div.WordSection1
{page:WordSection1;}
--></style><!--[if gte mso 9]><xml>
<o:shapedefaults v:ext="edit" spidmax="1026" />
</xml><![endif]--><!--[if gte mso 9]><xml>
<o:shapelayout v:ext="edit">
<o:idmap v:ext="edit" data="1" />
</o:shapelayout></xml><![endif]-->
</head>
<body lang="EN-US" link="blue" vlink="purple">
<div class="WordSection1">
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D">Yonghong, In our implementation (not open sourced), we don’t do outlining the Front-End. See my previous reply to Medhi’s email.
<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D">Xinmin<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D"><o:p> </o:p></span></p>
<div style="border:none;border-top:solid #B5C4DF 1.0pt;padding:3.0pt 0in 0in 0in">
<p class="MsoNormal"><b><span style="font-size:10.0pt;font-family:"Tahoma","sans-serif"">From:</span></b><span style="font-size:10.0pt;font-family:"Tahoma","sans-serif""> Yonghong Yan [mailto:yan@oakland.edu]
<br>
<b>Sent:</b> Friday, January 20, 2017 11:18 AM<br>
<b>To:</b> Mehdi Amini<br>
<b>Cc:</b> Tian, Xinmin; llvm-dev@lists.llvm.org; llvm-dev-request@lists.llvm.org; Adve, Vikram Sadanand<br>
<b>Subject:</b> Re: [llvm-dev] [RFC] IR-level Region Annotations<o:p></o:p></span></p>
</div>
<p class="MsoNormal"><o:p> </o:p></p>
<div>
<p class="MsoNormal">Xinmin,<o:p></o:p></p>
<div>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
<div>
<p class="MsoNormal">outlining turns a parallel program into a sequential one from compiler's perspective, and that is why most of the parallel-ignorant pass would hurt. In your IR description for Sanjoy's example, does the current approach of outlining impacting
the way of the IR should be enhanced for parallelism?<o:p></o:p></p>
</div>
<div>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
<div>
<p class="MsoNormal">For that specific example (or other analysis and optimization SPMD) and what is implemented in clang, I am not sure whether we are going to change the frontend so not to outline the parallel region, or allow to perform certain optimization
such as hoisting that alloca in clang which is not desired I believe. Or annotate the outlined function together with the intrinsic_a so that hoisting can be performed, in which case the instrisic_a would like this:<o:p></o:p></p>
</div>
<div>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
<div>
<p class="MsoNormal">tok = llvm.experimental.intrinsic_a()[ "DIR.PARALLEL"(), "QUAL.PRIVATE"(i32* %val, i32 %i), "QUAL.NUM_THREADS"(i32 4) plus info for OUTLINED_call.<o:p></o:p></p>
</div>
<div>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
<div>
<p class="MsoNormal">Mehdi, <o:p></o:p></p>
</div>
<div>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
<div>
<p class="MsoNormal">I think i am asking the same question as you asked me. <o:p></o:p></p>
</div>
<div>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
<div>
<p class="MsoNormal">Yonghong<o:p></o:p></p>
</div>
<div>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
<div>
<div>
<div>
<div>
<div>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
</div>
</div>
</div>
<p class="MsoNormal"><o:p> </o:p></p>
<div>
<p class="MsoNormal">On Fri, Jan 20, 2017 at 1:54 PM, Mehdi Amini via llvm-dev <<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a>> wrote:<o:p></o:p></p>
<p class="MsoNormal"><br>
> On Jan 20, 2017, at 10:44 AM, Tian, Xinmin via llvm-dev <<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a>> wrote:<br>
><br>
> Sanjoy, the IR would be like something below. It is ok to hoist alloca instruction outside the region. There are some small changes in optimizer to understand region-annotation intrinsic.<br>
><br>
> { void main() {<br>
> i32* val = alloca i32<br>
> tok = llvm.experimental.intrinsic_a()[ "DIR.PARALLEL"(), "QUAL.PRIVATE"(i32* val), "QUAL.NUM_THREADS"(i32 4)]<br>
><br>
> int i = omp_get_thread_num();<br>
> compute_something_into_val(val, i);<br>
> a[i] = val;<br>
><br>
> llvm.experimental.intrinsic_b(tok)["DIR.END.PARALLEL"()];<br>
> }<br>
><br>
> With above representation, we can do privatization and outlining as below<br>
><br>
> { void main() {<br>
> i32* val = alloca i32<br>
> i32* I = alloca 32<br>
> tok = llvm.experimental.intrinsic_a()[ "DIR.PARALLEL"(), "QUAL.PRIVATE"(i32* %val, i32 %i), "QUAL.NUM_THREADS"(i32 4)]<br>
><br>
> %ii = omp_get_thread_num();<br>
> compute_something_into_val(%val, %i);<br>
> a[i] = %val;<br>
><br>
> llvm.experimental.intrinsic_b(tok)["DIR.END.PARALLEL"()];<br>
> }<br>
><br>
<br>
Here we come to the interesting part: the hoisting of "i32* I = alloca 32” above the intrinsics required to update the intrinsics information “QUAL.PRIVATE”.<br>
This means that the optimizer has to be aware of it, I’m missing the magic here?<br>
I understand that an openmp specific optimization can do it, the question is how it an openmp agnostic supposed to behave in face of llvm.experimental.intrinsic_a?<br>
<br>
—<br>
Mehdi<o:p></o:p></p>
<div>
<div>
<p class="MsoNormal"><br>
<br>
<br>
> 1. create i32* priv_val = alloca i32 %priv_i = ...in the region, and replace all %val with %prv_val in the region.<br>
> 2. perform outlining.<br>
><br>
> Caller code<br>
> ....<br>
> omp_push_num_threads(4)<br>
> omp_fork_call( .... outline_par_region....) ....<br>
><br>
> Callee code:<br>
> Outlined_par_rgion {<br>
> I32* priv_val = alloca 32<br>
> I32* priv_i = ....<br>
><br>
> Ret<br>
> }<br>
><br>
> For OpenMP, we do support it at -O0, -O1, -O2 and -O3. We had to make sure it runs correctly w/ and w/o optimizations and advanced analysis. So we need to preserve all source information for BE.<br>
> You can take a look our LLVM-HPC paper for a bit some details. There are still tons of work to be done. Thanks.<br>
><br>
> Xinmin<br>
><br>
> -----Original Message-----<br>
> From: llvm-dev [mailto:<a href="mailto:llvm-dev-bounces@lists.llvm.org" target="_blank">llvm-dev-bounces@lists.llvm.org</a>] On Behalf Of Sanjoy Das via llvm-dev<br>
> Sent: Thursday, January 19, 2017 10:13 PM<br>
> To: Adve, Vikram Sadanand <<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>><br>
> Cc: llvm-dev <<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a>>;
<a href="mailto:llvm-dev-request@lists.llvm.org" target="_blank">llvm-dev-request@lists.llvm.org</a><br>
> Subject: Re: [llvm-dev] [RFC] IR-level Region Annotations<br>
><br>
> Hi Vikram,<br>
><br>
> On Thu, Jan 19, 2017 at 9:27 PM, Adve, Vikram Sadanand <<a href="mailto:vadve@illinois.edu" target="_blank">vadve@illinois.edu</a>> wrote:<br>
>> Hi Sanjoy,<br>
>><br>
>> Yes, that's exactly what we have been looking at recently here, but<br>
>> the region tags seem to make it possible to express the control flow<br>
>> as well, so I think we could start with reg ions+metadata, as Hal and<br>
><br>
> I'm not yet convinced that region tags are sufficient to model exotic control flow.<br>
><br>
> (I don't know OpenMP so this is a copy-pasted-edited example)<br>
><br>
> Say we have:<br>
><br>
> void main() {<br>
> #pragma omp parallel num_threads(4)<br>
> {<br>
> int i = omp_get_thread_num();<br>
> int val;<br>
> compute_something_into_val(&val, i);<br>
> a[i] = val;<br>
> }<br>
> }<br>
><br>
> I presume the (eventual) intended lowering is something like this (if the intended lowering is different than this, and avoids the issue I'm trying to highlight then my point is moot):<br>
><br>
> void main() {<br>
> tok = llvm.experimental.intrinsic_a();<br>
><br>
> int i = omp_get_thread_num();<br>
> i32* val = alloca i32<br>
> compute_something_into_val(val, i);<br>
> a[i] = val;<br>
><br>
> llvm.experimental.intrinsic_b(tok);<br>
> }<br>
><br>
> However, LLVM is free to hoist the alloca to the entry block:<br>
><br>
> void main() {<br>
> i32* val = alloca i32<br>
> tok = llvm.experimental.intrinsic_a();<br>
><br>
> int i = omp_get_thread_num();<br>
> compute_something_into_val(val, i);<br>
> a[i] = val;<br>
><br>
> llvm.experimental.intrinsic_b(tok);<br>
> }<br>
><br>
> and now you have a race between the four parallel forks.<br>
><br>
> The problem here is that nothing in the IR expresses that we have four copies of the region running "at the same time". In fact, such a control flow is alien to LLVM today.<br>
><br>
> For instance, another evil optimization may turn:<br>
><br>
> void main() {<br>
> int a[4];<br>
> #pragma omp parallel num_threads(4)<br>
> {<br>
> int i = omp_get_thread_num();<br>
> int val = compute_something_into_val(i);<br>
> a[i] = val;<br>
> }<br>
><br>
> return a[0] + a[1];<br>
> }<br>
><br>
> to<br>
><br>
> void main() {<br>
> int a[4];<br>
> #pragma omp parallel num_threads(4)<br>
> {<br>
> int i = omp_get_thread_num();<br>
> int val = compute_something_into_val(i);<br>
> a[i] = val;<br>
> }<br>
><br>
> return undef;<br>
> }<br>
><br>
> since a[i] = val could have initialized at most one element in a.<br>
><br>
> Now you could say that the llvm.experimental.intrinsic_a and llvm.experimental.intrinsic_b intrinsics are magic, and even such "obvious" optimizations are not allowed to happen across them; but then calls to these intrinsics is pretty fundamentally different
from "normal" calls, and are probably best modeled as new instructions.<br>
> You're going to have to do the same kind of auditing of passes either way, and the only extra cost of a new instruction is the extra bitcode reading / writing code.<br>
><br>
> I hope I made sense.<br>
><br>
> -- Sanjoy<br>
><br>
>> Xinmin proposed, and then figure out what needs to be first class<br>
>> instructions.<br>
><br>
>><br>
>> --Vikram Adve<br>
>><br>
>><br>
>><br>
>>> On Jan 19, 2017, at 11:03 PM, Sanjoy Das <<a href="mailto:sanjoy@playingwithpointers.com" target="_blank">sanjoy@playingwithpointers.com</a>> wrote:<br>
>>><br>
>>> Hi,<br>
>>><br>
>>> My bias is to use both (b) and (d), since they have complementary<br>
>>> strengths. We should use (b) for expressing concepts that can't be<br>
>>> semantically modeled as a call or invoke (this branch takes both its<br>
>>> successors), and (d) for expressing things that can be (this call may<br>
>>> never return), and annotation like things (this region (denoted by<br>
>>> def-use of a token) is a reduction).<br>
>>><br>
>>> I don't grok OpenMP, but perhaps we can come with one or two<br>
>>> "generalized control flow"-type instructions that can be used to<br>
>>> model the non-call/invoke like semantics we'd like LLVM to know<br>
>>> about, and model the rest with (d)?<br>
>>><br>
>>> -- Sanjoy<br>
>>><br>
>>> On Thu, Jan 19, 2017 at 8:28 PM, Hal Finkel via llvm-dev<br>
>>> <<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a>> wrote:<br>
>>>><br>
>>>> On 01/19/2017 03:36 PM, Mehdi Amini via llvm-dev wrote:<br>
>>>><br>
>>>><br>
>>>> On Jan 19, 2017, at 1:32 PM, Daniel Berlin <<a href="mailto:dberlin@dberlin.org" target="_blank">dberlin@dberlin.org</a>> wrote:<br>
>>>><br>
>>>><br>
>>>><br>
>>>>> On Thu, Jan 19, 2017 at 1:12 PM, Mehdi Amini <<a href="mailto:mehdi.amini@apple.com" target="_blank">mehdi.amini@apple.com</a>> wrote:<br>
>>>>><br>
>>>>><br>
>>>>> On Jan 19, 2017, at 12:04 PM, Daniel Berlin <<a href="mailto:dberlin@dberlin.org" target="_blank">dberlin@dberlin.org</a>> wrote:<br>
>>>>><br>
>>>>><br>
>>>>><br>
>>>>> On Thu, Jan 19, 2017 at 11:46 AM, Mehdi Amini via llvm-dev<br>
>>>>> <<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a>> wrote:<br>
>>>>>><br>
>>>>>><br>
>>>>>>> On Jan 19, 2017, at 11:36 AM, Adve, Vikram Sadanand via llvm-dev<br>
>>>>>>> <<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a>> wrote:<br>
>>>>>>><br>
>>>>>>> Hi Johannes,<br>
>>>>>>><br>
>>>>>>>> I am especially curious where you get your data from. Tapir [0]<br>
>>>>>>>> (and to some degree PIR [1]) have shown that,<br>
>>>>>>>> counterintuitively, only a few changes to LLVM passes are<br>
>>>>>>>> needed. Tapir was recently used in an MIT class with a lot of<br>
>>>>>>>> students and it seemed to work well with only minimal changes to<br>
>>>>>>>> analysis and especially transformation passes.<br>
>>>>>>><br>
>>>>>>> TAPIR is an elegant, small extension and, in particular, I think<br>
>>>>>>> the idea of asymmetric parallel tasks and control flow is a<br>
>>>>>>> clever way to express parallelism with serial semantics, as in<br>
>>>>>>> Cilk. Encoding the control flow extensions as explicit<br>
>>>>>>> instructions is orthogonal to that, though arguably more elegant than using region tags + metadata.<br>
>>>>>>><br>
>>>>>>> However, Cilk is a tiny language compared with the full<br>
>>>>>>> complexity of other languages, like OpenMP. To take just one<br>
>>>>>>> example, TAPIR cannot express the ORDERED construct of OpenMP. A<br>
>>>>>>> more serious concern, IMO, is that TAPIR (like Cilk) requires<br>
>>>>>>> serial semantics, whereas there are many parallel languages, OpenMP included, that do not obey that restriction.<br>
>>>>>>> Third, OpenMP has *numerous* clauses, e.g., REDUCTION or PRIVATE,<br>
>>>>>>> that are needed because without that, you’d be dependent on<br>
>>>>>>> fundamentally hard compiler analyses to extract the same<br>
>>>>>>> information for satisfactory parallel performance; realistic<br>
>>>>>>> applications cannot depend on the success of such analyses.<br>
>>>>>><br>
>>>>>> I agree with this, but I’m also wondering if it needs to be first<br>
>>>>>> class in the IR?<br>
>>>>>> For example we know our alias analysis is very basic, and C/C++<br>
>>>>>> have a higher constraint thanks to their type system, but we<br>
>>>>>> didn’t inject this higher level information that helps the<br>
>>>>>> optimizer as first class IR constructs.<br>
>>>>><br>
>>>>><br>
>>>>> FWIW, while i agree with the general point, i wouldn't use this example.<br>
>>>>> Because we pretty much still suffer to this day because of it (both<br>
>>>>> in AA, and devirt, and ...) :) We can't always even tell fields<br>
>>>>> apart<br>
>>>>><br>
>>>>><br>
>>>>> Is it inherent to the infrastructure, i.e. using metadata instead<br>
>>>>> of first class IR construct or is it just a “quality of implementation” issue?<br>
>>>><br>
>>>> Not to derail this conversation:<br>
>>>><br>
>>>> IMHO, At some point there is no real difference :)<br>
>>>><br>
>>>> Because otherwise, everything is a QOI issue.<br>
>>>><br>
>>>> IE if it's super tricky to get metadata that works well and works<br>
>>>> right, doesn't get lost, etc, and that's inherent to using metadata,<br>
>>>> that to me is not a QOI issue.<br>
>>>><br>
>>>> So could it be done with metadata? Probably?<br>
>>>> But at the same time, if it had been done with more first class<br>
>>>> constructs, it would have happened years ago and been much lower cost.<br>
>>>><br>
>>>><br>
>>>> This is what I meant by “inherent to the infrastructure”, thanks for<br>
>>>> clarifying.<br>
>>>><br>
>>>><br>
>>>> To clarify, we were proposing metadata that is used as arguments to<br>
>>>> the region-annotation intrinsics. This metadata has the nice<br>
>>>> property that it does not get dropped (so it is just being used as a<br>
>>>> way of encoding whatever data structures are necessary without predefining a syntactic schema).<br>
>>>><br>
>>>> -Hal<br>
>>>><br>
>>>><br>
>>>> —<br>
>>>> Mehdi<br>
>>>><br>
>>>><br>
>>>><br>
>>>><br>
>>>> _______________________________________________<br>
>>>> LLVM Developers mailing list<br>
>>>> <a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a><br>
>>>> <a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" target="_blank">
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br>
>>>><br>
>>>><br>
>>>> --<br>
>>>> Hal Finkel<br>
>>>> Lead, Compiler Technology and Programming Languages Leadership<br>
>>>> Computing Facility Argonne National Laboratory<br>
>>>><br>
>>>><br>
>>>> _______________________________________________<br>
>>>> LLVM Developers mailing list<br>
>>>> <a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a><br>
>>>> <a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" target="_blank">
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br>
>>>><br>
> _______________________________________________<br>
> LLVM Developers mailing list<br>
> <a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a><br>
> <a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" target="_blank">
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br>
> _______________________________________________<br>
> LLVM Developers mailing list<br>
> <a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a><br>
> <a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" target="_blank">
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br>
<br>
_______________________________________________<br>
LLVM Developers mailing list<br>
<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a><br>
<a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" target="_blank">http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><o:p></o:p></p>
</div>
</div>
</div>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
</div>
</div>
</body>
</html>