[QST] Why the pipeline wait before update tma descriptors are removed, at tma array gemm kernel

[Dixeran]

What is your question?

Before cutlass 4.0.0, there is a pipeline wait before update the tma descriptor (

cutlass/include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_cooperative.hpp

Lines 709 to 713 in 2b78c2f

	// Purpose of this pipeline state is to make sure TMA loads have finished before doing descriptor updates
	// Since this state is waiting for loads to finish, it must start in the inverted phase.
	typename CollectiveMainloop::PipelineState mainloop_pipe_tma_consumer_state =
	{mainloop_pipe_producer_state.index(), !mainloop_pipe_producer_state.phase(), mainloop_pipe_producer_state.count()};
	mainloop_pipeline.consumer_wait(mainloop_pipe_tma_consumer_state);

).

As commented:

// Purpose of this pipeline state is to make sure TMA loads have finished before doing descriptor updates
// Since this state is waiting for loads to finish, it must start in the inverted phase.

why it can be safely removed?

[depaulmillz]

The kernels now use cp.async.bulk.commit_group and cp.async.bulk.wait_group to ensure the tensormap has been read from. This ensures that we can safely update the tensormap.

Relevant PTX docs: https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-wait-group

[Dixeran]

The kernels now use cp.async.bulk.commit_group and cp.async.bulk.wait_group to ensure the tensormap has been read from. This ensures that we can safely update the tensormap.

Relevant PTX docs: https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-wait-group

Thanks for your reply! So the finishing of TMA loads for A and B can observed by both complete-tx (which comsumer wg waits on) and bulk async-group based completion mechanism (which tensormap update watis on) ?