<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8">
</head>
<body style="word-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;" class="">
Thanks Andrea,
<div class=""><br class="">
</div>
<div class="">This is helpful and hopefully I can make some progress on the complicated approach to see if it works.</div>
<div class=""><br class="">
</div>
<div class="">-Cannada<br class="">
<div><br class="">
<blockquote type="cite" class="">
<div class="">On Jan 7, 2020, at 6:05 AM, Andrea Di Biagio <<a href="mailto:andrea.dibiagio@gmail.com" class="">andrea.dibiagio@gmail.com</a>> wrote:</div>
<br class="Apple-interchange-newline">
<div class=""><br class="Apple-interchange-newline">
<br style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none;" class="">
<div class="gmail_quote" style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none;">
<div dir="ltr" class="gmail_attr">On Mon, Jan 6, 2020 at 4:51 PM Lewis, Cannada <<a href="mailto:canlewi@sandia.gov" class="">canlewi@sandia.gov</a>> wrote:<br class="">
</div>
<blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;">
<div style="overflow-wrap: break-word;" class="">Andrea, thanks for the advice. <br class="">
<div class=""><br class="">
<blockquote type="cite" class="">
<div class="">On Jan 2, 2020, at 8:09 AM, Andrea Di Biagio <<a href="mailto:andrea.dibiagio@gmail.com" target="_blank" class="">andrea.dibiagio@gmail.com</a>> wrote:</div>
<br class="">
<div class="">
<div dir="ltr" class="">
<div class="">Hi Lewis,</div>
<div class=""><br class="">
</div>
Basically - if I understand correctly - you want to design a pass that uses llvm-mca as a library to compute throughput indicators for your outlined functions. You would then use those indicators to classify outlined functions.<br class="">
</div>
</div>
</blockquote>
Yes basically, the idea is to build a performance model for that outlined function.<br class="">
<blockquote type="cite" class="">
<div class="">
<div dir="ltr" class="">
<div class=""><br class="">
</div>
<div class="">llvm-mca doesn't know how to evaluate branches or instructions that affect the control flow. That basically restricts the analysis to single basic blocks that are assumed to be hot. I am not sure if this would be a blocker for your particular
use case.<br class="">
</div>
</div>
</div>
</blockquote>
That would be okay and is something we would need to work around on our end anyways since we don’t know the branch probability. <br class="">
<blockquote type="cite" class="">
<div class="">
<div dir="ltr" class="">
<div class=""><br class="">
<div class="">llvm-mca only knows how to analyze/simulate a sequence of `mca::Instruction`. So, the expectation is that instructions in input have already been lowered into a sequence of mca::Instruction. The only way currently to obtain an `mca::Instruction`
is by calling method `mca::<span class="">InstrBuilder</span>::<span class="">createInstruction</span>()` [1] on every instruction in input (see for example how it is done in llvm-mca.cpp [2]).<br class="">
</div>
<br class="">
</div>
<div class="">Unfortunately method `createInstructions()` only works on `MCInst&`. This strongly limits the usability of llvm-mca as a library; the expectation/assumption is that instructions have already been lowered to a sequence of MCInst.<br class="">
Basically the only supported scenarios are:<br class="">
<div class=""> - We have reached code emission stage and instructions have already been lowered into a sequence of MCInst, or</div>
</div>
</div>
</div>
</blockquote>
<blockquote type="cite" class="">
<div class="">
<div dir="ltr" class="">
<div class="">
<div class=""> - We obtained an MCInst sequence by parsing an assembly code sequence with the help of other llvm libraries (this is what the llvm-mca tool does).<br class="">
<br class="">
<div class="">
<div class="">It is possible to implement a variant of `createInstruction()` that lowers directly from `MachineInstr` to`mca::Instruction`. That would make the mca library more usable. In particular, it would make it possible to use mca from a post regalloc
pass which runs before code emission. Unfortunately, that functionality doesn't exist today (we can definitely implement it though; it may unblock other interesting use cases). That being said, I am not sure if it could help your particular use case. When
would you want to run your new pass? Using llvm-mca to analyze llvm IR is unfortunately not possible.<br class="">
</div>
</div>
</div>
</div>
</div>
</div>
</blockquote>
I would want to run my pass as late as possible so that all optimizations have been run on the outlined function. The values passed into the capture function should never influence the outlined function so I could in principle do something with a python script
like: </div>
<div class=""><br class="">
</div>
<div class="">1. Compile to ASM with MCA enable comments on the function I care about. </div>
<div class="">2. Run llvm-mca on that region</div>
<div class="">3. Source to source the original code with the new MCA information. </div>
<div class=""><br class="">
</div>
<div class="">Since this solution is not awesome and we may have many functions in a given TU that we care about I was hoping to find a better way. <br class="">
</div>
</div>
</blockquote>
<div class=""><br class="">
</div>
<div class="">I agree that the current workflow is not great.</div>
<div class=""><br class="">
</div>
<div class="">Ideally you would want to add your pass directly before code emission. That pass would 'run on MachineFunction'; it would only simulate outlined functions using mca (other functions will be just skipped). Unfortunately this solution would require
an API to lower from MachineInstr to mca::Instruction...<br class="">
</div>
<div class=""><br class="">
</div>
<div class="">There is however another not-so-simple and ugly approach. I didn't try it myself, but it should work for your particual use case (see below):<br class="">
</div>
<div class=""><br class="">
--<br class="">
<br class="">
Another approach is to take advantage of the logic already available in AsmPrinter.<br class="">
<br class="">
</div>
<div class="">Pass AsmPrinter is a machine function pass responsible for lowering and emitting machine functions.<br class="">
The instruction lowering logic is implemented by method `AsmPrinter:: EmitInstruction()`. That method is redefined in override by the llvm targets, and it is mainly responsible for a) lowering a MachineInstr into an MCInst, and b) delegating code emission to
an output streamer (an helper class which derives from MCStreamer).</div>
<div class=""><br class="">
</div>
<div class="">You could try to reuse the logic already available in AsmPrinter to implement point a).<br class="">
</div>
<div class="">You would then need to customize the streamer used in point b). Basically, you may want to mimic what is done in the llvm::mca::CodeRegionGenerator: you need a custom MCStreamer similar to the llvm::mca::MCStreamerWrapper [1]. Your MCStreamer
object would not really emit any code; instead it would just store each lowered MCInst into a sequence (for example, a SmallVector which can be accessed via the ArrayRef interface by your code).<br class="">
</div>
<div class=""><br class="">
</div>
<div class="">Once all instructions are lowered into MCInst, you can convert them into mca::Instruction using method `mca::InstrBuilder::createInstruction(MI)`. See for example how it is done in llvm-mca.cpp [2]<br class="">
</div>
<div class=""><br class="">
</div>
<div class="">The trick is to obtain a custom AsmPrinter pass that lowers machine instructions into MCInst, and then delegates "code emission" to your own custom MCStreamer class (whose goal is to simply store the lowered MCInsts somewhere in memory).<br class="">
</div>
<div class="">Your new pass would literally act as a proxy for the existing AsmPrinter.<br class="">
</div>
<div class="">Method runOnMachineFunction() would:<br class="">
a) delegate to AsmPrinter::runOnMachineFunction(), and<span class="Apple-converted-space"> </span><br class="">
b) analyze the MCInst sequence (obtained with the help of your custom streamer) with mca.<br class="">
</div>
<div class=""><br class="">
</div>
<div class="">Pass AsmPrinter is currently added to the pass manager by method `LLVMTargetMachine::addAsmPrinter()` [3][4]<br class="">
<br class="">
</div>
<div class="">You could provide another `addMCAPass()` hook to that interface to instantiate your new pass. That hook would implement logic that is similar to the one in method `addAsmPrinter()`. The difference is that instructions would be "emitted" by your
custom MCStreamer. It would create an AsmPrinter which is then used to construct your own pass P. P would own that particular AsmPrinter instance and; only P would be added to the pass manager directly before code emission.<br class="">
</div>
<div class=""><br class="">
</div>
<div class="">Not sure if this makes sense.<br class="">
</div>
<div class="">It may be a bit too complicated for your particular use case. However, I can't think of any simple/quick way to do this differently.<br class="">
</div>
<div class="">As I wrote: ideally in future we will be able to lower directly from MachineInstr and solve the problem differently. That would simplify use cases like yours.<br class="">
</div>
<div class=""><br class="">
</div>
<div class="">I hope it helps<br class="">
</div>
<div class="">-Andrea<br class="">
</div>
<div class=""><br class="">
[1]<span class="Apple-converted-space"> </span><a href="https://github.com/llvm-mirror/llvm/blob/master/tools/llvm-mca/CodeRegionGenerator.cpp#L42" class="">https://github.com/llvm-mirror/llvm/blob/master/tools/llvm-mca/CodeRegionGenerator.cpp#L42</a><br class="">
[2]<span class="Apple-converted-space"> </span><a href="https://github.com/llvm-mirror/llvm/blob/master/tools/llvm-mca/llvm-mca.cpp#L462" class="">https://github.com/llvm-mirror/llvm/blob/master/tools/llvm-mca/llvm-mca.cpp#L462</a><span class="Apple-converted-space"> </span><br class="">
[3]<span class="Apple-converted-space"> </span><a href="https://github.com/llvm-mirror/llvm/blob/master/lib/CodeGen/LLVMTargetMachine.cpp#L182" class="">https://github.com/llvm-mirror/llvm/blob/master/lib/CodeGen/LLVMTargetMachine.cpp#L182</a><br class="">
[4]<span class="Apple-converted-space"> </span><a href="https://github.com/llvm-mirror/llvm/blob/master/lib/CodeGen/LLVMTargetMachine.cpp#L116" class="">https://github.com/llvm-mirror/llvm/blob/master/lib/CodeGen/LLVMTargetMachine.cpp#L116</a><br class="">
<br class="">
</div>
<blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;">
<div style="overflow-wrap: break-word;" class="">
<div class="">
<blockquote type="cite" class="">
<div class="">
<div dir="ltr" class="">
<div class="">
<div class=""><br class="">
To compute the throughput indicators you would need to implement a logic similar to the one implemented by class SummaryView [3]. Ideally, most of that logic could be factored out into a helper class in order to help your particular use case and possibly avoid
code duplication.<br class="">
</div>
</div>
</div>
</div>
</blockquote>
Thanks, I’ll look into it. <br class="">
<blockquote type="cite" class="">
<div class="">
<div dir="ltr" class="">
<div class="">
<div class=""><br class="">
</div>
<div class="">I hope it helps,<br class="">
</div>
<div class="">-Andrea<br class="">
</div>
<div class=""><br class="">
</div>
</div>
<div class="">[1]<span class="Apple-converted-space"> </span><a href="https://github.com/llvm-mirror/llvm/blob/master/include/llvm/MCA/InstrBuilder.h" target="_blank" class="">https://github.com/llvm-mirror/llvm/blob/master/include/llvm/MCA/InstrBuilder.h</a><br class="">
[2]<span class="Apple-converted-space"> </span><a href="https://github.com/llvm-mirror/llvm/blob/master/tools/llvm-mca/llvm-mca.cpp" target="_blank" class="">https://github.com/llvm-mirror/llvm/blob/master/tools/llvm-mca/llvm-mca.cpp</a><br class="">
[3]<span class="Apple-converted-space"> </span><a href="https://github.com/llvm-mirror/llvm/blob/master/tools/llvm-mca/Views/SummaryView.h" target="_blank" class="">https://github.com/llvm-mirror/llvm/blob/master/tools/llvm-mca/Views/SummaryView.h</a></div>
</div>
<br class="">
<div class="gmail_quote">
<div dir="ltr" class="gmail_attr">On Tue, Dec 24, 2019 at 5:08 PM Lewis, Cannada via llvm-dev <<a href="mailto:llvm-dev@lists.llvm.org" target="_blank" class="">llvm-dev@lists.llvm.org</a>> wrote:<br class="">
</div>
<blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;">
Hi,<span class="Apple-converted-space"> </span><br class="">
<br class="">
I am trying to generate performance models for specific pieces of code like an omp.outlined function. Lets say I have the following code:<br class="">
<br class="">
start_collect_parallel_for_data(-1.0,-1.0,-1.0, size, “tag for this region”);<br class="">
#pragma omp parallel for<br class="">
for(auto i = 0; i < size; ++i){<br class="">
// … do work<br class="">
}<br class="">
stop_collecting_parallel_for_data();<br class="">
<br class="">
The omp region will get outlined into a new function and what I would like to be be able to do in opt is compile just that function to assembly, for some target that I have chosen, run llvm-mca just on that function, and then replace the -1.0s with uOps Per
Cycle, IPC, and Block RThroughput so that my logging code has some estimate of the performance of that region.<span class="Apple-converted-space"> </span><br class="">
<br class="">
Is there any reasonable way to do this from inside opt? I already have everything in place to find the start_collect_parallel_for_data calls and find the functions called between start and stop, but I could use some help with the rest of my idea.<br class="">
<br class="">
Thanks<br class="">
-Cannada Lewis<br class="">
_______________________________________________<br class="">
LLVM Developers mailing list<br class="">
<a href="mailto:llvm-dev@lists.llvm.org" target="_blank" class="">llvm-dev@lists.llvm.org</a><br class="">
<a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" rel="noreferrer" target="_blank" class="">https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a></blockquote>
</div>
</div>
</blockquote>
</div>
</div>
</blockquote>
</div>
</div>
</blockquote>
</div>
<br class="">
</div>
</body>
</html>