[llvm-dev] [RFC] Late (OpenMP) GPU code "SPMD-zation"
Doerfert, Johannes via llvm-dev
llvm-dev at lists.llvm.org
Wed Jan 30 16:56:42 PST 2019
Hi Doru,
[+ llvm-dev and cfe-dev]
On 01/30, Gheorghe-Teod Bercea wrote:
> Hi Johannes,
>
> First of all thanks for looking into the matter of improving non-SPMD mode!
>
> 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?
There will not be a difference. Let me explain in some details as there
seems to be a lot of confusion on this state machine topic:
Now:
Build a state machine in the user code (module) with all the parallel
regions as explicit targets of the switch statement and a fallback
default that does a indirect call to the requested parallel region.
Proposed, after Clang:
Use the runtime state machine implementation [0] which reduces the
switch to the default case, thus an indirect call to the requested
parallel region. This will always work, regardless of the translation
unit that contained the parallel region (pointer).
Proposed, after OpenMP-Opt pass in LLVM (assuming SPMD wasn't achieved):
All reachable parallel regions in a kernel are collected and used to
create the switch statement in the user code (module) [1, line 111] with
a fallback if there are potentially [1, line 212] hidden parallel
regions.
Does that make sense?
[0] https://reviews.llvm.org/D57460#change-e9Ljd9RgdWYz
[1] https://reviews.llvm.org/D57460#change-8gnnGNfJVR4B
> 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.
See the answer above.
> 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:
>
> #pragma omp target teams distribute
> for(...){
> <code1>
> #pragma omp parallel for
> for(...) {
> <code2>
> }
> <code3>
> }
>
> which forces the use of the master-worker scheme (non-SPMD mode) without any
> other distractions.
The atomic stuff I used to determine correctness. I haven't yet looked
at performance. I will do so now and inform you on my results.
> It would then be interesting to understand how you plan to change the LLVM code
> generated for this,
The examples show how the LLVM-IR is supposed to look like, right?
> 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.
I can also compare register usage in addition to performance but there
is no difference in synchronization. The number and (relative) order of
original runtime library calls stays the same. The number of user code
-> runtime library calls is even decreased.
Please let me know if this helps and what questions remain.
Thanks,
Johannes
> From: "Doerfert, Johannes" <jdoerfert at anl.gov>
> To: Alexey Bataev <a.bataev at outlook.com>
> Cc: Guray Ozen <gozen at nvidia.com>, Gheorghe-Teod Bercea
> <gheorghe-teod.bercea at ibm.com>, "openmp-dev at lists.llvm.org"
> <openmp-dev at lists.llvm.org>, "Finkel, Hal J." <hfinkel at anl.gov>,
> "Gregory.Rodgers at amd.com" <Gregory.Rodgers at amd.com>, "kli at ca.ibm.com"
> <kli at ca.ibm.com>
> Date: 01/30/2019 04:14 PM
> Subject: Re: [RFC] Late (OpenMP) GPU code "SPMD-zation"
> ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━
>
>
>
> I don't really see "many ifs and maybes", actually none.
>
> Anyway, I will now work on a patch set that adds the new functionality under a
> cmd flag
> in order to showcase correctness and performance on real code.
>
> If you, or somebody else, have interesting examples, please feel free to point
> me at them.
>
> Thanks,
> Johannes
>
>
> ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━
>
> From: Alexey Bataev <a.bataev at outlook.com>
> Sent: Wednesday, January 30, 2019 2:18:19 PM
> To: Doerfert, Johannes
> Cc: Guray Ozen; Gheorghe-Teod Bercea; openmp-dev at lists.llvm.org; Finkel, Hal
> J.; Gregory.Rodgers at amd.com; kli at ca.ibm.com
> Subject: Re: [RFC] Late (OpenMP) GPU code "SPMD-zation"
>
> 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.
>
> -------------
> Best regards,
> Alexey Bataev
> 30.01.2019 14:49, Doerfert, Johannes :
> No, SPMD mode will not be affected at all.
>
> The "worse" part is the following:
> If we inline runtime library calls before the openmp-opt pass had a chance to
> look at the code,
> we will not have a customized state machine for the __non-SPMD__ case. That
> is, the if-cascade
> checking the work function pointer is not there.
>
> Avoiding this potential performance decline is actually very easy. While we do
> not have the "inline_late" capability,
> run the openmp-opt pass __before__ the inliner and we will not get "worse"
> code. We might however miss out on
> _new_ non-SPMD -> SPMD transformations.
>
>
> Does that make sense?
>
> ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━
>
> From: Alexey Bataev <a.bataev at outlook.com>
> Sent: Wednesday, January 30, 2019 1:44:10 PM
> To: Doerfert, Johannes
> Cc: Guray Ozen; Gheorghe-Teod Bercea; openmp-dev at lists.llvm.org; Finkel, Hal
> J.; Gregory.Rodgers at amd.com; kli at ca.ibm.com
> Subject: Re: [RFC] Late (OpenMP) GPU code "SPMD-zation"
>
> 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?)
>
> -------------
> Best regards,
> Alexey Bataev
> 30.01.2019 14:38, Doerfert, Johannes :
> The LLVM optimization (openmp-opt), which does non-SPMD -> SPMD and custom
> state machine generation, will not fire if
> the __kernel_general_... calls are "missing". Thus if we inline "to early", we
> are "stuck" with the non-SPMD choice (not worse than
> what we have now!) and the default library state machine ("worse" than what we
> have right now). Does that make sense?
>
> The second option described what I want to see us do "later" in order to avoid
> the above scenario and always get both,
> openmp-opt and inlining of the runtime and work functions.
>
>
> ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━
>
> From: Alexey Bataev <a.bataev at outlook.com>
> Sent: Wednesday, January 30, 2019 1:25:42 PM
> To: Doerfert, Johannes
> Cc: Guray Ozen; Gheorghe-Teod Bercea; openmp-dev at lists.llvm.org; Finkel, Hal
> J.; Gregory.Rodgers at amd.com; kli at ca.ibm.com
> Subject: Re: [RFC] Late (OpenMP) GPU code "SPMD-zation"
>
> Sorry, did not understand your answer correctly. But you wrote:
> for now, not doing the optimization is just fine.
> What do you mean?
>
> -------------
> Best regards,
> Alexey Bataev
> 30.01.2019 14:23, Doerfert, Johannes :
> Alexey,
>
> I'm not sure how to interpret "Bad idea!". but I think there is again a
> misunderstanding.
> To help me understand, could you try to elaborate a bit?
>
> To make my last email clear:
> I __do__ want inlining. Both answers to your earlier inlining questions do
> actually assume the runtime library calls __are eventually inlined__,
> that is why I mentioned LTO and the runtime as bitcode.
> .
> Cheers,
> Johannes
>
>
>
> ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━
>
> From: Alexey Bataev <a.bataev at outlook.com>
> Sent: Wednesday, January 30, 2019 1:14:56 PM
> To: Doerfert, Johannes
> Cc: Guray Ozen; Gheorghe-Teod Bercea; openmp-dev at lists.llvm.org; Finkel, Hal
> J.; Gregory.Rodgers at amd.com; kli at ca.ibm.com
> Subject: Re: [RFC] Late (OpenMP) GPU code "SPMD-zation"
>
> Bad idea!
>
> -------------
> Best regards,
> Alexey Bataev
> 30.01.2019 14:11, Doerfert, Johannes :
> Sure I do. Why do you think I don't?
>
> ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━
>
> From: Alexey Bataev <a.bataev at outlook.com>
> Sent: Wednesday, January 30, 2019 1:00:59 PM
> To: Doerfert, Johannes
> Cc: Guray Ozen; Gheorghe-Teod Bercea; openmp-dev at lists.llvm.org; Finkel, Hal
> J.; Gregory.Rodgers at amd.com; kli at ca.ibm.com
> Subject: Re: [RFC] Late (OpenMP) GPU code "SPMD-zation"
>
> You don't want to do the inlining?
>
> -------------
> Best regards,
> Alexey Bataev
> 30.01.2019 13:59, Doerfert, Johannes :
> - for now, not doing the optimization is just fine. The whole idea is that code
> is always valid.
>
>
--
Johannes Doerfert
Researcher
Argonne National Laboratory
Lemont, IL 60439, USA
jdoerfert at anl.gov
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 228 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20190131/2f00ed57/attachment-0001.sig>
More information about the llvm-dev
mailing list