Can wgmma.async and barrier.arrive Ensure GEMM Completion Before Moving Forward?

[ziyuhuang123]

Since the GEMM here uses wgmma.async, the subsequent barrier.arrive cannot guarantee that GEMM execution is complete, right? It can only ensure that the GEMM has been issued, correct? Otherwise, if strictly following the order of GEMM0, GEMM1, and SOFTMAX, it would be impossible to achieve the overlap within the warp group as shown in the figure below:

[ziyuhuang123]

Considering that GEMM is asynchronous (assuming) and softmax is synchronous (execution must complete before proceeding), I tried to combine the intra and inter illustrations to roughly draw something like this...

[tridao]

Yes barrier.arrive only guarantees that the gemm has been issued and not necessarily completed