[wgmma] Insert commit_group and wait_group after mma_async

[jacobhinkle]

We need to wait for the MMA to complete before arriving at the circular buffer mbarrier, to avoid overwriting the inputs while the wgmma is being computed. This PR inserts a commit and wait just after the wgmma, which makes the wgmma synchronous.

Note that so far I have not removed the WAR sync insertion for wgmma which is no longer needed.

Fixes #3561.

[jacobhinkle]

!test

[rdspring1]

@jacobhinkle I ran into #3561 when working with register sharing and warp specialization. I wonder if I can use this.

[rdspring1]

There can be multiple wgmma operations per circular buffer stage. e.g., when the k dimension is a multiple of mma macro.

It should be safe to wait for all operations per stage.

I made the following change in WarAsyncWaitInserter.
https://github.com/NVIDIA/Fuser/pull/3616/files#diff-49bea61a8cde014ec0396c89d0654813136e0bd3ffe8b3a5974ee9ccf3a5fbb8R1010-R1024

It didn't any performance difference. 🤷🏼

[jacobhinkle]

!test

[jacobhinkle]

!test

[jacobhinkle]

!test

[jacobhinkle]

Looks like the MmaTest/HopperRS.SingleTile/Hopper*NoSwizzle* tests are failing due to invalid reads from smem.

[jacobhinkle]

Looks like the MmaTest/HopperRS.SingleTile/Hopper*NoSwizzle* tests are failing due to invalid reads from smem.

These are failing the compute-sanitizer check. This happens because we have inserted the WAR sync twice. This PR adds the sync right after the wgmma, essentially making it synchronous. That means there is no need to also add a WAR sync for the wgmma.