<font size=2 face="sans-serif">Hi Johannes,</font><br><br><font size=2 face="sans-serif">First of all thanks for looking into
the matter of improving non-SPMD mode!</font><br><br><font size=2 face="sans-serif">I have a question regarding the state
machine that you said you'd like to replace/improve. There are cases (such
as target regions that span multiple compilation units) where the switch
statement is required. Is this something that your changes will touch in
any way?</font><br><br><font size=2 face="sans-serif">My next question is, for the workloads
which are in the same compilation unit there is a trick that code gen performs
(or could perform I'm not sure if this has been upstreamed) where it can
check for the specific name of an outlined function and then just call
it directly thus making that function inline-able (thus erasing most if
not all the overhead of having the state machine in the first place). In
other words the "worst" part of the switch statement will only
apply to outlined functions from other compilation units. With this in
mind what would the impact of your changes be in the end? If this part
isn't clear I can do some digging to find out how this actually works in
more details it's been too long since I've had to look at this part.</font><br><br><font size=2 face="sans-serif">Can you share some performance numbers
given an example you have been looking at? I see you have one that uses
"#pragma omp atomic". I would avoid using something like that
since it may have other overheads not related to your changes. I would
put together an example with this directive structure:</font><br><br><font size=2 face="sans-serif"><b>#pragma omp target teams distribute</b></font><br><font size=2 face="sans-serif"><b>for(...){</b></font><br><font size=2 face="sans-serif"> <code1></font><br><font size=2 face="sans-serif"> <b>#pragma omp parallel for</b></font><br><font size=2 face="sans-serif"><b> for(...) {</b></font><br><font size=2 face="sans-serif"> <code2></font><br><font size=2 face="sans-serif"> }</font><br><font size=2 face="sans-serif"> <code3></font><br><font size=2 face="sans-serif">}</font><br><br><font size=2 face="sans-serif">which forces the use of the master-worker
scheme (non-SPMD mode) without any other distractions.</font><br><br><font size=2 face="sans-serif">It would then be interesting to understand
how you plan to change the LLVM code generated for this, what the overheads
that you're targeting are (register usage, synchronization cost etc), and
then what the performance gain is compared to the current scheme.</font><br><br><font size=2 face="sans-serif">Thanks,</font><br><br><font size=2 face="sans-serif">--Doru</font><font size=1 face="Arial"><br></font><font size=2 face="sans-serif"><br></font><br><br><br><br><font size=1 color=#5f5f5f face="sans-serif">From:
</font><font size=1 face="sans-serif">"Doerfert, Johannes"
<jdoerfert@anl.gov></font><br><font size=1 color=#5f5f5f face="sans-serif">To:
</font><font size=1 face="sans-serif">Alexey Bataev <a.bataev@outlook.com></font><br><font size=1 color=#5f5f5f face="sans-serif">Cc:
</font><font size=1 face="sans-serif">Guray Ozen <gozen@nvidia.com>,
Gheorghe-Teod Bercea <gheorghe-teod.bercea@ibm.com>, "openmp-dev@lists.llvm.org"
<openmp-dev@lists.llvm.org>, "Finkel, Hal J." <hfinkel@anl.gov>,
"Gregory.Rodgers@amd.com" <Gregory.Rodgers@amd.com>, "kli@ca.ibm.com"
<kli@ca.ibm.com></font><br><font size=1 color=#5f5f5f face="sans-serif">Date:
</font><font size=1 face="sans-serif">01/30/2019 04:14 PM</font><br><font size=1 color=#5f5f5f face="sans-serif">Subject:
</font><font size=1 face="sans-serif">Re: [RFC] Late
(OpenMP) GPU code "SPMD-zation"</font><br><hr noshade><br><br><br><font size=3 face="Helvetica">I don't really see "many ifs and
maybes", actually none.</font><br><br><font size=3 face="Helvetica">Anyway, I will now work on a patch set
that adds the new functionality under a cmd flag</font><br><font size=3 face="Helvetica">in order to showcase correctness and
performance on real code. </font><br><br><font size=3 face="Helvetica">If you, or somebody else, have interesting
examples, please feel free to point me at them.</font><br><br><font size=3 face="Helvetica">Thanks,</font><br><font size=3 face="Helvetica"> Johannes</font><br><br><p><hr><br><font size=2 face="sans-serif"><b>From:</b> Alexey Bataev <a.bataev@outlook.com><b><br>Sent:</b> Wednesday, January 30, 2019 2:18:19 PM<b><br>To:</b> Doerfert, Johannes<b><br>Cc:</b> Guray Ozen; Gheorghe-Teod Bercea; openmp-dev@lists.llvm.org; Finkel,
Hal J.; Gregory.Rodgers@amd.com; kli@ca.ibm.com<b><br>Subject:</b> Re: [RFC] Late (OpenMP) GPU code "SPMD-zation"</font><font size=3></font><br><font size=3> </font><br><font size=3>Currently, there are too many "if"s and "maybe"s.
If you can provide solution that does not break anything and does not affect
the performance, does not require changes in the backend - then go ahead
with the patches.</font><p><tt><font size=3>-------------<br>Best regards,<br>Alexey Bataev</font></tt><br><font size=3>30.01.2019 14:49, Doerfert, Johannes пишет:</font><br><font size=3 face="Helvetica">No, SPMD mode will not be affected at
all.</font><br><br><font size=3 face="Helvetica">The "worse" part is the following:</font><br><font size=3 face="Helvetica"> If we inline runtime library calls
before the openmp-opt pass had a chance to look at the code,</font><br><font size=3 face="Helvetica"> we will not have a customized
state machine for the __non-SPMD__ case. That is, the if-cascade </font><br><font size=3 face="Helvetica"> checking the work function pointer
is not there.</font><br><br><font size=3 face="Helvetica">Avoiding this potential performance decline
is actually very easy. While we do not have the "inline_late"
capability,</font><br><font size=3 face="Helvetica">run the openmp-opt pass __before__ the
inliner and we will not get "worse" code. We might however miss
out on</font><br><font size=3 face="Helvetica">_new_ non-SPMD -> SPMD transformations.</font><p><br><font size=3 face="Helvetica">Does that make sense? </font><br><p><hr><br><font size=2 face="sans-serif"><b>From:</b> Alexey Bataev </font><a href=mailto:a.bataev@outlook.com><font size=2 color=blue face="sans-serif"><u><a.bataev@outlook.com></u></font></a><font size=2 face="sans-serif"><b><br>Sent:</b> Wednesday, January 30, 2019 1:44:10 PM<b><br>To:</b> Doerfert, Johannes<b><br>Cc:</b> Guray Ozen; Gheorghe-Teod Bercea; </font><a href="mailto:openmp-dev@lists.llvm.org"><font size=2 color=blue face="sans-serif"><u>openmp-dev@lists.llvm.org</u></font></a><font size=2 face="sans-serif">;
Finkel, Hal J.; </font><a href=mailto:Gregory.Rodgers@amd.com><font size=2 color=blue face="sans-serif"><u>Gregory.Rodgers@amd.com</u></font></a><font size=2 face="sans-serif">;
</font><a href=mailto:kli@ca.ibm.com><font size=2 color=blue face="sans-serif"><u>kli@ca.ibm.com</u></font></a><font size=2 face="sans-serif"><b><br>Subject:</b> Re: [RFC] Late (OpenMP) GPU code "SPMD-zation"</font><font size=3></font><br><font size=3> </font><br><font size=3>Any "worse" is not a good idea. We need to avoid
it. It would be good that the new code did not affect the performance,
especially for SPMD mode (I think, this "worse" will affect exactly
SPMD mode, no?)</font><p><tt><font size=3>-------------<br>Best regards,<br>Alexey Bataev</font></tt><br><font size=3>30.01.2019 14:38, Doerfert, Johannes пишет:</font><br><font size=3 face="Helvetica">The LLVM optimization (openmp-opt), which
does non-SPMD -> SPMD and custom state machine generation, will not
fire if</font><br><font size=3 face="Helvetica">the __kernel_general_... calls are "missing".
Thus if we inline "to early", we are "stuck" with the
non-SPMD choice (not worse than</font><br><font size=3 face="Helvetica">what we have now!) and the default library
state machine ("worse" than what we have right now). Does that
make sense?</font><br><br><font size=3 face="Helvetica">The second option described what I want
to see us do "later" in order to avoid the above scenario and
always get both, </font><br><font size=3 face="Helvetica">openmp-opt and inlining of the runtime
and work functions.</font><br><br><p><hr><br><font size=2 face="sans-serif"><b>From:</b> Alexey Bataev </font><a href=mailto:a.bataev@outlook.com><font size=2 color=blue face="sans-serif"><u><a.bataev@outlook.com></u></font></a><font size=2 face="sans-serif"><b><br>Sent:</b> Wednesday, January 30, 2019 1:25:42 PM<b><br>To:</b> Doerfert, Johannes<b><br>Cc:</b> Guray Ozen; Gheorghe-Teod Bercea; </font><a href="mailto:openmp-dev@lists.llvm.org"><font size=2 color=blue face="sans-serif"><u>openmp-dev@lists.llvm.org</u></font></a><font size=2 face="sans-serif">;
Finkel, Hal J.; </font><a href=mailto:Gregory.Rodgers@amd.com><font size=2 color=blue face="sans-serif"><u>Gregory.Rodgers@amd.com</u></font></a><font size=2 face="sans-serif">;
</font><a href=mailto:kli@ca.ibm.com><font size=2 color=blue face="sans-serif"><u>kli@ca.ibm.com</u></font></a><font size=2 face="sans-serif"><b><br>Subject:</b> Re: [RFC] Late (OpenMP) GPU code "SPMD-zation"</font><font size=3></font><br><font size=3> </font><br><font size=3>Sorry, did not understand your answer correctly. But you
wrote:</font><br><font size=3><u>for now, not doing the optimization is just fine.</u></font><br><font size=3>What do you mean?</font><p><tt><font size=3>-------------<br>Best regards,<br>Alexey Bataev</font></tt><br><font size=3>30.01.2019 14:23, Doerfert, Johannes пишет:</font><br><font size=3 face="Helvetica">Alexey, </font><br><br><font size=3 face="Helvetica">I'm not sure how to interpret "Bad
idea!". but I think there is again a misunderstanding.</font><br><font size=3 face="Helvetica">To help me understand, could you try
to elaborate a bit?</font><br><br><font size=3 face="Helvetica">To make my last email clear:</font><br><font size=3 face="Helvetica">I __do__ want inlining. Both answers
to your earlier inlining questions do actually assume the runtime library
calls __are eventually inlined__,</font><br><font size=3 face="Helvetica">that is why I mentioned LTO and the runtime
as bitcode.</font><br><font size=3 face="Helvetica">.</font><br><font size=3 face="Helvetica">Cheers,</font><br><font size=3 face="Helvetica"> Johannes</font><br><br><br><p><hr><br><font size=2 face="sans-serif"><b>From:</b> Alexey Bataev </font><a href=mailto:a.bataev@outlook.com><font size=2 color=blue face="sans-serif"><u><a.bataev@outlook.com></u></font></a><font size=2 face="sans-serif"><b><br>Sent:</b> Wednesday, January 30, 2019 1:14:56 PM<b><br>To:</b> Doerfert, Johannes<b><br>Cc:</b> Guray Ozen; Gheorghe-Teod Bercea; </font><a href="mailto:openmp-dev@lists.llvm.org"><font size=2 color=blue face="sans-serif"><u>openmp-dev@lists.llvm.org</u></font></a><font size=2 face="sans-serif">;
Finkel, Hal J.; </font><a href=mailto:Gregory.Rodgers@amd.com><font size=2 color=blue face="sans-serif"><u>Gregory.Rodgers@amd.com</u></font></a><font size=2 face="sans-serif">;
</font><a href=mailto:kli@ca.ibm.com><font size=2 color=blue face="sans-serif"><u>kli@ca.ibm.com</u></font></a><font size=2 face="sans-serif"><b><br>Subject:</b> Re: [RFC] Late (OpenMP) GPU code "SPMD-zation"</font><font size=3></font><br><font size=3> </font><br><font size=3>Bad idea!</font><p><tt><font size=3>-------------<br>Best regards,<br>Alexey Bataev</font></tt><br><font size=3>30.01.2019 14:11, Doerfert, Johannes пишет:</font><br><font size=3 face="Helvetica">Sure I do. Why do you think I don't?</font><br><p><hr><br><font size=2 face="sans-serif"><b>From:</b> Alexey Bataev </font><a href=mailto:a.bataev@outlook.com><font size=2 color=blue face="sans-serif"><u><a.bataev@outlook.com></u></font></a><font size=2 face="sans-serif"><b><br>Sent:</b> Wednesday, January 30, 2019 1:00:59 PM<b><br>To:</b> Doerfert, Johannes<b><br>Cc:</b> Guray Ozen; Gheorghe-Teod Bercea; </font><a href="mailto:openmp-dev@lists.llvm.org"><font size=2 color=blue face="sans-serif"><u>openmp-dev@lists.llvm.org</u></font></a><font size=2 face="sans-serif">;
Finkel, Hal J.; </font><a href=mailto:Gregory.Rodgers@amd.com><font size=2 color=blue face="sans-serif"><u>Gregory.Rodgers@amd.com</u></font></a><font size=2 face="sans-serif">;
</font><a href=mailto:kli@ca.ibm.com><font size=2 color=blue face="sans-serif"><u>kli@ca.ibm.com</u></font></a><font size=2 face="sans-serif"><b><br>Subject:</b> Re: [RFC] Late (OpenMP) GPU code "SPMD-zation"</font><font size=3></font><br><font size=3> </font><br><font size=3>You don't want to do the inlining?</font><p><tt><font size=3>-------------<br>Best regards,<br>Alexey Bataev</font></tt><br><font size=3>30.01.2019 13:59, Doerfert, Johannes пишет:</font><br><font size=3>- for now, not doing the optimization is just fine. The
whole idea is that code is always valid.</font><br><br><BR>