<div dir="ltr"><div dir="ltr"><span id="gmail-docs-internal-guid-0fbc8020-7fff-cfdc-c8b3-3fa5e2fff142"><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 for the reply!</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 sounds like LLVM IR is being considered for optimizations in OpenMP constructs. There seems to be plans regarding improvement of LLVM IR Framework for providing things required for OpenMP / flang(?)</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">Are there any design considerations which contain pros and cons about using the MLIR vs LLVM IR for various OpenMP related optimizations/ transformations?</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">The latest RFC [ (3) in my original post ] mentions that:</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">> </span><span style="font-size:11.5pt;font-family:Arial;font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">So there exist some questions regarding where the optimisations should be carried out.</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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">Could you please provide more details on this?</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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">I would like to quote Chris here:</span></p><br><span style="font-size:11pt;font-family:Arial;color:rgb(0,0,0);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">“if you ignore the engineering expense, it would clearly make sense to reimplement the mid-level LLVM optimizers on top of MLIR and replace include/llvm/IR with a dialect definition in MLIR instead.“ -- </span><span style="text-decoration-line:underline;font-size:11pt;font-family:Arial;font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap"><a href="http://lists.llvm.org/pipermail/llvm-dev/2020-January/138341.html" style="text-decoration-line:none">http://lists.llvm.org/pipermail/llvm-dev/2020-January/138341.html</a> </span></span></div><div dir="ltr"><br></div><div><b>Rest of the comment are inlined.</b></div><div dir="ltr"><br></div><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Thu, Feb 13, 2020 at 11:48 PM Johannes Doerfert <<a href="mailto:jdoerfert@anl.gov">jdoerfert@anl.gov</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">Hi Vinay,<br>
<br>
Thanks for taking an interest and the detailed discussion.<br>
<br>
To start by picking a few paragraph from your email to clarify a couple<br>
of things that lead to the current design or that might otherwise need<br>
clarification. We can talk about other points later as well.<br>
<br>
[<br>
Site notes:<br>
1) I'm not an MLIR person.<br>
2) It seems unfortnuate that we do not have a mlir-dev list.<br>
]<br>
<br>
<br>
> 1. With the current design, the number of transformations / optimizations<br>
> that one can write on OpenMP constructs would become limited as there can<br>
> be any custom loop structure with custom operations / types inside it.<br>
<br>
OpenMP, as an input language, does not make many assumptions about the<br>
code inside of constructs*.</blockquote><div><br></div><div><span id="gmail-docs-internal-guid-57299942-7fff-9101-fd47-eb81bf65f033"><span style="font-size:11pt;font-family:Arial;color:rgb(0,0,0);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">This isn’t entirely correct because the current OpenMP API specification (</span><a href="https://www.openmp.org/spec-html/5.0/openmpch1.html" style="text-decoration-line:none"><span style="font-size:11pt;font-family:Arial;font-variant-numeric:normal;font-variant-east-asian:normal;text-decoration-line:underline;vertical-align:baseline;white-space:pre-wrap">https://www.openmp.org/spec-html/5.0/openmpch1.html</span></a><span style="font-size:11pt;font-family:Arial;color:rgb(0,0,0);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">) assumes that the code inside the constructs belong to C, C++ and Fortran programs. </span></span><br></div><div> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"> So, inside a parallel can be almost anything<br>
the base language has to offer, both lexically and dynamically.<br></blockquote><div> </div><div><span id="gmail-docs-internal-guid-3b7dffe5-7fff-3211-3960-b6bbdf2bf41b"><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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">I am mostly concerned with the MLIR side of things for OpenMP representation.</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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap"> MLIR can not only support operations for General Purpose languages like C,C++, Fortran, etc but also various Domain Specific Language representations as dialects (Example, ML, etc.). Note that there is also SPIR V dialect which is again meant for “Parallel Compute”. </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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap"> It becomes important to define the scope of the dialects / operations / types supported inside OpenMP operations in MLIR.</span></p></span></div><div> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
Assuming otherwise is not going to work. Analyzing a "generic" OpenMP<br>
representation in order to determine if can be represented as a more<br>
restricted "op" seems at least plausible. You will run into various<br>
issue, some mentioned explicitly below.</blockquote><div> </div><div><span id="gmail-docs-internal-guid-8825957f-7fff-fd8f-4330-38754efa9057"><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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">Isn’t it the other way around? For example, it doesn’t make much sense to wrap OpenMP operations for SPIR-V operations / types. </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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">I think it is important to specify (in the design) which existing MLIR dialects are supported in this effort and the various lowerings / transformations / optimizations which are planned for them.</span></p></span></div><div> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"> For starters, you still have to<br>
generate proper OpenMP runtime calls, e.g., from your GPU dialect, even<br>
if it is "just" to make sure the OMPD/OMPT interfaces expose useful<br>
information.<br>
<br></blockquote><div><span id="gmail-docs-internal-guid-cf8622cb-7fff-2408-2f83-270b0479becb"><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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">You can have a well-defined call-like mlir::Operation which calls the GPU kernel. Perform all cross-device transformations in an easier way. Then, this operation can be lowered to OpenMP runtime calls during LLVM dialect conversion. I think this is much better than directly having calls to the OpenMP runtime library based on a kernel name mentioned in llvm::GlobalVariable.</span></p></span></div><div> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
<br>
* I preclude the `omp loop` construct here as it is not even implemented<br>
anywhere as far as I know.<br>
<br>
<br>
> 2. It would also be easier to transform the Loop nests containing OpenMP<br>
> constructs if the body of the OpenMP operations is well defined (i.e., does<br>
> not accept arbitrary loop structures). Having nested redundant "parallel" ,<br>
> "target" and "do" regions seems unnecessary.<br>
<br>
As mentioned above, you cannot start with the assumption OpenMP input is<br>
structured this this way. You have to analyze it first. This is the same<br>
reason we cannot simply transform C/C++ `for loops` into `affine.for`<br>
without proper analysis of the loop body.<br>
<br>
Now, more concrete. Nested parallel and target regions are not<br>
necessarily redundant, nor can/should we require the user not to have<br>
them. Nested parallelism can easily make sense, depending on the problem<br>
decomposition. Nested target will make a lot of sense with reverse<br>
offload, which is already in the standard, and it also should be allowed<br>
for the sake of a modular (user) code base.<br></blockquote><div><br></div><span id="gmail-docs-internal-guid-b8a31152-7fff-e3db-1487-577fca28b859"><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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">Just to be clear, having all three of “target”, “parallel” and “do” doesn’t represent “Nested parallelism” at all in the proposed design! ( 2(d) ).</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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">omp.target {</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);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);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);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);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);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);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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">Above invokes a call to the tgt_target() for the code inside omp.do as mentioned in the proposal.</span></p></span><div> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
<br>
> 3. There would also be new sets of loop structures in new dialects when<br>
> C/C++ is compiled to MLIR. It would complicate the number of possible<br>
> combinations inside the OpenMP region.<br>
<br>
Is anyone working on this? If so, what is the timeline? I personally was<br>
not expecting Clang to switch over to MLIR any time soon but I am happy<br>
if someone wants to correct me on this. I mention this only because it<br>
interacts with the arguments I will make below.<br>
<br>
<br>
> E. Lowering of target constructs mentioned in ( 2(d) ) specifies direct<br>
> lowering to LLVM IR ignoring all the advantages that MLIR provides. Being<br>
> able to compile the code for heterogeneous hardware is one of the biggest<br>
> advantages that MLIR brings to the table. That is being completely missed<br>
> here. This also requires solving the problem of handling target information<br>
> in MLIR. But that is a problem which needs to be solved anyway. Using GPU<br>
> dialect also gives us an opportunity to represent offloading semantics in<br>
> MLIR.<br>
<br>
I'm unsure what the problem with "handling target information in MLIR" is but<br>
whatever design we end up with, we need to know about the target<br>
(triple) in all stages of the pipeline, even if it is just to pass it<br>
down.<br>
<br>
<br>
> Given the ability to represent multiple ModuleOps and the existence of GPU<br>
> dialect, couldn't higher level optimizations on offloaded code be done at<br>
> MLIR level?. The proposed design would lead us to the same problems that we<br>
> are currently facing in LLVM IR.<br>
><br>
> Also, OpenMP codegen will automatically benefit from the GPU dialect based<br>
> optimizations. For example, it would be way easier to hoist a memory<br>
> reference out of GPU kernel in MLIR than in LLVM IR.<br>
<br>
While I agree with the premise that you can potentially reuse MLIR<br>
transformations, it might not be as simple in practice.<br>
<br>
As mentioned above, you cannot assume much about OpenMP codes, almost<br>
nothing for a lot of application codes I have seen. Some examples:<br>
<br>
If you have a function call, or any synchronization event for that<br>
matter, located between two otherwise adjacent target regions (see<br>
below), you cannot assume the two target regions will be offloaded to<br>
the same device.<br>
```<br>
#omp target<br>
{}<br>
foo();<br>
#omp target<br>
{}<br>
```<br></blockquote><div><br></div><div><span id="gmail-docs-internal-guid-42371700-7fff-de5a-808f-8b387f30199e"><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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">These kinds of optimizations are much easier to write 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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">LLVM IR for the above code would contain a series of instructions of OpenMP runtime call setup and foo() in the middle followed by another set of OpenMP runtime related instructions. The body of the two target constructs would be in two different outlined functions (if not modules).</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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">It takes quite a bit of code to do analysis / transformation to write any optimization on the generated 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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">vs.</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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">MLIR provides a way to represent the operations closer to the source. It is as simple as checking the next operation(s) in the mlir::Block. OpenMP target operation contains an inlined region which can easily be fused/ split / or any other valid transformation for that matter.</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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">Note that you can also perform various Control Structure Analysis / Transformations much easier in MLIR. For example, you can decide to execute foo() based on certain conditions, and you can merge the two target regions in the else path.</span></p></span></div><div> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
Similarly, you cannot assume a `omp parallel` is allowed to be executed<br>
with more than a single thread, or that a `omp [parallel] for` does not<br>
have loop carried data-dependences, ...<br></blockquote><div><br></div><div><span id="gmail-docs-internal-guid-15dd07d7-7fff-81a9-eccc-b6619f8ff688"><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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">With multi-dimensional index support for arrays, wouldn’t it be better to do the data dependence analysis in MLIR?</span></p><br><span style="font-size:11pt;font-family:Arial;color:rgb(0,0,0);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">LLVM IR has linearized subscripts for multi-dimensional arrays. llvm::DependenceAnalysis tries to “guess” the indices based on different patterns in SCEV. It takes an </span><a href="http://llvm.org/devmtg/2020-04/talks.html#LightningTalk_88" style="text-decoration-line:none"><span style="font-size:11pt;font-family:Arial;font-variant-numeric:normal;font-variant-east-asian:normal;text-decoration-line:underline;vertical-align:baseline;white-space:pre-wrap">intrinsic</span></a><span style="font-size:11pt;font-family:Arial;color:rgb(0,0,0);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap"> </span><span style="font-size:10.5pt;font-family:Roboto,sans-serif;color:rgb(0,0,0);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">or metadata or some other mechanism of communication from the front end</span><span style="font-size:10.5pt;font-family:Roboto,sans-serif;color:rgb(60,64,67);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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">(not the built-in set of instructions) to solve this problem. </span></span><br></div><div> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
Data-sharing attributes are also something that has to be treated<br>
carefully:<br>
```<br>
x = 5;<br>
#omp task<br>
x = 3;<br>
print(x);<br>
```<br>
Should print 5, not 3.<br></blockquote><div><br></div><div><span style="color:rgb(0,0,0);font-family:Arial;font-size:11pt;white-space:pre-wrap">You can have “x” as a locally defined variable inside the “task” contained region in MLIR OR custom data-sharing attributes in OpenMP dialect.</span> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
<br>
I hope I convinced you that OpenMP is not trivially mappable to existing<br>
dialects without proper analysis. If not, please let me know why you<br>
expect it to be.<br>
<br></blockquote><div><span id="gmail-docs-internal-guid-2ff4af67-7fff-991e-500b-6804de8abd8b"><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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">I do not see much reason why the issues you mentioned can’t trivially be mapped to the MLIR infrastructure. There is an easy way to define custom operations / types / attributes in OpenMP dialect and perform optimizations based on the *IR that is created especially for OpenMP*. The analysis / transformations required can be easily written on the custom operations defined rather than having a lowered form in the 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);font-variant-numeric:normal;font-variant-east-asian:normal;vertical-align:baseline;white-space:pre-wrap">The various dialects / transformations in MLIR are in development / early phase (Example, GPU dialect) waiting to be improved with use cases such as this!</span></p></span></div><div> </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
<br>
Now when it comes to code analyses, LLVM-IR offers a variety of<br>
interesting features, ranging from a mature set of passes to the<br>
cross-language LTO capabilities. We are working on the missing parts,<br>
e.g., heterogeneous llvm::Modules as we speak. Simple OpenMP<br>
optimizations are already present in LLVM and interesting ones are<br>
prototyped for a while now (let me know if you want to see more not-yet<br>
merged patches/optimizations). I also have papers, results, and<br>
talks that might be interesting here. Let me know if you need pointers<br>
to them.<br>
<br>
<br>
Cheers,<br>
Johannes<br>
<br>
<br>
<br>
On 02/13, Vinay Madhusudan via llvm-dev wrote:<br>
> Hi,<br>
> <br>
> I have few questions / concerns regarding the design of OpenMP dialect in<br>
> MLIR that is currently being implemented, mainly for the f18 compiler.<br>
> Below, I summarize the current state of various efforts in clang / f18 /<br>
> MLIR / LLVM regarding this. Feel free to add to the list in case I have<br>
> missed something.<br>
> <br>
> 1. [May 2019] An OpenMPIRBuilder in LLVM was proposed for flang and clang<br>
> frontends. Note that this proposal was before considering MLIR for FIR.<br>
> <br>
> a. llvm-dev proposal :<br>
> <a href="http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-May/000197.html" rel="noreferrer" target="_blank">http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-May/000197.html</a><br>
> <br>
> b. Patches in review: <a href="https://reviews.llvm.org/D70290" rel="noreferrer" target="_blank">https://reviews.llvm.org/D70290</a>. This also includes<br>
> the clang codegen changes.<br>
> <br>
> 2. [July - September 2019] OpenMP dialect for MLIR was discussed /<br>
> proposed with respect to the f18 compilation stack (keeping FIR in mind).<br>
> <br>
> a. flang-dev discussion link:<br>
> <a href="https://lists.llvm.org/pipermail/flang-dev/2019-September/000020.html" rel="noreferrer" target="_blank">https://lists.llvm.org/pipermail/flang-dev/2019-September/000020.html</a><br>
> <br>
> b. Design decisions captured in PPT:<br>
> <a href="https://drive.google.com/file/d/1vU6LsblsUYGA35B_3y9PmBvtKOTXj1Fu/view" rel="noreferrer" target="_blank">https://drive.google.com/file/d/1vU6LsblsUYGA35B_3y9PmBvtKOTXj1Fu/view</a><br>
> <br>
> c. MLIR google groups discussion:<br>
> <a href="https://groups.google.com/a/tensorflow.org/forum/#!topic/mlir/4Aj_eawdHiw" rel="noreferrer" target="_blank">https://groups.google.com/a/tensorflow.org/forum/#!topic/mlir/4Aj_eawdHiw</a><br>
> <br>
> d. Target constructs design:<br>
> <a href="http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-September/000285.html" rel="noreferrer" target="_blank">http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-September/000285.html</a><br>
> <br>
> e. SIMD constructs design:<br>
> <a href="http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-September/000278.html" rel="noreferrer" target="_blank">http://lists.flang-compiler.org/pipermail/flang-dev_lists.flang-compiler.org/2019-September/000278.html</a><br>
> <br>
> 3. [Jan 2020] OpenMP dialect RFC in llvm discourse :<br>
> <a href="https://llvm.discourse.group/t/rfc-openmp-dialect-in-mlir/397" rel="noreferrer" target="_blank">https://llvm.discourse.group/t/rfc-openmp-dialect-in-mlir/397</a><br>
> <br>
> 4. [Jan- Feb 2020] Implementation of OpenMP dialect in MLIR:<br>
> <br>
> a. The first patch which introduces the OpenMP dialect was pushed.<br>
> <br>
> b. Review of barrier construct is in progress:<br>
> <a href="https://reviews.llvm.org/D72962" rel="noreferrer" target="_blank">https://reviews.llvm.org/D72962</a><br>
> <br>
> I have tried to list below different topics of interest (to different<br>
> people) around this work. Most of these are in the design phase (or very<br>
> new) and multiple parties are interested with different sets of goals in<br>
> mind.<br>
> <br>
> I. Flang frontend and its integration<br>
> <br>
> II. Fortran representation in MLIR / FIR development<br>
> <br>
> III. OpenMP development for flang, OpenMP builder in LLVM.<br>
> <br>
> IV. Loop Transformations in MLIR / LLVM with respect to OpenMP.<br>
> <br>
> It looks like the design has evolved over time and there is no one place<br>
> which contains the latest design decisions that fits all the different<br>
> pieces of the puzzle. I will try to deduce it from the above mentioned<br>
> references. Please correct me If I am referring to anything which has<br>
> changed.<br>
> <br>
> A. For most OpenMP design discussions, FIR examples are used (as seen in<br>
> (2) and (3)). The MLIR examples mentioned in the design only talks about<br>
> FIR dialect and LLVM dialect.<br>
> <br>
> This completely ignores the likes of standard, affine (where most loop<br>
> transformations are supposed to happen) and loop dialects. I think it is<br>
> critical to decouple the OpenMP dialect development in MLIR from the<br>
> current flang / FIR effort. It would be useful if someone can mention these<br>
> examples using existing dialects in MLIR and also how the different<br>
> transformations / lowerings are planned.<br>
> <br>
> B. In latest RFC(3), it is mentioned that the initial OpenMP dialect<br>
> version will be as follows,<br>
> <br>
> omp.parallel {<br>
> <br>
> omp.do {<br>
> <br>
> fir.do %i = 0 to %ub3 : !fir.integer {<br>
> <br>
> ...<br>
> <br>
> }<br>
> <br>
> }<br>
> <br>
> }<br>
> <br>
> and then after the "LLVM conversion" it is converted as follows:<br>
> <br>
> omp.parallel {<br>
> <br>
> %ub3 =<br>
> <br>
> omp.do %i = 0 to %ub3 : !llvm.integer {<br>
> <br>
> ...<br>
> <br>
> }<br>
> <br>
> }<br>
> <br>
> <br>
> a. Is it the same omp.do operation which now contains the bounds and<br>
> induction variables of the loop after the LLVM conversion? If so, will the<br>
> same operation have two different semantics during a single compilation?<br>
> <br>
> b. Will there be different lowerings for various loop operations from<br>
> different dialects? loop.for and affine.for under omp operations would need<br>
> different OpenMP / LLVM lowerings. Currently, both of them are lowered to<br>
> the CFG based loops during the LLVM dialect conversion (which is much<br>
> before the proposed OpenMP dialect lowering).<br>
> <br>
> There would be no standard way to represent OpenMP operations (especially<br>
> the ones which involve loops) in MLIR. This would drastically complicate<br>
> lowering.<br>
> <br>
> C. It is also not mentioned how clauses like firstprivate, shared, private,<br>
> reduce, map, etc are lowered to OpenMP dialect. The example in the RFC<br>
> contains FIR and LLVM types and nothing about std dialect types. Consider<br>
> the below example:<br>
> <br>
> #pragma omp parallel for reduction(+:x)<br>
> <br>
> for (int i = 0; i < N; ++i)<br>
> <br>
> x += a[i];<br>
> <br>
> How would the above be represented in OpenMP dialect? and What type would<br>
> "x" be in MLIR? It is not mentioned in the design as to how the various<br>
> SSA values for various OpenMP clauses are passed around in OpenMP<br>
> operations.<br>
> <br>
> D. Because of (A), (B) and (C), it would be beneficial to have an omp.<br>
> parallel_do operation which has semantics similar to other loop structures<br>
> (may not be LoopLikeInterface) in MLIR. To me, it looks like having OpenMP<br>
> operations based on standard MLIR types and operations (scalars and memrefs<br>
> mainly) is the right way to go.<br>
> <br>
> Why not have omp.parallel_do operation with AffineMap based bounds, so as<br>
> to decouple it from Value/Type similar to affine.for?<br>
> <br>
> 1. With the current design, the number of transformations / optimizations<br>
> that one can write on OpenMP constructs would become limited as there can<br>
> be any custom loop structure with custom operations / types inside it.<br>
> <br>
> 2. It would also be easier to transform the Loop nests containing OpenMP<br>
> constructs if the body of the OpenMP operations is well defined (i.e., does<br>
> not accept arbitrary loop structures). Having nested redundant "parallel" ,<br>
> "target" and "do" regions seems unnecessary.<br>
> <br>
> 3. There would also be new sets of loop structures in new dialects when<br>
> C/C++ is compiled to MLIR. It would complicate the number of possible<br>
> combinations inside the OpenMP region.<br>
> <br>
> E. Lowering of target constructs mentioned in ( 2(d) ) specifies direct<br>
> lowering to LLVM IR ignoring all the advantages that MLIR provides. Being<br>
> able to compile the code for heterogeneous hardware is one of the biggest<br>
> advantages that MLIR brings to the table. That is being completely missed<br>
> here. This also requires solving the problem of handling target information<br>
> in MLIR. But that is a problem which needs to be solved anyway. Using GPU<br>
> dialect also gives us an opportunity to represent offloading semantics in<br>
> MLIR.<br>
> <br>
> Given the ability to represent multiple ModuleOps and the existence of GPU<br>
> dialect, couldn't higher level optimizations on offloaded code be done at<br>
> MLIR level?. The proposed design would lead us to the same problems that we<br>
> are currently facing in LLVM IR.<br>
> <br>
> Also, OpenMP codegen will automatically benefit from the GPU dialect based<br>
> optimizations. For example, it would be way easier to hoist a memory<br>
> reference out of GPU kernel in MLIR than in LLVM IR.<br>
</blockquote></div></div>