[QST] Why do we need one MMA before the main loop in the unit test of pipeline_tma_async.cu?

[hyhieu]

My question is regarding the file test/unit/pipeline/pipeline_tma_async.cu.

Specifically, why do we need these lines?

cutlass/test/unit/pipeline/pipeline_tma_async.cu

Lines 133 to 141 in a8f2c80

	// MMA Prologue (Compute) - modeling inflight MMAs
	for (int iter = 0; iter < K_TILE_MMAS; ++iter)
	{
	pipeline.consumer_wait(smem_pipe_read);
	warpgroup_arrive();
	// GMMA would typically happen here
	++smem_pipe_read;
	}

In my understanding, NumStages copies have been issued in these lines

cutlass/test/unit/pipeline/pipeline_tma_async.cu

Lines 123 to 131 in a8f2c80

	// DMA Prologue (Loads)
	CUTLASS_PRAGMA_UNROLL
	for(int i = 0; i < k_pipe_tma_prologue; ++i) {
	pipeline.producer_acquire(smem_pipe_write);
	// cp.async.bulk.tensor would typically happen here
	pipeline.producer_commit(smem_pipe_write, per_cta_bytes);
	++smem_pipe_write;
	}
	tma_k_iterations -= k_pipe_tma_prologue;

So, why can't we just go into the main loop, instead of issuing K_TILE_MMAS = 1 MMA beforehands? Perhaps I misunderstood something. In that case, please help correct me 😄

Thank you in advance!