说明

这里研究的cutlass版本是3.5

gemm讲解

  using CutlassGemm = cutlass::gemm::device::Gemm<float,        // Data-type of A matrixColumnMajor,  // Layout of A matrixfloat,        // Data-type of B matrixColumnMajor,  // Layout of B matrixfloat,        // Data-type of C matrixColumnMajor>; // Layout of C matrixCutlassGemm gemm_operator;CutlassGemm::Arguments args({M , N, K},  // Gemm Problem dimensions{A, lda},    // Tensor-ref for source matrix A{B, ldb},    // Tensor-ref for source matrix B{C, ldc},    // Tensor-ref for source matrix C{C, ldc},    // Tensor-ref for destination matrix D (may be different memory than source C matrix){alpha, beta}); // Scalars used in the Epiloguecutlass::Status status = gemm_operator(args);

上面是核心代码，可以看到首先要实例化一个类型CutlassGemm（编译期就要定下来），然后根据这个类型实例化一个对象gemm_operator（运行期），然后对象调用operator(args)做计算（运行期）。

编译期

  using CutlassGemm = cutlass::gemm::device::Gemm<float,        // Data-type of A matrixColumnMajor,  // Layout of A matrixfloat,        // Data-type of B matrixColumnMajor,  // Layout of B matrixfloat,        // Data-type of C matrixColumnMajor>; // Layout of C matrix

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

Gemm类

template <typename ElementA_,typename LayoutA_,typename OperatorClass_ = arch::OpClassSimt,typename ArchTag_ = arch::Sm70,typename ThreadblockShape_ = typename DefaultGemmConfiguration<OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_,ElementAccumulator_>::ThreadblockShape,//省略
> 
Gemm{}
//偏特化一个
template<省略>
Gemm<layoutC=layout::ColumnMajor,>

1. 这里偏特化很奇怪，单独给layoutC为列优先时候准备了一个类，具体什么原因也不深究，因为测试例子给的就是个ColumnMajor的layoutC，所以我们直接看这个偏特化类型。

这里增加了一个小知识，就是偏特化的模板不需要再传入默认值，会自动复用原始模板的默认值，此外由于偏特化实例化了一个值，导致在类里使用的时候没有了形参，为此可以看到源码里在类的开头搞了一堆的 类似using LayoutC = LayoutC_;即使偏特化实例化后，也能在类中再搞一个形参使用，CPP这搞得的是真恶心。