[PATCH] D123878: [AMDGPU] Add remarks to output some resource usage

[Scott Linder via Phabricator via llvm-commits]

scott.linder added inline comments.


================
Comment at: llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp:1178-1266
+
+void AMDGPUAsmPrinter::emitResourceUsageRemarks(
+    const MachineFunction &MF, const SIProgramInfo &CurrentProgramInfo,
+    bool isModuleEntryFunction, bool hasMAIInsts) {
+  if (!ORE)
+    return;
+
----------------



================
Comment at: llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp:596
+               MF.getFunction().getSubprogram(), &MF.front())
+           << "------------------------------";
+  });
----------------
vangthao wrote:
> scott.linder wrote:
> > This seems like the wrong place to segment the output; could this be made a part of the emitter itself?
> > 
> > Or maybe better yet, could all of these values be collected into a single remark? That seems to be how e.g. `llvm/lib/CodeGen/RegAllocGreedy.cpp` does it:
> > 
> > ```
> > Pass:            regalloc                                                                                Name:            SpillReloadCopies                                                                       Function:        f                                                                                       Args:                                                                                                      - NumSpills:       '1'                                                                                   - String:          ' spills '                                                                            - TotalSpillsCost: '1.000000e+00'                                                                        - String:          ' total spills cost '                                                                 - NumReloads:      '1'                                                                                   - String:          ' reloads '                                                                           - TotalReloadsCost: '1.000000e+00'                                                                       - String:          ' total reloads cost '                                                                - String:          generated in function
> > ```
> We also want to output this in a readable format for the frontend. Collecting it all into a single remark seems to break the output format since clang seems to ignore all newlines from a diagnostic remark.
Rather than use newlines, RegAllocGreedy uses spaces; we can debate aesthetics, but I feel like we should have a compelling reason before we choose our own format.

For example, with RegAllocGreedy you can get output along the lines of:

```
foo/bar/baz.cpp:42:1: remark: 10 spills 100 total spills cost 7 folded spills 70 total folded spills cost 22 reloads 44 total reloads cost 77 folded reloads 120 total folded reloads cost 1 zero cost folded reloads 78 virtual registers copies 111 total copies cost
void foo()
^
```

RegAllocFast appears to be the original use-case that caused the machine remarks to be invented, and it has been updated recently (https://reviews.llvm.org/D100020) so I don't suspect this is just some legacy cruft.

If we follow the same approach we get something like:

```
  void AMDGPUAsmPrinter::emitResourceUsageRemarks(
      const MachineFunction &MF, const SIProgramInfo &CurrentProgramInfo) {
    if (!ORE)
      return;

    ORE->emit([&]() {
      return MachineOptimizationRemarkAnalysis(
                 "kernel-resource-usage", "ResourceUsage",
                 MF.getFunction().getSubprogram(), &MF.front())
             << ore::NV("NumSGPR", CurrentProgramInfo.NumSGPR) << " SGPRs "
             << ore::NV("NumVGPR", CurrentProgramInfo.NumArchVGPR) << " VGPRs "
             << ore::NV("NumAGPR", CurrentProgramInfo.NumAccVGPR) << " AGPRs "
             << ore::NV("ScratchSize", CurrentProgramInfo.ScratchSize)
             << " scratch bytes/thread "
             << ore::NV("Occupancy", CurrentProgramInfo.Occupancy)
             << " occupancy waves/SIMD "
             << ore::NV("SGPRSpill", CurrentProgramInfo.SGPRSpill)
             << " SGPR spills "
             << ore::NV("VGPRSpill", CurrentProgramInfo.VGPRSpill)
             << " VGPR spills " << ore::NV("BytesLDS", CurrentProgramInfo.LDSSize)
             << " LDS size bytes/block ";
    });                                                                                                                                                                                                                                                                   }
```

which produces:

```
clang/test/Frontend/amdgcn-machine-analysis-remarks.cl:14:1: remark: 9 SGPRs 10 VGPRs 12 AGPRs 0 scratch bytes/thread 10 occupancy waves/SIMD 0 SGPR spills 0 VGPR spills 0 LDS size bytes/block  [-Rpass-analysis=kernel-resource-usage]
__kernel void foo() {
^
```

That seems reasonable to me, and avoids bloating the other formats like YAML with extra remarks, including some that have no actual content (i.e. the "KernelName" and "KernelEnd" remarks are meaningless).

Of course I'm just one opinion, and if others prefer the several-remark approach I'm fine with it.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D123878/new/

https://reviews.llvm.org/D123878