[llvm] [AMDGPU] Balance VM_CNT histories across branches (PR #221115)

[Dan Zimmerman via llvm-commits]

danzimm wrote:

Thanks for taking a look @Pierre-vh !

>  it should be in a separate file at the bare minimum, and broken up into multiple PRs

Definitely sounds reasonable, this is only my second AMDGPU contribution, so I'm learning the ropes! Let's sort out the specifics before I make anymore PRs.

> Only one specific kernel mentioned as the driver for the change

The example kernel was a minified repro of a performance issue seen in other kernels. This sort of pattern can appear in any ragged kernel. More precisely, it can appear in any kernel with either masked non-bufferload-able data (notably, this can also be data that technically can use buffer loads but static analysis can't determine it's safe to do), or buffer loads with non-zero `other`. A common example of the latter is [`softmax`](https://triton-lang.org/main/getting-started/tutorials/02-fused-softmax.html) since the `other` is nominally `-inf`.

> Change is costly, lots of new code, and adds multiple new dependencies to the pass that were previously optimized out.

I definitely understand the burden of maintenance here. I'm not trying to make your job more difficult. Just curious for my own learning (not trying to be nitty)- can you clarify the dependencies you're referring to?

> Super specific control knob to serve one user because we don't know if this is actually a useful thing in the grand scheme of things.

I'm happy to get rid of the knob. I originally included it because I'm used gating new optimizations for experimentation.

> Can't this be done in the kernel directly ? Or, can we help differently ? e.g. by adding a __pad_vmcnt built-in that the kernel can use and that'd inserts those "no-ops" (which also seem very hacky), but without burdening us with all the analysis ?

As (I think) you inferred, this isn't possible from the kernel side today, since there's no way to emit `buffer_inv 0` in a way that SIInsertWaitcnts recognizes (it doesn't parse through inline asm to update its scorecard). Introducing something like `__pad_vmcnt` could do the trick and avoid the analysis this PR introduces.

Do you think introducing `__builtin_amdgcn_pad_vmcnt() / llvm.amdgcn.pad.vmcnt()` or `__builtin_amdgcn_buffer_inv(arg) / llvm.amdgcn.buffer.inv(arg)` is better? I'm leaning towards the latter, but you'd know which is best.

For the sake of making sure I understand correctly (again not nitting, just for learning): we don't want to put the analysis/optimization in the AMDGPU backend to prevent introducing new complexity, right? If so, it appears that we're suggesting the tradeoff of pushing this logic to triton (or the kernel author in HIP/CK world) is worth it?

Just as devil's advocate (again for the sake of learning): I originally aimed at a backend pass because this trick is tied to specifically gfx942/gfx950, and is generic across any frontend. It sounds like this isn't enough to warrant the new analysis, is that right? I guess the risk here is a frontend/kernel author missing this opportunity because they don't know about the trick. I don't have the hardware in front of me, but gfx12+/CDNA5 appears it might have a similar opportunity, so the surface for the trick might expand (both in the sense of more hardware and more counters on CDNA5). As long as this is the right view of the world for you I'm happy to abide (to show my cards: I want explicit confirmation to refer to this thread in an inevitable conversation I'll have in a triton PR).

> Side note, I think this class of issue would likely be fixed by the currently-in-research-phase InsertWaitcnt rewrite

This is great to hear! Is the rewrite public somewhere that I can track?



https://github.com/llvm/llvm-project/pull/221115