CUTLASS的GEMM（low-bit 内存搬运过程记录）（简单版）

记录CUTLASS的low-bit矩阵乘法GEMM计算中，主机内存到设备内存的搬运过程。（简单版）

（防止自己忘记。）

一、float型

首先对于float型的输入矩阵A和B而言，在主机中使用cutlass gemm模版创建两个tenosr来装矩阵A和B，然后通过sync_device函数将主机中的矩阵A和B放进设备中。在这个过程中，主机中的矩阵A和B的内存大小等于在设备中的矩阵A和B的内存大小。最后进行矩阵A和B的矩阵乘法计算。

二、low-bit型

当输入矩阵A和B是low-bit数据类型时，会有一些小变化。

例如，当输入矩阵A和B是1bit类型时，且假设矩阵A和矩阵B大小分别是8*128（row-major），128*8（col-major）。这里是指假设1bit为一个数据，那么矩阵A有8*128个数据，矩阵B有128*8个数据。

如果使用gemm模版来计算矩阵乘法时，首先会使用

cutlass::HostTensor<typename Gemm::ElementA, typename Gemm::LayoutA> tensor_A;
cutlass::HostTensor<typename Gemm::ElementB, typename Gemm::LayoutB> tensor_B;

在主机中生成矩阵A和矩阵B，此时矩阵A和B的数据个数应该各自是8*128，128*8。但是由于CUTLASS中装载1bit数据的数据类型是CUTLASS自己定义的cutlass::uint1b_t，这个数据类型实际上是由uint8改装而成的。因此此时在主机中，矩阵A和矩阵B各自的内存大小应该是8*128*8bit，128*8*8bit。也就是说，每一个uint8数据类型中，就只有1个有效数据。

注意，其中的ElementA和ElementB都是指的矩阵A和B的各自的数据类型，在这里应该是

cutlass::uint1b_t （这个数据类型的特点以后再记录）。

在主机中完成1bit矩阵A和B的定义和生成后，需要将他们使用sync_device函数传输到GPU设备中，进行实际的矩阵乘法计算。

在传输到GPU设备后，矩阵A和B的实际内存大小变成了8*128bit，128*8bit。

这说明，在进行sync_device这个命令时，将矩阵A和B从主机传输到设备前，会先将矩阵A和B压缩。具体压缩过程是，将主机中的每一个cutlass::uint1b_t（uint8）的一个有效bit位提取出来，然后每8个为一组，放进GPU设备的cutlass::uint1b_t（uint8）中。所以矩阵A在设备中，应该就是有128个cutlass::uint1b_t（uint8）数据类型的数据。矩阵B同理。

如果，1bit矩阵A和B已经在GPU设备中了，那么只需要将1bit矩阵A和B的有效数据压缩到cutlass::uint1b_t（uint8）中，然后将数据传递到gemm矩阵计算模版中，就可以计算1bit矩阵乘法了。

（如有错误，请不吝指正，万分感谢。）