feng__shuai的博客

最朴素的思想是一个gemm起一个kernel，然后循环就可以n次就行，但是有些gemm尺寸比较小，这样的会sm无法都利用，为此用多stream来启动，这样就可以保证多个kernel可以一块执行，但是效果不是很好，可能是因为多次启动的原因。为此需要继续优化思路，这里利用。

这是一个经典cuda core 做gemm计算的例子，核心是掌握warp gemm级别的配置。

A矩阵举例，A是row major, 这里kContiguous =K，kStrided =M, 确实是的，连续数据的长度是k, 对于B矩阵，连续的值长度就是N，这里整明白了后可以看看其他DefaultMmaCore实例化的结果，验证发现是这个意思。这里A矩阵和B矩阵都是row-major，但是我们看到传入PitchLinearShape的尺寸和我们传统的行列顺序不一样，先传入的是列，然后才是行，为啥？

对于tensorcore的具体使用，需要配合文档才能使用，因为这里的mma是warp为概念的，对于每个线程只要做搬运数据搬运工作就行，程序员的作用就是根据文档，让每个线程搬运指定位置的数据到寄存器就行，得到的数据再放到指定位置就行，nv这块文档给的不详细，后来推荐用wmma的api，想搞细节可以看看tensorcore,不想搞的话用用wmma也行。和之前相比，区别在于做warpgemm的时候，上面的小绿色块是整个warp去计算得到的，之前是一个线程做。

这个case是介绍ampere架构的gemm，核心是看看multistage的逻辑。

可以看到，编译期时候，程序员必须要定下输入矩阵的layout和数据类型。事实上真的是这样吗?我们来深究一下这个cutlass::gemm::device::Gemm，从这个名字就可以看出来，cutlass实现了一个gemm，有device, threadblock, warp, thread几个级别gemm，这个sample里面用的是device级别, 所谓的device级别就是在cpu端的代码可以调用的，这个其实和cub中的逻辑是一样的。这里研究的cutlass版本是3.5。

矩阵的layout, 其实就是把coord包装一下，主要看看两者的区别, 其实layout中主要就是包含上面说的一个stride, 也确实，一旦固定是行优先还是列优先后，只要定了stride后就固定了格式。W], 至于为啥是这个顺序我也不知道，估计是约定俗成的。理解步长其实就是在内存中，N0距离N1的长度是C。不是很明白为啥RowMajor是class, 而其他类型是struct？真是让我强迫症犯了。

最近自己的服务器终于搞定了，虽然显卡是3060Ti, 但是基本目前nv主流的库和工具都支持，所以后续开始稳定研究cuda了，nv的库用过很多，目前玩过cub, thrust, cudnn，tensorrt等，但是只有少部分开源，比如cub和cutlass, cub是开源的但是都是一些经典并行算法，自己也移植过一个sort算法收获颇丰, cutlass更多是一些AI相关的算法，为了紧跟大模型时代，所以后面准备做个博客系列来分享学习经历，一方面记录，一方面看看能不能给国产GPU的高性能工程师一点优化参考，毕竟这