[MatMul] loop interleaving pass to interleave double buffered unrolled loops

[shmsong]

The loop interleaving optimization in this PR is needed on ampere with cp.async.

The transform this pass enables is the following:

original code:

for i0 in 0..4
  expr1

for i1 in 0..8
  expr2

for i2 in 0..4
  expr3

with some simple conservative checking that expr 1-3 have no direct dependencies, the pass transforms the above into:

for i0 in {0}
  expr1

for i1 in {0,1}
  expr2

for i2 in {0}
  expr3

for i0 in {1}
  expr1

for i1 in {2,3}
  expr2

for i2 in {1}
  expr3

for i0 in {2}
  expr1

for i1 in {4,5}
  expr2

for i2 in {2}
  expr3

...

The particular use case is the following:

for i0 in 0..4
  cp.async

for i1 in 0..8
  load.shared

// In here we are accumulating a lot of instructions
//  that either read or write shared memory and we will
//  see slow down due to congestions on the hardware

for i2 in 0..4
  mma

The interleaving essentially optimizes away the congestion mentioned above on the comment.

[naoyam]

Add comments on the pair

[naoyam]

Shouldn't a break be added here?

[naoyam]

Typo: Point 1

[naoyam]

Can't follow this case.

What I can't yet figure out is what the underlying generic condition this transformation must satisfy. Generally speaking, it seems safe if there's no data dependency between the subloop TVs, which basically corresponds to the Point 3. In the case of Point 2, it is also safe despite the data dependency because the dependency is constrained inside the sub loop, right? I can't wrap my head around the Point 1 yet, though.

[naoyam]

Does this mean the interleave main loop must also be a main loop of double buffering?

[naoyam]

Is this a generic bug fix or is it related to the interleaving transformation?

[naoyam]

Not completely following what should be done here, but the above comment on need_insert_commit indicates a commit should be inserted before the wait, but this seems to insert a commit after the wait inserted above. Am I missing something?