<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=iso-8859-1">
<meta name="Generator" content="Microsoft Word 15 (filtered medium)">
<!--[if !mso]><style>v\:* {behavior:url(#default#VML);}
o\:* {behavior:url(#default#VML);}
w\:* {behavior:url(#default#VML);}
.shape {behavior:url(#default#VML);}
</style><![endif]--><style><!--
/* Font Definitions */
@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:Verdana;
panose-1:2 11 6 4 3 5 4 4 2 4;}
@font-face
{font-family:"Lucida Console";
panose-1:2 11 6 9 4 5 4 2 2 4;}
/* Style Definitions */
p.MsoNormal, li.MsoNormal, div.MsoNormal
{margin:0cm;
margin-bottom:.0001pt;
font-size:11.0pt;
font-family:"Calibri",sans-serif;}
a:link, span.MsoHyperlink
{mso-style-priority:99;
color:blue;
text-decoration:underline;}
a:visited, span.MsoHyperlinkFollowed
{mso-style-priority:99;
color:#954F72;
text-decoration:underline;}
p
{mso-style-priority:99;
mso-margin-top-alt:auto;
margin-right:0cm;
mso-margin-bottom-alt:auto;
margin-left:0cm;
font-size:12.0pt;
font-family:"Times New Roman",serif;}
tt
{mso-style-priority:99;
font-family:"Courier New";}
span.EmailStyle19
{mso-style-type:personal-reply;
font-family:"Verdana",sans-serif;
color:#1F497D;}
.MsoChpDefault
{mso-style-type:export-only;
font-size:10.0pt;}
@page WordSection1
{size:612.0pt 792.0pt;
margin:72.0pt 72.0pt 72.0pt 72.0pt;}
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-GB" link="blue" vlink="#954F72">
<div class="WordSection1">
<p class="MsoNormal">The reason is that the user would expect the same outcome/performance from both,<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Verdana",sans-serif;color:#1F497D">The user should not expect the same performance.<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Verdana",sans-serif;color:#1F497D">Consider, if you write<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D">#pragma omp parallel for<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D">for (…)<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D"> …<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Verdana",sans-serif;color:#1F497D">and expand it as described by the standard, you have (where the implicit barrier at the end of omp for is made visible)<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D">#pragma omp parallel<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D">{<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D"># pragma omp for nowait<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D"> for (…)<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D"> …<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D"># pragma omp barrier<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D">}<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Verdana",sans-serif;color:#1F497D">That barrier is unnecessary since it immediately precedes a join which is itself a barrier (the mast thread cannot leave until all threads arrive).<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Verdana",sans-serif;color:#1F497D">Therefore, the code that is notmally generated elides the unnecessary barrier and expands into<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D">#pragma omp parallel<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D">{<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D"># pragma omp for nowait<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D"> for (…)<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D"> …<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Lucida Console";color:#1F497D">}<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Verdana",sans-serif;color:#1F497D"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Verdana",sans-serif;color:#1F497D;mso-fareast-language:EN-US">Thus it can easily be the case that omp parallel do/for is faster than omp parallel + omp do/for.<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Verdana",sans-serif;color:#1F497D;mso-fareast-language:EN-US"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Verdana",sans-serif;color:#1F497D;mso-fareast-language:EN-US">(Of course you can fix this if you expand to parallel; do/for
<b>nowait</b>…)<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Verdana",sans-serif;color:#1F497D;mso-fareast-language:EN-US"><o:p> </o:p></span></p>
<div>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Verdana",sans-serif;color:#1F497D">-- Jim<br>
<br>
Jim Cownie <james.h.cownie@intel.com><br>
SSG/DPD/TCAR (Technical Computing, Analyzers, and Runtimes)<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Verdana",sans-serif;color:#1F497D">Tel: +44 117 9071438</span><span style="color:#1F497D"><o:p></o:p></span></p>
</div>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Verdana",sans-serif;color:#1F497D;mso-fareast-language:EN-US"><o:p> </o:p></span></p>
<div>
<div style="border:none;border-top:solid #E1E1E1 1.0pt;padding:3.0pt 0cm 0cm 0cm">
<p class="MsoNormal"><b><span lang="EN-US">From:</span></b><span lang="EN-US"> Openmp-dev [mailto:openmp-dev-bounces@lists.llvm.org]
<b>On Behalf Of </b>Schürmann, Daniel via Openmp-dev<br>
<b>Sent:</b> Friday, June 2, 2017 9:58 PM<br>
<b>To:</b> Arpith C Jacob <acjacob@us.ibm.com><br>
<b>Cc:</b> openmp-dev@lists.llvm.org<br>
<b>Subject:</b> Re: [Openmp-dev] Proposal: Resolve combined directives in parsing phase<o:p></o:p></span></p>
</div>
</div>
<p class="MsoNormal"><o:p> </o:p></p>
<div>
<p class="MsoNormal">Thank you for your feedback, Arpith.<o:p></o:p></p>
<p class="MsoNormal">I see that the nvptx codegen relies in parts on combined directives and understand the possible difficulties.<o:p></o:p></p>
<p class="MsoNormal">However, codegen is already able to do this with the target and teams directive. Remember that the split directives are immediate child nodes in the AST.<o:p></o:p></p>
<p class="MsoNormal">But you are right, this change can only be made after changes to codegen.<o:p></o:p></p>
<p class="MsoNormal">In my opinion the code generation should in general be able to handle both cases equivally, combined directives and the separated version.
<o:p></o:p></p>
<p class="MsoNormal">The reason is that the user would expect the same outcome/performance from both, e.g.<o:p></o:p></p>
<p class="MsoNormal">-> #pragma omp target parallel for<o:p></o:p></p>
<p class="MsoNormal">as well as<o:p></o:p></p>
<p class="MsoNormal">-> #pragma omp target<o:p></o:p></p>
<p class="MsoNormal"> #pragma omp parallel<o:p></o:p></p>
<p class="MsoNormal"> #pragma omp for<o:p></o:p></p>
<p class="MsoNormal"> <o:p></o:p></p>
<p class="MsoNormal">like the spec states. (The only exception seem to be the distribute- combined directives.<o:p></o:p></p>
<p class="MsoNormal"> <o:p></o:p></p>
<p class="MsoNormal">Kind regards,<o:p></o:p></p>
<p class="MsoNormal">Daniel<o:p></o:p></p>
<p class="MsoNormal"> <o:p></o:p></p>
<div style="border:none;border-top:solid #E1E1E1 1.0pt;padding:3.0pt 0cm 0cm 0cm">
<p class="MsoNormal"><b>Von: </b><a href="mailto:acjacob@us.ibm.com">Arpith C Jacob</a><br>
<b>Gesendet: </b>Freitag, 2. Juni 2017 20:24<br>
<b>An: </b><a href="mailto:daniel.schuermann@campus.tu-berlin.de">Schürmann, Daniel</a><br>
<b>Cc: </b><a href="mailto:openmp-dev@lists.llvm.org">openmp-dev@lists.llvm.org</a><br>
<b>Betreff: </b>Re: [Openmp-dev] Proposal: Resolve combined directives in parsing phase<o:p></o:p></p>
</div>
<p class="MsoNormal"> <o:p></o:p></p>
</div>
<div>
<p><span style="font-size:10.0pt">Daniel,</span><br>
<br>
<span style="font-size:10.0pt">> </span><tt><span style="font-size:10.0pt">Cons: - Code optimizations for combined directives may be harder to implement</span></tt><br>
<br>
<span style="font-size:10.0pt">From the perspective of GPU code generation treating combined directives as a special case is *immensely* important for performance. Knowing that there is *no* serial section in the target region makes it possible to generate
simplified and low overhead code that resembles typical CUDA kernels. I am worried that splitting combined directives will make it much harder to get this performance back, particularly because we must do OpenMP lowering in Clang.</span><br>
<br>
<span style="font-size:10.0pt">> </span><tt><span style="font-size:10.0pt">the code generation is implemented as inlined function what results in
</span></tt><span style="font-size:10.0pt;font-family:"Courier New""><br>
<tt>> ignoring the semantic meaning of these directives.</tt><br>
</span><br>
<span style="font-size:10.0pt">Can the current code generation be altered to fix the issues that you see?</span><br>
<br>
<tt><span style="font-size:10.0pt">> Unrelated question:</span></tt><span style="font-size:10.0pt;font-family:"Courier New""><br>
<tt>> I don't understand the necessity of the __kmpc_fork_teams() run-time </tt><br>
<tt>> call as the __tgt_target_teams() implementation should be able to handle </tt>
<br>
<tt>> this case.</tt><br>
</span><br>
<span style="font-size:10.0pt">My understanding is that the __tgt* calls are implemented in the target offload library while the __kmpc* calls are in the OpenMP runtime. On CPUs, forking of teams is done in the __kmpc_fork_teams() call. On the GPU, the offload
call __tgt_target_teams() launches a kernel with multiple teams so the __kmpc_fork_teams() is a no-op.</span><br>
<br>
<span style="font-size:10.0pt">Regards,</span><br>
<span style="font-size:10.0pt">Arpith</span><br>
<br>
<img border="0" width="16" height="16" id="_x0000_i1025" src="cid:image001.gif@01D2DDDD.AA05E670" alt="Inactive hide details for Daniel Schürmann via Openmp-dev ---06/02/2017 09:08:11 AM---At the moment, combined directives have t"><span style="font-size:10.0pt;color:#424282">Daniel
Schürmann via Openmp-dev ---06/02/2017 09:08:11 AM---At the moment, combined directives have their own ast representation for type-checking and code gen</span><br>
<br>
<span style="font-size:10.0pt;color:#5F5F5F">From: </span><span style="font-size:10.0pt">Daniel Schürmann via Openmp-dev <<a href="mailto:openmp-dev@lists.llvm.org">openmp-dev@lists.llvm.org</a>></span><br>
<span style="font-size:10.0pt;color:#5F5F5F">To: </span><span style="font-size:10.0pt"><<a href="mailto:openmp-dev@lists.llvm.org">openmp-dev@lists.llvm.org</a>></span><br>
<span style="font-size:10.0pt;color:#5F5F5F">Date: </span><span style="font-size:10.0pt">06/02/2017 09:08 AM</span><br>
<span style="font-size:10.0pt;color:#5F5F5F">Subject: </span><span style="font-size:10.0pt">[Openmp-dev] Proposal: Resolve combined directives in parsing phase</span><br>
<span style="font-size:10.0pt;color:#5F5F5F">Sent by: </span><span style="font-size:10.0pt">"Openmp-dev" <<a href="mailto:openmp-dev-bounces@lists.llvm.org">openmp-dev-bounces@lists.llvm.org</a>></span><o:p></o:p></p>
<div class="MsoNormal"><span style="font-size:12.0pt;font-family:"Times New Roman",serif">
<hr size="2" width="100%" noshade="" style="color:#8091A5" align="left">
</span></div>
<p class="MsoNormal" style="margin-bottom:12.0pt"><span style="font-size:12.0pt;font-family:"Times New Roman",serif"><br>
<br>
<br>
</span><tt><span style="font-size:10.0pt">At the moment, combined directives have their own ast representation for
</span></tt><span style="font-size:10.0pt;font-family:"Courier New""><br>
<tt>type-checking and code generation. For some of the combined constructs, </tt>
<br>
<tt>the code generation is implemented as inlined function what results in </tt><br>
<tt>ignoring the semantic meaning of these directives.</tt><br>
<br>
<tt>This is true for e.g.</tt><br>
<tt>EmitOMPTargetParallelForSimdDirective</tt><br>
<tt>EmitOMPTargetSimdDirective</tt><br>
<tt>EmitOMPTeamsDistributeDirective</tt><br>
<tt>EmitOMPTargetTeamsDistributeDirective</tt><br>
<tt>EmitOMPTargetTeamsDistributeParallelForDirective</tt><br>
<tt>and more</tt><br>
<br>
<tt>One solution would be the proper codegen implementation for these </tt><br>
<tt>directives.</tt><br>
<tt>However, I would like to propose a simpler and closer-to-spec approach:</tt><br>
<tt>By resolving combined directives in the parsing phase into nested AST nodes.</tt><br>
<br>
<tt>E.g. an OMPTargetTeamsDistributeDirective would be resolved into</tt><br>
<tt>OMPTargetDirective</tt><br>
<tt> |- OMPTeamsDirective</tt><br>
<tt> |- OMPDistributeDirective</tt><br>
<br>
<tt>whereas type-checking and codegen for these single directives is already </tt>
<br>
<tt>implemented.</tt><br>
<tt>The advantages are:</tt><br>
<tt>- Much simpler type-checking and code generation</tt><br>
<tt>- We match the specification stating that combined directives have the </tt><br>
<tt>semantic meaning of one construct immediately followed by the other </tt><br>
<tt>construct</tt><br>
<tt>- All combined directives are fully supported if their derived </tt><br>
<tt>constructs are supported</tt><br>
<br>
<tt>Potential disadvantages:</tt><br>
<tt>- The AST representation differs from the input. However, this is </tt><br>
<tt>already the case due to inserted implicit parameters.</tt><br>
<tt>- Code optimizations for combined directives may be harder to implement</tt><br>
<br>
<tt>In my opinion the benefits outweigh the disadvantages, but I may not be </tt>
<br>
<tt>aware of some implications. Please let me know your thoughts about this </tt>
<br>
<tt>idea. And tell me if I missunderstood anything related that led to the </tt><br>
<tt>decision for the actual design.</tt><br>
<br>
<tt>Unrelated question:</tt><br>
<tt>I don't understand the necessity of the __kmpc_fork_teams() run-time </tt><br>
<tt>call as the __tgt_target_teams() implementation should be able to handle </tt>
<br>
<tt>this case.</tt><br>
<br>
<br>
<tt>Daniel</tt><br>
<tt>_______________________________________________</tt><br>
<tt>Openmp-dev mailing list</tt><br>
<tt><a href="mailto:Openmp-dev@lists.llvm.org">Openmp-dev@lists.llvm.org</a></tt><br>
<tt><a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev">http://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev</a></tt><br>
</span><span style="font-size:12.0pt;font-family:"Times New Roman",serif"><br>
<br>
<o:p></o:p></span></p>
</div>
</div>
<p>---------------------------------------------------------------------<br>
Intel Corporation (UK) Limited<br>
Registered No. 1134945 (England)<br>
Registered Office: Pipers Way, Swindon SN3 1RJ<br>
VAT No: 860 2173 47</p>
<p>This e-mail and any attachments may contain confidential material for<br>
the sole use of the intended recipient(s). Any review or distribution<br>
by others is strictly prohibited. If you are not the intended<br>
recipient, please contact the sender and delete all copies.</p></body>
</html>