[QST] depthwise conv operand copy

[yupatrick22]

What is your question?
In the sample code 46, it seems that cp.async is used for copy data from device memory to shared memory

Each thread calculate the dst address in outer loop, and src address in inner loop and issue the async copy requests.
The inner loop is always fully unrolled, the outer loop is also fully unrolled for Kfixedstridedilation, while for Koptimized, the outer loop is not unrolled.

Is the above understanding correct?

For hopper architecture, making use of TMA (i.e. cp.async.bulk) will give better performance than cp.async, and TMA can completely replace cp.async for depthwise conv.

will TMA be used for depthsie conv in the future?
@Ethan-Yan27

[hwu36]

@Ethan-Yan27

[Ethan-Yan27]

Yes. For Koptimized, the number of load iterations is a runtime value, so the compiler generates a runtime branch.

Regarding the TMA feature, as far as I know, there are no plans yet. If you are interested, feel free to contribute code and compare performance.

[yupatrick22]

TMA can generate address and handle OOB, it is a perfect solution to decrease these two overheads. I highly expect it can be enabled in depthwise conv.