good first issue
Repository metrics
- Stars
- (2,381 stars)
- PR merge metrics
- (PR metrics pending)
Description
For Grace Hopper GPU, since TMA requires a minimum size of 8X8, and WGMMA requires a minimum size of 64X8X16, we only enable these feature when the shape match to avoid padding, otherwise we use the cp.async and mma instructions for Ampere
Sometimes we can do padding if the required shared memory size does not exceed, and use TMA/WGMMA to utilize the performance. so we can add a memory check to try padding the shape if not fit in the instruction requirement. To use Hopper features as much as possible.