<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=Windows-1252">
<style type="text/css" style="display:none;"> P {margin-top:0;margin-bottom:0;} </style>
</head>
<body dir="ltr">
<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">Just to be clear, are you suggesting that if OpenMPIRBuilder is used, MLIR will have to be used as pass-through without the optimizations?</span><br>
</div>
</blockquote>
<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">No, I am not suggesting that. For e.g. We can perform optimizations in the MLIR layer like constant propagation into the OpenMP region,
barrier reduction, loop transformations etc and still use the OpenMPIRBuilder to generate IR for the OpenMP Constructs. With the collapse example, I was just pointing out that we need not always use the OpenMPIRBuilder.<br>
<br>
OpenMPIRBuilder provides us with a fast and easy to use mechanism to generate IR for OpenMP constructs which is as good as what Clang generates. It is our choice whether to use it or not. But if we lower all the OpenMP operations to the LLVM dialect then we
cannot use this mechanism.</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,<br>
Kiran</span></div>
<div id="appendonsend"></div>
<hr style="display:inline-block;width:98%" tabindex="-1">
<div id="divRplyFwdMsg" dir="ltr"><font face="Calibri, sans-serif" style="font-size:11pt" color="#000000"><b>From:</b> Vinay Madhusudan <vinay@compilertree.com><br>
<b>Sent:</b> 17 February 2020 18:36<br>
<b>To:</b> Kiran Chandramohan <Kiran.Chandramohan@arm.com><br>
<b>Cc:</b> Vinay Madhusudan via flang-dev <flang-dev@lists.llvm.org>; llvm-dev@lists.llvm.org <llvm-dev@lists.llvm.org><br>
<b>Subject:</b> Re: [llvm-dev] [flang-dev] About OpenMP dialect in MLIR</font>
<div> </div>
</div>
<div>
<div dir="ltr">
<div dir="ltr"><span id="x_gmail-docs-internal-guid-1217584d-7fff-e9a8-a4e2-9ffccc681ac3">
<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 point here is that we do not want to use MLIR just as a pass-through layer because MLIR has a lot of strengths</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 point here is that if we lower to LLVM dialect, we will not be able to reuse OpenMP codegen & optimisation code from Clang/LLVM.</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">Just
to be clear, are you suggesting that if OpenMPIRBuilder is used, MLIR will have to be used as pass-through without the optimizations?</span></p>
</span><br class="x_gmail-Apple-interchange-newline">
</div>
<br>
<div class="x_gmail_quote">
<div dir="ltr" class="x_gmail_attr">On Sun, Feb 16, 2020 at 6:49 AM Kiran Chandramohan <<a href="mailto:Kiran.Chandramohan@arm.com">Kiran.Chandramohan@arm.com</a>> wrote:<br>
</div>
<blockquote class="x_gmail_quote" style="margin:0px 0px 0px 0.8ex; border-left:1px solid rgb(204,204,204); padding-left:1ex">
<div dir="ltr">
<div style="font-family:Calibri,Arial,Helvetica,sans-serif; font-size:12pt; color:rgb(0,0,0)">
<p dir="ltr" style="color:rgb(32,31,30); font-size:15px; background-color:rgb(255,255,255); margin-top:0px; margin-bottom:0px; line-height:1.38">
Thanks, Vinay for further details and discussion.</p>
<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="color:rgb(32,31,30); font-size:15px; background-color:rgb(255,255,255); margin-top:0px; margin-bottom:0px; line-height:1.38">
<span style="margin:0px; font-size:11pt; font-family:Arial; color:black">> “</span><span style="margin:0px; font-size:11pt; font-family:Arial; color:black">If we decide that the OpenMP construct (for e.g. collapse) can be handled fully in MLIR and that is
the best place to do it (based on experiments) then we will not use the OpenMP IRBuilder for these constructs.</span><span style="margin:0px; font-size:11pt; font-family:Arial; color:black">” -- latest RFC in discourse</span></p>
<br style="color:rgb(32,31,30); font-size:15px; background-color:rgb(255,255,255)">
<p dir="ltr" style="color:rgb(32,31,30); font-size:15px; background-color:rgb(255,255,255); margin-top:0px; margin-bottom:0px; line-height:1.38">
<span style="margin:0px; font-size:11pt; font-family:Arial; color:black">If it is not finalized that the OpenMPIRBuilder will be used for all the constructs, wouldn’t it be better to delay the submission of “translation to LLVM IR” patch in MLIR? Lowering code
will become inconsistent if the OpenMPIRBuilder is used only for a few constructs and not for others.</span></p>
</blockquote>
<div style="font-family:Calibri,Arial,Helvetica,sans-serif; font-size:12pt; color:rgb(0,0,0)">
<br>
</div>
I was hoping that we can identify a set of constructs/clauses which can be fully handled inside the MLIR layer itself. As an example, I provided the collapse clause. This will include constructs/clauses which do not generate runtime API calls, which pass metadata
to LLVM to do some optimisation etc. Yes, the list is not finalized. The point here is that we do not want to use MLIR just as a pass-through layer because MLIR has a lot of strengths. Would it be OK to continue with only constructs that use OpenMPIRBuilder
before the final list is made?<br>
<br style="color:rgb(32,31,30); font-size:15px; background-color:rgb(255,255,255)">
<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="color:rgb(32,31,30); font-size:15px; background-color:rgb(255,255,255); margin-top:0px; margin-bottom:0px; line-height:1.38">
<span style="margin:0px; font-size:11pt; font-family:Arial; color:black">Also, the patch does OpenMP dialect lowering *alongside* LLVM Dialect to LLVM IR. This is different from most dialects which get directly lowered to LLVM Dialect. I think lowering to LLVM
Dialect would be a cleaner way if OpenMPIRBuilder is not being considered for all constructs.<br>
<span style="margin:0px; font-size:11pt; font-family:Arial; color:black">Mehdi also seems to have the same suggestion: “</span><span style="margin:0px; font-size:10pt; font-family:Arial; color:black">I agree that having dialect lowering would be cleaner” in<span> </span></span><a href="https://reviews.llvm.org/D72962" rel="noopener noreferrer" target="_blank" style="margin:0px; font-size:15px; text-decoration-line:none"><span style="margin:0px; font-size:10pt; font-family:Arial; text-decoration-line:underline">https://reviews.llvm.org/D72962</span></a><br>
</span></p>
</blockquote>
<br>
The point here is that if we lower to LLVM dialect, we will not be able to reuse OpenMP codegen & optimisation code from Clang/LLVM. </div>
<div style="font-family:Calibri,Arial,Helvetica,sans-serif; font-size:12pt; color:rgb(0,0,0)">
It was pointed out to me early on by the MLIR developers that there are a few dialects (like NVVM) which are lowered along with LLVM dialect.</div>
<div style="font-family:Calibri,Arial,Helvetica,sans-serif; font-size:12pt; color:rgb(0,0,0)">
<br>
<div style="font-family:Calibri,Arial,Helvetica,sans-serif; font-size:12pt; color:rgb(0,0,0)">
<br>
</div>
<font color="#000000" face="Calibri, Arial, Helvetica, sans-serif" style="background-color:rgb(255,255,255)">You ask a lot of specific questions about which types, dialects, memory access operations will be supported and also the lowering for parallel do. </font>Yes,
the RFC does not provide all this information. This will become clear only as we make progress with the OpenMP dialect. I would like and I am interested to provide answers to all your questions in the following weeks. Please allow some time.<br>
<br>
</div>
</div>
</blockquote>
<div><br>
</div>
<div><span id="x_gmail-docs-internal-guid-1217584d-7fff-e9a8-a4e2-9ffccc681ac3">
<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">Thank
you all for your responses so far. Awaiting your further responses.</span></p>
</span></div>
<div> </div>
<blockquote class="x_gmail_quote" style="margin:0px 0px 0px 0.8ex; border-left:1px solid rgb(204,204,204); padding-left:1ex">
<div dir="ltr">
<div style="font-family:Calibri,Arial,Helvetica,sans-serif; font-size:12pt; color:rgb(0,0,0)">
--Kiran</div>
<div><br>
<br>
</div>
<div id="x_gmail-m_-3649984070533980621appendonsend"></div>
<hr style="display:inline-block; width:98%">
<div id="x_gmail-m_-3649984070533980621divRplyFwdMsg" dir="ltr"><font face="Calibri, sans-serif" color="#000000" style="font-size:11pt"><b>From:</b> Vinay Madhusudan <<a href="mailto:vinay@compilertree.com" target="_blank">vinay@compilertree.com</a>><br>
<b>Sent:</b> 15 February 2020 16:22<br>
<b>To:</b> Kiran Chandramohan <<a href="mailto:Kiran.Chandramohan@arm.com" target="_blank">Kiran.Chandramohan@arm.com</a>><br>
<b>Cc:</b> Vinay Madhusudan via flang-dev <<a href="mailto:flang-dev@lists.llvm.org" target="_blank">flang-dev@lists.llvm.org</a>>;
<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a> <<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a>><br>
<b>Subject:</b> Re: [llvm-dev] [flang-dev] About OpenMP dialect in MLIR</font>
<div> </div>
</div>
<div>
<div dir="ltr"><span id="x_gmail-m_-3649984070533980621x_gmail-docs-internal-guid-14fd630f-7fff-4dad-dba0-24360cfdb7c0">
<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">Reply
to Kiran Chandramohan:</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">>
You are welcome to participate, provide feedback and criticism to change the design as well as to contribute to the implementation.</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">Thank
you Kiran.</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">>
But the latest is what is there in the RFC in discourse.</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 used this as reference for the response.</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">>
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></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: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">Yes, parallel and flush would be the next two constructs that we will do.”
-- from a comment in latest RFC</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">For the above
mentioned reasons, I will try to restrict my reply to how the “parallel (do)” construct would be lowered. </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">>
If it is OK we can have further discussions in discourse as River points 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); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Given
that the multiple components of the LLVM project, namely clang, flang, MLIR and LLVM are involved, llvm-dev is probably a better place, with a much wider audience, until it is clear how different components must interact.</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 is the review for translation to LLVM IR that is in progress.</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: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 we decide that the OpenMP construct (for e.g. collapse) can be handled
fully in MLIR and that is the best place to do it (based on experiments) then we will not use the OpenMP IRBuilder for these constructs.</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">”
-- latest RFC in discourse</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">If
it is not finalized that the OpenMPIRBuilder will be used for all the constructs, wouldn’t it be better to delay the submission of “translation to LLVM IR” patch in MLIR? Lowering code will become inconsistent if the OpenMPIRBuilder is used only for a few
constructs and not for others.</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,
the patch does OpenMP dialect lowering *alongside* LLVM Dialect to LLVM IR. This is different from most dialects which get directly lowered to LLVM Dialect. I think lowering to LLVM Dialect would be a cleaner way if OpenMPIRBuilder is not being considered
for all constructs.</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">Mehdi
also seems to have the same suggestion: “</span><span style="font-size:10pt; 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 agree that having dialect lowering
would be cleaner” in </span><a href="https://reviews.llvm.org/D72962" target="_blank" style="text-decoration-line:none"><span style="font-size:10pt; font-family:Arial; 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>
<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">>
Yes, the design has mildly changed over time to incorporate feedback. But the latest is what is there in the RFC in discourse.</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">RFC
fails to discuss the following (I have also mentioned some of them in my reply to Johannes):</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 proposed plan involves a) lowering F18 AST with OpenMP directly to a mix of OpenMP and FIR dialects. b) converting this finally to a mix of OpenMP and LLVM dialects.</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
is unclear in the RFC what other dialects are considered as supported for OpenMP dialect (std, affine, vector, loop, etc) and how it would be transformed, used and lowered from FIR to LLVM.</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
becomes important to list down the various dialects / operations / types supported for OpenMP (which is mainly defined for C, C++ and Fortran programs. MLIR has a much wider scope.</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
wouldn’t add much value for the proposed OpenMP dialect to be in the MLIR tree if it cannot support at least the relevant standard dialect types / 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-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">>
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).</span></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><a href="https://llvm.discourse.group/t/rfc-openmp-dialect-in-mlir/397/7?u=kiranchandramohan" target="_blank" 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/7?u=kiranchandramohan</span></a></p>
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><a href="https://fosdem.org/2020/schedule/event/llvm_flang/attachments/slides/3839/export/events/attachments/llvm_flang/slides/3839/flang_llvm_frontend.pdf" target="_blank" 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://fosdem.org/2020/schedule/event/llvm_flang/attachments/slides/3839/export/events/attachments/llvm_flang/slides/3839/flang_llvm_frontend.pdf</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">Although
it is mentioned that the affine/ loop.for is used, following things are unclear:</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
am assuming that there will be lowering / conversion code in f18 repo dialect from fir.do to loop.for / affine.for. Is it the case? If so, I think it is worth mentioning it in the “sequential code flow representation” in the RFC.</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
raises the following questions.</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"><br>
</span></p>
<ol style="margin-top:0px; margin-bottom:0px">
<li dir="ltr" style="list-style-type:decimal; 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">
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Which types are supported?
Standard dialect types and FIR types?</span></p>
</li></ol>
<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">For
example, what types are used for Fortran arrays used inside OpenMP regions? Is it std.memref OR Fortran array representation in FIR dialect (fir.array?) OR both? Note that Fortran has support for column major arrays. std.memref supports custom memory layouts.
What custom layouts are supported?</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"><br>
</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">How
would different non-scalar types in standard dialect be lowered to LLVM IR and passed to OpenMP runtime calls? Can you please elaborate 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); background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">The
example provided in slide 20 of the fosdem presentation contains </span></p>
<br>
<p dir="ltr" style="line-height:1.38; text-indent:36pt; 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">“loop.for %j = %lb2 to %ub2 : !integer {“ </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">But
loop.for accepts “index” type. Not sure what type “!integer” represents here.</span></p>
<br>
<ol start="2" style="margin-top:0px; margin-bottom:0px">
<li dir="ltr" style="list-style-type:decimal; 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">
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">What are the different
memory access operations which are supported inside the OpenMP region and lowered to proper OpenMP runtime calls in LLVM IR?</span></p>
</li></ol>
<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
possibilities are: </span></p>
<ol style="margin-top:0px; margin-bottom:0px">
<li dir="ltr" style="list-style-type:decimal; 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; margin-left:36pt">
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">affine.load / affine.store </span></p>
</li><li dir="ltr" style="list-style-type:decimal; 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; margin-left:36pt">
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">std.load / std.store </span></p>
</li><li dir="ltr" style="list-style-type:decimal; 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; margin-left:36pt">
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">FIR dialect memory
access operations.</span></p>
</li></ol>
<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 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.
</span><a href="http://lists.llvm.org/pipermail/flang-dev/2019-September/000042.html" target="_blank" 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.llvm.org/pipermail/flang-dev/2019-September/000042.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">Looks
like placement (MLIR / LLVM) of various transformations related to OpenMP has not been finalized, from what I could infer from Johannes’s reply and the below text in the latest RFC in discourse:</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: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">So
there exist some questions regarding where the optimisations should be carried out. We will decide on which framework to choose only after some experimentation.”</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) we need to keep the loops separately so as to take advantage of transformations that other dialects like affine/loop would provide.</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)
Keeping the loops separate from the OpenMP operations will expose them to the “regular” transformations passes in MLIR inside the OpenMP region. Most of them are invalid or in-efficient for OpenMP operations.</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">Examples: </span></p>
<ol style="margin-top:0px; margin-bottom:0px">
<li dir="ltr" style="list-style-type:decimal; 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">
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Constant propagation
example mentioned by Johannes in this thread. (omp task shared(x))</span></p>
</li><li dir="ltr" style="list-style-type:decimal; 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">
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Loop (nest) transformations
(permute / split / fuse / tile, etc) will happen ignoring the surrounding OpenMP operations.</span></p>
</li><li dir="ltr" style="list-style-type:decimal; 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">
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Hoisting and sinking
of various memory/ SSA values inside the OpenMP region. This goes against the likes of “map”, “firstprivate”, shared, etc clauses and more.</span></p>
</li></ol>
<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)
Various loop operations (loop.for, affine.for, fir.do) have (or will have) different transformations/ optimization passes which are different from one another. </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">Example:</span></p>
<ol style="margin-top:0px; margin-bottom:0px">
<li dir="ltr" style="list-style-type:decimal; 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">
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">AffineLoopInvariantCodeMotion.cpp
is different from LoopInvariantCodeMotion.cpp. </span></p>
</li><li dir="ltr" style="list-style-type:decimal; 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">
<p dir="ltr" style="line-height:1.38; margin-top:0pt; margin-bottom:0pt"><span style="font-size:11pt; background-color:transparent; font-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Other Loop transformation
passes for affine.for</span></p>
</li></ol>
<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">These
loops also use different Types and memory access operations in general for transformations. Example, most Affine dialect transformations (if not all) work on affine.load and affine.store 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-variant-numeric:normal; font-variant-east-asian:normal; vertical-align:baseline; white-space:pre-wrap">Supporting
different loop operations means that there would be *OpenMP specific transformations* for each one of them and also requires a way to restrict each of them from existing transformations (when nested in OpenMP constructs).</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 different lowerings for different loop operations as well. Example, affine.for and loop.for would have to be lowered to omp.do in different ways. </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">From
slide 20 of fosdem presentation you mentioned, the LLVM + OpenMP dialect representation is 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">------------------------------</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">Mlir.region(…)
{ </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"> 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"> omp.do
%i = 0 to %ub3 : !integer { </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>
<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">Currently,
the LLVM Dialect doesn’t contain a high level loop operation. It is all based on CFG implementation. </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">Will
omp.do follow the same structure (SizedRegion<1>) as loop.for? OR there would be CFG for LLVM Dialect based loop operation?</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">Can
you please mention how the OpenMP + LLVM dialect will look like for the below parallel do construct?</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">integer
:: i=1, k=10</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">integer
:: a(10), b(10), c(10)</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"> !$omp
parallel 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"> do
i = 1, k</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"> if
(i .ne. 1) *cycle*</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"> c(i)
= a(i) + b(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"> end
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"> !$omp
end parallel do</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">print
*,c</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"><br>
</span></p>
<p 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 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>
</span></div>
<br>
<div>
<div dir="ltr">On Fri, Feb 14, 2020 at 6:52 AM Kiran Chandramohan via llvm-dev <<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a>> wrote:<br>
</div>
<blockquote style="margin:0px 0px 0px 0.8ex; border-left:1px solid rgb(204,204,204); padding-left:1ex">
<div 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">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"><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"><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">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="x_gmail-m_-3649984070533980621x_gmail-m_-8677415914596354611LPlnk278665" target="_blank" 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="x_gmail-m_-3649984070533980621x_gmail-m_-8677415914596354611LPlnk452670" target="_blank" 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" target="_blank" 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" target="_blank" style="margin:0px"><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>
</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">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">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"><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"><span style="font-family:Arial; background-color:rgb(255,255,255); display:inline">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">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"><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"><span style="font-family:Arial; background-color:rgb(255,255,255); display:inline">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">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="x_gmail-m_-3649984070533980621x_gmail-m_-8677415914596354611LPNoLP808880" target="_blank">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="x_gmail-m_-3649984070533980621x_gmail-m_-8677415914596354611LPNoLP581920" target="_blank">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="x_gmail-m_-3649984070533980621x_gmail-m_-8677415914596354611LPlnk702758" target="_blank">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"><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">
<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">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><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">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">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"><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"><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">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 style="display:inline-block; width:98%">
<div id="x_gmail-m_-3649984070533980621x_gmail-m_-8677415914596354611divRplyFwdMsg" dir="ltr">
<font face="Calibri, sans-serif" color="#000000" style="font-size:11pt"><b>From:</b> flang-dev <<a href="mailto:flang-dev-bounces@lists.llvm.org" target="_blank">flang-dev-bounces@lists.llvm.org</a>> on behalf of Vinay Madhusudan via flang-dev <<a href="mailto:flang-dev@lists.llvm.org" target="_blank">flang-dev@lists.llvm.org</a>><br>
<b>Sent:</b> 13 February 2020 16:33<br>
<b>To:</b> <a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a> <<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a>>;
<a href="mailto:flang-dev@lists.llvm.org" target="_blank">flang-dev@lists.llvm.org</a> <<a href="mailto:flang-dev@lists.llvm.org" target="_blank">flang-dev@lists.llvm.org</a>><br>
<b>Subject:</b> [flang-dev] About OpenMP dialect in MLIR</font>
<div> </div>
</div>
<div>
<div dir="ltr"><span id="x_gmail-m_-3649984070533980621x_gmail-m_-8677415914596354611x_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" target="_blank" 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" target="_blank" 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" target="_blank" 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" target="_blank" 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" target="_blank" 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" target="_blank" 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" target="_blank" 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" target="_blank" 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" target="_blank" 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="x_gmail-m_-3649984070533980621x_gmail-m_-8677415914596354611LPlnk991155" target="_blank">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>
</div>
_______________________________________________<br>
LLVM Developers mailing list<br>
<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a><br>
<a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" rel="noreferrer" target="_blank">https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br>
</blockquote>
</div>
</div>
</div>
</blockquote>
</div>
</div>
</div>
</body>
</html>