[QST] `wgmma_sm90.cu` tutorial pipelining

[jeromeku]

What is your question?
In the wgmma_sm90 tutorial, why is the pipelined inner loop executed k_tile_count > -K_PIPE_MAX times?

The producers load K_PIPE_MAX tiles here before entering the main loop, so wouldn't there be k_tile_count - K_PIPE_MAX tiles remaining?

[ccecka]

Those tiles have been preloaded, but not computed on yet. We still have k_tile_count calls to gemm to issue, which the mainloop does.

The useless loads

    k_tile_next = (k_tile_next >= K_TILE_MAX) ? K_TILE_MAX-1 : k_tile_next;

at the end don't matter within a properly pipelined kernel wherein the compute ops perfectly overlaps the memory ops to achieve SoL. Unlike the SM80 examples, the SM90 examples do not achieve SoL, but dummy loads are still more efficient and simple than attempting to predicate off the tail of the pipeline.

[jeromeku]

@ccecka

Makes sense! Thanks.