本次博客的github链接：
从安培开始，nv引入了异步拷贝的概念，在hopper架构中则体现为tma（tensor memory accelerator），其原理为，用极短的时间把要拷贝的tensor的shape和stride发送给launch函数后，kernel就可以去干别的事情，同时launch的拷贝也在继续，因此实现了异步的效果。比方说对于读->写->存模式的kernel:
__global__ void kernel(...)
{
load_data();
compute();
store_data();
}
可以通过一些方式在其中加入for循环，比如本来一个kernel只处理一个数据，现在处理十六个数据，那么就可以引入十六次for循环：
__global__ void kernel(...)
{
for(int i=0; i<16; i++){
load_data(i);
compute(i);
store_data(i);
}
}
此时如果我们不使用普通的load data方式（a=global_memory[i]），而使用tma来load_data，那么自然我们可以期望如下程序可以通过计算-访存重叠来加速:
__global__ void kernel(...)
{
launch_tma(0);
wait_tma(0);
for(int i=0; i<16; i++){
if(i<15){
launch_tma(i+1);
}
compute(i);
store_data(i);
if(i<15){
wait_tma(i+1);
}
}
}
这个代码对应github中的study3，事与愿违，速度只有不用tma的kernel的70%到90%。
于是开始分析为什么用了tma流水线反而变慢了。
和各路大佬们讨论后，结论如下：
①该算子本身的occupancy很高(occupancy概念可见juejin.cn/post/731233… )，但是用了tma之后由于需要使用共享内存，因此occupancy下降，而自己写的tma的访存-计算重叠效果，不如sm自己切warp来重叠计算和访存的效果好，因此最后速度更慢。
②tma适合什么场景？从上面那一点可以看出，如果occupancy无法升高，那么就有tma的用武之处，因为occupancy无法升高意味着无法通过切换warp来实现计算访存重叠，sm始终只有极少的warps，因此此时只能通过warps自己来实现软件流水线从而实现计算访存重叠，而gemm用到的mma指令会使用大量的寄存器，因此造成gemm算子的occupancy都打不高，于是只能走tma来实现流水。
最后附上gtc的某截图：
