<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 style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important">Hello Vinay,<br>
</span></div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important"><br>
Thanks for your mail about the OpenMP dialect in MLIR. Happy to know that you and several other groups are interested in the OpenMP dialect. At the outset, I must point out that the design is not set in stone and will change as we make progress.<span style="font-family: Arial; background-color: rgb(255, 255, 255); display: inline !important"><span> </span>You
are welcome to participate, provide feedback and criticism to change the design as well as to contribute to the implementation. I provide some clarifications and replies to your comments below. If it is OK we can have further discussions in discourse as River
points out.</span></span></div>
<blockquote style="border-left: 3px solid rgb(200, 200, 200); border-top-color: rgb(200, 200, 200); border-right-color: rgb(200, 200, 200); border-bottom-color: rgb(200, 200, 200); padding-left: 1ex; margin-left: 0.8ex; color: rgb(102, 102, 102);">
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important">1. [May 2019] An OpenMPIRBuilder in LLVM was proposed for flang and clang frontends. Note that this proposal was before considering MLIR
for FIR.</span></div>
</blockquote>
<div style="color: rgb(0, 0, 0);"><font face="Calibri, Arial, Helvetica, sans-serif"><span style="font-size: 11pt; font-family: Arial, Helvetica, sans-serif;">A correction here. The proposal for OpenMPIRBuilder was made when MLIR was being considered for FIR. <br>
</span></font><font face="Calibri, Arial, Helvetica, sans-serif"><span style="font-size: 11pt; font-family: Arial, Helvetica, sans-serif;">(i) Gary Klimowicz's minutes for Flang call in April 2019 mentions considering MLIR for FIR.</span></font><br>
<a href="http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-April/000194.html" id="LPlnk278665" style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt;"><span style="font-size: 11pt; font-family: Arial, Helvetica, sans-serif;">http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-April/000194.html</span></a><br>
<font face="Calibri, Arial, Helvetica, sans-serif"><span style="font-size: 11pt; font-family: Arial, Helvetica, sans-serif;">(ii) My reply to Johaness's proposal in May 2019 mentions MLIR for FIR.</span></font><br>
<a href="http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-May/000220.html" id="LPlnk452670" style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt;"><span style="font-size: 11pt; font-family: Arial, Helvetica, sans-serif;">http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-May/000220.html</span></a><br>
<br>
<blockquote style="border-left: 3px solid rgb(200, 200, 200); border-top-color: rgb(200, 200, 200); border-right-color: rgb(200, 200, 200); border-bottom-color: rgb(200, 200, 200); padding-left: 1ex; margin-left: 0.8ex; color: rgb(102, 102, 102);">
<p dir="ltr" style="margin-top: 0pt; margin-bottom: 0pt; color: rgb(50, 49, 48); background-color: rgb(255, 255, 255); line-height: 1.38;">
<span style="font-family: Arial; margin: 0px; font-size: 11pt; color: rgb(0, 0, 0);">b. Review of barrier construct is in progress:
</span><span style="margin: 0px; text-decoration-line: underline;"><a href="https://reviews.llvm.org/D72962" style="font-family: Arial; font-size: 11pt; margin: 0px; text-decoration-line: none;">https://reviews.llvm.org/D72962</a></span></p>
</blockquote>
<p dir="ltr" style="margin-top: 0pt; margin-bottom: 0pt; color: rgb(50, 49, 48); background-color: rgb(255, 255, 255); line-height: 1.38;">
</p>
<div dir="ltr" style="margin-top: 0px; margin-bottom: 0px; background-color: rgb(255, 255, 255); line-height: 1.38;">
<span style="margin: 0px; font-family: Arial, Helvetica, sans-serif; font-size: 11pt; line-height: normal; color: rgb(0, 0, 0);">Minor correction here. The addition of barrier construct was accepted and has landed (</span><a href="https://reviews.llvm.org/D72400" style="margin: 0px; font-family: "Segoe UI", "Segoe UI Web (West European)", "Segoe UI", -apple-system, system-ui, Roboto, "Helvetica Neue", sans-serif"><span style="font-family: Arial, Helvetica, sans-serif; font-size: 11pt; line-height: normal; color: rgb(0, 0, 0);">https://reviews.llvm.org/D7240</span></a><span style="margin: 0px; font-family: Arial, Helvetica, sans-serif; font-size: 11pt; line-height: normal; color: rgb(0, 0, 0);">)</span><span style="margin: 0px; font-family: Arial, Helvetica, sans-serif; font-size: 11pt; line-height: normal; color: rgb(0, 0, 0);">.
It is the review for translation to LLVM IR that is in progress.</span><br style="font-family: "Segoe UI", "Segoe UI Web (West European)", "Segoe UI", -apple-system, system-ui, Roboto, "Helvetica Neue", sans-serif">
</div>
<p></p>
<blockquote style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; border-left: 3px solid rgb(200, 200, 200); border-top-color: rgb(200, 200, 200); border-right-color: rgb(200, 200, 200); border-bottom-color: rgb(200, 200, 200); padding-left: 1ex; margin-left: 0.8ex; color: rgb(102, 102, 102);">
<span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important">It looks like the design has evolved over time and there is no one place which contains the latest design decisions that fits all the different
pieces of the puzzle. I will try to deduce it from the above mentioned references. Please correct me If I am referring to anything which has changed.</span></blockquote>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt;">
<span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important">Yes, the design has mildly changed over time to incorporate feedback. But the latest is what is there in the RFC in discourse.</span></div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt;">
<span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important"><br>
</span></div>
<blockquote style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; border-left: 3px solid rgb(200, 200, 200); border-top-color: rgb(200, 200, 200); border-right-color: rgb(200, 200, 200); border-bottom-color: rgb(200, 200, 200); padding-left: 1ex; margin-left: 0.8ex; color: rgb(102, 102, 102);">
<div><span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important"><span style="font-family: Arial; background-color: rgb(255, 255, 255); display: inline !important">For most OpenMP design discussions,
FIR examples are used (as seen in (2) and (3)). The MLIR examples mentioned in the design only talks about FIR dialect and LLVM dialect.</span><br>
</span></div>
</blockquote>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt;">
<span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important">Our initial concern was how will all these pieces (FIR, LLVM Dialect, OpenMPIRBuilder, LLVM IR) fit together. Hence you see the prominence
of FIR and LLVM dialect and more information about lowering/translation than transformations/optimisations.</span></div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt;">
<span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important"><br>
</span></div>
<blockquote style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; border-left: 3px solid rgb(200, 200, 200); border-top-color: rgb(200, 200, 200); border-right-color: rgb(200, 200, 200); border-bottom-color: rgb(200, 200, 200); padding-left: 1ex; margin-left: 0.8ex; color: rgb(102, 102, 102);">
<div><span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important"><span style="font-family: Arial; background-color: rgb(255, 255, 255); display: inline !important">This completely ignores the likes
of standard, affine (where most loop transformations are supposed to happen) and loop dialects.</span><br>
</span></div>
</blockquote>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt;">
<span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important">Adding to the reply above. We would like to take advantage of the transformations in cases that are possible. FIR loops will be converted
to affine/loop dialect. So the loop inside an omp.do can be in these dialects as clarified in the discussion in discourse and also shown in slide 20 of the fosdem presentation (links to both below).<br>
<a href="https://llvm.discourse.group/t/rfc-openmp-dialect-in-mlir/397/7?u=kiranchandramohan" id="LPNoLP808880">https://llvm.discourse.group/t/rfc-openmp-dialect-in-mlir/397/7?u=kiranchandramohan</a><br>
<a href="https://fosdem.org/2020/schedule/event/llvm_flang/attachments/slides/3839/export/events/attachments/llvm_flang/slides/3839/flang_llvm_frontend.pdf" id="LPNoLP581920">https://fosdem.org/2020/schedule/event/llvm_flang/attachments/slides/3839/export/events/attachments/llvm_flang/slides/3839/flang_llvm_frontend.pdf</a><br>
<br>
I must also point out that the question of where to do loop transformations is a topic we have not fully converged on. See the following thread for discussions.<br>
<a href="http://lists.llvm.org/pipermail/flang-dev/2019-September/000042.html" id="LPlnk702758">http://lists.llvm.org/pipermail/flang-dev/2019-September/000042.html</a><br>
</span></div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt;">
<span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important"><br>
</span></div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt;">
<span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important">
<blockquote style="border-left: 3px solid rgb(200, 200, 200); border-top-color: rgb(200, 200, 200); border-right-color: rgb(200, 200, 200); border-bottom-color: rgb(200, 200, 200); padding-left: 1ex; margin-left: 0.8ex; color: rgb(102, 102, 102);">
<span style="font-family: Arial; background-color: rgb(255, 255, 255); display: inline !important">Is it the same omp.do operation which now contains the bounds and induction variables of the loop after the LLVM conversion?</span><br>
</blockquote>
The point here is that i) we need to keep the loops separately so as to take advantage of transformations that other dialects like affine/loop would provide. ii) We will need the loop information while lowering the OpenMP do operation. For implementation, if
reusing the same operation (in different contexts) is difficult then we can add a new operation.</span></div>
<blockquote style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; border-left: 3px solid rgb(200, 200, 200); border-top-color: rgb(200, 200, 200); border-right-color: rgb(200, 200, 200); border-bottom-color: rgb(200, 200, 200); padding-left: 1ex; margin-left: 0.8ex; color: rgb(102, 102, 102);">
<span style="margin: 0px; font-size: 11pt; font-family: Arial">It is also not mentioned how clauses like
</span><span style="margin: 0px; font-weight: 700; font-size: 11pt; font-family: Arial">firstprivate, shared, private, reduce, map, etc
</span><span style="margin: 0px; font-size: 11pt; font-family: Arial">are lowered to OpenMP dialect.</span></blockquote>
<font face="Calibri, Arial, Helvetica, sans-serif"><span style="font-size: 11pt; font-family: Arial, Helvetica, sans-serif;">Yes, it is not mentioned. We did a study of a few constructs and clauses which was shared as mails to flang-dev and the RFC. As we make
progress and before implementation, we will share further details.</span></font><br>
<br>
<blockquote style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; border-left: 3px solid rgb(200, 200, 200); border-top-color: rgb(200, 200, 200); border-right-color: rgb(200, 200, 200); border-bottom-color: rgb(200, 200, 200); padding-left: 1ex; margin-left: 0.8ex; color: rgb(102, 102, 102);">
<span style="margin: 0px; font-size: 11pt; font-family: Arial">it would be beneficial to have an omp.</span><span style="margin: 0px; font-weight: 700; font-size: 11pt; font-family: Arial">parallel_do</span><span style="margin: 0px; font-size: 11pt; font-family: Arial">
operation which has semantics similar to other loop structures (may not be LoopLikeInterface) in MLIR. </span></blockquote>
<div style=""><font face="Arial"><span style="font-size: 14.6667px;">I am not against adding parallel_do if it can help with transformations or reduce the complexity of lowering. Please share the details in discourse as a reply to the RFC or a separate thread.<br>
<blockquote style="border-left: 3px solid rgb(200, 200, 200); border-top-color: rgb(200, 200, 200); border-right-color: rgb(200, 200, 200); border-bottom-color: rgb(200, 200, 200); padding-left: 1ex; margin-left: 0.8ex; color: rgb(102, 102, 102);">
<span style="font-family: Arial; background-color: rgb(255, 255, 255); display: inline !important">it looks like having OpenMP operations based on standard MLIR types and operations (scalars and memrefs mainly) is the right way to go.</span></blockquote>
<div><span style="font-family: Arial; background-color: rgb(255, 255, 255); display: inline !important">This will definitely be the first version that we implement. But I do not understand why we should restrict to only the standard types and operations. To
ease lowering and translation and to avoid adding OpenMP operations to other dialects, I believe OpenMP dialect should also be able to exist with other dialects like FIR and LLVM.</span></div>
</span></font></div>
<blockquote style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; border-left: 3px solid rgb(200, 200, 200); border-top-color: rgb(200, 200, 200); border-right-color: rgb(200, 200, 200); border-bottom-color: rgb(200, 200, 200); padding-left: 1ex; margin-left: 0.8ex; color: rgb(102, 102, 102);">
<div><span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important"><span style="margin: 0px; font-weight: 700; font-size: 11pt; font-family: Arial">E.
</span><span style="margin: 0px; font-size: 11pt; font-family: Arial">Lowering of target constructs mentioned in ( 2(d) ) specifies direct lowering to LLVM IR ignoring all the advantages that MLIR provides.
</span><br>
</span></div>
</blockquote>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt;">
<span style="font-family: Arial; font-size: 14.6667px; background-color: rgb(255, 255, 255); display: inline !important"><span style="margin: 0px; font-size: 11pt; font-family: Arial">
<blockquote style="border-left: 3px solid rgb(200, 200, 200); border-top-color: rgb(200, 200, 200); border-right-color: rgb(200, 200, 200); border-bottom-color: rgb(200, 200, 200); padding-left: 1ex; margin-left: 0.8ex; color: rgb(102, 102, 102);">
<span style="font-family: Arial; background-color: rgb(255, 255, 255); display: inline !important">Also, OpenMP codegen will automatically benefit from the GPU dialect based optimizations. For example, it would be way easier to hoist a memory reference out
of GPU kernel in MLIR than in LLVM IR.</span><br>
</blockquote>
I might not have fully understood you here. But the dialect lives independently of the translation to LLVM IR. If there are optimisations (like hoisting that you mention here) I believe they can be performed as transformation passes on the dialect. It is not
ruled out.<br>
<br>
--Kiran</span></span></div>
</div>
<hr tabindex="-1" style="display:inline-block; width:98%">
<div id="divRplyFwdMsg" dir="ltr"><font face="Calibri, sans-serif" color="#000000" style="font-size:11pt"><b>From:</b> flang-dev <flang-dev-bounces@lists.llvm.org> on behalf of Vinay Madhusudan via flang-dev <flang-dev@lists.llvm.org><br>
<b>Sent:</b> 13 February 2020 16:33<br>
<b>To:</b> llvm-dev@lists.llvm.org <llvm-dev@lists.llvm.org>; flang-dev@lists.llvm.org <flang-dev@lists.llvm.org><br>
<b>Subject:</b> [flang-dev] About OpenMP dialect in MLIR</font>
<div> </div>
</div>
<div>
<div dir="ltr"><span id="x_gmail-docs-internal-guid-a31980cc-7fff-b136-d820-5164bcfb4113">
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Hi,</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">I
have few questions / concerns regarding the design of OpenMP dialect in MLIR that is currently being implemented, mainly for the f18 compiler. Below, I summarize the current state of various efforts in clang / f18 / MLIR / LLVM regarding this. Feel free to
add to the list in case I have missed something.</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">1.
[May 2019] An OpenMPIRBuilder in LLVM was proposed for flang and clang frontends. Note that this proposal was before considering MLIR for FIR.</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">a.
llvm-dev proposal : </span><a href="http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-May/000197.html" style="text-decoration-line:none"><span style="font-size:11pt; font-family:Arial; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; text-decoration-line:underline; vertical-align:baseline; white-space:pre-wrap">http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-May/000197.html</span></a></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">b.
Patches in review: </span><a href="https://reviews.llvm.org/D70290" style="text-decoration-line:none"><span style="font-size:11pt; font-family:Arial; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; text-decoration-line:underline; vertical-align:baseline; white-space:pre-wrap">https://reviews.llvm.org/D70290</span></a><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">.
This also includes the clang codegen changes.</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">2.
[July - September 2019] OpenMP dialect for MLIR was discussed / proposed with respect to the f18 compilation stack (keeping FIR in mind).</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">a.
flang-dev discussion link: </span><a href="https://lists.llvm.org/pipermail/flang-dev/2019-September/000020.html" style="text-decoration-line:none"><span style="font-size:11pt; font-family:Arial; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; text-decoration-line:underline; vertical-align:baseline; white-space:pre-wrap">https://lists.llvm.org/pipermail/flang-dev/2019-September/000020.html</span></a></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">b.
Design decisions captured in PPT: </span><a href="https://drive.google.com/file/d/1vU6LsblsUYGA35B_3y9PmBvtKOTXj1Fu/view" style="text-decoration-line:none"><span style="font-size:11pt; font-family:Arial; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; text-decoration-line:underline; vertical-align:baseline; white-space:pre-wrap">https://drive.google.com/file/d/1vU6LsblsUYGA35B_3y9PmBvtKOTXj1Fu/view</span></a></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">c.
MLIR google groups discussion: </span><a href="https://groups.google.com/a/tensorflow.org/forum/#!topic/mlir/4Aj_eawdHiw" style="text-decoration-line:none"><span style="font-size:11pt; font-family:Arial; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; text-decoration-line:underline; vertical-align:baseline; white-space:pre-wrap">https://groups.google.com/a/tensorflow.org/forum/#!topic/mlir/4Aj_eawdHiw</span></a></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">d.
Target constructs design: </span><a href="http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-September/000285.html" style="text-decoration-line:none"><span style="font-size:11pt; font-family:Arial; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; text-decoration-line:underline; vertical-align:baseline; white-space:pre-wrap">http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-September/000285.html</span></a></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">e.
SIMD constructs design: </span><a href="http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-September/000278.html" style="text-decoration-line:none"><span style="font-size:11pt; font-family:Arial; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; text-decoration-line:underline; vertical-align:baseline; white-space:pre-wrap">http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-September/000278.html</span></a></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">3.
[Jan 2020] OpenMP dialect RFC in llvm discourse : </span><a href="https://llvm.discourse.group/t/rfc-openmp-dialect-in-mlir/397" style="text-decoration-line:none"><span style="font-size:11pt; font-family:Arial; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; text-decoration-line:underline; vertical-align:baseline; white-space:pre-wrap">https://llvm.discourse.group/t/rfc-openmp-dialect-in-mlir/397</span></a></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">4.
[Jan- Feb 2020] Implementation of OpenMP dialect in MLIR:</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">a.
The first patch which introduces the OpenMP dialect was pushed.</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">b.
Review of barrier construct is in progress: </span><a href="https://reviews.llvm.org/D72962" style="text-decoration-line:none"><span style="font-size:11pt; font-family:Arial; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; text-decoration-line:underline; vertical-align:baseline; white-space:pre-wrap">https://reviews.llvm.org/D72962</span></a></p>
<a href="https://reviews.llvm.org/D72400" id="LPlnk991155">https://reviews.llvm.org/D72400</a><br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">I
have tried to list below different topics of interest (to different people) around this work. Most of these are in the design phase (or very new) and multiple parties are interested with different sets of goals in mind.</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">I.
Flang frontend and its integration</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">II.
Fortran representation in MLIR / FIR development</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">III.
OpenMP development for flang, OpenMP builder in LLVM.</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">IV.
Loop Transformations in MLIR / LLVM with respect to OpenMP.</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">It
looks like the design has evolved over time and there is no one place which contains the latest design decisions that fits all the different pieces of the puzzle. I will try to deduce it from the above mentioned references. Please correct me If I am referring
to anything which has changed.</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-weight:700; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">A.
</span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">For most OpenMP design discussions, FIR examples
are used (as seen in (2) and (3)). The MLIR examples mentioned in the design only talks about FIR dialect and LLVM dialect.</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">This
completely ignores the likes of standard, affine (where most loop transformations are supposed to happen) and loop dialects. I think it is critical to decouple the OpenMP dialect development in MLIR from the current flang / FIR effort. It would be useful if
someone can mention these examples using existing dialects in MLIR and also how the different transformations / lowerings are planned. </span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-weight:700; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">B.
</span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">In latest RFC(3), it is mentioned that the initial
OpenMP dialect version will be as follows,</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> omp.parallel
{</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> </span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-weight:700; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">omp.do
{</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> fir.do
%i = 0 to %ub3 : !</span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-weight:700; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">fir.integer
</span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">{</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> ...</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> }</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> }</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> }</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">and
then after the "LLVM conversion" it is converted as follows:</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> omp.parallel
{</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> %ub3
=</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> </span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-weight:700; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">omp.do
%i = 0 to %ub3 : !llvm.integer </span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">{</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> ...</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> }</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> }</span></p>
<br>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">a.
Is it the same omp.do operation which now contains the bounds and induction variables of the loop after the LLVM conversion? If so, will the same operation have two different semantics during a single compilation?</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">b.
Will there be different lowerings for various loop operations from different dialects? loop.for and affine.for under omp operations would need different OpenMP / LLVM lowerings</span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-weight:700; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">.</span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">
Currently, both of them are lowered to the CFG based loops during the LLVM dialect conversion (which is much before the proposed OpenMP dialect lowering).</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">There
would be no standard way to represent OpenMP operations (especially the ones which involve loops) in MLIR. This would drastically complicate lowering.</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-weight:700; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">C.</span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">
It is also not mentioned how clauses like </span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-weight:700; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">firstprivate,
shared, private, reduce, map, etc </span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">are lowered
to OpenMP dialect. The example in the RFC contains FIR and LLVM types and nothing about std dialect types. Consider the below example:</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">#pragma
omp parallel for </span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-weight:700; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">reduction(+:x)</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">for
(int i = 0; i < N; ++i)</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> x
+= a[i];</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">How
would the above be represented in OpenMP dialect? and What type would "x" be in MLIR? It is not mentioned in the design as to how the various SSA values for various OpenMP clauses are passed around in OpenMP operations.</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-weight:700; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">D.
</span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Because of (A), (B) and (C), it would be beneficial
to have an omp.</span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-weight:700; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">parallel_do</span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">
operation which has semantics similar to other loop structures (may not be LoopLikeInterface) in MLIR. To me, it looks like having OpenMP operations based on standard MLIR types and operations (scalars and memrefs mainly) is the right way to go.
</span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-weight:700; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap"> </span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Why
not have omp.parallel_do operation with AffineMap based bounds, so as to decouple it from Value/Type similar to affine.for?</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">1.
With the current design, the number of transformations / optimizations that one can write on OpenMP constructs would become limited as there can be any custom loop structure with custom operations / types inside it. </span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">2.
It would also be easier to transform the Loop nests containing OpenMP constructs if the body of the OpenMP operations is well defined (i.e., does not accept arbitrary loop structures). Having nested redundant "parallel" , "target" and "do" regions seems unnecessary. </span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">3.
There would also be new sets of loop structures in new dialects when C/C++ is compiled to MLIR. It would complicate the number of possible combinations inside the OpenMP region.</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-weight:700; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">E.
</span><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Lowering of target constructs mentioned in ( 2(d)
) specifies direct lowering to LLVM IR ignoring all the advantages that MLIR provides. Being able to compile the code for heterogeneous hardware is one of the biggest advantages that MLIR brings to the table. That is being completely missed here. This also
requires solving the problem of handling target information in MLIR. But that is a problem which needs to be solved anyway. Using GPU dialect also gives us an opportunity to represent offloading semantics in MLIR.</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Given
the ability to represent multiple ModuleOps and the existence of GPU dialect, couldn't higher level optimizations on offloaded code be done at MLIR level?. The proposed design would lead us to the same problems that we are currently facing in LLVM IR. </span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Also,
OpenMP codegen will automatically benefit from the GPU dialect based optimizations. For example, it would be way easier to hoist a memory reference out of GPU kernel in MLIR than in LLVM IR.</span></p>
<br>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Thanks,</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; font-family:Arial; color:rgb(0,0,0); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Vinay</span></p>
<br>
</span></div>
</div>
</body>
</html>