Meaning of the num_stages variable.

[DanielHesslow]

Hi 👋

Is there any reference someone can point me towards about what exactly the number of stages means?

Shared memory requirements appears to increase linearly with number of stages, so I guess it's some kind of buffering to increase memory bandwidth but if there's some more documentation about it it'd be helpful.

Thanks

[ptillet]

Hello! The Triton compiler pipelines some loops using the new asynchronous copy instruction on A100. num_stages refers to the depth of this pipeline. Because data from the next iteration(s) is asynchronously prefetched into shared memory, the shared memory requirements of the kernel increases linearly

[DanielHesslow]

Thanks!

Had completely missed that A100s have async copy, then everything makes much more sense.

[Shoreshen]

Hi does it means that one of the example can be the number of cp.async-group in ptx? which is a group of cp.async instruction. these instructions are not ordered within the cp.async-group and ordered between cp.async-group